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. : 당신이 5 인쇄 할 경우

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

   값을 다시 복사해야합니다.

   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)); 
   

: 대신에이 같은 나는 현명한 결과를 얻는다. 다음은 완전히 작동하는 코드입니다.

$ 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이 있습니다.