CUDA与TensorRT学习二:CUDA硬件信息获取、Nsight system和Nsight compute、共享内存和bank conflict，预处理后处理、stream和event、双线性插

一、理解CUDA的grid和Block

• 目标
  理解Cuda中一维、二维、三维的grid、block的写法，以及遍历thread的方法

1）第一个cuda项目

• 修改项目的Makefile.config
  
• 总体文件目录
  
• 代码

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


__global__ void print_idx_kernel(){
    printf("block idx: (%3d, %3d, %3d), thread idx: (%3d, %3d, %3d)\n",
         blockIdx.z, blockIdx.y, blockIdx.x,
         threadIdx.z, threadIdx.y, threadIdx.x);
}

__global__ void print_dim_kernel(){
    printf("grid dimension: (%3d, %3d, %3d), block dimension: (%3d, %3d, %3d)\n",
         gridDim.z, gridDim.y, gridDim.x,
         blockDim.z, blockDim.y, blockDim.x);
}

__global__ void print_thread_idx_per_block_kernel(){
    int index = threadIdx.z * blockDim.x * blockDim.y + \
              threadIdx.y * blockDim.x + \
              threadIdx.x;

    printf("block idx: (%3d, %3d, %3d), thread idx: %3d\n",
         blockIdx.z, blockIdx.y, blockIdx.x,
         index);
}

__global__ void print_thread_idx_per_grid_kernel(){
    int bSize  = blockDim.z * blockDim.y * blockDim.x;

    int bIndex = blockIdx.z * gridDim.x * gridDim.y + \
               blockIdx.y * gridDim.x + \
               blockIdx.x;

    int tIndex = threadIdx.z * blockDim.x * blockDim.y + \
               threadIdx.y * blockDim.x + \
               threadIdx.x;

    int index  = bIndex * bSize + tIndex;

    printf("block idx: %3d, thread idx in block: %3d, thread idx: %3d\n", 
         bIndex, tIndex, index);
}

__global__ void print_cord_kernel(){
    int index = threadIdx.z * blockDim.x * blockDim.y + \
              threadIdx.y * blockDim.x + \
              threadIdx.x;

    int x  = blockIdx.x * blockDim.x + threadIdx.x;
    int y  = blockIdx.y * blockDim.y + threadIdx.y;

    printf("block idx: (%3d, %3d, %3d), thread idx: %3d, cord: (%3d, %3d)\n",
         blockIdx.z, blockIdx.y, blockIdx.x,
         index, x, y);
}

void print_one_dim(){
    int inputSize = 8;
    int blockDim = 4;
    int gridDim = inputSize / blockDim;

    dim3 block(blockDim);
    dim3 grid(gridDim);

    /* 这里建议大家吧每一函数都试一遍*/
    // print_idx_kernel<<<grid, block>>>();
    // print_dim_kernel<<<grid, block>>>();
    // print_thread_idx_per_block_kernel<<<grid, block>>>();
    print_thread_idx_per_grid_kernel<<<grid, block>>>();

    cudaDeviceSynchronize();
}

void print_two_dim(){
    int inputWidth = 4;

    int blockDim = 2;
    int gridDim = inputWidth / blockDim;

    dim3 block(blockDim, blockDim);
    dim3 grid(gridDim, gridDim);

    /* 这里建议大家吧每一函数都试一遍*/
    // print_idx_kernel<<<grid, block>>>();
    // print_dim_kernel<<<grid, block>>>();
    // print_thread_idx_per_block_kernel<<<grid, block>>>();
    print_thread_idx_per_grid_kernel<<<grid, block>>>();

    cudaDeviceSynchronize();
}

void print_cord(){
    int inputWidth = 4;

    int blockDim = 2;
    int gridDim = inputWidth / blockDim;

    dim3 block(blockDim, blockDim);
    dim3 grid(gridDim, gridDim);

    print_cord_kernel<<<grid, block>>>();

    cudaDeviceSynchronize();
}

