您的位置:首页 > 编程语言

(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的第一个地址,就可以动态分配空间

float*sh_data2=(float*)&sh_data[shared_size];//这里的shared_size的大小为sh_data的大小;[/code]
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的使用;

写的内容一直都是文字比较多,代码比较少,其实学习的过程更重要的思想,实践的代码,最好是自己写,唯一可以学习的是思想,学习更重要的也是思想的交流,知识的传播,最好的是思想的传播,代码,方法,都是只是一些工具而已。但是工具的熟练层度,就得靠自己下来多练习。
内容来自用户分享和网络整理,不保证内容的准确性,如有侵权内容,可联系管理员处理 点击这里给我发消息
标签: