1

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

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

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;
}
4

1 に答える 1

2

回答：threadIdx.xをthreadIdx.yと交換してください。最終的には、セマンティクスの問題になります。テクスチャは、私たちがよく知っているx軸とy軸に沿ったオフセットとしてインデックスを使用します。行列は、インデックスを使用して行と列のインデックスを参照します。基本的に、基底ベクトルは交換されます。

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

于 2012-05-07T01:14:59.427 に答える