(1) 有学弟反应 CUDA 内容有点繁杂, 很多概念容易搞混, 而且希望多点范例, 所以这两个礼拜把之前的文稿整理成【新手速成篇】,希望对他们有所帮助. (2) 顺便帮国网打个广告: CUDA 中文教学 DVD (免费线上版) 出现了 请至国网的教育训练网登入 https://edu.nchc.org.tw 详情请看编号 18026 一文 ※ 第十章 新手速成篇(上) ============================================================================ 前言 ============================================================================ 因为 CUDA 的一些延伸语法太繁杂，容易让人混淆 (例如记忆体种类就有4~5种， 同样的 global memory 又有两种写法)，所以针对这个问题，写成了速成篇， 去除那些枝枝节节，只讲最重要的，并佐以范例，务求让初学者【七招闯天下】。 第一招 主机、装置 第二招 使用 API (配置装置记忆体 & 主机和装置间资料搬移) 第三招 函式 & 呼叫 (主机、装置) 第四招 网格、区块、执行绪 (线程群组) 第五招 记忆体 (主机、装置、共享) 第六招 执行绪同步 (网格、区块) 第七招 合并读取 (最佳化) 函式部份，只介绍 __global__ 标签，记忆体部份，只介绍 __shared__ 标签， 配置显示记忆体以及资料搬移的方式，也只使用一种，简单来说，这份速成篇 并不是完整的 CUDA，只是删减後的正交子集合，用来突显主要概念，以及避免 初学者常犯的错误，熟悉之後，务必再深入了解其它延伸语法。 ============================================================================ 第一招 主机、装置 ============================================================================ (1) 区分主机和装置的不同： 【主机】就是PC。 【装置】就是显示卡。 (2) 两者皆有【中央处理器】，主机上为 CPU，装置上为 GPU，指令集不同： 主机上的程式码使用传统 C/C++ 语法撰写成，实作与呼叫和一般函式无异， 装置上的程式码称为【核心】(kernel)，需使用 CUDA 的延伸语法 (函式前加 __global__ 等标签) 来撰写，并於呼叫时指定执行绪群组大小 (详见第三招) (3) 两者皆有【各自的记忆体】(DRAM)，拥有独立的定址空间： 主机上的透过 malloc()、free()、new、delete 等函式配置与释放， 装置上的透过 cudaMalloc()、cudaFree() 等 API 配置与释放， 主机和装置之间的资料搬移，使用 cudaMemcpy() 这个 API (详见第二招) (4) 因为主机和装置的不同，C/C++ 的标准函式库不能在 kernel 中直接使用， 例如要秀出计算结果，必需使用 cudaMemcpy() 先将资料搬移至主机， 再呼叫 printf 或 cout 等标准输出函式。 (5) 使用时先在主机记忆体设好资料的初始值，然後传入装置记忆体，接着执行核心， 如果可以的话就尽量让资料保留在装置中，进行一连串的 kernel 操作， 避免透过 PCI-E 搬移造成效能下降，最後再将结果传回主机中显示。 ============================================================================ 第二招 使用 API (配置装置记忆体 & 主机和装置间资料搬移) ============================================================================ 最基本的 API 有 5 个 (1)配置装置记忆体 cudaMalloc() [cuda.h] (2)释放装置记忆体 cudaFree() [cuda.h] (3)记忆体复制 cudaMemcpy() [cuda.h] (4)错误字串解译 cudaGetErrorString() [cuda.h] (5)同步化 cudaThreadSynchronize() [cuda.h] 用法如下 -------------------------------------------------------- (1)配置显示记忆体 cudaMalloc() [cuda.h] -------------------------------------------------------- cudaError_t cudaMalloc(void** ptr, size_t count); ptr 指向目的指位器之位址 count 欲配置的大小(单位 bytes) 传回值 cudaError_t 是个 enum, 执行成功时传回 0, 其它的错误代号可用 cudaGetErrorString() 来解译. -------------------------------------------------------- (2)释放显示记忆体 cudaFree() [cuda.h] -------------------------------------------------------- cudaError_t cudaFree(void* ptr); ptr 指向欲释放的位址 (device memory) -------------------------------------------------------- (3)记忆体复制 cudaMemcpy() [cuda.h] -------------------------------------------------------- cudaError_t cudaMemcpy(void* dst, const void* src, size_t count, enum cudaMemcpyKind kind); dst 指向目的位址 src 指向来源位址 count 拷贝区块大小 (单位 bytes) kind 有四种拷贝流向 cudaMemcpyHostToHost 主机 -> 主机 cudaMemcpyHostToDevice 主机 -> 装置 cudaMemcpyDeviceToHost 装置 -> 主机 cudaMemcpyDeviceToDevice 装置 -> 装置 -------------------------------------------------------- (4)错误字串解译 cudaGetErrorString() [cuda.h] -------------------------------------------------------- const char* cudaGetErrorString(cudaError_t error); 传回错误代号(error)所代表的字串 -------------------------------------------------------- (5)同步化 cudaThreadSynchronize() [cuda.h] -------------------------------------------------------- cudaError_t cudaThreadSynchronize(void); 使前後两个核心时序上分离, 确保资料的前後相依性正确 //------------------------------------------------------------------------- //范例(1): 透过装置记忆体进行复制 [081112-api.cu] // PCI-E PCI-E // 主机记忆体 a[] --------> 装置记忆体 g[] --------> 主机记忆体 b[] //------------------------------------------------------------------------- #include<stdio.h> #include<cuda.h> int main(){ const int num=100; int* g; cudaError_t r; //主机阵列 & 初始化 int a[num], b[num]; for(int k=0; k<num; k++){ a[k]=k; b[k]=0; } //配置装置记忆体 & 显示错误讯息 r=cudaMalloc((void**) &g, sizeof(int)*num); printf("cudaMalloc : %s\n",cudaGetErrorString(r)); //复制记忆体: 主机记忆体 a[] ------> 装置记忆体 g[] r=cudaMemcpy(g, a, sizeof(int)*num, cudaMemcpyHostToDevice); printf("cudaMemcpy a => g : %s\n",cudaGetErrorString(r)); //复制记忆体: 装置记忆体 g[] ------> 主机记忆体 b[] r=cudaMemcpy(b, g, sizeof(int)*num, cudaMemcpyDeviceToHost); printf("cudaMemcpy g => b : %s\n",cudaGetErrorString(r)); //结果比对 bool ooo=true; for(int k=0; k<num; k++){ if(a[k]!=b[k]){ ooo=false; break; } } printf("check a==b? : %s\n",ooo?"pass":"wrong"); //释放装置记忆体 r=cudaFree(g); printf("cudaFree : %s\n",cudaGetErrorString(r)); return 0; } ------------------------------------------------------------- 范例(1)执行结果: ------------------------------------------------------------- cudaMalloc : no error cudaMemcpy a => g : no error cudaMemcpy g => b : no error check a==b? : pass cudaFree : no error ============================================================================ 第三招 函式 & 呼叫 (主机、装置) ============================================================================ CUDA 中，主机函式的写法与呼叫和传统 C/C++ 无异，而装置核心 (kernel) 要使用 延伸语法： __global__ void 函式名称 (函式引数...){ ...函式内容... }; 多了 __global__ 这标签来标明这道函式是核心程式码，要编译器特别照顾一下， 注意事项如下： (1) 传回值只能是 void (要传东西出来请透过引数) (2) 里面不能呼叫主机函式或 global 函式 (这两者皆是主机用的) (3) 输入的资料若是位址或参考时，必需指向装置记忆体。 呼叫 kernel 函式的语法比一般 C 函式多了指定网格和区块大小的手序： 函式名称 <<<网格大小, 区块大小>>> (函式引数...); 网格和区块详见第四招 //----------------------------------------------------------------------- //范例(2): hello CUDA 函式 (使用 global 函式填入字串) [081112-hello.cu] //----------------------------------------------------------------------- #include<stdio.h> #include<cuda.h> //装置函式(核心) 在显示卡记忆体中填入 hello CUDA 字串 __global__ void hello(char* s){ char w[50]="hello CUDA ~~~ =^.^="; int k; for(k=0; w[k]!=0; k++) s[k]=w[k]; s[k]=0; }; //主机函式 int main(){ char* d; char h[100]; //配置装置记忆体 cudaMalloc((void**) &d, 100); //呼叫装置核心 (只使用单一执行绪) hello<<<1,1>>>(d); //下载装置记忆体内容到主机上 cudaMemcpy(h, d, 100, cudaMemcpyDeviceToHost); //显示内容 printf("%s\n", h); //释放装置记忆体 cudaFree(d); return 0; } ------------------------------------------------------------- 范例(2)执行结果: ------------------------------------------------------------- hello CUDA ~~~ =^.^= ============================================================================ 第四招 网格、区块、执行绪 (线程群组) ============================================================================ 网格、区块、执行绪是 CUDA 中最重要的部份, 必需熟悉 (1) GPU 是具备超多核心，能行大量平行化运算的晶片，执行绪众多，要分群组管理： 最基本的执行单位是【执行绪】(thread)， 数个执行绪组成【区块】(block)， 数个区块组成【网格】(grid)， 整个网格就是所谓的【核心】(kernel)。 (2)【执行绪】是最基本的执行单位，程式设计师站在执行绪的角度，透过内建变数， 定出执行绪的位置，对工作进行主动切割。 (3)【区块】为执行绪的群组，一个区块可包含 1~512 个执行绪， 每个执行绪在区块中拥有唯一的索引编号，记录於内建变数 threadIdx。 每个区块中包含的执行绪数目，记录於内建变数 blockDim。 相同区块内的执行绪可同步化，而且可透过共享记忆体交换资料 (详见第五、六招) (4)【网格】为区块的群组，一个网格可包含 1~65535 个区块， 每个区块在网格中拥有唯一的索引编号，记录於内建变数 blockIdx。 每个网格中包含的区块数目，记录於内建变数 gridDim。 网格中的区块可能会同时或分散在不同时间执行，视硬体情况而定。 (5) 内建唯读变数 gridDim, blockDim, blockIdx, threadIdx 皆是 3D 正整数的结构体 uint3 gridDim ：网格大小 (网格中包含的区块数目) uint3 blockIdx ：区块索引 (网格中区块的索引) uint3 blockDim ：区块大小 (区块中包含的执行绪数目) uint3 threadIdx：执行绪索引 (区块中执行绪的索引) 其中 uint3 为 3D 的正整数型态，定义如下 struct uint3{ unsigned int x,y,z; }; 这些唯读变数只能在核心中使用。 (6) 核心呼叫时指定的网格和区块大小对应的就是其中 gridDim 和 blockDim 两变数 uint3 gridDim ：网格大小 (网格中包含的区块数目) uint3 blockDim ：区块大小 (区块中包含的执行绪数目) 可以在呼叫时只指定一维，此时变数里面的 y 和 z 成员都等於 1： 核心名称<<<int 网格大小, int 区块大小>>>(引数...); 也可以指定三维的呼叫: 核心名称<<<dim3 网格大小, dim3 区块大小>>>(引数...); 或者混合使用: 核心名称<<<dim3 网格大小, int 区块大小>>>(引数...); 核心名称<<<int 网格大小, dim3 区块大小>>>(引数...); 其中 dim3 等於 uint3，只是有写好 constructor 而己。 (7) 网格和区块大小在设定时有一定的限制 网格: max(gridDim) = 65535 区块: max(blockDim) = 512 实际在用的时候 blockDim 还会有资源上的限制, 主要是暂存器数目, 所以有时达不到 512 这个数量, 在 3 维的情况还会有其它的限制, 建议使用 1 维的方式呼叫, 到核心中再去切, 执行绪组态比较简单, 而且 bug 和限制也会比较少. //----------------------------------------------------------------- //范例(3): 列出在各执行绪中看到的区块和执行绪索引 [081112-idx.cu] // 【使用一维结构】 //----------------------------------------------------------------- #include<stdio.h> #include<cuda.h> //索引用到的绪构体 struct Index{ int block, thread; }; //核心:把索引写入装置记忆体 __global__ void prob_idx(Index id[]){ int b=blockIdx.x; //区块索引 int t=threadIdx.x; //执行绪索引 int n=blockDim.x; //区块中包含的执行绪数目 int x=b*n+t; //执行绪在阵列中对应的位置 //每个执行绪写入自己的区块和执行绪索引. id[x].block=b; id[x].thread=t; }; //主函式 int main(){ Index* d; Index h[100]; //配置装置记忆体 cudaMalloc((void**) &d, 100*sizeof(Index)); //呼叫装置核心 int g=3, b=4, m=g*b; prob_idx<<<g,b>>>(d); //下载装置记忆体内容到主机上 cudaMemcpy(h, d, 100*sizeof(Index), cudaMemcpyDeviceToHost); //显示内容 for(int i=0; i<m; i++){ printf("h[%d]={block:%d, thread:%d}\n", i,h[i].block,h[i].thread); } //释放装置记忆体 cudaFree(d); return 0; } ------------------------------------------------------------- 范例(3)执行结果: ------------------------------------------------------------- h[0]={block:0, thread:0} h[1]={block:0, thread:1} h[2]={block:0, thread:2} h[3]={block:0, thread:3} h[4]={block:1, thread:0} h[5]={block:1, thread:1} h[6]={block:1, thread:2} h[7]={block:1, thread:3} h[8]={block:2, thread:0} h[9]={block:2, thread:1} h[10]={block:2, thread:2} h[11]={block:2, thread:3} //------------------------------------------------------------------- //范例(4): 列出在各执行绪中看到的区块和执行绪索引 [081112-idx_3d.cu] // 【使用三维结构】 //------------------------------------------------------------------- #include<stdio.h> #include<cuda.h> //索引用到的绪构体 struct Index{ uint3 block, thread; }; //核心:把索引写入装置记忆体 __global__ void prob_idx_3d(Index* id){ //计算区块索引 int b=(blockIdx.z*gridDim.y+blockIdx.y)*gridDim.x+blockIdx.x; //计算执行绪索引 int t=(threadIdx.z*blockDim.y+threadIdx.y)*blockDim.x+threadIdx.x; //计算区块中包含的执行绪数目 int n=blockDim.x*blockDim.y*blockDim.z; //执行绪在阵列中对应的位置 int x=b*n+t; //每个执行绪写入自己的区块和执行绪索引. id[x].block=blockIdx; id[x].thread=threadIdx; } //主函式 int main(){ //网格和区块大小设定 dim3 grid=dim3(4,1,1); dim3 block=dim3(2,3,1); printf("gridDim = dim3(%d,%d,%d)\n", grid.x,grid.y,grid.z); printf("blockDim = dim3(%d,%d,%d)\n", block.x,block.y,block.z); //计算总执行绪数 int num=grid.x*grid.y*grid.z*block.x*block.y*block.z; printf("total num of threads = %d\n", num); //配置主机记忆体 & 清空 Index* h=(Index*)malloc(num*sizeof(Index)); memset(h,0,num*sizeof(Index)); //配置装置记忆体 & 清空 Index* d; cudaMalloc((void**) &d, num*sizeof(Index)); cudaMemcpy(d, h, num*sizeof(Index), cudaMemcpyHostToDevice); //呼叫装置核心. prob_idx_3d<<<grid,block>>>(d); //测试是否执行成功. cudaError_t r=cudaGetLastError(); printf("prob_idx_3d: %s\n", cudaGetErrorString(r)); if(r!=0) goto end; //下载装置记忆体内容到主机上. cudaMemcpy(h, d, num*sizeof(Index), cudaMemcpyDeviceToHost); //显示内容 for(int i=0; i<num; i++){ printf("h[%d]={block:(%d,%d,%d), thread:(%d,%d,%d)}\n", i, h[i].block.x, h[i].block.y, h[i].block.z, h[i].thread.x, h[i].thread.y, h[i].thread.z ); } end:; //释放装置记忆体. cudaFree(d); free(h); return 0; } ------------------------------------------------------------- 范例(4)执行结果: ------------------------------------------------------------- gridDim = dim3(4,1,1) blockDim = dim3(2,3,1) total num of threads = 24 prob_idx_3d: no error h[0]={block:(0,0,0), thread:(0,0,0)} h[1]={block:(0,0,0), thread:(1,0,0)} h[2]={block:(0,0,0), thread:(0,1,0)} h[3]={block:(0,0,0), thread:(1,1,0)} h[4]={block:(0,0,0), thread:(0,2,0)} h[5]={block:(0,0,0), thread:(1,2,0)} h[6]={block:(1,0,0), thread:(0,0,0)} h[7]={block:(1,0,0), thread:(1,0,0)} h[8]={block:(1,0,0), thread:(0,1,0)} h[9]={block:(1,0,0), thread:(1,1,0)} h[10]={block:(1,0,0), thread:(0,2,0)} h[11]={block:(1,0,0), thread:(1,2,0)} h[12]={block:(2,0,0), thread:(0,0,0)} h[13]={block:(2,0,0), thread:(1,0,0)} h[14]={block:(2,0,0), thread:(0,1,0)} h[15]={block:(2,0,0), thread:(1,1,0)} h[16]={block:(2,0,0), thread:(0,2,0)} h[17]={block:(2,0,0), thread:(1,2,0)} h[18]={block:(3,0,0), thread:(0,0,0)} h[19]={block:(3,0,0), thread:(1,0,0)} h[20]={block:(3,0,0), thread:(0,1,0)} h[21]={block:(3,0,0), thread:(1,1,0)} h[22]={block:(3,0,0), thread:(0,2,0)} h[23]={block:(3,0,0), thread:(1,2,0)} 我们可以由范例(3)和(4)看出执行绪索引的配置方式. =========================================================================== 待续... -- 。o O　○。o O　○。o O　○。o O　○。o O　○。o 国网 CUDA 中文教学 DVD 影片 (免费线上版) 请至国网的教育训练网登入 https://edu.nchc.org.tw --

※ 发信站: 批踢踢实业坊(ptt.cc) ◆ From: 114.45.214.93 ※ a5000ml:转录至看板 C_and_CPP 11/12 23:10 ※ uf2000uf:转录至某隐形看板 11/13 09:43