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

? 纹理内存使用

● 纹理内存使用有两套 API，称为 Object API 和 Reference API 。纹理对象（texture object）在运行时被 Object API 创建，同时指定了纹理单元。纹理引用（Tezture Reference）在编译时被 Reference API 创建，但是在运行时才指定纹理单元，并将纹理引用绑定到纹理单元上面去。

● 不同的纹理引用可能绑定到相同或内存上有重叠的的纹理单元上，纹理单元可能是 CUDA 线性内存或CUDA array 的任意部分。

● 可以定义一维到三维的数组作为纹理内存。数组中的元素简称 texel （texture element）。

● 纹理内存的数据类型可以是 char, int, long, long long, float, double，即所有基本的占用1、2、4字节的数据类型。
● 纹理内存的访问模式有 cudaReadModeNormalizedFloat 和 cudaReadModeElementType 两种。前者读取 4 字节整数时会除以 0x8fff（有符号整数）或 0xffff（无符号整数），从而把值线性映射到 [-1.0, 1.0] 区间（有符号整数）或 [0, 1] 区间（无符号整数），读取 2 字节整数时也会发生类似变换，除以 0x8f 或 0xff 。后者则不会发生这种转换。

● 纹理数组使用浮点坐标进行引用。长为 N 的一维数组，默认纹理坐标中，下标范围是 [0.0, N-1] 之间的浮点数；正规化纹理坐标中，下标范围是 [0.0, 1-1/N] 之间的浮点数。二维或三维数组每一维上的坐标也遵循这个原则。

● 寻址模式，可以对数组范围以外的坐标进行访问（越界访问），不同的寻址模式定义了这种操作的效果。默认寻址模式  cudaAddressModeClamp 下，越界访问取各维上的边界值；边界模式 cudaAddressModeBorder 下，越界访问会返回 0 。使用正规化坐标时，可以选用束模式和镜像模式，束模式 cudaAddressModeWrap（想象成左右边界相连）下，越界坐标做变换 x‘ = x - floor(x) （教程上少了 减号）；镜像模式  cudaAddressModeMirror（想象成 左 ~ 右 → 右 ~ 左 → 左 ~ 右）下，越界坐标做变换 x‘ = x（floor(x) 为偶数）或 x‘ = 1 - x（floor(x) 为奇数）。

● 滤波模式，决定了如何把整数坐标的数组数据值转化为浮点坐标的引用值。最临近插值 cudaFilterModePoint 使用最接近访问坐标的整数坐标点数据，可以返回整数值（若纹理数组本身是整数型）；线性插值  cudaFilterModeLinear 使用每维度上最接近访问坐标的两个整数坐标点数据进行插值，可以单线性（一维，2 点）、双线性（二维，4 点）和三线性（三维，8 点），只能返回浮点数值。

