Когда вы используете вариант pitch2D для текстурной операции, базовое распределение должно быть правильным Pitch распределением. Я думаю, что обычно люди создают это с cudaMallocPitch. Однако заявленное требование :
cudaResourceDesc :: res :: pitch2D :: pitchInBytes указывает шаг между двумя строками в байтах и должен быть выровнен по cudaDeviceProp :: texturePitchAlignment.
В моем графическом процессоре это последнее свойство равно 32. Я не знаю о вашем графическом процессоре, но держу пари, что это свойство не равно 4 для вашего графического процессора. Однако вы указываете 4 здесь:
resDesc.res.pitch2D.pitchInBytes = 4;
Опять же, я думаю, что люди обычно используют для этого выделенное распределение, созданное с помощью cudaMallocPitch. Однако мне представляется возможным пройти обычное линейное распределение, если размерность от строки к строке (в байтах) делится на texturePitchAlignment (32 в моем случае).
Еще одно изменение, которое я сделал, - это использование cudaCreateChannelDesc<>() вместо ручной настройки параметров, как вы сделали. Это создает другой набор параметров desc и, похоже, также влияет на результат. Не должно быть трудно изучить различия.
Когда я настраиваю ваш код для решения этих проблем, я получаю результаты, которые кажутся мне разумными:
$ cat t30.cu
#include <stdio.h>
#include <stdint.h>
typedef uint8_t mt; // use an integer type
__global__ void kernel(cudaTextureObject_t tex)
{
int x = threadIdx.x;
int y = threadIdx.y;
mt val = tex2D<mt>(tex, x, y);
printf("%d, ", val);
}
int main(int argc, char **argv)
{
cudaDeviceProp prop;
cudaGetDeviceProperties(&prop, 0);
printf("texturePitchAlignment: %lu\n", prop.texturePitchAlignment);
cudaTextureObject_t tex;
const int num_rows = 4;
const int num_cols = prop.texturePitchAlignment*1; // should be able to use a different multiplier here
const int ts = num_cols*num_rows;
const int ds = ts*sizeof(mt);
mt dataIn[ds];
for (int i = 0; i < ts; i++) dataIn[i] = i;
mt* dataDev = 0;
cudaMalloc((void**)&dataDev, ds);
cudaMemcpy(dataDev, dataIn, ds, cudaMemcpyHostToDevice);
struct cudaResourceDesc resDesc;
memset(&resDesc, 0, sizeof(resDesc));
resDesc.resType = cudaResourceTypePitch2D;
resDesc.res.pitch2D.devPtr = dataDev;
resDesc.res.pitch2D.width = num_cols;
resDesc.res.pitch2D.height = num_rows;
resDesc.res.pitch2D.desc = cudaCreateChannelDesc<mt>();
resDesc.res.pitch2D.pitchInBytes = num_cols*sizeof(mt);
struct cudaTextureDesc texDesc;
memset(&texDesc, 0, sizeof(texDesc));
cudaCreateTextureObject(&tex, &resDesc, &texDesc, NULL);
dim3 threads(4, 4);
kernel<<<1, threads>>>(tex);
cudaDeviceSynchronize();
printf("\n");
return 0;
}
$ nvcc -o t30 t30.cu
$ cuda-memcheck ./t30
========= CUDA-MEMCHECK
texturePitchAlignment: 32
0, 1, 2, 3, 32, 33, 34, 35, 64, 65, 66, 67, 96, 97, 98, 99,
========= ERROR SUMMARY: 0 errors
$
Как и было сказано в комментариях, если бы я собирался сделать что-то подобное, но с использованием cudaMallocPitch и cudaMemcpy2D, это могло бы выглядеть примерно так:
$ cat t1421.cu
#include <stdio.h>
#include <stdint.h>
typedef uint8_t mt; // use an integer type
__global__ void kernel(cudaTextureObject_t tex)
{
int x = threadIdx.x;
int y = threadIdx.y;
mt val = tex2D<mt>(tex, x, y);
printf("%d, ", val);
}
int main(int argc, char **argv)
{
cudaDeviceProp prop;
cudaGetDeviceProperties(&prop, 0);
printf("texturePitchAlignment: %lu\n", prop.texturePitchAlignment);
cudaTextureObject_t tex;
const int num_rows = 4;
const int num_cols = prop.texturePitchAlignment*1; // should be able to use a different multiplier here
const int ts = num_cols*num_rows;
const int ds = ts*sizeof(mt);
mt dataIn[ds];
for (int i = 0; i < ts; i++) dataIn[i] = i;
mt* dataDev = 0;
size_t pitch;
cudaMallocPitch((void**)&dataDev, &pitch, num_cols*sizeof(mt), num_rows);
cudaMemcpy2D(dataDev, pitch, dataIn, num_cols*sizeof(mt), num_cols*sizeof(mt), num_rows, cudaMemcpyHostToDevice);
struct cudaResourceDesc resDesc;
memset(&resDesc, 0, sizeof(resDesc));
resDesc.resType = cudaResourceTypePitch2D;
resDesc.res.pitch2D.devPtr = dataDev;
resDesc.res.pitch2D.width = num_cols;
resDesc.res.pitch2D.height = num_rows;
resDesc.res.pitch2D.desc = cudaCreateChannelDesc<mt>();
resDesc.res.pitch2D.pitchInBytes = pitch;
struct cudaTextureDesc texDesc;
memset(&texDesc, 0, sizeof(texDesc));
cudaCreateTextureObject(&tex, &resDesc, &texDesc, NULL);
dim3 threads(4, 4);
kernel<<<1, threads>>>(tex);
cudaDeviceSynchronize();
printf("\n");
return 0;
}
$ nvcc -o t1421 t1421.cu
$ cuda-memcheck ./t1421
========= CUDA-MEMCHECK
texturePitchAlignment: 32
0, 1, 2, 3, 32, 33, 34, 35, 64, 65, 66, 67, 96, 97, 98, 99,
========= ERROR SUMMARY: 0 errors
$
Несмотря на то, что мы имеем здесь текстурные объекты, достаточно просто продемонстрировать, что подобные проблемы возникают с ссылками на текстуры. Вы не можете создать сколь угодно малую ссылку на 2D-текстуру так же, как вы не можете создать сколь угодно маленький объект 2D-текстуры. Я не собираюсь приводить демонстрацию этого также, поскольку это в значительной степени дублирует вышеприведенное, и людям больше не следует использовать ссылки на текстуры для новых разработок - лучше всего использовать объекты текстуры.