《CUDA By Example》【Chapter 07】纹理内存 ？

最新推荐文章于 2024-08-25 23:33:46 发布
最新推荐文章于 2024-08-25 23:33:46 发布
阅读量265 收藏
点赞数
分类专栏： AI GPU
版权声明：本文为博主原创文章，遵循 CC 4.0 BY-SA 版权协议，转载请附上原文出处链接和本声明。
本文链接：https://blog.csdn.net/wondervictor/article/details/80596156
版权
AI 同时被 2 个专栏收录
18 篇文章 1 订阅
订阅专栏
GPU
9 篇文章 0 订阅
订阅专栏

7.1 本章目标

纹理内存和常量内存一样，都是只读内存，用来减少带宽流量，提升性能。
1，了解纹理内存的特性
2，了解如何在CUDA C中使用一维纹理内存
3，了解如何在CUDA C中使用二维纹理内存

7.2 纹理内存简介（Texture Memory）

7.3 热传导模拟

heat.cu

#include "cuda.h"
#include "../common/book.h"
#include "../common/cpu_anim.h"

#define DIM 1024
#define PI 3.1415926535897932f
#define MAX_TEMP 1.0f
#define MIN_TEMP 0.0001f
#define SPEED   0.25f

// these exist on the GPU side
texture<float>  texConstSrc;
texture<float>  texIn;
texture<float>  texOut;



// this kernel takes in a 2-d array of floats
// it updates the value-of-interest by a scaled value based
// on itself and its nearest neighbors
__global__ void blend_kernel( float *dst,
                              bool dstOut ) {
    // map from threadIdx/BlockIdx to pixel position
    int 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++;
    if (x == DIM-1) right--; 

    int top = offset - DIM;
    int bottom = offset + DIM;
    if (y == 0)   top += DIM;
    if (y == DIM-1) bottom -= DIM;

    float   t, l, c, r, b;
    if (dstOut) {
        t = tex1Dfetch(texIn,top);
        l = tex1Dfetch(texIn,left);
        c = tex1Dfetch(texIn,offset);
        r = tex1Dfetch(texIn,right);
        b = tex1Dfetch(texIn,bottom);

    } else {
        t = tex1Dfetch(texOut,top);
        l = tex1Dfetch(texOut,left);
        c = tex1Dfetch(texOut,offset);
        r = tex1Dfetch(texOut,right);
        b = tex1Dfetch(texOut,bottom);
    }
    dst[offset] = c + SPEED * (t + b + r + l - 4 * c);
}

// NOTE - texOffsetConstSrc could either be passed as a
// parameter to this function, or passed in __constant__ memory
// if we declared it as a global above, it would be
// a parameter here: 
// __global__ void copy_const_kernel( float *iptr,
//                                    size_t texOffset )
__global__ void copy_const_kernel( float *iptr ) {
    // map from threadIdx/BlockIdx to pixel position
    int x = threadIdx.x + blockIdx.x * blockDim.x;
    int y = threadIdx.y + blockIdx.y * blockDim.y;
    int offset = x + y * blockDim.x * gridDim.x;

    float c = tex1Dfetch(texConstSrc,offset);
    if (c != 0)
        iptr[offset] = c;
}

// globals needed by the update routine
struct DataBlock {
    unsigned char   *output_bitmap;
    float           *dev_inSrc;
    float           *dev_outSrc;
    float           *dev_constSrc;
    CPUAnimBitmap  *bitmap;

    cudaEvent_t     start, stop;
    float           totalTime;
    float           frames;
};

