CUDA by example Chapter7 纹理内存

最新推荐文章于 2022-11-17 19:57:01 发布
最新推荐文章于 2022-11-17 19:57:01 发布
阅读量414 收藏 5
点赞数
分类专栏： CUDA 文章标签： cuda by example 纹理内存
版权声明：本文为博主原创文章，遵循 CC 4.0 BY-SA 版权协议，转载请附上原文出处链接和本声明。
本文链接：https://blog.csdn.net/wolfcsharp/article/details/100301221
版权

纹理内存

纹理内存简介

和常量内存一样，纹理内存是另一种类型的只读内存。纹理内存是专门为那些在内存访问模式中存在大量空间局部性的图形应用程序而设计的。假如在某个计算应用程序中，一个线程读取的位置与邻近线程读取的位置"非常接近"，使用纹理内存将会减少内存通信量，从而提高性能。
在这里插入图片描述

热传导模拟

简单的传热模型

假设一个矩形网格，热源分布如下。假设热源单元本身的温度保持不变，热量可以在相邻单元之间流动，从高温单元传导到低温单元。
在这里插入图片描述
温度更新的计算方法：
将单元与其邻接单元的温差加起来，然后加上原有温度，等于新时刻的温度：
TNEW = TOLD+sum_NEIGHBOR(k(TNEIGHBOR-TOLD))
由于邻接单元只有上下左右四个，将上式化简，得到：
TNEW=TOLD+k(TTOP+TBOTTOm+TLEFT+TRIGHT-4TOLD)

代码中温度更新的计算

更新流程：

  • 给定包含初始温度的网格，它的大部分单元都是0，少部分是初始温度值。将其中作为热源的初始温度值复制到当前时间的网格对应单元中。确保"加热单元将保持恒温"，这个复制操作时在copy_const_kernal()中执行的。
  • 给定一个输入温度网格，根据温度更新的公式计算输出温度网格。这个更新操作是在blend_kernal()中执行的。
  • 将输入温度网格和输出温度网格交换，为下一个步骤的计算做好准备。当模拟下一个时间步时，步骤2中计算得到的输出温度网格将成为步骤1中的输入温度网格。

GPU不使用纹理内存的热传导模拟计算

完整代码

// 代码7.3不使用纹理内存的热传导模型计算
//时间：2019.07.28
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <iostream>
#include "cpu_anim.h"
#include "book.h"

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



//更新函数中需要的全局变量
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 copy_const_kernal(float *iptr, const float *cptr)
{
	//将threadIdx/blockIdx映射到像素位置
	int x = threadIdx.x + blockIdx.x*blockDim.x;
	int y = threadIdx.y + blockIdx.y*blockDim.y;
	int offset = x + y*blockDim.x*gridDim.x;

	if (cptr[offset] != 0)iptr[offset] = cptr[offset];

}

__global__ void blend_kernal(float *outSrc, const float *inSrc)
{
	//将threadIdx/blockIdx映射到像素位置
	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;

	outSrc[offset] = inSrc[offset] + SPEED*(inSrc[top] + inSrc[bottom] + inSrc[left] + inSrc[right] - inSrc[offset] * 4);

}


//在swap这里栽了两天的跟头
void swap(float **a, float **b)
{
	float *c;
	c = *a;
	*a = *b;
	*b = c;
}



void anim_gpu(DataBlock *d, int ticks)
{
	cudaEventRecord(d->start, 0);
	dim3 grids(DIM / 16, DIM / 16);
	dim3 blocks(16, 16);
	CPUAnimBitmap *bitmap = d->bitmap;
	float *c;
	for (int i = 0; i < 90; i++)
	{
		//printf("faf");
		copy_const_kernal << <grids, blocks >> > ( d->dev_inSrc, d->dev_constSrc);
		blend_kernal << <grids, blocks >> > ( d->dev_outSrc, d->dev_inSrc );

		swap(&d->dev_inSrc,&d->dev_outSrc);
	}
	//将float数值映射成颜色，以便用图像显示它
	float_to_color << <grids, blocks >> >(d->output_bitmap,d->dev_inSrc);
	cudaMemcpy(bitmap->get_ptr(), d->output_bitmap, bitmap->image_size(), cudaMemcpyDeviceToHost);
	cudaEventRecord(d->stop,0);
	cudaEventSynchronize(d->stop);
	float elapsedTime;
	cudaEventElapsedTime(&elapsedTime, d->start, d->stop);
	d->totalTime += elapsedTime;
	++d->frames;
	printf("Average Time per frame:%3.1fms\n", d->totalTime / d->frames);
}

void anim_exit(DataBlock *d)
{
	cudaFree(d->dev_constSrc);
	cudaFree(d->dev_inSrc);
	cudaFree(d->dev_outSrc);

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

int main()
{
	DataBlock data;
	CPUAnimBitmap bitmap(DIM, DIM, &data);
	data.bitmap = &bitmap;
	data.totalTime = 0;
	data.frames = 0;
	cudaEventCreate(&data.start);
	cudaEventCreate(&data.stop);
	cudaMalloc((void **)&data.output_bitmap, bitmap.image_size());
	cudaMalloc((void**)&data.dev_constSrc, bitmap.image_size());
	cudaMalloc((void **)&data.dev_inSrc, bitmap.image_size());
	cudaMalloc((void **)&data.dev_outSrc, bitmap.image_size());

	//在CPU上分配临时内存，对其填充一些内容，用于初始化data.dev_constSrc和data.dev_inSrc
	float *temp = (float *)malloc(bitmap.image_size());
	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;
		}
	}
	//用当前填充的内容初始化data.dev_constSrc
	cudaMemcpy(data.dev_constSrc, temp, bitmap.image_size(), cudaMemcpyHostToDevice);

	//继续向temp中填充一些内容
	for (int y = 800; y < DIM; y++)
	{
		for (int x = 0; x < 200; x++)
		{
			temp[x + y*DIM] = MAX_TEMP;
		}
	}
	//用给当前填充的内容初始化data.dev_inSrc
	cudaMemcpy(data.dev_inSrc, temp, bitmap.image_size(), cudaMemcpyHostToDevice);

	//释放CPU端的临时内存
	free(temp);
	//用anim_gpu生成动画的每一帧，用anim_exit退出动画显示并释放内存
	bitmap.anim_and_exit((void(*) (void *, int))anim_gpu, (void(*)(void *))anim_exit);
}

运行结果
在这里插入图片描述
在这里插入图片描述
需要极其注意，更新步骤中的第三步，交换输入网格和输出网格：
因为dev_inSrc和dev_outSrc分别是两幅图像的入口指针，将二者交换只需要：

float *swap;
swap = dev_inSrc;
dev_inSrc = dev_outSrc;
dev_outSrc = swap;

但是，如果使用函数swap来完成这个工作，由于函数中传入的是形参，形参的改变并不会改变传入参数的原始值，因此我们需要将&dev_inSrc和&dev_outSrc传入swap函数中，才能完成交换任务

//在swap这里栽了两天的跟头
void swap(float **a, float **b)
{
	float *c;
	c = *a;
	*a = *b;
	*b = c;
}

//使用方法
swap(&d->dev_inSrc,&d->dev_outSrc);

GPU使用一维纹理内存的热传导模拟计算

1. 纹理内存的使用场景
在温度更新计算的内存访问模式中存在着巨大的内存空间局部性，这种访问模式可以通过GPU纹理内存来加速。
2.一维纹理内存变量引用的声明

//声明纹理内存变量的引用，这些变量将位于GPU上
texture<float> texConstSrc;
texture<float> texIn;
texture<float> texOut;

3.一维纹理内存引用变量与实际内存缓冲区的绑定
通过cudaBindTexture()可以将纹理内存引用变量与内存缓冲区进行绑定。

	cudaMalloc((void **)&data.output_bitmap, bitmap.image_size());
	cudaMalloc((void**)&data.dev_constSrc, bitmap.image_size());
	cudaMalloc((void **)&data.dev_inSrc, bitmap.image_size());
	cudaMalloc((void **)&data.dev_outSrc, bitmap.image_size());

	//分配内存之后，需要将这些变量绑定到之前声明的纹理引用上去
	cudaBindTexture(NULL, texConstSrc, data.dev_constSrc, bitmap.image_size());
	cudaBindTexture(NULL, texIn, data.dev_inSrc,  bitmap.image_size());
	cudaBindTexture(NULL, texOut, data.dev_outSrc,  bitmap.image_size());

4.读取一维纹理内存方法
使用tex1Dfetch()读取一维纹理内存,传给blend_kernal的dstOut是一个bool变量，用于指明读取哪一个缓冲区作为输入。

__global__ void blend_kernal(float *dst, bool dstOut)
{
	//将threadIdx/blockIdx映射到像素位置
	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, b, l, r, c;
	if (dstOut)
	{
		t = tex1Dfetch(texIn,top);
		b = tex1Dfetch(texIn, bottom);
		l = tex1Dfetch(texIn, left);
		r = tex1Dfetch(texIn, right);
		c = tex1Dfetch(texIn, offset);
	}
	else
	{
		t = tex1Dfetch(texOut, top);
		b = tex1Dfetch(texOut, bottom);
		l = tex1Dfetch(texOut, left);
		r = tex1Dfetch(texOut, right);
		c = tex1Dfetch(texOut, offset);
	}

	dst[offset] = c + SPEED*(t + b + l + r - 4 * c);
}

5.完整代码

// 代码7.3.4使用一维纹理内存的热传导模型计算
//时间：2019.07.29
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <iostream>
#include "cpu_anim.h"
#include "cuda.h"
#include "book.h"

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

//声明纹理内存变量的引用，这些变量将位于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 copy_const_kernal(float *iptr)
{
	//将threadIdx/blockIdx映射到像素位置
	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;

}

__global__ void blend_kernal(float *dst, bool dstOut)
{
	//将threadIdx/blockIdx映射到像素位置
	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, b, l, r, c;
	if (dstOut)
	{
		t = tex1Dfetch(texIn,top);
		b = tex1Dfetch(texIn, bottom);
		l = tex1Dfetch(texIn, left);
		r = tex1Dfetch(texIn, right);
		c = tex1Dfetch(texIn, offset);
	}
	else
	{
		t = tex1Dfetch(texOut, top);
		b = tex1Dfetch(texOut, bottom);
		l = tex1Dfetch(texOut, left);
		r = tex1Dfetch(texOut, right);
		c = tex1Dfetch(texOut, offset);
	}

	dst[offset] = c + SPEED*(t + b + l + r - 4 * c);
}





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

	//由于tex是全局并且是有界的，因此我们必须通过一个标志来来选择每次迭代中哪个是输入/输出
	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
		{
			in = d->dev_outSrc;
			out = d->dev_inSrc;
		}
		copy_const_kernal << <grids, blocks >> > (in);
		blend_kernal << <grids, blocks >> > (out,dstOut);
		dstOut = !dstOut;
	}
	//将float数值映射成颜色，以便用图像显示它
	float_to_color << <grids, blocks >> >(d->output_bitmap, d->dev_inSrc);
	cudaMemcpy(bitmap->get_ptr(), d->output_bitmap, bitmap->image_size(), cudaMemcpyDeviceToHost);
	cudaEventRecord(d->stop, 0);
	cudaEventSynchronize(d->stop);
	float elapsedTime;
	cudaEventElapsedTime(&elapsedTime, d->start, d->stop);
	d->totalTime += elapsedTime;
	++d->frames;
	printf("Average Time per frame:%3.1fms\n", d->totalTime / d->frames);
}

void anim_exit(DataBlock *d)
{
	cudaUnbindTexture(texIn);
	cudaUnbindTexture(texOut);
	cudaUnbindTexture(texConstSrc);

	cudaFree(d->dev_constSrc);
	cudaFree(d->dev_inSrc);
	cudaFree(d->dev_outSrc);

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

int main()
{
	DataBlock data;
	CPUAnimBitmap bitmap(DIM, DIM, &data);
	data.bitmap = &bitmap;
	data.totalTime = 0;
	data.frames = 0;
	cudaEventCreate(&data.start);
	cudaEventCreate(&data.stop);
	cudaMalloc((void **)&data.output_bitmap, bitmap.image_size());
	cudaMalloc((void**)&data.dev_constSrc, bitmap.image_size());
	cudaMalloc((void **)&data.dev_inSrc, bitmap.image_size());
	cudaMalloc((void **)&data.dev_outSrc, bitmap.image_size());

	//分配内存之后，需要将这些变量绑定到之前声明的纹理引用上去
	//cudaChannelFormatDesc desc = cudaCreateChannelDesc<float>();
	cudaBindTexture(NULL, texConstSrc, data.dev_constSrc, bitmap.image_size());
	cudaBindTexture(NULL, texIn, data.dev_inSrc,  bitmap.image_size());
	cudaBindTexture(NULL, texOut, data.dev_outSrc,  bitmap.image_size());


	//在CPU上分配临时内存，对其填充一些内容，用于初始化data.dev_constSrc和data.dev_inSrc
	float *temp = (float *)malloc(bitmap.image_size());
	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;
		}
	}
	//用当前填充的内容初始化data.dev_constSrc
	cudaMemcpy(data.dev_constSrc, temp, bitmap.image_size(), cudaMemcpyHostToDevice);

	//继续向temp中填充一些内容
	for (int y = 800; y < DIM; y++)
	{
		for (int x = 0; x < 200; x++)
		{
			temp[x + y*DIM] = MAX_TEMP;
		}
	}
	//用给当前填充的内容初始化data.dev_inSrc
	cudaMemcpy(data.dev_inSrc, temp, bitmap.image_size(), cudaMemcpyHostToDevice);

	//释放CPU端的临时内存
	free(temp);
	//用anim_gpu生成动画的每一帧，用anim_exit退出动画显示并释放内存
	bitmap.anim_and_exit((void(*) (void *, int))anim_gpu, (void(*)(void *))anim_exit);
}

6.运行结果
在这里插入图片描述
在这里插入图片描述
bang!!!性能并没有得到提升，反而急剧恶化，不知所措，如果有哪位同学遇到跟我同样的问题，欢迎联系我讨论

GPU使用二维纹理内存的热传导模拟计算

1. 二维纹理内存变量引用的声明

//声明纹理内存变量的引用，这些变量将位于GPU上
texture<float,2> texConstSrc;
texture<float,2> texIn;
texture<float,2> texOut;

2.二维纹理内存引用变量与实际内存缓冲区的绑定
通过cudaBindTexture2D()可以将纹理内存引用变量与内存缓冲区进行绑定。
绑定二维纹理时，CUDA运行时要求提供一个cudaChannelFormatDesc，保持默认设置就可以了。

	cudaMalloc((void **)&data.output_bitmap, bitmap.image_size());
	cudaMalloc((void**)&data.dev_constSrc, bitmap.image_size());
	cudaMalloc((void **)&data.dev_inSrc, bitmap.image_size());
	cudaMalloc((void **)&data.dev_outSrc, bitmap.image_size());

	//分配内存之后，需要将这些变量绑定到之前声明的纹理引用上去
	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);

3.读取二维纹理内存方法
通过tex2D()读取二维纹理内存

__global__ void blend_kernal(float *dst, bool dstOut)
{
	//将threadIdx/blockIdx映射到像素位置
	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, b, l, r, c;
	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 + l + r - 4 * c);
}

4.完整代码

// 代码7.3.4使用二维纹理内存的热传导模型计算
//时间：2019.07.29
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <iostream>
#include "cpu_anim.h"
#include "cuda.h"
#include "book.h"

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

//声明纹理内存变量的引用，这些变量将位于GPU上
texture<float,2> texConstSrc;
texture<float,2> texIn;
texture<float,2> 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 copy_const_kernal(float *iptr)
{
	//将threadIdx/blockIdx映射到像素位置
	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;

}

__global__ void blend_kernal(float *dst, bool dstOut)
{
	//将threadIdx/blockIdx映射到像素位置
	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, b, l, r, c;
	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 + l + r - 4 * c);
}





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

	//由于tex是全局并且是有界的，因此我们必须通过一个标志来来选择每次迭代中哪个是输入/输出
	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
		{
			in = d->dev_outSrc;
			out = d->dev_inSrc;
		}
		copy_const_kernal << <grids, blocks >> > (in);
		blend_kernal << <grids, blocks >> > (out, dstOut);
		dstOut = !dstOut;
	}
	//将float数值映射成颜色，以便用图像显示它
	float_to_color << <grids, blocks >> >(d->output_bitmap, d->dev_inSrc);
	cudaMemcpy(bitmap->get_ptr(), d->output_bitmap, bitmap->image_size(), cudaMemcpyDeviceToHost);
	cudaEventRecord(d->stop, 0);
	cudaEventSynchronize(d->stop);
	float elapsedTime;
	cudaEventElapsedTime(&elapsedTime, d->start, d->stop);
	d->totalTime += elapsedTime;
	++d->frames;
	printf("Average Time per frame:%3.1fms\n", d->totalTime / d->frames);
}

