c++高性能：多进程 cuda编程

        GPU(Device)用于异构计算，CPU(Host)控制GPU。极简入门。API(CUDA Driver API+CUDA Runtime API+CUDA Libraries)

hello world

// main.cu
#include <iostream>
#include <stdio.h>

void cpu_hello(){
    printf("cpu\n");
}
__global__ void cuda_hello(){
    printf("gpu\n");
}

int main() {
    cpu_hello();
    cuda_hello<<<1,10>>>();
    cudaDeviceSynchronize();//witout this no output

    std::cout << "Hello, World!" << std::endl;
    return 0;
}

编译

• nvcc -arch=sm_86 -o CUDATEST main.cu -run
• 30**系类显卡用sm_86编译，如果不通过可以减小

# https://docs.nvidia.com/cuda/cuda-compiler-driver-nvcc/index.html#gpu-compilation
nvcc x.cu
 -gencode arch=compute_50,code=sm_50
 -gencode arch=compute_60,code=sm_60
 -gencode arch=compute_70,code=\"compute_70,sm_70\"

嵌入与计算能力5.0和6.0兼容的二进制代码以及与计算能力7.0兼容的PTX和二进制代码。

全局函数，设备，主机函数标识 __global__，__device__，__host__

https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#cuda-enabled-gpus		–
__global__ void fun(){}	在设备上执行，可从主机（计算能力> 5.0设备 ）调用	必须无返回值
__device__	在设备上执行，只能从设备调用
__host__	在主机上执行，只能从主机调用。	可省略

块和线程

cuda的<<< >>>语法糖是用于并行计算的一种语法，它是cuda中的特有语法，用于表示在GPU中启动的线程块和线程的数量。其中<<<表示启动的线程块的数量，表示每个线程块中的线程数>>>。

例如，如果我们想启动一个包含64个线程的线程块，我们可以使用以下语法：

kernel<<<1, 64>>>(args);

这表示我们只启动了一个线程块，每个线程块中有64个线程。如果我们想启动多个线程块，我们可以使用以下语法：

kernel<<<N, 64>>>(args);

其中N是启动的线程块的数量，每个线程块中有64个线程。通过这种语法糖，我们可以轻松地启动大量的线程块和线程，实现高效的并行计算。

2块，3线程(执行2*3次 global)	add<<<2,3>>>(a,b,c)	和mpi类似，使用块号或者线程号进行细节操作
girdDim.x	块的个数
blockDim.x	一个块的线程个数
blockIdx.x	块的索引
threadIdx.x	一个块中线程的索引
在所有线程中的index	blockIdx.x*blockDim.x+threadIdx.x	blockIdx.x*3+threadIdx.x

dim3

dim3是一个CUDA中的结构体，用于表示三维中的向量。它包含三个unsigned int类型的成员变量x、y和z，分别表示向量在三个维度上的分量。

在CUDA编程中，dim3通常用于表示线程块和线程的数量。我们可以使用dim3类型来指定启动的线程块和线程的数量，例如：

dim3 block(16, 16, 1);
dim3 grid(32, 32, 1);
kernel<<<grid, block>>>(args);

这里我们创建了一个包含16 x 16个线程的线程块，以及32 x 32个线程块的网格。在启动kernel函数时，我们将grid作为第一个参数传递给<<<>>>语法糖，表示启动32 x 32个线程块，而block作为第二个参数传递给kernel函数，表示每个线程块包含16 x 16个线程。

在实际编程中，dim3还可以用于表示其他三维向量，例如图像的分辨率、数据的维度等。

	dim3 dimBlock(x,y);Kernel<<<2,dimBlock>>>(argv)	blockIdx.x * blockDim.x * blockDim.y + threadIdx.y * blockDim.x + threadIdx.x;
	dim3 dimGrid(x,y,z)；Kernel<<<dimGrid,2>>>(argv)；

全为3D的索引（第几块的第几个线程）：

• int blockId = blockIdx.x+ blockIdx.y * gridDim.x+ blockIdx.z * gridDim.x * gridDim.y; （第几块（对应最初始公式中的blockIdx.x）：blockIdx.*为系数， gridDim.*为次数）

