GPU高性能编程CUDA实战

书中摘录+稍微的补充

1、CPU及系系统的内存称为主机，将GPU及其内存称为设备，在GPU设备上执行的函数通常称为核函数(kernel)；

2、CPU并行线程结构

《1》、线程 –> 线程块 –>（线程格）grid ；


《2》、关键的内置变量：
threadIdx：3维向量，标识线程。每个线程块内的threadIdx是唯一的；
blockIdx：3维向量，标识线程块，每个grid内的blockIdx是唯一的；
blockDim：3维向量，存储线程块每个维度线程的数量信息；
gridDim: 3维向量，存储线程格每个维度线程块的数量信息

实际执行核函数的线程的tid计算(一维情况)： tid = threadIdx.x + blockIdx.x * blockDim.x

3、核函数带有修饰符: _ global _

__global__ void kernel(int a)
{
    a;
}

核函数的调用带有修饰符<<<>>>,在主机端调用kernel如下：

kennel_name<<<Dg, Db, Ns, S>>>([kernel arguments])

Dg: dim3类型，表示使用的grid的纬度和大小信息，既是设备在执行该函数时使用的并行线程块的数量；
Db: dim3类型，存储为每个线程块的纬度及大小信息，既是设备在执行该函数时每个线程块内的线程数量；
Ns: size_t类型，可选项，默认为0；
S: cudaStream_t类型，cuda使用的流。可选项，默认是0；

例子：

__global__ void add(int a, int b, int *c)
{
    *c = a + b;
}

int main(void)
{
    int *dev_c;
    .....
    add<<<1,1>>>(2,7,dev_c);
    .....
    return 0;
}

4、CUDA架构上GPU标准内存（又称linear memory）分配（还有一些特殊内存如常量内存、页锁定内存等不涉及）

《1》、标准内存分配与释放：cudaMalloc()/cudaFree()
[还有其它的分配方式如： cudaMallocPitch() 、 cudaMalloc3D() ]
《2》、主机与设备之间内存数据拷贝函数：cudaMemcpy()
cudaMemcpy函数需要传进标志位来识别数据的流动方向：
cudaMemcpyHostToDevice — 数据从主机到设备
cudaMemcpyDeviceToHost — 数据从设备到主机

// Allocate vectors in device memory
float* d_A;
cudaMalloc(&d_A, size);
float* d_B;
cudaMalloc(&d_B, size);
float* d_C;
cudaMalloc(&d_C, size);

// Copy vectors from host memory to device memory
cudaMemcpy(d_A, h_A, size, cudaMemcpyHostToDevice);
cudaMemcpy(d_B, h_B, size, cudaMemcpyHostToDevice);
cudaMemcpy(h_C, d_C, size, cudaMemcpyDeviceToHost);

// Free device memory
cudaFree(d_A);
cudaFree(d_B);
cudaFree(d_C);

5、设备指针的使用限制

设备指针：指向GPU内存地址的指针；
限制：只能在GPU上操作其所指内存，不能再主机上操作设备指针所指内容
，具体如下：
《1》、可以将cudaMalloc()分配的指针传达给在设备上执行的函数；
《2》、可以在设备代码中使用cudeMalloc()分配的指针进行内存读/写操作；
《3》、可以将cudaMalloc()分配的指针传递给在主机上执行的函数；
《4》、不能再主机代码使用cudaMalloc()分配的指针进行内存读/写操作；

6、共享内存与同步

《1》、共享内存声明关键字：_share_
共享内存能被线程块内的所有线性访问,编译器将为每个线程块生成共享内存的一个副本。
《2》、共享内存缓存区位于物理GPU上而不是驻留在GPU之外，故在访问共享内存时的延迟要远远低于访问普通缓冲区的延迟；
《3》、共享内存使用时需要进行线程之间的同步: __syncthreads()函数
该函数将确保线程块中的每个线程都执行完__syncthreads()前面的语句后，才会执行下一条语句。

const int threadsPerBlock = 256;

__global__ void test(float *a , float* b, float *c)
{
    __share__ float cache[threadsPerBlock];
    int tid = threadIdx.x + blockIdx.x * blockDim.x;

    int cacheIndex = threadIdx.x;   
    //共享内存缓存中的偏移就等于线性索引，因为每个线程块都拥有该共享内存的私有副本

    float temp = 0;
     while（tid < N）
    {    
        temp += a[i] + b[i];
        tid += blockDim.x * gridDim.x;
    }

    cache[cacheIndex] = temp;

    //对线程块中的线程进行同步
    __syncthreads();

}

线程发散(Thread Divergence): 当某些线程需要执行一条指令而其他线程不需要执行是，这种情况称为线程发散。
* 如果__syncthreads()位于线程发散的分支，那GPU程序会一直保持等待，不会结束
*

7、两类只读内存

a、常量内存（constant memory spaces）– 只读

《1》、 常量内存一般用于保存在核函数执行期间不会发生变化的数据，NVIDIA硬件提供了64KB的常量内存，并且对常量内存采取了不同于标准全局内存的处理方式。在某些情况中，用常量内存来替换全局内存能有效的减少内存带宽；
减少带宽原因：
对常量内存的单次读操作可以广播到其它的“邻近（Nearby）”线程，这将节约15次读取操作；
常量内存的数据将缓存起来，因此对相同的地址的连续读操作将不会产生额外的内存通信量。
“邻近（Nearby）”线程：线程束概念【一个包含32个线程的集合–cuda中定义】

《2》、常量内存声明修饰符关键字: _constant_

#define SPHERES 20

struct Sphere {
     int a;
    float b;
};

__constant__ Sphere s[SPHERES];

