Cuda By Example - 12 (Texture Memory)

《Cuda By Exmaple》文中的Texture Reference方法，CUDA 12已经不再支持了。为了试试Texture内存模式的加速功能，补充利用Texture Object API的例子。

Texture Object API

Texture Object API里，仍然有tex1Dfetch和tex2D两个函数，用于获取Texture里数据。tex1Dfetch和tex2D在新版本中，定义为两个模板函数。

template<class T>
T tex1Dfetch(cudaTextureObject_t texObj, int x);template<class T>
T tex2D(cudaTextureObject_t texObj, float x, float y);

模板参数T为Texture里的数据类型，texObj指代的Texture对象，可以存储各种类型的数据，比如常用的浮点数类型float。tex1Dfetch的第二个参数 int x为数据在Texture中的坐标。由坐标类型为整数可知，tex1Dfetch只支持非归一化的坐标值。tex2D用于从二维Texture对象中取值，x, y分别为在x轴和y轴的坐标值。float类型允许使用归一化后的坐标值。要从一个存储数据类型为float的Texture对象texObj中取出 x 处的值，可以用下面的语句：

float v = tex2D<float>(texObj, x);

因此，前文中的copy_const_kernel和blend_kernel可以改写成：

__global__ void copy_const_kernel(float* iptr, cudaTextureObject_t texConst) {// map threadIdx/blockIdx to x, yint x = threadIdx.x + blockIdx.x * blockDim.x;int y = threadIdx.y + blockIdx.y * blockDim.y;int offset = x + y * blockDim.x * gridDim.x;float v = tex1Dfetch<float>(texConst, offset);if (v != 0)iptr[offset] = v;
}

除了引入了参数cudaTextureObject_t texConst作为热源输入外，跟前文几乎一样。

__global__ void blend_kernel(float* outSrc, cudaTextureObject_t texIn)
{// map threadIdx/blockIdx to x, yint x = threadIdx.x + blockIdx.x * blockDim.x;int y = threadIdx.y + blockIdx.y * blockDim.y;int offset = x + y * blockDim.x * gridDim.x;int left = offset - 1;int right = offset + 1;if (x == 0)left = left + 1;if (x == DIM - 1)right = right - 1;int top = offset - DIM;int bottom = offset + DIM;if (y == 0)top += DIM;if (y == DIM - 1)bottom -= DIM;outSrc[offset] = tex1Dfetch<float>(texIn, offset) +  SPEED * (tex1Dfetch<float>(texIn, top) + tex1Dfetch<float>(texIn, bottom)+ tex1Dfetch<float>(texIn, left) + tex1Dfetch<float>(texIn, right) - 4 * tex1Dfetch<float>(texIn, offset));
}

同样，因为tex1Dfetch不是像之前那样，在编译时就跟Texture Reference静态捆绑在一起，我们可以把输入Texture作为函数参数传进来。

Texture 对象

为了跟前文做直观对比，tex1Dfetch 和 tex2D放在前面讲了。同时可能也引发了我们的好奇心，Texture对象是怎么得来的。

Texture对象在CUDA中的类型名是：cudaTextureObject_t。创建一个Texture对象，我们先声明一个cudaTextureObject_t变量，然后使用cudaCreateTextureObject对变量进行初始化。

先来看一下cudaCreateTextureObject的定义：

cudaError_t cudaCreateTextureObject(cudaTextureObject_t* pTexObject,const cudaResourceDesc* pResDesc,const cudaTextureDesc* pTexDesc,const cudaResourceViewDesc* pResViewDesc
);
  • pTexObject：返回的CUDA纹理对象。
  • pResDesc：纹理资源描述符，描述了纹理在内存中的位置和格式。
  • pTexDesc：纹理描述符，描述了纹理的滤波和寻址模式等属性。
  • pResViewDesc：资源视图描述符，描述了纹理的视图。

因此要创建一个Texture对象，我们必须提供三方面的信息。纹理视图跟我们的例子无关，可以先不管。剩下的cudaResourceDesc结构描述的是Texture对象存放数据的内存信息，cudaTextureDesc描述了访问内存时的一些规定。

