《
CUDA by example
》
中的第六章讲解了在
GPU
上实现光线跟踪的一个例子,旨在介绍常量内存(
constant memory
)和事件,下面给出这个例子的详细解读(
http://code2.us/2012/02/cuda_learning_11-constant_memory_and_events/
)。
#include <stdio.h>
#include "common/cpu_bitmap.h"
//是否使用__constan__常量内存的开关
#define CONSTANT
#define INF 2e10f
#define rnd(x) (x*rand()/RAND_MAX)
#define SPHERES 200
#define DIM 800
//求结构体
struct
Sphere {
float
r, g, b;
//球的颜色
float
radius;
//球的半径
float
x, y, z;
//球心坐标
//判断从像素点(ox, oy)射出的射线是否与该球相交,并返回交点的z坐标
__device__
float
hit(
float
ox,
float
oy,
float
*n)
{
float
dx = ox - x;
float
dy = oy - y;
if
(dx*dx + dy*dy < radius*radius)
{
float
dz = sqrtf(radius*radius - dx*dx - dy*dy);
*n = dz/sqrtf(radius*radius);
return
dz + z;
}
return
-INF;
}
};
#ifdef CONSTANT
//__constant__ 常量内存
__constant__ Sphere s[SPHERES];
#else
Sphere *s;
#endif
#ifdef CONSTANT
//使用constant常量内存时,不能将其当作参数传到global函数
__global__
void
kernel(unsigned
char
* ptr)
#else
//普通全局变量必须用传参的形式传递到global函数
__global__
void
kernel(unsigned
char
* ptr, Sphere *s)
#endif
{
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
ox = (x - DIM/2);
float
oy = (y - DIM/2);
float
r=0,g=0,b=0;
//获得最近的交点
float
maxz = -INF;
for
(
int
i=0; i<SPHERES; i++)
{
float
n;
float
t = s[i].hit(ox, oy, &n);
if
(t>maxz)
{
float
fscale = n;
r = s[i].r * fscale;
g = s[i].g * fscale;
b = s[i].b * fscale;
maxz = t;
}
}
ptr[offset*4 + 0] = (
int
)(r*255);
ptr[offset*4 + 1] = (
int
)(g*255);
ptr[offset*4 + 2] = (
int
)(b*255);
ptr[offset*4 + 3] = 255;
}
int
main(
void
)
{
//使用cuda事件来测试性能
cudaEvent_t start, stop;
cudaEventCreate(&start);
cudaEventCreate(&stop);
cudaEventRecord(start, 0);
CPUBitmap bitmap(DIM, DIM);
unsigned
char
* dev_bitmap;
cudaMalloc((
void
**) &dev_bitmap, bitmap.image_size());
#ifdef CONSTANT
// __constant__常量内存不需要动态分配内存
#else
// 在GPU设备上分配内存给球数组
cudaMalloc((
void
**) &s,
sizeof
(Sphere) * SPHERES);
#endif
// 在CPU上生成求数据数据
Sphere *temp_s = (Sphere *)
malloc
(
sizeof
(Sphere) * SPHERES);
for
(
int
i=0; i<SPHERES; i++)
{
temp_s[i].r = rnd(1.0f);
temp_s[i].g = rnd(1.0f);
temp_s[i].b = rnd(1.0f);
temp_s[i].x = rnd(1000.f) - 500;
temp_s[i].y = rnd(1000.f) - 500;
temp_s[i].z = rnd(1000.f) - 500;
temp_s[i].radius = rnd(100.f) + 20;
}
#ifdef CONSTANT
//从CPU拷贝到__constant__常量内存
cudaMemcpyToSymbol(s, temp_s,
sizeof
(Sphere) * SPHERES);
#else
//从CPU拷贝到GPU
cudaMemcpy(s, temp_s,
sizeof
(Sphere) * SPHERES, cudaMemcpyHostToDevice);
#endif
free
(temp_s);
dim3 grids(DIM/16, DIM/16);
dim3 threads(16, 16);
#ifdef CONSTANT
kernel<<<grids, threads>>>(dev_bitmap);
#else
kernel<<<grids, threads>>>(dev_bitmap, s);
#endif
cudaMemcpy(bitmap.get_ptr(), dev_bitmap, bitmap.image_size(), cudaMemcpyDeviceToHost);
cudaEventRecord(stop, 0);
//事件同步
cudaEventSynchronize(stop);
float
elapseTime;
cudaEventElapsedTime(&elapseTime, start, stop);
printf
(
"Time to generate: %3.1f ms\n"
, elapseTime);
cudaEventDestroy(start);
cudaEventDestroy(stop);
bitmap.display_and_exit();
cudaFree(dev_bitmap);
#ifdef CONSTANT
// __constant__ 常量内存不需要free
#else
cudaFree(s);
#endif
return
1;
}
与标准的全局常量内存相比,常量内存存在着一些限制,但在某些情况中,实用常量内存将提升应用程序的性能。
特别是,当线程束中的所有线程都访问相同的只读数据时,将获得额外的性能提升。在这种数据访问模式中实用常量
内存可以节约内存带宽,不仅是因为这种模式可以将读取操作在半线程束中广播,而且还因为在芯片上包含了常量内存
缓存。在许多算法中,内存带宽都是一种瓶颈,因此采用一些机制来改善这种情况是非常有用的。
知识点:
1. 使用__constant__修饰符来声明变量为常量内存;
2. 常量内存为静态分配空间,所以不需要调用 cudaMalloc(), cudaFree();
3. CUDA 中的时间本质上是一个 GPU 时间戳,这个时间戳是在用户指定的时间点上记录的。
|