CUDA外部紋理聲明

我想申報我的紋理一次，並在我的所有內核和文件中使用它。因此，我宣佈它作爲一個頭extern幷包括所有其他文件頭（以下的SO How do I use extern to share variables between source files?）CUDA外部紋理聲明

我有一個包含我的紋理頭cudaHeader.cuh文件：

extern texture<uchar4, 2, cudaReadModeElementType> texImage;

在我file1.cu，我分配我的CUDA數組，並將其綁定到紋理：

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中使用它時，我的紋理不包含任何內容或損壞的值。即使我在file1.cu中直接使用功能kernel，數據也不正確。

如果我在file1.cu和file2.cu中加上texture<uchar4, 2, cudaReadModeElementType> texImage;，編譯器會說有重新定義。

編輯：

我想同樣的事情CUDA版本5.0，但仍然出現同樣的問題。如果我在file1.cu和file2.cu中輸入texImage的地址，我沒有相同的地址。變量texImage的聲明必須存在問題。

來源

2012-10-12 Seltymar

什麼CUD你使用的是哪一個版本？ – talonmies

@talonmies CUDA是4.2，我使用CC 2.0的quadro 5000。 – Seltymar

好吧，一行答案是否定的，你不能這樣做。您正在使用的CUDA版本沒有設備代碼鏈接器。所有設備符號 - 內核，紋理和其他內存聲明*必須在同一個編譯單元中進行。不允許外部鏈接。 – talonmies

這是一個非常古老的問題，答案在talonmies和Tom的評論中提供。在前CUDA 5.0場景中，extern紋理是不可行的，因爲缺乏真正的鏈接器導致連接的可能性。因此，和湯姆提到，

你可以有不同的編譯單元，但他們不能相互引用

在後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，以使外部連桿

來源

2014-11-03 20:44:08 JackOLantern

CUDA外部紋理聲明

回答

相關問題