文章目录

  • 前言
  • 1. 纹理内存的理解
  • 2. C++ CUDA实现常量内存使用
  • 3. 执行结果
  • 总结
  • 学习资料

VS2017 CUDA编程学习1:CUDA编程两变量加法运算
VS2017 CUDA编程学习2:在GPU上执行线程
VS2017 CUDA编程学习3:CUDA获取设备上属性信息
VS2017 CUDA编程学习4:CUDA并行处理初探 - 向量加法实现
VS2017 CUDA编程学习5:CUDA并行执行-线程
VS2017 CUDA编程学习6: GPU存储器架构
VS2017 CUDA编程学习7:线程同步-共享内存
VS2017 CUDA编程学习8:线程同步-原子操作
VS2017 CUDA编程学习9:常量内存


前言

这里继续CUDA编程的学习,今天学习了设备纹理内存的使用,这里分享给大家!


1. 纹理内存的理解

纹理内存,是设备中一种只读存储器,在设备内存中也有对应的缓冲器Cache来加速IO操作。纹理内存对于程序中有空间邻近性访问的数据很高效。这里的空间邻近性是指每个线程的读取位置和其他线程的读取位置邻近,这在图像处理中非常有用,因为图像处理中经常需要进行邻域比较,比如4邻域和8邻域。

当然全局内存也可以进行这种有空间邻近性特性数据的存储与访问,但是速度上要慢很多,需要进行大量的显存读取操作。但是纹理内存则只需要从显存读取一次就可以(不理解这种说法,书中这么说的),纹理内存支持2维和3维纹理读取操作。


2. C++ CUDA实现常量内存使用

这里以一个简单的例子:从纹理内存中读取数据并赋值。

纹理内存需要和“纹理引用”与CUDA数组配套使用来实现。
纹理引用通过texture<param1,param2,param3>texture<param1, param2, param3>texture<param1,param2,param3>类型变量来定义,param1表示纹理元素的类型,param2表示纹理引用的类型(1-表示1D, 2-表示2D, 3-表示3D),param3则表示读取模式,这是一个可选参数,说明是否要执行自动类型转换。

注意,纹理引用变量要确保被定义为全局静态变量(个人理解,这样才能在内核函数中使用),同时还要确保这个变量不能作为参数传递给任何其他函数。

下面的例子中通过纹理引用读取线程ID作为索引位置的数据,然后复制到d_out指针指向全局内存中。

在main函数中则使用CUDA数组cudaArraycudaArraycudaArray并通过cudaMemcpyToArray(param1,param2,param3,param4,param5,param6)cudaMemcpyToArray(param1, param2, param3, param4, param5, param6)cudaMemcpyToArray(param1,param2,param3,param4,param5,param6)对CUDA数组进行赋值,param1表示目标CUDA数组变量, param2和param3表示将主机数据赋值到CUDA数组横向和纵向的偏移量,(0, 0)表示CUDA数组的左上角开始赋值。

使用cudaBindTextureToArray()cudaBindTextureToArray()cudaBindTextureToArray()函数将纹理引用和CUDA数组绑定,这样就可以通过访问纹理引用来获取纹理内存中存储的CUDA数组数据。这里使用texture相应的函数进行访问。当使用完纹理内存后,要执行解除纹理引用与CUDA数组的绑定的代码,即调用cudaUnbindTexture()cudaUnbindTexture()cudaUnbindTexture()。

详细代码如下所示:

#include <stdio.h>
#include <iostream>
#include <cuda.h>
#include <cuda_runtime.h>
#include <cuda_runtime_api.h>
#include <device_launch_parameters.h>#ifndef __CUDACC__
#define __CUDACC__
#endif
#include <cuda_texture_types.h>//texture<> depend head file
#include <texture_fetch_functions.h>//tex1D() depend head file#define NUM_THREADS 10
#define N 10//定义1维的纹理引用
texture<float, 1, cudaReadModeElementType> textureRef;
//textureReference textureRef;//定义内核函数:从纹理内存中获取数据并赋值给设备内存
__global__ void gpu_texture_memory(int n, float* d_out)
{int idx = blockIdx.x *blockDim.x + threadIdx.x;if (idx < n){float temp = tex1D(textureRef, float(idx));d_out[idx] = temp;}
}int main()
{int num_blocks = N / NUM_THREADS + ((N % NUM_THREADS) ? 1 : 0);float* d_out;cudaMalloc(&d_out, sizeof(float) * N);float* h_out = (float*)malloc(sizeof(float) * N);float h_in[N];for (int i = 0; i < N; i++){h_in[i] = float(i);}//定义CUDA数组cudaArray *cu_Array;cudaMallocArray(&cu_Array, &textureRef.channelDesc, N, 1);cudaMemcpyToArray(cu_Array, 0, 0, h_in, sizeof(float)*N, cudaMemcpyHostToDevice);//绑定CUDA数组到纹理引用变量cudaBindTextureToArray(&textureRef, cu_Array, &textureRef.channelDesc);//调用内核函数gpu_texture_memory << <num_blocks, NUM_THREADS >> > (N, d_out);//拷贝结果到主机cudaMemcpy(h_out, d_out, sizeof(float)*N, cudaMemcpyDeviceToHost);printf("在GPU上使用纹理内存:\n");//打印结果for (int i = 0; i < N; i++){printf("%f\n", h_out[i]);}//回收分配的资源free(h_out);cudaFree(d_out);cudaFreeArray(cu_Array);cudaUnbindTexture(&textureRef);system("pause");return 0;
}

3. 执行结果


总结