● 使用 Texture Object API 。

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

  1 // texture_types.h
  2 struct __device_builtin__ cudaTextureDesc
  3 {
  4     enum cudaTextureAddressMode addressMode[3];     // 寻址模式，cudaResourceDesc::resType == cudaResourceTypeLinear 时无效
  5     enum cudaTextureFilterMode  filterMode;         // 滤波模式，cudaResourceDesc::resType == cudaResourceTypeLinear 时无效
  6     enum cudaTextureReadMode    readMode;           // 访问模式
  7     int                         sRGB;               // ？读取时将sRGB范围正规化
  8     float                       borderColor[4];     // ？文理边界颜色
  9     int                         normalizedCoords;   // 是否使用正规化坐标
 10     unsigned int                maxAnisotropy;      // ？
 11     enum cudaTextureFilterMode  mipmapFilterMode;   // ？
 12     float                       mipmapLevelBias;    // ？
 13     float                       minMipmapLevelClamp;// ？
 14     float                       maxMipmapLevelClamp;// ？
 15 };
 16
 17 enum __device_builtin__ cudaTextureAddressMode
 18 {
 19     cudaAddressModeWrap = 0,
 20     cudaAddressModeClamp = 1,
 21     cudaAddressModeMirror = 2,
 22     cudaAddressModeBorder = 3
 23 };
 24
 25 enum __device_builtin__ cudaTextureFilterMode
 26 {
 27     cudaFilterModePoint = 0,
 28     cudaFilterModeLinear = 1
 29 };
 30
 31 enum __device_builtin__ cudaTextureReadMode
 32 {
 33     cudaReadModeElementType = 0,
 34     cudaReadModeNormalizedFloat = 1
 35 };
 36
 37 typedef __device_builtin__ unsigned long long cudaTextureObject_t;
 38
 39 // driver_types.h
 40 enum __device_builtin__ cudaChannelFormatKind
 41 {
 42     cudaChannelFormatKindSigned = 0,    // 有符号整数模式
 43     cudaChannelFormatKindUnsigned = 1,  // 无符号整数模式
 44     cudaChannelFormatKindFloat = 2,     // 浮点模式
 45     cudaChannelFormatKindNone = 3       // 无通道模式
 46 };
 47
 48 struct __device_builtin__ cudaChannelFormatDesc
 49 {
 50     int                        x;   // 通道 0 数据位深度
 51     int                        y;   // 通道 1 数据位深度
 52     int                        z;   // 通道 2 数据位深度
 53     int                        w;   // ？
 54     enum cudaChannelFormatKind f;   // 通道模式
 55 };
 56
 57 typedef struct cudaArray *cudaArray_t;
 58 typedef struct cudaMipmappedArray *cudaMipmappedArray_t;
 59
 60 enum __device_builtin__ cudaResourceType
 61 {
 62     cudaResourceTypeArray = 0x00,           // 数组资源
 63     cudaResourceTypeMipmappedArray = 0x01,  // 映射数组资源
 64     cudaResourceTypeLinear = 0x02,          // 线性资源
 65     cudaResourceTypePitch2D = 0x03          // 对齐二维资源
 66 };
 67
 68 struct __device_builtin__ cudaResourceDesc
 69 {
 70     enum cudaResourceType resType;              // 资源类型
 71
 72     union res
 73     {
 74         struct array                            // cuda数组
 75         {
 76             cudaArray_t array;
 77         };
 78         struct mipmap                           // mipmap 数组
 79         {
 80             cudaMipmappedArray_t mipmap;
 81         };
 82         struct linear                           // 一维数组
 83         {
 84             void *devPtr;                       // 设备指针，符合 cudaDeviceProp::textureAlignment 的对齐要求
 85             struct cudaChannelFormatDesc desc;  // texel 的属性描述
 86             size_t sizeInBytes;                 // 数组字节数
 87         };
 88         struct pitch2D                          // 二位数组
 89         {
 90             void *devPtr;                       // 设备指针，符合 cudaDeviceProp::textureAlignment 的对齐要求
 91             struct cudaChannelFormatDesc desc;  // texel 的属性描述
 92             size_t width;                       // 数组列数
 93             size_t height;                      // 数组行数
 94             size_t pitchInBytes;                // 数组行字节数
 95         };
 96     };
 97 };
 98
 99 // cuda_runtime_api.h
100
101 extern __host__ struct cudaChannelFormatDesc CUDARTAPI cudaCreateChannelDesc(int x, int y, int z, int w, enum cudaChannelFormatKind f);
102
103 extern __host__ cudaError_t CUDARTAPI cudaMallocArray(cudaArray_t *array, const struct cudaChannelFormatDesc *desc, size_t width, size_t height __dv(0), unsigned int flags __dv(0));
104
105 extern __host__ cudaError_t CUDARTAPI cudaMemcpyToArray(cudaArray_t dst, size_t wOffset, size_t hOffset, const void *src, size_t count, enum cudaMemcpyKind kind);
106
107 extern __host__ cudaError_t CUDARTAPI cudaCreateTextureObject(cudaTextureObject_t *pTexObject, const struct cudaResourceDesc *pResDesc, const struct cudaTextureDesc *pTexDesc, const struct cudaResourceViewDesc *pResViewDesc);
108
109 extern __host__ cudaError_t CUDARTAPI cudaDestroyTextureObject(cudaTextureObject_t texObject);

