并行编程实战——CUDA编程的纹理内存

原创于 2025-09-06 08:36:08 发布 · 780 阅读

18 ·

CC 4.0 BY-SA版权

文章标签：

#c++ #CUDA

C++ 同时被 3 个专栏收录

464 篇文章

订阅专栏

C++11

349 篇文章

订阅专栏

并行编程

48 篇文章

订阅专栏

部署运行你感兴趣的模型镜像

一、纹理内存

我们在学习C++中内存可以分成堆和栈，其实纹理内存和这个差不多，它是CUDA中的一种内存。在CUDA中，纹理即提取的纹理内存，纹理对象在运行时创建，并在创建纹理对象指定纹理。纹理可以是线性内存的任何区域或CUDA数组。
所谓纹理内存其实就是一种只读的、能缓存（硬件缓存）优化、自动插值并可以自动进行数据转换的内存，主要应用于具有空间局部性读取模式的设计。
其主要的特点包括：
1、只读的：数据不可修改，适合于数据的查找等
2、缓存的：符合计算机的局部性访问原理
3、多维支持的：支持1D/2D/3D纹理
4、自动边界处理的：自动处理越界的数据
纹理内存的数据存储在显存中，但通过独立的纹理硬件单元进行管理，其需要显式的将数据绑定到纹理对象或纹理引用。

二、使用方法和相关说明

纹理内存的应用一般有四个步骤：
1、声明纹理内存
2、通过API将纹理内存绑定到纹理引用
3、在 CUDA 内核中使用纹理引用读取纹理内存
4、从纹理引用中解绑纹理内存

在实际声明纹理对象后，需要对纹理内存的一些重要属性进行设置，主要有：
1、纹理维度：主要对纹理是 1D、2D 还是 3D 数组寻址。纹理中的元素也称为纹素。深度、宽度和高度也被设置以定义每个维度。需要注意的是：不同的GPU架构都定义了每个维度的最大尺寸。
2、纹理类型：基本整数或浮点纹素的大小。
3、纹理读取模式：即元素的读取方式。它们可以以NormalizedFloat或ModeElement格式读取。标准化浮点模式期望在[0.0 1.0]和[-1.0 1.0]范围内的索引，对于无符号整数和有符号整数类型。
4、纹理坐标归一化。默认情况下，纹理函数通过[0, N-1]范围内的浮点坐标进行访问（N为对应维度的纹理尺寸）。例如64x32大小的纹理在x/y维度上分别使用[0,63]和[0,31]坐标范围。启用归一化后，坐标范围转换为[0.0, 1.0-1/N]，此时前述64x32纹理在x/y维度均使用[0, 1-1/N]范围坐标。归一化坐标能自动适应不同纹理尺寸的需求。
5、纹理寻址模式：纹理的一个独特特性是它如何进行超范围寻址。可能不些不可思议，但在许多图像算法中很常见。例如，如果正在通过平均相邻像素来应用插值，那么边界像素的行为应该是什么？纹理为开发人员提供了这个选项，以便他们可以选择将超出范围视为夹紧、包裹或镜像。在调整大小的示例中，已将其设置为夹紧模式，这基本上意味着超出范围的访问被夹紧到边界。
6、纹理过滤模式：设置模式定义在获取纹理（根据输入纹理坐标）时如何计算返回值。支持两类的过滤模式：cudaFilterModePoint和cudaFilterModeLinear。设置为线性模式时，可以进行插值（1D 的简单线性，2D 的双线性和 3D 的三线性）。仅当返回类型为浮点类型时，线性模式才有效。另一方面，ModePoint不执行插值，而是返回最近坐标的纹素。

三、纹理内存分类

