CUDA优化实例（五）纹理内存与常量内存

最新推荐文章于 2024-08-07 10:51:33 发布
最新推荐文章于 2024-08-07 10:51:33 发布
阅读量1.3k 收藏 6
点赞数 1
分类专栏： CUDA 文章标签： 笔记
版权声明：本文为博主原创文章，遵循 CC 4.0 BY-SA 版权协议，转载请附上原文出处链接和本声明。
本文链接：https://blog.csdn.net/fb_help/article/details/79614999
版权

CUDA优化实例（五）纹理内存与常量内存

做了个卷积核的实验：
最基本情况：

结果：

这里写图片描述

使用常量内存：
这里写图片描述

纹理内存：
这里写图片描述

结果分析

没有使用texture memory 核函数花费80us，使用了纹理内存的，不管是线性内存还是CUDA Array都是140us

实例

源码：
CMakeLists.txt:

# CMakeLists.txt for G4CU project
project(test_cuda_project)
# required cmake version
cmake_minimum_required(VERSION 2.8)
# packages
find_package(CUDA)
# nvcc flags
set(CUDA_NVCC_FLAGS -arch=compute_60)
#set(CUDA_NVCC_FLAGS -gencode arch=compute_52,code=sm_52;-G;-g)

file(GLOB_RECURSE CURRENT_HEADERS  *.h *.hpp *.cuh)
file(GLOB CURRENT_SOURCES  *.cpp *.cu)
source_group("Include" FILES ${CURRENT_HEADERS}) 
source_group("Source" FILES ${CURRENT_SOURCES}) 



# Find OpenCV, you may need to set OpenCV_DIR variable
# to the absolute path to the directory containing OpenCVConfig.cmake file
# via the command line or GUI

set(OpenCV_DIR /home/jie/third_party/opencv-3.4.0/build)

find_package(OpenCV REQUIRED)



# If the package has been found, several variables will
# be set, you can find the full list with descriptions
# in the OpenCVConfig.cmake file.
# Print some message showing some of them
message(STATUS "OpenCV library status:")
message(STATUS "    version: ${OpenCV_VERSION}")
message(STATUS "    libraries: ${OpenCV_LIBS}")
message(STATUS "    include path: ${OpenCV_INCLUDE_DIRS}")

# Add OpenCV headers location to your include paths
include_directories(${OpenCV_INCLUDE_DIRS})


#CUDA_ADD_EXECUTABLE(test_cuda_project ${CURRENT_HEADERS} ${CURRENT_SOURCES})
CUDA_ADD_EXECUTABLE(test_cuda_project convolution.cu)


# Link your application with OpenCV libraries
target_link_libraries(test_cuda_project ${OpenCV_LIBS})

CUDA_ADD_EXECUTABLE(tex tex.cu)
# Link your application with OpenCV libraries
target_link_libraries(tex ${OpenCV_LIBS})

CUDA_ADD_EXECUTABLE(qq texture.cu)
# Link your application with OpenCV libraries
target_link_libraries(qq ${OpenCV_LIBS})

texture.cu:

#include <stdio.h>  
#include <iostream>
#include <stdlib.h>  
#include <cuda_runtime.h>  
#include "device_launch_parameters.h"
#include "opencv2/imgproc/imgproc.hpp"
#include "opencv2/highgui/highgui.hpp"
#include <time.h>
#include <math.h>

using namespace cv;


#define DEGRE_TO_RADIAN(x) ((x) * acos(-1) / 180)
#define CEIL(x,y) ((((x) + (y) - 1))/ (y) )
#define CHECK(res) { if(res != cudaSuccess){printf("Error ：%s:%d , ", __FILE__,__LINE__);   \
printf("code : %d , reason : %s \n", res,cudaGetErrorString(res));exit(-1);}}

//texture<uchar, cudaTextureType2D, cudaReadModeElementType> texRef;
texture<uchar, 1, cudaReadModeElementType> texRef;

texture<uchar, 2, cudaReadModeElementType> texRef2D;

__constant__ int dev_H[9];

