[菜鸟每天来段CUDA_C]CUDA实现简单热传导动态模拟

原创 2013年12月03日 15:14:09

本文利用CUDA实现简单二维平面上的热传导模拟。假设有一个矩形房间,将其分成一个网格,在网格中随机散布一些热源,热源有不同的固定温度,然后就可以计算网格中每个单元的温度随时间的变化情况。本文将热传导模型做如下简化:某单元的新时刻温度等于原有温度加其与邻接单元的温差,邻接单元取上下左右四个单元。

温度更新算法:

Step1: 给定一个包含初始输入温度的网格,将其作为热源温度值复制到网格对应的单元;

Step2: 规定一个输入温度网格,根据热传导模型更新温度到输出网格;

Step3: 交换输入温度网格和输出温度网格的数据,转step1.


实验结果如下:下图显示按先后顺序的四个不同时刻的热量分布及变化情况。



主要程序代码如下:

/********************************************************************
*  heatTransfer.cu
*********************************************************************/

#include <stdio.h>
#include <stdlib.h>
#include <cuda_runtime.h>
#include <cutil_inline.h>
#include "CPUAnim.h"

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

struct DataBlock 
{
	unsigned char	*outputBitmap;
	float			*dev_inSrc;
	float			*dev_outSrc;
	float			*dev_constSrc;
	CPUAnimBitmap	*bitmap;

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

template< typename T >
void swap( T& a, T& b ) {
	T t = a;
	a = b;
	b = t;
}


/************************************************************************/
/* Init CUDA                                                            */
/************************************************************************/
bool InitCUDA(void)
{
    ......
}

/************************************************************************/
__global__ void copyConstKernel(float *iPtr, const float *cPtr)
{
	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 blendKernel(float* outSrc, const float* inSrc)
{
	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] - 4 * inSrc[offset]);
}

__device__ unsigned char value( float n1, float n2, int hue )
{
      ......
}
__global__ void floatToColor( unsigned char *optr,
							   const float *outSrc ) 
{
    ......
}

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

	for (int i=0; i<90; i++)
	{
		copyConstKernel<<<blocks, threads>>>(d->dev_inSrc, d->dev_constSrc);
		blendKernel<<<blocks, threads>>>(d->dev_outSrc, d->dev_inSrc);
		swap(d->dev_inSrc, d->dev_outSrc);
	}
	floatToColor<<<blocks, threads>>>(d->outputBitmap, d->dev_inSrc);

	cudaMemcpy(bitmap->get_ptr(), d->outputBitmap, bitmap->image_size(), cudaMemcpyDeviceToHost);
	cutilSafeCall(cudaEventRecord(d->stop, 0));
	cutilSafeCall(cudaEventSynchronize(d->stop));

	float elapsedTime;
	cutilSafeCall(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 animExit(DataBlock *d)
{
	cudaFree(d->dev_inSrc);
	cudaFree(d->dev_outSrc);
	cudaFree(d->dev_constSrc);

	cutilSafeCall(cudaEventDestroy(d->start));
	cutilSafeCall(cudaEventDestroy(d->stop));
}
/************************************************************************/

int main(int argc, char* argv[])
{

	if(!InitCUDA()) {
		return 0;
	}

	DataBlock	data;
	CPUAnimBitmap bitmap(DIM, DIM, &data);
	data.bitmap = &bitmap;
	data.totalTime = 0;
	data.frames = 0;

	cutilSafeCall(cudaEventCreate(&data.start));
	cutilSafeCall(cudaEventCreate(&data.stop));

	cutilSafeCall(cudaMalloc((void**)&data.outputBitmap, data.bitmap->image_size()));

	cutilSafeCall(cudaMalloc((void**)&data.dev_inSrc, data.bitmap->image_size()));
	cutilSafeCall(cudaMalloc((void**)&data.dev_outSrc, data.bitmap->image_size()));
	cutilSafeCall(cudaMalloc((void**)&data.dev_constSrc, data.bitmap->image_size()));

	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>301) && (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;
		}
	}

	cutilSafeCall(cudaMemcpy(data.dev_constSrc, temp, bitmap.image_size(), cudaMemcpyHostToDevice));
	
	for(int y=800; y<DIM; y++)
	{
		for (int x=0; x<200; x++)
		{
			temp[x+y*DIM] = MAX_TEMP;
		}
	}
	cutilSafeCall(cudaMemcpy(data.dev_inSrc, temp, bitmap.image_size(), cudaMemcpyHostToDevice));
	free(temp);
	bitmap.anim_and_exit((void(*)(void*, int))animGPU, (void(*)(void*))animExit);
	return 0;
}


相关文章推荐

在没有nvidia显卡的环境下学些cuda编程

https://developer.nvidia.com/cuda-toolkit-archive

OpenCUDA-基于CUDA的图像并行算法开源程序库

开新坑……OpenCUDACUDA(Compute Unified Device Architecture),是显卡厂商NVIDIA推出的运算平台。 随着GPU的发展,CUDA使用人数也越来越多。但...

图像分割之(六)CUDA实现GrabCut

CUDA提供了GrabCut测试案例: NVIDIA Corporation\CUDA Samples\v5.0\7_CUDALibraries\grabcutNPP 原图: 结果图: htt...

[菜鸟每天来段CUDA_C]GPU上通过常量内存实现光线跟踪

光线跟踪是从三维场景生成二维图像的一种方式。主要思想为:在场景中选择一个位置放上一台假想的相机,该相机包含一个光传感器来生成图像,需要判断那些光将接触到这个传感器。图像中每个像素与命中传感器的光线有相...

[菜鸟每天来段CUDA_C] GPU上实现直方图计算

本文通过在GPU上计算直方图说明GPU计算中的原子操作。

[菜鸟每天来段CUDA_C] 利用页锁定内存提高运算效率

本文通过使用malloc分配内存和cudaHostAlloc分配页锁定内存,说明使用页锁定内存可提高运算效率,并指出哪些场合适合使用页锁定内存。 malloc分配的是标准的可分页的(pagable)的...

[菜鸟每天来段CUDA_C]基于GPU的Julia集

Julia集是一个著名的分形集,它是复数经过迭代得到的,是满足某个复数计算函数的所有点构成的边界。 算法思想: 通过一个简单的迭代等式对复平面中的点求值,如果在计算某个点时迭代的结果是发散...

基于纹理内存的CUDA热传导模拟

  • 2014年09月24日 16:25
  • 14.93MB
  • 下载

[菜鸟每天来段CUDA_C]CppIntegration在C++程序中引用CUDA程序

本文主要实现在C++程序中引用CUDA程序,主要意义是使顺序定义的数据能在CUDA程序中并行执行,然后返回结果。 程序主要包括main.cpp                           ...

[菜鸟每天来段CUDA_C]使用多个CUDA流提高程序执行效率

实验通过把一组数据分块复制到GPU执行,返回执行结果,来说明使用cuda流的使用能提高程序的执行效率...
内容举报
返回顶部
收藏助手
不良信息举报
您举报文章:[菜鸟每天来段CUDA_C]CUDA实现简单热传导动态模拟
举报原因:
原因补充:

(最多只允许输入30个字)