void anim_gpu( DataBlock *d, int ticks ) {
    HANDLE_ERROR( cudaEventRecord( d->start, 0 ) );
    dim3    blocks(DIM/16,DIM/16);
    dim3    threads(16,16);
    CPUAnimBitmap  *bitmap = d->bitmap;

    // since tex is global and bound, we have to use a flag to
    // select which is in/out per iteration
    volatile bool dstOut = true;
    for (int i=0; i<90; i++) {
        float   *in, *out;
        if (dstOut) {
            in  = d->dev_inSrc;
            out = d->dev_outSrc;
        } else {
            out = d->dev_inSrc;
            in  = d->dev_outSrc;
        }
        copy_const_kernel<<<blocks,threads>>>( in );
        blend_kernel<<<blocks,threads>>>( out, dstOut );
        dstOut = !dstOut;
    }
    float_to_color<<<blocks,threads>>>( d->output_bitmap,
                                        d->dev_inSrc );

    HANDLE_ERROR( cudaMemcpy( bitmap->get_ptr(),
                              d->output_bitmap,
                              bitmap->image_size(),
                              cudaMemcpyDeviceToHost ) );

    HANDLE_ERROR( cudaEventRecord( d->stop, 0 ) );
    HANDLE_ERROR( cudaEventSynchronize( d->stop ) );
    float   elapsedTime;
    HANDLE_ERROR( cudaEventElapsedTime( &elapsedTime,
                                        d->start, d->stop ) );
    d->totalTime += elapsedTime;
    ++d->frames;
    printf( "Average Time per frame:  %3.1f ms\n",
            d->totalTime/d->frames  );
}

// clean up memory allocated on the GPU
void anim_exit( DataBlock *d ) {
    cudaUnbindTexture( texIn );
    cudaUnbindTexture( texOut );
    cudaUnbindTexture( texConstSrc );
    HANDLE_ERROR( cudaFree( d->dev_inSrc ) );
    HANDLE_ERROR( cudaFree( d->dev_outSrc ) );
    HANDLE_ERROR( cudaFree( d->dev_constSrc ) );

    HANDLE_ERROR( cudaEventDestroy( d->start ) );
    HANDLE_ERROR( cudaEventDestroy( d->stop ) );
}


int main( void ) {
    DataBlock   data;
    CPUAnimBitmap bitmap( DIM, DIM, &data );
    data.bitmap = &bitmap;
    data.totalTime = 0;
    data.frames = 0;
    HANDLE_ERROR( cudaEventCreate( &data.start ) );
    HANDLE_ERROR( cudaEventCreate( &data.stop ) );

    int imageSize = bitmap.image_size();

    HANDLE_ERROR( cudaMalloc( (void**)&data.output_bitmap,
                               imageSize ) );

    // assume float == 4 chars in size (ie rgba)
    HANDLE_ERROR( cudaMalloc( (void**)&data.dev_inSrc,
                              imageSize ) );
    HANDLE_ERROR( cudaMalloc( (void**)&data.dev_outSrc,
                              imageSize ) );
    HANDLE_ERROR( cudaMalloc( (void**)&data.dev_constSrc,
                              imageSize ) );

    HANDLE_ERROR( cudaBindTexture( NULL, texConstSrc,
                                   data.dev_constSrc,
                                   imageSize ) );

    HANDLE_ERROR( cudaBindTexture( NULL, texIn,
                                   data.dev_inSrc,
                                   imageSize ) );

    HANDLE_ERROR( cudaBindTexture( NULL, texOut,
                                   data.dev_outSrc,
                                   imageSize ) );

    // intialize the constant data
    float *temp = (float*)malloc( imageSize );
    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;
    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;
        }
    }
    HANDLE_ERROR( cudaMemcpy( data.dev_constSrc, temp,
                              imageSize,
                              cudaMemcpyHostToDevice ) );    

    // initialize the input data
    for (int y=800; y<DIM; y++) {
        for (int x=0; x<200; x++) {
            temp[x+y*DIM] = MAX_TEMP;
        }
    }
    HANDLE_ERROR( cudaMemcpy( data.dev_inSrc, temp,
                              imageSize,
                              cudaMemcpyHostToDevice ) );
    free( temp );

    bitmap.anim_and_exit( (void (*)(void*,int))anim_gpu,
                           (void (*)(void*))anim_exit );
}
一，声明数据类型为texture
// these exist on the GPU side
texture<float>  texConstSrc;
texture<float>  texIn;
texture<float>  texOut;

