编程 / CUDA
15 分钟 读完 (大约 2232 个字)

cudanote4

第七章 纹理内存

纹理内存、纹理内存实现热传导模拟

GitHub

建议下载下来用Typora软件阅读markdown文件

作者github:littlebearsama 原文链接

(建议下载Typora来浏览markdown文件)

第七章 纹理内存

  • 纹理内存(Texture Memory)和常量内存一样，纹理内存是另外一种类型的只读内存，在特定的访问模式中，纹理内存同样能够提升性能并减少内存流量。
  • 虽然纹理内存最初是针对传统的图形处理应用程序而设计的，但在某些GPU计算应用程序中同样非常有用。
  • 与常量内存类似的是，纹理内存同样缓存在芯片上（利用了芯片上的缓存加速）！！！，因此在某些情况中，它能够减少对内存的请求并提供更高效的内存带宽。
  • 纹理缓存是专门为那些在内存访问模式中存在大量空间局部性（Spatial Locality）的图形应用程序而设计的。在某个计算应用程序中，这意味着一个线程读取的位置可能与邻近的线程的读取位置非常接近

img

上图中，从数学角度来看，图中的四个地址并非连续的，在一般的CPU缓存模式中，这些地址将不会缓存。但由于GPU纹理内存是专门为了加速这种访问模式而设计的，因此如果在这种情况中使用纹理内存而不是全局内存，那么将获得性能提升。

使用纹理内存实现热传导模拟

1.算法描述：

  1. 环境是一个矩形网格，在网格中随机散布一些”热源“，热源有着不同的固定温度（该点处的温度不会变）
  2. 在随时间递进的每个步骤中，我们假设热量在某个单元机器邻接单元之间”流动“/如果某个单元的温度比邻接单元的温度更高，那么热量将从邻接单元传导到该单元。
  3. 我们对新单元中心温度的计算方法为，将单元与邻接单元的温差相加起来，加上原有温度：

  1. 常量k表示模拟过程中热量的流动速率。k值越大，表示系统会更快地达到稳定温度，而k值越小，则温度梯度将存在更长时间。
  2. 只考虑上下左右四个邻域的话讲上述式子展开有

2.实现流程：

  1. 给定一个包含初始输入温度的网格，将其中作为热源的单元温度值复制到网格的相应单元中。这将覆盖这些单元之前计算出的温度，因此也就确保了”加热单元将保持恒温“这个条件。用下面代码中的copy_const_kernel()实现；
  2. 给定一个输入网格，用上面公式计算出输出网格。用下面代码中的blend_kernel()实现；
  3. 将输入网格和输出网格交换，为下一个计算步骤做好准备。当模拟下一个时间步时，在步骤2中计算得到的输出温度网格将成为步骤1中的输入温度网格。

3.代码：（使用的是二维纹理内存）

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
93
94
95
96
97
98
99
100
101
102
103
104
105
106
107
108
109
110
111
112
113
114
115
116
117
118
119
120
121
122
123
124
125
126
127
128
129
130
131
132
133
134
135
136
137
138
139
140
141
142
143
144
145
146
147
148
149
150
151
152
153
154
155
156
157
158
159
160
161
162
163
164
165
166
167
168
169
170
171
172
173
174
175
176
177
178
179
180
181
182
183
184
185
186
187
188
189
190
191
192
193
194
195
196
197
198
199
200
201
202
203
204
205
206
#include "cuda.h"
#include "../common/book.h"
#include "../common/image.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,2>  texConstSrc;
texture<float,2>  texIn;
texture<float,2>  texOut;

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

    float   t, l, c, r, b;
    //根据dstOut标志来看读取的输出部分的内存还是输出部分的内存
    if (dstOut) {              
        t = tex2D(texIn,x,y-1);
        l = tex2D(texIn,x-1,y);
        c = tex2D(texIn,x,y);
        r = tex2D(texIn,x+1,y);
        b = tex2D(texIn,x,y+1);
    } else {
        t = tex2D(texOut,x,y-1);
        l = tex2D(texOut,x-1,y);
        c = tex2D(texOut,x,y);
        r = tex2D(texOut,x+1,y);
        b = tex2D(texOut,x,y+1);
    }
    dst[offset] = c + SPEED * (t + b + r + l - 4 * c);
}

