CUDA：tex2D（）を使用する場合の問題

Question

2D CUDA配列から行列要素A（m、n）を取得しようとすると、多くの問題が発生しています。興味深いことに、m = nのときに正しい要素を得る。要素は対角線に沿っている。それ以外の場合は、予期しない動作が発生します。たとえば、要素A（13,12）をフェッチして、tex2D（tex、row + 0.5f、col + 0.5f）を使用して取得しようとすると、A （14,11）。私が知る限り、私はすべきことをすべてやっているので、私は本当にどこに間違っていたのか知​​りたいと思っています。

カーネルは以下のとおりです。間違いは、最初の2回の呼び出しの直後に発生します。したがって、残りの部分は実際には関係ありません。

texture<float, 2, cudaReadModeElementType> tex_a; 
texture<float, 2, cudaReadModeElementType> tex_b; 

// Assume that BinaryFunc is multiplication, and AccumulationFunc is addition. 
// Then this kernel computes the standard matrix product, and uses prefetching 
// with tile sizes given by the template parameter TileSize. 
template <unsigned TileSize, class T, class SizeType, class BinaryFunc, 
     class AccumulationFunc> 
    __global__ void 
matrix_prod_tex_prefetch(T* c, const SizeType dim, BinaryFunc binary_func, 
     AccumulationFunc accum_func) 
{ 
    __shared__ T as[TileSize][TileSize]; 
    __shared__ T bs[TileSize][TileSize]; 
    SizeType row = blockIdx.y * TileSize + threadIdx.y; 
    SizeType col = blockIdx.x * TileSize + threadIdx.x; 
    T p = 0; 

    T l = tex2D(tex_a, row + 0.5f, threadIdx.x + 0.5f); 
    T m = tex2D(tex_b, threadIdx.y + 0.5f, col + 0.5f); 
    __syncthreads(); 

    for (SizeType i = 1; i != dim/TileSize; ++i) { 
     as[threadIdx.y][threadIdx.x] = l; 
     bs[threadIdx.y][threadIdx.x] = m; 
     __syncthreads(); 
     l = tex2D(tex_a, row + 0.5f, i * TileSize + threadIdx.x + 0.5f); 
     m = tex2D(tex_b, i * TileSize + threadIdx.y + 0.5f, col + 0.5f); 
     for (SizeType k = 0; k != TileSize; ++k) { 
      p = accum_func(p, binary_func(
         as[threadIdx.y][k], 
         bs[k][threadIdx.x] 
         )); 
     } 
     __syncthreads(); 
    } 

    as[threadIdx.y][threadIdx.x] = l; 
    bs[threadIdx.y][threadIdx.x] = m; 
    __syncthreads(); 
    for (SizeType k = 0; k != TileSize; ++k) { 
     p = accum_func(p, binary_func(
        as[threadIdx.y][k], 
        bs[k][threadIdx.x] 
        )); 
    } 
    c[dim * row + col] = p; 
} 

Answer 1

回答：threadIdx.xをthreadIdx.yでスワップします。最終的には、それはセマンティクスの問題になります。テクスチャは、我々がよく知っているx軸とy軸に沿ってオフセットをインデックスとして使用します。マトリックスは、インデックスを使用して行と列のインデックスを参照します。本質的に、基底ベクトルは入れ替えられる。

1DのメモリレイアウトでthreadIdx.xとthreadIdx.yを使用すると、同等の結果が得られる可能性があるため、プロセス内での合体メモリアクセスパターンが失われる可能性があることに注意してください。