二，为三个缓冲区分配GPU内存
    // assume float == 4 chars in size (ie rgba)
    HANDLE_ERROR( cudaMalloc( (void**)&data.dev_inSrc,
                              imageSize ) );
    HANDLE_ERROR( cudaMalloc( (void**)&data.dev_outSrc,
                              imageSize ) );
    HANDLE_ERROR( cudaMalloc( (void**)&data.dev_constSrc,
                              imageSize ) );

三，通过cudaBindTexture()将这些变量绑定到内存缓冲区

 - 希望将制定的缓冲区作为纹理来使用
 - 希望将纹理引用作为纹理的“名字”

    HANDLE_ERROR( cudaBindTexture( NULL, texConstSrc,
                                   data.dev_constSrc,
                                   imageSize ) );

    HANDLE_ERROR( cudaBindTexture( NULL, texIn,
                                   data.dev_inSrc,
                                   imageSize ) );

    HANDLE_ERROR( cudaBindTexture( NULL, texOut,
                                   data.dev_outSrc,
                                   imageSize ) );

heat_2d.cu

#include "cuda.h"
#include "../common/book.h"
#include "../common/cpu_anim.h"

#define DIM 1024
#define PI 3.1415926535897932f
#define MAX_TEMP 1.0f
#define MIN_TEMP 0.0001f
#define SPEED   0.25f

// these exist on the GPU side
texture<float>  texConstSrc;
texture<float>  texIn;
texture<float>  texOut;



// this kernel takes in a 2-d array of floats
// it updates the value-of-interest by a scaled value based
// on itself and its nearest neighbors
__global__ void blend_kernel( float *dst,
                              bool dstOut ) {
    // map from threadIdx/BlockIdx to pixel position
    int 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++;
    if (x == DIM-1) right--; 

    int top = offset - DIM;
    int bottom = offset + DIM;
    if (y == 0)   top += DIM;
    if (y == DIM-1) bottom -= DIM;

    float   t, l, c, r, b;
    if (dstOut) {
        t = tex1Dfetch(texIn,top);
        l = tex1Dfetch(texIn,left);
        c = tex1Dfetch(texIn,offset);
        r = tex1Dfetch(texIn,right);
        b = tex1Dfetch(texIn,bottom);

    } else {
        t = tex1Dfetch(texOut,top);
        l = tex1Dfetch(texOut,left);
        c = tex1Dfetch(texOut,offset);
        r = tex1Dfetch(texOut,right);
        b = tex1Dfetch(texOut,bottom);
    }
    dst[offset] = c + SPEED * (t + b + r + l - 4 * c);
}

// NOTE - texOffsetConstSrc could either be passed as a
// parameter to this function, or passed in __constant__ memory
// if we declared it as a global above, it would be
// a parameter here: 
// __global__ void copy_const_kernel( float *iptr,
//                                    size_t texOffset )
__global__ void copy_const_kernel( float *iptr ) {
    // map from threadIdx/BlockIdx to pixel position
    int x = threadIdx.x + blockIdx.x * blockDim.x;
    int y = threadIdx.y + blockIdx.y * blockDim.y;
    int offset = x + y * blockDim.x * gridDim.x;

    float c = tex1Dfetch(texConstSrc,offset);
    if (c != 0)
        iptr[offset] = c;
}

// globals needed by the update routine
struct DataBlock {
    unsigned char   *output_bitmap;
    float           *dev_inSrc;
    float           *dev_outSrc;
    float           *dev_constSrc;
    CPUAnimBitmap  *bitmap;

    cudaEvent_t     start, stop;
    float           totalTime;
    float           frames;
};