■ 完整的应用样例代码。初始化一个 32×32 的矩阵，利用纹理对其进行平移和旋转，输出调整之后的矩阵。

  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 DEGRE_TO_RADIAN(x) ((x) * 3.1416f / 180)
  8 #define CEIL(x,y) (((x) + (y) - 1) / (y) + 1)
  9
 10 // 简单的线性变换
 11 __global__ void transformKernel(float* output, cudaTextureObject_t texObj, int width, int height, float theta)
 12 {
 13     // 计算正规化纹理坐标
 14     unsigned int idx = blockIdx.x * blockDim.x + threadIdx.x;
 15     unsigned int idy = blockIdx.y * blockDim.y + threadIdx.y;
 16
 17     // 正规化和平移
 18     float u = idx / (float)width - 0.5f;
 19     float v = idy / (float)height - 0.5f;
 20
 21     // 旋转
 22     float tu = u * __cosf(theta) - v * __sinf(theta) + 0.5f;
 23     float tv = v * __cosf(theta) + u * __sinf(theta) + 0.5f;
 24
 25     //printf("\n(%2d,%2d,%2d,%2d)->(%f,%f,%f)",
 26     //    blockIdx.x, blockIdx.y, threadIdx.x, threadIdx.y, tu, tv,tex2D<float>(texObj, tu, tv));
 27
 28     // 纹理内存写入全局内存
 29     output[idy * width + idx] = tex2D<float>(texObj, tu, tv);
 30 }
 31
 32 int main()
 33 {
 34     // 基本数据
 35     int i;
 36     float *h_data, *d_data;
 37     int width = 32;
 38     int height = 32;
 39     float angle = DEGRE_TO_RADIAN(30);
 40
 41     int size = sizeof(float)*width*height;
 42     h_data = (float *)malloc(size);
 43     cudaMalloc((void **)&d_data, size);
 44
 45     for (i = 0; i < width*height; i++)
 46         h_data[i] = (float)i;
 47
 48     printf("\n\n");
 49     for (i = 0; i < width*height; i++)
 50     {
 51         printf("%6.1f ", h_data[i]);
 52         if ((i + 1) % width == 0)
 53             printf("\n");
 54     }
 55
 56     // 申请 cuda 数组并拷贝数据
 57     cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc(32, 0, 0, 0,cudaChannelFormatKindFloat);
 58     cudaArray* cuArray;
 59     cudaMallocArray(&cuArray, &channelDesc, width, height);
 60     cudaMemcpyToArray(cuArray, 0, 0, h_data, size, cudaMemcpyHostToDevice);
 61
 62     // 指定纹理资源
 63     struct cudaResourceDesc resDesc;
 64     memset(&resDesc, 0, sizeof(resDesc));
 65     resDesc.resType = cudaResourceTypeArray;
 66     resDesc.res.array.array = cuArray;
 67
 68     // 指定纹理对象参数
 69     struct cudaTextureDesc texDesc;
 70     memset(&texDesc, 0, sizeof(texDesc));
 71     texDesc.addressMode[0] = cudaAddressModeWrap;
 72     texDesc.addressMode[1] = cudaAddressModeWrap;
 73     texDesc.filterMode = cudaFilterModeLinear;
 74     texDesc.readMode = cudaReadModeElementType;
 75     texDesc.normalizedCoords = 1;
 76
 77     // 创建文理对象
 78     cudaTextureObject_t texObj = 0;
 79     cudaCreateTextureObject(&texObj, &resDesc, &texDesc, NULL);
 80
 81     // 运行核函数
 82     dim3 dimBlock(16, 16);
 83     dim3 dimGrid(CEIL(width, dimBlock.x), CEIL(height, dimBlock.y));
 84     transformKernel << <dimGrid, dimBlock >> > (d_data, texObj, width, height, angle);
 85     cudaDeviceSynchronize();
 86
 87     // 结果回收和检查结果
 88     cudaMemcpy(h_data, d_data, size, cudaMemcpyDeviceToHost);
 89
 90     printf("\n\n");
 91     for (i = 0; i < width*height; i++)
 92     {
 93         printf("%6.1f ", h_data[i]);
 94         if ((i + 1) % width == 0)
 95             printf("\n");
 96     }
 97
 98     // 回收工作
 99     cudaDestroyTextureObject(texObj);
100     cudaFreeArray(cuArray);
101     cudaFree(d_data);
102
103     getchar();
104     return 0;
105 }

