常量内存
使用常量内存的原因：
GPU中含有数量庞大的计算单元，因此性能瓶颈通常并不在于芯片的数学计算吞吐量，而是在于芯片的内存带宽。
在某些情况下，使用常量内存可以显著减少内存通信量
光线跟踪
光信跟踪是从三维对象场景中生成二维图像的一种方式。将跟踪从像素中投射出的光线穿过场景，直到光线命中某个物体，然后计算这个像素的颜色。
在GPU上实现光线追踪
- 构建球面数据结构
struct Sphere
{
float r, g, b;//颜色值
float radius;//半径
float x, y, z;//球心坐标
//运行在GPU上的函数，计算相机到光线命中球面处的距离
__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 z + dz;
}
return -INF;
}
};
其中方法hit(float ox,float oy,float *n)对来自(ox,oy)处像素的光线，计算光线是否与这个球面相交。如果光线与球面相交，这个方法将计算从相机到光线命中球面处的距离。当光线命中多个球面时，由于遮挡，只有最接近相机的球面才会被看见。
详细解释示意图：
之所以return dz+z是因为：去前球面的距离与取后球面的距离，本质上并没有什么差别。所以return dz+z和return -dz+z效果应该是一样的。
2. 思路说明
先随机生成一些球，然后在核函数中，对每个像素都遍历一遍这些球，找到最短的hit距离，根据这个hit距离与其球心的距离差值，确定这个像素的透明度，然后根据球的颜色和计算出的像素透明度对图像着色。
3.完整代码
// 代码6.2.2在GPU上实现光线跟踪
//时间：2019.07.28
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <iostream>
#include "cpu_bitmap.h"
#define INF 2e10f
#define SPHERES 20
#define DIM 1024
#define rnd(x) (x*rand()/RAND_MAX)
struct Sphere
{
float r, g, b;//颜色值
float radius;//半径
float x, y, z;//球心坐标
//运行在GPU上的函数，计算相机到光线命中球面处的距离
__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 z + dz;
}
return -INF;
}
};
struct DataBlock
{
unsigned char *dev_bitmap;
Sphere *s;
};
__global__ void kernal(Sphere *s,unsigned char *ptr)
{
//将threadIDX/blockIdx映射到像素位置
int x = threadIdx.x + blockIdx.x*blockDim.x;
int y = threadIdx.y + blockIdx.y*blockDim.y;
int offset = x + y*DIM;
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()
{
DataBlock data;
//记录起始时间
cudaEvent_t start, stop;
cudaEventCreate(&start);
cudaEventCreate(&stop);
cudaEventRecord(start, 0);
CPUBitmap bitmap(DIM, DIM, &data);
unsigned char *dev_bitmap;
Sphere *s;
//在GPU上分配内存以计算输出位图
cudaMalloc((void **)&dev_bitmap, bitmap.image_size());
//为Sphere数据集分配内存
cudaMalloc((void **)&s, sizeof(Sphere)*SPHERES);
//在CPU上分配Shpere的临时内存，对其初始化，并复制到GPU内存中，然后释放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.0f) - 500;
temp_s[i].y = rnd(1000.0f) - 500;
temp_s[i].z = rnd(1000.0f) - 500;
temp_s[i].radius = rnd(100.0f) + 20;
}
cudaMemcpy(s, temp_s, sizeof(Sphere)*SPHERES, cudaMemcpyHostToDevice);
free(temp_s);
//从球面数据中生成一张位图
dim3 grids(DIM / 16, DIM / 16);
dim3 blocks(16, 16);
kernal << <grids, blocks >> >(s,dev_bitmap);
//将位图从GPU复制到CPU中
cudaMemcpy(bitmap.get_ptr(), dev_bitmap, bitmap.image_size(), cudaMemcpyDeviceToHost);
// get stop time, and display the timing results
cudaEventRecord(stop, 0);
cudaEventSynchronize(stop);
float elapsedTime;
cudaEventElapsedTime(&elapsedTime,start, stop);
printf("Time to generate: %3.1f ms\n", elapsedTime);
cudaEventDestroy(start);
cudaEventDestroy(stop);
//释放GPU内存
cudaFree(dev_bitmap);
cudaFree(s);
//显示
bitmap.display_and_exit();
}
4. 运行结果(return z+dz)
5. 运行结果(return z-dz)
通过常量内存来实现光线追踪
1. 常量内存定义
**常量内存是指那些在程序运行过程中不能修改的内存空间。**光线追踪示例中只有一个输入数据，即球面数组，它在程序运行过程中不需要修改，因此将这个数据保存到常量内存中，以期待更少的内存通信量和更高效的运行性能。
2. 常量内存的使用方法
定义方法：
由原来的Sphere *s变成：
__constant__ Sphere s[SPHERES]
分配内存方法：
常量内存的分配不需要调用cudaMalloc()和cudaFree(),而是有专门的copy函数cudaMemcpyToSymbol，这个函数将会把CPU的内存内容copy到GPU中的常量内存中。
cudaMemcpyToSymbol(s,temp_s,sizeof(Sphere)*SPHERES);
其中temp_s是CPU内存指针。
cudaMemcpyToSymbol会复制到常量内存，而cudaMemcpy()会复制到全局内存。
3. 常量内存带来的性能提升
与从全局内存中读取数据相比，从常量内存中读取相同的数据可以节约内存带宽，主要有两个原因：
- 对常量内存的单次读操作可以广播到其他的"邻近"线程(半个warp)，这将节约15次读取操作。在CUDA架构中，线程束是指一个包含32个线程的集合，这个线程集合被"编织在一起"并且"步调一致"的形式执行。
- 常量内存的数据将缓存起来，因此对相同地址的连续读操作将不会产生额外的内存通信量。
在光线追踪示例中，使用常量内存后，硬件只需要请求这个数据一次。在缓存数据后，其他每个线程将不会产生内存流量，原因有两个：
- 线程将在半线程束的广播中收到这个数据
- 从常量内存缓存中收到数据
讨论：
只有当16个线程每次都只需要相同的读取请求时，才值得将这个读取操作广播到16个线程，如果半线程束中的所有16个线程需要访问常量内存中不同的数据，那么这16次不同的读取操作会被串行化，从而需要16倍的时间来发出请求。但如果从全局内存中读取，那么这些请求会同时发出。在这种情况中，从常量内存读取就慢于从全局内存中读取。
4. 完整代码
// 代码6.2.3在GPU上通过常量内存实现光线跟踪
//时间：2019.07.28
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <iostream>
#include "cpu_bitmap.h"
#define INF 2e10f
#define SPHERES 20
#define DIM 1024
#define rnd(x) (x*rand()/RAND_MAX)
struct Sphere
{
float r, g, b;
float radius;
float x, y, z;
//运行在GPU上的函数，计算相机到光线命中球面处的距离
__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 z + dz;
}
return -INF;
}
};
//Sphere *s;
__constant__ Sphere s[SPHERES];//GPU常量内存大小直接分配
//__global__ void kernal(Sphere *s, unsigned char *ptr)//也不需要在核函数中传入常量内存的地址，常量内存在核函数中可以直接使用
__global__ void kernal(unsigned char *ptr)
{
//将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 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;
}
struct DataBlock
{
unsigned char *dev_bitmap;
};
int main()
{
DataBlock data;
//记录起始时间
cudaEvent_t start, stop;
cudaEventCreate(&start);
cudaEventCreate(&stop);
cudaEventRecord(start, 0);
CPUBitmap bitmap(DIM, DIM,&data);
unsigned char *dev_bitmap;
//在GPU上分配内存以计算输出位图
cudaMalloc((void **)&dev_bitmap, bitmap.image_size());
//不需要再为Sphere数据集显式分配内存
//cudaMalloc((void **)&s, sizeof(Sphere)*SPHERES);
//在CPU上分配Shpere的临时内存，对其初始化，并复制到GPU内存中，然后释放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.0f) - 500;
temp_s[i].y = rnd(1000.0f) - 500;
temp_s[i].z = rnd(1000.0f) - 500;
temp_s[i].radius = rnd(100.0f) + 20;
}
//cudaMemcpy(s, temp_s, sizeof(Sphere)*SPHERES, cudaMemcpyHostToDevice);
cudaMemcpyToSymbol(s,temp_s,sizeof(Sphere)*SPHERES);//这将把CPU内存内容copy到GPU的常量内存上
free(temp_s);
//从球面数据中生成一张位图
dim3 grids(DIM / 16, DIM / 16);
dim3 blocks(16, 16);
kernal << <grids, blocks >> >(dev_bitmap);
//将位图从GPU复制到CPU中
cudaMemcpy(bitmap.get_ptr(), dev_bitmap, bitmap.image_size(), cudaMemcpyDeviceToHost);
// get stop time, and display the timing results
cudaEventRecord(stop, 0);
cudaEventSynchronize(stop);
float elapsedTime;
cudaEventElapsedTime(&elapsedTime, start, stop);
printf("Time to generate: %3.1f ms\n", elapsedTime);
cudaEventDestroy(start);
cudaEventDestroy(stop);
//释放GPU内存
cudaFree(dev_bitmap);
//cudaFree(s);//因为不需要显式分配常量内存，所以也不需要显式释放常量内存
//显示
bitmap.display_and_exit();
}
- 运行结果
bang!!!性能并没有提升，如果有同学看到这里，想深究原因的，请联系我一起交流讨论
使用事件来测量性能
为了测量GPU在某个任务上花费的时间，我们将使用CUDA的事件API。CUDA中的事件本质上是一个GPU时间戳，获得一个时间戳只需要两个步骤：1.创建一个事件 2.记录一个事件
使用方法：
cudaEvent_t start,stop;
cudaEventCreate(&start);
cudaEventCreate(&stop);
cudaEventRecord(start,0);
//在GPU上执行一些工作
cudaEventRecord(stop,0);
//同步stop事件，即只有所有线程的stop事件都发生时，才会执行下面的操作。
//当cudaEventSynchronize返回时，表明在stop事件之前的所有GPU工作已经完成了，
//此时可以安全的读取stop中保存的时间戳
cudaEventSynchronize(stop);
//获得结束事件，并显示计时结果
float elapsedTime;
cudaEventElapsedTime(&elapsedTime,start,stop);
printf("Time to generate: %3.1f ms\n",elapsedTime);
//销毁cuda事件
cudaEventDestroy(start);
cudaEventDestroy(stop);
通过常量内存来实现光线追踪一节的完整代码中已经包含了事件的完整代码，这里不再赘述。
扫一扫
515
被折叠的 条评论
为什么被折叠?
到【灌水乐园】发言
