/perf/kseta

#define A(i, j) shmem[i * BLOCK_SIZE + j]

#define B(i, j) shmem[BLOCK_SIZE * BLOCK_SIZE + i * BLOCK_SIZE + j]

 

__kernel void multiply(__global float *res, __global float *a, __global float *b, unsigned long size, __local float *shmem) {

    float sum = 0;

    

    int tx = get_local_id(0);

    int ty = get_local_id(1);

 

    int i = get_global_id(1);

    int j = get_global_id(0);

 

 

    int k, l;

 

    for(k = 0; k < size; k += BLOCK_SIZE) {

        A(ty, tx) = a[i * size + (k + tx)];

        B(ty, tx) = b[(k + ty) * size + j];

        

        barrier(CLK_LOCAL_MEM_FENCE);

        

#pragma unroll BLOCK_SIZE

        for (l = 0; l < BLOCK_SIZE; ++l) {

            sum  += A(ty, l) * B(l, tx);

        }

 

        barrier(CLK_LOCAL_MEM_FENCE);

    }

    

    res[i * size +j] = sum;

}