__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 = tex2D(texConstSrc,x,y);
    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;
    IMAGE           *bitmap;

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


// clean up memory allocated on the GPU
void cleanup( 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;
    IMAGE bitmap_image( DIM, DIM );
    data.bitmap = &bitmap_image;
    data.totalTime = 0;
    data.frames = 0;
    HANDLE_ERROR( cudaEventCreate( &data.start ) );
    HANDLE_ERROR( cudaEventCreate( &data.stop ) );

    int imageSize = bitmap_image.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 ) );
    //通道格式描述符
    cudaChannelFormatDesc desc = cudaCreateChannelDesc<float>();
    HANDLE_ERROR( cudaBindTexture2D( NULL, texConstSrc,
                                   data.dev_constSrc,
                                   desc, DIM, DIM,
                                   sizeof(float) * DIM ) );

    HANDLE_ERROR( cudaBindTexture2D( NULL, texIn,
                                   data.dev_inSrc,
                                   desc, DIM, DIM,
                                   sizeof(float) * DIM ) );

    HANDLE_ERROR( cudaBindTexture2D( NULL, texOut,
                                   data.dev_outSrc,
                                   desc, DIM, DIM,
                                   sizeof(float) * DIM ) );

    // initialize 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 );

    int ticks=0;
    bitmap_image.show_image(30);
    while(1)
    {
        HANDLE_ERROR( cudaEventRecord( data.start, 0 ) );
        dim3    blocks(DIM/16,DIM/16);
        dim3    threads(16,16);
        IMAGE  *bitmap = data.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  = data.dev_inSrc;
                out = data.dev_outSrc;
            } else {
                out = data.dev_inSrc;
                in  = data.dev_outSrc;
            }
            copy_const_kernel<<<blocks,threads>>>( in );
            blend_kernel<<<blocks,threads>>>( out, dstOut );
            dstOut = !dstOut;
        }
        float_to_color<<<blocks,threads>>>( data.output_bitmap,
                                            data.dev_inSrc );

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

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

        ticks++;
        char key = bitmap_image.show_image(30);
        if(key==27)
        {
            break;
        }
    }

    cleanup(&data);
    return 0;
}

4.代码解析（下面是使用一维纹理内存的解析）

  • 1.申请纹理内存：使用了浮点类型纹理内存的引用；纹理内存必须声明为文件作用域内的全局变量！

    1
    2
    3
    4
    //这些变量位于GPU上
    texture<float>  texConstSrc;
    texture<float>  texIn;
    texture<float>  texOut;

  • 2.申请GPU全局内存：下面代码为这三个缓存区分配了GPU内存（全局内存）,data.dev_inSrc等三个指针已经在结构对象data中声明了。

1
2
3
4
5
6
HANDLE_ERROR( cudaMalloc( (void**)&data.dev_inSrc,
                          imageSize ) );
HANDLE_ERROR( cudaMalloc( (void**)&data.dev_outSrc,
                          imageSize ) );
HANDLE_ERROR( cudaMalloc( (void**)&data.dev_constSrc,
                          imageSize ) );

  • 3.纹理内存与GPU全局内存绑定：需要通过cudaBindTexture()将这些变量（上面的纹理内存引用）绑定到内存缓冲区。相当于告诉CUDA运行时两件事情：

    a. 指定的缓冲区作为纹理来使用

    b.纹理引用作为纹理的”名字”

