kernel3.cu

__global__ void kernel3_char(uint8_t *current, uint8_t *matrix, char *text, int textLen, int N, int matrixWidth, int currWidth, int fullArea) {
    int thid = threadIdx.x + blockIdx.x * blockDim.x;
    chunk_t cc;
    extern __shared__ uint8_t s_matrix[];

    if (!thid) {
        for (int i = 0; i < (sizeof(CHARS_STR) - 1) * fullArea; i++) {
            s_matrix[i] = matrix[i];
        }
    }

    __syncthreads();

    int start = thid * N;
    int max = start + N;

    for (int offset = 0, j = 0; j < textLen; j++, offset += matrixWidth) {
        int idx;
        for (int i = 0; i < (sizeof(CHARS_STR) - 1); i++) {
            if (CHARS_STR[i] == text[j]) {
                idx = i;
                break;
            }
        }

        uint8_t *morpheus = s_matrix + idx * fullArea;

        for (int i = start; i < max; i++) {

            int reali = i * sizeof cc;
            int x = offset + reali % matrixWidth;
            int y = reali / matrixWidth;

            int idx = (y * currWidth + x) / sizeof cc;
            cc = ((chunk_t *)current)[idx];

            #pragma unroll
            for (int j = 0; j < sizeof cc; j++) {
                ((uint8_t *)&cc)[j] = morpheus[i * sizeof cc + j];
            }

            ((chunk_t *)current)[idx] = cc;
        } 
    }

}