NVIDIA-CUDA HPC 编程模型与内存管理初探

高性能计算异步编程模型

CUDA计算是跨越主机和设备的并行计算，计算过程包含四个阶段，分别是：

1. 主机上的计算.
2. 主机到设备的数据传输.
3. 设备上的计算.
4. 设备将数据回传给主机.

如下图所示：

CUDA的全称是Compute Unified Device Architecture,是显卡厂商NVIDIA推出的运算平台，开发者可以使用C语言来编写CUDA代码，使用NVCC编译器可以在支持CUDA的GPU处理器上以高速运行。虽然AMD也做显卡，但是CUDA是老黄自家提出的标准，没带AMD一起玩儿，所以，提到基于CUDA的高性能计算，使用的都是Nvidia的显卡。

首先安装CUDA环境，具体方式参考博客：

FairMOT Cuda环境搭建并进行推理_tugouxp的专栏-CSDN博客环境准备1.PC Host Ubuntu 18.04.6,Linux Kernel 5.4，内核版本关系不大，记录下来备查。2.安装基础工具,比如GCC，CMAKE，VIM，GIT等等，工具尽量完备, 如果做不到，遇到问题临时下载也可。3.安装python3发行版，我用的是anaconda发行版，具体版本是 Anaconda3-2020.11-Linux-x86_64.sh下载地址在如下链接，选择对应的版本即可。https://repo.anaco...https://blog.csdn.net/tugouxp/article/details/121248457环境安装后，可以用如下方法验证环境：

nvidia-smi命令枚举了系统中的所有显卡支持信息

nvcc工具是CUDA编译器，用nvcc -V 验证编译器是否可以工作:

cuda编程

编辑helloworld.cu文件，编码内容:

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

int main(void)
{
	printf("hellow world!\n");
	return 0;
}

之后执行 nvcc helloworld.cu -o helloworld,并运行

可以看到，运行程序后打印除了helloworld.

但是，这个程序用到显卡了吗？很遗憾，没有。如果非要用显卡做点什么的化，可以改成这个样子：

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

__global__ void kernel(void)
{

}

int main(void)
{
	kernel<<<1,1>>>();

	printf("hellow world!\n");
	return 0;
}

我们定义了一个空函数送给GPU跑，函数是空函数，什么也没做，白嫖一下GPU就退出，编译并运行：

生成的helloworld文件是ELF格式的目标文件，与GCC产生的无异，可以通过objdump反编译一把：

来看一下main函数的片段：

粗略一看，首先给人的印象是NVCC不是一个人在战斗，毕竟我们的代码才短短几行，反编译后却有这么多条指令，而且貌似有些指令是没有出现在源码层面调用的。还能看出一点的就是源码是按照C++编译的，因为看到了明显的名字改编。

那就是编译器做的手脚咯,幸好我们有办法确认这一点，方式就是在nvcc编译的时候加上--verbose选项:

#$ _NVVM_BRANCH_=nvvm
#$ _SPACE_= 
#$ _CUDART_=cudart
#$ _HERE_=/usr/local/cuda-11.5/bin
#$ _THERE_=/usr/local/cuda-11.5/bin
#$ _TARGET_SIZE_=
#$ _TARGET_DIR_=
#$ _TARGET_DIR_=targets/x86_64-linux
#$ TOP=/usr/local/cuda-11.5/bin/..
#$ NVVMIR_LIBRARY_DIR=/usr/local/cuda-11.5/bin/../nvvm/libdevice
#$ LD_LIBRARY_PATH=/usr/local/cuda-11.5/bin/../lib::/usr/local/cuda-11.5/lib64
#$ PATH=/usr/local/cuda-11.5/bin/../nvvm/bin:/usr/local/cuda-11.5/bin:/home/caozilong/anaconda3/bin:/home/caozilong/anaconda3/condabin:/home/caozilong/.local/bin:/usr/local/sbin:/usr/local/bin:/usr/sbin:/usr/bin:/sbin:/bin:/usr/games:/usr/local/games:/snap/bin:/usr/local/cuda-11.5/bin
#$ INCLUDES="-I/usr/local/cuda-11.5/bin/../targets/x86_64-linux/include"  
#$ LIBRARIES=  "-L/usr/local/cuda-11.5/bin/../targets/x86_64-linux/lib/stubs" "-L/usr/local/cuda-11.5/bin/../targets/x86_64-linux/lib"
#$ CUDAFE_FLAGS=
#$ PTXAS_FLAGS=
#$ gcc -D__CUDA_ARCH__=520 -D__CUDA_ARCH_LIST__=520 -E -x c++  -DCUDA_DOUBLE_MATH_FUNCTIONS -D__CUDACC__ -D__NVCC__  "-I/usr/local/cuda-11.5/bin/../targets/x86_64-linux/include"    -D__CUDACC_VER_MAJOR__=11 -D__CUDACC_VER_MINOR__=5 -D__CUDACC_VER_BUILD__=50 -D__CUDA_API_VER_MAJOR__=11 -D__CUDA_API_VER_MINOR__=5 -D__NVCC_DIAG_PRAGMA_SUPPORT__=1 -include "cuda_runtime.h" -m64 "helloworld.cu" -o "/tmp/tmpxft_0000596f_00000000-9_helloworld.cpp1.ii" 
#$ cicc --c++14 --gnu_version=70500 --display_error_number --orig_src_file_name "helloworld.cu" --orig_src_path_name "/home/caozilong/cuda/helloworld.cu" --allow_managed   -arch compute_52 -m64 --no-version-ident -ftz=0 -prec_div=1 -prec_sqrt=1 -fmad=1 --include_file_name "tmpxft_0000596f_00000000-3_helloworld.fatbin.c" -tused --gen_module_id_file --module_id_file_name "/tmp/tmpxft_0000596f_00000000-4_helloworld.module_id" --gen_c_file_name "/tmp/tmpxft_0000596f_00000000-6_helloworld.cudafe1.c" --stub_file_name "/tmp/tmpxft_0000596f_00000000-6_helloworld.cudafe1.stub.c" --gen_device_file_name "/tmp/tmpxft_0000596f_00000000-6_helloworld.cudafe1.gpu"  "/tmp/tmpxft_0000596f_00000000-9_helloworld.cpp1.ii" -o "/tmp/tmpxft_0000596f_00000000-6_helloworld.ptx"
#$ ptxas -arch=sm_52 -m64  "/tmp/tmpxft_0000596f_00000000-6_helloworld.ptx"  -o "/tmp/tmpxft_0000596f_00000000-10_helloworld.sm_52.cubin" 
#$ fatbinary -64 --cicc-cmdline="-ftz=0 -prec_div=1 -prec_sqrt=1 -fmad=1 " "--image3=kind=elf,sm=52,file=/tmp/tmpxft_0000596f_00000000-10_helloworld.sm_52.cubin" "--image3=kind=ptx,sm=52,file=/tmp/tmpxft_0000596f_00000000-6_helloworld.ptx" --embedded-fatbin="/tmp/tmpxft_0000596f_00000000-3_helloworld.fatbin.c" 
#$ rm /tmp/tmpxft_0000596f_00000000-3_helloworld.fatbin
#$ gcc -D__CUDA_ARCH_LIST__=520 -E -x c++ -D__CUDACC__ -D__NVCC__  "-I/usr/local/cuda-11.5/bin/../targets/x86_64-linux/include"    -D__CUDACC_VER_MAJOR__=11 -D__CUDACC_VER_MINOR__=5 -D__CUDACC_VER_BUILD__=50 -D__CUDA_API_VER_MAJOR__=11 -D__CUDA_API_VER_MINOR__=5 -D__NVCC_DIAG_PRAGMA_SUPPORT__=1 -include "cuda_runtime.h" -m64 "helloworld.cu" -o "/tmp/tmpxft_0000596f_00000000-5_helloworld.cpp4.ii" 
#$ cudafe++ --c++14 --gnu_version=70500 --display_error_number --orig_src_file_name "helloworld.cu" --orig_src_path_name "/home/caozilong/cuda/helloworld.cu" --allow_managed  --m64 --parse_templates --gen_c_file_name "/tmp/tmpxft_0000596f_00000000-6_helloworld.cudafe1.cpp" --stub_file_name "tmpxft_0000596f_00000000-6_helloworld.cudafe1.stub.c" --module_id_file_name "/tmp/tmpxft_0000596f_00000000-4_helloworld.module_id" "/tmp/tmpxft_0000596f_00000000-5_helloworld.cpp4.ii" 
#$ gcc -D__CUDA_ARCH__=520 -D__CUDA_ARCH_LIST__=520 -c -x c++  -DCUDA_DOUBLE_MATH_FUNCTIONS "-I/usr/local/cuda-11.5/bin/../targets/x86_64-linux/include"   -m64 "/tmp/tmpxft_0000596f_00000000-6_helloworld.cudafe1.cpp" -o "/tmp/tmpxft_0000596f_00000000-11_helloworld.o" 
#$ nvlink -m64 --arch=sm_52 --register-link-binaries="/tmp/tmpxft_0000596f_00000000-7_helloworld_dlink.reg.c"    "-L/usr/local/cuda-11.5/bin/../targets/x86_64-linux/lib/stubs" "-L/usr/local/cuda-11.5/bin/../targets/x86_64-linux/lib" -cpu-arch=X86_64 "/tmp/tmpxft_0000596f_00000000-11_helloworld.o"  -lcudadevrt  -o "/tmp/tmpxft_0000596f_00000000-12_helloworld_dlink.sm_52.cubin"
#$ fatbinary -64 --cicc-cmdline="-ftz=0 -prec_div=1 -prec_sqrt=1 -fmad=1 " -link "--image3=kind=elf,sm=52,file=/tmp/tmpxft_0000596f_00000000-12_helloworld_dlink.sm_52.cubin" --embedded-fatbin="/tmp/tmpxft_0000596f_00000000-8_helloworld_dlink.fatbin.c" 
#$ rm /tmp/tmpxft_0000596f_00000000-8_helloworld_dlink.fatbin
#$ gcc -D__CUDA_ARCH_LIST__=520 -c -x c++ -DFATBINFILE="\"/tmp/tmpxft_0000596f_00000000-8_helloworld_dlink.fatbin.c\"" -DREGISTERLINKBINARYFILE="\"/tmp/tmpxft_0000596f_00000000-7_helloworld_dlink.reg.c\"" -I. -D__NV_EXTRA_INITIALIZATION= -D__NV_EXTRA_FINALIZATION= -D__CUDA_INCLUDE_COMPILER_INTERNAL_HEADERS__  "-I/usr/local/cuda-11.5/bin/../targets/x86_64-linux/include"    -D__CUDACC_VER_MAJOR__=11 -D__CUDACC_VER_MINOR__=5 -D__CUDACC_VER_BUILD__=50 -D__CUDA_API_VER_MAJOR__=11 -D__CUDA_API_VER_MINOR__=5 -D__NVCC_DIAG_PRAGMA_SUPPORT__=1 -m64 "/usr/local/cuda-11.5/bin/crt/link.stub" -o "/tmp/tmpxft_0000596f_00000000-13_helloworld_dlink.o" 
#$ g++ -D__CUDA_ARCH_LIST__=520 -m64 -Wl,--start-group "/tmp/tmpxft_0000596f_00000000-13_helloworld_dlink.o" "/tmp/tmpxft_0000596f_00000000-11_helloworld.o"   "-L/usr/local/cuda-11.5/bin/../targets/x86_64-linux/lib/stubs" "-L/usr/local/cuda-11.5/bin/../targets/x86_64-linux/lib"  -lcudadevrt  -lcudart_static  -lrt -lpthread  -ldl  -Wl,--end-group -o "helloworld" 