cudaResourceDesc

cudaResourceDesc包含了一个枚举类型enum cudaResourceType，用于描述资源的类型。当前CUDA支持4中资源，分别是：Array，Mipmapped Array，Linear 和 Pitch 2D。cudaResourceDesc的其余部分是一个联合体，对应每一种资源类型，定义了描述该类资源结构。本例中我们使用Linear类型资源，即一维数组。需要提供的信息包含在下面的结构中。

        struct {void *devPtr;                      /**< Device pointer */struct cudaChannelFormatDesc desc; /**< Channel descriptor */size_t sizeInBytes;                /**< Size in bytes */} linear;

devPtr为跟Texture关联的内存的指针，通过cudaMalloc分配得到。sizeInBytes指定以byte为单位，内存的长度，我们的内存大小为显示网格的图像大小。desc描述内存数据包含的通道数量及类型，我们使用单通道的float类型。

我们使用下面的代码初始化热源的cudaResourceDesc：

cudaMalloc((void**)&d.dev_constSrc, bitmap.image_size());
cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc<float>();
struct cudaResourceDesc resDesc;
memset(&resDesc, 0, sizeof(resDesc));
resDesc.resType = cudaResourceTypeLinear;
resDesc.res.linear.devPtr = d.dev_constSrc;
resDesc.res.linear.desc = channelDesc;
resDesc.res.linear.sizeInBytes = bitmap.image_size();

cudaTextureDesc

cudaTextureDesc包含了很多变量，addressMode指定坐标超出边界时，如何处理坐标，我们自己维护了坐标值确保了它不会出界。 不过在2D Texture的示例中，我们需要把它设置为cudaAddressModeClamp。 readMode 和 normalizedCoords 也需要指定。 我们需要读取的数组元素，readMode 设置为cudaReadModeElementType。坐标为原始坐标，不需要归一化，设置为0.

由此我们使用下面的代码初始化热源的cudaTextureDesc：

	struct cudaTextureDesc texDesc;memset(&texDesc, 0, sizeof(texDesc));texDesc.addressMode[0] = cudaAddressModeClamp;texDesc.addressMode[1] = cudaAddressModeClamp;//texDesc.filterMode = cudaFilterModePoint;texDesc.readMode = cudaReadModeElementType;//texDesc.normalizedCoords = 0;

最后创建Texture对象，我们的Texture对象声明在DataBlock里，所以有以下的代码：

cudaCreateTextureObject(&d.texConstSrc, &resDesc, &texDesc, NULL);

完整代码

In 和 Out 的Texture对象的声明跟热源类型，只是用到不同内存，其他都一样。下面是完整的代码，在CUDA 12 编译通过。遗憾的是性能还不如最早使用Global Momory的版本。