说实话,这里还是无法体会到纹理内存的提速,这里只是简单介绍了纹理内存的用法,由于这里只是记录个人的学习笔记,所以如果有误,还请原谅,如果能够指正,万分感谢!

学习资料

《基于GPU加速的计算机视觉编程》

VS2017 CUDA编程学习10:纹理内存相关推荐

  1. VS2017 CUDA编程学习11:CUDA性能测量

    文章目录 前言 1. CUDA事件API 2. C++ 实现CUDA事件例子 3. CUDA事件例子的执行结果 4. NVIDIA Visual Profiler工具 总结 学习资料 VS2017 C ...

  2. VS2017 CUDA编程学习8:线程同步-原子操作

    文章目录 前言 1. 原子操作的理解 2. C++ CUDA实现原子操作 3. 执行结果 总结 学习资料 VS2017 CUDA编程学习1:CUDA编程两变量加法运算 VS2017 CUDA编程学习2 ...

  3. VS2017 CUDA编程学习4:CUDA并行处理初探 - 向量加法实现

    文章目录 前言 CUDA实现向量加法 2. 结果展示 总结 学习资料 VS2017 CUDA编程学习1:CUDA编程两变量加法运算 VS2017 CUDA编程学习2:在GPU上执行线程 VS2017 ...

  4. VS2017 CUDA编程学习5:CUDA并行执行-线程

    文章目录 前言 1. 线程 2. 代码实现 总结 学习资料 VS2017 CUDA编程学习1:CUDA编程两变量加法运算 VS2017 CUDA编程学习2:在GPU上执行线程 VS2017 CUDA编 ...

  5. CUDA 编程学习

    0 简单的CUDA简介 1.简单教程 CUDA C ++只是使用CUDA创建大规模并行应用程序的方法之一.它允许您使用功能强大的C ++编程语言来开发由GPU上运行的数千个并行线程加速的高性能算法.许 ...

  6. CUDA编程学习3——并行计算初窥CUDA的软硬件架构

    目录 并行计算概述 查看GPU相关信息 软硬件架构基础 物理层(物理结构) 逻辑层(kernel组织) 物理层和逻辑层的总结 参考 并行计算概述 所谓并行计算的概念定义 同时多个计算资源一起工作(逻辑 ...

  7. CUDA编程学习0——环境搭建环境详解

    目录 环境配置 软件安装 1.支持最高的cuda版本查询,下载cuda开发软件: 3.配置环境(~/.bashrc添加环境变量) 4.后续维护查询 补:关于windows下的cuda环境配置 一.Vi ...

  8. Cuda编程学习(一)

    使用Cuda进行GPU编程 --GPU高性能编程Cuda实战    第四章C并行编程总结 开始学习Cuda和GPU编程的相关知识啦. 感觉GPU编程会越来越重视,尤其是在移动端的计算复杂度优化方面.掌 ...

  9. GPU编程自学10 —— 流并行

    深度学习的兴起,使得多线程以及GPU编程逐渐成为算法工程师无法规避的问题.这里主要记录自己的GPU自学历程. 目录 <GPU编程自学1 -- 引言> <GPU编程自学2 -- CUD ...

  10. 2023年的深度学习入门指南(10) - CUDA编程基础

    2023年的深度学习入门指南(10) - CUDA编程基础 上一篇我们走马观花地看了下SIMD和GPGPU的编程.不过线条太粗了,在开发大模型时遇到问题了肯定还会晕. 所以我们还是需要深入到CUDA中 ...

最新文章

  1. 格式化硬盘 FAT32 和NTFS 什么区别?
  2. Android JNI 第三篇 Java参数类型与本地参数类型对照
  3. OPENCV中的数据结构总结
  4. 【控制】《多智能体机器人系统信息融合与协调》范波老师-第8章-Agent 技术在机器人智能控制系统的应用
  5. 隔壁大爷都会网络布线啦、搞起来搞起来
  6. Serverless 时代前端避坑指南
  7. xmlrpc.php 漏洞利用
  8. mybatis注解开发_Spring Boot 中集成 MyBatis
  9. 第十篇 JVM核心机制之JVM运行和类加载全过程(五)
  10. 机器学习之路: python 支持向量机 LinearSVC 手写字体识别
  11. Java编程:排序算法——归并排序
  12. 计算机软件在哪里建文本文档,电脑点击右键的新建文本文档不见了的解决方法 怎么解决电脑点击右键的新建文本文档不见了...
  13. 用setuna截图时自动放大_setuna截图软件下载
  14. python 相关系数显著检验_Python dataframe 算相关系数用corr(),算不出结果
  15. 七种寻址方式(立即寻址、寄存器寻址)
  16. echarts markline X轴 Y轴 添加标识线
  17. (不打广告)推荐这款永久免费内网穿透软件-神卓互联
  18. Android Content Providers(三)——Contacts Provider
  19. 江苏,浙江一带的人有的那么富裕,他们大都是做什么的?
  20. FCoE协议详细解析

热门文章

  1. 论一只爬虫的自我修养11:Scrapy框架之初窥门径
  2. oracle excel更新,excel表格的数据如何更新-如何将excel表格更新oracle数据库中数据...
  3. viewUrl 终结者
  4. 六大手机软件商店分析报告
  5. Firefox扩展开发
  6. 手机wap网站制作教程
  7. 二、博客首页完成《iVX低代码仿CSDN个人博客制作》
  8. 画出典型计算机控制系统的方框图,计算机控制系统作业参考答案
  9. 一、从0开始——黑客学习路线
  10. bt种子磁力播放器 android,(安卓)种子磁力在线播放器 — 可离线至本地