/*
    Marly da Cruz Cláudio 8936885
    Danilo Marques Araujo dos Santos 8598670

    Particionamento: por dados, onde cada elemento da matriz é processado em paralelo
    Comunicação: o cálculo da nova matriz depende da leitura da matriz e do seu limiar,
    e a exibição da matriz depende do cálculo estar pronto
    Aglomeração: cada bloco é responsável por checar o limiar de um elemento da matriz
    Mapeamento: cada bloco (1 dimensão = x) possui uma thread executando
*/

#include <stdlib.h>
#include <stdio.h>

__global__ void filter(int *matriz, int *matriz_limiar, int N, int limiar) {
    int tid = blockIdx.x;

    if(tid < N) {
        if(matriz[tid] <= limiar)
            matriz_limiar[tid] = 0;
        else
            matriz_limiar[tid] = 1;
    }
}

int main(int argc,char **argv) {

    int lin, col, i, j, limiar, *matriz, *matriz_limiar, elementos;
    int *dev_matriz, *dev_matriz_limiar;

    fscanf(stdin, "%d ", &lin);
    fscanf(stdin, "%d\n", &col);
    fscanf(stdin, "%d\n", &limiar);

    elementos = lin * col;

    matriz=(int *)malloc(elementos * sizeof(int));
    matriz_limiar=(int *)malloc(elementos * sizeof(int));

    cudaMalloc((void**)&dev_matriz, elementos * sizeof(int));
    cudaMalloc((void**)&dev_matriz_limiar, elementos * sizeof(int));

    for(i = 0; i < lin; i++) {
        for(j = 0; j < col; j++) {
            fscanf(stdin, "%d ", &(matriz[i*col+j]));
        }
    }

    cudaMemcpy(dev_matriz, matriz, elementos * sizeof(int),cudaMemcpyHostToDevice);

    filter<<<elementos,1>>>(dev_matriz, dev_matriz_limiar, elementos, limiar);

    cudaMemcpy(matriz_limiar, dev_matriz_limiar, elementos * sizeof(int),cudaMemcpyDeviceToHost);

    for(i = 0; i < lin; i++) {
        for(j = 0; j < col; j++) {
            printf("%d ", matriz_limiar[i*col+j]);
        }
        printf("\n");
    }

    cudaFree(matriz);
    cudaFree(matriz_limiar);

    free(matriz);
    free(matriz_limiar);

    return 0;
}