#include "book.h"
#include "cpu_anim.h"#define  SPEED  0.25f
const int DIM = 1024;const float MAX_TEMP = 1.0f;
const float MIN_TEMP = 0.0001f;__global__ void copy_const_kernel(float* iptr, cudaTextureObject_t texConst) {// map threadIdx/blockIdx to x, yint x = threadIdx.x + blockIdx.x * blockDim.x;int y = threadIdx.y + blockIdx.y * blockDim.y;int offset = x + y * blockDim.x * gridDim.x;float v = tex1Dfetch<float>(texConst, offset);if (v != 0)iptr[offset] = v;
}__global__ void blend_kernel(float* outSrc, cudaTextureObject_t texIn)
{// map threadIdx/blockIdx to x, yint x = threadIdx.x + blockIdx.x * blockDim.x;int y = threadIdx.y + blockIdx.y * blockDim.y;int offset = x + y * blockDim.x * gridDim.x;int left = offset - 1;int right = offset + 1;if (x == 0)left = left + 1;if (x == DIM - 1)right = right - 1;int top = offset - DIM;int bottom = offset + DIM;if (y == 0)top += DIM;if (y == DIM - 1)bottom -= DIM;outSrc[offset] = tex1Dfetch<float>(texIn, offset) +  SPEED * (tex1Dfetch<float>(texIn, top) + tex1Dfetch<float>(texIn, bottom)+ tex1Dfetch<float>(texIn, left) + tex1Dfetch<float>(texIn, right) - 4 * tex1Dfetch<float>(texIn, offset));
}struct DataBlock {unsigned char* output_bitmap;float * dev_inSrc;float * dev_outSrc;float * dev_constSrc;cudaTextureObject_t texConstSrc;cudaTextureObject_t texIn;cudaTextureObject_t texOut;CPUAnimBitmap* bitmap;cudaEvent_t start, stop;float totalTime;float frames;
};void anim_gpu(DataBlock* d, int ticks) {cudaEventRecord(d->start, 0);dim3 blocks(DIM / 16, DIM / 16);dim3 threads(16, 16);// tranfer 90 times for a good displayvolatile bool dstOut = true;for (int i = 0; i < 90; i++) {if (dstOut) {copy_const_kernel <<<blocks, threads>>> (d->dev_inSrc, d->texConstSrc);blend_kernel <<<blocks, threads>>> (d->dev_outSrc, d->texIn);}else {copy_const_kernel <<<blocks, threads>>> (d->dev_outSrc, d->texConstSrc);blend_kernel <<<blocks, threads>>> (d->dev_inSrc, d->texOut);}dstOut = !dstOut;}float_to_color << <blocks, threads >> > (d->output_bitmap, d->dev_outSrc);//Copy image from devicecudaMemcpy(d->bitmap->get_ptr(), d->output_bitmap, d->bitmap->image_size(), cudaMemcpyDeviceToHost);cudaEventRecord(d->stop, 0);cudaEventSynchronize(d->stop);float elapsedTime;cudaEventElapsedTime(&elapsedTime, d->start, d->stop);d->totalTime += elapsedTime;++d->frames;printf("Average Time per frame: %3.1f ms\n", d->totalTime / d->frames);
}void anim_exit(DataBlock* d) {cudaDestroyTextureObject(d->texConstSrc);cudaDestroyTextureObject(d->texIn);cudaDestroyTextureObject(d->texOut);cudaFree(d->dev_constSrc);cudaFree(d->dev_inSrc);cudaFree(d->dev_outSrc);cudaEventDestroy(d->start);cudaEventDestroy(d->stop);
}int main(int argc, char* argv[]) {DataBlock d;CPUAnimBitmap bitmap(DIM, DIM, &d);d.bitmap = &bitmap;d.frames = 0;d.totalTime = 0.0f;cudaEventCreate(&d.start);cudaEventCreate(&d.stop);HANDLE_ERROR(cudaMalloc((void**)&d.output_bitmap, bitmap.image_size()));//Allocate CUDA Array in device memory.cudaMalloc((void**)&d.dev_constSrc, bitmap.image_size());cudaMalloc((void**)&d.dev_inSrc, bitmap.image_size());cudaMalloc((void**)&d.dev_outSrc, bitmap.image_size());// Create Texture obj for const,in and outcudaChannelFormatDesc channelDesc = cudaCreateChannelDesc<float>();struct cudaResourceDesc resDesc;memset(&resDesc, 0, sizeof(resDesc));resDesc.resType = cudaResourceTypeLinear;resDesc.res.linear.devPtr = d.dev_constSrc;resDesc.res.linear.desc = channelDesc;resDesc.res.linear.sizeInBytes = bitmap.image_size();struct cudaTextureDesc texDesc;memset(&texDesc, 0, sizeof(texDesc));texDesc.addressMode[0] = cudaAddressModeClamp;texDesc.addressMode[1] = cudaAddressModeClamp;//texDesc.filterMode = cudaFilterModePoint;texDesc.readMode = cudaReadModeElementType;//texDesc.normalizedCoords = 0;cudaCreateTextureObject(&d.texConstSrc, &resDesc, &texDesc, NULL);resDesc.res.linear.devPtr = d.dev_inSrc;cudaCreateTextureObject(&d.texIn, &resDesc, &texDesc, NULL);resDesc.res.linear.devPtr = d.dev_outSrc;cudaCreateTextureObject(&d.texOut, &resDesc, &texDesc, NULL);//create constant heaterfloat* temp = (float*)malloc(bitmap.image_size());for (int i = 0; i < DIM * DIM; i++) {temp[i] = 0;int x = i % DIM;int y = i / DIM;if ((x > 300) && (x < 600) && (y > 310) && (y < 601))temp[i] = MAX_TEMP;}temp[DIM * 100 + 100] = (MAX_TEMP + MIN_TEMP) / 2.0;temp[DIM * 700 + 100] = MIN_TEMP;temp[DIM * 300 + 300] = MIN_TEMP;temp[DIM * 200 + 700] = MIN_TEMP;for (int y = 800; y < 900; y++) {for (int x = 400; x < 500; x++) {temp[x + y * DIM] = MIN_TEMP;}}// Copy data located at address h_data in host memory to device memorycudaMemcpy(d.dev_constSrc, temp, bitmap.image_size(), cudaMemcpyHostToDevice);for (int y = 800; y < DIM; y++) {for (int x = 0; x < 200; x++) {temp[x + y * DIM] = MAX_TEMP;}}cudaMemcpy(d.dev_inSrc, temp, bitmap.image_size(), cudaMemcpyHostToDevice);free(temp);bitmap.anim_and_exit((void(*)(void*, int))anim_gpu, (void(*)(void*))anim_exit);cudaFree(d.output_bitmap);
}

