nvidia-smi -q -i 0 #只显示0卡信息
nvidia-smi -q -i 0 -d MEMORY | tail -n 5 #只显示0卡内存信息
nvidia-smi -q -i 0 -d UTILIZATION | tail -n 4 #只显示0卡使用率
nvidia-smi -a -q -d CLOCK | fgrep -A 3 "Max Clocks" | fgrep "Memory" #检测设备的内存频率
#-d参数：
# MEMORY
# UTILIZATION
# ECC
# TEMPERATURE
# POWER
# CLOCK
# COMPUTE
# PIDS
# PERFORMANCE
# SUPPORTED_CLOCKS
# PAGE_RETIREMENT
# ACCOUNTING

#nvidia驱动程序将只使用ID为2和3的设备，并且会将设备ID分别映射为0和1
export CUDA_VISIBLE_DEVICES=2,3

//Nsight加注释
#include "nvToolsExt.h"
nvtxRangePushA(__FUNCTION__);
nvtxRangePop();

#nsight system性能分析
nsys profile -t cuda, osrt, nvtx, cudnn, cublas \ #需要追踪的API
    -y 60 \                                       #延迟分析(单位：秒)
    -d 20 \                                       #分析持续的区间(单位：秒)
    -o baseline \                                 #输出文件名
    -f true \                                     #是否覆盖源文件
    -w true \                                     #是否显示
    ./exe                                         #执行程序命令
#nvprof 性能分析，SM7.0以上用 ncu
nvporf --metrics branch_efficiency ./exe #检查内核线程束分化
nvprof --events branch,divergent_branch ./exe #获得分支和分化分支的事件计数器
nvprof --metrics achieved_occupancy ./exe #获得一个内核的可实现占用率：每周期内活跃线程束的平均数量与一个SM支持的线程束最大数量的比值
nvprof --metrics gld_throughput ./exe #检查内核的内存读取效率
nvprof --metrics gld_efficiency ./exe #检测全局加载效率，即被请求的全局加载吞吐量占所需的全局加载吞吐量的比值。它衡量了应用程序的加载操作利用设备内存带宽的程度
nvprof --devices 0 --metrics gld_transactions ./exe $OFFSET #显示全局加载事务数量，OFFSET去掉显示的是全局内存加载效率
nvprof --devices 0 --metrics gld_efficiency,gst_efficiency ./exe #获取全局加载效率和全局存储效率指标

const char* cudaGetErrorString ( cudaError_t error );
//cuda函数调用时，增加这个函数打印错误
#define CHECK(call) \
{ \
 const cudaError_t error = call; \
 if (error != cudaSuccess) \
 { \
 printf("Error: %s:%d, ", __FILE__, __LINE__); \
 printf("code:%d, reason: %s\n", error, cudaGetErrorString(error)); \
 exit(1); \
 } \
}


//核函数定义
kernel_name <<<grid, block>>>(argument list);
__global__ void kernel_name(argument list);

//CUDA核函数的限制：
//·只能访问设备内存
//·必须具有void返回类型
//·不支持可变数量的参数
//·不支持静态变量
//·显示异步行为

//获取GPU显卡信息
cudaError_t cudaGetDeviceProperties(cudaDeviceProp* prop, int device);

const char* cudaGetErrorString ( cudaError_t error );
//cuda函数调用时，增加这个函数打印错误
#define CHECK(call) \
{ \
 const cudaError_t error = call; \
 if (error != cudaSuccess) \
 { \
 printf("Error: %s:%d, ", __FILE__, __LINE__); \
 printf("code:%d, reason: %s\n", error, cudaGetErrorString(error)); \
 exit(1); \
 } \
}
//核函数定义
kernel_name <<<grid, block>>>(argument list);
__global__ void kernel_name(argument list);
//CUDA核函数的限制：
//·只能访问设备内存
//·必须具有void返回类型
//·不支持可变数量的参数
//·不支持静态变量
//·显示异步行为
//获取GPU显卡信息
cudaError_t cudaGetDeviceProperties(cudaDeviceProp* prop, int device);