__global__ void test()
{
    //do something
}
int main( void ) 
{
    // allocate temp memory, initialize it, copy to constant
    // 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;
    }

    cudaMemcpyToSymbol( s, temp_s, sizeof(Sphere) * SPHERES);
    free( temp_s );

    return 0;
}

《3》、常量内存的拷贝使用cudaMemcpyToSymbol()函数；

b、纹理内存( texture and surface memory spaces)【书上简单介绍，详细可看《cuda c programing guide》的3.11部分】

《1》、只读内存，在特定的访问模式中，纹理内存同样能够提升性能并减少内存流量。纹理内存时专门为那些在内存访问模式中存在大量空间局部性(Spatial Locality)的图形应用程序而设计的。在某个计算应用程序中，这意味着这一个线程读取的位置可能与邻近线程读取的位置“非常接近”。
纹理内存分3种：一维/二维/三维纹理内存；
《2》、声明texture类型的引用API：texture

texture<float> texConstSrc;
texture<int> textIn;

texture 引用只能被声明为全局静态变量，不能将其作为参数传递给函数。

例子：

// these exist on the GPU side
texture<float>  texConstSrc;

__global__ void copy_const_kernel( float *iptr ) {
    // 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 c = tex1Dfetch(texConstSrc,offset);//读取一维纹理内存的方式
    if (c != 0)
        iptr[offset] = c;
}

int main( void )
{
   int imageSize = 1024;
   float* dev_inSrc = NULL;

   cudaMalloc( (void**)&dev_inSrc, imageSize );
   cudaBindTexture( NULL, texConstSrc,dev_inSrc, imageSize );//必须绑定

    return 0;
｝

至此CUDA 内存包括：全局内存，共享内存，常量内存，纹理内存 共四类。

8、事件机制用于测量性能（耗时）

cuda中的事件本质上是一个GPU时间戳，这个时间戳实在用户指定的事件点上的记录。GPU本身支持记录时间戳，避免了很多与CPU定时器一起统计的麻烦。

事件的创建及使用：

cudaEvent_t start, stop;
cudaEventCreate(&start);
cudaEventCreate(&stop);

cudaEventRecord(start,0);  //第二参数指定使用的流，默认0

//kernel function

cudaEventRecord(stop,0);
cudaEventSynchronize(stop);  //同步  

float elapsedTime = 0.0f; 
cudaEventElapsedTime(&elapsedTime, start, stop);
printf("Cost Time: %3.1f ms \n", elapsedTime);  //毫秒级

9、计算功能集(Compute Capability)：NVIDIA GPU支持的各种功能的统称。

10、原子操作(像Linux 下的原子操作)

11、页锁定内存(Page-locked memory)

《1》、页锁定内存的一个重要属性：操作系统不会把这块内存分页，也不会将其交换到磁盘上，从而确保了该内存始终驻留在物理内存中。该内存的物理地址是能访问的，因为这块内存将不会被破坏或重新定位。
《2》、GPU知道内存的物理地址，因此可以通过DMA技术来在GPU和主机页锁定内存之间复制数据。在GPU和主机之间进行DMA数据传输时，使用页锁定内存会比使用标准内存的性能快很多（大约2倍）。
使用页锁定内存的一个不足是：使用页锁定内存会更快的耗尽系统内存。所以使用页锁定内存应该有针对性不能随意使用，一般来说仅对cudaMemcpy()函数调中的源内存或目标内存才使用页锁定内存
。

《3》、通过cudaHostAlloc函数分配， cudaFreeHost()释放；

float cuda_host_alloc_test( int size, bool up ) {
    cudaEvent_t    start, stop;
    int         *a, *dev_a;
    float        elapsedTime;

    HANDLE_ERROR( cudaEventCreate( &start ) );
    HANDLE_ERROR( cudaEventCreate( &stop ) );

    HANDLE_ERROR( cudaHostAlloc( (void**)&a,size * sizeof( *a ),cudaHostAllocDefault ) );
    HANDLE_ERROR( cudaMalloc( (void**)&dev_a,size * sizeof( *dev_a ) ) );

    HANDLE_ERROR( cudaEventRecord( start, 0 ) );
    for (int i=0; i<100; i++) {
        if (up)
            HANDLE_ERROR( cudaMemcpy( dev_a, a,size * sizeof( *a ),cudaMemcpyHostToDevice ) );
        else
            HANDLE_ERROR( cudaMemcpy( a, dev_a,size * sizeof( *a ),cudaMemcpyDeviceToHost ) );
    }
    HANDLE_ERROR( cudaEventRecord( stop, 0 ) );
    HANDLE_ERROR( cudaEventSynchronize( stop ) );
    HANDLE_ERROR( cudaEventElapsedTime( &elapsedTime,start, stop ) );

    HANDLE_ERROR( cudaFreeHost( a ) );
    HANDLE_ERROR( cudaFree( dev_a ) );
    HANDLE_ERROR( cudaEventDestroy( start ) );
    HANDLE_ERROR( cudaEventDestroy( stop ) );

    return elapsedTime;
}

12、cuda流

《1》、cuda流表示一个GPU操作队列，并且该队列中的操作将以指定的顺序执行。cuda流在加速应用程序方面起着重要作用，可以将每个流视为GPU上的一个任务，并且这些任务是可以并行执行的。将操作添加到流的顺序也就是它们的执行顺序。
《2》、设备重叠功能（Device Overlap）:支持设备重叠功能的GPU能够在执行一个CUDA C核函数的同事，还能在设备与主机之间执行复制操作。

《3》、从逻辑上看，不同流之间是相互独立的，但事实上这种理解并不完全符合GPU的队列机制。在硬件中并没有流的概念而是包含一个或多个引擎来执行内存复制操作，以及一个引擎来执行核函数。


编程的时候需要按照硬件处理逻辑来进行，否则无法实现并行。