本文来自互联网用户投稿，该文观点仅代表作者本人，不代表本站立场。本站仅提供信息存储空间服务，不拥有所有权，不承担相关法律责任。如若转载，请注明出处：http://www.rhkb.cn/news/461914.html

如若内容造成侵权/违法违规/事实不符，请联系长河编程网进行投诉反馈email:809451989@qq.com，一经查实，立即删除！

相关文章

在服务器运维过程中，发现服务器时间倒退以及DNS无法解析域名造成yum不可用的问题解决

在服务器运维过程中，发现服务器时间倒退以及DNS无法解析域名造成yum不可用的问题解决

目录 一.问题描述 二.问题排查过程 2.1yum下载NTP 2.2排查DNS 三.问题解决过程 3.1修复DNS 3.2更新yum源 3.3下载ntp 四.问题解决结果 4.1ntp服务情况检查 4.2服务器时间检查 4.3软件系统时间检查 一.问题描述 对服务器进行运维的过程中&#xff0c;发现服务器时间…
阅读更多...
Redis高频面试题

Redis高频面试题

一、Redis有什么好处？ 高性能：Redis是一个基于内存的数据存储系统，相比于传统的基于磁盘的数据库系统，它能够提供更高的读写性能。支持丰富的数据类型：Redis支持多种数据结构，包括字符串、哈希、列表、集合、有序集合等，这使得它可以用于多种不同的应用场景。持久化：Re…
阅读更多...
[POI2014] PTA-Little Bird（单调队列优化 DP）

[POI2014] PTA-Little Bird（单调队列优化 DP）

luogu 传送门https://www.luogu.com.cn/problem/P3572 解题思路 先设 表示到 的最小劳累值。 很容易得出转移&#xff1a; 其中 由 和 的大小关系决定&#xff0c;并且 。 很显然&#xff0c;直接暴力是 的&#xff0c;会超时。 于是&#xff0c;考虑优化。 我们发现…
阅读更多...
如何在Linux系统中使用Apache HTTP Server

如何在Linux系统中使用Apache HTTP Server

如何在Linux系统中使用Apache HTTP Server Apache简介 安装Apache 在Debian/Ubuntu系统中安装 在CentOS/RHEL系统中安装 启动Apache服务 验证Apache是否正在运行 访问Apache默认页面 配置Apache虚拟主机 创建虚拟主机配置文件 示例虚拟主机配置 创建网站根目录 准备静态网站内…
阅读更多...
ISME Comm | 西南大学时伟宇团队在功能基因水平揭示植被演替过程中磷限制对土壤微生物碳代谢潜力的抑制作用机制