void anim_gpu( DataBlock *d, int ticks ) {
    HANDLE_ERROR( cudaEventRecord( d->start, 0 ) );
    dim3    blocks(DIM/16,DIM/16);
    dim3    threads(16,16);
    CPUAnimBitmap  *bitmap = d->bitmap;

    // since tex is global and bound, we have to use a flag to
    // select which is in/out per iteration
    volatile bool dstOut = true;
    for (int i=0; i<90; i++) {
        float   *in, *out;
        if (dstOut) {
            in  = d->dev_inSrc;
            out = d->dev_outSrc;
        } else {
            out = d->dev_inSrc;
            in  = d->dev_outSrc;
        }
        copy_const_kernel<<<blocks,threads>>>( in );
        blend_kernel<<<blocks,threads>>>( out, dstOut );
        dstOut = !dstOut;
    }
    float_to_color<<<blocks,threads>>>( d->output_bitmap,
                                        d->dev_inSrc );

    HANDLE_ERROR( cudaMemcpy( bitmap->get_ptr(),
                              d->output_bitmap,
                              bitmap->image_size(),
                              cudaMemcpyDeviceToHost ) );

    HANDLE_ERROR( cudaEventRecord( d->stop, 0 ) );
    HANDLE_ERROR( cudaEventSynchronize( d->stop ) );
    float   elapsedTime;
    HANDLE_ERROR( cudaEventElapsedTime( &elapsedTime,
                                        d->start, d->stop ) );
    d->totalTime += elapsedTime;
    ++d->frames;
    printf( "Average Time per frame:  %3.1f ms\n",
            d->totalTime/d->frames  );
}

// clean up memory allocated on the GPU
void anim_exit( DataBlock *d ) {
    cudaUnbindTexture( texIn );
    cudaUnbindTexture( texOut );
    cudaUnbindTexture( texConstSrc );
    HANDLE_ERROR( cudaFree( d->dev_inSrc ) );
    HANDLE_ERROR( cudaFree( d->dev_outSrc ) );
    HANDLE_ERROR( cudaFree( d->dev_constSrc ) );

    HANDLE_ERROR( cudaEventDestroy( d->start ) );
    HANDLE_ERROR( cudaEventDestroy( d->stop ) );
}


int main( void ) {
    DataBlock   data;
    CPUAnimBitmap bitmap( DIM, DIM, &data );
    data.bitmap = &bitmap;
    data.totalTime = 0;
    data.frames = 0;
    HANDLE_ERROR( cudaEventCreate( &data.start ) );
    HANDLE_ERROR( cudaEventCreate( &data.stop ) );

    int imageSize = bitmap.image_size();

    HANDLE_ERROR( cudaMalloc( (void**)&data.output_bitmap,
                               imageSize ) );

    // assume float == 4 chars in size (ie rgba)
    HANDLE_ERROR( cudaMalloc( (void**)&data.dev_inSrc,
                              imageSize ) );
    HANDLE_ERROR( cudaMalloc( (void**)&data.dev_outSrc,
                              imageSize ) );
    HANDLE_ERROR( cudaMalloc( (void**)&data.dev_constSrc,
                              imageSize ) );

    HANDLE_ERROR( cudaBindTexture( NULL, texConstSrc,
                                   data.dev_constSrc,
                                   imageSize ) );

    HANDLE_ERROR( cudaBindTexture( NULL, texIn,
                                   data.dev_inSrc,
                                   imageSize ) );

    HANDLE_ERROR( cudaBindTexture( NULL, texOut,
                                   data.dev_outSrc,
                                   imageSize ) );

    // intialize the constant data
    float *temp = (float*)malloc( imageSize );
    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;
    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;
        }
    }
    HANDLE_ERROR( cudaMemcpy( data.dev_constSrc, temp,
                              imageSize,
                              cudaMemcpyHostToDevice ) );    

    // initialize the input data
    for (int y=800; y<DIM; y++) {
        for (int x=0; x<200; x++) {
            temp[x+y*DIM] = MAX_TEMP;
        }
    }
    HANDLE_ERROR( cudaMemcpy( data.dev_inSrc, temp,
                              imageSize,
                              cudaMemcpyHostToDevice ) );
    free( temp );

    bitmap.anim_and_exit( (void (*)(void*,int))anim_gpu,
                           (void (*)(void*))anim_exit );
}

