我想声明一次我的纹理并在我的所有内核和文件中使用它。因此,我将其声明为 extern
在 header 中,并在所有其他文件中包含该 header (遵循 SO How do I use extern to share variables between source files? )
我有一个标题 cudaHeader.cuh
包含我的纹理的文件:
extern texture<uchar4, 2, cudaReadModeElementType> texImage;
在我的file1.cu
,我分配 CUDA 数组并将其绑定(bind)到纹理:
cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc< uchar4 >( );
cudaStatus=cudaMallocArray( &cu_array_image, &channelDesc, width, height );
if (cudaStatus != cudaSuccess) {
fprintf(stderr, "cudaMallocArray failed! cu_array_image couldn't be created.\n");
return cudaStatus;
}
cudaStatus=cudaMemcpyToArray( cu_array_image, 0, 0, image, size_image, cudaMemcpyHostToDevice);
if (cudaStatus != cudaSuccess) {
fprintf(stderr, "cudaMemcpyToArray failed! Copy from the host memory to the device texture memory failed.\n");
return cudaStatus;
}
// set texture parameters
texImage.addressMode[0] = cudaAddressModeWrap;
texImage.addressMode[1] = cudaAddressModeWrap;
texImage.filterMode = cudaFilterModePoint;
texImage.normalized = false; // access with normalized texture coordinates
// Bind the array to the texture
cudaStatus=cudaBindTextureToArray( texImage, cu_array_image, channelDesc);
if (cudaStatus != cudaSuccess) {
fprintf(stderr, "cudaBindTextureToArray failed! cu_array couldn't be bind to texImage.\n");
return cudaStatus;
}
在 file2.cu
,我使用kernel
中的纹理功能如下:
__global__ void kernel(int width, int height, unsigned char *dev_image) {
int x = blockIdx.x*blockDim.x + threadIdx.x;
int y = blockIdx.y*blockDim.y + threadIdx.y;
if(y< height) {
uchar4 tempcolor=tex2D(texImage, x, y);
//if(tempcolor.x==0)
// printf("tempcolor.x %d \n", tempcolor.x);
dev_image[y*width*3+x*3]= tempcolor.x;
dev_image[y*width*3+x*3+1]= tempcolor.y;
dev_image[y*width*3+x*3+2]= tempcolor.z;
}
}
问题是当我在 file2.cu
中使用它时,我的纹理不包含任何内容或损坏的值。 。即使我使用函数 kernel
直接在file1.cu
,数据不正确。
如果我添加:texture<uchar4, 2, cudaReadModeElementType> texImage;
在file1.cu
和file2.cu
,编译器说有一个重新定义。
编辑:
我在 CUDA 版本 5.0
上尝试了同样的操作但出现同样的问题。如果我打印texImage
的地址在file1.cu
和file2.cu
,我没有相同的地址。变量 texImage
的声明肯定有问题.
最佳答案
这是一个非常古老的问题,talonmies 和 Tom 的评论中提供了答案。在 CUDA 5.0
之前的场景中,extern
纹理不可行,因为缺乏真正的链接器,导致 extern
链接的可能性。因此,正如汤姆提到的,
you can have different compilation units, but they cannot reference each other
在后 CUDA 5.0
场景中,extern
纹理是可能的,我想在下面提供一个简单的示例,展示这一点,希望它对其他用户。
kernel.cu编译单元
#include <stdio.h>
texture<int, 1, cudaReadModeElementType> texture_test;
/********************/
/* CUDA ERROR CHECK */
/********************/
#define gpuErrchk(ans) { gpuAssert((ans), __FILE__, __LINE__); }
inline void gpuAssert(cudaError_t code, const char *file, int line, bool abort=true)
{
if (code != cudaSuccess)
{
fprintf(stderr,"GPUassert: %s %s %d\n", cudaGetErrorString(code), file, line);
if (abort) exit(code);
}
}
/*************************/
/* LOCAL KERNEL FUNCTION */
/*************************/
__global__ void kernel1() {
printf("ThreadID = %i; Texture value = %i\n", threadIdx.x, tex1Dfetch(texture_test, threadIdx.x));
}
__global__ void kernel2();
/********/
/* MAIN */
/********/
int main() {
const int N = 16;
// --- Host data allocation and initialization
int *h_data = (int*)malloc(N * sizeof(int));
for (int i=0; i<N; i++) h_data[i] = i;
// --- Device data allocation and host->device memory transfer
int *d_data; gpuErrchk(cudaMalloc((void**)&d_data, N * sizeof(int)));
gpuErrchk(cudaMemcpy(d_data, h_data, N * sizeof(int), cudaMemcpyHostToDevice));
gpuErrchk(cudaBindTexture(NULL, texture_test, d_data, N * sizeof(int)));
kernel1<<<1, 16>>>();
gpuErrchk(cudaPeekAtLastError());
gpuErrchk(cudaDeviceSynchronize());
kernel2<<<1, 16>>>();
gpuErrchk(cudaPeekAtLastError());
gpuErrchk(cudaDeviceSynchronize());
gpuErrchk(cudaUnbindTexture(texture_test));
}
kernel2.cu编译单元
#include <stdio.h>
extern texture<int, 1, cudaReadModeElementType> texture_test;
/**********************************************/
/* DIFFERENT COMPILATION UNIT KERNEL FUNCTION */
/**********************************************/
__global__ void kernel2() {
printf("Texture value = %i\n", tex1Dfetch(texture_test, threadIdx.x));
}
记得编译生成可重定位设备代码,即-rdc = true
,以启用外部链接
关于CUDA 外部纹理声明,我们在Stack Overflow上找到一个类似的问题: https://stackoverflow.com/questions/12852416/