//CUDA编程的一个基本原则应是尽可能 地减少主机与设备之间的传输。

3. 显卡性能计算，以Tesla K10为例：

单精度峰值浮点运算次数： 745 MHz核心频率*2 GPU/芯片*（8个多处理器*192个浮点单元*32核心/多处理器）*2 OPS/周期＝4.58 TFLOPS（FLOPS表示每秒浮点运算次数）

内存带宽峰值：2 GPU/芯片*256位*2500 MHz内存时钟*2 DDR/8位/字节＝320 GB/s

指令比:字节 ： 4.58 TFLOPS/320 GB/s，也就是13.6个指令:1个字节

如果你的应用程序每访问一个字节所产生的指令数多于13.6，那 么你的应用程序受算法性能限制。大多数HPC工作负载受内存带宽的限制

二. 线程结构

　　一个内核启动所产生的所有线程统称为一个网格(Grid)，同一网格中的所有线程共享相同的全局内存空间;

       一个网格(Grid)由多个线程块(Block)构成，一个线程块(Block)包含一组线程(Thread);

       同一线程块(Block)内的线程(Thread)协作可以通过同步和共享内存来实现，不同块内的线程不能协作。

       blockIdx：线程块Block在网格Grid内的索引，blockIdx.x表示block的x坐标。

　　threadIdx：线程块Block内的线程Thread索引，同理blockIdx。

　　blockDim：线程块block维度，用每个线程块中的线程数来表示，对应的是threadIdx.

　　gridDim：线程网格grid维度，用每个线程网格中的线程数来表示，对应的是blockIdx。

一般会把grid组织成2D，block为3D。grid和block都使用dim3作为声明，例如：

　　dim3 block(3);

       int grid_sizie = (nElem+block.x-1)/block.x; //[nElem/block.x] 向上取整

　　dim3 grid(grid_sizie); 

dim3 block(dimx, dimy); 

dim3 grid((nx + block.x - 1) / block.x, (ny + block.y - 1) / block.y);

需要注意的是，dim3仅为host端可见，其对应的device端类型为uint3。

线程块和线程网格的划分方法：

由二维线程块构成的二维网格

    ix = threadIdx.x + blockIdx.x * blockDim.x 

    iy = threadIdx.y + blockIdx.y * blockDim.y

    idx = iy * nx + ix

由一维线程块构成的一维网格 

        ix = threadIdx.x + blockIdx.x * blockDim.x

        idx = ix

由一维线程块构成的二维网格

    ix = threadIdx.x + blockIdx.x * blockDim.x 

    iy = blockIdx.y

    idx = iy * nx + ix

三. GPU执行模型

//当__syncthreads被调用时，在同一个线程块中每个线程都必须等待直至该线程块中所有其他线程都已经达到这个同步点
__device__ void __syncthreads(void);

cudaError_t cudaFuncSetCacheConfig(const void* func, enum cudaFuncCache cacheConfig);
//常量内存初始化 将count个字节从src指向的内存复制到symbol指向的内存中
cudaError_t cudaMemcpyToSymbol(const void* symbol, const void* src, size_t count);
//读取常量内存 从常量内存中拷贝数据
cudaError_t cudaMemcpyFromSymbol(void *dst, const void *symbol, size_t count, size_t offset = 0, cudaMemcpyKind kind = cudaMemcpyDeviceToHost);

//如何Malloc函数执行失败则返回cudaErrorMemoryAllocation。在已分配的全局内存中的值不会被清除，需要调用cudaMemset来初始化显存
cudaError cudaMalloc(void** devPtr, size_t size);
 //cudaMemcpyKind:
 //cudaMemcpyHostToHost       //CPU->CPU
 //cudaMemcpyHostToDevice     //CPU->GPU
 //cudaMemcpyDeviceToHost     //GPU->CPU
 //cudaMemcpyDeviceToDevice   //GPU->GPU