void anim_exit(DataBlock *d)
{
	cudaUnbindTexture(texIn);
	cudaUnbindTexture(texOut);
	cudaUnbindTexture(texConstSrc);

	cudaFree(d->dev_constSrc);
	cudaFree(d->dev_inSrc);
	cudaFree(d->dev_outSrc);

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

int main()
{
	DataBlock data;
	CPUAnimBitmap bitmap(DIM, DIM, &data);
	data.bitmap = &bitmap;
	data.totalTime = 0;
	data.frames = 0;
	cudaEventCreate(&data.start);
	cudaEventCreate(&data.stop);
	cudaMalloc((void **)&data.output_bitmap, bitmap.image_size());
	cudaMalloc((void**)&data.dev_constSrc, bitmap.image_size());
	cudaMalloc((void **)&data.dev_inSrc, bitmap.image_size());
	cudaMalloc((void **)&data.dev_outSrc, bitmap.image_size());

	//分配内存之后，需要将这些变量绑定到之前声明的纹理引用上去
	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);


	//在CPU上分配临时内存，对其填充一些内容，用于初始化data.dev_constSrc和data.dev_inSrc
	float *temp = (float *)malloc(bitmap.image_size());
	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;
		}
	}
	//用当前填充的内容初始化data.dev_constSrc
	cudaMemcpy(data.dev_constSrc, temp, bitmap.image_size(), cudaMemcpyHostToDevice);

	//继续向temp中填充一些内容
	for (int y = 800; y < DIM; y++)
	{
		for (int x = 0; x < 200; x++)
		{
			temp[x + y*DIM] = MAX_TEMP;
		}
	}
	//用给当前填充的内容初始化data.dev_inSrc
	cudaMemcpy(data.dev_inSrc, temp, bitmap.image_size(), cudaMemcpyHostToDevice);

	//释放CPU端的临时内存
	free(temp);
	//用anim_gpu生成动画的每一帧，用anim_exit退出动画显示并释放内存
	bitmap.anim_and_exit((void(*) (void *, int))anim_gpu, (void(*)(void *))anim_exit);
}

5.运行结果
在这里插入图片描述
在这里插入图片描述
bang!!!性能并没有得到提升，反而急剧恶化，不知所措，如果有哪位同学遇到跟我同样的问题，欢迎联系我讨论

Leo-Ma
关注
  • 0
    点赞
  • 踩
  • 5
    收藏
    觉得还不错? 一键收藏
  • 0
    评论
