GPU编程前的准备工作

1、 看看自己计算机上有几块显卡

int deviceCount;

cudaGetDeviceCount(&deviceCount);

int device;

for(device = 0; device < deviceCount; ++device) 

{

    cudaDeviceProp deviceProp; 

    cudaGetDeviceProperties(&deviceProp, device);

    printf("Device %d has compute capability %d.%d. \n", device, deviceProp.major, deviceProp.minor);

}

2、指定自己用哪一块显卡

size_t size = 1024*sizeof(float);

cudaSetDevice(0); //切换到设备0

float * p0;

cudaMalloc(&p0, size); //在设备0 上分配global内存

MyKernel<<<1000, 128>>>(p0); //在设备0 上执行kernel函数

cudaSetDevice(1); //切换到设备 1

float * p1;

cudaMalloc(&p1, size); //在设备1 上分配global内存

MyKernel<<<1000, 128>>>(p1); //在设备1 上执行kernel函数

一个 Host线程可以在任何时候使用 cudaSetDevice() 来指配设备进行运算.并切换所有执行环境.分配内存,kernel launch,streams,events等,都在最近指定的设备GPU上运行. 如果没有指定则当前选择设备号 = 0.

在多GPU设备的条件下,耗时的任务可以指派给多个GPU进行运算.这是很好的.(SLI技术是多GPU完成单个任务,与这个不同)

3、 认识GPU的基本属性：

CUDA Runtime API :: CUDA Toolkit Documentation

struct cudaDeviceProp 

{

    char name[256]; // 识别设备的ASCII字符串（比如，"GeForce GTX 940M"）

    size_t totalGlobalMem; // 全局内存大小

    size_t sharedMemPerBlock; // 每个block内共享内存的大小

    int regsPerBlock; // 每个block 32位寄存器的个数

    int warpSize; // warp大小

    size_t memPitch; // 内存中允许的最大间距字节数

    int maxThreadsPerBlock; // 每个Block中最大的线程数是多少

    int maxThreadsDim[3]; // 一个块中每个维度的最大线程数

    int maxGridSize[3]; // 一个网格的每个维度的块数量

    size_t totalConstMem; // 可用恒定内存量

    int major; // 该设备计算能力的主要修订版号

    int minor; // 设备计算能力的小修订版本号

    int clockRate; // 时钟速率

    size_t textureAlignment; // 该设备对纹理对齐的要求

    int deviceOverlap; // 一个布尔值，表示该装置是否能够同时进行cudamemcpy()和内核执行

    int multiProcessorCount; // 设备上的处理器的数量

    int kernelExecTimeoutEnabled; // 一个布尔值，该值表示在该设备上执行的内核是否有运行时的限制

    int integrated; // 返回一个布尔值，表示设备是否是一个集成的GPU（即部分的芯片组、没有独立显卡等）

    int canMapHostMemory; // 表示设备是否可以映射到CUDA设备主机内存地址空间的布尔值

    int computeMode; // 一个值，该值表示该设备的计算模式：默认值，专有的，或禁止的

    int maxTexture1D; // 一维纹理内存最大值

    int maxTexture2D[2]; // 二维纹理内存最大值

    int maxTexture3D[3]; // 三维纹理内存最大值

    int maxTexture2DArray[3]; // 二维纹理阵列支持的最大尺寸