● 使用 Texture Reference API。

■ 纹理引用的一些只读属性需要在声明的时候指定，以便编译时提前确定，只能在全局作用域内静态指定，不能作为参数传递给函数。使用 texture 指定纹理引用属性，Datatype 为 texel 的数据类型，Type 为纹理引用类型，有 7 种，默认 cudaTextureType1D，ReadMode 为访问类型，默认 cudaReadModeElementType 。其他属性可以在主机运行时动态的修改。

 1 texture<DataType, Type, ReadMode> texRef;
 2
 3 // cuda_texture_types.h
 4 template<class T, int texType = cudaTextureType1D, enum cudaTextureReadMode mode = cudaReadModeElementType>
 5 struct __device_builtin_texture_type__ texture : public textureReference
 6 {
 7 #if !defined(__CUDACC_RTC__)
 8     __host__ texture(int norm = 0, enum cudaTextureFilterMode  fMode = cudaFilterModePoint, enum cudaTextureAddressMode aMode = cudaAddressModeClamp)
 9     {
10         normalized = norm;
11         filterMode = fMode;
12         addressMode[0] = aMode;
13         addressMode[1] = aMode;
14         addressMode[2] = aMode;
15         channelDesc = cudaCreateChannelDesc<T>();
16         sRGB = 0;
17     }
18     __host__ texture(int norm, enum cudaTextureFilterMode   fMode, enum cudaTextureAddressMode  aMode, struct cudaChannelFormatDesc desc)
19     {
20         normalized = norm;
21         filterMode = fMode;
22         addressMode[0] = aMode;
23         addressMode[1] = aMode;
24         addressMode[2] = aMode;
25         channelDesc = desc;
26         sRGB = 0;
27     }
28 #endif
29 };
30
31 //texture_types.h
32 #define cudaTextureType1D              0x01
33 #define cudaTextureType2D              0x02
34 #define cudaTextureType3D              0x03
35 #define cudaTextureTypeCubemap         0x0C
36 #define cudaTextureType1DLayered       0xF1
37 #define cudaTextureType2DLayered       0xF2
38 #define cudaTextureTypeCubemapLayered  0xFC

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

 1 // texture_types.h
 2 struct __device_builtin__ textureReference
 3 {
 4     int                          normalized;            // 是否使用正规化坐标
 5     enum cudaTextureFilterMode   filterMode;            // 滤波模式
 6     enum cudaTextureAddressMode  addressMode[3];        // 寻址模式
 7     struct cudaChannelFormatDesc channelDesc;           // texel 的格式，其元素数据类型与声明 texture 时的 Datatype 一致
 8     int                          sRGB;                  // ？读取时将sRGB范围正规化
 9     unsigned int                 maxAnisotropy;         // ？
10     enum cudaTextureFilterMode   mipmapFilterMode;      // ？
11     float                        mipmapLevelBias;       // ？
12     float                        minMipmapLevelClamp;   // ？
13     float                        maxMipmapLevelClamp;   // ？
14     int                          __cudaReserved[15];    // ？
15 };
16
17 // cuda_runtime_api.h
18 extern __host__ cudaError_t CUDARTAPI cudaBindTexture(size_t *offset, const struct textureReference *texref, const void *devPtr, const struct cudaChannelFormatDesc *desc, size_t size __dv(UINT_MAX));
19
20 extern __host__ cudaError_t CUDARTAPI cudaBindTexture2D(size_t *offset, const struct textureReference *texref, const void *devPtr, const struct cudaChannelFormatDesc *desc, size_t width, size_t height, size_t pitch);
21
22 extern __host__ cudaError_t CUDARTAPI cudaBindTextureToArray(const struct textureReference *texref, cudaArray_const_t array, const struct cudaChannelFormatDesc *desc);
23
24 extern __host__ cudaError_t CUDARTAPI cudaUnbindTexture(const struct textureReference *texref);
25
26 extern __host__ cudaError_t CUDARTAPI cudaGetTextureReference(const struct textureReference **texref, const void *symbol);