int main() {
    /*
    synchronize是同步的意思，有几种synchronize

    cudaDeviceSynchronize: CPU与GPU端完成同步，CPU不执行之后的语句，知道这个语句以前的所有cuda操作结束
    cudaStreamSynchronize: 跟cudaDeviceSynchronize很像，但是这个是针对某一个stream的。只同步指定的stream中的cpu/gpu操作，其他的不管
    cudaThreadSynchronize: 现在已经不被推荐使用的方法
    __syncthreads:         线程块内同步
    */
    // print_one_dim();
    // print_two_dim();
    print_cord();
    return 0;
}

• 注意
  __global__ 表示核函数kernel

• 需求：找到某个block下面的thread
  

代码如下（先走z，然后y，最后z）

一般的优化

这里的找坐标表示的是：找到线程在哪个grid、对应grid下哪个block下的，block里面的位置坐标

二、理解.cu和.cpp的相互引用及Makefile

• 编译器
  不再是gcc或g++，而是nvcc，这样才不会编译报错
• 编译项目一指令

nvcc print_index.cu  -o app -I  /usr/local/cuda/include/

• cuda_check作用
  发生错误的时候告诉你错误发生在哪里

#define CUDA_CHECK(call) {                                                 \
    cudaError_t error = call;                                              \
    if (error != cudaSuccess) {                                            \
        printf("ERROR: %s:%d, ", __FILE__, __LINE__);                      \
        printf("CODE:%d, DETAIL:%s\n", error, cudaGetErrorString(error));  \
        exit(1);                                                           \
    }                                                                      \
}

三、利用CUDA矩阵乘法(matmul)计算、Error Handle 及硬件信息获取

1）矩阵乘法

• 目的
  理解使用cuda进行矩阵运算的加速方法，tile的用意
• 项目目录
  
• 生成打印的效果（对比CPU\GPU(warm up)\GPU（正常操作general下的计时））
  
• 项目目的
  利用cpu作为host端，gpu作为device端做矩阵计算（有涉及到数据传输）
  

额外注意点：
1）在cuda中，一个block可分配的thread数量最大是1024个线程，
若大于1024则会显示配置错误

函数执行说明：
启动核函数是异步启动GPU端，需要CPU同步等待GPU结果

参数说明：
grid和block一定配置，其他两个shared mem和stream不是必须一定配置

补充api说明：
1）cudaMalloc：在device端分配空间，是cuda runtime api
2）cudaMallocHost：在host端的pinned memory上分配空间
3）cudaMemcpy：以同步的方式，将数据在host->device，device->device ，device->host进行传输
4）cudaMemcpyAsync：以异步的方式，进行数据传输

• cuda各种api及libraries的关系（补充说明）
  
  

• Cuda Core的矩阵乘法计算是怎么做的？
  C = A * B
  ①案例一，只有一个block参与计算，这里需要8个clk来计算（涉及FMA，也就是需要乘法和加法混合计算结果，乘法和加法混合算一次）
  
  ②案例二，若计算全部block，也就是16个block，则需要8*16=128个clk才可以完成
  
  ③优化,这里每个block计算是没有关系的，可以每个block分配一个thread去计算，只需要8个clk
  

• 主体代码

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

#include "utils.hpp"
#include "timer.hpp"
#include "matmul.hpp"