    int concurrentKernels; // 一个布尔值，该值表示该设备是否支持在同一上下文中同时执行多个内核

｝

//获取相关代码

#include  "cuda_runtime.h"

#include <iostream>

int main()

{

        cudaDeviceProp prop;

        int deviceID;

        cudaGetDevice(&deviceID);

        cudaGetDeviceProperties(&prop, deviceID);

        std::cout<<prop.deviceOverlap<<endl;

        return 0;

}

可以根据实际使用查看自己使用的GPU卡的相关情况

4、不同设备切换时出现的问题

cudaSetDevice(0); //切换到设备0

cudaStream_t s0;

cudaSreamCreate(&s0); //在当前设备0 中创建 stream s0

MyKernel<<<100,64,0,s0>>>(); //在当前设备0 中的 stream s0 中加入(异步) kernel launch指令

cudaSetDevice(1); //切换到设备1 

cudaStream_t s1; 

cudaSreamCreate(&s1); //在当前设备1 中创建 stream s1 

MyKernel<<<100,64,0,s1>>>(); //在当前设备1 中的 stream s1 中加入(异步) kernel launch指令 

// 上述代码是正确的 

// 下面这行代码会失败 

MyKernal<<<100,64,0,s0>>>(); #Error //在当前设备1 中试图往设备0 中的 stream s0加入kernel launch指令

而内存拷贝指令却与当前设备选择无关:

// 下述代码是正确的

cudaSetDevice(0); //切换到设备0

cudaStream_t s0;

cudaSreamCreate(&s0); //在当前设备0 中创建 stream s0

cudaSetDevice(1); //切换到设备1 

cudaMemcpyAsync(devMemPtr, hostMemPtr, size, cudaMemcpyHostToDevice, s0); //This is OK

udaStreamWaitEvent() 可以在多个GPU设备之间做同步.每个设备拥有自己的默认 stream (see Default Stream).所以不同 GPU设备之间的任务执行是独立无序的,你需要自己控制设备间的同步问题.

自己理解，当你切换好当下设备下，你只能启动当前设备，而其他设备你无权启动，不过其他设备你可以操作内存拷贝，cudaMemcpyAsync

5、单个GPPU使用的多任务

CUDA Stream是一个任务队列，单个CUDA Stream内部保序执行，即根据发送过来的任务依次执行；不同CUDA Stream的任务执行顺序无保证。

//在host调用，在device会按着0,1执行

Kernel_0<<<100, 512, 0, stream>>>(args...)

Kernel_1<<<200, 1024, 0, stream>>>(args...)

//这种情况下虽然在Host端是先Launch的Kernel_0，然后Launch Kernel_1，

//但是在Device端执行的时候却无法保证这两个CUDA Kernel的执行顺序，也就是cuda stream之间不保序。

Kernel_0<<<100, 512, 0, stream1>>>(args...)

Kernel_1<<<200, 1024, 0, stream2>>>(args...)

为了使用多任务，尽量使用asyncudaMemcpy(),，各自的steam可以并行操作，也就是说不同的流可以拷贝各自的内存从主机上，而

如果使用cudaMemcpy()就会出现阻塞，由于没有参数stream，所以等于阻塞host，说白了就是A流在拷贝内存的时候，B流绝对不能拷贝内存。此外单卡多任务还有以下的一些操作：

单块卡的性能：使用前要看三个并行情况

1.data transfer and kernel execution

2.concurrent kernel execution

3.concurrent data transfers

6.内存的配置

线程除了拥有local memory, share memory, global memory之外还有constant和 texture memory两个只读内存，

实际上，warp 也是 CUDA 中，每一个 SM 执行的最小单位;如果 GPU 有 16 组 SM 的话，也就代表他真正在执行的 thread 数目会是 32x16 个。不过由于 CUDA 是要透过 warp 的切换来隐藏 thread 的延迟、等待，来达到大量平行化的目的，所以会用所谓的 active thread 这个名词来代表一个 SM 里同时可以处理的 thread 数目。

　　而在 block 的方面，一个 SM 可以同时处理多个 thread block，当其中有 block 的所有 thread 都处理完后，他就会再去找其他还没处理的 block 来处理。假设有 16 个 SM、64 个 block、每个 SM 可以同时处理三个 block 的话，那一开始执行时，device 就会同时处理 48 个 block;而剩下的 16 个 block 则会等 SM 有处理完 block 后，再进到 SM 中处理，直到所有 block 都处理结束。

　　为一个多处理器指定了一个或多个要执行的线程块时，它会将其分成warp块，并由SIMT单元进行调度。将块分割为warp的方法总是相同的，每个warp都包含连续的线程，递增线程索引，第一个warp中包含全局线程过索引0-31。每发出一条指令时，SIMT单元都会选择一个已准备好执行的warp块，并将指令发送到该warp块的活动线程。Warp块每次执行一条通用指令，因此在warp块的全部32个线程执行同一条路径时，可达到最高效率。如果一个warp块的线程通过独立于数据的条件分支而分散，warp块将连续执行所使用的各分支路径，而禁用未在此路径上的线程，完成所有路径时，线程重新汇聚到同一执行路径下，其执行时间为各时间总和。分支仅在warp块内出现，不同的warp块总是独立执行的–无论它们执行的是通用的代码路径还是彼此无关的代码路径。

建议的数值?

　　在 Compute Capability 1.0/1.1 中，每个 SM 最多可以同时管理 768 个 thread(768 active threads)或 8 个 block(8 active blocks);而每一个 warp 的大小，则是 32 个 thread，也就是一个 SM 最多可以有 768 / 32 = 24 个 warp(24 active warps)。到了 Compute Capability 1.2 的话，则是 active warp 则是变为 32，所以 active thread 也增加到 1024。

　　在这里，先以 Compute Capability 1.0/1.1 的数字来做计算。根据上面的数据，如果一个 block 里有 128 个 thread 的话，那一个 SM 可以容纳 6 个 block;如果一个 block 有 256 个 thread 的话，那 SM 就只能容纳 3 个 block。不过如果一个 block 只有 64 个 thread 的话，SM 可以容纳的 block 不会是 12 个，而是他本身的数量限制的 8 个。

　　因此在 Compute Capability 1.0/1.1 的硬件上，决定 block 大小的时候，最好让里面的 thread 数目是 warp 数量(32)的倍数(可以的话，是 64 的倍数会更好);而在一个 SM 里，最好也要同时存在复数个 block。如果再希望能满足最多 24 个 warp 的情形下，block 里的 thread 数目似乎会是 96(一个 SM 中有 8 个 block)、128(一个 SM 中有 6 个 block)、192(一个 SM 中有 4 个 block)、256(一个 SM 中有 3 个 block) 这些数字了～

　　Compute Capability 1.0/1.1 的硬件上，每个grid最多可以允许65535×65535个block。每个block最多可以允许512个thread，但是第三维上的最大值为64。而官方的建议则是一个 block 里至少要有 64 个 thread，192 或 256 个也是通常比较合适的数字(请参考 Programming Guide)。

　　但是是否这些数字就是最合适的呢?其实也不尽然。因为实际上，一个 SM 可以允许的 block 数量，还要另外考虑到他所用到 SM 的资源：shared memory、registers 等。在 G80 中，每个 SM 有 16KB 的 shared memory 和 8192 个 register。而在同一个 SM 里的 block 和 thread，则要共享这些资源;如果资源不够多个 block 使用的话，那 CUDA 就会减少 Block 的量，来让资源够用。在这种情形下，也会因此让 SM 的 thread 数量变少，而不到最多的 768 个。

　　比如说如果一个 thread 要用到 16 个 register 的话(在 kernel 中宣告的变量)，那一个 SM 的 8192 个 register 实际上只能让 512 个 thread 来使用;而如果一个 thread 要用 32 个 register，那一个 SM 就只能有 256 个 thread 了～而 shared memory 由于是 thread block 共享的，因此变成是要看一个 block 要用多少的 shread memory、一个 SM 的 16KB 能分给多少个 block 了。

　　所以虽然说当一个 SM 里的 thread 越多时，越能隐藏 latency，但是也会让每个 thread 能使用的资源更少。因此，这点也就是在优化时要做取舍的了。