■ 将 2D 纹理引用绑定到 2D 数组上的范例代码

 1 // 准备工作
 2 texture<float, cudaTextureType2D, cudaReadModeElementType> texRef;
 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 textureReference* texRefPtr;
13 cudaGetTextureReference(&texRefPtr, &texRef);
14 cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc<float>();
15 size_t offset;
16 cudaBindTexture2D(&offset, texRefPtr, d_data, &channelDesc, width, height, pitch);
17
18 // 第二种方法，高层 API
19 cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc<float>();
20 size_t offset;
21 cudaBindTexture2D(&offset, texRef, d_data, channelDesc, width, height, pitch);

■ 将 2D 纹理引用绑定到 cuda 数组上的范例代码

 1 // 准备工作
 2 texture<float, cudaTextureType2D, cudaReadModeElementType> texRef;
 3
 4 //...
 5
 6 cudaArray* cuArray;
 7 cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc(32, 0, 0, 0, cudaChannelFormatKindFloat);
 8 cudaMallocArray(&cuArray, &channelDesc, width, height);
 9
10 // 第一种方法，低层 API
11 textureReference* texRefPtr;
12 cudaGetTextureReference(&texRefPtr, &texRef);
13 memset(&channelDesc, 0, sizeof(cudaChannelFormatDesc));
14 cudaChannelFormatDesc channelDesc;
15 cudaGetChannelDesc(&channelDesc, cuArray);
16 cudaBindTextureToArray(texRef, cuArray, &channelDesc);
17
18 // 第二种方法，高层 API
19 cudaBindTextureToArray(texRef, 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 DEGRE_TO_RADIAN(x) ((x) * 3.1416f / 180)
 8 #define CEIL(x,y) (((x) + (y) - 1) / (y) + 1)
 9
10 // 声明纹理引用
11 texture<float, cudaTextureType2D, cudaReadModeElementType> texRef;
12
13 // 简单的线性变换
14 __global__ void transformKernel(float* output, int width, int height, float theta)
15 {
16     // 计算正规化纹理坐标
17     unsigned int idx = blockIdx.x * blockDim.x + threadIdx.x;
18     unsigned int idy = blockIdx.y * blockDim.y + threadIdx.y;
19
20     // 正规化和平移
21     float u = idx / (float)width;
22     float v = idy / (float)height;
23
24     // 旋转
25     float tu = u * __cosf(theta) - v * __sinf(theta) + 0.5f;
26     float tv = v * __cosf(theta) + u * __sinf(theta) + 0.5f;
27
28     //printf("\n(%2d,%2d,%2d,%2d)->(%f,%f,%f)",
29     //    blockIdx.x, blockIdx.y, threadIdx.x, threadIdx.y, tu, tv,tex2D<float>(texObj, tu, tv));
30
31     // 纹理内存写入全局内存
32     output[idy * width + idx] = tex2D(texRef, tu, tv);
33 }
34
35 int main()
36 {
37     // 基本数据
38     int i;
39     float *h_data, *d_data;
40     int width = 32;
41     int height = 32;
42     float angle = DEGRE_TO_RADIAN(30);
43
44     int size = sizeof(float)*width*height;
45     h_data = (float *)malloc(size);
46     cudaMalloc((void **)&d_data, size);
47
48     for (i = 0; i < width*height; i++)
49         h_data[i] = (float)i;
50
51     printf("\n\n");
52     for (i = 0; i < width*height; i++)
53     {
54         printf("%6.1f ", h_data[i]);
55         if ((i + 1) % width == 0)
56             printf("\n");
57     }
58
59     // 申请 cuda 数组并拷贝数据
60     cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc(32, 0, 0, 0, cudaChannelFormatKindFloat);
61     cudaArray* cuArray;
62     cudaMallocArray(&cuArray, &channelDesc, width, height);
63     cudaMemcpyToArray(cuArray, 0, 0, h_data, size, cudaMemcpyHostToDevice);
64
65     // 指定纹理引用参数，注意与纹理对象的使用不一样
66     texRef.addressMode[0] = cudaAddressModeWrap;
67     texRef.addressMode[1] = cudaAddressModeWrap;
68     texRef.filterMode = cudaFilterModeLinear;
69     texRef.normalized = 1;
70
71     // 绑定纹理引用
72     cudaBindTextureToArray(texRef, cuArray, channelDesc);
73
74     // 运行核函数
75     dim3 dimBlock(16, 16);
76     dim3 dimGrid(CEIL(width, dimBlock.x), CEIL(height, dimBlock.y));
77     transformKernel << <dimGrid, dimBlock >> > (d_data, width, height, angle);
78     cudaDeviceSynchronize();
79
80     // 结果回收和检查结果
81     cudaMemcpy(h_data, d_data, size, cudaMemcpyDeviceToHost);
82
83     printf("\n\n");
84     for (i = 0; i < width*height; i++)
85     {
86         printf("%6.1f ", h_data[i]);
87         if ((i + 1) % width == 0)
88             printf("\n");
89     }
90
91     // 回收工作
92     cudaFreeArray(cuArray);
93     cudaFree(d_data);
94
95     getchar();
96     return 0;
97 }

? 半精度浮点数。

■ CUDA 没有原生支持半精度浮点数据类型，可以把半精度数据存储在 short 数据类型中，在需要计算的时候用内建函数将其与浮点类型进行转换。

■ 这些函数只能在设备代码中使用，可以在 OpenEXR 库中找到其等价的函数。

■ 在纹理计算过程中，半精度浮点数默认转化为单精度浮点数。

 1 // cuda_fp16.h
 2 __CUDA_FP16_DECL__ float __half2float(const __half h)
 3 {
 4     float val;
 5     asm volatile("{  cvt.f32.f16 %0, %1;}\n" : "=f"(val) : "h"(h.x));
 6     return val;
 7 }
 8 __CUDA_FP16_DECL__ __half2 __float2half2_rn(const float f)
 9 {
10     __half2 val;
11     asm("{.reg .f16 low;\n"
12         "  cvt.rn.f16.f32 low, %1;\n"
13         "  mov.b32 %0, {low,low};}\n" : "=r"(val.x) : "f"(f));
14     return val;
15 }

? 分层纹理 Layered Texture

● 在 Direct3D 中叫 texture array，在 OpenGL 中叫 array texture 。

● 一组分层纹理是由若干相同维度、尺寸和数据类型的的纹理内存构成的，相当于多了一个整数下标的维度。支持一维分层纹理和二维分层纹理。

● 一维分层纹理使用一个整数和一个浮点数作为坐标进行访问；二维分层纹理使用一个整数和两个浮点数作为坐标进行访问。

● 分层纹理只能使用函数 cudaMalloc3DArray() 加上 cudaArrayLayered 标志来声明，使用函数 tex1DLayered() 和 tex2DLayered() 来进行访问。滤波只在同一层内部进行，不会跨层执行。

● 分层纹理在书《CUDA专家手册》中讲的稍微详细一点，看完再来填坑。

? 立方体贴图纹理 Cubemap Textures

● 一种特殊的二维分层纹理，共六个尺寸相同的、宽度等于高度的二维纹理构成，代表了正方体的六个面。

● 使用三个浮点数的有序组 (x, y, z) 来定义立方体贴图纹理的层号和表面坐标，按照以下表格分情况讨论。各表面坐标按照（(s / m + 1) / 2, (t / m + 1) / 2）计算。

● 立方体贴图纹理只能使用函数 cudaMalloc3DArray() 加上 cudaArrayCubemap 标志来声明，使用函数 texCubemap() 来访问。

? 分层立方体贴图纹理 Cubemap Layered Textures

● 一种分层纹理内存，由若干尺寸相同的立方体贴图纹理构成。使用一个整数下标和三个浮点数有序组来定义层号和面号、表面坐标。

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

? 纹理汇集。

● 使用函数 tex2Dgather() 来抽取二维纹理内存的特定内容，没看懂。

1 // texture_fetch_functions.h
2 template <typename T>
3 static __device__ typename __nv_tex2dgather_ret<T>::type tex2Dgather(texture<T, cudaTextureType2D, cudaReadModeElementType>, float, float, int = 0) {  }

? 压缩版的 texture_types.h。所有内容在本文中都有体现。

 1 #if !defined(__TEXTURE_TYPES_H__)
 2 #define __TEXTURE_TYPES_H__
 3
 4 #include "driver_types.h"
 5
 6 #define cudaTextureType1D              0x01
 7 #define cudaTextureType2D              0x02
 8 #define cudaTextureType3D              0x03
 9 #define cudaTextureTypeCubemap         0x0C
10 #define cudaTextureType1DLayered       0xF1
11 #define cudaTextureType2DLayered       0xF2
12 #define cudaTextureTypeCubemapLayered  0xFC
13
14 // CUDA texture address modes
15 enum __device_builtin__ cudaTextureAddressMode
16 {
17     cudaAddressModeWrap   = 0,    // Wrapping address mode
18     cudaAddressModeClamp  = 1,    // Clamp to edge address mode
19     cudaAddressModeMirror = 2,    // Mirror address mode
20     cudaAddressModeBorder = 3     // Border address mode
21 };
22
23 // CUDA texture filter modes
24 enum __device_builtin__ cudaTextureFilterMode
25 {
26     cudaFilterModePoint  = 0,     // Point filter mode
27     cudaFilterModeLinear = 1      // Linear filter mode
28 };
29
30 // CUDA texture read modes
31 enum __device_builtin__ cudaTextureReadMode
32 {
33     cudaReadModeElementType     = 0,  // Read texture as specified element type
34     cudaReadModeNormalizedFloat = 1   // Read texture as normalized float
35 };
36
37 // CUDA texture reference
38 struct __device_builtin__ textureReference
39 {
40     // Indicates whether texture reads are normalized or not
41     int                          normalized;
42     // Texture filter mode
43     enum cudaTextureFilterMode   filterMode;
44     // Texture address mode for up to 3 dimensions
45     enum cudaTextureAddressMode  addressMode[3];
46     // Channel descriptor for the texture reference
47     struct cudaChannelFormatDesc channelDesc;
48     // Perform sRGB->linear conversion during texture read
49     int                          sRGB;
50     // Limit to the anisotropy ratio
51     unsigned int                 maxAnisotropy;
52     // Mipmap filter mode
53     enum cudaTextureFilterMode   mipmapFilterMode;
54     // Offset applied to the supplied mipmap level
55     float                        mipmapLevelBias;
56     // Lower end of the mipmap level range to clamp access to
57     float                        minMipmapLevelClamp;
58     // Upper end of the mipmap level range to clamp access to
59     float                        maxMipmapLevelClamp;
60     int                          __cudaReserved[15];
61 };
62
63 // CUDA texture descriptor
64 struct __device_builtin__ cudaTextureDesc
65 {
66     // Texture address mode for up to 3 dimensions
67     enum cudaTextureAddressMode addressMode[3];
68     // Texture filter mode
69     enum cudaTextureFilterMode  filterMode;
70     // Texture read mode
71     enum cudaTextureReadMode    readMode;
72     // Perform sRGB->linear conversion during texture read
73     int                         sRGB;
74     // Texture Border Color
75     float                       borderColor[4];
76     // Indicates whether texture reads are normalized or not
77     int                         normalizedCoords;
78     // Limit to the anisotropy ratio
79     unsigned int                maxAnisotropy;
80     // Mipmap filter mode
81     enum cudaTextureFilterMode  mipmapFilterMode;
82     // Offset applied to the supplied mipmap level
83     float                       mipmapLevelBias;
84     // Lower end of the mipmap level range to clamp access to
85     float                       minMipmapLevelClamp;
86     // Upper end of the mipmap level range to clamp access to
87     float                       maxMipmapLevelClamp;
88 };
89
90 // An opaque value that represents a CUDA texture object
91 typedef __device_builtin__ unsigned long long cudaTextureObject_t;
92
93 #endif