1

我在尝试从 2D CUDA 数组中检索矩阵元素 A(m,n) 时遇到了很多麻烦。有趣的是，当 m = n 时，我得到了正确的元素；即元素沿对角线。否则，我会得到一些意想不到的行为：例如，如果我想获取元素 A(13,12)，并尝试使用 tex2D(tex, row + 0.5f, col + 0.5f) 检索它，我会得到 A (14,11)。据我所知，我正在做我应该做的一切，所以我真的很想知道我哪里出错了。

内核如下。错误发生在前两个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 回答