extern __shared__
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_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("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的一些隱藏的性質；
總結
以上是生活随笔為你收集整理的extern __shared__的全部內容，希望文章能夠幫你解決所遇到的問題。
- 上一篇： caffe blob操作
- 下一篇： CUDA的global内存访问的问题