• int Idx = blockId * (blockDim.x * blockDim.y * blockDim.z) + (threadIdx.z * (blockDim.x * blockDim.y)) + (threadIdx.y * blockDim.x)+ threadIdx.x; （后边部分对应于最初始公式的threadIdx.x）

• 矩阵加法 https://b23.tv/AWa1Pml

内存限制

硬件约束：
这是易于量化的部分。 当前 CUDA 编程指南的附录 F 列出了许多硬限制，这些限制限制了内核启动的每个块可以有多少线程。如果超过其中任何一个，您的内核将永远不会运行。它们可以大致概括为：

1. 每个块总共不能超过 512/1024 个线程（计算能力分别为 1.x 或 2.x 及更高版本）
2. 每个块的最大尺寸限制为 [1024,1024,64]/[2^31-1,65535,65535]（计算 2.x/2.x+）
3. 每个块的总容量不能超过 8k/16k/32k/64k/32k/64k/32k/64k/32k/64k （计算 1.0,1.1/1.2,1.3/2.x-/3.0/3.2/3.5-5.2/5.3/6-6.1/6.2/7.0）
4. 每个块不能消耗超过 16kb/48kb/96kb 的共享内存（计算 1.x/2.x-6.2/7.0）

如果您保持在这些限制范围内，则可以成功编译的任何内核都将启动而不会出错。

Grid-Stride循环

• CUDA编程-Grid-Stride循环实现的灵活Kernel函数
• 之前总是假设单个grid里的线程可以一次性处理整个数组，但实际上难以实现。
• 循环的步长是blockDim.x * gridDim.x，这是Grid中线程的数量，

__global__
void saxpy(int n, float a, float *x, float *y)
{
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) 
        y[i] = a * x[i] + y[i];
}

__global__
void saxpy(int n, float a, float *x, float *y)
{
    for (int i = blockIdx.x * blockDim.x + threadIdx.x; 
         i < n; 
         i += blockDim.x * gridDim.x) 
      {
          y[i] = a * x[i] + y[i];
      }
}

错误处理

• How to do error checking in CUDA(如何在CUDA里做错误检查)

• cudaError err = cudaGetLastError();可以处理没有返回值的global报错

内存

统一内存

• cudaMallocManaged(); 统一内存，cpu gpu都能访问,使用时会自动切页，但是程序降速
• int* x;cudaMallocManaged(&x, sizeof(int) * 2);

设备内存

c	cuda c	备注
malloc	cudaMalloc	申请)显存
memcpy	cudaMemcpy	同步执行函数，且具有方向参数
free	cudaFree	释放现存
	cudaMallocHost

• __host__ ​cudaError_t cudaMemcpy ( void* dst, const void* src, size_t count, cudaMemcpyKind kind )
• cudaMemcpy用于在主机（Host）和设备（Device）之间往返的传递数据，用法如下：
  主机到设备：cudaMemcpy(d_A,h_A,nBytes,cudaMemcpyHostToDevice)
  设备到主机：cudaMemcpy(h_A,d_A,nBytes,cudaMemcpyDeviceToHost)
  注意：该函数是同步执行函数，在未完成数据的转移操作之前会锁死并一直占有CPU进程的控制权，所以不用再添加cudaDeviceSynchronize()函数

example:

// https://www.olcf.ornl.gov/tutorials/cuda-vector-addition/
#include <stdio.h>
#include <cuda_runtime.h>

#define N 10

// CUDA kernel. Each thread takes care of one element of c
__global__ void vecAdd(double *a, double *b, double *c, int n)
{
    // Get our global thread ID
    int id = blockIdx.x*blockDim.x+threadIdx.x;
 
    // Make sure we do not go out of bounds
    if (id < n)
        c[id] = a[id] + b[id];
}

