GPU架构理解与代码编写

参考文档: https://blog.csdn.net/wangbowj123/article/details/105514084

首先是对 CUDA 编程中可能会遇到的各种概念进行简单总结。如对 Kernel、Grid、Device、Host、Thread、Thread、Block、SM 等部件进行梳理，并且牵涉到CUDA编程的基本理念与基本方法。

CUDA 编程之软硬件结构的相关概念

首先是对 CUDA 编程中可能会遇到的各种概念进行简单总结。

下面这个图是 CUDA 编程中常见到的软硬件的结构图。

Kernel核: 可以理解为C/C++中的一个函数function。不过这样的理解其实不够准确，个人认为更准确的理解方式是将其视为一次对在 device 上运行函数的调用，每次调用 kernel 核，都需要指定一些参数，参数的形式也有很多种。下图中绿色的框便是一个 kernel 核（只有一个线程块）。一个 kernel 核可以具备很多个线程块。

所以说，Kernel 对应于需要在GPU上执行的程序，并且一个Kernel对应一个 Grid 。

SM（stream multiprocessor）: 流处理器

GPU :每个GPU有若干个SM，最少有1个，每个SM并行而独立运行

从上面的图可以看出：CUDA最大的特点：对线程块将在何处、何时运行不作保证。

优点：

硬件真正有效的运行，灵活
无需要线程间互相等待
可扩展性强

后果：

对于那个块在那个SM上运行无法进行任何假设
无法获取块之间的明确通讯（hard to get communications between blocks）
dead lock（并行死锁）
线程退出

不过也有一些确定的概念：

所有在同一个线程块上的线程必然会在同一时间运行在同一个SM上
同一个内核的所有线程块必然会全部完成了后，才会运行下一个内核如下图所示：

下图是 GPU 工作的 内存模型 。GPU 硬件架构具有三层的存储结构，分别是： 线程私有存储、线程块共享存储、全局存储 。访问速度自然是由快到慢。CUDA 程序的编写的一个重要理念也是尽可能少的访问全局缓存。

同步性synchronisation和屏障barrier：不同的线程在共享和全局内存中读写数据需要有先后的控制，所以引入了同步性的概念。

屏障的作用: 用来控制多个线程的停止与等待，当所有线程都到达了屏障点，程序才继续进行。

CUDA程序中CPU是主导地位，负责完成以下的事情：

从CPU同步数据到GPU
从GPU同步数据到CPU（1、2使用cudaMemcpy）
给GPU分配内存（cudaMalloc）
加载Kernel到GPU上，launch kernel on GPU

CUDA 代码

接下来便是对 CUDA 代码的编写。

CPU分配空间给GPU（cudaMalloc）
CPU复制数据给GPU（cudaMemcpy）
CPU加载kernels给GPU做计算
CPU把GPU计算结果复制回来

过程中，一般要尽量降低数据通讯的消耗，所以如果程序需要复制大量的数据到GPU，显然不是很合适使用GPU运算，最理想的情况是，每次复制的数据很小，然后运算量很大，输出的结果还是很小，复制回CPU。

第一个 CUDA 程序——并行地求平方

main.cu

#include <stdio.h>

__global__ void square(float* d_out, float* d_in) {
  int idx = threadIdx.x;
  float f = d_in[idx];
  d_out[idx] = f * f;
}

int main(int argc, char** argv) {
  const int ARRAY_SIZE = 8;
  const int ARRAY_BYTES = ARRAY_SIZE * sizeof(float);

  // generate the input array on the host
  float h_in[ARRAY_SIZE];
  for(int i=0; i<ARRAY_SIZE; i++) {
    h_in[i] = float(i);
  }
  float h_out[ARRAY_SIZE];

  // declare GPU memory pointers
  float* d_in;
  float* d_out;


  // allocate GPU memory
  cudaMalloc((void**) &d_in, ARRAY_BYTES);
  cudaMalloc((void**) &d_out, ARRAY_BYTES);

  // transfer the array to GPU
  cudaMemcpy(d_in, h_in, ARRAY_BYTES, cudaMemcpyHostToDevice);

  // launch the kernel
  square<<<1, ARRAY_SIZE>>>(d_out, d_in);

  // copy back the resoult array to the GPU
  cudaMemcpy(h_out, d_out, ARRAY_BYTES, cudaMemcpyDeviceToHost);

  // print out the resulting array
  for(int i=0; i<ARRAY_SIZE; i++) {
    printf("%f", h_out[i]);
    printf(((i%4) != 3)? "\t":"\n");
  }

  // free GPU memory allocation
  cudaFree(d_in);
  cudaFree(d_out);

  return 0;
}

