前言

GPU介绍(从GPU诞生之日起，GPU设计逻辑与CPU设计逻辑差别很大。GPU自诞生之日起，其定位为3D图形渲染设备。在设计GPU从其功能出发，在数据处理中使用更多的晶体管。这使得GPU相比CPU单精度浮点操作能力更强。充分利用人GPU性能，用了很多方法。加速处理比较常见。

在而GPU在实现加速软件时，我们可能会听到 opengl opencl cuda这些名词。让我们分别介绍这些差异：

OpenGL（英语：Open Graphics Library，翻译：开放图形库或开放图形库D、3D矢量图形的跨语言、跨平台的应用程序编程接口（API）。该接口由近350个不同的函数调用组成，用于绘制从简单的图形比特到复杂的三维场景。另一种程序接口系统仅用于Microsoft Windows上的Direct3D。OpenGL常用于CAD、电子游戏开发虚拟现实、科学可视化程序。

OpenCl(是由苹果（Apple）由公司发起，业内众多知名厂商共同制定的开放免费标准，也是统一的编程环境。软件开发人员为高性能计算服务器、桌面计算系统和手持设备编写高效轻便的代码，广泛应用于多核心处理器(CPU)、图形处理器(GPU)、Cell类型架构和数字信号处理器(DSP)其他并行处理器在游戏、娱乐、科研、医疗等领域具有广阔的发展前景。)

自2007年以来，基于CUDA和OpenCL这些新语法特征与高级语言相似GPGPU语言减少了人们的使用GPGPU难度平缓了开始时的学习曲线。GPGPU领域，OpenGL中的GLSL逐渐退出人们的视线。

简单说OpenCL与OpenGL都是基于硬件的API的编程。OpenGL是针对图形的，而OpenCL并行计算API，下面还有一个并行计算cuda。

CUDA（Compute Unified Device Architecture，英伟达统一计算架构NVIDIA该公司推出的集成技术是该公司GPGPU正式名称。用户可以使用此技术NVIDIA的GeForce 8以后的GPU和较新的Quadro GPU进行计算。也是第一次使用GPU作为C-开发环境的编译器。NVIDIA在营销过程中，编译器和架构经常混合推广，造成混乱。实际上，CUDA可以相容OpenCL或者自家的C-编译器。无论是CUDA C-语言或是OpenCL，指令最终都会被驱动程式转换成PTX代码，交给显示核心计算。

通俗介绍，cuda是nvidia公司生态，前者为单一供应商提供完整的工具包(NVIDIA)成熟的开发平台，opencl是开源标准。

CUDA是NVIDIA GPU编程语言， OpenCL是异构计算库。CUDA和C 虽然都可以用nvcc编译，但C 只能在CPU上跑，CUDA只能在GPU上跑；而OpenCL不局限于计算设备，旨在通过其提供的抽象接口在各种硬件上运行相同的任务（CPU，GPU，FPGA，etc）

跨平台性和通用性OpenCL有很大的优势(这也是很多)National Laboratory使用OpenCL科学计算的主要原因)。OpenCL支持包括ATI,NVIDIA,Intel,ARM在内的多类处理器，并能支持运行在CPU并行代码也是独一无二的Task-Parallel Execution Mode，更好的支持Heterogeneous Computing。这只支持数据级并行，只能NVIDIA在众核处理器上运行CUDA做不到。

友好程度的开发者CUDA在这方面显然受更多开发者青睐。原因是套件的统一开发(CUDA Toolkit, NVIDIA GPU Computing SDK以及NSight等等)，非常丰富的库(cuFFT, cuBLAS, cuSPARSE, cuRAND, NPP, Thrust)以及NVCC(NVIDIA的CUDA有编译器)PTX(一种SSA中间表示不同NVIDIA GPU设备提供统一的静态ISA)更成熟的编译特性，如代码生成和离线编译。相比之下，使用OpenCL只有开发AMD对OpenCL驱动相对成熟。相比之下，使用OpenCL只有开发AMD对OpenCL驱动相对成熟。 来源

