2

テクスチャを一度宣言して、すべてのカーネルとファイルで使用したいと思います。したがって、ヘッダーとして宣言し、extern他のすべてのファイルにヘッダーを含めます（SOに従って、externを使用してソースファイル間で変数を共有するにはどうすればよいですか？）

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です。で関数kernelを直接使用してもfile1.cu、データが正しくありません。

を追加texture<uchar4, 2, cudaReadModeElementType> texImage;するfile1.cuとfile2.cu、コンパイラは再定義があると言います。

編集：

CUDAバージョンでも同じことを試しました5.0が、同じ問題が発生します。とのアドレスを印刷するtexImageとfile1.cu、file2.cu同じアドレスになりません。変数の宣言に問題があるはずですtexImage。

4

1 に答える 1

3

これは非常に古い質問であり、タロンミーとトムのコメントで回答が提供されました。CUDA 以前の5.0シナリオでは、リンケージの可能性externにつながる真のリンカーがないため、テクスチャは実現できませんでした。externその結果、Tom が述べたように、

異なるコンパイル単位を持つことはできますが、相互に参照することはできません

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-03T20:44:08.327 に答える