亦庄做网站,网站怎么做付费项目,根据 我司申请 网站建设,张店网站建设yx718http://blog.csdn.net/openhero/article/details/3890578
首先是讲一下shared的memory的两种使用方法#xff0c;然后讲解一下shared memory的bank conflict的问题#xff0c;这个是shared memory访问能否高效的问题所在#xff1b;
Shared memory的常规使用#xff1a;…http://blog.csdn.net/openhero/article/details/3890578
首先是讲一下shared的memory的两种使用方法然后讲解一下shared memory的bank conflict的问题这个是shared memory访问能否高效的问题所在
Shared memory的常规使用
1. 使用固定大小的数组
/************************************************************************/
/* Example */
/************************************************************************/
__global__ void shared_memory_1(float* result, int num, float* table_1)
{
__shared__ float sh_data[THREAD_SIZE];
int idx threadIdx.x;
float ret 0.0f;
sh_data[idx] table_1[idx];
for (int i 0; i num; i)
{
ret sh_data[idx %BANK_CONFLICT];
}
result[idx] ret;
}
这里的sh_data就是固定大小的数组 2. 使用动态分配的数组
extern __shared__ char array[];
__global__ void shared_memory_1(float* result, int num, float* table_1, int shared_size)
{
float* sh_data (float*)array; // 这里就让sh_data指向了shared memory的第一个地址就可以动态分配空间
float* sh_data2 (float*)sh_data[shared_size]; // 这里的shared_size的大小为sh_data的大小
int idx threadIdx.x;
float ret 0.0f;
sh_data[idx] table_1[idx];
for (int i 0; i num; i)
{
ret sh_data[idx %BANK_CONFLICT];
}
result[idx] ret;
}
这里是动态分配的空间extern __shared__ char array[];指定了shared的第一个变量的地址这里其实是指向shared memory空间地址后面的动态分配float* sh_data (float*)array;让sh_data指向array其实就是指向shared memory上的第一个地址
后面的float* sh_data2 (float*)sh_data[shared_size];这里的sh_data2是指向的第一个sh_data的shared_size的地址就是sh_data就是有了shared_size的动态分配的空间 3. 下面是讲解bank conflict
我们知道有每一个half-warp是16个thread然后shared memory有16个bank怎么分配这16个thread分别到各自的bank去取shared memory如果大家都到同一个bank取款就会排队这就造成了bank conflict上面的代码可以用来验证一下bank conflict对代码性能造成的影响
/************************************************************************/
/* Example */
/************************************************************************/
__global__ void shared_memory_1(float* result, int num, float* table_1)
{
__shared__ float sh_data[THREAD_SIZE];
int idx threadIdx.x;
float ret 0.0f;
sh_data[idx] table_1[idx];
for (int i 0; i num; i)
{
//并行的思想
ret sh_data[idx %BANK_CONFLICT];
}
result[idx] ret;
}
// 1,2,3,4,5,6,7.....16
#define BANK_CONFLICT 16
这里的BANK_CONFLICT 定义为从1到16的大小可以自己修改来看看bank conflict对性能的影响当BANK_CONFLICT为2的时候就会通用有8个thread同时访问同一个bank因为idx%2的取值只有2个0和1所以16个都会访问bank0和bank1以此类推就可以测试整个的性能
当然我们还可以利用16bank conflict大家都访问同一个bank的同一个数据的时候就可以形成一个broadcast那样就会把数据同时广播给16个thread这样就可以合理利用shared memory的broadcast的机会。
下面贴出代码最好自己测试一下
/********************************************************************
* shared_memory_test.cu
* This is a example of the CUDA program.
* Author: zhao.kaiyong(at)gmail.com
* http://blog.csdn.NET/openhero
* http://www.comp.hkbu.edu.hk/~kyzhao/
*********************************************************************/
#include
#include
#include
#include
// 1,2,3,4,5,6,7.....16
#define BANK_CONFLICT 16
#define THREAD_SIZE 16
/************************************************************************/
/* static */
/************************************************************************/
__global__ void shared_memory_static(float* result, int num, float* table_1)
{
__shared__ float sh_data[THREAD_SIZE];
int idx threadIdx.x;
float ret 0.0f;
sh_data[idx] table_1[idx];
for (int i 0; i num; i)
{
ret sh_data[idx%BANK_CONFLICT];
}
result[idx] ret;
}
/************************************************************************/
/* dynamic */
/************************************************************************/
extern __shared__ char array[];
__global__ void shared_memory_dynamic(float* result, int num, float* table_1, int shared_size)
{
float* sh_data (float*)array; // 这里就让sh_data指向了shared memory的第一个地址就可以动态分配空间
float* sh_data2 (float*)sh_data[shared_size]; // 这里的shared_size的大小为sh_data的大小
int idx threadIdx.x;
float ret 0.0f;
sh_data[idx] table_1[idx];
for (int i 0; i num; i)
{
ret sh_data[idx%BANK_CONFLICT];
}
result[idx] ret;
}
/************************************************************************/
/* Bank conflict */
/************************************************************************/
__global__ void shared_memory_bankconflict(float* result, int num, float* table_1)
{
__shared__ float sh_data[THREAD_SIZE];
int idx threadIdx.x;
float ret 0.0f;
sh_data[idx] table_1[idx];
for (int i 0; i num; i)
{
ret sh_data[idx % BANK_CONFLICT];
}
result[idx] ret;
}
/************************************************************************/
/* HelloCUDA */
/************************************************************************/
int main(int argc, char* argv[])
{
if ( cutCheckCmdLineFlag(argc, (const char**) argv, device))
{
cutilDeviceInit(argc, argv);
}else
{
int id cutGetMaxGflopsDeviceId();
cudaSetDevice(id);
}
float *device_result NULL;
float host_result[THREAD_SIZE] {0};
CUDA_SAFE_CALL( cudaMalloc((void**) device_result, sizeof(float) * THREAD_SIZE));
float *device_table_1 NULL;
float host_table1[THREAD_SIZE] {0};
for (int i 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));
unsigned int timer 0;
CUT_SAFE_CALL( cutCreateTimer( timer));
CUT_SAFE_CALL( cutStartTimer( timer));
shared_memory_static1, THREAD_SIZE(device_result, 1000, device_table_1);
//shared_memory_dynamic1, THREAD_SIZE(device_result, 1000, device_table_1, 16);
//shared_memory_bankconflict1, THREAD_SIZE(device_result, 1000, device_table_1);
CUT_CHECK_ERROR(Kernel execution failed/n);
CUDA_SAFE_CALL( cudaMemcpy(host_result, device_result, sizeof(float) * THREAD_SIZE, cudaMemcpyDeviceToHost));
CUT_SAFE_CALL( cutStopTimer( timer));
printf(Processing time: %f (ms)/n, cutGetTimerValue( timer));
CUT_SAFE_CALL( cutDeleteTimer( timer));
for (int i 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大家可以测试一下。下一章节会将一些shared memory的更多的特性更深入的讲解shared memory的一些隐藏的性质
