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 memoryfloat *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 memorycudaMemcpy(x_d, x, N*sizeof(float), cudaMemcpyHostToDevice);cudaMemcpy(y_d, y, N*sizeof(float), cudaMemcpyHostToDevice);// Perform computation on GPUconst 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 memorycudaMemcpy(z, z_d, N*sizeof(float), cudaMemcpyDeviceToHost);// Deallocate GPU memorycudaFree(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()