※ [本文转录自 VideoCard 看板] 作者: a5000ml (咖啡里的海洋蓝) 看板: VideoCard 标题: [分享] CUDA 程式设计(13) -- 材质快取 (一) 时间: Thu Dec 11 01:38:43 2008 终於要步入最佳化的部份, 实在令人兴奋... XDD =========================================================================== 材质参考 (part.1) =========================================================================== ◆简介 -------------------------------------------------------------------- 「材质参考」(Texture Reference) 其实就是开启 GPU 的「快取机制」(cache)。 传统 CPU 在记忆体的读取上，预设是有快取的，并提供一些指令来控制後续的指令 资料流的快取 (例如: x86 的预先拮取 prefetch)，然而这种方式，程式设计师必需 事先计算资料从 DRAM 预先被拮取所需的时间，并在适当的地方穿插 prefetch 指令， 可是硬体的进展使得预先拮取的时间变得不确定，所以 prefetch 的效果并不具备 向下相容性，使得在某型号的机器上好不容易调到最佳效能，换新机器後就不见得。 而「材质参考」是完全不一样的设计，源自於绘图所需的材质贴图管线 (texture pipeline)，因为在贴材质时需要快速的位址计算，以及快速拮取 2D 相邻的资料， 所以显示卡中早已内建多个位址计算单元和 2D 的硬体拮取来应付材质贴图的需求， 到了 CUDA 时代才把它化为快取机制而己，在硬体的指令设计上，也和传统 CPU 完全 不同，采用读取材质和直接读取使用不同指令集的方式，去除了 prefetch 这种 效果不佳的指令，改由程式执行之初的快取方式设定，支援多维度和归一化浮点数的 定址，且让多个 multiprocessors 可以共用，增加资料在快取中的运用率， 这些设计保留未来硬体调整的弹性，使得「材质参考」效果具备向下相容性。 ◆两种不同的最佳化风格 (合并存取 vs. 材质参考) -------------------------------------------------------------------- 「合并存取」可视为是 RISC 的延伸，因为资料要合并存取才能最佳化，会迫使 程式设计师先把一块连续资料载入处理器中，处理完後再把这些资料做一次输出， 形成了所谓 Load/Save 的程式风格，以及大量暂存器的需求，而这种风格与需求 正是 RISC 的本意。 「材质参考」本质为快取机制，如果资料刚好在快取中，那就 Lucky 载入很快， 如果没有，那就要靠快取处理单元，将资料从记忆体中读入，然後再喂给指令， 因为使用上太便利了，程式设计师不用管这些细节，会养成在需要资料时才载入， 那就是 CISC 的架构所推崇的方式。 事实上 GPU 本身是 RISC 架构，具备可控制的快取机制，只是在 CUDA 这个以 C/C++ 为基础的延伸语言中，提供了两种不同的最佳化方式让使用者自行调整，可完全使用 「材质参考」写程式 (会比较便利，并具备向下相容性)，也可选择「合并存取」来做 进阶最佳化，或是两者的混合达成平衡。 ◆效率因素 -------------------------------------------------------------------- 使用材质参考时最重要的是避免 __syncthreads()。 单纯的材质参考会比直接读取慢，因为材质位址单元还要经过一些运算，以及载入时 会一并载入相邻的资料，但合并读取是由执行绪分开发出的，在许多情况下 (执行绪 非独立串流) 还要透过共享记忆体来交换相邻的资料，所以必需 __syncthreads()， 这样会使得执行绪停摆率增加，且资料载入无法和运算重叠 (除非能塞入足够的运算 在 __syncthreads() 之前)，失去了使用多执行绪来隐藏资料载入延迟的好处。 而材质快取是跨执行绪的(甚至跨区块)，不需透过共享记忆体交换相邻的资料， 理论上可以完全不用 __syncthreads()，使得资料载入可以很好的和运算重叠， 多执行绪也因此可以有效的隐藏资料载入延迟，而且少用共享记忆体，也可以 增加多处理器同时执行的区块数。 实际使用时，两者要配合，因为 cache 空间有限，有些不需和相邻执行绪交换的资料， 那就用直接载入，原则是尽量不触发 __syncthreads()，如此可避免过多的材质参考 使得快取效能大幅下降。 ◆注意事项 -------------------------------------------------------------------- (1)材质记忆体是「唯读快取」，本身在硬体上并不是的独立的记忆晶片，而是 显示卡的 DRAM (global memory)，配合 GPU 上的 cache 来存取而己。 (2)在 CUDA 中预设的全域记忆体存取是没有快取的，只有透过材质参考来存取， 才能使用 GPU 的读取快取，两者的读取指令并不相同。 (3)材质参考可被设定成较适合的方式以增进存取效率，现阶段 nv 提供 1D、2D、3D 三种模式，1D 的可以直接绑在全域记忆体上，其余 2D 和 3D 必需绑 CUDA Array， 这东西其实和也和全域记忆体一样位在显示卡的 DRAM 上，只是有特殊的位址对齐 符合 2D 和 3D 的要求，使得材质快取可以有效率的进行而己。 (4)材质参考提供多种定址方式，除了整数位址之外，还提供归一化後的浮点数位址 (速度和整数位址差不多)，在某些应用上可节省位址计算的时间，并达到 fuzzy 定址的效果。 (5)材质参考可以跨多个 multiprocessors 共享，增加资料的利用率 (和硬体版本有关 G80/G90 系列的为两个多处理器共用，而 G200 为 3 个共用)。 ◆使用方式 -------------------------------------------------------------------- 使用材质快取只需三个步骤： (1)宣告材质的参考方式 (2)在呼叫 kernel 前的 host 程序中，做材质参考的绑定 (3)在 kernel 中使用材质参考。 先介绍 1D 对全域记忆体的材质参考，然後在范例中测试它的效能： (1)宣告材质参考的物件 【语法】 texture<资料型态, 维度, 定址模式> 材质物件名称; 资料型态：必需是基本的资料型态，或位址对齐的基本向量 (含 2 或 4 个元素) 例如 float, int, int4, float2, uint4 等 维度： 现阶段可以指定 1、2、3 维。 全域记忆体只能指定 1 维。 定址模式：可以是 cudaReadModeElementType 或 cudaReadModeNormalizedFloat 前者代表传统的整数位址，後者代表归一化过的浮点数位址， 全域记忆体只能用整数位址。 【范例】 texture<float, 1, cudaReadModeElementType> tex; 宣告材质快取对应的阵列元素之资料型态为 float，快取的方式为 1D， 使用整数位址存取，材质物件名称为 tex (2)在 host 程序中使用 API 对材质参考做绑定 【语法】 cudaBindTexture(偏移量, 材质物件, 全域记忆体位址, 记忆体大小); 偏移量：传回所需的偏移，和记忆体对齐有关。(这里可以先不管它) 材质物件：填入前面宣告的材质物件 全域记忆体位址：填入要绑定的记忆体位址 记忆体大小：填入要绑定的记忆体大小 (byte) 【范例】 float* aaa; cudaMalloc((void**)&aaa,N*sizeof(float)); ... cudaBindTexture(0, tex, aaa, N*sizeof(float)); 把全域记忆体 aaa 绑定 N 个 float 到 tex 这个材质物件。 (3)在 kernel 中使用材质参考。 【语法】 tex1Dfetch(材质名称, x); //全域记忆体只能用一维整数位址 tex2Dfetch(材质名称, x, y); tex3Dfetch(材质名称, x, y, z); 【范例】 比较 直接读取全域记忆体 b=aaa[k]; 透过材质参考读取 (tex 绑定到 aaa) b=tex1Dfetch(tex, x); ◆范例一：1D Laplace (d/dx)^2 的差分算符 -------------------------------------------------------------------- #include <stdlib.h> #include <stdio.h> #include <math.h> #include <cuda.h> //对照函数. void laplace(float* y, float* x, int n){ for(int i=1; i<n-1; i++){ y[i]=x[i+1]-2*x[i]+x[i-1]; } //periodic boundary y[0]=x[1]-2*x[0]+x[n-1]; y[n-1]=x[0]-2*x[n-1]+x[n-2]; } //使用合并读取+共享记忆体来做. __global__ void ker_laplace_shared(float* y, float* x, int n){ int t=threadIdx.x; int b=blockIdx.x*blockDim.x; int i=b+t; __shared__ float sm[512+2]; float* s=sm+1; if(i<n){ s[t]=x[i]; } //载入区块边界点 (使用两个 warp 来分散). if(t==0){ if(blockIdx.x==0){ s[-1]=x[n-1]; } else{ s[-1]=x[b-1]; } } if(t==32){ if(n-b<=blockDim.x){ s[n-blockIdx.x*blockDim.x]=x[0]; } else{ s[blockDim.x]=x[(blockIdx.x+1)*blockDim.x]; } } __syncthreads(); if(i<n){ y[i]=s[t+1]-2*s[t]+s[t-1]; } } //最原始的 kernel. __global__ void ker_laplace_naive(float* y, float* x, int n){ int i=blockIdx.x*blockDim.x+threadIdx.x; if(i==0){ y[0]=x[1]-2*x[0]+x[n-1]; } else if(i<n-1){ y[i]=x[i+1]-2*x[i]+x[i-1]; } else if(i==n-1){ y[n-1]=x[0]-2*x[n-1]+x[n-2]; } } //宣告材质物件. texture<float, 1, cudaReadModeElementType> texRefX; //使用材质快取的 kernel. __global__ void ker_laplace_texture(float* y, int n){ int i=blockIdx.x*blockDim.x+threadIdx.x; //用 macro 节省篇打字. #define xx(k) tex1Dfetch(texRefX, k) if(i==0){ y[0]=xx(1)-2*xx(0)+xx(n-1); } else if(i<n-1){ y[i]=xx(i+1)-2*xx(i)+xx(i-1); } else if(i==n-1){ y[n-1]=xx(0)-2*xx(n-1)+xx(n-2); } #undef xx } //乱数向量产生器. void vec_gen(float* vec, int size){ for(int i=0; i<size; i++){ vec[i]=(float)rand()/RAND_MAX*2-1; } } //比对两向量的相对误差. double diff(float* v1, float* v2, int size){ double sd=0; double sv=0; for(int i=0; i<size; i++){ double d=v1[i]-v2[i]; double v=v1[i]; sd+=d*d; sv+=v*v; } return sqrt(sd/sv); } //计萛时间差. double diff(timespec& t1, timespec& t2){ return (double)(t1.tv_sec-t2.tv_sec) + (double)(t1.tv_nsec-t2.tv_nsec)*1e-9; } //主函数. int main(){ int loops=503; //test loops int size=1024*1024; srand(time(0)); timespec ts1,ts2; float* v1=new float[size]; float* v2=new float[size]; float* v3=new float[size]; printf("------------------------------\n"); printf("1D Laplace Operator (periodic)\n"); printf(" vector size : %dK\n",size/1024); printf(" average loop : %d\n",loops); printf("------------------------------\n"); //---- generate profile in host----- vec_gen(v1,size); memset(v2,0,size*sizeof(float)); memset(v3,0,size*sizeof(float)); //---- set device memory ----- float *gv1,*gv3; cudaMalloc((void**)&gv1,(size+2)*sizeof(float)); cudaMalloc((void**)&gv3,(size+2)*sizeof(float)); cudaMemcpy(gv1,v1, size*sizeof(float), cudaMemcpyHostToDevice); //绑定材质参考 cudaBindTexture(0, texRefX, gv1, size*sizeof(float)); double dtdev,dthost; //---- test host performance ----- clock_gettime(0,&ts1); for(int k=0; k<loops; k++){ laplace(v2, v1, size); } clock_gettime(0,&ts2); dthost=diff(ts2,ts1)*1000/loops; printf("time(host): %g ms\n",dthost); //---- test naive device performance ----- memset(v3,0,size*sizeof(float)); cudaMemset(gv3,0,size*sizeof(float)); cudaThreadSynchronize(); clock_gettime(0,&ts1); for(int k=0; k<loops; k++){ ker_laplace_naive<<<size/512+1,512>>>(gv3,gv1,size); } cudaThreadSynchronize(); clock_gettime(0,&ts2); dtdev=diff(ts2,ts1)*1000/loops; cudaMemcpy(v3,gv3, size*sizeof(float), cudaMemcpyDeviceToHost); printf("time(naive): %g ms (%g x) error:%g\n",dtdev, dthost/dtdev,diff(v2,v3,size)); //---- test texture device performance ----- memset(v3,0,size*sizeof(float)); cudaMemset(gv3,0,size*sizeof(float)); cudaThreadSynchronize(); clock_gettime(0,&ts1); for(int k=0; k<loops; k++){ ker_laplace_texture<<<size/512+1,512>>>(gv3, size); } cudaThreadSynchronize(); clock_gettime(0,&ts2); dtdev=diff(ts2,ts1)*1000/loops; cudaMemcpy(v3,gv3, size*sizeof(float), cudaMemcpyDeviceToHost); printf("time(texture): %g ms (%g x) error:%g\n",dtdev, dthost/dtdev,diff(v2,v3,size)); //---- test shared device performance ----- memset(v3,0,size*sizeof(float)); cudaMemset(gv3,0,size*sizeof(float)); cudaThreadSynchronize(); clock_gettime(0,&ts1); for(int k=0; k<loops; k++){ ker_laplace_shared<<<size/512+1,512>>>(gv3,gv1,size); } cudaThreadSynchronize(); clock_gettime(0,&ts2); dtdev=diff(ts2,ts1)*1000/loops; cudaMemcpy(v3,gv3, size*sizeof(float), cudaMemcpyDeviceToHost); printf("time(shared): %g ms (%g x) error:%g\n",dtdev, dthost/dtdev,diff(v2,v3,size)); cudaFree(gv1); cudaFree(gv3); delete [] v1; delete [] v2; delete [] v3; return 0; } ◆结果一 -------------------------------------------------------------------- 这次测试使用 GTX260 vs. Intel E8400，其结果如下： ------------------------------ 1D Laplace Operator (periodic) vector size : 1024K average loop : 503 ------------------------------ time(host): 5.62482 ms time(naive): 0.246039 ms (22.8615 x) error:4.02879e-08 time(texture): 0.117252 ms (47.972 x) error:4.02879e-08 time(shared): 0.166139 ms (33.856 x) error:4.02879e-08 ------------------------------ ker_laplace_texture() 基本上只把 ker_laplace_naive() 改成材质快取而己， 里面没有 __syncthread()，大家可以看到它飙得超快，将近 host 的 50x， 而 shared 版本写得超辛苦，结果还是比开启快取的版本慢上许多。 ◆范例二：在 ker_laplace_texture() 插入 __syncthreads() -------------------------------------------------------------------- __global__ void ker_laplace_texture(float* y, int n){ int i=blockIdx.x*blockDim.x+threadIdx.x; #define xx(k) tex1Dfetch(texRefX, k) //先载入资料. float a,b,c; if(i==0){ a=xx(1); b=xx(0); c=xx(n-1); } else if(i<n-1){ a=xx(i+1); b=xx(i); c=xx(i-1); } else if(i==n-1){ a=xx(0); b=xx(n-1); c=xx(n-2); } //像 shared 版一样在资料读取和运算中间插入一个同步指令. __syncthreads(); //计算. y[i]=a-2*b+c; #undef xx } ◆结果二：在 ker_laplace_texture() 插入 __syncthreads() -------------------------------------------------------------------- 结果如下，我们可看出 texture 效果显着变差了，几乎和 shared 一样慢， 可见 __syncthread() 的确会对材质的效能造成很大的冲击，这里会稍微比 shared 快是因为 texture+sync 版少了载入边界的这个动作。 ------------------------------ 1D Laplace Operator (periodic) vector size : 1024K average loop : 503 ------------------------------ time(host): 5.62484 ms time(naive): 0.24602 ms (22.8634 x) error:4.02361e-08 time(texture): 0.156428 ms (35.958 x) error:4.02361e-08 (sync 版) time(shared): 0.166084 ms (33.8675 x) error:4.02361e-08 ------------------------------ 这个测试证实了我们前面的理论分析，所以要尽量避免 __syncthreads()， 才能让 texture 发挥最大功效。 -- 好像国网光碟大家都下载得差不多了, 现在用的人好像还不多, 没关系~~ 让我们继续聊, 好东西值得继续发展 ^^y -- 。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 http://rapidshare.de/files/41036559/NCHC_CUDA_video.torrent.html http://www.btghost.com/link/54915319/ --

※ 发信站: 批踢踢实业坊(ptt.cc) ◆ From: 114.45.209.244 ※ 编辑: a5000ml 来自: 114.45.209.244 (12/11 01:43) --

※ 发信站: 批踢踢实业坊(ptt.cc) ◆ From: 114.45.209.244