__global__ void Matrixcopy(float *odata, const float *idata)
{
// threadblock size = (TILE_DIM, BLOCK_ROWS) = (32, 8)
// each block copies a 32 * 32 tile
int x = blockIdx.x * TILE_DIM threadIdx.x;
int y = blockIdx.y * TILE_DIM threadIdx.y;
int width = gridDim.x * TILE_DIM;
for (int j = 0; j < TILE_DIM; j = BLOCK_ROWS)
odata[(y j)*width x] = idata[(y j)*width x];
}
我對多維陣列的合并訪問的概念感到非常困惑。合并的全域記憶體訪問的定義是
順序記憶體訪問相鄰
根據文獻Learn CUDA Programming。對于一維陣列，很容易理解執行緒被索引為threadIdx.x blockDim.x * blockIdx.x，它可以很容易地映射到真正的一維陣列：warp 中的相鄰執行緒訪問一維陣列的相鄰物理地址。
然而對于二維陣列或矩陣和二維執行緒塊，如上面的代碼，我不確定我是否理解正確：經線中的相鄰執行緒位于同一行，即相同的 y 值，不同的x 值。如果 y = 0，則連續執行緒為 x = [1, 2, 3, 4, 5, 6...]，并且它們訪問連續地址 [1, 2, 3, 4, 5, 6...] 如果j = 0。所以這段代碼是合并訪問的。我理解正確嗎？而這只是一個簡單的cuda代碼，如果我們有一個復雜的cuda內核，我們如何快速判斷一個訪問是否合并？
uj5u.com熱心網友回復：
所以這段代碼是合并訪問的。我理解正確嗎？
是的，差不多。我會說執行緒是 x = [ 0 , 1, 2, 3, 4, 5, 6...]，并且它們訪問連續地址 [ 0 , 1, 2, 3, 4, 5, 6 ... ] 但基本上我們是一致的。
如果我們有一個復雜的 cuda 內核，我們如何快速確定訪問是否合并？
您可以查看任何索引構造并使用以下測驗：如果threadIdx.x變數作為加法因子包含在索引中，并且它沒有乘法因子，那么訪問將在典型用法中合并（如果您有平方執行緒塊）。任何可以表示為的索引：
idx = f threadIdx.x
wheref是任意的，但不包括threadIdx.x將導致合并訪問。“x”中的相鄰執行緒將訪問記憶體中的相鄰位置。對于“非方形”執行緒塊，您可以使用threadIdx.y. 例如，尺寸為 (1,32) 的螺紋塊將要求將threadIdx.y其作為僅可添加的因素包括在內。
轉載請註明出處，本文鏈接：https://www.uj5u.com/gongcheng/397873.html
上一篇：對容器命名要求的澄清
下一篇：右值從哪里來？[復制]