▶ 本章介绍了常量内存的使用，并给光线追踪的一个例子。介绍了结构cudaEvent_t及其在计时方面的使用。
● 章节代码，大意是有SPHERES个球分布在原点附近，其球心坐标在每个坐标轴方向上分量绝对值不大于500，其半径介于20到120；观察者（画面平面）位于z正半轴充分远处（z>500），现将所有的球体平行投影到画面平面上，考虑遮挡关系，并考虑球面与画面平面的夹角给球体绘制阴影。使用常量内存时球数组定义在所有函数外部，核函数只需图形参数就够了；不使用常量内存时球数组定义在结构DataBlock内部，核函数需要球数组和图形参数。
#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_bitmap.h" #define DIM 1024
#define rnd( x ) (x * rand() / RAND_MAX)
#define INF 2e10f
#define SPHERES 40
#define USE_CONSTANT_MEMORY false struct Sphere
{
float r, b, g;
float radius;
float x, y, z;
__device__ float hit(float ox, float oy, float *n)//计算球体上一点处的法向量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;
}
}; #if USE_CONSTANT_MEMORY
__constant__ Sphere s[SPHERES];
struct DataBlock
{
unsigned char *dev_bitmap;
};
#else
struct DataBlock
{
unsigned char *dev_bitmap;
Sphere *s;
};
#endif #if USE_CONSTANT_MEMORY
__global__ void kernel(unsigned char *ptr)
#else
__global__ void kernel(Sphere *s, unsigned char *ptr)
#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 / );
float oy = (y - DIM / ); float r = , g = , b = ;//计算该像素应该显示什么颜色
float maxz = -INF;
for (int i = ; i<SPHERES; i++)
{
float n;
float t = s[i].hit(ox, oy, &n);//o理解成"observation"，即当前线程代表的坐标
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 * + ] = (int)(r * );
ptr[offset * + ] = (int)(g * );
ptr[offset * + ] = (int)(b * );
ptr[offset * + ] = ; return;
} int main(void)
{
DataBlock data;
cudaEvent_t start, stop;// 计时器
cudaEventCreate(&start);
cudaEventCreate(&stop);
cudaEventRecord(start, ); CPUBitmap bitmap(DIM, DIM, &data);
unsigned char *dev_bitmap;
#if !USE_CONSTANT_MEMORY
Sphere *s;
#endif
cudaMalloc((void**)&dev_bitmap,bitmap.image_size());
cudaMalloc((void**)&s,sizeof(Sphere) * SPHERES);
Sphere *temp_s = (Sphere*)malloc(sizeof(Sphere) * SPHERES); for (int i = ; 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.0f) - ;
temp_s[i].y = rnd(1000.0f) - ;
temp_s[i].z = rnd(1000.0f) - ;
temp_s[i].radius = rnd(100.0f) + ;
} #if USE_CONSTANT_MEMORY
cudaMemcpyToSymbol(s, temp_s, sizeof(Sphere) * SPHERES);
kernel << < dim3(DIM / , DIM / ), dim3(, ) >> > (dev_bitmap);
#else
cudaMemcpy(s, temp_s, sizeof(Sphere) * SPHERES, cudaMemcpyHostToDevice);
kernel << < dim3(DIM / , DIM / ), dim3(, ) >> > (s, dev_bitmap);
#endif cudaMemcpy(bitmap.get_ptr(), dev_bitmap,bitmap.image_size(),cudaMemcpyDeviceToHost); cudaEventRecord(stop, );//测量计算耗时
cudaEventSynchronize(stop);
float elapsedTime;
cudaEventElapsedTime(&elapsedTime,start, stop);
printf("Time to generate: %3.1f ms\n", elapsedTime);
cudaEventDestroy(start);
cudaEventDestroy(stop); free(temp_s);
cudaFree(dev_bitmap);
cudaFree(s); bitmap.display_and_exit(); getchar();
return;
}
● 使用了结构cudaEvent_t用于计时，并介绍了与此相关的时间控制函数，按顺序使用如下。
cudaEvent_t start, stop;
cudaEventCreate(&start);
cudaEventCreate(&stop); cudaEventRecord(start, ); //Do something cudaEventRecord(stop, );
cudaEventSynchronize(stop); float elapsedTime;
cudaEventElapsedTime(&elapsedTime, start, stop);
cudaEventDestroy(start);
cudaEventDestroy(stop);
● 使用cudaMemcpyToSymbol()函数复制内存到到常量内存中（可以反向从显存复制到内存中，但由于只读一般没有情况会这样做）。其定义于cuda_runtime.h中
template<class T>
static __inline__ __host__ cudaError_t cudaMemcpyToSymbol(
const T &symbol,
const void *src,
size_t count,
size_t offset = ,
enum cudaMemcpyKind kind = cudaMemcpyHostToDevice)
{
return ::cudaMemcpyToSymbol((const void*)&symbol, src, count, offset, kind);
}