正文

cudacprogrammingguide在线教程学习笔记part3

爨爨爨好  爨爨爨好  2022-09-30  279

关键词：

? 表面内存使用

● 创建 cuda 数组时使用标志 cudaArraySurfaceLoadStore 来创建表面内存，可以用表面对象（surface object）或表面引用（surface reference）来对其进行读写。

● 使用 Surface Object API

■ 涉及的结构定义、接口函数。

1 // vector_types.h
2 struct __device_builtin__ __align__(4) uchar4
3 {
4     unsigned char x, y, z, w;
5 };
6 
7 // surface_types.h
8 typedef __device_builtin__ unsigned long long cudaSurfaceObject_t;

■ 完整的测试代码，使用表面内存进行简单的读写。

 1 #include <stdio.h>
 2 #include <stdlib.h>
 3 #include <malloc.h>
 4 #include <cuda_runtime_api.h>
 5 #include "device_launch_parameters.h"
 6 
 7 #define CEIL(x,y) (((x) + (y) - 1) / (y) + 1)
 8 
 9 __global__ void myKernel(cudaSurfaceObject_t inputSurfObj, cudaSurfaceObject_t outputSurfObj, int width, int height)
10 {
11     unsigned int idx = blockIdx.x * blockDim.x + threadIdx.x;
12     unsigned int idy = blockIdx.y * blockDim.y + threadIdx.y;
13     if (idx < width && idy < height)
14     {
15         uchar4 data;
16         // 简单的表面内存读写，使用了字节地址，而不是简单的线程编号
17         surf2Dread(&data, inputSurfObj, sizeof(float) * idx, idy);
18         surf2Dwrite(data, outputSurfObj, sizeof(float) * idx, idy);
19     }
20     cudaBindSurfaceToArray();
21 }
22 
23 int main()
24 {
25     // 基本数据
26     int i;
27     float *h_data, *d_data; 
28     int width = 32;
29     int height = 32;
30 
31     int size = sizeof(float)*width*height;
32     h_data = (float *)malloc(size);
33     cudaMalloc((void **)&d_data, size);
34 
35     for (i = 0; i < width*height; i++)
36         h_data[i] = (float)i;
37 
38     printf("

");
39     for (i = 0; i < width*height; i++)
40     {
41         printf("%6.1f ", h_data[i]);
42         if ((i + 1) % width == 0)
43             printf("
");
44     }
45     
46     // 申请 cuda 数组
47     cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc(8, 8, 8, 8, cudaChannelFormatKindUnsigned);
48     cudaArray* cuInputArray;
49     cudaMallocArray(&cuInputArray, &channelDesc, width, height, cudaArraySurfaceLoadStore); 
50     cudaArray* cuOutputArray;
51     cudaMallocArray(&cuOutputArray, &channelDesc, width, height, cudaArraySurfaceLoadStore);
52     cudaMemcpyToArray(cuInputArray, 0, 0, h_data, size,cudaMemcpyHostToDevice);
53 
54     // 指定表面内存
55     struct cudaResourceDesc resDesc;
56     memset(&resDesc, 0, sizeof(resDesc));
57     resDesc.resType = cudaResourceTypeArray;
58 
59     // 创建表面对象
60     resDesc.res.array.array = cuInputArray;
61     cudaSurfaceObject_t inputSurfObj = 0;
62     cudaCreateSurfaceObject(&inputSurfObj, &resDesc);
63     resDesc.res.array.array = cuOutputArray;
64     cudaSurfaceObject_t outputSurfObj = 0;
65     cudaCreateSurfaceObject(&outputSurfObj, &resDesc);
66 
67     // 运行核函数
68     dim3 dimBlock(16, 16);
69     dim3 dimGrid(CEIL(width, dimBlock.x), CEIL(height, dimBlock.y));
70     myKernel << <dimGrid, dimBlock >> > (inputSurfObj, outputSurfObj, width, height);
71 
72     // 结果回收和检查结果
73     memset(h_data,0,size);// 刷掉原来的 h_data，再用 cuOutputArray 的数据写入     
74     cudaMemcpyFromArray(h_data, cuOutputArray, 0, 0, size, cudaMemcpyDeviceToHost);
75 
76     printf("

");
77     for (i = 0; i < width*height; i++)
78     {
79         printf("%6.1f ", h_data[i]);
80         if ((i + 1) % width == 0)
81             printf("
");
82     }
83 
84     // 回收工作
85     cudaDestroySurfaceObject(inputSurfObj);
86     cudaDestroySurfaceObject(outputSurfObj);
87     cudaFreeArray(cuInputArray);
88     cudaFreeArray(cuOutputArray);
89 
90     getchar();
91     return 0;
92 }

● 使用 Surface Reference API。

■ 表面引用的一些只读属性需要在声明的时候指定，以便编译时提前确定，只能在全局作用域内静态指定，不能作为参数传递给函数。使用 surface 指定纹理引用属性，Datatype 为数据类型，Type 为纹理引用类型，有 7 种，默认 cudaSurfaceType1D。

 1 surface<void, Type> surfRef;
 2 
 3 // cuda_texture_types.h
 4 template<class T, int dim = 1>
 5 struct __device_builtin_surface_type__ surface : public surfaceReference
 6 {
 7 #if !defined(__CUDACC_RTC__)
 8     __host__ surface(void)
 9     {
10         channelDesc = cudaCreateChannelDesc<T>();
11     }
12 
13     __host__ surface(struct cudaChannelFormatDesc desc)
14     {
15         channelDesc = desc;
16     }
17 #endif /* !__CUDACC_RTC__ */  
18 };
19 
20 //surface_types.h
21 #define cudaSurfaceType1D              0x01
22 #define cudaSurfaceType2D              0x02
23 #define cudaSurfaceType3D              0x03
24 #define cudaSurfaceTypeCubemap         0x0C
25 #define cudaSurfaceType1DLayered       0xF1
26 #define cudaSurfaceType2DLayered       0xF2
27 #define cudaSurfaceTypeCubemapLayered  0xFC
28 
29 // 访问边界模式
30 enum __device_builtin__ cudaSurfaceBoundaryMode
31 {
32     cudaBoundaryModeZero = 0,    // 0 边界模式
33     cudaBoundaryModeClamp = 1,   // 挤压模式
34     cudaBoundaryModeTrap = 2     // 陷阱模式
35 };
36 
37 // ？表面格式模式
38 enum __device_builtin__  cudaSurfaceFormatMode
39 {
40     cudaFormatModeForced = 0,   // 强制模式
41     cudaFormatModeAuto = 1      // 自动模式
42 };
43 
44 // 表面引用的通道描述
45 struct __device_builtin__ surfaceReference
46 {
47     struct cudaChannelFormatDesc channelDesc;
48 };
49 
50 // cuda_runtime_api.h
51 extern __host__ cudaError_t CUDARTAPI cudaBindSurfaceToArray(const struct surfaceReference *surfref, cudaArray_const_t array, const struct cudaChannelFormatDesc *desc);

■ 表面引用使用字节地址来定位访问（而不是像纹理那样使用 fetch 函数），如以上代码中 surf1Dread(surfRef, sizeof(float) * idx) 或是 surf1Dread(surfRef, sizeof(float) * idx) 。

■ 表面引用必须用函数 cudaBindSurfaceToArray() 绑定到 cuda 数组上才能使用，要求表面引用的维度、数据类型与该数组匹配，否则操作时未定义的，使用完后不需要特殊函数来解除绑定。

■ 将表面引用绑定到 cuda 数组上的范例代码。

 1 // 准备工作
 2 surface<void, Type>surfRef;
 3 
 4 ...
 5 
 6 int width, height;
 7 size_t pitch;
 8 float *d_data;
 9 cudaMallocPitch((void **)&d_data, &pitch, sizeof(float)*width, height);
10 
11 // 第一种方法，低层 API
12 surfaceReference* surfRefPtr;
13 cudaGetSurfaceReference(&surfRefPtr, "surfRef");
14 cudaChannelFormatDesc channelDesc;
15 cudaGetChannelDesc(&channelDesc, cuArray);
16 cudaBindSurfaceToArray(surfRef, cuArray, &channelDesc);
17 
18 // 第二种方法，高层 API
19 cudaBindSurfaceToArray(surfRef, cuArray);

■ 完整的应用样例代码。与前面表面对象代码的功能相同。

 1 #include <stdio.h>
 2 #include <stdlib.h>
 3 #include <malloc.h>
 4 #include <cuda_runtime_api.h>
 5 #include "device_launch_parameters.h"
 6 
 7 #define CEIL(x,y) (((x) + (y) - 1) / (y) + 1)
 8 
 9 // 声明表面引用
10 surface<void, 2> inputSurfRef;
11 surface<void, 2> outputSurfRef;
12 
13 __global__ void myKernel(int width, int height)
14 {
15     unsigned int idx = blockIdx.x * blockDim.x + threadIdx.x;
16     unsigned int idy = blockIdx.y * blockDim.y + threadIdx.y;
17     if (idx < width && idy < height)
18     {
19         uchar4 data;
20         // 简单的表面内存读写，使用了字节地址，而不是简单的线程编号
21         surf2Dread(&data, inputSurfRef, sizeof(float) * idx, idy);
22         surf2Dwrite(data, outputSurfRef, sizeof(float) * idx, idy);
23     }
24 }
25 
26 int main()
27 {
28     // 基本数据
29     int i;
30     float *h_data, *d_data; 
31     int width = 32;
32     int height = 32;
33 
34     int size = sizeof(float)*width*height;
35     h_data = (float *)malloc(size);
36     cudaMalloc((void **)&d_data, size);
37 
38     for (i = 0; i < width*height; i++)
39         h_data[i] = (float)i;
40 
41     printf("

");
42     for (i = 0; i < width*height; i++)
43     {
44         printf("%6.1f ", h_data[i]);
45         if ((i + 1) % width == 0)
46             printf("
");
47     }
48     
49     // 申请 cuda 数组
50     cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc(8, 8, 8, 8, cudaChannelFormatKindUnsigned);
51     cudaArray* cuInputArray;
52     cudaMallocArray(&cuInputArray, &channelDesc, width, height, cudaArraySurfaceLoadStore); 
53     cudaArray* cuOutputArray;
54     cudaMallocArray(&cuOutputArray, &channelDesc, width, height, cudaArraySurfaceLoadStore);
55     cudaMemcpyToArray(cuInputArray, 0, 0, h_data, size,cudaMemcpyHostToDevice);
56 
57     // 绑定表面引用，注意与表面对象的使用不一样
58     cudaBindSurfaceToArray(inputSurfRef, cuInputArray);
59     cudaBindSurfaceToArray(outputSurfRef, cuOutputArray);
60 
61     // 运行核函数
62     dim3 dimBlock(16, 16);
63     dim3 dimGrid(CEIL(width, dimBlock.x), CEIL(height, dimBlock.y));
64     myKernel << <dimGrid, dimBlock >> > (width, height);
65 
66     // 结果回收和检查结果
67     memset(h_data,0,size);// 刷掉原来的 h_data，再用 cuOutputArray 的数据写入     
68     cudaMemcpyFromArray(h_data, cuOutputArray, 0, 0, size, cudaMemcpyDeviceToHost);
69 
70     printf("

");
71     for (i = 0; i < width*height; i++)
72     {
73         printf("%6.1f ", h_data[i]);
74         if ((i + 1) % width == 0)
75             printf("
");
76     }
77 
78     // 回收工作
79     cudaFreeArray(cuInputArray);
80     cudaFreeArray(cuOutputArray);
81 
82     getchar();
83     return 0;
84 }

? 立方体表面 Cubemap Surface 。（想象成一个正方体的外表面）

● 一种特殊的二维分层表面。函数 surfCubemapread() 和函数 surfCubemapwrite() 来对其进行读写，使用一个整数下标和两个浮点数有序组来定义层号和表面坐标。

? 分层立方体表面 Cubemap Layered Surfaces 。（想象成一个多层的正方体的各外表面）

● 一种特殊的二维分层表面。函数 surfCubemapread() 和函数 surfCubemapwrite() 来对齐进行读写。使用一个整数下标和两个浮点数有序组来定义层号和表面坐标。

● 分层立方体贴图纹理只能使用函数 cudaMAlloc3DArray() 加上 cudaArrayLayered 和 cudaArrayCubemap 标志来声明，使用函数 texCubemapLayered() 来进行访问滤波只在同一层内部进行，不会跨层执行。

? cuda 数组。

● cuda 优化的数组类型，可以有一维或二维或三维，每个元素可以有 1 个或 2 个或 4 个分量，各分量可以是 1 B 或 2 B 或 4 B 尺寸的有符号或无符号整数，或 2 B 或 4 B 尺寸的浮点数。cuda 数组只能用纹理访问函数来访问，或表面函数来进行读写。

● 纹理内存和表面内存都是可缓存的，且不能保证缓存和内存的一致性。同一个核函数中，用纹理访问或表面访问来读取“已经全局写入或表面写入的内存”是未定义的。

cudacprogrammingguide在线教程学习笔记part2

? 纹理内存使用● 纹理内存使用有两套API，称为ObjectAPI和ReferenceAPI。纹理对象（textureobject）在运行时被ObjectAPI创建，同时指定了纹理单元。纹理引用（TeztureReference）在编译时被ReferenceAPI创建，但是在运行时才指定纹理单... 查看详情

cudacprogrammingguide在线教程学习笔记part5

附录A，CUDA计算设备附录B，C语言扩展?函数的标识符●__device__，__global__和__host__　　●宏__CUDA_ARCH__可用于区分代码的运行位置.1__host____device__voidfun()2{3#if__CUDA_ARCH__>=6004//代码运行于计算能力6.x设备5#elif__CUDA_ARCH__>=5006//代码运... 查看详情

CUDA 加法与移位指令性能

...我写了两个小程序来比较加法和移位指令的吞吐量。根据CUDACProgrammingGuide，移位指令的吞吐量是加法指令的一半。但是，当我在TeslaM2070上测量跟随两个程序的时间时，时间是完查看详情

CUDA 真的没有类似 calloc() 的 API 调用吗？

...间】：2014-01-2012:10:51【问题描述】：从CUDA5.5APIReference和CUDACProgrammingGuide看来，似乎没有cudaCalloc()，在GPU上相当于标准C库的calloc()。真的没有用于分配初始化为全零的缓冲区的API功能吗？有查看详情

CUDA 中的矩阵向量乘法：基准测试和性能

...表述了问题以更具体并更新了代码）...我使用共享内存在CUDACProgrammingGuide之后在CUDAC中实现了一个用于矩阵向量乘法的查看详情

软件测试资料

...享给你的朋友和同事。欢迎推荐您看到的好资源。w3school在线教程易百在线教程 QTP教程JMeter教程 Robotframework教程TestNG教程Ranorex教程SilkTest教程Selenium教程qualitycenter教程（QC）Python教程Perl教程SQL教程Shell教程Tcl教程Lua教... 查看详情

在线使用在线软件来完成任务

在线画画ProcessOn 　　百度脑图在线办公软件user.me在线写作ShowDoc 　　看云在线获取编程知识菜鸟教程　　MobDevGroup 　　在线设计APP原型图墨刀查看详情

学习反应钩子的最佳在线教程是啥？

】学习反应钩子的最佳在线教程是啥？【英文标题】：Whatisthebesttutorialonlinetolearnreacthooks?学习反应钩子的最佳在线教程是什么？【发布时间】：2020-02-2821:54:19【问题描述】：我刚刚完成了React课程，其中包括React的基础知识和基... 查看详情

jquery在线教程

http://www.runoob.com/jquery/jquery-slide.htmlhttp://www.w3school.com.cn/jquery/http://www.phpstudy.net/e/jquery/http://www.runoob.com/jqueryui/jqueryui-tutorial.html 查看详情

springmvc在线教程

Springhelloworld实例Spring松耦合实例SpringJavaConfig实例SpringJavaConfig@Import实例Spring依赖注入（DI）Spring使用Setter依赖注入Spring通过构造方法依赖注入SpringBean引用例子如何注入值到Springbean属性Springbean加载多个配置文件Spring内部bean实例Sp... 查看详情

推介一个学习java的系列教程-狗鱼it教程

...个学JAVA的零基础学习JAVA的网站，推介一个学习JAVA的系列教程-狗鱼IT教程下面是java的系教程：1?[java教程]Java教程2?[java教程]Java简介3?[java教程]Java开发环境配置4?[java教程]Java基础语法5?[java教程]Java对象和类6?[java教程]Java基本... 查看详情

java微服务架构视频教程

JAVA微服务架构视频教程教程目录：┣━JAVA微服务架构视频教程┃┣━Java教程：第1章微服务简介4┃┃┣━Java教程：001构建单体应用┃┃┣━Java教程：002微服务解决复杂问题┃┃┣━Java教程：003微服务的优点┃┃┣━Java教程：... 查看详情

jasperreports教程_编程入门自学教程_菜鸟教程-免费教程分享

教程简介JasperReports入门教程-使用包含从环境设置，报告设计，编译报告设计，填充报告，查看和打印报告，导出，参数，数据源开始的基础知识到高级知识的初学者教程，简单易学地设计和创建JasperReports，字段，表达式，变量... 查看详情

关于Oracle的java在线教程中使用HashMap存储字谜的例子

】关于Oracle的java在线教程中使用HashMap存储字谜的例子【英文标题】：RegardingtheexampleinOracle\'sjavaonlinetutorialthatusesHashMaptostoreanagram【发布时间】：2012-11-1406:11:29【问题描述】：我正在阅读Oracle在线java教程中使用HashMap存储字谜的示... 查看详情

在线考试系统视频教程和源码

...sp;夜鹰教程的忠实粉丝们，经过几周的努力，夜鹰教程网在线考试系统视频教程终于录制完毕了，本套教程重点讲解了在线考试系统的数据库设计和程序开发，主要实现的功能包括的试题的新增和维护，题库的管理，考试主体信... 查看详情

pycharm教程

Pycharm教程（1）——定制外观Pycharm教程（2）——代码风格Pycharm教程（3）——代码的调试、运行Pycharm教程（4）——有关Python解释器的相关配置Pycharm教程（5）——Python快捷键相关设置Pycharm教程（6）——将Pycharm作为Vim编辑器使... 查看详情

extjs入门教程

ExtJs入门教程一[学习方法]ExtJs入门教程二[HelloWorld]ExtJs入门教程三[窗体：Window组件]ExtJs入门教程四[表单：FormPanel]ExtJs入门教程五[文本框：TextField]ExtJs入门教程六[按钮：Button]ExtJs入门教程七[登陆窗体Demo：Login]ExtJs入门教程八[脚... 查看详情

php语言在线代码运行编译工具推荐

PHP语言在线运行编译,是一款可在线编程编辑器,在编辑器上输入PHP语言代码,点击运行,可在线编译运行PHP语言,PHP语言代码在线运行调试,PHP语言在线编译,可快速在线测试您的PHP语言代码,在线编译PHP语言代码发现是否存在错误,如果... 查看详情

正文