ISME Comm | 西南大学时伟宇团队在功能基因水平揭示植被演替过程中磷限制对土壤微生物碳代谢潜力的抑制作用机制

本文首发于“生态学者”微信公众号&#xff01; 植被群落长期演替过程中&#xff0c;生态系统普遍受养分限制&#xff0c;微生物群落代谢功能在生态系统物质循环中尤为关键。西南大学时伟宇教授团队联合国内外学者&#xff0c;在功能基因水平&#xff0c;将微生物群落功能纳入生…
阅读更多...
Unity控制物体透明度的改变

Unity控制物体透明度的改变

目录标题 效果图代码调用注意事项 效果图 代码 注意&#xff1a;在控制全部的模型进行透视时&#xff0c;已经隐藏的子物体仍然要处理。 using System.Collections; using System.Collections.Generic; using UnityEngine; using DG.Tweening; public class FadeModel {priva…
阅读更多...
工业网络监控中的IP保护与软件授权革新

工业网络监控中的IP保护与软件授权革新

未来的智能工厂离不开稳定而高效的通信网络&#xff0c;这些网络在支撑生产流程的同时&#xff0c;也面临着复杂的管理与安全挑战。PROCENTEC推出了一系列硬件和软件产品&#xff0c;如Atlas、Mercury和Osiris&#xff0c;以提供全面的网络监控和故障排除能力。然而&#xff0c…
阅读更多...
springboot 整合 抖音 移动应用 授权

springboot 整合 抖音 移动应用 授权

后端开发&#xff0c;因为没有JavaSDK&#xff0c;maven依赖&#xff0c;用到的是API接口去调用 抖音API开发文档 开发前先申请好移动应用&#xff0c;抖音控制台-移动应用 之后还需要开通所有能开通的能力 拿到应用的 clientKey 和 clientSecret&#xff0c;就可以进入开发了 …
阅读更多...
后台管理系统的通用权限解决方案(七)SpringBoot整合SpringEvent实现操作日志记录（基于注解和切面实现）

后台管理系统的通用权限解决方案(七)SpringBoot整合SpringEvent实现操作日志记录（基于注解和切面实现）

1 Spring Event框架 除了记录程序运行日志&#xff0c;在实际项目中一般还会记录操作日志&#xff0c;包括操作类型、操作时间、操作员、管理员IP、操作原因等等&#xff08;一般叫审计&#xff09;。 操作日志一般保存在数据库&#xff0c;方便管理员查询。通常的做法在每个…
阅读更多...
视频设备一体化监控运维方案

视频设备一体化监控运维方案

随着平安城市、雪亮工程等项目建设的号召&#xff0c;视频监控系统的建设如火如荼地开展。无论在公共场所、企业单位、住宅小区、矿山工地还是交通枢纽&#xff0c;视频监控系统已成为保障安全、维护秩序和提升管理效率的重要工具。但由于对视频监控系统中的前端设备&#xff0…
阅读更多...
二十八、Python基础语法（面向对象-下）

二十八、Python基础语法（面向对象-下）

一、self 从函数的语法上来看, self 是形参 , 是一个普通的参数,那么在调用的时候,就需要传递实参值。从调用上看, 我们没有给 self 这个形参传递实参值, 但是 Python 解释器会自动的将调用这个方法的对象&#xff0c;作为实参值传递给 self。 class Dog:def eat(self):print…
阅读更多...
【Leecode】Leecode刷题之路第37天之解数独

【Leecode】Leecode刷题之路第37天之解数独

题目出处 37-解数独-题目出处 题目描述 个人解法 思路&#xff1a; todo代码示例&#xff1a;&#xff08;Java&#xff09; todo复杂度分析 todo官方解法 37-解数独-官方解法 方法1&#xff1a;回溯 思路&#xff1a; 代码示例&#xff1a;&#xff08;Java&#xff09; p…
阅读更多...
【golang/navmesh】使用recast navigation进行寻路

