第六章 常量内存与事件
光线追踪、常量内存与事件、常量内存带来的性能提升、线程束
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()
将变量从主机内存复制到GPU上的常量内存。(cudaMencpyHostToDevice()的cudaMemcpy()之间的唯一差异在于,cudaMemcpyToSymbol()
会复制到常量内存,而cudaMemcpy()会复制到全局内存 )
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事件对核函数和设备内存复制之外的代码进行计时,将得到不可靠的结果 。