二维分层纹理

▶ 源代码。用纹理方法把元素按原顺序从 CUDA3D 数组中取出来，求个相反数再加上层数放入全局内存，输出。

 #include <stdio.h>

 #include "cuda_runtime.h"

 #include "device_launch_parameters.h"

 #include <helper_functions.h>

 #include <helper_cuda.h>

 #define MIN_EPSILON_ERROR 5e-3f

 #define OUTPUT 5

 texture<float, cudaTextureType2DLayered> tex;

 __global__ void transformKernel(float *g_odata, int width, int height, int layer)

 {

     unsigned int x = blockIdx.x*blockDim.x + threadIdx.x;

     unsigned int y = blockIdx.y*blockDim.y + threadIdx.y;

     float u = (x + 0.5f) / (float)width;

     float v = (y + 0.5f) / (float)height;

     g_odata[layer*width*height + y*width + x] = - tex2DLayered(tex, u, v, layer) + layer;

 }

 int main(int argc, char **argv)

 {

     unsigned int width = , height = , num_layers = ;

     unsigned int size = width * height * num_layers * sizeof(float);

     float *h_data = (float *)malloc(size);

     float *h_data_ref = (float *)malloc(size);

     float *d_data = NULL;

     cudaMalloc((void **)&d_data, size);

     for (unsigned int layer = ; layer < num_layers; layer++)

     {

         for (int i = ; i < (int)(width * height); i++)

             h_data[layer*width*height + i] = (float)i;

     }

     for (unsigned int layer = ; layer < num_layers; layer++)

     {

         for (int i = ; i < (int)(width * height); i++)

             h_data_ref[layer*width*height + i] = - h_data[layer*width*height + i] + layer;

     }

     printf("\n\t\Input data\n\t");

     for (int i = ; i < num_layers; i++)

     {

         for (int j = ; j < OUTPUT; j++)

         {

             for(int k=;k<OUTPUT;k++)

                 printf("%2.1f ", h_data[i*width*height+j*width+k]);

             printf("\n\t");

         }

         printf("\n\t");

     }

     printf("\n\t\Ideal output data\n\t");

     for (int i = ; i < num_layers; i++)

     {

         for (int j = ; j < OUTPUT; j++)

         {

             for (int k = ; k<OUTPUT; k++)

                 printf("%2.1f ", h_data_ref[i*width*height + j*width + k]);

             printf("\n\t");

         }

         printf("\n\t");

     }

     // 设置 CUDA 3D 数组参数和数据拷贝

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

     cudaArray *cu_3darray;

     cudaMalloc3DArray(&cu_3darray, &channelDesc, make_cudaExtent(width, height, num_layers), cudaArrayLayered);

     cudaMemcpy3DParms myparms = {  };

     myparms.srcPos = make_cudaPos(, , );

     myparms.dstPos = make_cudaPos(, , );

     myparms.srcPtr = make_cudaPitchedPtr(h_data, width * sizeof(float), width, height);

     myparms.dstArray = cu_3darray;

     myparms.extent = make_cudaExtent(width, height, num_layers);

     myparms.kind = cudaMemcpyHostToDevice;

     cudaMemcpy3D(&myparms);

     // 设置纹理参数并绑定

     tex.addressMode[] = cudaAddressModeWrap;

     tex.addressMode[] = cudaAddressModeWrap;

     tex.filterMode = cudaFilterModeLinear;

     tex.normalized = true;

     cudaBindTextureToArray(tex, cu_3darray, channelDesc);

     dim3 dimBlock(, , );

     dim3 dimGrid(width / dimBlock.x, height / dimBlock.y, );

     printf("Covering 2D data of %d * %d * %d: Grid size is %d x %d, each block has 8 x 8 threads\n", width, height, num_layers, dimGrid.x, dimGrid.y);

     transformKernel << < dimGrid, dimBlock >> >(d_data, width, height, );// 预跑

     cudaDeviceSynchronize();

     StopWatchInterface *timer = NULL;

     sdkCreateTimer(&timer);

     sdkStartTimer(&timer);

     for (unsigned int layer = ; layer < num_layers; layer++)// 启用多个核，每个核完成一层

         transformKernel << < dimGrid, dimBlock,  >> >(d_data, width, height, layer);

     cudaDeviceSynchronize();

     sdkStopTimer(&timer);

     printf("\n\Time: %.3f msec, %.2f Mtexlookups/sec\n", sdkGetTimerValue(&timer), (width *height *num_layers / (sdkGetTimerValue(&timer) / 1000.0f) / 1e6));

     sdkDeleteTimer(&timer);

     // 返回计算结果并检验

     memset(h_data, , size);

     cudaMemcpy(h_data, d_data, size, cudaMemcpyDeviceToHost);

     if (checkCmdLineFlag(argc, (const char **)argv, "regression"))

         sdkWriteFile<float>("./data/regression.dat", h_data, width * width, 0.0f, false);

     else

         printf("Comparing kernel output to expected data return %d\n", compareData(h_data, h_data_ref, width * height * num_layers, MIN_EPSILON_ERROR, 0.0f));

     printf("\n\tActual output data\n\t");

     for (int i = ; i < num_layers; i++)

     {

         for (int j = ; j < OUTPUT; j++)

         {

             for (int k = ; k<OUTPUT; k++)

                 printf("%2.1f ", h_data[i*width*height + j*width + k]);

             printf("\n\t");

         }

         printf("\n\t");

     }

     free(h_data);

     free(h_data_ref);

     cudaFree(d_data);

     cudaFreeArray(cu_3darray);

     getchar();

     return ;

 }