__global__ void kernel(uchar* dout_data,uchar *d_data,int width,int height)
{
    int x = threadIdx.x + blockIdx.x*blockDim.x;
    int y = threadIdx.y + blockIdx.y*blockDim.y;
    int coreIndex = y*width+x;

    if (x > 0 && x < width - 1 && y >0 && y < height - 1)
    {
        // dout_data[coreIndex] = (
        //  d_data[coreIndex - 1] * 2  + d_data[coreIndex] *4 + d_data[coreIndex+1]*2+
        //  d_data[coreIndex - width - 1] + d_data[coreIndex - width ]*2 + d_data[coreIndex - width + 1]+
        //  d_data[coreIndex + width - 1] + d_data[coreIndex + width] *2 + d_data[coreIndex+ width + 1] )/16 ;


        dout_data[coreIndex] =(
            d_data[coreIndex - width - 1]* 1  + d_data[coreIndex - width] *  2 + d_data[coreIndex - width + 1] * 1+
            d_data[coreIndex - 1] *        2  + d_data[coreIndex        ] *  4   + d_data[coreIndex+1        ] * 2+
            d_data[coreIndex + width - 1]* 1  + d_data[coreIndex + width] *  2 + d_data[coreIndex + width + 1] * 1 
        )/16;

    }

}
__global__ void kernelUseCon(uchar* dout_data,uchar *d_data,int width,int height)
{
    int x = threadIdx.x + blockIdx.x*blockDim.x;
    int y = threadIdx.y + blockIdx.y*blockDim.y;
    int coreIndex = y*width+x;

    if (x > 0 && x < width - 1 && y >0 && y < height - 1)
    {           
        dout_data[coreIndex] =(
            d_data[coreIndex - width - 1]* dev_H[0]  + d_data[coreIndex - width] *  dev_H[1] + d_data[coreIndex - width + 1] * dev_H[2]+
            d_data[coreIndex - 1] *        dev_H[3]  + d_data[coreIndex        ] *  dev_H[4]   + d_data[coreIndex+1        ] * dev_H[5]+
            d_data[coreIndex + width - 1]* dev_H[6]  + d_data[coreIndex + width] *  dev_H[7] + d_data[coreIndex + width + 1] * dev_H[8] 
        )/16;

    }

}
__global__ void kernelUseTex(uchar* dout_data,int width,int height)
{
    int x = threadIdx.x + blockIdx.x*blockDim.x;
    int y = threadIdx.y + blockIdx.y*blockDim.y;
    int coreIndex = y*width+x;

    if (x > 0 && x < width - 1 && y >0 && y < height - 1)
    {       
        dout_data[coreIndex] =(
            tex1Dfetch(texRef,coreIndex - width - 1)* 1  + tex1Dfetch(texRef,coreIndex - width) *  2 + tex1Dfetch(texRef,coreIndex - width + 1) * 1+
            tex1Dfetch(texRef,coreIndex - 1) *        2  + tex1Dfetch(texRef,coreIndex        ) *  4 + tex1Dfetch(texRef,coreIndex+1          ) * 2+
            tex1Dfetch(texRef,coreIndex + width - 1)* 1  + tex1Dfetch(texRef,coreIndex + width) *  2 + tex1Dfetch(texRef,coreIndex + width + 1) * 1 
        )/16;

    }

}
__global__ void kernelUseTex2D(uchar* dout_data,int width,int height)
{
    int x = threadIdx.x + blockIdx.x*blockDim.x;
    int y = threadIdx.y + blockIdx.y*blockDim.y;
    int coreIndex = y*width+x;

    if (x > 0 && x < width - 1 && y >0 && y < height - 1)
    {       
        dout_data[coreIndex] =(
            tex2D(texRef2D,x-1,y-1)* 1  + tex2D(texRef2D,x,y-1) *  2 + tex2D(texRef2D,x+1,y-1) * 1+
            tex2D(texRef2D,x-1,y)  * 2  + tex2D(texRef2D,x,y)   *  4 + tex2D(texRef2D,x+1,y)   * 2+
            tex2D(texRef2D,x-1,y+1)* 1  + tex2D(texRef2D,x,y+1) *  2 + tex2D(texRef2D,x+1,y+1) * 1 
        )/16;

    }

}
int main()
{
    Mat src = imread("../road.png",CV_LOAD_IMAGE_GRAYSCALE);
    Mat dst;
    dst.create(src.size(), src.type());
    int height = src.size().height;
    int width = src.size().width;
    int size = width*height*sizeof(uchar);
    uchar *d_data,*dout_data;

    if (src.empty())
    {
        fprintf(stderr, "Can not load image %s\n", "road.png");
        return -1;
    }

    /*  common */
    CHECK(cudaMalloc((void **)&d_data, size));
    CHECK(cudaMalloc((void **)&dout_data, size));
    CHECK(cudaMemcpy(d_data, src.data, size, cudaMemcpyHostToDevice));

    dim3 dimBlock(32, 4);
    dim3 dimGrid(CEIL(width, dimBlock.x), CEIL(height, dimBlock.y));

    kernel<<<dimGrid,dimBlock>>>(dout_data,d_data,width,height);
    CHECK(cudaDeviceSynchronize());


    CHECK(cudaMemcpy(dst.data, dout_data, size, cudaMemcpyDeviceToHost));
    /*  common */

    /*  use constance  */
    // int H[9];
    // H[0]=1;H[1]=2;H[2]=1;
    // H[3]=2;H[4]=4;H[5]=2;
    // H[6]=1;H[7]=2;H[8]=1;

    // cudaMemcpyToSymbol(dev_H,H,sizeof(int)*9);

    // CHECK(cudaMalloc((void **)&d_data, size));
    // CHECK(cudaMalloc((void **)&dout_data, size));
    // CHECK(cudaMemcpy(d_data, src.data, size, cudaMemcpyHostToDevice));

    // dim3 dimBlock(32, 4);
    // dim3 dimGrid(CEIL(width, dimBlock.x), CEIL(height, dimBlock.y));

    // kernelUseCon<<<dimGrid,dimBlock>>>(dout_data,d_data,width,height);
    // CHECK(cudaDeviceSynchronize());


    // CHECK(cudaMemcpy(dst.data, dout_data, size, cudaMemcpyDeviceToHost));
    /*  use constance  */


    /*  use texture Linar 1D */
    // CHECK(cudaMalloc((void **)&d_data, size));
    // CHECK(cudaMalloc((void **)&dout_data, size));
    // CHECK(cudaMemcpy(d_data, src.data, size, cudaMemcpyHostToDevice));

    // CHECK(cudaBindTexture(NULL,texRef,d_data,size));

    // dim3 dimBlock(32, 4);
    // dim3 dimGrid(CEIL(width, dimBlock.x), CEIL(height, dimBlock.y));

    // kernelUseTex<<<dimGrid,dimBlock>>>(dout_data,width,height);
    // CHECK(cudaDeviceSynchronize());


    // CHECK(cudaMemcpy(dst.data, dout_data, size, cudaMemcpyDeviceToHost));
    /*  use texture Linar 1D */

    /*  use texture CUDA Array 2D */
    // CHECK(cudaMalloc((void **)&dout_data, size));

    // cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc<uchar>();
    // cudaArray* cuArray;
    // CHECK(cudaMallocArray(&cuArray, &channelDesc, width, height));
    // CHECK(cudaMemcpyToArray(cuArray,0,0,src.data,size,cudaMemcpyHostToDevice));


    // CHECK(cudaBindTextureToArray(texRef2D,cuArray,channelDesc));

    // dim3 dimBlock(32, 4);
    // dim3 dimGrid(CEIL(width, dimBlock.x), CEIL(height, dimBlock.y));

    // kernelUseTex2D<<<dimGrid,dimBlock>>>(dout_data,width,height);
    // CHECK(cudaDeviceSynchronize());


    // CHECK(cudaMemcpy(dst.data, dout_data, size, cudaMemcpyDeviceToHost));
    /*  use texture CUDA Array 2D */




    imshow("Linear Blend",dst);
    waitKey(0);



    return 0;
}