【golang/navmesh】使用recast navigation进行寻路

目录 说在前面安装使用可视化 说在前面 go version&#xff1a;1.20.2 linux/amd64操作系统&#xff1a;wsl2detour-go版本&#xff1a;v0.2.0github&#xff1a;这里&#xff0c;求star! 安装 使用go mod安装即可go get github.com/o0olele/detour-go使用 使用场景模型构建n…
阅读更多...
qt QFormLayout详解

qt QFormLayout详解

QFormLayout 是 Qt 框架中用于创建表单布局的一个类&#xff0c;适合于将标签和输入控件整齐地排列在一起。它可以帮助开发者轻松构建用户输入界面&#xff0c;尤其是在处理表单时。 QFormLayout以两列的形式展示其子项&#xff0c;常用于创建“标签-字段”对的布局。其中&…
阅读更多...
电脑小白必看｜电脑安装常用软件简单小技巧

电脑小白必看｜电脑安装常用软件简单小技巧

前言 最近同事换了新电脑&#xff0c;问我怎么下载常用软件&#xff1f; 我反问了一下&#xff1a;什么常用软件呢&#xff1f; 她说&#xff1a;微信、QQ、钉钉、酷狗、wps这种类型的软件。 哦豁&#xff0c;那其实很简单&#xff0c;但很多人还是没学会。小白之前分享过一…
阅读更多...
RocketMQ 消息消费失败的处理机制

RocketMQ 消息消费失败的处理机制

在分布式消息系统中&#xff0c;处理消费失败的消息是非常关键的一环。 RocketMQ 提供了一套完整的消息消费失败处理机制&#xff0c;下面我将简要介绍一下其处理逻辑。 截图代码版本&#xff1a;4.9.8 步骤1 当消息消费失败时&#xff0c;RocketMQ会发送一个code为36的请求到…
阅读更多...
数据结构算法学习方法经验总结

数据结构算法学习方法经验总结

DSA:Data Structures, Algorithms, and Problem-Solving Techniques 三大核心支柱 一次学习一个主题&#xff0c;按照如下顺序学习 如何开始学习新的主题 学习资源 https://www.youtube.com/playlist?listPLDN4rrl48XKpZkf03iYFl-O29szjTrs_O (Algorithms) https://ww…
阅读更多...
Linux 操作系统的诞生与发展历程

Linux 操作系统的诞生与发展历程

目录 背景与起源 诞生过程 特点与影响 背景与起源 历史背景&#xff1a; 1980年代末至1990年代初&#xff0c;计算机操作系统市场主要由商业软件主导&#xff0c;如DOS、Windows以及Unix的各种版本。然而&#xff0c;这些系统往往价格昂贵&#xff0c;且源代码不开放&#…
阅读更多...
第三届北京国际水利科技博览会将于25年3月在国家会议中心召开

第三届北京国际水利科技博览会将于25年3月在国家会议中心召开

由中国农业节水和农村供水技术协会、北京水利学会、振威国际会展集团等单位联合主办的第三届北京国际水利科技博览会暨供水技术与设备展&#xff08;北京水利展&#xff09;将于2025年3月31日至4月2日在北京•国家会议中心举办&#xff01; 博览会以“新制造、新服务、新业态”…
阅读更多...
贪心算法习题其二【力扣】【算法学习day.19】

贪心算法习题其二【力扣】【算法学习day.19】

前言 ###我做这类文档一个重要的目的还是给正在学习的大家提供方向&#xff08;例如想要掌握基础用法&#xff0c;该刷哪些题&#xff1f;&#xff09;我的解析也不会做的非常详细&#xff0c;只会提供思路和一些关键点&#xff0c;力扣上的大佬们的题解质量是非常非常高滴&am…
阅读更多...
最新文章
推荐文章

Copyright @ 2022~2023