【CUDA】《CUDA编程：基础与实践》Hello CUDA

CUDA Hello World!

简介

本文主要实现CUDA的 Hello World 。介绍了nvcc CUDA 编译时虚拟架构算力指定和真实架构算力指定，最后使用xmake编译CUDA程序。

CUDA 代码

• 一个真正利用GPU的CUDA程序既有主机Host（CPU）代码，也有设备Devie(GPU)代码。
• 主机对设备的调用时通过 核函数来实现。
  • 核函数可以使用 global 修饰
  • 核函数必须是 void

  // src/main.cu
  #include <stdio.h>
  __global__ void hello_from_gpu() {
  	printf("Hello World the GPU! \n");
  }
  
  int main() {
  	hello_from_gpu<<<1,1>>>();
  	cudaDeviceSynchronize();
  	return 0;
  }
  

• 核函数调用 hello_from_gpu<<<线程块数(), >>>

编译

虚拟算力和实际算力设置

CUDA 编译过程

• CUDA的编译器驱动 nvcc 先将全部源码分离为主机代码和设备代码。主机代码(CPU)完整地支持C++语法，但设备代码(GPU)只部分支持C++。
• nvcc先将设备代码编译为PTX(Parallel Thread eXecution)伪汇编代码。
• 将PTX代码编译为二进制cubin目标代码。

虚拟架构和真实架构算力指定

• 在将源码编译为PTX代码时，需要用选项 -arch=comput_XY 指定一个虚拟架构的计算能力，用以确定代码中能够使用的CUDA功能。
• 将PTX代码编译为cubin代码时，需要用选项-code=sm_ZW指定一个真实架构的计算能力。
• 真实架构的算力必须大于或等于指定算力的虚拟架构算力。

比如这里我们可以选择 虚拟架构算力为3.5 真实架构算力为6.0：

-arch=compute_35 -code=sm_60

但是不允许 虚拟架构算力大于真实架构算力

-arch=compute_60 -code=sm_35

如何设置？

如我们指定算力选项为：

-arch=compute_35 -code=sm_35

编译出来的可执行文件只能在帕斯卡架构的GPU中执行。如果希望编译出来的可执行文件能够在更多的GPU中执行，可以同时指定多组计算能力，每一组用如下形式的编译选项：

-gencode arch=compute_35 code=sm_35
-gencode arch=compute_50 code=sm_50
-gencode arch=compute_60 code=sm_60
-gencode arch=compute_70 code=sm_70

这样在不同GPU中运行会自动选择对应的二进制版本。

xmake 编译PTX编译设置

• 指定虚拟架构算力 add_cugencodes(“sm_35”)
• add_cugencodes(“native”) 根据当前CUDA 架构指定最低的算力
• 指定实际架构算力 add_cugencodes(“compute_35”)

 -- generate SASS code for SM architecture of current host
    add_cugencodes("native")

    -- generate PTX code for the virtual architecture to guarantee compatibility
    add_cugencodes("compute_35")

    -- -- generate SASS code for each SM architecture
    -- add_cugencodes("sm_35", "sm_37", "sm_50", "sm_52", "sm_60", "sm_61", "sm_70", "sm_75")

    -- -- generate PTX code from the highest SM architecture to guarantee forward-compatibility
    -- add_cugencodes("compute_75")

xmake 编译脚本

add_rules("mode.debug", "mode.release")

target("LearningCUDA")
    set_kind("binary")
    add_files("src/*.cu")

    -- generate relocatable device code for device linker of dependents.
    -- if __device__ or __global__ functions will be called cross file,
    -- or dynamic parallelism will be used,
    -- this instruction should be opted in.
    -- add_cuflags("-rdc=true")

    -- generate SASS code for SM architecture of current host
    add_cugencodes("native")

    -- generate PTX code for the virtual architecture to guarantee compatibility
    add_cugencodes("compute_35")

    -- -- generate SASS code for each SM architecture
    -- add_cugencodes("sm_35", "sm_37", "sm_50", "sm_52", "sm_60", "sm_61", "sm_70", "sm_75")

    -- -- generate PTX code from the highest SM architecture to guarantee forward-compatibility
    -- add_cugencodes("compute_75")
--