cudaError cudaMemcpy(void* dst, const void* src, size_t count, cudaMemcpyKind kind);
//异步数据拷贝
cudaError_t cudaMemcpyAsync(void* dst, const void* src, size_t count, cudaMemcpyKind kind, cudaStream_t stream = 0);
cudaError_t cudaMemset(void *devPtr, int value, size_t count);
cudaError_t cudaFree(void *devPtr);

/*静态声明一个全局变量并使用*/
__device__ float devData;
//初始化全局变量
cudaError_t cudaMemcpyToSymbol(devData, &value, sizeof(float));
//获取GPU全局变量地址
cudaError_t cudaGetSymbolAddress(void** devPtr, const void* symbol);
//拷贝全局变量到host端
cudaError_t cudaMemcpyFromSymbol(&value, devData, sizeof(float));

//创建一个零拷贝内存, 用cudaFreeHost释放
    //flag:
    //cudaHostAllocDefalt:			使cudaHostAlloc函数的行为与cudaMallocHost函数一致
    //cudaHostAllocPortable:		返回能被所有CUDA上下文使用的固定内存，而不仅是执行内存分配的那一个
    //cudaHostAllocWriteCombined:	返回写结合内存，该内存可以在某些系统配置上通过PCIe总线上更快地传输，但是它在大多数主机上不能被有效地读取
    //cudaHostAllocMapped:			返回可以实现主机写入和设备读取被映射到设备地址空间中的主机内存
cudaError_t cudaHostAlloc(void **pHost, size_t count, unsigned int flags);
//获取映射到固定内存的device指针
cudaError_t cudaHostGetDevicePointer(void **pDevice, void *pHost, unsigned int flags);

//托管内存，托管变量的声明
__device__ __managed__ int y;
//cuda6.0以后废弃，使用托管内存的程序可以利用自动数据传输和重复指针消除功能
cudaError_t cudaMallocManaged(void **devPtr, size_t size, unsigned int flags=0);

//分配固定主机内存
cudaError_t cudaMallocHost(void **devPtr, size_t count);
//释放固定主机内存
cudaError_t cudaFreeHost(void *ptr);

// allocate zero-copy memory at the host side
cudaHostAlloc((void **)&h_A, nBytes, cudaHostAllocMapped);
cudaHostAlloc((void **)&h_B, nBytes, cudaHostAllocMapped);
// initialize data at the host side
initialData(h_A, nElem);
initialData(h_B, nElem);
//不需要调用cudaHostGetDevicePointer
// invoke the kernel with zero-copy memory
sumArraysZeroCopy<<<grid, block>>>(h_A, h_B, d_C, nElem);

//Cuda Stream使用
cudaError_t cudaStreamCreate(cudaStream_t* pStream);
//当执行异步数据传输时，必须使用固定（或非分页的）主机内存
cudaError_t cudaMemcpyAsync(void* dst, const void* src, size_t count, cudaMemcpyKind kind, cudaStream_t stream = 0);
kernel_name<<<grid, block, sharedMemSize, stream>>>(argument list);
cudaError_t cudaStreamDestroy(cudaStream_t stream);
//强制阻塞主机，直到在给定流中所有的操作都完成
cudaError_t cudaStreamSynchronize(cudaStream_t stream);
//检查流中所有操作是否都已经完成
//当所有操作都完成时cudaStreamQuery函数会返回cudaSuccess
//当一个或多个操作仍在执行或等待执行时返回cudaErrorNotReady
cudaError_t cudaStreamQuery(cudaStream_t stream);
//创建带有优先级的流
//流优先级不会影响数据传输操作，只对计算内核有影响
cudaError_t cudaStreamCreateWithPriority(cudaStream_t* pStream, unsigned int flags, int priority);
//查询流可设置的优先级
//如果当前的设备不支持流优先级，cudaDeviceGetStreamPriorityRange将0返回给这两个参数
cudaError_t cudaDeviceGetStreamPriorityRange(int *leastPriority, int *greatestPriority);
//创建同步流和异步流
//flags:
//cudaStreamDefault: 默认流 (阻塞)
//cudaStreamNonBlocking: 异步流 (非阻塞)，使得非空流对于空流的阻塞行为失效
cudaError_t cudaStreamCreateWithFlags(cudaStream_t* pStream, unsigned int flags);