纹理的维数（DIM x DIM）以及通道格式描述符（desc）将这三个输入缓冲区绑定为二维纹理。

常量内存和纹理内存的区别？？？

CUDA学习5 常量内存与纹理内存
https://www.cnblogs.com/qw12/p/6418857.html

wondervictor
关注
专栏目录
《GPU高性能编程》——cpu_anim.h
weixin_45699916的博客
11-01 567
《GPU 高性能编程 CUDA 实战》 书中 cpu_anim.h 头文件分享 《GPU高性能编程CUDA实战》是一本不错的好书，其中所含示例代码，经常包含有 "cpu_anim.h"等头文件, 为方便学习者查找，将头文件 “cpu_anim.h” 分享在这里。 /* * Copyright 1993-2010 NVIDIA Corporation. All rights reserved. ...
CUDA By Examples 8 - 纹理内存Texture Memory
Full_Speed_Turbo
06-28 1387
可以使用一维和二维纹理内存. 纹理内存是只读的.
CUDA by example Chapter7 纹理内存
wolfcsharp的博客
09-02 494
纹理内存纹理内存简介热传导模拟简单的传热模型代码中温度更新的计算GPU不使用纹理内存的热传导模拟计算GPU使用一维纹理内存的热传导模拟计算GPU使用二维纹理内存的热传导模拟计算 纹理内存简介 和常量内存一样，纹理内存是另一种类型的只读内存。纹理内存是专门为那些在内存访问模式中存在大量空间局部性的图形应用程序而设计的。假如在某个计算应用程序中，一个线程读取的位置与邻近线程读取的位置"非常接近"，使用...
详解cuda by example中第七章纹理内存（热量图）代码
最新发布
m0_57813464的博客
08-25 786
这两天学到第七章了，发现这几章代码对我来说都有点小复杂于是就把我对第七章的代码理解也写一下，仅供学习，如果大家有发现不对的地方欢迎指正。（文章代码顺序按照官方示例代码顺序给出，可按顺序食用）cuda by example这本书及其代码下载链接如下个人认为这段代码在纹理内存上的应用主要在blend_kernel，copy_const_kernel这两个函数上，因此我会对这两段代码进行详细注释这部分代码不在赘述，定义了一些变量，引入头文件。
Visual Studio中编译.cu +.cpp文件报错：C++ 未定义标识符 “clock64“....
weixin_43686259的博客
03-26 2971
Visual Studio中编译.cu +.cpp文件报错：C++ 未定义标识符 "clock64"....
CUDA_BY_EXAMPLE自带的几个头文件
09-27
CUDA_BY_EXAMPLE自带的几个头文件，包括 book.h cpu_anim.h cpu_bitmap.h gl_helper.h gpu_anim.h glext.h glut.h
CUDA by Example（英文原书+自带源代码）
01-02
《CUDA by Example》是一本专为想要学习GPU编程，尤其是CUDA技术的读者量身打造的实战指南。CUDA，全称Compute Unified Device Architecture，是NVIDIA公司推出的一种并行计算平台和编程模型，它允许开发者利用图形...
cuda by example.pdf
07-10
《CUDA by Example》是一本介绍通用目的图形处理单元（GPU）编程的教程书籍，作者是Jason Sanders和Edward Kandrot。本书不仅为读者提供了通过实例学习CUDA编程的机会，也是NVIDIA官方推荐的GPU学习教程。它深入浅出...
CUDA by Example (2010).pdf
10-29
由于给定文件的【标题】和【描述】均重复了"CUDA by Example (2010).pdf"，而【标签】是"CUDA by Example"，这表明文件的内容应该与这本书——《CUDA by Example》——紧密相关。该书由Jason Sanders和Edward ...
《CUDA By Example》中文译名《GPU高性能编程CUDA实战》是研究GPGPU异构并行计算非常不错的工具书。
04-04
《CUDA By Example》中文译名《GPU高性能编程CUDA实战》是研究GPGPU异构并行计算非常不错的工具书。书中给出的代码，非常个别的地方有失误，但是都有人为标注了，而且对不同的编程工具可能需要自己配置链接库。...
cuda by example中文版
04-10
3. **内存管理**：CUDA提供了多种类型的内存，如全局内存、共享内存、常量内存和纹理内存等。理解它们的区别和使用场景，能有效提升数据传输效率。 4. **流与同步**：CUDA流（stream）允许异步执行计算任务，提高...
GPU高性能编程CUDA实战—book.h CPUBitmap.h等头文件
04-09
GPU高性能编程CUDA实战—book.h CPUBitmap.h等头文件
《CUDA BY EXAMPLE》（GPU高性能编程CUDA实战）书中所有例子源码
03-29
《CUDA BY EXAMPLE》（GPU高性能编程CUDA实战）书中所有EXAMPLE的代码，包括书中提供的库。
CUDA By Example笔记--线程协作
牵一只蜗牛去散步
03-07 344
CUDA 架构确保除非线程块中的每个线程都执行了 __syncthreads()，否则没有线程能够执行 __syncthreads() 之后的指令，因此 __syncthreads() 不能位于线程发散的分支中（即只有满足条件的部分线程才执行指令的代码分支，并不是所有线程都会执行该分支的指令，这种分支称为发散分支）计算每一个线程在所有线程中对应的索引值，threadIdx.x表示线程的索引，blockIdx.x表示线程块的索引，blockDim.x表示一个线程块的x长度；
使用二维纹理内存模拟热传导
wufenxia的专栏
06-19 1608
#include #include "book.h" #include "cpu_anim.h" #define DIM 1024 #define PI 3.141592653f #define MAX_TEMP 1.0f #define MIN_TEMP 0.0001f #define SPEED 0.25f /*使用二维维纹理内存*/ //在GPU上定义纹理变量 text
 CPUAnimBitmap 类型的定义
