admin 管理员组文章数量: 893558
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关系
版权声明:本文标题:CUDA学习第三天:Kernel+grid+block关系 内容由网友自发贡献,该文观点仅代表作者本人, 转载请联系作者并注明出处:http://www.freenas.com.cn/jishu/1687329745h90241.html, 本站仅提供信息存储空间服务,不拥有所有权,不承担相关法律责任。如发现本站有涉嫌抄袭侵权/违法违规的内容,一经查实,本站将立刻删除。
发表评论