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型輸出，如圖;
代碼如下： 如果不好理解的話，假設我們的blockDim=4，取4個float3型變量，我們會發現，每一個thread中輸入操作（輸出操作一樣）為： Thread 0： S_data[0]=g_in[0];?S_data[4]=g_in[4];?S_data[8]=g_in[8]; Thread 1： S_data[1]=g_in[1];?S_data[5]=g_in[5];?S_data[9]=g_in[9]; Thread 2： S_data[2]=g_in[2];?S_data[6]=g_in[6];?S_data[10]=g_in[10]; Thread 3： S_data[3]=g_in[3];?S_data[7]=g_in[7];?S_data[11]=g_in[11]; 可以看出，對于每個thread同一時刻（similar step）的數據讀入，地址均是連續，這樣就達到了coalescing。 2）使用數組的結構體(SOA)來取代結構體的數組(AOS) 3）使用alignment specifiers __align__(X), where X = 4, 8, or 16 struct __align__(16) {float x;?float y; ?float z;?}; 盡管這損失了比較多的空間： 2）第二個實例：矩陣轉置 Matrix Transpose. 一般做法：Uncoalesced Transpose，GMEM為Global memory 我們發現一般的做法，在寫output時，地址是不連續的，即uncoalesced，因此我們利用shared memory存儲輸入數據，根據轉置的關系，來實現coalescing，SMEM為shared memory，如下圖： 代碼如下： __global__ void transpose(float *odata, float *idata, int width, int height) { __shared__ float block[BLOCK_DIM*BLOCK_DIM]; unsigned int xBlock = blockDim.x * blockIdx.x; unsigned int yBlock = blockDim.y * blockIdx.y; unsigned int xIndex = xBlock + threadIdx.x; unsigned int yIndex = yBlock + threadIdx.y; unsigned int index_out, index_transpose; if (xIndex < width && yIndex < height) { unsigned int index_in = width * yIndex + xIndex; unsigned int index_block = threadIdx.y * BLOCK_DIM + threadIdx.x; block[index_block] = idata[index_in]; index_transpose = threadIdx.x * BLOCK_DIM + threadIdx.y; index_out = height * (xBlock + threadIdx.y) + yBlock + threadIdx.x; } __syncthreads(); if (xIndex < width && yIndex < height) odata[index_out] = block[index_transpose]; 程序的邏輯關系有時還挺繞的，我們以一個4*4矩陣為例，將邏輯關系展示如下： 設dim3 gridDim(4,1), dim3 blockDim(1,4)，以橙色block為例，如輸入數據時，將其放入到sharememory中，代碼體現在： unsigned int index_in = width * yIndex + xIndex; unsigned int index_block = threadIdx.y * BLOCK_DIM + threadIdx.x; block[index_block] = idata[index_in]; 接下來的代碼實際上是將block的區域給換了,如左下圖所示，block換成了一列四種不同顏色的，最終轉置的矩陣如右下圖所示，從圖示可以看出，最終結果的坐標系Height、Width、blockIdx.x、blockIdx.y均對位變換了，這時我們只需要找threadIdx.x'、threadIdx.y'與threadIdx.x、threadIdx.y之間的關系，其實可以看出，一個block里面的坐標系沒有發生變換，則threadIdx.x'=threadIdx.x，threadIdx.y'=threadIdx.y，所以代碼如下： index_transpose = threadIdx.x * BLOCK_DIM + threadIdx.y; index_out = height * (xBlock + threadIdx.y) + yBlock + threadIdx.x; odata[index_out] = block[index_transpose];

總體來說，Global memory中coalescing就是保證其在數據讀取或者寫入時，使用連續的地址，且地址所存儲的變量尺寸為32、64、128 bit，我們常常使用share memory來解決coalescing問題。

總結

以上是生活随笔為你收集整理的Cuda中Global memory中coalescing例程解释的全部內容，希望文章能夠幫你解決所遇到的問題。

如果覺得生活随笔網站內容還不錯，歡迎將生活随笔推薦給好友。