本章目标
- 了解如何在CUDA C中使用常量内存
- 了解常量内存的性能特性
-
学习如何使用CUDA事件来测量应用程序的性能
GPU上包含有数百个数学计算单元,因此性能瓶颈通常并不在于芯片的数学计算吞吐量,而是在于芯片的内存带宽。由于在图形处理器上包含了非常多的数学逻辑单元(ALU),因此有时输入数据的速率甚至无法维持如此高的计算速率。因此,有必要研究一些手段来减少计算问题时的内存通信量。 CUDA中可以使用全局内存和共享内存,常量内存。NVIDIA硬件提供了64KB的常量内存。在某些情况下,用常量内存来替换全局内存能有效地减少内存带宽。
光线跟踪使用常量内存的实例
光线跟踪是从三维对象场景中生成二维图像的一种方式。 方法:跟踪从像素中投射出的光线穿过场景,直到光线命中某个物体,然后计算这个像素的颜色。我们说像素将看到这个物体,并根据它所看到物体的颜色来设置它的颜色。光线跟踪中的大部分计算都是光线与场景中物体的相交 运算。
在GPU上实现光线跟踪
// Created by zhangjian on 19-6-17.
//
#include "../common/book.h"
#include "../common/image.h"
#include <iostream>
#define rnd(x) (x * rand() / RAND_MAX)
#define INF 2e10f
#define SPHERES 20
const int DIM = 1024;
struct Sphere{
float r, b, g;
float radius;
float x, y, 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;
}
};
__constant__ Sphere s[SPHERES];
__global__ void kernel(unsigned char *ptr){
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){
cudaEvent_t start, stop;
cudaEventCreate(&start);
cudaEventCreate(&stop);
cudaEventRecord(start, 0);
IMAGE bitmap(DIM, DIM);
unsigned char * dev_bitmap;
// Sphere *s;
cudaMalloc((void**)&dev_bitmap, bitmap.image_size());
// cudaMalloc((void**)&s, sizeof(Sphere) * SPHERES);
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);
free(temp_s);
dim3 grids(DIM / 16, DIM / 16);
dim3 threads(16, 16);
kernel<<<grids, threads >>>(dev_bitmap);
cudaMemcpy(bitmap.get_ptr(), dev_bitmap, bitmap.image_size(), cudaMemcpyDeviceToHost);
cudaEventRecord(stop, 0);
cudaEventSynchronize(stop);
float elapsedTime;
cudaEventElapsedTime(&elapsedTime, start, stop);
std::cout << "Time to generate: %3.1f ms\n" << elapsedTime << std::endl;
cudaEventDestroy(start);
cudaEventDestroy(stop);
cudaFree(dev_bitmap);
// cudaFree(s);
bitmap.show_image();
}
- 常量内存的声明是在变量前加上__constant__修饰符。
- 对于常量内存,不需要再调用cudaMalloc()和cudaFree()。
- 当从主机内存复制到GPU上的常量内存时,需要使用cudaMemcpyToSymbol()。cudaMemcpy()复制全局内存。
常量内存带来的性能提升
__constant__将把变量的访问限制为只读,从常量内存中读取 相同数据可以节约内存带宽。是因为线程束(包含32个线程)中,如果半线程束中的每个线程都从常量内存中读取相同数据时,GPU只会产生一次读取请求并在随后将那个数据广播到每个线程。此方式产生的内存流量为使用全局内存的4%。 当16个线程读取相同地址时,可以极大地提升性能,但是当他们分别读取不同地址时,实际上会降低性能。
使用CUDA事件来测量性能
CUDA事件本质上是一个时间戳。获得一个时间戳需要两个步骤:首先创建一个事件,然后记录一个事件。
cudaEvent_t start;
cudaEventCreate(&start);
cudaEventRecord(start, 0);
起始事件创建完毕,然后创建结束事件。
cudaEvent_t start, stop;
cudaEventCreate(&start);
cudaEventCreate(&stop);
cudaEventRecord(start, 0);
// 在GPU上执行工作
cudaEventRecord(stop, 0);
在GPU执行完之前,CPU会继续执行程序中的下一行代码,所以需要两者同步。
添加
cudaEventSynchronize(stop);
最后cudaEventElapsedTime()和cudaEventDestroy(),前者是一个工具函数,用于记录两个事件之间经历的时间,第一个参数为某个浮点变量的地址,单位为毫秒。最后用后者进行销毁。