今天的主题来自介绍cuda使用，概念介绍的前一部分（来源是网络，只有一小部分是自己写的）只是为了帮助你理解今天要使用的cuda，然后开始进入正题。(自己带电脑NVIDIA手里有显卡，手里有显卡jetson nano的NVIDIA这就是为什么使用板卡cuda，请注意您手中的硬件是否有其他朋友使用NVIDIA的显卡）。

接下来，我将大致分为三个部分：一、cuda并行介绍计算原理 ，二、cuda三、opencv中cuda安装和使用

作者：良心依存

转载授权及旁观者：欢迎关注微信公众号：羽林君

或添加作者个人微信：become_me

cuda并行介绍计算原理

GPU和CPU不同的硬件特性决定了它们的应用场景，CPU它是计算机操作和控制的核心，GPU主要用于图形图像处理。计算机中图像呈现的形式是矩阵。我们对图像的处理实际上是操作各种矩阵进行计算，而许多矩阵的操作实际上可以并行化，这使得图像处理非常快，因此GPU也在图形图像领域了大展拳脚的机会。下图表示的就是一个多GPU计算机硬件系统，可以看出，一个GPU内存就有很多个SP和各类内存，这些硬件都是GPU进行高效并行计算的基础。

现在再从数据处理的角度来对比CPU和GPU的特点。CPU需要很强的通用性来处理各种不同的数据类型，比如整型、浮点数等，同时它又必须擅长处理逻辑判断所导致的大量分支跳转和中断处理，所以CPU其实就是一个能力很强的伙计，他能把很多事处理得妥妥当当，当然啦我们需要给他很多资源供他使用（各种硬件），这也导致了CPU不可能有太多核心（核心总数不超过16）。而GPU面对的则是类型高度统一的、相互无依赖的大规模数据和不需要被打断的纯净的计算环境，GPU有非常多核心（费米架构就有512核），虽然其核心的能力远没有CPU的核心强，但是胜在多， 在处理简单计算任务时呈现出“人多力量大”的优势，这就是并行计算的魅力。内容来源

整理一下两者特点就是：

• 1.CPU：擅长流程控制和逻辑处理，不规则数据结构，不可预测存储结构，单线程程序，分支密集型算法
• 2.GPU：擅长数据并行计算，规则数据结构，可预测存储模式

CUDA存储器类型：
每个线程拥有自己的 register寄存器 and loacal memory 局部内存
每个线程块拥有一块 shared memory 共享内存
所有线程都可以访问 global memory 全局内存

还有，可以被所有线程访问的
     只读存储器：
     constant memory (常量内容) and texture memory
    
a. 寄存器Register
   寄存器是GPU上的高速缓存器，其基本单元是寄存器文件，每个寄存器文件大小为32bit.
   Kernel中的局部(简单类型)变量第一选择是被分配到Register中。
       特点：每个线程私有，速度快。

b. 局部存储器 local memory

   当register耗尽时，数据将被存储到local memory。
   如果每个线程中使用了过多的寄存器，或声明了大型结构体或数组，
   或编译器无法确定数组大小，线程的私有数据就会被分配到local memory中。
   
       特点：每个线程私有；没有缓存，慢。
       注：在声明局部变量时，尽量使变量可以分配到register。如：
       unsigned int mt[3];
       改为：　unsigned int mt0, mt1, mt2;

c. 共享存储器 shared memory
       可以被同一block中的所有线程读写
       特点：block中的线程共有；访问共享存储器几乎与register一样快.

d. 全局存储器 global memory
　 特点：所有线程都可以访问；没有缓存

e. 常数存储器constant memory
   用于存储访问频繁的只读参数
   特点：只读；有缓存；空间小(64KB)
   注：定义常数存储器时，需要将其定义在所有函数之外，作用于整个文件

f. 纹理存储器 texture memory
       是一种只读存储器，其中的数据以一维、二维或者三维数组的形式存储在显存中。
   在通用计算中，其适合实现图像处理和查找，对大量数据的随机访问和非对齐访问也有良好的加速效果。
       特点：具有纹理缓存，只读。
threadIdx，blockIdx, blockDim, gridDim之间的区别与联系
 在启动kernel的时候，要通过指定gridsize和blocksize才行
 dim3 gridsize(2,2);   // 2行*2列*1页 形状的线程格，也就是说 4个线程块
    gridDim.x，gridDim.y，gridDim.z相当于这个dim3的x，y，z方向的维度，这里是2*2*1。
序号从0到3，且是从上到下的顺序，就是说是下面的情况:
 具体到 线程格 中每一个 线程块的 id索引为：
 
    grid 中的 blockidx 序号标注情况为：      0     2 
                                          1     3
				       
   dim3 blocksize(4,4);  // 线程块的形状，4行*4列*1页，一个线程块内部共有 16个线程
 
   blockDim.x，blockDim.y，blockDim.z相当于这个dim3的x，y，z方向的维度，
   这里是4*4*1.序号是0-15，也是从上到下的标注：

block中的 threadidx 序号标注情况   0       4       8      12 
				  1       5       9      13
				  2       6       10     14
				  3       7       11     15
1.  1维格子，1维线程块，N个线程======

 实际的线程id tid =  blockidx.x * blockDim.x + threadidx.x
 块id   0 1 2 3 
 线程id 0 1 2 3 4
2. 1维格子，2D维线程块
 块id   1 2 3 
 线程id  0  2
        1  3
		       块id           块线程总数
 实际的线程id tid =  blockidx.x * blockDim.x * blockDim.y + 
		       当前线程行数    每行线程数
		    threadidx.y * blockDim.x   + 
		       当前线程列数
		    threadidx.x

更加详细信息，大家可以参考这两篇文章：

https://www.cnblogs.com/skyfsm/p/9673960.html

https://github.com/Ewenwan/ShiYanLou/tree/master/CUDA

安装cuda使用环境

直接使用cuda

wget https://developer.download.nvidia.com/compute/cuda/repos/ubuntu2004/x86_64/cuda-ubuntu2004.pin
sudo mv cuda-ubuntu2004.pin /etc/apt/preferences.d/cuda-repository-pin-600
wget https://developer.download.nvidia.com/compute/cuda/11.6.0/local_installers/cuda-repo-ubuntu2004-11-6-local_11.6.0-510.39.01-1_amd64.deb
sudo dpkg -i cuda-repo-ubuntu2004-11-6-local_11.6.0-510.39.01-1_amd64.deb
sudo apt-key add /var/cuda-repo-ubuntu2004-11-6-local/7fa2af80.pub
sudo apt-get update
sudo apt-get -y install cuda

#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <stdio.h>
#include <iostream>
using namespace std;

__global__ void kernel(void) { //带有了__global__这个标签，表示这个函数是在GPU上运行
  printf("hello world gpu \n");
}
int main() {
  kernel<<<1, 1>>>();//调用除了常规的参数之外，还增加了<<<>>>修饰,其中第一个1，代表线程格里只有一个线程块；第二个1，代表一个线程块里只有一个线程。
    cudaError_t cudaStatus;

    cudaStatus = cudaDeviceSynchronize();
    if (cudaStatus != cudaSuccess) {
        fprintf(stderr, "cudaDeviceSynchronize returned error code %d after launching addKernel!\n", cudaStatus);
    }

  return 0;
}

这里面记得使用cudaDeviceSynchronize函数，等待GPU完成运算。如果不加的话，你将看不到printf的输出

编译选项 nvcc test_cuda_hello.cu

这个程序中GPU调用的函数和普通的C程序函数的区别

• 函数的调用除了常规的参数之外，还增加了<<<>>>修饰。
• 调用通过<<<参数1,参数2>>>，用于说明内核函数中的线程数量，以及线程是如何组织的。
• 以线程格（Grid）的形式组织，每个 线程格 由若干个 线程块（block）组成，
• 而每个线程块又由若干个线程（thread）组成。
• 是以block为单位执行的。

接下来，我们再来使用一个demo，使用一个简单的计算，将其放置于GPU进行计算

// 输入变量 以指针方式传递===================================
#include <iostream>
#include <cuda.h>
#include <cuda_runtime.h>
#include <stdio.h>

// 输入变量 全部为指针  类似数据的引用
__global__ void gpuAdd(int *d_a, int *d_b, int *d_c) 
{
	*d_c = *d_a + *d_b;
}

int main(void) 
{
	// CPU变量
	int h_a,h_b, h_c;
	// CPU  指针 变量 指向 GPU数据地址
	int *d_a,*d_b,*d_c;
	// 初始化CPU变量
	h_a = 1;
	h_b = 4;
  
	// 分配GPU 变量内存
	cudaMalloc((void**)&d_a, sizeof(int));
	cudaMalloc((void**)&d_b, sizeof(int));
	cudaMalloc((void**)&d_c, sizeof(int));
  
	// 输入变量 CPU 拷贝到 GPU   右 到 左
	cudaMemcpy(d_a, &h_a, sizeof(int), cudaMemcpyHostToDevice);
	cudaMemcpy(d_b, &h_b, sizeof(int), cudaMemcpyHostToDevice);
  
	// 调用核函数
	gpuAdd << <1, 1 >> > (d_a, d_b, d_c);
  
	// 拷贝GPU数据结果 d_c 到 CPU变量
	cudaMemcpy(&h_c, d_c, sizeof(int), cudaMemcpyDeviceToHost);
	printf("Passing Parameter by Reference Output: %d + %d = %d\n", h_a, h_b, h_c);
  
	// 清理GPU内存 Free up memory 
	cudaFree(d_a);
	cudaFree(d_b);
	cudaFree(d_c);
	return 0;
}

CUDA代码中比较重要的函数:

cudaMalloc 与C语言中的malloc函数一样，只是此函数在GPU的内存你分配内存。

addKernel<<<1, size>>>

这里就涉及了GPU和主机之间的内存交换了，cudaMalloc是在GPU的内存里开辟一片空间， 然后通过操作之后，这个内存里有了计算出来内容，再通过cudaMemcpy这个函数把内容从GPU复制出来。可以在设备代码中使用cudaMalloc()分配的指针进行设备内存读写操作；cudaMalloc()分配的指针传递给在设备上执行的函数；但是不可以在主机代码中使用cudaMalloc()分配的指针进行主机内存读写操作（即不能进行解引用）。 cudaFree 与c语言中的free()函数一样，只是此函数释放的是cudaMalloc()分配的内存。

cudaMemcpy 与c语言中的memcpy函数一样，只是此函数可以在主机内存和GPU内存之间互相拷贝数据。此外与C中的memcpy()一样，以同步方式执行，即当函数返回时，复制操作就已经完成了，并且在输出缓冲区中包含了复制进去的内容。 相应的有个异步方式执行的函数cudaMemcpyAsync().

该函数第一个参数是目的指针，第二个参数是源指针，第三个参数是复制内存的大小，第四个参数告诉运行时源指针，这是一个什么类型的指针，即把内存从哪里复制到哪里。第四个参数可以选用以下的形式：参考

• cudaMemcpyHostToDevice：从主机复制到设备；
• cudaMemcpyDeviceToHost：从设备复制到主机；
• cudaMemcpyDeviceToDevice：从设备复制到设备；
• cudaMemcpyHostToHost：从主机复制到主机；

调用的核函数前面已经介绍过，此处不再赘述。还有其他函数大家可以自行参考官网的API介绍。

大家也可以看《GPU高性能编程CUDA实战》这本书

opencv中cuda库安装与使用

OpenCV Contrib库是非官方的第三方开发扩充库。通过这个库，我们能使用如dnn、相机标注、3D成像、ArUco、物体追踪，甚至是需付费的SURF、SIFT特征点提取算法

安装OpenCV Contrib也是需要和opencv进行版本匹配的。大家可以在下面位置进行对应版本下载：

• contrib库：https://github.com/opencv/opencv_contrib/tags
• opencv版本：https://opencv.org/releases.html

对与我个人而言，opencv版本是 4.4.5

$ pkg-config opencv4 --modversion 
4.4.5

下载好对应的opencv和opencv_contrib库，开始进行编译，和之前单独opencv编译类似，解压下载好的文件，然后在opencv建立一个build目录，在里面进行cmake 编译。

在这里给大家在注解一些相关选项的含义：

 cmake -D CMAKE_BUILD_TYPE=Release\
 			-D ENABLE_CXX11=ON\
 	        -D CMAKE_INSTALL_PREFIX=/usr/local\
 	        -D WITH_CUDA=ON\
 	         -D CUDA_ARCH_BIN=${cuda_compute}\
 	         -D CUDA_ARCH_PTX=""\
 	         -D ENABLE_FAST_MATH=ON\
 	         -D CUDA_FAST_MATH=ON\
 	         -D WITH_CUBLAS=ON\
 	          -D WITH_LIBV4L=ON\
 	          -D WITH_GSTREAMER=ON\
 	           -D WITH_GSTREAMER_0_10=OFF\
 	           -D WITH_QT=ON\
 	           -D WITH_OPENGL=ON\
 	           -D CUDA_NVCC_FLAGS="--expt-relaxed-constexpr" \
 	           -D CUDA_TOOLKIT_ROOT_DIR=/usr/local/cuda-11.6\
 	           -D WITH_TBB=ON\
 	           -D OPENCV_EXTRA_MODULES_PATH=/home/lyn/Documents/application/opencv_contrib-4.5.5/modules          ../

• CUDA_ARCH_BIN=${cuda_compute} 显卡算力，大家可以对照我最后给算力查询自行查询，也可以进行系统自己查找匹配版本。

• WITH_QT=ON\ 这是是QT的编译选项 大家可以按照自己需要选择是否用QT进行编译开与关。

• CUDA_TOOLKIT_ROOT_DIR 这是自己电脑安装的cuda 版本

• OPENCV_EXTRA_MODULES_PATH=/home/lyn/Documents/application/opencv_contrib-4.5.5/modules 这是解压好后的opencv_contrib目录

安装时候可能会遇到如下问题：

In file included from /home/lyn/Documents/application/opencv-4.5.5/build/modules/python_bindings_generator/pyopencv_custom_headers.h:7,
                 from /home/lyn/Documents/application/opencv-4.5.5/modules/python/src2/cv2.cpp:88:
/home/lyn/Documents/application/opencv_contrib-4.5.5/modules/phase_unwrapping/misc/python/pyopencv_phase_unwrapping.hpp:2:13: error: ‘phase_unwrapping’ in namespace ‘cv’ does not name a type
    2 | typedef cv::phase_unwrapping::HistogramPhaseUnwrapping::Params HistogramPhaseUnwrapping_Params;

介绍一个简单例子：

# 声明要求的 cmake 最低版本
cmake_minimum_required( VERSION 2.8 )
# 声明一个 cmake 工程
project(opencv_cuda)
# 设置编译模式
set( CMAKE_BUILD_TYPE "Debug" )
set( CMAKE_CXX_FLAGS "-std=c++11")
#添加OPENCV库
find_package(OpenCV REQUIRED)
#添加OpenCV头文件
include_directories(${OpenCV_INCLUDE_DIRS})
#显示OpenCV_INCLUDE_DIRS的值
message(${OpenCV_INCLUDE_DIRS})
add_executable(target test_cuda_cv.cpp)
# 将库文件链接到可执行程序上
target_link_libraries(target  ${OpenCV_LIBS} ) 

#include <iostream>
#include "opencv2/opencv.hpp"
#include <opencv2/core/cuda.hpp> 
#include <opencv2/cudaarithm.hpp>
#include <opencv2/core/version.hpp>

int main (int argc, char* argv[])
{
    //Read Two Images 
    cv::Mat h_img1 = cv::imread( "/home/lyn/Documents/work-data/test_code/opencv/learn_code/"
      "opencv_tutorial_data-master/images/sp_noise.png");
    cv::Mat h_img2 = cv::imread( "/home/lyn/Documents/work-data/test_code/opencv/learn_code/"
      "opencv_tutorial_data-master/images/sp_noise.png");      
    // cv::Mat h_img2 = cv::imread("/home/lyn/Documents/work-data/test_code/opencv/learn_code/c++/build/filter_line.png");
  
  // cv::namedWindow( "1", cv::WINDOW_FREERATIO); 
  // cv::imshow("1",h_img1);  

  // cv::namedWindow( "2", cv::WINDOW_FREERATIO); 
  // cv::imshow("2",h_img2);  

    cv::Mat h_result1;


    #if 0
    // cv::add(h_img1, h_img2,h_result1);//加法操作 api
    cv::multiply(h_img1,cv::Scalar(2, 2, 2),h_result1);//乘法操作 api

    #else
    // 定义GPU数据
    cv::cuda::GpuMat d_result1,d_img1, d_img2;
    
    // CPU 到 GPU   
    d_img1.upload(h_img1);
    d_img2.upload(h_img2);
    
    // 调用GPU接执行 mat substract
    // cv::cuda::add(d_img1, d_img2,d_result1);
    cv::cuda::multiply(d_img1,cv::Scalar(2, 2, 2),d_result1);//乘法操作 api

    // gpu 结果 到 cpu
    d_result1.download(h_result1);

    #endif 
 

    // 显示，保存
  // cv::namedWindow( "Image1", cv::WINDOW_FREERATIO); 
    cv::imshow("Image1 ", h_img1);

  // cv::namedWindow( "Image2", cv::WINDOW_FREERATIO); 
    cv::imshow("Image2 ", h_img2);

  // cv::namedWindow( "Result_Subtraction", cv::WINDOW_FREERATIO); 
    cv::imshow("Result_Subtraction ", h_result1);
    cv::imwrite("/home/lyn/Documents/work-data/test_code/opencv/learn_code/c++/result_add.png", h_result1);
    cv::waitKey();
    return 0;
}

附录：这是官网提供算力表，大家可以参考选择自己的要求算力的显卡。

结语

                              ‧‧‧‧‧‧‧‧‧‧‧‧‧‧‧‧  END  ‧‧‧‧‧‧‧‧‧‧‧‧‧‧‧‧