wufenxia的专栏
06-14 1903
/*  * Copyright 1993-2010 NVIDIA Corporation.  All rights reserved.  *  * NVIDIA Corporation and its licensors retain all intellectual property and  * proprietary rights in and to this software an
CUDA中共享内存、常量内存和纹理内存的概念和应用（小白入门）
小P的博客
03-09 4528
刚接触CUDA，结合教材和资料整理了这篇共享内存、常量内存和纹理内存的基础内容，希望能帮到路过的你:P
NVIDIA CUDA 学习 (2) CUDA Parallel Programming
weixin_47658743的博客
09-06 274
CUDA Parallel ProgrammingCPU并发编程GPU并发编程尖括号的含义块索引 CPU并发编程 #include "../common/book.h" #define N 10 void add( int *a, int *b, int *c ) { int tid = 0; // this is CPU zero, so we start at zero while (tid < N) { c[tid] = a[tid] + b[tid]; tid += 1; //
CUDA-By-Example: Julia 修复包详解与应用
资源摘要信息:"CUDA-By-Example里的Julia修复包（patched）" 知识点概述： CUDA（Compute Unified Device Architecture）是由NVIDIA推出的一种通用并行计算架构。它允许开发者使用NVIDIA的GPU（图形处理器）来执行...
评论
添加红包

请填写红包祝福语或标题

个

红包个数最小为10个

元

红包金额最低5元

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

抵扣说明：

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

余额充值