第六章 常量内存与事件 
光线追踪、常量内存与事件、常量内存带来的性能提升、线程束
GitHub 
建议下载下来用Typora软件阅读markdown文件
作者github:littlebearsama 原文链接 
(建议下载Typora来浏览markdown文件) 
第六章 常量内存与事件 0.光线追踪 
常量内存用于保存在核函数执行期间不会发生变化的数据 。Nvidia硬件提供了64KB的常量内存,并且对常量内存采取了不同于标准全局内存的处理方式。在某些情况中,用常量内存来替换全局内存能有效地减少内存带宽。 
在光线跟踪的例子中,没有利用常量内存的代码运行时间为1.8ms,利用了常量内存的代码运行时间为0.8ms 
将球面数组存入常量内存中。 
 
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 42 43 44 45 46 47 48 49 50 51 52 53 54 55 56 57 58 59 60 61 62 63 64 65 66 67 68 69 70 71 72 73 74 75 76 77 78 79 80 81 82 83 84 85 86 87 88 89 90 91 92 93 94 95 96 97 98 99 100 101 102 103 104 105 106 107 108 109 110 111 112 113 114 115 116 117 118 119 #include  "cuda.h"  #include  "../common/book.h"  #include  "../common/image.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 20 __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 ; } struct  DataBlock  {    unsigned  char    *dev_bitmap; }; int  main ( void  )      DataBlock   data;          cudaEvent_t     start, stop;     HANDLE_ERROR( cudaEventCreate( &start ) );     HANDLE_ERROR( cudaEventCreate( &stop ) );     HANDLE_ERROR( cudaEventRecord( start, 0  ) );     IMAGE bitmap ( DIM, DIM )  ;     unsigned  char    *dev_bitmap;          HANDLE_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 );          dim3    grids (DIM/16 ,DIM/16 )  ;     dim3    threads (16 ,16 )  ;     kernel<<<grids,threads>>>( dev_bitmap );          HANDLE_ERROR( cudaMemcpy( bitmap.get_ptr(), dev_bitmap,                               bitmap.image_size(),                               cudaMemcpyDeviceToHost ) );          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 ) );          bitmap.show_image(); } 
变量前面加上__constant____constant__ Sphere s[SPHERES];常量内存为静态分配空间,所以不需要调用 cudaMalloc(), cudaFree(); 
在主机端分配临时内存,对其初始化Sphere *temp_s = (Sphere*)malloc( sizeof(Sphere) * SPHERES );在把变量复制到常量内存后释放内存free( temp_s ); 
使用函数cudaMemcpyToSymbol()cudaMemcpyToSymbol()全局内存 ) 
 
1.常量内存带来的性能提升 与从全局内存中读取数据相比,从常量内存中读取相同的数据可以节约带宽,原因有二:
对常量内存的单次读操作可以广播到其他的“邻近”线程,这将节约15次读取操作。 
常量内存的数据将缓存(cache)起来,因此对相同地址的连续读操作将不会产生额外的内存通信量。 
 
2.线程束Warp 在CUDA架构中,线程束是指一个包含32个线程的集合,这个线程集合被“编织在一起”并且以“步调一致(Lockstep)”的形式执行。在程序中的每一行,线程束中的每个线程都将在不同数据上执行相同的指令。
线程束 当处理常量内存是,NVIDIA硬件将把单次内存读取操作 广播到每半个线程束(Half-Warp) 。在半线程束中包含了16和线程。如果在半线程束中的每个线程都从常量内存的相同地址上读取数据 ,那么GPU只会产生一次读取请求 并在随后将数据广播到每个线程 。如果从常量内存中读取大量的数据,那么这种方式生产的内存流量只是全局内存的1/16(大约6%)。
常量内存与缓存 但在读取常量内存是,所节约的并不只限于减少94%的带宽。由于这块内存的内容是不会发生变化的,因此硬件将主动把这个常量数据缓存在GPU上。 在第一次从常量内存的某个地址上读取后,当其他半线程束请求同一地址是,那么将命中缓存(cahce) ,这同样减少了额外的内存流量。在光线追踪程序中,将球面数据保存在常量内存后,硬件只需要请求这个数据一次。在缓存数据后,其他每个线程将不会产生内存流量,原因有两个: 1. 线程将在半线程结束的广播中收到这个数据。 2. 从常量内存缓存中收到数据。
负面影响 当使用常量内存是,也可能对性能产生负面影响。半线程束广播功能实际是把双刃剑 。虽然当所有16个线程地址都读取相同地址是,这个功能可以极大地提高性能,但当所有16个线程分别读取不同地址时,它实际上会降低性能。
3.使用事件来测试性能 代码:
1 2 3 4 5 6 7 8 9 10 11 12 13 14 cudaEvent_t     start, stop; cudaEventCreate( &start ); cudaEventCreate( &stop ); cudaEventRecord( start, 0  ); cudaEventRecord( stop, 0  ); cudaEventSynchronize( stop ); float    elapsedTime;cudaEventElapsedTime( &elapsedTime,start, stop ); printf ( "Time to generate:  %3.1f ms\n" , elapsedTime );cudaEventDestroy( start ); cudaEventDestroy( stop ); 
运行记录事件start时,还指定了第二个参数。cudaEventRecord( start, 0 );在上面代码中为0,流(Stream)的编号。 
当且仅当GPU完成了之间的工作并且记录了stop事件后,才能安全地读取stop时间值。幸运的是,还有一种方式告诉CPU在某个事件上同步,这个时间API函数就是cudaEventSynchronize();, 当cudaEventSynchronize返回时,我们知道stop事件之前的所有GPU工作已经完成了,因此可以安全地读取在stop保存的时间戳。 
由于CUDA事件是直接在GPU上实现的,因此它们不适用于对同时包含设备代码和主机代码的混合代码计时。也就是说,你通过CUDA事件对核函数和设备内存复制之外的代码进行计时,将得到不可靠的结果 。