详解cuda by example中第七章纹理内存（热量图）代码-CSDN博客

本文链接：https://blog.csdn.net/m0_57813464/article/details/141535210

前言

这两天学到第七章了，发现这几章代码对我来说都有点小复杂于是就把我对第七章的代码理解也写一下，仅供学习，如果大家有发现不对的地方欢迎指正。（文章代码顺序按照官方示例代码顺序给出，可按顺序食用）

cuda by example这本书及其代码下载链接如下https://github.com/yottaawesome/cuda-by-example

个人认为这段代码在纹理内存上的应用主要在blend_kernel，copy_const_kernel这两个函数上，因此我会对这两段代码进行详细注释

这部分代码不在赘述，定义了一些变量，引入头文件

#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

在gpu上分配三个纹理内存，texConstSrc存储发热源，texIn用来存储输入，texOut用来存储输出

texture<float>  texConstSrc;
texture<float>  texIn;
texture<float>  texOut;

blend_kernel函数

定义blend_kernel函数，其中有两个参数float*dst用来指定要修改数据的内存区域，在本文中分别是dev_inSrc与dev_outSrc，一个bool类型的dstOut，主要用来变换输入与输出（后面会详细解释这个变量的用处）

__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) {//dstout为true读取inSrc中的数据
        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);
}

下面详细拆分这段代码

1.x，y，offest表示的含义

threadIdx.x：表示线程的序号有x

blockIdx：表示线程块在线程格中的索引有x，y两个方向，可以类比为一个二维数组
blockDim.x 和 blockDim.y是线程块的尺寸，分别表示每个线程块在 x 方向和 y 方向上有多少个线程。

如下图一个中方块代表一个线程块里面有3*3个线程，最外面的大方块是一个线程格，里面有3*3个线程块。

    int x = threadIdx.x + blockIdx.x * blockDim.x;
    int y = threadIdx.y + blockIdx.y * blockDim.y;
    int offset = x + y * blockDim.x * gridDim.x;

我们举一个例子可以更好的理解下x,y,offset假设我们计算途中红色点实际在数组中的位置

(x是列，y是行)我们发现在线程块中threadIdx.x=0,threadIdx.y=1;在线程格中（中方块）blockIdx.x=0,blockIdx.y=1;block.Dim.x=3，block.Dim.y=3(每个线程块中x，y方向上有3个线程),gridDim.x=3(线程格中x，y方向上有三个线程块)
x=threadIdx.x + blockIdx.x * blockDim.x(0+0*3=0)
y=threadIdx.y + blockIdx.y * blockDim.y(1+1*3=4)
offset = x + y * blockDim.x * gridDim.x(0+4*3*3=36)即可算出偏移量

2.left，right，top，bottom表示的含义

    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;

根据红色坐标计算它的上下左右的坐标，过程跟算offset相似这里不再赘述

3.读取数据

    float   t, l, c, r, b;
    if (dstOut) {//dstout为true读取inSrc中的数据
        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);

tex1Dfetch()函数用来从纹理内存中读取数据前一个参数值要读取数据的内存地址（数组），后一个是序号，在获取数据后给dst数组中的offset序号赋值

copy_const_kernel函数

将texConstSrc中的热源像素的位置取出来，并赋值给iptr中的相应位置

__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;
}

数据块

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;
};

定义dev_inSrc,dev_outSrc,dev_constSrc,CPUAnimBitmap* bitmap分别在gpu上为输入，输出，热源常量以及位图分配指针（可以理解为分配起始地址，直白理解为数组也行）

anim_gpu函数

1.初始化线程块线程

设定线程块个数以及线程块里线程的个数，指明位图地址

dim3    blocks(DIM / 16, DIM / 16);
dim3    threads(16, 16);
CPUAnimBitmap* bitmap = d->bitmap;

2.理解 dstOut变量以及如何变换输入输出交替进行计算

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;//这一步相当于对调dev_inSrc与dev_outSrc，简化操作，否则还需要在全局内存上对dev_inSrc与dev_outSrc进行交换浪费时间
    }

这里主要理解dstOut这个变量的含义，我们知道我们为输入和输出分别分配了dev_inSrc与dev_outSrc两块gpu上的内存

当dstOut为true时in=dev_inSrc，out=dev_outSrc，此时调用copy_const_kernel函数相当于对dev_inSrc上的数据进行调用，blend_kernel函数对texIn上的数据
原因如下图 dstout为true执行if后的语句

此时float*dst等于out=dev_outSrc 原因如下图函数传递的参数

待blend_kernel函数执行完后dstout=！dstout此时dstout为False，第二次进入for循环
此时执行else后的代码，此时out=dev_inSrc，in=dev_outSrc，刚好将输入与输出对调，先前经过第一次for循环我们将第一次计算的输出放到了dev_outSrc中，执行copy_const_kernel（in）这个函数时此时in代指的是dev_outSrc内存的数据（也就是第一次计算的输出），然后执行blend_kernel函数(out, dstOut)此时out=dev_inSrc内存的数据，dstOut=false，将texOut中的数据当作输入，计算后输出到dev_inSrc中。

然后再进入for循环依次进行上面的过程