CUDA

CUDA-常量内存提升性能-例-RayTracing

  1. constants 常量必须在函数外声明

  2. cudaMemcpyToSymbol()

    这个函数是特殊版本的cudaMemcpy(), 唯一的不同是，cudaMemcpyToSymbol()将数据复制到常量内存constant memory，而cudaMemcpy()将数据复制到global memory。

  3. 常量内存为什么能提升性能-Ray tracing

    __constant__ 把变量的访问限制为read-only，有了这个限制，必然获得某种回报。这个回报是：与从global memory中读取数据相比，从constant memory中读取相同的数据可以节约内存带宽，有两个原因：

    • 对constant memory的单次读操作可以广播到其他临近的threads。

    • constant memory的数据将会缓存起来，因此相同地址的连续读操作不会产生额外的内存通信。

      half-warp的 广播是一把双刃剑。当16 个threads都读取constant memory相同的地址的时，性能极大地提升。但是当16个threads分别读取从constant memory不同的地址时，读操作会被串行化，性能会下降，这种情况就不如从global中读取。

      从Ray-Tracing实例中体会：

      要使用到的数据结构：

      1
      2
      3
      4
      5
      6
      7
      8
      9
      10
      11
      12
      13
      14
      15
      16
      17
      18
      19
      20
      #define INF 2e10f
      struct Sphere {
          float r,b,g;
          float radius;
          float x,y,z;
          // 假设有一个观察平面，那么（ox，oy）是这个观察平面中一个像素的坐标，
          // 这个方法将计算从这个像素中发射出的光线是否与这个球面相交。
          // 如果相交，则计算从这个像素点到这个球面的距离。当光线命中多个球时，
          // 只有最近的球面才会被看到。
          __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;
          }
      };

      kernel函数输入一个bitmap：

      1
      2
      3
      4
      5
      6
      7
      8
      9
      10
      11
      12
      13
      14
      15
      16
      17
      18
      19
      20
      21
      22
      23
      24
      25
      __global__ void kernel( unsigned char *ptr ) {
          // map from threadIdx/BlockIdx to pixel position
          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;
              }
          }
          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;
      }

      for循环中，每一个thread循环球集合Sphere *s中的所有球s[i]，得到每个像素的一个颜色值。这里的模式是：所有threads都会只读相同的存储地址。考虑上述的关于constant的特点，使用constant memory 进行优化。

      优化前的main函数如下：其中s是声明在global memory中的。用constant memory优化只需要将下面code中的<1><2>分别修改为

    • <1> __constant__ Sphere s[SPHERES];

    • <2> HANDLE_ERROR( cudaMemcpyToSymbol( s, temp_s, sizeof(Sphere) * SPHERES) );

      1
      2
      3
      4
      5
      6
      7
      8
      9
      10
      11
      12
      13
      14
      15
      16
      17
      18
      19
      20
      21
      22
      23
      24
      25
      26
      27
      28
      29
      30
      31
      32
      33
      34
      35
      36
      37
      38
      39
      40
      41
      #include "cpu_bitmap.h"
      #define rnd( x ) (x * rand() / RAND_MAX)
      #define SPHERES 20
      Sphere *s; // <1>

      int main( void ) {
          CPUBitmap bitmap( DIM, DIM );
          unsigned char *dev_bitmap;
          // allocate memory on the GPU for the output bitmap
          HANDLE_ERROR( cudaMalloc( (void**)&dev_bitmap, bitmap.image_size() ) );
          // allocate memory for the Sphere dataset
          HANDLE_ERROR( cudaMalloc( (void**)&s, sizeof(Sphere) * SPHERES ) ); // <2>

          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;
          }

          HANDLE_ERROR( cudaMemcpy( s, temp_s, sizeof(Sphere) * SPHERES,
                                      cudaMemcpyHostToDevice ) );
          free( temp_s );

          // generate a bitmap from our sphere data
          dim3 grids(DIM/16,DIM/16);
          dim3 threads(16,16);
          kernel<<<grids,threads>>>( dev_bitmap );

          // copy our bitmap back from the GPU for display
          HANDLE_ERROR( cudaMemcpy( bitmap.get_ptr(), dev_bitmap, bitmap.image_size(),
                                      cudaMemcpyDeviceToHost ) );
          bitmap.display_and_exit();
          // free our memory
          cudaFree( dev_bitmap );
          cudaFree( s );
      }

      上述过程使用constant memory保存只读对象。感受这个模式：每个threads都访问相同的只读数据时，将获得额外的性能提升. 前面说了的两个原因：第一，这种模式将读取操作在半个warp中广播，第二，芯片上包含了常量内存缓存。

      在许多算法中，内存带宽都是瓶颈，因此要时刻想着改善这种情况，