Блог

Я написал следующий код, чтобы узнать, как использовать текстурную память для одномерного массива. но функция tex1D не извлекает значение из массива для соответствующего идентификатора потока.Пожалуйста, исправьте этот код и скажите мне, как эффективно и эффективно использовать текстурную память для одномерного массива.

 __global__ void sum(float *b,cudaTextureObject_t texObj)

    {
    b[threadIdx.x]=tex1D<float>(texObj,threadIdx.x);
    //printf("n%fn",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(amp;cuArray, amp;channelDesc, 5, 0);

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


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


      struct cudaTextureDesc texDesc;
        memset(amp;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(amp;texObj, amp;resDesc, amp;texDesc, NULL);


    cudaMalloc(amp;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("%ft",b[i]);
      cudaDestroyTextureObject(texObj); 
    cudaFreeArray(cuArray);
    cudaFree(d_b);

        return 0;

    }
  

Существует как минимум 2 проблемы:

1. В конце вы копируете только одну плавающую величину с устройства на хост:

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

   если вы хотите напечатать 5 значений, вам следует скопировать 5 значений обратно:

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

2. Вы выбрали нормализованные координаты:

    texDesc.normalizedCoords = 1;
     

   это означает, что вы должны передавать координату с плавающей запятой между 0 и 1 в качестве индекса, а не целочисленную координату от 0 до 4:

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

   вместо этого используйте что-то вроде этого:

     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%fn",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(amp;cuArray, amp;channelDesc, 5, 0);

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


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


      struct cudaTextureDesc texDesc;
        memset(amp;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(amp;texObj, amp;resDesc, amp;texDesc, NULL);


    cudaMalloc(amp;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("%ft",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) общий объем производства составляет 4 количества (0,5, 1,5, 2,5, 3,5).представляет средние точки между вашими 5 точками данных.

Если вы хотите узнать больше о нормализованном индексировании координат, это описано в руководстве по программированию, как и другие понятия, такие как режимы границ и тому подобное. Кроме того, существуют различные примеры кодов CUDA, которые демонстрируют правильное использование текстур.