CUDA: Проблемы с использованием tex2D()

Question

У меня возникли проблемы с попыткой извлечь матричный элемент A (m, n) из массива 2D CUDA. Интересно, что я получаю правильный элемент при 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; 
} 

Answer 1

Ответа на этот вопрос: swap threadIdx.x with threadIdx.y. В конечном счете, это сводится к вопросу о семантике: текстуры используют индексы в качестве смещений вдоль оси x и y, с которой мы все знакомы. Матрицы используют индексы для обращения к индексу строки и столбца. По существу, базовые векторы меняются местами.

Следует помнить, что при замене использования threadIdx.x и threadIdx.y для макетов памяти 1D могут быть получены эквивалентные результаты, вы можете потерять объединенные шаблоны доступа к памяти в процессе.