(CUDA 编程9).CUDA shared memory使用------GPU的革命

作者：赵开勇 来源:http://www.hpctech.com/2009/0818/207.html

这一章节准备写一下shared memory的访问的问题，首先是讲一下shared的memory的两种使用方法，然后讲解一下shared memory的bank conflict的问题，这个是shared memory访问能否高效的问题所在

9. CUDA shared memory使用------GPU的革命

书接上文《8. CUDA内存使用 global 二------GPU的革命》 讲了global内存访问的时候，需要对齐的问题，只有在对齐的情况下才能保证global内存的高效访问。这一章节准备写一下shared memory的访问的问题，首先是讲一下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 <stdio.h>

#include <stdlib.h>

#include <cutil.h>

#include <cutil_inline.h>

// 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 ");

CUDA_SAFE_CALL( cudaMemcpy(host_result, device_result, sizeof(float) * THREAD_SIZE, cudaMemcpyDeviceToHost));

CUT_SAFE_CALL( cutStopTimer( timer));

printf("Processing time: %f (ms) ", 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的一些隐藏的性质；

再在接下来的章节会讲一些constant和texture的使用；

写的内容一直都是文字比较多，代码比较少，其实学习的过程更重要的思想，实践的代码，最好是自己写，唯一可以学习的是思想，学习更重要的也是思想的交流，知识的传播，最好的是思想的传播，代码，方法，都是只是一些工具而已。但是工具的熟练层度，就得靠自己下来多练习。