作者：Rain雨露Dew | 来源：互联网 | 2023-10-13 16:58
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
__global__ void accessFloat3(float3 *d_in,
float3* d_out)
{
int index &＃61; blockIdx.x * blockDim.x &＃43; threadIdx.x;
float3 a &＃61; d_in[index];
a.x &＃43;&＃61; 2;
a.y &＃43;&＃61; 2;
a.z &＃43;&＃61; 2;
d_out[index] &＃61; a;
}
在这段代码中&＃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];
Thread 1&＃xff1a;
S_data[1]&＃61;g_in[1]; S_data[5]&＃61;g_in[5]; S_data[9]&＃61;g_in[9];
Thread 2&＃xff1a;
S_data[2]&＃61;g_in[2]; S_data[6]&＃61;g_in[6]; S_data[10]&＃61;g_in[10];
Thread 3&＃xff1a;
S_data[3]&＃61;g_in[3]; S_data[7]&＃61;g_in[7]; S_data[11]&＃61;g_in[11];
可以看出&＃xff0c;对于每个thread同一时刻&＃xff08;similar step&＃xff09;的数据读入&＃xff0c;地址均是连续&＃xff0c;这样就达到了coalescing。
2&＃xff09;使用数组的结构体(SOA)来取代结构体的数组(AOS)
3&＃xff09;使用alignment specifiers
__align__(X), where X &＃61; 4, 8, or 16
struct __align__(16) {float x; float y; float z; };
尽管这损失了比较多的空间&＃xff1a;
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;
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];
接下来的代码实际上是将block的区域给换了,如左下图所示&＃xff0c;block换成了一列四种不同颜色的&＃xff0c;最终转置的矩阵如右下图所示&＃xff0c;从图示可以看出&＃xff0c;最终结果的坐标系Height、Width、blockIdx.x、blockIdx.y均对位变换了&＃xff0c;这时我们只需要找threadIdx.x&＃39;、threadIdx.y&＃39;与threadIdx.x、threadIdx.y之间的关系&＃xff0c;其实可以看出&＃xff0c;一个block里面的坐标系没有发生变换&＃xff0c;则threadIdx.x&＃39;&＃61;threadIdx.x&＃xff0c;threadIdx.y&＃39;&＃61;threadIdx.y&＃xff0c;所以代码如下&＃xff1a;
index_transpose &＃61; threadIdx.x * BLOCK_DIM &＃43; threadIdx.y;
index_out &＃61; height * (xBlock &＃43; threadIdx.y) &＃43; yBlock &＃43; threadIdx.x;
odata[index_out] &＃61; block[index_transpose];
总体来说&＃xff0c;Global memory中coalescing就是保证其在数据读取或者写入时&＃xff0c;使用连续的地址&＃xff0c;且地址所存储的变量尺寸为32、64、128 bit&＃xff0c;我们常常使用share memory来解决coalescing问题。
京公网安备 11010802041100号