GPU 计算 CMPS224 2021 学习笔记 02

• 并行类型

（1）任务并行

（2）数据并行

• CPU & GPU

CPU和GPU拥有相互独立的内存空间，需要在两者之间相互传输数据。

（1）分配GPU内存

（2）将CPU上的数据复制到GPU上

（3）在GPU上对数据进行计算操作

（4）将计算结果从GPU复制到CPU上

（5）释放GPU内存

• CUDA内存管理API

（1）分配内存

cudaError_t cudaMalloc(void **devPtr, size_t size)

 （2）释放内存

cudaError_t cudaFree(void *devPtr)

（3）内存复制

cudaError_t cudaMemcpy(void *dst, const void *src, size_t count, enum cudaMemcpyKind kind)

• grid & block & thread

• 向量加法

__global__ void vecadd_kernel(float* x, float* y, float* z, int N) {
    int i = blockDim.x*blockIdx.x + threadIdx.x;
    if(i < N) {
        z[i] = x[i] + y[i];
    }
}

void vecaddGPU(float* x, float* y, float* z, int N) {

    // Allocate GPU memory
    float *x_d, *y_d, *z_d;
    cudaMalloc((void**) &x_d, N*sizeof(float));
    cudaMalloc((void**) &y_d, N*sizeof(float));
    cudaMalloc((void**) &z_d, N*sizeof(float));

    // Copy data to GPU memory
    cudaMemcpy(x_d, x, N*sizeof(float), cudaMemcpyHostToDevice);
    cudaMemcpy(y_d, y, N*sizeof(float), cudaMemcpyHostToDevice);

    // Perform computation on GPU
    const unsigned int numThreadsPerBlock = 512;
    const unsigned int numBlocks = (N + numThreadsPerBlock – 1)/numThreadsPerBlock;
    vecadd_kernel <<< numBlocks, numThreadsPerBlock >>> (x_d, y_d, z_d, N);

    // Copy data from GPU memory
    cudaMemcpy(z, z_d, N*sizeof(float), cudaMemcpyDeviceToHost);

    // Deallocate GPU memory
    cudaFree(x_d);
    cudaFree(y_d);
    cudaFree(z_d);

}

• 编译

• 函数声明

 

__host__ __device__ float f(float a, float b) {
    return a + b;
}
void vecaddCPU(float* x, float* y, float* z, int N) {
    for(unsigned int i = 0; i < N; ++i) {
        z[i] = f(x[i], y[i]);
    }
}
__global__ void vecadd_kernel(float* x, float* y, float* z, int N) {
    int i = blockDim.x*blockIdx.x + threadIdx.x;
    if (i < N) {
        z[i] = f(x[i], y[i]);
    }
}

• 核函数的异步启动

GPU上核函数的调用是异步的，GPU上的核函数启动后，立即返回，CPU会继续执行下面的程序，不会等待核函数执行完成。

cudaError_t cudaDeviceSynchronize()

可以使用上述函数API来同步CPU和GPU之间的操作，CPU调用cudaDeviceSynchronize()后，会等待GPU上的所有核函数执行完成后才会执行下面的程序。

• 错误检查

所有 CUDA API 调用都会返回错误代码 cudaError_t，可用于检查是否发生任何错误。

cudaError_t err = ...;
if(err != cudaSuccess) {
    printf("Error: %s\n", cudaGetErrorString(err));
    exit(0);
}

对于内核调用，可以检查 cudaDeviceSynchronize() 返回的错误或调用以下 API 函数：

cudaError_t cudaGetLastError()