号外~~ BT 牌的国网光碟已经完成了, 感谢 b 君和 c 君帮忙 ^^ 种子下载点 http://www.badongo.com/file/12156676 请大家帮忙传播~~~ 【修正】 (1) 上次忘了说 GRID 的大小虽然是 uint3 的结构，但只能使用 2D 而已 (其 z 成员只能是 1)，BLOCK 才能完整支援 3D 结构。 (2) API 部份 cudaThreadSynchronize() 是用来进行核心和主机程序的同步， 把它错打成 __syncthreads() 的功能，大概是那天太累了 @_@ 详见第六招。 ============================================================================ 第五招 记忆体 (主机、装置、共享、暂存器) ============================================================================ 基本的记忆体部份，最重要的是区分主机、装置、共享记忆体以及暂存器，硬体位置 和存取速度列於下表： ---------------------------------------------------------------- 中文名称 英文名称 硬体位置 存取速度 ---------------------------------------------------------------- 主机记忆体 Host Memory PC 上的 DRAM 透过 PCI-E, 很慢 装置记忆体 Device Memory 显示卡上的 DRAM 400~600 cycles 共享记忆体 Shared Memory 显示晶片 4 cycles issue 暂存器 Register 显示晶片 立即 ---------------------------------------------------------------- 这些记忆体的大小、功能和使用方式如下 ---------------------------------------------------------------- 名称 大小 功能 使用方式 ---------------------------------------------------------------- 主机记忆体 0.5~10GB 存放传统的主机资料 透过 API 装置记忆体 0.5~10GB 存放给 GPU 使用资料 kernel 直接存取 共享记忆体 16KB/BLOCK 执行绪之间交换资料 kernel 直接存取 暂存器 32~64KB/BLOCK 执行绪的区域变数 kernel 直接存取 ---------------------------------------------------------------- 相关细节如下： (1) 【主机记忆体】我们都很熟了，就是传统 C/C++ 使用的变数，可以透过 malloc()、 free()、new、delete 等来配置或释放。 (2) 【装置记忆体】的地位和主机记忆体很像，只是主机记忆体是对付 CPU，而 装置记忆体是对付 GPGPU，两者间的资料传送要透过 CUDA 的 API 达成， 可以透过 cudaMalloc()、cudaFree() 等函式来配置或释放 (详见第二招)， 在 CUDA 中，这类的记忆体称为【全域记忆体 (global memory)】。 (3) 【共享记忆体】是比较特殊的，在传统的序列化程式设计里没有直接的对应， 初学者可能要花多一点的时间在这上面，它用在『区块内执行绪间交换资料』， 只能在 kernel 中使用，宣告时使用 __shared__ 这个标签，使用时必需同步， 以确保资料读写时序上的正确性，现阶段在每个区块上最大的容量为 16KB， 超过便无法执行或编译，因为 on chip 的关系，它属於快速记忆体。 (4) 【暂存器】kernel 中大部份的区域变数都是以暂存器的型式存放，不需做额外的 宣告手序，这些暂存器是区块中执行绪共享的，也就是如果每个执行绪使用到 8 个 暂存器，呼叫这个 kernel 时区块大小指定为 10，则整个区块使用到 8*10=80 个 暂存器，若呼叫时指定区块大小为 50，则整个区块使用 8*50=400 个暂存器。 (5) 当使用过多的暂存器时 (>32~64KB/BLOCK，看是那个世代的 GPU)，系统会自动把 一些资料置换到全域记忆体中，导致执行绪变多，但效率反而变慢 (类似作业系统 虚拟记忆体的 swap)；另一个会引发这种 swap 的情况是在使用动态索引存取阵列， 因为此时需要阵列的顺序性，而暂存器本身是没有所谓的顺序的，所以系统会自动 把阵列置於全域记忆体中，再按索引存取，这种情况建议使用共享记忆体手动避免。 (6) 【暂存器】和【共享记忆体】的使用量会限制执行绪的数目，在开发复杂 kernel 时 宜注意，可使用 nvcc --ptxas-options=-v 这个选项在设计时期监控，或使用 nvcc --maxrregcount 选项限制每个执行绪的暂存器使用量。 //---------------------------------------------------------------------------- //范例(5): 平滑处理 (使用相邻的三点做加权平均,使资料变平滑)[081119-smooth.cu] // 执行绪同步 __syncthreads() 和 cudaThreadSynchronize()，详见第六招 //--------------------------------------------------------------------------- #include<stdio.h> #include<time.h> #include<cuda.h> //设定区块大小 (shared 版本会用到, 所以先宣告). #define BLOCK 512 //-------------------------------------------------------- //(1) 对照组 (host 版). //-------------------------------------------------------- void smooth_host(float* b, float* a, int n){ for(int k=1; k<n-1; k++){ b[k]=(a[k-1]+2*a[k]+a[k+1])*0.25; } //边界为0 b[0]=(2*a[0]+a[1])*0.25; b[n-1]=(a[n-2]+2*a[n-1])*0.25; } //-------------------------------------------------------- //(2) 装置核心(global 版). //-------------------------------------------------------- __global__ void smooth_global(float* b, float* a, int n){ int k = blockIdx.x*blockDim.x+threadIdx.x; if(k==0){ b[k]=(2*a[0]+a[1])*0.25; } else if(k==n-1){ b[k]=(a[n-2]+2*a[n-1])*0.25; } else if(k<n){ b[k]=(a[k-1]+2*a[k]+a[k+1])*0.25; } } //-------------------------------------------------------- //(3) 装置核心(shared 版). //-------------------------------------------------------- __global__ void smooth_shared(float* b, float* a, int n){ //---------------------------------------- //计算区块的基底 //---------------------------------------- int base = blockIdx.x*blockDim.x; int t = threadIdx.x; //---------------------------------------- //宣告共享记忆体. //---------------------------------------- __shared__ float s[BLOCK+2]; //---------------------------------------- //载入主要资料 s[1]~s[BLOCK] //---------------------------------------- // s[0] <-- a[base-1] (左边界) // s[1] <-- a[base] // s[2] <-- a[base+1] // s[3] <-- a[base+2] // ... // s[BLOCK] <-- a[base+BLOCK-1] // s[BLOCK+1] <-- a[base+BLOCK] (右边界) //---------------------------------------- if(base+t<n){ s[t+1]=a[base+t]; } //---------------------------------------- //载入边界资料 s[0] & s[BLOCK+1] (只用两个执行绪处理) //---------------------------------------- if(t==0){ //左边界. if(base==0){ s[0]=0; } else{ s[0]=a[base-1]; } } //*** 使用独立的 warp 让 branch 更快 *** if(t==32){ //右边界. if(base+BLOCK>=n){ s[n-base+1]=0; } else{ s[BLOCK+1] = a[base+BLOCK]; } } //---------------------------------------- //同步化 (确保共享记忆体已写入) //---------------------------------------- __syncthreads(); //---------------------------------------- //输出三点加权平均值 //---------------------------------------- if(base+t<n){ b[base+t]=(s[t]+2*s[t+1]+s[t+2])*0.25; } }; //-------------------------------------------------------- //主函式. //-------------------------------------------------------- int main(){ //-------------------------------------------------- //参数. //-------------------------------------------------- int num=10*1000*1000; int loop=130; //测试回圈次数 (量时间用) //-------------------------------------------------- //配置主机记忆体. //-------------------------------------------------- float* a=new float[num]; float* b=new float[num]; float* bg=new float[num]; float* bs=new float[num]; //-------------------------------------------------- //配置装置记忆体. //-------------------------------------------------- float *GA, *GB; cudaMalloc((void**) &GA, sizeof(float)*num); cudaMalloc((void**) &GB, sizeof(float)*num); //-------------------------------------------------- //初始化(乱数) & 复制资料到显示卡的 DRAM. //-------------------------------------------------- for(int k=0; k<num; k++){ a[k]=(float)rand()/RAND_MAX; b[k]=bg[k]=bs[k]=0; } cudaMemcpy(GA, a, sizeof(float)*num, cudaMemcpyHostToDevice); //-------------------------------------------------- //Test(1): smooth_host //-------------------------------------------------- double t_host=(double)clock()/CLOCKS_PER_SEC; for(int k=0; k<loop; k++){ smooth_host(b,a,num); } t_host=((double)clock()/CLOCKS_PER_SEC-t_host)/loop; //-------------------------------------------------- //Test(2): smooth_global (GRID*BLOCK 必需大於 num). //-------------------------------------------------- double t_global=(double)clock()/CLOCKS_PER_SEC; cudaThreadSynchronize(); for(int k=0; k<loop; k++){ smooth_global<<<num/BLOCK+1,BLOCK>>>(GB,GA,num); } cudaThreadSynchronize(); t_global=((double)clock()/CLOCKS_PER_SEC-t_global)/loop; //下载装置记忆体内容到主机上. cudaMemcpy(bg, GB, sizeof(float)*num, cudaMemcpyDeviceToHost); //-------------------------------------------------- //Test(3): smooth_shared (GRID*BLOCK 必需大於 num). //-------------------------------------------------- double t_shared=(double)clock()/CLOCKS_PER_SEC; cudaThreadSynchronize(); for(int k=0; k<loop; k++){ smooth_shared<<<num/BLOCK+1,BLOCK>>>(GB,GA,num); } cudaThreadSynchronize(); t_shared=((double)clock()/CLOCKS_PER_SEC-t_shared)/loop; //下载装置记忆体内容到主机上. cudaMemcpy(bs, GB, sizeof(float)*num, cudaMemcpyDeviceToHost); //-------------------------------------------------- //比较正确性 //-------------------------------------------------- double sum_dg2=0, sum_ds2=0, sum_b2=0; for(int k=0; k<num; k++){ double dg=bg[k]-b[k]; double ds=bs[k]-b[k]; sum_b2+=b[k]*b[k]; sum_dg2+=dg*dg; sum_ds2+=ds*ds; } //-------------------------------------------------- //报告 //-------------------------------------------------- //组态. printf("vector size: %d \n", num); printf("\n"); //时间. printf("Smooth_Host: %g ms\n", t_host*1000); printf("Smooth_Global: %g ms\n", t_global*1000); printf("Smooth_Shared: %g ms\n", t_shared*1000); printf("\n"); //相对误差. printf("Diff(Smooth_Global): %g \n", sqrt(sum_dg2/sum_b2)); printf("Diff(Smooth_Shared): %g \n", sqrt(sum_ds2/sum_b2)); printf("\n"); //-------------------------------------------------- //释放装置记忆体. //-------------------------------------------------- cudaFree(GA); cudaFree(GB); delete [] a; delete [] b; delete [] bg; delete [] bs; return 0; } //-------------------------------------------------------- //范例(5): 执行结果 (测试 10M 个 float) //-------------------------------------------------------- vector size: 10000000 Smooth_Host: 36.9231 ms Smooth_Global: 14.4615 ms Smooth_Shared: 5.07692 ms Diff(Smooth_Global): 3.83862e-08 Diff(Smooth_Shared): 3.83862e-08 (1) 这次测试的机器比较烂: P4-3.2 (prescott) vs. 9600GT 不过我们仍可看到共享记忆体使载入的资料量变为 1/3 所得到的增速 (2) 在 smooth_shared() 里我们用 2 个 warp 使得条件判断可以独立, 如果第 2 个 if(t==32) 改成 if(t==16) 或其它小於 32 的值, 也就是和第 1 个 if 使用同一个 warp, 则速度会变慢, 有兴趣的朋友 可以去试试看，warp 不打算在新手篇讲，之後硬体时才详细讨论。 (3) 测试效能时使用 cudaThreadSynchronize() 同步主机和装置核心， 以免量到错误的时间 (4) 在 smooth_shared() 里使用 __syncthreads() 同步化执行绪， 以免在计算 output 时仍有共享记忆体还没完成写入动作， 却有执行绪已经需要使用它的资料。 ============================================================================ 第六招 执行绪同步 (网格、区块) ============================================================================ 同步执行绪有两个函式，分别是 __syncthreads() 和 cudaThreadSynchronize() ----------------------------------------------------------------- 同步化函式 使用地点 功能 ----------------------------------------------------------------- __syncthreads() 核心程序中 同步化【区块内的执行绪】 cudaThreadSynchronize() 主机程序中 同步化【核心和主机程序】 ----------------------------------------------------------------- (1) 在 kernel 中，使用 __syncthreads() 来进行区块内的执行绪的同步， 避免资料时序上的问题 (来自不同 threads)，时常和共享记忆体一起使用， 在范例(5)中示范了使用 __syncthreads() 来隔开共享记忆体的【写入周期】 和【读取周期】，避免 WAR 错误 (write after read)。 (2) 在主机程序中，使用 cudaThreadSynchronize() 来进行核心和主机程序的同步， 范例(5)中示范了用它来避免量到不正确的主机时间 (kernel仍未完成就量时间)， 因为主机的程序和装置程序预设是不同步的 (直到下载结果资料之前)，这个 API 可以强迫它们同步。 -- 。o O　○。o O　○。o O　○。o O　○。o O　○。o 国网 CUDA 中文教学 DVD 影片 (免费线上版) 请至国网的教育训练网登入 https://edu.nchc.org.tw BT 牌的种子下载点 http://www.badongo.com/file/12156676 --

※ 发信站: 批踢踢实业坊(ptt.cc) ◆ From: 114.45.214.33 ※ 编辑: a5000ml 来自: 114.45.214.33 (11/20 02:28) ※ a5000ml:转录至看板 C_and_CPP 11/20 02:30 ※ 编辑: a5000ml 来自: 114.45.214.33 (11/20 02:55)