Cuda中Globalmemory中coalescing例程解释

Global memory是cuda中最常见的存储类型&＃xff0c;又叫做Device memory&＃xff0c;位于Host主机区域上&＃xff0c;它的生命周期是在整个Grid里面&＃xff0c;大约具有500个cycle latency。在cuda并行程序中&＃xff0c;尽量用Coalesing accessing的策略来最大化带宽bandwidth。什么是Coalesing accessing呢&＃xff1f;如图所示&＃xff1a;

当半个Warp的16个threads在一次memory transaction中coalesced时&＃xff0c;Global memory中的带宽得到了最大的利用。其中&＃xff0c;需要注意的是&＃xff0c;Device在一次transaction中&＃xff0c;从global memory中可以一次读取32-bit&＃xff0c;64-bit&＃xff0c;128-bit&＃xff0c;例如

64 bytes - each thread reads a word: int, float, …

128 bytes - each thread reads a double-word: int2, float2, …

32 bytes (compute capability 1.2&＃43;) - each thread reads a short  int.

下面有两个实例来说明Global memory中的coalescing问题&＃xff1a;

1&＃xff09;float3型Uncoalesced

在这段代码中&＃xff0c;float3有12个bytes&＃xff0c;不等于要求的4,8,16 bytes&＃xff0c;半个warp读取3个64bytes中非连续区域&＃xff0c;如图&＃xff1a;

有三种方法可以解决这个问题

1:使用shared memory&＃xff0c;也叫做3-step approach

假如每个block中使用256个threads&＃xff0c;这样一个thread block需要 sizeof(float3)*256 bytes的share memory空间&＃xff0c;每个thread读取3个单独的float型&＃xff0c;这实质上是指讲输入定义为float型&＃xff0c;在核函数里面讲读取在share memory中的float变量转换为float3型并进行操作&＃xff0c;最后再转换成float型输出&＃xff0c;如图;

代码如下&＃xff1a;

如果不好理解的话&＃xff0c;假设我们的blockDim&＃61;4&＃xff0c;取4个float3型变量&＃xff0c;我们会发现&＃xff0c;每一个thread中输入操作&＃xff08;输出操作一样&＃xff09;为&＃xff1a;

Thread 0&＃xff1a;

S_data[0]&＃61;g_in[0]; S_data[4]&＃61;g_in[4]; S_data[8]&＃61;g_in[8];

3&＃xff09;使用alignment specifiers

__align__(X), where X &＃61; 4, 8, or 16

struct __align__(16) {float x; float y;  float z; };

2&＃xff09;第二个实例&＃xff1a;矩阵转置 Matrix Transpose.

一般做法&＃xff1a;Uncoalesced Transpose&＃xff0c;GMEM为Global memory

我们发现一般的做法&＃xff0c;在写output时&＃xff0c;地址是不连续的&＃xff0c;即uncoalesced&＃xff0c;因此我们利用shared memory存储输入数据&＃xff0c;根据转置的关系&＃xff0c;来实现coalescing&＃xff0c;SMEM为shared memory&＃xff0c;如下图&＃xff1a;

代码如下&＃xff1a;

__global__ void transpose(float *odata, float *idata, int width, int height)

{

__shared__ float block[BLOCK_DIM*BLOCK_DIM];

unsigned int xBlock &＃61; blockDim.x * blockIdx.x;

unsigned int yBlock &＃61; blockDim.y * blockIdx.y;

unsigned int xIndex &＃61; xBlock &＃43; threadIdx.x;

unsigned int yIndex &＃61; yBlock &＃43; threadIdx.y;

unsigned int index_out, index_transpose;

if (xIndex

{

unsigned int index_in &＃61; width * yIndex &＃43; xIndex;

unsigned int index_block &＃61; threadIdx.y * BLOCK_DIM &＃43; threadIdx.x;

block[index_block] &＃61; idata[index_in];

index_transpose &＃61; threadIdx.x * BLOCK_DIM &＃43; threadIdx.y;

index_out &＃61; height * (xBlock &＃43; threadIdx.y) &＃43; yBlock &＃43; threadIdx.x;

}

__syncthreads();

if (xIndex

odata[index_out] &＃61; block[index_transpose];

程序的逻辑关系有时还挺绕的&＃xff0c;我们以一个4*4矩阵为例&＃xff0c;将逻辑关系展示如下&＃xff1a;

设dim3 gridDim(4,1), dim3 blockDim(1,4)&＃xff0c;以橙色block为例&＃xff0c;如输入数据时&＃xff0c;将其放入到sharememory中&＃xff0c;代码体现在&＃xff1a;