专栏目录
【CUDA学习笔记】1.纹理内存
dgh_dean的博客
11-09 1264
并没有一个专门的显存空间作为纹理内存，而是数据以一维、二维或者三维数组的形式存储在显存中，可以通过缓存加速访问的内存，并且可以声明大小比常数存储器要大的多。 调用纹理获取函数读取纹理的过程称为纹理获取（texture fetch）。每一个纹理获取都需要指定一个参数，如果使用纹理对象API则该参数称为纹理对象，使用纹理引用API时称为纹理引用参数。 ·纹理是被获取的一块纹理内存。纹理对象在运行时
CUDA By Examples 8 - 纹理内存Texture Memory
Full_Speed_Turbo
06-28 1344
可以使用一维和二维纹理内存. 纹理内存是只读的.
《CUDA By Example》【Chapter 07】纹理内存 ？
wondervictor的博客
06-06 243
7.1 本章目标 纹理内存和常量内存一样，都是只读内存，用来减少带宽流量，提升性能。 1，了解纹理内存的特性 2，了解如何在CUDA C中使用一维纹理内存 3，了解如何在CUDA C中使用二维纹理内存 7.2 纹理内存简介（Texture Memory） 7.3 热传导模拟...
[菜鸟每天来段CUDA_C]CUDA实现简单热传导动态模拟
Running Jonny
12-03 3349
本文利用CUDA实现简单二维平面上的热传导模拟。假设有一个矩形房间，将其分成一个网格，在网格中随机散布一些热源，热源有不同的固定温度，然后就可以计算网格中每个单元的温度随时间的变化情况。本文将热传导模型做如下简化：某单元的新时刻温度等于原有温度加其与邻接单元的温差，邻接单元取上下左右四个单元。
CUDA by Example（英文原书+自带源代码）
01-02
《CUDA by Example》是一本专为想要学习GPU编程，尤其是CUDA技术的读者量身打造的实战指南。CUDA，全称Compute Unified Device Architecture，是NVIDIA公司推出的一种并行计算平台和编程模型，它允许开发者利用图形...
CUDA by Example (2010).pdf
10-29
CUDA by Example (2010).pdf CUDA by Example (2010).pdf
cuda by example 中英文及代码
03-24
cuda by example 中英文版及代码，中文版名称是 GPU高性能编程CUDA实战，含书签
cuda by example.pdf
07-10
通过实例学习cuda编程，nvidia GPU学习教程...........
cuda by example中文版
04-10
3. **内存管理**：CUDA提供了多种类型的内存，如全局内存、共享内存、常量内存和纹理内存等。理解它们的区别和使用场景，能有效提升数据传输效率。 4. **流与同步**：CUDA流（stream）允许异步执行计算任务，提高...
简单的传热模型的动画模拟
wufenxia的专栏
06-14 4102
/* 简单的二维热传导模拟： 首先，假设有一个矩形房间，将其分成一个网格，在网格中随机散布一些热源， 他们有着不同的固定温度，在给定了矩形网格及热源分布后可以计算网格中每 个单元的温度随时间的变化情况，为了简单热源单元本身的温度将保持不变， 在时间递进的每个步骤中，假设热量在某个单干及其邻接单元之间流动，如果 某个单元的临接单元的温度更高，那么热量将从临接单元传导到该单元，相反的 如
使用用一维纹理存储器的热传导模拟
wufenxia的专栏
06-19 2034
#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 /*在温度更新计算的内存访问模式中存在着巨大的内存空间局部性，这种访
基于纹理内存的CUDA热传导模拟
weixin_30677475的博客
09-24 184
原文链接 项目中有三个，第一个是全局内存，其余两个分别是基于1d和2d纹理内存。项目打包下载。 纹理内存是只读内存，与常量内存相同的是，纹理内存也缓存在芯片中，因此某些情况下，它能减少对内存的请求并提供更高效的内存宽带。纹理内存专门为那些内存访问模式中存在大量空间局部性的图形应用程序而设计的。在某个计算应用程序中，这意味着一个线程读取的位置可能与邻近线程读取的位置“非常接近”。纹理缓存为了加速...
CUDA Samples: Long Vector Add
网络资源是无限的
07-20 1134
CUDA Samples: Long Vector Add
使用二维纹理内存模拟热传导
wufenxia的专栏
06-19 1589
#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
CUDA总结：纹理内存
热门推荐
Kelvin_Yan的专栏
01-04 1万+
来源：《CUDA_Runtime_API》、《CUDA_C_Programming_Guide-V8.0》 纹理内存和表面内存（surface memory）实质上是全局内存的一个特殊形态，全局内存被绑定为纹理内存（表面内存），对其的读（写）操作将通过专门的texture cache（纹理缓存）进行，其实称为纹理缓存更加贴切。 纹理缓存的优势：纹理缓存具备硬件插值特性，可以实现最近邻插值和线性
CUDA By Example(六)——纹理内存
hhhcbw的博客，欢迎各位来访
11-17 1050
CUDA C中纹理内存使用的说明，并以热传导模拟为例
CUDA内存--纹理内存的说明与使用
学习使我快乐
09-22 7548
纹理介绍      纹理存储器（texture memory）是一种只读存储器，由GPU用于纹理渲染的的图形专用单元发展而来，因此也提供了一些特殊功能。      纹理缓存的优势：纹理缓存具备硬件插值特性，可以实现最近邻插值和线性插值。纹理缓存针对二维空间的局部性访问进行了优化，所以通过纹理缓存访问二维矩阵的邻域会获得加速。纹理缓存不需要满足全局内存的合并访问条件。      纹理可以是一段
NVIDIA CUDA 学习 (5) Texture Memory
weixin_47658743的博客
09-09 720
Texture Memory纹理内存热传导模型 纹理内存 热传导模型
《GPU高性能编程CUDA实战》学习笔记（八）
博客
09-13 5638
第8章 图形互操作性 GPU既执行渲染计算，又执行通用计算。。。 CUDA C应用可以无缝地与OpenGL和DirectX这两种实时渲染API进行交互。 本章中包含大量OpenGL和GLUT（OpenGL Utility Toolkit）代码，要有些思想准备。 8.1 本章目标 了解图形互操作性是什么，以及为什么需要使用它。了解如何设置某个CUDA设备的图形互操作性。了解如何在C
cuda by example
最新发布
07-30
《CUDA by Example》是一本深入介绍CUDA编程技术的书籍。 《CUDA by Example》这本书详细地介绍了CUDA平台的基本概念和编程模型，并通过实际的示例代码来演示如何使用CUDA编写高效的并行计算程序。 首先，书中详尽...
评论
添加红包

请填写红包祝福语或标题

个

红包个数最小为10个

元

红包金额最低5元

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

抵扣说明：

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

余额充值