分析：

常量内存不怎么影响程序的性能，但它可以减少本地内存的使用，有利于占用率的提升。
纹理内存没有增加系统的性能，反而降级，原因可能是此例的内存访问，在线程束中是合并的，且独立的内存请求较多（类似展开）。

fb_help
关注
专栏目录
第八章 CUDA内存应用与性能优化篇(上篇)
weixin_38252409的博客
08-13 2291
学习我的教程专栏，你将绝对能实现CUDA工程化，实现环境安装、index计算、kernel核函数编程、内存优化与steam性能优化、原子操作、nms的cuda算子、yolov5的cuda部署等内容，并开源教程源码。以上章节中，我们已经比较熟练掌握如何使用cuda编写自己想要的计算逻辑，已能成功编写cuda代码了。那么，另外一个重要问题值得我们关注，如何优化其性能，使其工程部署能加速运行了。而这种性能优化与cuda内存密切相关。为此，我们在本节中介绍cuda内存相关内容，并附其源码。
第八章 CUDA内存应用与性能优化篇(下篇)
weixin_38252409的博客
08-13 772
学习我的教程专栏，你将绝对能实现CUDA工程化，实现环境安装、index计算、kernel核函数编程、内存优化与steam性能优化、原子操作、nms的cuda算子、yolov5的cuda部署等内容，并开源教程源码。以上章节中，我们已经比较熟练掌握如何使用cuda编写自己想要的计算逻辑，已能成功编写cuda代码了。那么，另外一个重要问题值得我们关注，如何优化其性能，使其工程部署能加速运行了。而这种性能优化与cuda内存密切相关。为此，我们在本节中介绍cuda内存相关内容，并附其源码。
CUDA纹理内存入门详解
hgkjbm的博客
03-31 3181
CUDA 进阶技巧 纹理内存 入门指南
Cuda核函数优化方法
最新发布
千人斩的博客
08-07 389
Cuda核函数优化方法
VS2017 CUDA编程学习10：纹理内存
DU_YULIN的博客
11-03 1214
文章目录前言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：线程同步-共享内
CUDA一维纹理内存
weixin_30467087的博客
02-06 199
纹理一词来源于GPU图形世界，GPU通用并行计算“盗用”了纹理一词，定义了一个纹理内存的概念。纹理内存缓存在 设备上，在某些情况下能减少对内存的请求并降低内存带宽的使用，是专门为那些在内存访问模式中存在大量空间局部性的图形应用而设计，意味着一个线程读取的位置可能与邻近线程读取的位置“非常接近”。对于GPU内核而言，纹理内存是只读内存，并且只有通过特殊的纹理API...
CUDA实例系列五:In-Place矩阵转置(含源码)
专注于人工智能领域的小何尚
05-23 610
CUDA优化实例: In-Place矩阵转置 所谓In-Place矩阵转置, 简单点说就是我们在同一个内存空间内做矩阵转置的操作.这也是很多计算需要的步骤之一, 这样操作可以有效减少对存储单元的需求. 利用CUDA优化矩阵转置,可以参考我之前的blog: https://blog.csdn.net/kunhe0512/article/details/124884614 而In-Place矩阵转置简单点说, 我们可以使用一半的线程来完成任务: 这里使用一半的线程的意义在于即使使用全部线程也不能增加速度,
CUDA优化实例（四）纹理内存
fb_help的专栏
03-15 2000
CUDA优化实例（四）纹理内存 本文参考：http://www.cnblogs.com/cuancuancuanhao/p/7809713.html 本节的内容可能和标题不服，本节主要将纹理内存的使用，它到底有什么有什么速度的提升，优化体现在哪里，我下节会写，本节主要写一个纹理内存的例子。 引言 纹理内存与全局内存一样，都在DRAM上，所以纹理内存的容量是很大的 纹理内存是专为图像而...
CUDA——纹理内存（纹理参考）
weixin_44444450的博客
02-09 1285
纹理内存 1.1纹理类型 CUDA支持GPU上的一部分纹理硬件（它们原本是为图形处理而设计的），所以，纹理内存一般会用于加速图片的滤波和图片的resize操作上。 纹理内存实质上市全局内存的特殊形态，全局内存被绑定为纹理内存，对其的读，写操作将通过专门的texture cache（纹理缓存）进行。 纹理缓存的优势：纹理缓存具备硬件插值特性，可以实现最近邻插值和线性插值。纹理缓存访问二维矩阵的领域会...
CUDA纹理内存与常量内存在并行计算中的应用
CUDA内存模型包括全局内存、共享内存、寄存器和常量/纹理内存。其中，全局内存是所有线程都可以访问的内存，共享内存用于线程块内部的共享数据访问，寄存器用于存储线程私有数据，而常量/纹理内存则是特殊的只读内存...
第八章 CUDA内存应用与性能优化篇(中篇)
weixin_38252409的博客
08-13 1396
学习我的教程专栏，你将绝对能实现CUDA工程化，实现环境安装、index计算、kernel核函数编程、内存优化与steam性能优化、原子操作、nms的cuda算子、yolov5的cuda部署等内容，并开源教程源码。以上章节中，我们已经比较熟练掌握如何使用cuda编写自己想要的计算逻辑，已能成功编写cuda代码了。那么，另外一个重要问题值得我们关注，如何优化其性能，使其工程部署能加速运行了。而这种性能优化与cuda内存密切相关。为此，我们在本节中介绍cuda内存相关内容，并附其源码。
CUDA（Ⅷ）：纹理内存（Texture Memory）
每天进步一点点！
04-27 6682
1.学习目标： 了解纹理内存的性能特性 了解如何在CUDA C中使用一维纹理内存 了解如何在CUDA C中使用二维纹理内存 2.背景提出： 特定环境中使用特殊的内存将极大地提升应用程序的性能。与Constant Memory相似，Texture Memory同样缓存在芯片上，只读内存，在某些情况下，能减少对内存的请求并提供更高效的内存带宽。Texture Mmeory是专门为那些在内存访...
CUDA总结：纹理内存
热门推荐
Kelvin_Yan的专栏
01-04 1万+
来源：《CUDA_Runtime_API》、《CUDA_C_Programming_Guide-V8.0》 纹理内存和表面内存（surface memory）实质上是全局内存的一个特殊形态，全局内存被绑定为纹理内存（表面内存），对其的读（写）操作将通过专门的texture cache（纹理缓存）进行，其实称为纹理缓存更加贴切。 纹理缓存的优势：纹理缓存具备硬件插值特性，可以实现最近邻插值和线性
CUDA内存(五) 纹理内存
Full_Speed_Turbo
02-18 956
文章目录纹理内存的特性API最简单的用法：例子分析：Simpletexture1. 声明CUDA数组之前，必须先用结构体channelDesc描述CUDA数组中的数据类型。2. 声明纹理参照系3.设置运行时纹理参照系属性4. 纹理绑定5 纹理拾取 纹理内存的特性 纹理内存缓存在芯片上，因此在某些情况中，它能够减少对内存的请求并提供更高效的内存带宽。纹理缓存是专门为那些在内存访问模式中存在大量空间局...
CUDA的纹理内存讲解
Dezeming的博客
03-28 7096
CUDA的纹理内存讲解，来自翻译和补充官方文档 CUDA programming Guide Texture and Surface Memory CUDA支持供GPU处理图形用的纹理硬件和表面（surface）内存，但是只支持一部分功能。从纹理和表面内存中读取数据可以得到很多好处（以后再说）。 有两个不同的API可以来获取纹理内存和表面内存：1.纹理参考API，支持所有设备。2.纹理对象...
CUDA 学习——纹理内存（CUDA实战 第七章）
ygf666的博客
05-18 2704
1. 纹理内存使用 纹理内存的使用依赖于API函数。下面直接给出常见的使用流程： 1.1 声明纹理变量 texture<Type, Dim, ReadMode> VarName; //Type: 前面提到的基本的整型和浮点类型，以及其它的对齐类型 //Dim： 纹理数组的维度，值为1或2或3，默认缺省为1 //ReadMode：cudaReadModelNormalizedFloat 或 cudaReadModelElementType(默认) NOTE： cudaReadModelNor
CUDA By Example(六)——纹理内存
hhhcbw的博客，欢迎各位来访
11-17 1109
CUDA C中纹理内存使用的说明，并以热传导模拟为例
CUDA 学习（十四）、纹理内存
10-27 1318
一、概述         使用纹理内存主要基于以下两个目的：（1）计算能力为1.x 与3.x 硬件上的缓存机制；（2）基于硬件的内存读取操作。 二、纹理缓存        由于在计算能力为1.x 的硬件上根本没有高速缓存，因此，每个SM上的 6-8KB 的纹理内存为此设备提供了唯一真正缓存数据的方法。然而，随着费米架构的硬件出现，一级缓存最大可达到48KB，共享二级缓存最大可达到76
CUDA编程指南：内存管理与多线程分配
- **存储器层次**：包括全局内存、共享内存、常量内存和纹理内存等，每种内存类型有其特定的访问速度和使用场景。 - **异构编程**：CUDA允许在CPU和GPU之间进行混合编程，开发者可以利用CPU进行管理和初始化，而将...
评论
添加红包

请填写红包祝福语或标题

个

红包个数最小为10个

元

红包金额最低5元

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

抵扣说明：

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

余额充值