CUDA使用的一些Tips

1.查询GPU的配置参数

直接运行下面代码获得各项参数：

#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <stdio.h>

int main() {
    int deviceCount;
    cudaGetDeviceCount(&deviceCount);

    int dev;
    for (dev = 0; dev < deviceCount; dev++)
    {
        int driver_version(0), runtime_version(0);
        cudaDeviceProp deviceProp;
        cudaGetDeviceProperties(&deviceProp, dev);
        if (dev == 0)
            if (deviceProp.minor = 9999 && deviceProp.major == 9999)
                printf("\n");
        printf("\nDevice%d:\"%s\"\n", dev, deviceProp.name);
        cudaDriverGetVersion(&driver_version);
        printf("CUDA驱动版本:                                   %d.%d\n", driver_version / 1000, (driver_version % 1000) / 10);
        cudaRuntimeGetVersion(&runtime_version);
        printf("CUDA运行时版本:                                 %d.%d\n", runtime_version / 1000, (runtime_version % 1000) / 10);
        printf("设备计算能力:                                   %d.%d\n", deviceProp.major, deviceProp.minor);
        printf("Total amount of Global Memory:                  %u bytes\n", deviceProp.totalGlobalMem);
        printf("Number of SMs:                                  %d\n", deviceProp.multiProcessorCount);
        printf("Total amount of Constant Memory:                %u bytes\n", deviceProp.totalConstMem);
        printf("Total amount of Shared Memory per block:        %u bytes\n", (int)deviceProp.sharedMemPerBlock);
        printf("Total number of registers available per block:  %d\n", deviceProp.regsPerBlock);
        printf("Warp size:                                      %d\n", deviceProp.warpSize);
        printf("Maximum number of threads per SM:               %d\n", deviceProp.maxThreadsPerMultiProcessor);
        printf("Maximum number of threads per block:            %d\n", deviceProp.maxThreadsPerBlock);
        printf("Maximum size of each dimension of a block:      %d x %d x %d\n", deviceProp.maxThreadsDim[0],
            deviceProp.maxThreadsDim[1],
            deviceProp.maxThreadsDim[2]);
        printf("Maximum size of each dimension of a grid:       %d x %d x %d\n", deviceProp.maxGridSize[0], deviceProp.maxGridSize[1], deviceProp.maxGridSize[2]);
        printf("Maximum memory pitch:                           %u bytes\n", deviceProp.memPitch);
        printf("Texture alignmemt:                              %u bytes\n", deviceProp.texturePitchAlignment);
        printf("Clock rate:                                     %.2f GHz\n", deviceProp.clockRate * 1e-6f);
        printf("Memory Clock rate:                              %.0f MHz\n", deviceProp.memoryClockRate * 1e-3f);
        printf("Memory Bus Width:                               %d-bit\n", deviceProp.memoryBusWidth);
        
    }

    return 0;
}

2.性能提升的Tips

1.取消Debug information

进入到项目的属性选项中，选择CUDA C/C++中Device选项，反选Generate GPU Debug information

2.修改计算能力

首先查看自己设备所支持的计算能力（我的设备是6.0），接着在上图的属性页更改code generation为compute_60, sm_60，数字与设备所支持的最高计算能力相匹配，这样我们就可以使用一些更高性能的指令。

对SM_60的补充说明：
版本越高说明流处理器同时处理的流越多，如下图所示。
SM可以同时拥有多个处理流

注意：我尝试过设置为6.0以上的其他版本，例如：“compute_80, sm_80”，虽然代码能运行，但是GPU的计算结果会出问题，会有计算错误。所以还是根据自己的设备来设置。

3.计算GPU的运行时间

通常的CPU获取时间的函数不能应用在GPU上，比如下面的例子是不行的：

DWORD s1 = GetTickCount();
vecAdd <<<(N + 255) / 256, 256 >>> (N, devA, devB, devC);
DWORD e1 = GetTickCount();
cputime = (double)(e1 - s1) * 1.00 ;
printf("time for GPU is %f ms \n", cputime);

第二行代码只是CPU去通知GPU开始进行运算，通过在函数两端获取时间只会得到一个几乎是一瞬间的值，无论GPU将要做多大的运算。
正确的方法是利用CUDA的API函数，创建时间戳并记录时间，计算时间差，完整代码如下：

	cudaEvent_t time1, time2;
	cudaEventCreate(&time1);	cudaEventCreate(&time2);//创建事件
	
	cudaEventRecord(time1, 0);//记录时间戳
	vecAdd << <(N + 255) / 256, 256 >> > (N, devA, devB, devC);
	cudaEventRecord(time2, 0);
	
	cudaEventSynchronize(time1);	cudaEventSynchronize(time2);//等待时间戳前的线程同步
	
	cudaEventElapsedTime(&kernalExecutionTime, time1, time2);//计算时间差
	printf("time for GPU is %f ms \n", kernalExecutionTime);//输出
	
	cudaEventDestroy(time1);	cudaEventDestroy(time2);//销毁事件

4.在核函数中调用核函数

在项目属性中, 设定CUDA C/C++的Generate Relocatable Device Code为True(-rdc=true).