int main()
{
    int h_a[N], h_b[N], h_c[N];
    int *d_a, *d_b, *d_c;

    // Allocate memory on the device
    cudaMalloc(&d_a, N * sizeof(int));
    cudaMalloc(&d_b, N * sizeof(int));
    cudaMalloc(&d_c, N * sizeof(int));

    // Initialize host arrays
    for (int i = 0; i < N; i++) {
        h_a[i] = i;
        h_b[i] = i * i;
    }

    // Copy data from host to device
    cudaMemcpy(d_a, h_a, N * sizeof(int), cudaMemcpyHostToDevice);
    cudaMemcpy(d_b, h_b, N * sizeof(int), cudaMemcpyHostToDevice);

    // Perform vector addition on the device
    int blockSize = 256;
    int numBlocks = (N + blockSize - 1) / blockSize;
    addVectors<<<numBlocks, blockSize>>>(d_a, d_b, d_c, N);

    // Copy data from device to host
    cudaMemcpy(h_c, d_c, N * sizeof(int), cudaMemcpyDeviceToHost);

    // Print results
    for (int i = 0; i < N; i++) {
        printf("%d + %d = %d\n", h_a[i], h_b[i], h_c[i]);
    }

    // Free memory on the device
    cudaFree(d_a);
    cudaFree(d_b);
    cudaFree(d_c);

    return 0;
}

• 在本例中，我们将两个数组“h_a”和“h_b”从主机复制到设备，在设备上执行向量加法，然后将结果从设备复制回主机。

cudaMemPrefetchAsync(prefeatch预取)

• UM 分配存在切页过程，cudaMallocManaged 一般搭配使用cudaMemPrefetchAsync

• 后边的参数设置是放到cpu还是gpu

• 写在传入global函数之前int deviceId; cudaGetDevice(&deviceId); cudaMemPrefetchAsync(pointerToSomeUMData, size, deviceId); // Prefetch to GPU device.

• cudaMemPrefetchAsync(pointerToSomeUMData, size, cudaCpuDeviceId); // Prefetch to host.

shared memory

• __shared__ float sh_arr[128]
  
• https://www.jianshu.com/p/8d17817a7488

constant

• __constant__ 只读，全局

SM（stream multiprocessor）: 流处理器

• GPU:每个GPU有若干个SM，每个SM并行而独立运行

• 定义流：cudaStream_t s1;

• 创建流：cudaStreamCreate(&s1);

• 使用流：func_kernel<<< blocks,threads,0,s1 >>>

• 销毁流：cudaStreamDestory(s1);

• CUDA随笔之Stream的使用

CUDA同步操作

原子操作函数

• 和多线程的原子含义相同，自动加去锁保持线程的独占操作

__global__ void increment_atomic(int *g)
{
	int i = blockIdx.x * blockDim.x + threadIdx.x; 
	i = i % ARRAY_SIZE;  
	atomicAdd(& g[i], 1);//原来为g[i] = g[i] + 1;  
}

_syncthreads（）

• 线程块内线程同步，保证线程块内所有线程都执行到统一位置

_threadfence()
一个线程调用__threadfence后，该线程在该语句前对全局存储器或共享存储器的访问已经全部完成，执行结果对grid中的所有线程可见。

_threadfence_block()
一个线程调用__threadfence_block后，该线程在该语句前对全局存储器或者共享存储器的访问已经全部完成，执行结果对block中的所有线程可见。

以上两个函数的重要作用是，及时通知其他线程，全局内存或者共享内存内的结果已经读入或写入完成了。

cudaStreamSynchronize()/cudaEventSynchronize()
主机端代码中使用cudaThreadSynchronize():实现CPU和GPU线程同步
kernel启动后控制权将异步返回，利用该函数可以确定所有设备端线程均已运行结束

跟踪分析器

NVIDIA Nsight GPU 跟踪分析器

https://developer.nvidia.com/nvidia-visual-profiler

归约算法 规约∑

• CUDA专家手册：GPU编程权威指南
  

矩阵转置

#include <stdio.h>

#define N 16

__global__ void transpose(int *input, int *output) {
    __shared__ int tile[N][N+1];

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

    tile[threadIdx.y][threadIdx.x] = input[y * N + x];

    __syncthreads();

    x = blockIdx.y * N + threadIdx.x;
    y = blockIdx.x * N + threadIdx.y;

    output[y * N + x] = tile[threadIdx.x][threadIdx.y];
}

