-
并行类型
(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()