CUDA学习笔记（4）光线追踪（常量内存）

来源：互联网发布：手机熊猫网络设置编辑：程序博客网时间：2024/06/05 17:19

根据书上内容，本节通过一个实验“光线追踪”来了解和测试常量内存带来的效果提升。

常量内存是NVIDIA提供的一个64KB大小的内存空间，它的处理方式和普通的全局内存和共享内存都不一样，是有cuda专门提供的。

线程束的概念：线程束是指一个包含32个线程的集合，在程序中的每一行，线程束中的每个线程都将在不同的数据上执行相同的指令。

因此，常量内存的作用是，能够将单次内存的读取操作广播到每个半线程束（即16个线程），所以如果在半线程束中的每个线程都从常量内存的相同地址上读取数据，那么GPU只会产生一次读取请求，并将其广播，显而易见，这种方式的内存流量只是使用全局内存流量的1/16。这是常量内存的第一个好处，第二个好处则是由于这块内存的内容是不会发生变化的，因此硬件将主动把这个常量内存数据缓存到GPU上，这样第一次从敞亮内存的某个地址上读取后，其他半线程束请求同一个地址时，将直接在GPU上命中缓存，因此也减少了额外的内存流量。

使用常量内存只需加上:__constant__修饰符，当从主机内存复制内存到GPU上常量内存时，不用cudaMemcpy()而用cudaMemcpyToSymbol()，这样就复制到常量内存里了。

在测试程序中涉及到了cuda事件，我们通过这个API去记录程序运行的时间，以此来将是否用到常量内存这两个程序进行比较。这部分代码大致为：

cudaEvent_t start,stop;

cudaEventCreat(&start);

cudaEventCreat(&stop);

cudaEventRecord(start,0);

//执行工作

cudaEventRecord(stop,0);

cudaEventSynchronize(stop);

大致含义就是建立两个事件，并且记录第一个start时间，此时开始记录时间，工作执行完之后，再记录结束时间。然而我们为了让GPU完全执行完语句，再记录Stop里的准确时间，我们不得不加入cudaEventSynchronize()这个函数，否则将得到不可靠的时间结果。

以下代码为没有使用常量内存的光线追踪代码，测试球面为30个

#include "cuda_runtime.h"#include "device_launch_parameters.h"#include "H:\cuda_by_example\common\book.h"#include "H:\cuda_by_example\common\cpu_bitmap.h"#include "device_functions.h"#include <stdio.h>  #define DIM 1024    #define rnd( x ) (x * rand() / RAND_MAX)  #define INF     2e10f    //数据结构对球面建模struct Sphere {      float   r,b,g;      float   radius;      float   x,y,z;  //hit方法，计算光线是否与球面相交，若相交则返回光线到命中球面处的距离    __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;      }  };  #define SPHERES 30      //核函数内容__global__ void kernel( 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 * blockDim.x * gridDim.x;  //让图像坐标偏移DIM/2，使z轴穿过图像中心    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;  }      // globals needed by the update routine  struct DataBlock {      unsigned char   *dev_bitmap;      Sphere          *s;  };    int main( void ) {      DataBlock   data;      //记录起始时间    cudaEvent_t     start, stop;      HANDLE_ERROR( cudaEventCreate( &start ) );      HANDLE_ERROR( cudaEventCreate( &stop ) );      HANDLE_ERROR( cudaEventRecord( start, 0 ) );        CPUBitmap bitmap( DIM, DIM, &data );      unsigned char   *dev_bitmap;      Sphere          *s;          // 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 ) );        // allocate temp memory, initialize it, copy to      // memory on the GPU, then free our temp memory  //生成球面的中心坐标颜色和半径    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>>>( s, dev_bitmap );        // copy our bitmap back from the GPU for display      HANDLE_ERROR( cudaMemcpy( bitmap.get_ptr(), dev_bitmap,                                bitmap.image_size(),                                cudaMemcpyDeviceToHost ) );        // get stop time, and display the timing results      HANDLE_ERROR( cudaEventRecord( stop, 0 ) );      HANDLE_ERROR( cudaEventSynchronize( stop ) );      float   elapsedTime;      HANDLE_ERROR( cudaEventElapsedTime( &elapsedTime,                                          start, stop ) );      printf( "Time to generate:  %3.1f ms\n", elapsedTime );        HANDLE_ERROR( cudaEventDestroy( start ) );      HANDLE_ERROR( cudaEventDestroy( stop ) );        HANDLE_ERROR( cudaFree( dev_bitmap ) );      HANDLE_ERROR( cudaFree( s ) );        // display       bitmap.display_and_exit();  }

结果为：

接下来是使用常量内存后的代码，同样30个球面

#include "cuda_runtime.h"#include "device_launch_parameters.h"#include "H:\cuda_by_example\common\book.h"#include "H:\cuda_by_example\common\cpu_bitmap.h"#include "device_functions.h"#include <stdio.h>#define DIM 1024    #define rnd( x ) (x * rand() / RAND_MAX)  #define INF     2e10f    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;      }  };  #define SPHERES 30  __constant__ Sphere s[SPHERES];    __global__ void kernel( unsigned char *ptr ) {      // map from threadIdx/BlockIdx to pixel posiytion      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;// capture the start time and start to record itcudaEvent_t start,stop;HANDLE_ERROR(cudaEventCreate(&start));HANDLE_ERROR(cudaEventCreate(&stop));HANDLE_ERROR(cudaEventRecord(start,0));CPUBitmap bitmap(DIM,DIM,&data);unsigned char *dev_bitmap;//allocate the memory on the GPU for the output bitmapHANDLE_ERROR(cudaMalloc((void**)&dev_bitmap,bitmap.image_size()));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( cudaMemcpyToSymbol( s, temp_s,sizeof(Sphere) * SPHERES) );  free(temp_s);//generate a bitmap from our sphere datadim3 grids(DIM/16,DIM/16);dim3 threads(16,16);kernel<<<grids,threads>>>(dev_bitmap);//copy the bitmap back from GPU to CPU for displayHANDLE_ERROR(cudaMemcpy(bitmap.get_ptr(),dev_bitmap,bitmap.image_size(),cudaMemcpyDeviceToHost));HANDLE_ERROR(cudaEventRecord(stop,0));//stop the time recordHANDLE_ERROR(cudaEventSynchronize(stop));float elapsedTime;HANDLE_ERROR(cudaEventElapsedTime(&elapsedTime,start,stop));printf( "Time to generate:  %3.1f ms\n", elapsedTime ); HANDLE_ERROR(cudaEventDestroy(start));HANDLE_ERROR(cudaEventDestroy(stop));HANDLE_ERROR(cudaFree(dev_bitmap));bitmap.display_and_exit();}

其时间结果为：

结论：经过多次测试，发现两者的时间差距非常小，常量内存似乎没有多大作用。但是通过查阅资料发现，这应该是由于现在测试的GPU是GTX660M，已经远远好于当时书中所用的280GPU，因此在全局内存下的读取操作也非常快，而且此代码的工作量太小。但就算如此，使用常量内存的运行时间还是每次都稳稳地比不用常量内存的运行时间少几ms。

阅读全文

0 0