▶ 本章介绍了纹理内存的使用，并给出了热传导的两个个例子。分别使用了一维和二维纹理单元。
● 热传导（使用一维纹理）
#include <stdio.h>
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include "D:\Code\CUDA\book\common\book.h"
#include "D:\Code\CUDA\book\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中
texture<float> texConstSrc;
texture<float> texIn;
texture<float> texOut; 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;
}; __global__ void blend_kernel(float *dst, bool dstOut)
{
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 - ;//找到上下左右的块
int right = offset + ;
int top = offset - DIM;
int bottom = offset + DIM;
if (x == )
left++;
if (x == DIM - )
right--;
if (y == )
top += DIM;
if (y == DIM - )
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 - * c); return;
} __global__ void copy_const_kernel(float *iptr)// 将恒温常量矩阵覆盖输入矩阵
{
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 != )
iptr[offset] = c; return;
} void anim_gpu(DataBlock *d, int ticks)
{
cudaEventRecord(d->start, );
dim3 blocks(DIM / , DIM / );
dim3 threads(, );
CPUAnimBitmap *bitmap = d->bitmap; volatile bool dstOut = true;//确定输入矩阵是哪一个，true代表dev_inSrc，false代表ev_outSrc
for (int i = ; i < ; i++)
{
float *in, *out;
if (dstOut)
{
in = d->dev_inSrc;
out = d->dev_outSrc;
}
else
{
in = d->dev_outSrc;
out = d->dev_inSrc;
} 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); cudaMemcpy(bitmap->get_ptr(), d->output_bitmap, bitmap->image_size(), cudaMemcpyDeviceToHost); cudaEventRecord(d->stop, );
cudaEventSynchronize(d->stop);
float elapsedTime;
cudaEventElapsedTime(&elapsedTime, d->start, d->stop);
d->totalTime += elapsedTime;
++d->frames;
printf("Average Time per frame: %3.1f ms\n", d->totalTime / d->frames);
} void anim_exit(DataBlock *d)// 收拾申请的内存
{
cudaUnbindTexture(texIn);
cudaUnbindTexture(texOut);
cudaUnbindTexture(texConstSrc);
cudaFree(d->dev_inSrc);
cudaFree(d->dev_outSrc);
cudaFree(d->dev_constSrc); cudaEventDestroy(d->start);
cudaEventDestroy(d->stop);
return;
} int main(void)
{
DataBlock data;
CPUAnimBitmap bitmap(DIM, DIM, &data);
data.bitmap = &bitmap;
data.totalTime = ;
data.frames = ;
cudaEventCreate(&data.start);
cudaEventCreate(&data.stop); int imageSize = bitmap.image_size(); cudaMalloc((void**)&data.output_bitmap, imageSize); cudaMalloc((void**)&data.dev_inSrc, imageSize);
cudaMalloc((void**)&data.dev_outSrc, imageSize);
cudaMalloc((void**)&data.dev_constSrc, imageSize);
cudaBindTexture(NULL, texConstSrc, data.dev_constSrc, imageSize);//将内存绑定到之前声明的纹理引用中去
cudaBindTexture(NULL, texIn, data.dev_inSrc, imageSize);
cudaBindTexture(NULL, texOut, data.dev_outSrc, imageSize); float *temp = (float*)malloc(imageSize);
for (int i = ; i < DIM*DIM; i++)// 恒温格点数据
{
temp[i] = ;
int x = i % DIM;
int y = i / DIM;
if ((x >= ) && (x < ) && (y >= ) && (y < ))
temp[i] = MAX_TEMP;
if ((x >= ) && (x < ) && (y >= ) && (y < ))
temp[i] = MIN_TEMP;
}
cudaMemcpy(data.dev_constSrc, temp, imageSize, cudaMemcpyHostToDevice); for (int i = ; i < DIM*DIM; i++)// 初始温度场数据
{
temp[i] = 0.5;
int x = i % DIM;
int y = i / DIM;
if ((x >= ) && (x < ) && (y >= ) && (y < ))
temp[i] = MAX_TEMP;
}
cudaMemcpy(data.dev_inSrc, temp, imageSize, cudaMemcpyHostToDevice); free(temp); bitmap.anim_and_exit((void(*)(void*, int))anim_gpu, (void(*)(void*))anim_exit); getchar();
return;
}
● 输出结果（左侧为恒高温，中间为恒低温，右侧为初始高温点）
● 使用一维纹理内存的过程浓缩一下就变成了以下过程
texture<float> texSrc;// 在全局位置上声明纹理引用 float *dev_Src;
cudaMalloc((void**)&dev_Src, sizeof(float)*DIM);// 申请和绑定纹理内存
cudaBindTexture(NULL, texSrc, dev_Src, NULL); float *temp = (float *)malloc(sizeof(float)*DIM);// 初始化该内存中的内容
//Initalize data in temp and then free(temp) cudaMemcpy(dev_Src, temp, sizeof(float)*DIM, cudaMemcpyHostToDevice); //Do something cudaUnbindTexture(texSrc);// 解绑和释放内存
cudaFree(dev_Src);
● 访问纹理内存不用中括号下标，而是
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(texSrc, offset);
● 热传导（使用二维纹理），输出结果同一维纹理的的情况，速度上没有明显差别
#include <stdio.h>
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include "D:\Code\CUDA\book\common\book.h"
#include "D:\Code\CUDA\book\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 texture<float, > texConstSrc;
texture<float, > texIn;
texture<float, > texOut; 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;
}; __global__ void blend_kernel(float *dst,bool dstOut)
{
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;
if (dstOut)//不需要自己处理边界情况
{
t = tex2D(texIn, x, y - );
l = tex2D(texIn, x - , y);
c = tex2D(texIn, x, y);
r = tex2D(texIn, x + , y);
b = tex2D(texIn, x, y + );
}
else
{
t = tex2D(texOut, x, y - );
l = tex2D(texOut, x - , y);
c = tex2D(texOut, x, y);
r = tex2D(texOut, x + , y);
b = tex2D(texOut, x, y + );
}
dst[offset] = c + SPEED * (t + b + r + l - * c); return;
} __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 != )
iptr[offset] = c; return;
} void anim_gpu(DataBlock *d, int ticks)
{
cudaEventRecord(d->start, );
dim3 blocks(DIM / , DIM / );
dim3 threads(, );
CPUAnimBitmap *bitmap = d->bitmap; volatile bool dstOut = true;
for (int i = ; i < ; 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); cudaMemcpy(bitmap->get_ptr(), d->output_bitmap, bitmap->image_size(), cudaMemcpyDeviceToHost); cudaEventRecord(d->stop, );
cudaEventSynchronize(d->stop); float elapsedTime;
cudaEventElapsedTime(&elapsedTime, d->start, d->stop);
d->totalTime += elapsedTime;
++d->frames;
printf("Average Time per frame: %3.1f ms\n", d->totalTime / d->frames); return;
} void anim_exit(DataBlock *d)
{
cudaUnbindTexture(texIn);
cudaUnbindTexture(texOut);
cudaUnbindTexture(texConstSrc);
cudaFree(d->dev_inSrc);
cudaFree(d->dev_outSrc);
cudaFree(d->dev_constSrc); cudaEventDestroy(d->start);
cudaEventDestroy(d->stop);
return;
} int main(void)
{
DataBlock data;
CPUAnimBitmap bitmap(DIM, DIM, &data);
data.bitmap = &bitmap;
data.totalTime = ;
data.frames = ;
cudaEventCreate(&data.start);
cudaEventCreate(&data.stop); int imageSize = bitmap.image_size(); cudaMalloc((void**)&data.output_bitmap, imageSize); cudaMalloc((void**)&data.dev_inSrc, imageSize);
cudaMalloc((void**)&data.dev_outSrc, imageSize);
cudaMalloc((void**)&data.dev_constSrc, imageSize); cudaChannelFormatDesc desc = cudaCreateChannelDesc<float>();
cudaBindTexture2D(NULL, texConstSrc, data.dev_constSrc, desc, DIM, DIM, sizeof(float) * DIM);
cudaBindTexture2D(NULL, texIn, data.dev_inSrc, desc, DIM, DIM, sizeof(float) * DIM);
cudaBindTexture2D(NULL, texOut, data.dev_outSrc, desc, DIM, DIM, sizeof(float) * DIM); float *temp = (float*)malloc(imageSize);
for (int i = ; i<DIM*DIM; i++) {
temp[i] = ;
int x = i % DIM;
int y = i / DIM;
if ((x >= ) && (x < ) && (y >= ) && (y < ))
temp[i] = MAX_TEMP;
if ((x >= ) && (x < ) && (y >= ) && (y < ))
temp[i] = MIN_TEMP;
}
cudaMemcpy(data.dev_constSrc, temp, imageSize, cudaMemcpyHostToDevice); for (int i = ; i < DIM*DIM; i++)// 初始温度场数据
{
temp[i] = 0.5;
int x = i % DIM;
int y = i / DIM;
if ((x >= ) && (x < ) && (y >= ) && (y < ))
temp[i] = MAX_TEMP;
}
cudaMemcpy(data.dev_inSrc, temp, imageSize, cudaMemcpyHostToDevice);
free(temp); bitmap.anim_and_exit((void(*)(void*, int))anim_gpu, (void(*)(void*))anim_exit); getchar();
return ;
}
● 使用纹理内存的过程浓缩一下就变成了以下过程
texture<float, > texSrc;// 在全局位置上声明纹理引用 float *dev_Src;
cudaMalloc((void**)&dev_Src, DIM*DIM);// 申请和绑定纹理内存
cudaChannelFormatDesc desc = cudaCreateChannelDesc<float>();
cudaBindTexture2D(NULL, texSrc, dev_Src, desc, DIM, DIM, sizeof(float) * DIM*DIM); float *temp = (float*)malloc(sizeof(float)*DIM*DIM);// 初始化该内存中的内容
//Initalize data in temp and then free(temp) cudaMemcpy(dev_Src, temp, sizeof(float)*DIM*DIM, cudaMemcpyHostToDevice); //Do something cudaUnbindTexture(texSrc);// 解绑和释放内存
cudaFree(dev_Src);
● 访问纹理内存不用中括号下标，而是
int x = threadIdx.x + blockIdx.x * blockDim.x;
int y = threadIdx.y + blockIdx.y * blockDim.y;
float c = tex2D(texSrc, x, y);