▶ 输出结果

    Input data

    0.0 1.0 2.0 3.0 4.0

    512.0 513.0 514.0 515.0 516.0

    1024.0 1025.0 1026.0 1027.0 1028.0

    1536.0 1537.0 1538.0 1539.0 1540.0

    2048.0 2049.0 2050.0 2051.0 2052.0

    0.0 1.0 2.0 3.0 4.0

    512.0 513.0 514.0 515.0 516.0

    1024.0 1025.0 1026.0 1027.0 1028.0

    1536.0 1537.0 1538.0 1539.0 1540.0

    2048.0 2049.0 2050.0 2051.0 2052.0

    0.0 1.0 2.0 3.0 4.0

    512.0 513.0 514.0 515.0 516.0

    1024.0 1025.0 1026.0 1027.0 1028.0

    1536.0 1537.0 1538.0 1539.0 1540.0

    2048.0 2049.0 2050.0 2051.0 2052.0

    0.0 1.0 2.0 3.0 4.0

    512.0 513.0 514.0 515.0 516.0

    1024.0 1025.0 1026.0 1027.0 1028.0

    1536.0 1537.0 1538.0 1539.0 1540.0

    2048.0 2049.0 2050.0 2051.0 2052.0

    0.0 1.0 2.0 3.0 4.0

    512.0 513.0 514.0 515.0 516.0

    1024.0 1025.0 1026.0 1027.0 1028.0

    1536.0 1537.0 1538.0 1539.0 1540.0

    2048.0 2049.0 2050.0 2051.0 2052.0

    Ideal output data

    0.0 -1.0 -2.0 -3.0 -4.0

    -512.0 -513.0 -514.0 -515.0 -516.0

    -1024.0 -1025.0 -1026.0 -1027.0 -1028.0

    -1536.0 -1537.0 -1538.0 -1539.0 -1540.0

    -2048.0 -2049.0 -2050.0 -2051.0 -2052.0

    1.0 0.0 -1.0 -2.0 -3.0

    -511.0 -512.0 -513.0 -514.0 -515.0

    -1023.0 -1024.0 -1025.0 -1026.0 -1027.0

    -1535.0 -1536.0 -1537.0 -1538.0 -1539.0

    -2047.0 -2048.0 -2049.0 -2050.0 -2051.0

    2.0 1.0 0.0 -1.0 -2.0

    -510.0 -511.0 -512.0 -513.0 -514.0

    -1022.0 -1023.0 -1024.0 -1025.0 -1026.0

    -1534.0 -1535.0 -1536.0 -1537.0 -1538.0

    -2046.0 -2047.0 -2048.0 -2049.0 -2050.0

    3.0 2.0 1.0 0.0 -1.0

    -509.0 -510.0 -511.0 -512.0 -513.0

    -1021.0 -1022.0 -1023.0 -1024.0 -1025.0

    -1533.0 -1534.0 -1535.0 -1536.0 -1537.0

    -2045.0 -2046.0 -2047.0 -2048.0 -2049.0

    4.0 3.0 2.0 1.0 0.0

    -508.0 -509.0 -510.0 -511.0 -512.0

    -1020.0 -1021.0 -1022.0 -1023.0 -1024.0

    -1532.0 -1533.0 -1534.0 -1535.0 -1536.0

    -2044.0 -2045.0 -2046.0 -2047.0 -2048.0

    Covering 2D data of  *  * : Grid size is  x , each block has  x  threads

