1

我试图从 CUDA C 中的 3D 纹理中获取数据，但函数调用 tex3D() 总是返回零。以下是相关代码：

主持人：

#define L 64
typedef uint2 splitspin_t;
texture<splitspin_t, 3> texRef;

cudaArray *arrayPointer;
cudaExtent LLLextent = make_cudaExtent(L, L, L);    
cudaChannelFormatDesc cf = cudaCreateChannelDesc<splitspin_t>();
cudaChk(cudaMalloc3DArray( &arrayPointer, &cf, LLLextent ));

cudaMemcpy3DParms params = {0};
params.extent = LLLextent;
params.kind = cudaMemcpyHostToDevice;

params.srcPtr.ptr = h; // size L*L*L*sizeof(splitspin_t) allocated by malloc
params.srcPtr.pitch = sizeof(splitspin_t) * L;
params.srcPtr.xsize = L;
params.srcPtr.ysize = L;
params.srcPos.x = 0;
params.srcPos.y = 0;
params.srcPos.z = 0;

params.dstArray = arrayPointer;
params.dstPos.x = 0;
params.dstPos.y = 0;
params.dstPos.z = 0;

cudaChk(cudaMemcpy3D( &params ));

texRef.normalized = 0;                     
texRef.filterMode = cudaFilterModePoint;      
texRef.addressMode[0] = cudaAddressModeClamp; 
texRef.addressMode[1] = cudaAddressModeClamp;
texRef.addressMode[2] = cudaAddressModeClamp;
cudaChk(cudaBindTextureToArray( texRef, arrayPointer, cf ));

cudaFreeArray(arrayPointer);

设备：

 #define GX (threadIdx.x + blockIdx.x*blockDim.x)
 #define GY (threadIdx.y + blockIdx.y*blockDim.y)
 #define GZ (threadIdx.z + blockIdx.z*blockDim.z)

 printf("%lX %lx\n", tex3D(texRef, GX, GY, GZ).y, tex3D(texRef, GX, GY, GZ).x); // always prints zeros

我已经验证了 h 指向的内存被初始化为非零。我还通过在第一个 cudaMemcpy3D 之后将 h 归零，使用第二个 cudaMemcpy3D 从 arrayPointer 复制回 h 来验证 cudaMemcpy3D 是成功的，然后检查 h 然后包含与以前相同的数据。我想也许问题也可能是由于我使用了非标准类型（uint2），但是 typedef-ing splitspin_t to float 并没有解决问题。

因此，我怀疑 cudaBindTextureToArray 函数调用，但我看不到我到目前为止所犯的任何错误。

提前致谢。

4

1 回答 1

1

我不认为你想这样做：

cudaFreeArray(arrayPointer);

直到您的程序（或至少执行纹理访问的内核）完成。

如果您查看像simpleCubemapTexture这样的 cuda 示例之一， 您会看到典型的序列是：

  1. 创建数组
  2. 复制到设备
  3. 绑定到纹理
  4. 调用正在处理纹理的内核
  5. 释放设备阵列

此外，tex3D(...).x 和 .y 返回的值是 int 类型。如果将长格式说明符 ( l) 与 一起使用printf，您可能会得到令人费解的结果。

以下代码对我有用，以上是我对您发布的内容所做的仅有的两个重大更改：

#include <stdio.h>

#define cudaCheckErrors(msg) \
    do { \
        cudaError_t __err = cudaGetLastError(); \
        if (__err != cudaSuccess) { \
            fprintf(stderr, "Fatal error: %s (%s at %s:%d)\n", \
                msg, cudaGetErrorString(__err), \
                __FILE__, __LINE__); \
            fprintf(stderr, "*** FAILED - ABORTING\n"); \
            exit(1); \
        } \
    } while (0)


#define L 64
typedef uint2 splitspin_t;
texture<splitspin_t, 3> texRef;

__global__ void my_kernel(){


 printf("%X %x\n", tex3D(texRef, 4, 4, 4).y, tex3D(texRef, 4, 4, 4).x);

}

int main(){

splitspin_t *h, temp;
temp.x = 16;
temp.y = 65536;
h=(splitspin_t *)malloc(L*L*L*sizeof(splitspin_t));
if (h==0) {printf("malloc fail\n"); return 1;}
for (int i=0; i< (L*L*L); i++)
  h[i] = temp;

cudaArray *arrayPointer;
cudaExtent LLLextent = make_cudaExtent(L, L, L);
cudaChannelFormatDesc cf = cudaCreateChannelDesc<splitspin_t>();
cudaMalloc3DArray( &arrayPointer, &cf, LLLextent );
cudaCheckErrors("cudaMalloc3DArray");

cudaMemcpy3DParms params = {0};
params.extent = LLLextent;
params.kind = cudaMemcpyHostToDevice;

params.srcPtr.ptr = h; // size L*L*L*sizeof(splitspin_t) allocated by malloc
params.srcPtr.pitch = sizeof(splitspin_t) * L;
params.srcPtr.xsize = L;
params.srcPtr.ysize = L;
params.srcPos.x = 0;
params.srcPos.y = 0;
params.srcPos.z = 0;

params.dstArray = arrayPointer;
params.dstPos.x = 0;
params.dstPos.y = 0;
params.dstPos.z = 0;

cudaMemcpy3D( &params );
cudaCheckErrors("cudaMemcpy3D");

texRef.normalized = 0;
texRef.filterMode = cudaFilterModePoint;
texRef.addressMode[0] = cudaAddressModeClamp;
texRef.addressMode[1] = cudaAddressModeClamp;
texRef.addressMode[2] = cudaAddressModeClamp;
cudaBindTextureToArray( texRef, arrayPointer, cf );
cudaCheckErrors("cudaBind");

my_kernel<<<1,1>>>();
cudaDeviceSynchronize();
cudaCheckErrors("kernel");
cudaFreeArray(arrayPointer);
return 0;
}

当我编译并运行它时，我得到的打印输出是：

$ ./t192

10000 10

我认为这是正确的。

于 2013-07-05T17:15:53.177 回答