CUDA纹理内存入门详解

CUDA 进阶技巧 纹理内存 入门指南
摘要由CSDN通过智能技术生成

Texture Memory

作为图形世界里的一个特征,纹理是在多边形上拉伸、旋转和粘贴的图像,以形成我们熟悉的3D图形。而在GPU计算中使用纹理内存是对CUDA程序员的进阶建议,纹理内存允许像数组一样随机访问并且通过cache提高带宽。

总体来说,CUDA提供了两套不同的API访问纹理内存:

  • texture reference API(支持所有设备)
  • texture object API(仅支持capacity ≥ 3. x \ge3.\text x 3.x

PStexture reference在CUDA 5.0后已被弃用,但考虑到很多历史代码仍然使用并且理解了texture reference后可以快速上手texture object,所以本文一并介绍

设备执行kernel时,通过Texture Functions(显然属于device functions)访问纹理内存,引出一个称呼texture fetch——调用texture function之一访问纹理内存的过程。

每一个texture fetch都规定了一个参数texture object(如果使用texture object API)或texture reference(如果使用texture reference API)

texture object / texture reference

  1. texture:希望访问的一段纹理内存,可以是设备中一段线性内存或是一个CUDA数组

    • texture object:运行时创建,texture在其创建时也随之指定
    • texture reference:编译时创建,通过运行时绑定texture reference到某段纹理内存完成texture的指定;不同的texture reference可以绑定到相同texture或是内存有重叠的texture
  2. dimensionality:定义texture的维度,将texture看作一维、二维或三维数组。

    • texels:texture elements的简称,即数组中的元素
    • texture width & texture height & texture depth:数组各维度的大小
  3. typetexel的类型,只能为以下几种类型

    • 基本整数类型(8位、16位、32位和64位整型)

    • 单精度浮点数类型

    • 由基本整型和单精度度浮点数生成的内置向量类型,且组件个数只能为1、2和4

      PS:内置向量类型实际为结构体,CUDA允许其有2、3和4个成员,并且通过xyzw分别访问第一,第二,第三和第四个成员;构造函数形如make_<type name>,例如int2 make_int2(int x, int y)生成一个int2类型的值(x,y)

      详细用法参见官方文档Built-in Vector Types

  4. read mode:读数据模式

    • cudaReadModeElementType:不发生转换
    • cudaReadModeNormalizedFloat:当type为16-bit或8-bit整数(short/ushortchar/uchar,i.e.不适用于32位整数)时,texture fetch返回的值均匀地映射为 [ 0.0 , 1.0 ] [0.0, 1.0] [0.0,1.0](对于无符号类型,i.e. ushort/uchar)或 [ − 1.0 , 1.0 ] [-1.0, 1.0] [1.0,1.0](对于有符号数,i.e. short/char)中的浮点数
  5. 是否标准化texture坐标

    • 非标准化坐标: [ 0 , N − 1 ] [0, N-1] [0,<
  • 4
    点赞
  • 23
    收藏
    觉得还不错? 一键收藏
  • 0
    评论
CUDA共享内存是一种位于GPU上的高速存,可以用于在同一个线程块内的线程之间共享数据。使用共内存可以显著提内存访问的效,并且减少对全局内存访问次数。 以下使用CUDA共享内存的一般步: 1. 声明共内存:在GPU核函数中,可以使用`__shared__`关键字来声明共享内存。共享内存的大小需要在编译时确定,并且是所有线程块中的线程共享的。 ```cuda __shared__ float sharedData[SIZE]; ``` 2. 将数据从全局内存复制到共享内存:在GPU核函数中,使用线程块中的线程来将数据从全局内存复制到共享内存中。可以使用线程块索引和线程索引来确定数据的位置。 ```cuda int tid = threadIdx.x; int blockId = blockIdx.x; int index = blockId * blockDim.x + tid; sharedData[tid] = globalData[index]; ``` 3. 同步线程块:在将数据复制到共享内存后,需要使用`__syncthreads()`函数来同步线程块中的线程。这样可以确保所有线程都已经将数据复制到共享内存中。 ```cuda __syncthreads(); ``` 4. 使用共享内存:一旦所有线程都已经将数据复制到共享内存中,可以使用共享内存进行计算。由于共享内存位于GPU的高速缓存中,所以访问速度较快。 ```cuda sharedData[tid] += 1.0f; ``` 5. 将数据从共享内存复制回全局内存:在计算完成后,可以使用线程块中的线程将数据从共享内存复制回全局内存。 ```cuda globalData[index] = sharedData[tid]; ``` 需要注意的是,共享内存的大小是有限的,不同的GPU架构及型号都有不同的限制。因此,在使用共享内存时,需要确保不超过设备的限制,并且合理地利用共享内存,以提高性能。此外,需要注意线程同步的位置和使用方法,以避免数据竞争和错误的结果。 以上是使用CUDA共享内存的基本步骤,具体的实现方式会根据具体问题而有所不同。
评论
添加红包

请填写红包祝福语或标题

红包个数最小为10个

红包金额最低5元

当前余额3.43前往充值 >
需支付:10.00
成就一亿技术人!
领取后你会自动成为博主和红包主的粉丝 规则
hope_wisdom
发出的红包
实付
使用余额支付
点击重新获取
扫码支付
钱包余额 0

抵扣说明:

1.余额是钱包充值的虚拟货币,按照1:1的比例进行支付金额的抵扣。
2.余额无法直接购买下载,可以购买VIP、付费专栏及课程。

余额充值