编译: nvcc -o main.out main.cu

输出:

0.000000        1.000000        4.000000        9.000000
16.000000       25.000000       36.000000       49.000000

说明:

square<<<1, ARRAY_SIZE>>>(d_out, d_in); 对应参数: <<<blk, thr>>>

CUDA编译

使用nvcc编译cuda

参考文档: https://blog.csdn.net/fb_help/article/details/79283032

nvcc介绍

nvcc是编译cuda程序的编译器，CDUA C是在C语言上的扩展，所以它依赖C编译器（C编译器在window下是cl.exe，在Linux下是gcc）。因此我们编译CUDA程序必须依靠编译器nvcc。

其实，nvcc编译cuda程序和g++编译c++程序是差不多的。在我的其它博客中也写了有关g++编译c++程序的内容: g++ 命令的使用，可以作为参考。

示例

示例中创建了一个main.cu作为主程序入口，foo.cuh和foo.cu定义了一个函数实现。注意到文件数的后缀都是cuda程序可识别的后缀。

foo.cuh

#ifndef FOO_CUH
#define FOO_CUH

#include <stdio.h>

__global__ void foo();

extern "C" 
void useCUDA();

#endif // FOO_CUH

foo.cu

#include "foo.cuh"

#define CHECK(res) { if(res != cudaSuccess){printf("Error %s:%d , ", __FILE__,__LINE__);   \
printf("code : %d , reason : %s \n", res,cudaGetErrorString(res));exit(-1);}}

__global__ void foo() {
  printf("CUDA!\n");
}

void useCUDA() {
  foo<<<1, 5>>>();
  CHECK(cudaDeviceSynchronize())
}

main.cu:

#include <stdio.h>
#include <iostream>

#include "foo.cuh"

int main() {
  std::cout<<"Hello NVCC"<<std::endl;
  useCUDA();
  return 0;
}

编译: nvcc -o main.out src/main.cu src/foo.cu

运行 main.out 输出:

Hello NVCC
CUDA!
CUDA!
CUDA!
CUDA!
CUDA!

用CMake混合编译C++与cuda

参考文档: https://blog.csdn.net/fb_help/article/details/79330815

引言

许多c/c++的项目会使用cuda加速其算法，c/c++有其编译器：gcc/g++，cuda有其编译器nvcc。为了实现我们的目的，我们一般采用gcc/g++编译c/c++部分代码，nvcc编译cuda代码部分，即分离式编译。

分离式编译的基本思路：将cuda部分写成接口: void API(){...} 在c/c++代码中通过添加接口声明： extend "C" void API（）{...}

这样g++编译器就可以只处理C++有关代码，含有cuda代码的部分就由nvcc处理。只添加声明是不够的，要真正实现调用并完成功能，需要有真正的函数实现，即要将cuda接口制作成静态库或动态库，c++通过调用声明并调用库中的实现来完成功能。同时，我们常常需要构建一个具体的项目，CMake工具便是一个很好的选择。我们可以使用CMake工具，将cuda部分作为一个项目并制作为库，由其主项目调用。

示例

目录树:

