(CUDA 编程9).CUDA shared memory使用------GPU的革命
2012-07-29 09:20
387 查看
(CUDA编程9).CUDAsharedmemory使用------GPU的革命
作者:赵开勇来源:http://www.hpctech.com/2009/0818/207.html
这一章节准备写一下sharedmemory的访问的问题,首先是讲一下shared的memory的两种使用方法,然后讲解一下sharedmemory的bankconflict的问题,这个是sharedmemory访问能否高效的问题所在
9.CUDAsharedmemory使用------GPU的革命书接上文《8.CUDA
内存使用global二------GPU的革命》讲了global内存访问的时候,需要对齐的问题,只有在对齐的情况下才能保证global内存的高效访问。这一章节准备写一下sharedmemory的访问的问题,首先是讲一下shared的memory的两种使用方法,然后讲解一下shared
memory的bankconflict的问题,这个是sharedmemory访问能否高效的问题所在;
Sharedmemory的常规使用:
1.使用固定大小的数组:
/************************************************************************/
/*Example*/
/************************************************************************/
__global__voidshared_memory_1(float*result,intnum,float*table_1)
{
__shared__floatsh_data[THREAD_SIZE];
intidx=threadIdx.x;
floatret=0.0f;
sh_data[idx]=table_1[idx];
for(inti=0;i<num;i++)
{
ret+=sh_data[idx%BANK_CONFLICT];
}
result[idx]=ret;
}
这里的sh_data就是固定大小的数组;
2.使用动态分配的数组:
extern__shared__chararray[];
__global__voidshared_memory_1(float*result,intnum,float*table_1,intshared_size)
{
float*sh_data=(float*)array;//这里就让sh_data指向了sharedmemory的第一个地址,就可以动态分配空间
intidx=threadIdx.x;
floatret=0.0f;
sh_data[idx]=table_1[idx];
for(inti=0;i<num;i++)
{
ret+=sh_data[idx%BANK_CONFLICT];
}
result[idx]=ret;
}
这里是动态分配的空间,extern
__shared__char
array[];指定了shared的第一个变量的地址,这里其实是指向sharedmemory空间地址;后面的动态分配float*
sh_data=(float*)array;让sh_data指向array其实就是指向sharedmemory上的第一个地址;
后面的float*
sh_data2=(float*)&sh_data[shared_size];这里的sh_data2是指向的第一个sh_data的shared_size的地址,就是sh_data就是有了shared_size的动态分配的空间;
入下图:
3.下面是讲解bankconflict
我们知道有每一个half-warp是16个thread,然后sharedmemory有16个bank,怎么分配这16个thread,分别到各自的bank去取sharedmemory,如果大家都到同一个bank取款,就会排队,这就造成了bankconflict,上面的代码可以用来验证一下bankconflict对代码性能造成的影响:
/************************************************************************/
/*Example*/
/************************************************************************/
__global__voidshared_memory_1(float*result,intnum,float*table_1)
{
__shared__floatsh_data[THREAD_SIZE];
intidx=threadIdx.x;
floatret=0.0f;
sh_data[idx]=table_1[idx];
for(inti=0;i<num;i++)
{
ret+=sh_data[idx%BANK_CONFLICT];
}
result[idx]=ret;
}
//1,2,3,4,5,6,7.....16
#defineBANK_CONFLICT16
这里的BANK_CONFLICT定义为从1到16的大小,可以自己修改,来看看bankconflict对性能的影响;当BANK_CONFLICT为2的时候,就会通用有8个thread同时访问同一个bank,因为idx%2的取值只有2个0和1,所以16个都会访问bank0和bank1,以此类推,就可以测试整个的性能;
下面为示意图:
当然我们还可以利用16bankconflict,大家都访问同一个bank的同一个数据的时候,就可以形成一个broadcast,那样就会把数据同时广播给16个thread,这样就可以合理利用sharedmemory的broadcast的机会。
下面贴出代码,最好自己测试一下;
/********************************************************************
*shared_memory_test.cu
*ThisisaexampleoftheCUDAprogram.
*Author:zhao.kaiyong(at)gmail.com
*'target='_blank'>http://blog.csdn.net/openhero[/code] *'target='_blank'>http://www.comp.hkbu.edu.hk/~kyzhao/[/code] *********************************************************************/#include<stdio.h>#include<stdlib.h>#include<cutil.h>#include<cutil_inline.h>//1,2,3,4,5,6,7.....16#defineBANK_CONFLICT16#defineTHREAD_SIZE16/************************************************************************//*static*//************************************************************************/__global__voidshared_memory_static(float*result,intnum,float*table_1){__shared__floatsh_data[THREAD_SIZE];intidx=threadIdx.x;floatret=0.0f;sh_data[idx]=table_1[idx];for(inti=0;i<num;i++){ret+=sh_data[idx%BANK_CONFLICT];}result[idx]=ret;}/************************************************************************//*dynamic*//************************************************************************/extern__shared__chararray[];__global__voidshared_memory_dynamic(float*result,intnum,float*table_1,intshared_size){float*sh_data=(float*)array;//这里就让sh_data指向了sharedmemory的第一个地址,就可以动态分配空间float*sh_data2=(float*)&sh_data[shared_size];//这里的shared_size的大小为sh_data的大小;intidx=threadIdx.x;floatret=0.0f;sh_data[idx]=table_1[idx];for(inti=0;i<num;i++){ret+=sh_data[idx%BANK_CONFLICT];}result[idx]=ret;}/************************************************************************//*Bankconflict*//************************************************************************/__global__voidshared_memory_bankconflict(float*result,intnum,float*table_1){__shared__floatsh_data[THREAD_SIZE];intidx=threadIdx.x;floatret=0.0f;sh_data[idx]=table_1[idx];for(inti=0;i<num;i++){ret+=sh_data[idx%BANK_CONFLICT];}result[idx]=ret;}/************************************************************************//*HelloCUDA*//************************************************************************/intmain(intargc,char*argv[]){if(cutCheckCmdLineFlag(argc,(constchar**)argv,"device")){cutilDeviceInit(argc,argv);}else{intid=cutGetMaxGflopsDeviceId();cudaSetDevice(id);}float*device_result=NULL;floathost_result[THREAD_SIZE]={0};CUDA_SAFE_CALL(cudaMalloc((void**)&device_result,sizeof(float)*THREAD_SIZE));float*device_table_1=NULL;floathost_table1[THREAD_SIZE]={0};for(inti=0;i<THREAD_SIZE;i++){host_table1[i]=rand()%RAND_MAX;}CUDA_SAFE_CALL(cudaMalloc((void**)&device_table_1,sizeof(float)*THREAD_SIZE));CUDA_SAFE_CALL(cudaMemcpy(device_table_1,host_table1,sizeof(float)*THREAD_SIZE,cudaMemcpyHostToDevice));unsignedinttimer=0;CUT_SAFE_CALL(cutCreateTimer(&timer));CUT_SAFE_CALL(cutStartTimer(timer));shared_memory_static<<<1,THREAD_SIZE>>>(device_result,1000,device_table_1);//shared_memory_dynamic<<<1,THREAD_SIZE>>>(device_result,1000,device_table_1,16);//shared_memory_bankconflict<<<1,THREAD_SIZE>>>(device_result,1000,device_table_1);CUT_CHECK_ERROR("Kernelexecutionfailed");CUDA_SAFE_CALL(cudaMemcpy(host_result,device_result,sizeof(float)*THREAD_SIZE,cudaMemcpyDeviceToHost));CUT_SAFE_CALL(cutStopTimer(timer));printf("Processingtime:%f(ms)",cutGetTimerValue(timer));CUT_SAFE_CALL(cutDeleteTimer(timer));for(inti=0;i<THREAD_SIZE;i++){printf("%f",host_result[i]);}CUDA_SAFE_CALL(cudaFree(device_result));CUDA_SAFE_CALL(cudaFree(device_table_1));cutilExit(argc,argv);}
这里只是一个简单的demo,大家可以测试一下。下一章节会将一些sharedmemory的更多的特性,更深入的讲解sharedmemory的一些隐藏的性质;
再在接下来的章节会讲一些constant和texture的使用;
写的内容一直都是文字比较多,代码比较少,其实学习的过程更重要的思想,实践的代码,最好是自己写,唯一可以学习的是思想,学习更重要的也是思想的交流,知识的传播,最好的是思想的传播,代码,方法,都是只是一些工具而已。但是工具的熟练层度,就得靠自己下来多练习。
相关文章推荐
- (CUDA 编程8).CUDA 内存使用 global 二------GPU的革命
- (CUDA 编程10).CUDA cosnstant使用(一)------GPU的革命
- 9. CUDA shared memory使用------GPU的革命
- (CUDA 编程1).CUDA 线程执行模型分析(一)招兵 ------ GPU的革命
- CUDA GPU编程中使用结构体传递函数参数
- CUDA编程接口(一)------一十八般武器------GPU的革命 推荐
- CUDA 编程1).CUDA 线程执行模型分析(一)招兵 ------ GPU的革命(转)
- 《GPU高性能编程-CUDA实战》中例子头文件使用
- 使用GPU进行字符串匹配--cuda编程实现
- 《GPU高性能编程-CUDA实战》中例子头文件使用 (2012-12-25 20:45:48)
- 使用NSight进行CUDA调试,只能进行GPU代码调试,不能进入CPU端代码断点
- CUDA下的GPU编程--线程和变量
- CUDA下的GPU编程--线程和变量
- CUDA硬件实现分析(二)------规行矩步------GPU的革命
- 【并行计算-CUDA开发】GPU并行编程方法
- 《GPU高性能编程 CUDA实战》(CUDA By Example)读书笔记
- 看cuda初级教程视频笔记(周斌讲的)--CUDA、GPU编程模型
- 第一.CUDA的存在意义与其它GPU编程环境之个人愚见
- GPU 编程入门到精通(一)之 CUDA 环境安装
- 【并行计算-CUDA开发】CUDA编程——GPU架构,由sp,sm,thread,block,grid,warp说起