CUDA入门：CUDA简介习题

1. 矩阵加法接受两个矩阵 B 和 C 产生输出矩阵 A。输出矩阵 A 的每个元素是输入矩阵 B 和 C 相应元素之和，即 A[i][j] = B[i][j] + C[i][j]。方便起见只考虑方阵，元素都是单精度浮点值。写一个矩阵加法 kernel 函数和一个 stub 函数，可以用如下 4 个参数调用：指向输出矩阵的指针、指向第一个输入矩阵的指针、指向第二个输入矩阵的指针和每个维度上的元素个数。根据如下说明编写：

主机存根函数为输入和输出矩阵分配内存、传输输入数据到设备上、启动 kernel 函数、将输出数据传输至主机以及回收为输入和输出矩阵分配的设备存储器。

int matrixAdd(float *A, float *B, float *C, int n){
    int size = n *n * sizeof(float);
    float *d_A, *d_B, *d_C;
    cudaMalloc((void **)&d_A, size);
    cudaMemcpy(d_A, A, size, cudaMemcpyHostToDevice);
    cudaMalloc((void **)&d_B, size);
    cudaMemcpy(d_B, B, size, cudaMemcpyHostToDevice);
    cudaMalloc((void **)&d_C, size);

    // matrixAddKernel1<<<ceil(n * n / 512.0), 512>>>(d_A, d_B, d_C, n);
    // matrixAddKernelRow<<<ceil(n / 512.0), 512>>>(d_A, d_B, d_C, n);
    // matrixAddKernelCol<<<ceil(n / 512.0), 512>>>(d_A, d_B, d_C, n);
    
    cudaMemcpy(C, d_C, size, cudaMencpyDeviceToHost);
	cudaFree(d_A);
	cudaFree(d_B);
	cudaFree(d_C);
}

编写 kernel 函数，每个线程生成一个输出矩阵元素。并为此种方案填写执行配置参数。

//compute vector sum C = A + B
//Each thread performs one one pair-wise addition
__global__ void matrixAddKernel(float *A, float *B, float *C, int n){
	int i = threadId.x + blockDim.x * blockIdx.x;
	if(i < n) C[i] = A[i] + B[i];
}

编写 kernel 函数，每个线程生成一行输出矩阵元素。并为此种方案填写执行配置参数。

//compute matrix sum C = A + B
//Each thread performs 
__global__ void matrixAddKernelRow(float *A, float *B, float *C, int n){
    int i = threadIdx.x + blockDim.x * blockIdx.x;
    if(i < n){
        for(int j = 0;j < n; ++j){
            C[i + j * n] = A[i + j * n] + B[i + j * n];
        }
    }   
}

- 编写 kernel 函数，每个线程生成一列输出矩阵元素。并为此种方案填写执行配置参数。

__global__ void matrixAddKernelCol(float *A, float *B, float *C, int n){
    int i = threadIdx.x + blockDim.x * blockIdx.x;
    if(i < n){
        for(int j = 0;j < n; ++j){
            C[j * n + i] = A[j * n + i] + B[j * n + i];
        }
    }
}

分析以上几种 kernel 方案的优劣
我的理解是第一个比较快，一个线程算一个元素和，对 GPU 的利用率较高，另外两个一个线程算一行或一列，利用率较低。

main函数：

int main(int argc, char ** argv){
    int N = atoi(argv[1]);
    float *A, *B, *C;
    A = (float *)malloc(sizeof(float) * N * N);
    B = (float *)malloc(sizeof(float) * N * N);
    C = (float *)malloc(sizeof(float) * N * N);
    for(int i = 0;i < N;++i){
        for(int j = 0;j < N;++j){
            A[i * N + j] = i * N + j;
            B[i * N + j] = 1;
        }
    }
    matrixAdd(A, B, C, N);
    for(int i = 0;i < N;++i){
        for(int j = 0;j < N; ++j){
            printf("%f ",C[i * N + j]);
        }
        printf("\n");
    }
    free(A),free(B),free(C);
    return 0;
}

2. 写一个计算矩阵与向量相乘的 kernel 函数。主机存根函数可以通过指向输出矩阵上的指针、指向输入矩阵上的指针、指向输入向量的指针、以及每个维度上元素的数目。

//kernel function
//A is the input matrix, B in the input vector, C is the output vector, n is the dimension and length of matrix
//for clearity, we only consider the condition of phalanx
__global__ void MatVecMulKernel(float *A, float *B, float *C, int n){
    int i = int i = threadIdx.x + blockDim.x * blockIdx.x;
    if(i < n){
        for(int j = 0;j < n;++j){
            C[i] += A[i * n + j] * B[j];
        }
    }
}

//stub function
void MatVecMul(float *A, float *B, float *C, int n){
    int sizeM = n * n * sizeof(float),
        sizeV = n * sizeof(float);
    float *d_A, *d_B, *d_C;
    cudaMalloc((void **)&d_A, sizeM);
    cudaMemcpy(d_A, A, sizeM, cudaMemcpyHostToDevice);
    cudaMalloc((void **)&d_B, sizeV);
    cudaMemcpy(d_B, B, sizeV, cudaMemcpyHostToDevice);
    cudaMalloc((void **)&d_C, sizeV);
    
    MatVecMulKernel<<<ceil(n / 512.0), 512>>>(d_A, d_B, d_C, n);
    
    cudaMemcpy(C, d_C, sizeV, cudaMemcpyDeviceToHost);

	cudaFree(d_A);
	cudaFree(d_B);
	cudaFree(d_C);
}

3.我们希望使用每个线程来计算向量加法的两个(相邻的)输出元素。假设变量i应该是线程要处理的第一个元素的索引。将线程/块索引映射到第一个元素的数据索引的表达式是什么?

a) i = blockIdx.x * blockDim.x + threadIdx.x + 2
b) i = blockIdx.x * threadIdx.x * 2
c) i = (blockIdx.x * blockDim.x + threadIdx.x) * 2
d) i = blockIdx.x * blockDim.x * 2 + threadIdx.x

答案：一个线程两个元素，线程处理的第一个元素在向量中的索引为线程索引的二倍，选C

4.对于向量加法，假定向量长度2000，每个线程计算一个输出元素，线程块的大小为512,。网格中将有多少线程？

a) 2000
b) 2024
c) 2048
d) 2096
答案：ceil(2000/512) = 4, 4 * 512 = 2048。选c

刚开始学，有问题欢迎一起交流。