Question

Consider a C function negMat(), that negates each element of a K x K matrix y[][], and stores each result into the matrix x[][] :

void negMat(float *x, float *y, int K) {

int i, j;
for (i=0; i<K; i++) {

for (j=0; j<K; j++) {
x[i * K + j] = - y[i * K + j];

}

}

}

negMat() runs on the CPU (obviously), and x[][]and y[][] are stored in row-major order.

Write a CUDA kernel negMatGPU(), that negates each element of a K x K matrix yG[][], and stores each result into the matrix xG[][] (see the given prototype). xG[][]and yG[][] are in the GPU’s global memory, and also stored in row-major order. In your kernel, assume that there are K threads total, and each thread takes care of all the elements in a single row. Remember that the K threads may be in multiple blocks. Make sure your code is efficient; otherwise points will be deducted. (Check Problem 3 for useful CUDA information.)

__global__ void negMatGPU(float *xG, float *yG, int K) {

// your code follows here

Answer 1

__global__

void negMatGPU(float *xG, float *yG, int K) {

int column = blockDim.x * blockIdx.x + threadIdx.x;           //calculating column number

int row = blockDim.y * blockIdx.y + threadIdx.y;           //calculating row number

if (row < K && column < K)

{

int thread_id = row * K + column;            //finding exact location using row and column number in row major storage

xG[thread_id] = -yG[thread_id];

}

}

*/row and column depends on block size which can exceed K so we have to check if row and column lies within 0 to K or not and then thread_id is calculated which varies from 0 to K*K-1/*