CUDA - 常量内存与事件

CUDA学习

Posted by BY J on June 17, 2019

本章目标

  • 了解如何在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(),前者是一个工具函数,用于记录两个事件之间经历的时间,第一个参数为某个浮点变量的地址,单位为毫秒。最后用后者进行销毁。

效果展示

SPHRER