admin管理员组

文章数量:821300

CUDA学习第三天:Kernel+grid+block关系

1. 理一理前两天学到的概念之间的关系

  • CUDA && GPU

    CUDA: NIVID的CPUs上的一个通用并行计算平台和编程模型;
    GPU: CPU+GPU的异构计算架构,CPU所在位置为主机端(host), 而GPU所在位置为设备端(device).

  • SP && SM

    SP: 最基本的处理单元,streaming processor,也称为CUDA core。最后具体的指令和任务都是在SP上处理的。GPU进行并行计算,也就是很多个SP同时做处理。

    SM: 多个SP加上其他的一些资源组成一个streaming multiprocessor。也叫GPU大核,其他资源如:warp scheduler,register,shared memory等。SM可以看做GPU的心脏(对比CPU核心),register和shared memory是SM的稀缺资源。CUDA将这些资源分配给所有驻留在SM中的threads。因此,这些有限的资源就使每个SM中active warps有非常严格的限制,也就限制了并行能力。

    SP是线程执行的硬件单位,SM中包含多个SP,一个GPU可以有多个SM(比如16个),最终一个GPU可能包含有上千个SP

  • Host && Device

    Host与Device的内存交互模型:

    明确Host与Device分别位于什么位置:

2. Kernel与网格(grid)与线程块(block)与线程(Thread)

在实际编写CUDA程序时,我们需要设计block等的大小,也就是合理安排线程的数量。

3. CUDA实例程序二

  • 程序文件结构

    sample_01_plus.cu (cpp文件无法运行)

  • CMakeLists添加执行文件

    cuda_add_executable(sample_01 src/sample_01_plus.cu)

  • 程序源代码

		#include <iostream>#include <cuda.h>#include <cuda_runtime.h>#include <cuda_runtime_api.h>#include <device_launch_parameters.h>__global__ void addCuda(float* x, float* y, float* z, int n){int index = threadIdx.x+blockIdx.x*blockDim.x;int stride =  blockDim.x*gridDim.x;for(int i=index; i<n; i+=stride){z[i] = x[i] + y[i];}}int main(){int N =20;int nBytes = N*sizeof(float);float *x, *y, *z;/***********************************************  在device上分配内存函数:* *  cudaError_t cudaMalloc(void** devPtr, size_t size);* * 	devPtr: 指向所分配内存的指针,* * *****************************************************************/x = (float*)malloc(nBytes);y = (float*)malloc(nBytes);z = (float*)malloc(nBytes);for(int i=0; i<N; i++){x[i] = 10.0;y[i] = 20.0;}float *dx, *dy, *dz;/***********************************************  负责host与device之间数据通讯的cudaMemcpy函数:* *  cudaError_t cudaMemcpy(void* dst, const void* src, size_t count, cudaMemcpyKind kind)* * 	其中src指向数据源,而dst是目标区域,count是复制的字节数,其中kind控制复制的方向:* * 	cudaMemcpyHostToHost, cudaMemcpyHostToDevice, cudaMemcpyDeviceToHost, cudaMemcpyDeviceToDevice.* 	如cudaMemcpyHostToDevice将host上数据拷贝到device上* * *****************************************************************/cudaMalloc((void**)&dx, nBytes);cudaMalloc((void**)&dy, nBytes);cudaMalloc((void**)&dz, nBytes);// x-y为原始数据, 将其复制到gpu上计算;cudaMemcpy((void*)dx, (void*)x, nBytes, cudaMemcpyHostToDevice);cudaMemcpy((void*)dy, (void*)y, nBytes, cudaMemcpyHostToDevice);dim3 blockSize(256);dim3 gridsize((N + blockSize.x - 1) / blockSize.x);// 在gpu上计算的结果保存在dz上;addCuda<<<gridsize, blockSize>>>(dx, dy, dz, N);// dz为得到的结果,将其复制出来放到z上, 然后检测z与实际值的差距。cudaMemcpy((void*)z, (void*)dz, nBytes, cudaMemcpyDeviceToHost);float maxError = 0.0;for (int i = 0; i < N; i++)maxError = fmax(maxError, fabs(z[i] - 30.0));std::cout << "最大误差: " << maxError << std::endl;cudaFree(dx);cudaFree(dy);cudaFree(dz);free(x);free(y);free(z);return 0;}

本文标签: CUDA学习第三天Kernelgridblock关系