CUDAExample: 波纹示例

张大饼的最爱

于 2024-09-27 16:33:51 发布

阅读量142

点赞数 1

分类专栏： CUDA 文章标签： CUDA

本文链接：https://blog.csdn.net/Jane_yuhui/article/details/142596672

版权

CUDA 专栏收录该内容

4 篇文章 1 订阅

订阅专栏

需要用到的已经封装好的辅助类： CPUAnimBitmap

定义个grid 和 block 的维度都是二维的

#include"cpu_anim.h"
#include "cuda_runtime.h"
#include <device_launch_parameters.h>

#define DIM 512
static void CheckCudaErrorAux(const char*, unsigned, const char*, cudaError_t);
#define CUDA_CHECK_RETURN(value) CheckCudaErrorAux(__FILE__,__LINE__, #value, value)


struct DataBlock 
{
	unsigned char* dev_bitmap;
	CPUAnimBitmap* bitmap;	
};

void cleanup(DataBlock *d) 
{
	cudaFree(d->dev_bitmap);
}

// 第一个指针来指向北村输出像素值的设备内存，是一个全局变量
// ticks 表示当前的动画时间
__global__ void kernel(unsigned char * imageptr, int ticks)
{
	// 每个线程都得到它在线程块中的索引，以及这个线程块在线程格中的索引
	// 并将这两个值转换为图形中的唯一索引（x,y）
	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 fx = x - DIM / 2;
	float fy = y - DIM / 2;
	float d = sqrtf(fx * fx + fy * fy);

	unsigned char grey = (unsigned char)(128.0f + 127.0f * cos(d / 10.0f - ticks / 7.0f) / (d / 10.0f + 1.0f));

	imageptr[offset * 4 + 0] = grey;
	imageptr[offset * 4 + 1] = 0;
	imageptr[offset * 4 + 2] = 255 - grey;
	imageptr[offset * 4 + 3] = 255;

}

// 根据时间戳ticks生成新的图像
void generate_frame(DataBlock *d, int ticks) 
{
	// 线程格中包含的并行线程块的数量
	dim3 blocks(DIM / 16, DIM / 16);
	// 每个线程块中包含的线程数量
	dim3 threads(16, 16);
	kernel <<<blocks, threads>>> (d->dev_bitmap, ticks);
	CUDA_CHECK_RETURN(cudaMemcpy(d->bitmap->get_ptr(),
		d->dev_bitmap,
		d->bitmap->image_size(),
		cudaMemcpyDeviceToHost));
}


int main(void)
{
	DataBlock data;
	CPUAnimBitmap	bitmap(DIM, DIM, &data);
	data.bitmap = &bitmap;
	std::cout << "image size : " << data.bitmap->image_size() << std::endl;
	CUDA_CHECK_RETURN(cudaMalloc((void**)&data.dev_bitmap, bitmap.image_size()));
	// 将指向generate_frame()的函数指针传递给anim_and_exit()
	// 每当要生生成一帧新的动画都会调用generate_frame()
	bitmap.anim_and_exit((void(*)(void *,int))generate_frame, (void(*)(void*))cleanup);

	return 0;
}

static void CheckCudaErrorAux(const char* file, unsigned line, const char* statement, cudaError_t err)
{
	if (err == cudaSuccess)
		return;
	std::cerr << statement << " returned " << cudaGetErrorString(err) << "(" << err << ") at " << file << ":" << line << std::endl;
	exit(1);
}