int seed;
int main(){
    Timer timer;
    int width     = 1<<10; // 1,024
    int min       = 0;
    int max       = 1;
    int size      = width * width;
    int blockSize = 1;

    float* h_matM = (float*)malloc(size * sizeof(float));
    float* h_matN = (float*)malloc(size * sizeof(float));
    float* h_matP = (float*)malloc(size * sizeof(float));
    float* d_matP = (float*)malloc(size * sizeof(float));
    
    //1\生成不同矩阵，A和B
    seed = 1;
    initMatrix(h_matM, size, min, max, seed);
    seed += 1;
    initMatrix(h_matN, size, min, max, seed);
    
    //2、cpu开始计算/* CPU */
    timer.start();
    MatmulOnHost(h_matM, h_matN, h_matP, width);
    timer.stop();
    timer.duration<Timer::ms>("matmul in cpu");

    /* GPU warmup */
    timer.start();
    MatmulOnDevice(h_matM, h_matN, d_matP, width, blockSize);
    timer.stop();
    timer.duration<Timer::ms>("matmul in gpu(warmup)");

    /* GPU general implementation, bs = 16*/
    blockSize = 16;
    timer.start();
    MatmulOnDevice(h_matM, h_matN, d_matP, width, blockSize);
    timer.stop();
    timer.duration<Timer::ms>("matmul in gpu(bs = 16)");
    compareMat(h_matP, d_matP, size);

    /* GPU general implementation, bs = 1*/
    blockSize = 1;
    timer.start();
    MatmulOnDevice(h_matM, h_matN, d_matP, width, blockSize);
    timer.stop();
    timer.duration<Timer::ms>("matmul in gpu(bs = 1)");
    compareMat(h_matP, d_matP, size);

    /* GPU general implementation, bs = 32*/
    blockSize = 32;
    timer.start();
    MatmulOnDevice(h_matM, h_matN, d_matP, width, blockSize);
    timer.stop();
    timer.duration<Timer::ms>("matmul in gpu(bs = 32)");
    compareMat(h_matP, d_matP, size);
    return 0;
}

2）Error Handle

• 目标
  用cuda的error handler进行良好的编程

• 项目目录
  

• 修改的地方，添加了cudaCheck和kernel_check

#ifndef __UTILS_HPP__
#define __UTILS_HPP__

#include <cuda_runtime.h>
#include <system_error>

#define CUDA_CHECK(call)             __cudaCheck(call, __FILE__, __LINE__)
#define LAST_KERNEL_CHECK()          __kernelCheck(__FILE__, __LINE__)
#define BLOCKSIZE 16

inline static void __cudaCheck(cudaError_t err, const char* file, const int line) {
    if (err != cudaSuccess) {
        printf("ERROR: %s:%d, ", file, line);
        printf("CODE:%s, DETAIL:%s\n", cudaGetErrorName(err), cudaGetErrorString(err));
        exit(1);
    }
}

inline static void __kernelCheck(const char* file, const int line) {
    /* 
     * 在编写CUDA是，错误排查非常重要，默认的cuda runtime API中的函数都会返回cudaError_t类型的结果，
     * 但是在写kernel函数的时候，需要通过cudaPeekAtLastError或者cudaGetLastError来获取错误
     */
    cudaError_t err = cudaPeekAtLastError();
    if (err != cudaSuccess) {
        printf("ERROR: %s:%d, ", file, line);
        printf("CODE:%s, DETAIL:%s\n", cudaGetErrorName(err), cudaGetErrorString(err));
        exit(1);
    }
}

void initMatrix(float* data, int size, int low, int high, int seed);
void printMat(float* data, int size);
void compareMat(float* h_data, float* d_data, int size);

#endif //__UTILS_HPP__//

• 故意犯错打印
  