//流回调
cudaError_t cudaStreamAddCallback(cudaStream_t stream, cudaStreamCallback_t callback, void *userData, unsigned int flags);

    2. cuda事件

    CUDA中事件本质上是CUDA流中的标记，它与该流内操作流中的特定点相关联，可以使用事件来执行以下两个基本任务： ·同步流的执行 ·监控设备的进展

    只有当一个给定CUDA流中先前的所有操作都执行结束后，记录在该流内的事件才会起作用

//cuda事件
// 事件在流执行中标记了一个点。它们可以用来检查正在执行的流操作是否已经到达了给定点
cudaEvent_t event;
//创建事件
cudaError_t cudaEventCreate(cudaEvent_t* event);
//销毁事件
cudaError_t cudaEventDestroy(cudaEvent_t event);
//一个事件排队进入CUDA流
cudaError_t cudaEventRecord(cudaEvent_t event, cudaStream_t stream = 0);
//等待一个事件并阻塞主机线程的调用
cudaError_t cudaEventSynchronize(cudaEvent_t event);
//检查事件中所有操作是否都已经完成
cudaError_t cudaEventQuery(cudaEvent_t event);
//计算事件启动和停止之间的运行时间，以毫秒为单位
cudaError_t cudaEventElapsedTime(float* ms, cudaEvent_t start, cudaEvent_t stop);
//
//flags:
//    cudaEventDefault
//    cudaEventBlockingSync  #指定使用cudaEventSynchronize函数同步事件会阻塞调用的线程
//    cudaEventDisableTiming #创建的事件只能用来进行同步，不需要记录时序数据
//    cudaEventInterprocess	 #创建的事件可能被用作进程间事件
cudaError_t cudaEventCreateWithFlags(cudaEvent_t* event, unsigned int flags);

//cuda 同步
//强制主机端程序等待所有的核函数执行结束
cudaError_t cudaDeviceSynchronize(void);
//阻塞主机线程直到流中所有的操作完成
cudaError_t cudaStreamSynchronize(cudaStream_t stream);
cudaError_t cudaStreamQuery(cudaStream_t stream);
//阻塞主机线程直到等待一个事件的调用
cudaError_t cudaEventSynchronize(cudaEvent_t event);
cudaError_t cudaEventQuery(cudaEvent_t event);
//使指定的流等待指定的事件,该事件可能与同一个流相关，也可能与不同的流相关
cudaError_t cudaStreamWaitEvent(cudaStream_t stream, cudaEvent_t event);

本栏目推荐文章
• ubuntu20.04 opencv 4.2 + opencv_contrib 4.2.0 安装笔记 和 opencv 4.5.4 with cuda 安装笔记
• Windows10安装cuda失败
• Cannot re-initialize CUDA in forked subprocess.
• CUDA驱动深度学习发展 - 技术全解与实战
• slam相关-Sophus库cmake报错target "CUDA::cusolver" not found.
• 简易机器学习笔记（十）Windows下 PaddlePaddle配置CUDA加速环境
• AI_GPU-硬件-驱动-cuda
• cuda流的使用
• cuda\cudnn\torch安装
• cuda版本切换

Cudacuda9 cuda_home 案例cuda cuda_call cuda-kernels 模型cuda pytorch-cuda 笔记cuda cuda_visible_devices 深度cuda