其常见的分类有：
1、16位浮点类型（16-Bit Floating-Point Textures）
CUDA数组支持的16位浮点或half格式与IEEE 754-2008 binary2格式相同。CUDA C++提供了相关的转换函数用来处理不支持匹配的相关数据类型。可以通过调用 cudaCreateChannelDescHalf*() 函数来创建16位浮点格式的通道描述。
2、分层纹理（Layered Textures）
一维或二维分层纹理（在 Direct3D 中也称为纹理数组，在 OpenGL 中也称为数组纹理）是由一系列层组成的纹理，这些层都是具有相同维度、大小和数据类型的常规纹理.
其实很好理解，就是一个图层的堆叠，每个图层是一个纹理。大家可以认为是多个多维数组组成的一个多维数组（绕口啊，可以理解为二维数组就是两个一维数组组成），最终就是一个一维数组。
CUDA中提供了一系列的函数如cudaArrayLayered、tex2DLayered等来操作相关的纹理，需要注意的是，纹理过滤仅在层内完成而不能跨层。
注意，其仅在计算能力 2.0 及更高版本的设备上受支持。
3、立方体纹理(Cubemap Textures)
如果能理解分层纹理，就能更好的理解立方体纹理，它就好像一个立方体有六个面，然后用二级的分层纹理描述一下。再使用人们熟知的x,y,z来处理相关的纹理获取即可。
计算能力 2.0 及更高版本的设备上受支持。
4、分层的立方体纹理内存(Cubemap Layered Textures)
这个在理解上面的3和4后就容易理解了，可以认为是他们两个的组合即立方体贴图分层纹理是一种分层纹理，其层是相同维度的立方体贴图。
最后，说明一下纹理收集。作为一种特殊的纹理提取，它只适用于二维纹理，提取的函数为tex2Dgather()。纹理收集仅支持使用 cudaArrayTextureGather 标志创建的 CUDA 数组，其宽度和高度小于表 15 中为纹理收集指定的最大值，该最大值小于常规纹理提取。
纹理收集仅在计算能力 2.0 及更高版本的设备上受支持。

四、例程

下面的代码来自于官方代码：

#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <texture_fetch_functions.h>
#include <stdio.h>
#include <stdlib.h>
#include <cmath>
// Simple transformation kernel
__global__ void transformKernel(float* output,
                                cudaTextureObject_t texObj,
                                int width, int height,
                                float theta)
{
    // Calculate normalized texture coordinates
    unsigned int x = blockIdx.x * blockDim.x + threadIdx.x;
    unsigned int y = blockIdx.y * blockDim.y + threadIdx.y;

    float u = x / (float)width;
    float v = y / (float)height;

    // Transform coordinates
    u -= 0.5f;
    v -= 0.5f;
    float tu = u * cosf(theta) - v * sinf(theta) + 0.5f;
    float tv = v * cosf(theta) + u * sinf(theta) + 0.5f;

    // Read from texture and write to global memory：在CUDA内核中从纹理引用中读取纹理内存
    output[y * width + x] = tex2D<float>(texObj, tu, tv);
}
// Host code
int main()
{
    const int height = 1024;
    const int width = 1024;
    float angle = 0.5;

    // Allocate and set some host data
    float *h_data = (float *)std::malloc(sizeof(float) * width * height);
    for (int i = 0; i < height * width; ++i)
        h_data[i] = i;

    // Allocate CUDA array in device memory：创建一个通道描述并在链接纹理内存时使用
    cudaChannelFormatDesc channelDesc =
        cudaCreateChannelDesc(32, 0, 0, 0, cudaChannelFormatKindFloat);
    cudaArray_t cuArray;
    cudaMallocArray(&cuArray, &channelDesc, width, height);

    // Set pitch of the source (the width in memory in bytes of the 2D array pointed
    // to by src, including padding), we dont have any padding
    const size_t spitch = width * sizeof(float);
    // Copy data located at address h_data in host memory to device memory
    cudaMemcpy2DToArray(cuArray, 0, 0, h_data, spitch, width * sizeof(float),
                        height, cudaMemcpyHostToDevice);

    // Specify texture：指定纹理对象的参数
    struct cudaResourceDesc resDesc;
    memset(&resDesc, 0, sizeof(resDesc));
    resDesc.resType = cudaResourceTypeArray;
    resDesc.res.array.array = cuArray;

    // Specify texture object parameters
    struct cudaTextureDesc texDesc;
    memset(&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(&texObj, &resDesc, &texDesc, NULL);

    // Allocate result of transformation in device memory
    float *output;
    cudaMalloc(&output, width * height * sizeof(float));

    // Invoke kernel
    dim3 threadsperBlock(16, 16);
    dim3 numBlocks((width + threadsperBlock.x - 1) / threadsperBlock.x,
                    (height + threadsperBlock.y - 1) / threadsperBlock.y);
    transformKernel<<<numBlocks, threadsperBlock>>>(output, texObj, width, height,
                                                    angle);
    // Copy data from device back to host
    cudaMemcpy(h_data, output, width * height * sizeof(float),
                cudaMemcpyDeviceToHost);

    // Destroy texture object
    cudaDestroyTextureObject(texObj);

    // Free device memory
    cudaFreeArray(cuArray);
    cudaFree(output);

    // Free host memory
    free(h_data);

    return 0;
}