Time: 0.995 msec, 1317.00 Mtexlookups/sec

Comparing kernel output to expected data return 

    Actual output data

    0.0 -1.0 -2.0 -3.0 -4.0

    -512.0 -513.0 -514.0 -515.0 -516.0

    -1024.0 -1025.0 -1026.0 -1027.0 -1028.0

    -1536.0 -1537.0 -1538.0 -1539.0 -1540.0

    -2048.0 -2049.0 -2050.0 -2051.0 -2052.0

    1.0 0.0 -1.0 -2.0 -3.0

    -511.0 -512.0 -513.0 -514.0 -515.0

    -1023.0 -1024.0 -1025.0 -1026.0 -1027.0

    -1535.0 -1536.0 -1537.0 -1538.0 -1539.0

    -2047.0 -2048.0 -2049.0 -2050.0 -2051.0

    2.0 1.0 0.0 -1.0 -2.0

    -510.0 -511.0 -512.0 -513.0 -514.0

    -1022.0 -1023.0 -1024.0 -1025.0 -1026.0

    -1534.0 -1535.0 -1536.0 -1537.0 -1538.0

    -2046.0 -2047.0 -2048.0 -2049.0 -2050.0

    3.0 2.0 1.0 0.0 -1.0

    -509.0 -510.0 -511.0 -512.0 -513.0

    -1021.0 -1022.0 -1023.0 -1024.0 -1025.0

    -1533.0 -1534.0 -1535.0 -1536.0 -1537.0

    -2045.0 -2046.0 -2047.0 -2048.0 -2049.0

    4.0 3.0 2.0 1.0 0.0

    -508.0 -509.0 -510.0 -511.0 -512.0

    -1020.0 -1021.0 -1022.0 -1023.0 -1024.0

    -1532.0 -1533.0 -1534.0 -1535.0 -1536.0

    -2044.0 -2045.0 -2046.0 -2047.0 -2048.0

▶ 涨姿势

● 与前面立方体贴图纹理不同的地方：申请 CUDA3D 数组的时候使用标志 cudaArrayLayered 而不是 cudaArrayCubemap，并注意调整相关的维度参数。

最新文章

  1. python学习之路 第二天
  2. 【转】MySQL性能优化的最佳21条经验
  3. ASP.NET MVC 6 一些不晓得的写法
  4. [翻译]类型双关不好玩：C中使用指针重新解释是坏的
  5. 虚拟化平台cloudstack（4）——几个异常
  6. enmo_day_02
  7. mysql索引使用笔记
  8. [C#学习]在多线程中如何调用Winform[转]
  9. Eclipse图标含义
  10. Problem A: A + B
  11. [WPF疑难]如何禁用WPF窗口的系统菜单（SystemMenu)
  12. J2ee技术难点
  13. python3之正则表达式
  14. redis笔记总结之redis安装
  15. ORACLE ORA-01653: unable to extend table 的错误
  16. C/C++（static）
  17. tornado-版本迁移工具alembic
  18. sass：常用备忘
  19. 基于Window10搭建android开发环境
  20. mongodb学习(二)

热门文章

  1. Crazy Computer
  2. CTF-练习平台-WEB之 签到题
  3. j2ee使用struts实现用户的登陆及注冊
  4. jsfl读取xml,图片，并生成swf
  5. async（await）知识点
  6. cookie 知识点
  7. Spring 集成开发工具（STS）安装及配置
  8. POJ1325机器重启次数——二分图匈牙利算法模板
  9. CC2530中串口波特率改为9600时单个数据包来不及接收的解决方案
  10. 【appium】keyevent的keycode