CUDA Texture Memory
在執行 CUDA 程式前,都要把資料先從 Host 的記憶體,複製一份到 device 的記憶體中;一般來說,這樣的部分,都是使用 device 的 global memory 來直接進行存取。不過實際上,有的時候還有別的選擇的~在《nVidia CUDA 簡介》中一文就有提到,除了 global memory 外,還可以透過 constant memory 或 texture memory 的形式,來對 device memory 資料的存取。
texture 是一般 graphics 裡的名詞,2D texture 大致上可以理解為一張圖片,一般應該是翻譯成材質(維基百科是稱為「紋理」);而由於在傳統的 render pipeline 中,texture 佔了很重要的地位,所以在顯示卡的部分,也會對這部分做特殊的最佳化。而在 nVidia 的 CUDA 中,也把 texture 這項元素保留下來了!
如果在 CUDA 中把 device memory 的資料,當作使用 texture 的話,那資料會變成是唯讀的,要透過特殊的函式來讀取,沒有辦法進行修改;不過相對的,和 global memory 或 constant memory 比起來,也有不少優點~(詳細資料請參考《CUDA Programming Guide 1.1》的 5.4)
- 有快取,在某些狀況下頻寬會比較大
- 不像 global/constant memory 要依照某些存取模式下才有比較好的效能
- 位置計算的延遲被隱藏得更好,對於隨機存取資料的效能可能更好
- 被封包過的資料可以在一個運算中被個別的變數使用
- 8bit 和 16bit 整數可以簡單的轉換成 [0.0, 1.0] 或 [-1.0, 1.0]
除了上述的優點外,如果使用 CUDA Array 來當 texture 的話,更可以套用 filter、使用內建的功能來做內插取值、並且設定要用 clamping 或 repeat 模式處理邊界。
而在 CUDA 中,要使用 texture 的話,是要使用所謂的「texture reference」;下面就大概來介紹 CUDA 中 texture 的使用方法。首先,他的大致流程會是:
- Host 的 texture 建立部分
- 宣告出 texture reference
- 透過 Bind Texture 的函式,將 texture reference 和現有的 device memory 上的變數(linear memory 或 CUDA array)做連結
- Device 使用
- 透過 CUDA 提供的 texture 讀取函式(tex1Dfetch, tex1D, tex2D)來讀取 texture 的內容
- Host 刪除 texture
- 呼叫 unbind texture 的函式,將 texture reference 的資源釋放
不過,在 CUDA 中,對於 texture 的控制,有 low-level 和 high level 兩種,Heresy 在這就先針對比較簡單的 high-level 方法來做一些簡單的說明。
CUDA 的 texture 型別
在 CUDA 中,有提供名為 texture 的 template 型別,他的形式是:
texture<Type, Dim, ReadMode> texRef;复制代码
其中:
- Type
texture 中元素的資料型別;可以是一般的基本的 int, float 型別,也可以是 CUDA 中提供的 vector 型別。
- Dim
代表這個 texture 的維度,在 CUDA 中只有 1 或 2 兩種值,並不支援 3D texture。
- ReadMode
讀取 texture 的模式,有 cudaReadModeNormalizedFloat 和 cudaReadModeElementType 兩種模式。當模式是 cudaReadModeElementType 時,資料會以原來的方式讀取出來;當模式是 cudaReadModeNormalizedFloat、且資料型別是整數型別時,他則會對資料進行 normalize,傳回 [0,1] 或 [-1,1] 之間的數(視原始型別是否為 unsigned 決定)。舉個例子,如果我們要宣告一個資料型別是 int 的 1D texture,就可以寫成:
texture<int, 1, cudaReadModeElementType> texRef;复制代码
而 texRef 的資料,則還要再透過 BindTexture 的函式,來和 device 上的變數做連結,這樣才算完成 texture 的使用前準備。
不過另外一個要注意的是,目前的 CUDA 似乎只允許將 texture reference 宣告成 globalvariable,而無法將它宣告在函式內,用參數的方法傳遞給 kernel function(Heresy 這樣寫會使得 nvcc產生內部錯誤,要寫在 file-scope 是參考 ISI 的課程後才知道的)。
兩種不同的 texture 資料
在 CUDA 中,可以接受兩種資料:linear memory 和 CUDA array。其中,linear memory 就是之前提過,用 cudaMalloc()cudaMallocArray(),來宣告出一塊連續的 1D/2D 陣列。 宣告出的連續記憶體空間(一維陣列);而 CUDA array 則是透過 由兩種不同的資料所建立出來的 texture reference,在使用上有一些不同的性質:
Texture with linear memory Texture with CUDA Array 維度 只有一維 一維或二維 讀取的索引 整數 整數或浮點數 filter 不支援 支援 (cudaFilterModeLinear / cudaFilterModePoint) 邊界值 n/a cudaAddressModeClamp / cudaAddressModeWrap 讀取函式 tex1Dfetch() tex1D() / tex2D() 透過 CUDA array 的 texture,可以讓 CUDA 直接幫忙做好內插(雖然目前看來只有線性內插)、邊值處理的問題,其實在很多時候都是相當方便的~
而由於使用 linear memory 和 CUDA array 的 texture 在使用上有不少的差異,所以接下來就分開寫吧
转载于:https://blog.51cto.com/general/239036
CUDA Texture Memory相关推荐
- CUDA ---- Shared Memory
CUDA SHARED MEMORY shared memory在之前的博文有些介绍,这部分会专门讲解其内容.在global Memory部分,数据对齐和连续是很重要的话题,当使用L1的时候,对齐问题 ...
- workaround for %33 texture memory bug
原帖链接:http://www.cocos2d-iphone.org/forum/topic/29121 PS: 为什么要关心 NPOT 呢? 因为苹果的OpenGL驱动有一个bug,导致如果使用 P ...
- CUDA内存类型memory
CUDA内存类型memory 2013-09-30 10:34 4224人阅读 评论(1) 收藏 举报 分类: CUDA(106) http://www.cnblogs.com/traceorigin ...
- cuda的global memory介绍
CUDA Memory Model 对于程序员来说,memory可以分为下面两类: Programmable:我们可以灵活操作的部分. Non-programmable:不能操作,由一套自动机制来达到 ...
- RuntimeError: CUDA error: an illegal memory access was encountered
文章目录 4. 我的解决 1. 错误描述 2. 自我尝试 2.1 减小batch_size 2.2 换卡改代码 3. 调研情况 4. 我的解决 后续发现其实是某张卡有问题, 0~3一共4个GPU,只在 ...
- CUDA系列学习(一)An Introduction to GPU and CUDA
本文从软硬件层面讲一下CUDA的结构,应用,逻辑和接口.分为以下章节: (一).GPU与CPU (二).CUDA硬件层面 (三).CUDA安装 (四).CUDA 结构与接口 4.1 Kernels 4 ...
- 提升CUDA程序运行效率的几个关键点
目录 1.明确计算机中GPU卡片的计算资源,决定变量的性质(constant,share还是global)以及Grid,Block的维度,充分并合理利用GPU显卡的资源 2.提高PCI接口与GPU显卡 ...
- 深入浅出CUDA编程
标签: cuda编程threadfloatconflictexpress 2010-12-10 13:29 44960人阅读 评论(7) 收藏 举报 CUDA 是 NVIDIA 的 GPGPU 模型, ...
- CUDA学习笔记之 CUDA存储器模型
CUDA学习笔记之 CUDA存储器模型 标签: cuda存储bindingcache编程api 2010-12-14 01:33 1223人阅读 评论(0) 收藏 举报 分类: CUDA(26) GP ...
最新文章
- 为什么数据科学不值得?
- telnet 退出命令
- 42、谈谈你对闭包的理解?
- Warning: Attempt to present on whose view is not in模态跳转问题
- 庚顿数据:实时数据库赋能工业互联网
- 《C++高级编程(第3版)》
- linux成功mysql数据直接拷贝_mysql数据库数据从一个linux系统移植到另一个linux系统的方法...
- Linux(2) vi和vim编辑器
- OpenCV——读取图片的数据类型必须是int8类型吗?
- 基础Git操作与GitHub协作吐血整理,收好!| 原力计划
- plsa的java实现_java在注解中绑定方法参数的解决方案
- android 开发 目标绑定,Hippy: Hippy 是一个新生的跨端开发框架,目标是使开发者可以只写一套代码就直接运行于三个平台(iOS、Android 和 Web)...
- Dalvik和Java运行环境的区别
- Rbf神经网络使用Tensorflow实现
- Android SVG矢量图/矢量动画、Vector和VectorDrawable矢量图及动画,减少App Size
- 发票管理小工具(三):PDFMiner vs pdfminer3k vs Pdfminer.six
- 物联网和边缘部署的5大嵌入式工控机设计要求
- C++复习炒剩饭(1)心一意
- 融会贯通,并行不悖丨2022年8月《中国数据库行业分析报告》发布!
- 动手试试!手把手教你如何适配 iPhone X
热门文章
- /var/spool/postfix/maildrop小文件太多造成inode索引使用完解决
- Android Eclipse ADT使用Tips
- Entity Framework加载相关实体——Explicit Loading
- ExtTabMenu 控件
- 32个程序员泪(méng)流(fān)满(quán)面(chǎng)的瞬间
- css 解析 开源库_干货 | python库大全,全面高效
- 刚写的代码,就变成了遗留系统?
- 程序写不好,总理当到老!
- Jeecg-P3 1.0.1版本发布,轻量级微服务框架
- PLSQL创建用户及权限分配