Cuda中Global memory中coalescing例程解释

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

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

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+) - each thread reads a short   int.

下面有两个实例来说明Global memory中的coalescing问题：

1）float3型Uncoalesced

__global__ void accessFloat3(float3 *d_in,

float3* d_out)

{

int index = blockIdx.x * blockDim.x + threadIdx.x;

float3 a = d_in[index];

a.x += 2;

a.y += 2;

a.z += 2;

d_out[index] = a;

}

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

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

1:使用shared memory，也叫做3-step approach

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

代码如下：