int main() {
    int *h_input, *h_output;
    int *d_input, *d_output;

    // 分配和初始化主机内存
    h_input = (int*)malloc(N * N * sizeof(int));
    h_output = (int*)malloc(N * N * sizeof(int));
    for (int i = 0; i < N * N; i++) {
        h_input[i] = i;
    }

    // 分配设备内存
    cudaMalloc((void**)&d_input, N * N * sizeof(int));
    cudaMalloc((void**)&d_output, N * N * sizeof(int));

    // 将输入数据从主机内存复制到设备内存
    cudaMemcpy(d_input, h_input, N * N * sizeof(int), cudaMemcpyHostToDevice);

    // 定义网格和块的大小
    dim3 grid(N/N, N/N);
    dim3 block(N, N);

    // 调用内核函数
    transpose<<<grid, block>>>(d_input, d_output);

    // 将结果从设备内存复制到主机内存
    cudaMemcpy(h_output, d_output, N * N * sizeof(int), cudaMemcpyDeviceToHost);

    // 打印结果
    for (int i = 0; i < N; i++) {
        for (int j = 0; j < N; j++) {
            printf("%d ", h_output[i * N + j]);
        }
        printf("\n");
    }

    // 释放内存
    free(h_input);
    free(h_output);
    cudaFree(d_input);
    cudaFree(d_output);

    return 0;
}

在上述代码中，首先定义了一个transpose的CUDA内核函数。该函数使用了共享内存tile来存储输入矩阵的一个小块。然后，根据线程的索引计算输入矩阵的位置，并将数据存储到共享内存中。接下来，通过交换坐标x和y，将共享内存中的数据写入到输出矩阵中。最后，使用cudaMemcpy函数将结果从设备内存复制到主机内存，并打印输出结果。

在主函数中，首先分配并初始化主机内存，然后分配设备内存，并将输入数据从主机内存复制到设备内存。接着，定义了网格和块的大小，并调用内核函数。最后，将结果从设备内存复制到主机内存，并打印输出结果。

需要注意的是，上述代码中的矩阵大小为16x16，可以根据实际需求进行修改。

CG

• 异构计算的核心点在于“异构”二字，说白了就是用不同制程架构、不同指令集、不同功能的硬件组合起来提高算力水平。

tips

• 因为GPU流处理器原因,线程个数为32倍数最好

• 以__device__ shared__为关键词声明的变量会被分配至SM上的shared memory， 可以由block内的全部线程所共享，生命周期也随着block的结束而结束。（下图位置），为了编程方便，可以直接写为__shared

• 获取GPU温度：https://github.com/jordanbonilla/Read_GPU_Temperature_CUDA

• https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html

• 获取显卡算力

• Hillis Steele Scan

programming class

• https://github.com/teknoman117/cuda/blob/master/stuff/vectoradd.cu
• torch cuda
• NMS

image processing

• CUDA编程《二》用GPU写一个卷积 https://www.bilibili.com/video/BV1c4411z7o2/?
• BGR转灰度图的程序，希望能让大家稍微明白一点CUDA程序如何写
• Image Filtering using CUDA
• https://github.com/teknoman117/cuda/tree/master/imgproc_example
• https://github.com/LitLeo/OpenCUDA
• The CMake version of cuda_by_example
• https://github.com/paramhanji/CUDA-CNN
• https://github.com/cheesinglee/cuda-PHDSLAM

teaching class

• https://github.com/PacktPublishing/Learn-CUDA-Programming

• https://github.com/csc-training/CUDA

• 玩游戏学cuda,可这个为啥是python代码呀[飙泪笑],写的就是python呀，这个项目用的是numba

• 车道线识别之——增强黄色车道线 https://blog.csdn.net/YaoJiawei329/article/details/111032256?

• cuda编程资源
  

• CUDA编程-《Professional CUDA C Programming》第1章-读书笔记（文字+图解）https://zhuanlan.zhihu.com/p/628911759