现在总结一下CUDA编程的规则：

• 核函数，在GPU上执行的函数通常成为核函数，如上面程序中的kernel函数。
• 核函数一般通过标识符__global__修饰，通过<<<参数1，参数2>>>调用，用于说明内核函数中的线程数量，以及线程是如何组织的。
• 以线程格(Grid)的形式组织，每个线程格有若干个线程块(block)组成，而每个线程块又由若干个线程(thread)组成。
• 以Block为单位执行
• 能在主机端代码中调用
• 调用时必须声明内核函数的执行参数
• 在编程时，必须先为kernel函数中用到的数组或者变量分配好足够的空间，再调用kernel函数，否则在GPU计算时会发生错误，例如越界或者报错，甚至导致蓝屏和死机。

CUDA的变成模型如下图所示：

上面例子中,kernel函数恰好叫kernel是一种巧合，实际上你可以改成任何有意义的名字，只要按照CUDA要求的方式调用即可

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

__global__ void dummy(void)
{

}

int main(void)
{
	dummy<<<1,1>>>();

	printf("hellow world!\n");
	return 0;
}

对于上如上的例子，我们探究一下它的控制流是如何进行的,首先我们看到反编译文件中，首先main函数调用了_Z5dummyv

 不难看出这个函数名是经过C++名字改编的，我们用c++filt工具将其还原：

可以看到它就是dummy,我们继续追踪