• 错误原因（dim3 dimBlock(a, b））
  在定义 dimBlock 时，需要确保 a 和 b 的乘积不超过 1024，以避免超出线程块的线程数限制。在你的代码中，如果 blockSize 大于 32，那么 dimBlock(blockSize, blockSize) 将创建一个线程数超过 1024 的线程块，这就是导致错误的原因

3）硬件信息获取

• 目标
  学习使用cuda runtime api显示GPU硬件信息，以及理解GPU硬件信息重要性
• 项目布局
  
• 打印效果
  
• 相关代码

int main(){
    int count;
    int index = 0;
    cudaGetDeviceCount(&count);
    while (index < count) {
        cudaSetDevice(index);
        cudaDeviceProp prop;
        cudaGetDeviceProperties(&prop, index);
        LOG("%-40s",             "*********************Architecture related**********************");
        LOG("%-40s%d%s",         "Device id: ",                   index, "");
        LOG("%-40s%s%s",         "Device name: ",                 prop.name, "");
        LOG("%-40s%.1f%s",       "Device compute capability: ",   prop.major + (float)prop.minor / 10, "");
        LOG("%-40s%.2f%s",       "GPU global meory size: ",       (float)prop.totalGlobalMem / (1<<30), "GB");
        LOG("%-40s%.2f%s",       "L2 cache size: ",               (float)prop.l2CacheSize / (1<<20), "MB");
        LOG("%-40s%.2f%s",       "Shared memory per block: ",     (float)prop.sharedMemPerBlock / (1<<10), "KB");
        LOG("%-40s%.2f%s",       "Shared memory per SM: ",        (float)prop.sharedMemPerMultiprocessor / (1<<10), "KB");
        LOG("%-40s%.2f%s",       "Device clock rate: ",           prop.clockRate*1E-6, "GHz");
        LOG("%-40s%.2f%s",       "Device memory clock rate: ",    prop.memoryClockRate*1E-6, "Ghz");
        LOG("%-40s%d%s",         "Number of SM: ",                prop.multiProcessorCount, "");
        LOG("%-40s%d%s",         "Warp size: ",                   prop.warpSize, "");

        LOG("%-40s",             "*********************Parameter related************************");
        LOG("%-40s%d%s",         "Max block numbers: ",           prop.maxBlocksPerMultiProcessor, "");
        LOG("%-40s%d%s",         "Max threads per block: ",       prop.maxThreadsPerBlock, "");
        LOG("%-40s%d:%d:%d%s",   "Max block dimension size:",     prop.maxThreadsDim[0], prop.maxThreadsDim[1], prop.maxThreadsDim[2], "");
        LOG("%-40s%d:%d:%d%s",   "Max grid dimension size: ",     prop.maxGridSize[0], prop.maxGridSize[1], prop.maxGridSize[2], "");
        index ++;
        printf("\n");
    }
    return 0;
}

• 知道参数的重要性
  很多时候编译.cu代码需要在nvcc之后加上编译信息，就需要打印GPU信息出来方便编译（比如共享内存的使用对cuda程序的加速很重要，可以动态修改共享内存和L1 Cache，而且知道作为调度thread的warp是由多少个thread组成的，也可以提高利用率）

• 错误分类
  kernelcheck主要捕捉核函数报错的检查，get会reset错误，错误不会传播下去；而peek不会reset错误，后面的正确使用的API也会产生同样错误

四、安装Nsight system and Nsight compute

• 目标
  学会如何使用Nvidia提供的Profiling tool进行性能分析

• Nsight system
  
  ①能看到核函数情况
  ②内存读写情况
  ③核函数之间调度情况
  ④SM Warp的占有率
  ⑤能判断CPU和GPU异步执行的情况

• Nsight compute（更偏向于核函数内部操作的情况）
  
  ①体现出核函数不同的blocksize和gridsize
  ②体现出不同核函数的执行时间，执行的吞吐量
  ③带宽分析
  ④访存走向

1）比较常用的分析

①Nsight system

1）对kernel核函数的timeline分析，下面是memory的timeline分析

1）可以看出PCIe带宽和DRAM带宽使用率（可以看出这里没使用共享内存，PCIe的带宽使用不高不低）
2）可以看五个核函数，第四个的资源利用率最高

②Nsight compute

1）后面能根据roofline能看出核函数是计算密集型还是访存密集型，可以设置base line比较分析（这里看到蓝色点比绿色点在更上方，可以看出来是更好的）

1）可以看到不同block下、不同共享内存下warp的占有率是否能达到更高的效率

1）可以看到共享内存没有使用，流量都是往global memory跑，那以后就可以往下也就是共享内存那边优化

2）看有没有关于bank的冲突

2）如何使用这些分析

方法三可以用nsys或ncu把统计数据打出来传到host端去用

3）两者更明显的不同

①NS

②NC

4）安装方法

①下载安装包

②创建远程ssh链接

③链接上去然后选择要监测的指标

把当前执行文件目录拷贝到working-directory，还有如下修改命令行，带参数的话就在执行程序后面加，这里程序执行文件就叫trt-cuda

④展开，选中区域右键zoom in放大可以看到百分之70多在执行核函数，百分之20多在访问内存

这里可以看到内存的访问比例，host到device和device到host，在不同的时间段里面，后面优化能放到一个重叠，能提高效率

这里看到warp调用线程的利率用不高，也是优化的地方

鼠标停顿能看到配置信息

--------------------------------------------------（后面是NC）
①创建ssh链接，填入执行文件名字trt-cuda和work directory当前目录

②replay Mode记得改为Application

③根据目标勾选监测指标

④先点开Summary，可以看到各个核函数的执行情况

下面是NC给出的各个建议
⑤可以点开特定核函数，查看detail信息

⑥可以设置某个核函数的roofline为标准线（baseline）去对比其他核函数

更高会更好
⑦可以查看吞吐量的流向，可以看到核函数瓶颈

⑧可以查看warp线程调度器的效率

五、共享内存、Bank Conflict原因和解决方法、TRT用Cuda进行预处理/后处理来加速、Stream 与Event（用Cuda写流提高并发性）

1）共享内存

• 目标
  理解如何使用shared memory，为什么使用shared memory会有加速效果，以及在shared memory中使用动态/静态变量的注意事项
• 项目布局（多了matmul_gpu_shared.cu作为共享内存使用对比）
  
• 使用时间对比（4096x4096大小的矩阵mamtul计算对比）
  ①热身
  ②普通内存
  ③静态固定分配的共享内存
  ④动态调整分配的共享内存cu代码
  
• 代码

#include "cuda_runtime_api.h"
#include "utils.hpp"

#define BLOCKSIZE 16

/* 
    使用shared memory把计算一个tile所需要的数据分块存储到访问速度快的memory中
*/
__global__ void MatmulSharedStaticKernel(float *M_device, float *N_device, float *P_device, int width){
    __shared__ float M_deviceShared[BLOCKSIZE][BLOCKSIZE];
    __shared__ float N_deviceShared[BLOCKSIZE][BLOCKSIZE];
    /* 
        对于x和y, 根据blockID, tile大小和threadID进行索引
    */
    int x = blockIdx.x * blockDim.x + threadIdx.x;
    int y = blockIdx.y * blockDim.y + threadIdx.y;

    float P_element = 0.0;

    int ty = threadIdx.y;
    int tx = threadIdx.x;
    /* 对于每一个P的元素，我们只需要循环遍历width / tile_width 次就okay了，这里有点绕，画图理解一下*/
    for (int m = 0; m < width / BLOCKSIZE; m ++) {
        M_deviceShared[ty][tx] = M_device[y * width + (m * BLOCKSIZE + tx)];
        N_deviceShared[ty][tx] = N_device[(m * BLOCKSIZE + ty)* width + x];
        __syncthreads();

        for (int k = 0; k < BLOCKSIZE; k ++) {
            P_element += M_deviceShared[ty][k] * N_deviceShared[k][tx];
        }
        __syncthreads();
    }

    P_device[y * width + x] = P_element;
}

__global__ void MatmulSharedDynamicKernel(float *M_device, float *N_device, float *P_device, int width, int blockSize){
    /* 
        声明动态共享变量的时候需要加extern，同时需要是一维的 
        注意这里有个坑, 不能够像这样定义： 
            __shared__ float M_deviceShared[];
            __shared__ float N_deviceShared[];
        因为在cuda中定义动态共享变量的话，无论定义多少个他们的地址都是一样的。
        所以如果想要像上面这样使用的话，需要用两个指针分别指向shared memory的不同位置才行
    */

    extern __shared__ float deviceShared[];
    int stride = blockSize * blockSize;
    /* 
        对于x和y, 根据blockID, tile大小和threadID进行索引
    */
    int x = blockIdx.x * blockSize + threadIdx.x;
    int y = blockIdx.y * blockSize + threadIdx.y;

    float P_element = 0.0;

    int ty = threadIdx.y;
    int tx = threadIdx.x;
    /* 对于每一个P的元素，我们只需要循环遍历width / tile_width 次就okay了 */
    for (int m = 0; m < width / blockSize; m ++) {
        deviceShared[ty * blockSize + tx] = M_device[y * width + (m * blockSize + tx)];
        deviceShared[stride + (ty * blockSize + tx)] = N_device[(m * blockSize + ty)* width + x];
        __syncthreads();

        for (int k = 0; k < blockSize; k ++) {
            P_element += deviceShared[ty * blockSize + k] * deviceShared[stride + (k * blockSize + tx)];
        }
        __syncthreads();
    }

    if (y < width && x < width) {
        P_device[y * width + x] = P_element;
    }
}

/*
    使用Tiling技术
    一个tile处理的就是block, 将一个矩阵分为多个小的tile，这些tile之间的执行独立，并且可以并行
*/
void MatmulSharedOnDevice(float *M_host, float *N_host, float* P_host, int width, int blockSize, bool staticMem){
    /* 设置矩阵大小 */
    int size = width * width * sizeof(float);
    long int sMemSize = blockSize * blockSize * sizeof(float) * 2;

    /* 分配M, N在GPU上的空间*/
    float *M_device;
    float *N_device;
    CUDA_CHECK(cudaMalloc((void**)&M_device, size));
    CUDA_CHECK(cudaMalloc((void**)&N_device, size));

    /* 分配M, N拷贝到GPU上*/
    CUDA_CHECK(cudaMemcpy(M_device, M_host, size, cudaMemcpyHostToDevice));
    CUDA_CHECK(cudaMemcpy(N_device, N_host, size, cudaMemcpyHostToDevice));

    /* 分配P在GPU上的空间*/
    float *P_device;
    CUDA_CHECK(cudaMalloc((void**)&P_device, size));;

    /* 调用kernel来进行matmul计算, 在这个例子中我们用的方案是：使用一个grid，一个grid里有width*width个线程 */
    dim3 dimBlock(blockSize, blockSize);
    dim3 dimGrid(width / blockSize, width / blockSize);
    if (staticMem) {
        MatmulSharedStaticKernel <<<dimGrid, dimBlock>>> (M_device, N_device, P_device, width);
    } else {
        MatmulSharedDynamicKernel <<<dimGrid, dimBlock, sMemSize, nullptr>>> (M_device, N_device, P_device, width, blockSize);
    }

    /* 将结果从device拷贝回host*/
    CUDA_CHECK(cudaMemcpy(P_host, P_device, size, cudaMemcpyDeviceToHost));
    CUDA_CHECK(cudaDeviceSynchronize());

    /* 注意要在synchronization结束之后排查kernel的错误 */
    LAST_KERNEL_CHECK(); 

    /* Free */
    cudaFree(P_device);
    cudaFree(N_device);
    cudaFree(M_device);
}

• 额外补充
  ①cuda中的event事件来统计时间比std的chrono更加准确
  ②cuda中的event是用cuda的stream来标记某一个执行点
  
• 对比之前，共享内存的优化点
  ①理由：一直访问global memory有点慢，而有些数据可以多次被复用，就可以放在共享内存里面
  
  ②shared memory和global memory对比
  L1和L2还有共享内存都是 on-chip memory，memory load/store的overhea会比较小，是高速访问的memory；而global memory的延迟是最高的，一般在cudaMalloc都是在global memory里面进行的

• 不同架构的带宽不同，越往右边越新（同一块block共享一块共享内存）
  
• 备注
  ①动态申请的时候需要是一维的
  ②动态申请的变量地址都是一样的，会比静态申请的速度会慢一点（因为动态只能一维）
• 代码讲解
  ①__syncthread和shared memory是绑定一起使用的，同步线程的操作
  

2）bank conflict

• 目标
  理解什么在shared memory中的bank是什么，什么时候发生bank conflict并且如何减缓
• 项目布局
  
• 对比执行速度
  
• bank定义
  把共享内存分为每一个thread占用的warp（一个warp32线程，32个bank）

• bank宽度
  
• bank带来的结果（A【0】和A【32】共享一个bank）

• padding方法环节bank conflict（方法：申请共享内存多添加一列，那么就算所有线程同时访问一列的2也不会冲突）
  

3）stream与event

• 目标
  什么是stream，cuda编程中的显示隐式同步，以及如何利用多流进行隐藏访存和核函数执行延迟的调度

• 项目目录
  

• 为了不想让核函数跑太快，所以核函数没做什么业务
  

• 这里也做了一个GELU的一个cuda的实现
  （以后可能会发现有些算子tensorRT不支持），只能用cuda自己写一个核函数对算子进行加速，让tensorRT去识别它

• 流stteam定义
  在GPU里面一串成顺序的指令，只要资源没被占用,不同流之间的执行是可以overlap的（同一个流里面的执行顺序和各个kerne以及memcpy operation的启动的顺序是一致的）
• 注意
  
  演变流程
  

①一开始是D2H拆分到hernel函数里面，
②然后就是kernel和D2H拆分到H2D里面
③加上k4的CPU操作，体现了GPU和CPU异步操作
④核函数太长，可以分割为K1,1\K1.2\K1.3多个步骤去处理

• 默认流
  

• 指定流（可以看到这里D2H和H2D包括核函数都指定流）
  

• 补充（Host端分配pin内存）
  ①pin内存分配函数方法
  
  ②内存分配区别
  1）Pageble memory：可分页内存（也就是物理内存）
  2）Pinned memory/paged-locked memory：页锁定内存
  3）由此引发的思考：当需要写回数据的时候，有可能数据不在有可能会被换到磁盘里面，所以会需要页锁定内存，如下图所示，省去从可分页内存到页锁定内存的拷贝操作
  

• 当资源被占满的情况和资源没有被占满的情况（多流）
  

多流和单流区别

可以看到SM Warp多流的利用率很高
当多核的时候，资源被占满了核函数不能overlapping，只能一个个等，效率就不咋高

• 多流的重点
  在核函数期间做内存的访问

• 如何隐藏延迟（不同颜色代表不同流）
  ①memory
  
  ②kernel（CPU启动核函数是异步的）
  

六、双线性插值(bilinear interpolatian)与仿射变换

1）双线性插值

• 目的：
  理解如何使用cuda进行opencv的图像处理的加速，理解双线性插值进行图像大小调整的算法流程
• 项目目录
  

preprocess.cu就是实现了很多个核函数进行双线性插值

①普通双线性插值核函数：forx_resized_blinear，完全转变为24x24图片
②forx_resized_blinear_letterbox：让规格不变做一个resize
③forx_resized_blinear_letterbox_cetner:像yolo里面letterbox，让分辨率不变，图片居中
显示效果比较，GPU执行时间：

• 一些代码展示
  bilinear resize的opencv实现
  
  bilinear resize的cuda实现（接口部分）
  
• 双线性插值的解释
  一种对图像进行缩放/放大的一种计算方法。opencv默认的resize方式

代码解释原理（与四个角的关系）

缩放比例

图片中心化操作（先往上移动再往中间移动）
格式转换（若需要把图片传给DNN，需要把opencv的默认格式BGR转换，也就是改变channle方向）

• 把很多操作弄成一个核函数
  
  ①能避免多次调佣kenrl
  ②实现内存复用，提高效率

• 涉及的tatics列表
  

2）仿射变换