CUDAの1D配列にテクスチャメモリを使用する方法

Question

テクスチャメモリを1D配列に使用する方法を知るために、次のコードを記述しました.tex1D関数は対応するスレッドIDの配列から値を取り出していません。効率的かつ効果的に1Dアレイのテクスチャメモリを使用する方法。

__global__ void sum(float *b,cudaTextureObject_t texObj) 

    { 
    b[threadIdx.x]=tex1D<float>(texObj,threadIdx.x); 
    //printf("\n%f\n",tex1Dfetch<float>(texObj,threadIdx.x)); 
    } 
    int main() 
    { 
    float *a,*b; 
    float *d_a,*d_b; 
    int i; 
    a=(float*)malloc(sizeof(float)*5); 
    b=(float*)malloc(sizeof(float)*5); 

    for(i=0;i<5;i++) 
     a[i]=i; 

    cudaChannelFormatDesc channelDesc =cudaCreateChannelDesc(32, 0, 0, 0,cudaChannelFormatKindFloat); 

    cudaArray* cuArray; 
    cudaMallocArray(&cuArray, &channelDesc, 5, 0); 

    cudaMemcpyToArray(cuArray, 0, 0, a,sizeof(float)*5,cudaMemcpyHostToDevice); 


    struct cudaResourceDesc resDesc; 
     memset(&resDesc, 0, sizeof(resDesc)); 
     resDesc.resType = cudaResourceTypeArray; 
     resDesc.res.array.array = cuArray; 


     struct cudaTextureDesc texDesc; 
     memset(&texDesc, 0, sizeof(texDesc)); 
     texDesc.addressMode[0] = cudaAddressModeWrap; 
     texDesc.addressMode[1] = cudaAddressModeWrap; 
     texDesc.filterMode  = cudaFilterModeLinear; 
     texDesc.readMode   = cudaReadModeElementType; 
     texDesc.normalizedCoords = 1; 

     // Create texture object 
     cudaTextureObject_t texObj = 0; 
     cudaCreateTextureObject(&texObj, &resDesc, &texDesc, NULL); 


    cudaMalloc(&d_b, 5* sizeof(float)); 

    sum<<<1,5>>>(d_b,texObj); 



     // Free device memory 
    cudaMemcpy(b,d_b,sizeof(float),cudaMemcpyDeviceToHost); 

    for(i=0;i<5;i++) 
     printf("%f\t",b[i]); 
     cudaDestroyTextureObject(texObj); 
    cudaFreeArray(cuArray); 
    cudaFree(d_b); 

     return 0; 

    } 

Answer 1

は、少なくとも2の問題があります：あなただけの最後でホストするために戻ってデバイスから1フロート量をコピーしている

1. ：あなたは5を印刷したい場合は

   cudaMemcpy(b,d_b,sizeof(float),cudaMemcpyDeviceToHost); 
          ^^^^^^^^^^^^^ 
   

   値を5つ戻してコピーする必要があります。

   cudaMemcpy(b,d_b,5*sizeof(float),cudaMemcpyDeviceToHost); 
   

2. は、あなたがnormalized coordinates選択されました：

   b[threadIdx.x]=tex1D<float>(texObj,threadIdx.x); 
               ^^^^^^^^^^^ 
   

   使用の何か：

   texDesc.normalizedCoords = 1; 
   

   これはあなたのインデックスではなく、整数は0〜4から座標として0と1の間の座標浮動点を通過しなければならないことこれらの変化に対応し

   b[threadIdx.x]=tex1D<float>(texObj, ((float)threadIdx.x/5.0f)); 
   

：代わりに、このようなs、私は賢明な結果を得る。

$ cat t3.cu 
#include <stdio.h> 

__global__ void sum(float *b,cudaTextureObject_t texObj) 

    { 
    b[threadIdx.x]=tex1D<float>(texObj,((float)(threadIdx.x+1)/5.0f)); 

    //printf("\n%f\n",tex1Dfetch<float>(texObj,threadIdx.x)); 
    } 


int main() 
    { 
    float *a,*b; 
    float *d_b; 
    int i; 
    a=(float*)malloc(sizeof(float)*5); 
    b=(float*)malloc(sizeof(float)*5); 

    for(i=0;i<5;i++) 
     a[i]=i; 

    cudaChannelFormatDesc channelDesc =cudaCreateChannelDesc(32, 0, 0, 0,cudaChannelFormatKindFloat); 

    cudaArray* cuArray; 
    cudaMallocArray(&cuArray, &channelDesc, 5, 0); 

    cudaMemcpyToArray(cuArray, 0, 0, a,sizeof(float)*5,cudaMemcpyHostToDevice); 


    struct cudaResourceDesc resDesc; 
     memset(&resDesc, 0, sizeof(resDesc)); 
     resDesc.resType = cudaResourceTypeArray; 
     resDesc.res.array.array = cuArray; 


     struct cudaTextureDesc texDesc; 
     memset(&texDesc, 0, sizeof(texDesc)); 
     texDesc.addressMode[0] = cudaAddressModeWrap; 
     texDesc.addressMode[1] = cudaAddressModeWrap; 
     texDesc.filterMode  = cudaFilterModeLinear; 
     texDesc.readMode   = cudaReadModeElementType; 
     texDesc.normalizedCoords = 1; 

     // Create texture object 
     cudaTextureObject_t texObj = 0; 
     cudaCreateTextureObject(&texObj, &resDesc, &texDesc, NULL); 


    cudaMalloc(&d_b, 5* sizeof(float)); 

    sum<<<1,4>>>(d_b,texObj); 



     // Free device memory 
    cudaMemcpy(b,d_b,5*sizeof(float),cudaMemcpyDeviceToHost); 

    for(i=0;i<4;i++) 
     printf("%f\t",b[i]); 
     printf("\n"); 
     cudaDestroyTextureObject(texObj); 
    cudaFreeArray(cuArray); 
    cudaFree(d_b); 

     return 0; 

    } 
$ nvcc -arch=sm_61 -o t3 t3.cu 
$ cuda-memcheck ./t3 
========= CUDA-MEMCHECK 
0.500000  1.500000  2.500000  3.500000 
========= ERROR SUMMARY: 0 errors 
$ 

私はいくつかの変更を行ったことに注意してください。特に、サンプルポイントとサンプル数を調整して、5つのデータポイント（0,1,2,3,4）のそれぞれの途中で線形補間されたサンプルポイントを選択し、合計出力あなたの5つのデータポイント間の中点を表す4つの量（0.5,1.5,2.5,3.5）。

正規化座標索引付けの詳細については、the programming guideに記載されています（境界モードなど）。さらに、テクスチャの適切な使用を実証する様々なものがある（CUDA sample codes）。