cmake_cuda_demo
|-- CMakeLists.txt
`-- src
    |-- cuda
    |   |-- CMakeLists.txt
    |   |-- foo.cu
    |   `-- foo.cuh
    `-- main.cpp

2 directories, 5 files

main.cpp

#include <stdio.h>
#include <iostream>

#include "cuda/foo.cuh"

int main() {
  std::cout<<"Hello C++"<<std::endl;
  useCUDA();
  return 0;
}

CMakeLists.txt

cmake_minimum_required(VERSION 3.5)

project(test)

add_subdirectory(src/cuda)

set (EXTRA_LIBS ${EXTRA_LIBS} gpu)

add_executable(main src/main.cpp)
target_link_libraries(main ${EXTRA_LIBS})

cuda/foo.cuh

#ifndef FOO_CUH
#define FOO_CUH

#include <stdio.h>

extern "C"
void useCUDA();

#endif // FOO_CUH

foo.cu

#include "foo.cuh"

#define CHECK(res) { if(res != cudaSuccess){printf("Error %s:%d , ", __FILE__,__LINE__);   \
printf("code : %d , reason : %s \n", res,cudaGetErrorString(res));exit(-1);}}

__global__ void foo() {
  printf("CUDA!\n");
}

void useCUDA() {
  foo<<<1, 5>>>();
  CHECK(cudaDeviceSynchronize())
}

cuda/CMakeLists.txt

cmake_minimum_required(VERSION 3.5)

project(gpu)

find_package(CUDA REQUIRED)

#include_directories ("${PROJECT_SOURCE_DIR}")

# nvcc flags -g for debug
#set(CUDA_NVCC_FLAGS -O3;-G;-g)

#set(CUDA_NVCC_FLAGS -gencode arch=compute_20,code=sm_20;-G;-g)
#set(CUDA_NVCC_FLAGS -gencode arch=compute_52,code=sm_52;-G;-g)

file(GLOB_RECURSE CURRENT_HEADERS  *.h *.hpp *.cuh)
file(GLOB CURRENT_SOURCES  *.cpp *.cu)

source_group("Include" FILES ${CURRENT_HEADERS}) 
source_group("Source" FILES ${CURRENT_SOURCES}) 

#cuda_add_library(gpu SHARED ${CURRENT_HEADERS} ${CURRENT_SOURCES})
cuda_add_library(gpu STATIC ${CURRENT_HEADERS} ${CURRENT_SOURCES})

CUDA编程: zero copy

参考文档: https://blog.csdn.net/junparadox/article/details/50633641

零复制

zero copy（零复制）是一种特殊形式的内存映射，它允许你将host内存直接映射到设备内存空间上。其实就是设备可以通过直接内存访问（direct memory access，DMA）方式来访问主机的锁页内存。

锁页主机内存

现代操作系统都支持虚拟内存，操作系统实现虚拟内存的主要方法就是通过分页机制。操作系统将内存中暂时不使用的内容换出到外存（硬盘等大容量存储）上，从而腾出空间存放将要调入内存的信息。这样，系统好像为用户提供了一个比实际内存大得多的存储器，称为虚拟存储器。

锁页就是将内存页面标记为不可被操作系统换出的内存。所以设备驱动程序给这些外设编程时，可以使用页面的物理地址直接访问内存（DMA），从而避免从外存到内存的复制操作。CPU 仍然可以访问上述锁页内存，但是此内存是不能移动或换页到磁盘上的。CUDA 中把锁页内存称为pinned host memory 或者page-locked host memory。

锁页主机内存的优势

使用锁页内存（page-locked host memory）有一些优势：

锁页内存和GPU内存之间的拷贝可以和内核程序同时执行，也就是异步并发执行。
在一些设备上锁页内存的地址可以从主机地址空间映射到CUDA 地址空间，免去了拷贝开销。
在拥有前线总端的系统上，如果主机内存被分配为锁页内存，主机内存和GPU 内存带宽可以达到更高，如果主机内存被分配为Write-Combining Memory，带宽会进一步提升。

然而锁页主机存储器是稀缺资源，所以锁页内存分配得多的话，分配会失败。另外由于减少了系统可分页的物理存储器数量，分配太多的分页锁定内存会降低系统的整体性能。

使用锁页主机内存

在GPU 上分配的内存默认都是锁页内存，这只是因为GPU 不支持将内存交换到磁盘上。在主机上分配的内存默认都是可分页，如果需要分配锁页内存，则需要使用cudaMallocHost() 或者cudaHostAlloc()。释放时需要使用cudaFreeHost() 释放这一块内存。调用常规的C函数释放，可能会崩溃或者出现一些不常见的错误。也可以通过函数cudaHostRegister() 把可分页内存标记为锁页内存。

__host__ cudaError_t cudaMallocHost ( void** ptr, size_t size )

__host__ cudaError_t cudaHostAlloc ( void** pHost, size_t size, unsigned int  flags )

__host__ cudaError_t cudaFreeHost ( void* ptr )

cudaHostAlloc() 多了一个可选形参flags ，功能更强大。flags 的值可以取如下值。

#define cudaHostAllocDefault 0x00
Default page-locked allocation flag

#define cudaHostAllocMapped 0x02
Map allocation into device space

#define cudaHostAllocPortable 0x01
Pinned memory accessible by all CUDA contexts

#define cudaHostAllocWriteCombined 0x04
Write-combined memory

cudaHostRegister() 函数用于把已经的存在的可分页内存注册为分页锁定的。

__host__ cudaError_t cudaHostRegister ( void* ptr, size_t size, unsigned int  flags )

flags 是一个可选形参，可以取如下值。

#define cudaHostRegisterDefault 0x00
Default host memory registration flag

#define cudaHostRegisterIoMemory 0x04
Memory-mapped I/O space

#define cudaHostRegisterMapped 0x02
Map registered memory into device space

#define cudaHostRegisterPortable 0x01
Pinned memory accessible by all CUDA contexts

下面分别介绍这些flags 的作用。

Portable Memory

一块锁页内存可被系统中的所有设备使用（一个系统中有多个CUDA设备时）。启用这个特性需要在调用cudaHostAlloc() 时使用cudaHostAllocPortable 选项，或者在调用cudaHostRegister() 使用cudaHostRegisterPortable 选项。

Write-Combining Memory

默认情况下，锁页主机存储是可缓存的。可以在调用cudaHostAlloc() 时传入cudaHostAllocWriteCombined 标签使其被分配为写结合的（Write-Combining Memory）。写结合存储不使用L1 和L2 cache，所以程序的其它部分就有更多的缓存可用。此外，写结合内存通过PCI-E 传输数据时不会被监视（snoop），这能够获得高达40%的传输加速。从主机读取写结合存储非常慢（因为没有使用L1、L2cache），所以写结合存储应当只用于那些主机只写的存储。

Mapped Memory

一块锁页内存可以在调用cudaHostAlloc() 分配时传入cudaHostAllocMapped 标签或者在使用cudaHostRegister() 注册时使用cudaHostRegisterMapped 标签，把锁页内存地址映射到设备地址空间。这样，这块存储会有两个地址：一个是从cudaHostAlloc() 或malloc() 返回的在主机内存地址空间上；另一个在设备存储器上，可以通过cudaHostGetDevicePointer() 取得。内核函数可以使用这个指针访问这块存储。 cudaHostAlloc() 返回的地址指针一个的例外情况是，主机和设备使用统一地址空间（Unified Virtual Address Space）。

内核直接存取主机内存有很多优势：

无需在设备上分配内存，也无需在主机内存和设备内存之间拷贝数据。数据传输是在内核需要的时候隐式进行的。
无须使用流（cuda stream）就可以并发数据传输和内核执行；数据传输和内核执行自动并发执行。

因为映射的锁页主机内存是主机和设备之间共享的，所以在使用cuda stream 或者cuda event 时必须对内存读写同步；避免潜在的写后读，读后写或者写后写等多线程同步问题。

为了能够对任何映射的锁页主机内存解引用设备指针，必须在调用任何cuda 运行时函数前调用cudaSetDeviceFlags()，并传入cudaDeviceMapHost 标签。否则，cudaHostGetDevicePointer() 将会返回错误。

如果设备不支持被映射分页锁定存储，cudaHostGetDevicePointer() 将会返回错误。程序员可以检查canMapHostMemory 属性，如果设备支持映射锁页主机内存，将会返回1。

注意：使用映射锁页主机内存看，原子操作将不再保证原子性。cudaHostRegisterIoMemory 是cudaHostRegister() 特有的选项，可以把主机内存映射到IO 地址空间。

参考文献 [1]https://en.wikipedia.org/wiki/CUDA_Pinned_memory [2] Cook, Shane (2013). CUDA Programming: A Developer’s Guide to Parallel Computing with GPUs (1st ed.). Morgan Kaufmann Publishers Inc. pp. 334–335. ISBN 9780124159334.

vscode 配置

支持 cuda 语法

参考文档: https://blog.csdn.net/u011622208/article/details/113754328

插件：vscode插件商店搜索cudacpp进行安装后，可支持语法高亮以及«<»>等cuda专用符号。
参考《VScode 为 *.cu文件添加高亮及c++ intelligence相关操作的方法》，设置settings文件，添加文件cu后缀文件与cpp的关联：“files.associations”:{"*.cu":“cpp”}。设置完成后可支持cpp的语法高亮与跳转。

.vscode/settings.json

{
  "files.associations": {
    "*.cu":"cpp"
  }
}

vscode 配置 cuda 环境

参考文档: https://www.cnblogs.com/DLCannotBeAccelerated/p/15612820.html

安装扩展

vscode-cudacpp 代码高亮
Nsight Visual Studio Code Edition debug

配置文件 c_cpp_properties.json

{
    "configurations": [
        {
            "name": "Linux",
            "includePath": [
                "${workspaceFolder}/**",
                "/usr/local/cuda-10.2/include"
            ],
            "defines": [],
            "compilerPath": "/usr/bin/clang",
            "cStandard": "c11",
            "cppStandard": "c++14",
            "intelliSenseMode": "linux-clang-x64"
        }
    ],
    "version": 4
}

launch.json

{
    // Use IntelliSense to learn about possible attributes.
    // Hover to view descriptions of existing attributes.
    // For more information, visit: https://go.microsoft.com/fwlink/?linkid=830387
    "version": "0.2.0",
    "configurations": [
        {
            "name": "(gdb) Launch",
            "type": "cppdbg",
            "request": "launch",
            "program": "${fileDirname}/${fileBasenameNoExtension}.out",
            "args": [],
            "stopAtEntry": false,
            "cwd": "${fileDirname}",
            "environment": [],
            "externalConsole": false,
            "preLaunchTask": "build",
            "MIMode": "gdb",
            "setupCommands": [
                {
                    "description": "Enable pretty-printing for gdb",
                    "text": "-enable-pretty-printing",
                    "ignoreFailures": true
                }
            ],
            "sourceFileMap": {"/build/glibc-S9d2JN": "/usr/src/glibc"}
        }
    ]
}

tasks.json

{
    // See https://go.microsoft.com/fwlink/?LinkId=733558
    // for the documentation about the tasks.json format
    "version": "2.0.0",
    "tasks": [
        {
            "label": "build",
            "type": "shell",
            "command": "nvcc",
            "args":["-g","${file}","-o","${fileDirname}/${fileBasenameNoExtension}.out",
                // include 头文件
                "-I", "/usr/local/cuda/include",
                "-I", "/usr/local/cuda-10.2/samples/common/inc",
                // lib 库文件地址
                "-L", "/usr/local/cuda/lib64",   
                "-L", "/usr/local/cuda-10.2/samples/common/lib",  
                "-l", "cudart",                     
                "-l", "cublas",
                "-l", "cudnn",
                "-l", "curand",
                "-D_MWAITXINTRIN_H_INCLUDED"  
            ]
        }
    ]
}

安装glibc 这是因为调试cuda时，最后提示 vscode 无法打开 libc-start.c

sudo apt install glibc-source  
cd /usr/src/glibc/  
sudo tar -xvf glibc-2.27.tar.xz