1
2
3
4
5
6
7
8
9
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 ) );
  • 4.使用内置函数tex1Dfetch()：当读取核函数中的纹理时，需要通过特殊的函数来告诉GPU将读取请求转发到纹理内存而不是标准的全局内存。tex1Dfetch()它是一个编译器内置函数（Instrinsic）。
  • 5.使用二维纹理内存：性能与一维的基本相同，但代码更简单。在使用内置函数tex2Dfetch()，读取缓存区中的数据时，不用计算缓存区中的线性偏置，而是可以直接用计算的像素索引x，y，这样使得代码更为简洁，并且能自动处理边界问题。
  • 6.通道格式描述符：在绑定二维纹理内存时，CUDA运行时要求提供一个cudaChanelFormatDesc()。在二维纹理内存的代码包含了一个对通道格式描述符的声明(Channel Format Descriptor)。在这里可以使用默认的参数，并且只要指定需要的是一个浮点描述符。然后我们通过1.cudaBindTexture2D(),2.纹理内存的位数（DIMXDIM）以及3.通道格式描述（desc）将这三个输入缓冲区绑定为二维纹理，main()函数其他部分保持不变。

纹理采样器（Texture Sampler），找不到该部分的内容？

如果使用了纹理采样器自动执行某种转换，那么纹理内存还能带来额外的加速。

# CUDA

喜欢这篇文章？打赏一下作者吧

cudanote5
cudanote3

评论

关注我

链接

分类

标签云

3D Reconstruction 3D scene Reconstruction 4PCS C++ CUDA CV-Source Deep Learning Despriptors Efficient Sparse ICP FastICP GOICP Generalized-ICP Generic Programming Hash Map ICP LM Memory Management NDT ObjectDetection OpenGL PCL Path Planning Performance PoseEstimation RANSAC Reconstruction Registration SAC-IA SFM SSE STL SVD Segmentation Sparse ICP Super4PCS books cs231n git glfw hexo keypoints mv-LM-ICP std surface Reconstruction tools 图像 编程知识与VS错误

最新文章

归档

标签

3D Reconstruction 3
3D scene Reconstruction 1
4PCS 1
C++ 13
CUDA 16
CV-Source 4
Deep Learning 3
Despriptors 1
Efficient Sparse ICP 1
FastICP 1
GOICP 1
Generalized-ICP 2
Generic Programming 1
Hash Map 1
ICP 7
LM 1
Memory Management 1
NDT 2
ObjectDetection 5
OpenGL 3
PCL 8
Path Planning 1
Performance 1
PoseEstimation 7
RANSAC 1
Reconstruction 5
Registration 11
SAC-IA 1
SFM 1
SSE 1
STL 2
SVD 1
Segmentation 1
Sparse ICP 1
Super4PCS 1
books 1
cs231n 1
git 2
glfw 2
hexo 2
keypoints 1
mv-LM-ICP 1
std 1
surface Reconstruction 1
tools 1
图像 2
编程知识与VS错误 1

最新文章

归档

标签

3D Reconstruction 3
3D scene Reconstruction 1
4PCS 1
C++ 13
CUDA 16
CV-Source 4
Deep Learning 3
Despriptors 1
Efficient Sparse ICP 1
FastICP 1
GOICP 1
Generalized-ICP 2
Generic Programming 1
Hash Map 1
ICP 7
LM 1
Memory Management 1
NDT 2
ObjectDetection 5
OpenGL 3
PCL 8
Path Planning 1
Performance 1
PoseEstimation 7
RANSAC 1
Reconstruction 5
Registration 11
SAC-IA 1
SFM 1
SSE 1
STL 2
SVD 1
Segmentation 1
Sparse ICP 1
Super4PCS 1
books 1
cs231n 1
git 2
glfw 2
hexo 2
keypoints 1
mv-LM-ICP 1
std 1
surface Reconstruction 1
tools 1
图像 2
编程知识与VS错误 1
Your browser is out-of-date!

Update your browser to view this website correctly. Update my browser now

×