CUDA C Programming Guide 在线教程学习笔记 Part 3

Posted 2020-10-14 爨爨爨好

? 表面内存使用

● 创建 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("\n\n");
39     for (i = 0; i < width*height; i++)
40     {
41         printf("%6.1f ", h_data[i]);
42         if ((i + 1) % width == 0)
43             printf("\n");
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("\n\n");
77     for (i = 0; i < width*height; i++)
78     {
79         printf("%6.1f ", h_data[i]);
80         if ((i + 1) % width == 0)
81             printf("\n");
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("\n\n");
42     for (i = 0; i < width*height; i++)
43     {
44         printf("%6.1f ", h_data[i]);
45         if ((i + 1) % width == 0)
46             printf("\n");
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("\n\n");
71     for (i = 0; i < width*height; i++)
72     {
73         printf("%6.1f ", h_data[i]);
74         if ((i + 1) % width == 0)
75             printf("\n");
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 数组只能用纹理访问函数来访问，或表面函数来进行读写。

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