NVCC

目录

• 0x00 基础知识
• 0x01 EXAMPLE
  • 整体编译
  • 分离编译
  • makefile
• 0x03 例子3 - 分离编译
• 0x02 NVCC编译选项
  • cuda_arch/cuda_code
  • example1
  • example2

0x00 基础知识

Prior to the 5.0 release, CUDA did not support separate compilation, so CUDA code could not call device functions or access variables across files. Such compilation is referred to as whole program compilation. We have always supported the separate compilation of host code, it was just the device CUDA code that needed to all be within one file. Starting with CUDA 5.0, separate compilation of device code is supported, but the old whole program mode is still the default, so there are new options to invoke separate compilation.
CUDA5.0版本前没有设备端的独立编译，所以CUDA的代码不能跨文件访问函数和变量，这在当时称为整体编译，值得注意的是，无论CUDA什么版本，都支持host端的独立编译(例如g++编译器，要注意区分设备端独立编译和host端独立编译是不同的概念)，在这种情况下，cuda的代码要放在一个文件里。从5.0以后，开始支持device端的独立编译(nvcc编译器)。nvcc默认还是整体编译，使用独立编译需要设置nvcc参数。

CUDA works by embedding device code into host objects. In whole program compilation, it embeds executable device code into the host object. In separate compilation, we embed relocatable device code into the host object, and run nvlink, the device linker, to link all the device code together. The output of nvlink is then linked together with all the host objects by the host linker to form the final executable.
cuda编译一般将设备代码嵌入到host端代码完成编译。在whole program compilation中，嵌入host的是可执行的设备端代码，在separate compilation中，把可重定位设备端代码嵌入到host，再使用device linker完成链接。device linker的输出文件交给host linker完成最后的链接并声称可执行文件。

The generation of relocatable vs executable device code is controlled by the --relocatable-device-code option.
设备端在编译cu文件时生成的是可执行文件还是可重定位代码是由nvcc选项 --relocatable-device-code 决定的。
The --compile option is already used to control stopping a compile at a host object, so a new option --device-c is added that simply does --relocatable-device-code=true --compile.
nvcc的编译选项--compile已经使得编译截止到host。因此使用nvcc的选项--device-c等价于--relocatable-device-code=true --compile这个组合参数，意义是截止到host前，将cu编译成可重定位文件

To invoke just the device linker, the --device-link option can be used, which emits a host object containing the embedded executable device code. The output of that must then be passed to the host linker. Or:
nvcc
can be used to implicitly call both the device and host linkers. This works because if the device linker does not see any relocatable code it does not do anything.
使用--device-link选项生成一个包含可执行device代码的host文件。再将这个host文件用host的linker链接。

关于CUDA的编译过程可以参考这个文档，https://docs.nvidia.com/cuda/cuda-compiler-driver-nvcc/index.html#using-separate-compilation-in-cuda

0x01 EXAMPLE

整体编译

举个例子，在cmu15418的作业中，有写好的makefile，如下

  1
   2 EXECUTABLE := cudaSaxpy
   3
   4 CU_FILES   := saxpy.cu
   5
   6 CU_DEPS    :=
   7
   8 CC_FILES   := main.cpp
   9
  10 ###########################################################
  11
  12 ARCH=$(shell uname | sed -e 's/-.*//g')
  13
  14 OBJDIR=objs
  15 CXX=g++ -m64
  16 CXXFLAGS=-O3 -Wall
  17 ifeq ($(ARCH), Darwin)
  18 # Building on mac
  19 LDFLAGS=-L/usr/local/depot/cuda-8.0/lib/ -lcudart
  20 else
  21 # Building on Linux
  22 # LDFLAGS=-L/usr/local/cuda-11.2/lib64/ -lcudart
  23 LDFLAGS=-L/data/cuda/cuda-11.1/cuda/lib64 -lcudart
  24 endif
  25 NVCC=nvcc
  26 NVCCFLAGS=-O3 -m64 #--gpu-architecture compute_35
  27
  28
  29 OBJS=$(OBJDIR)/main.o  $(OBJDIR)/saxpy.o
  30
  31
  32 .PHONY: dirs clean
  33
  34 default: $(EXECUTABLE)
  35
  36 dirs:
  37         mkdir -p $(OBJDIR)/
  38
  39 clean:
  40         rm -rf $(OBJDIR) *.ppm *~ $(EXECUTABLE)
  41
  42 $(EXECUTABLE): dirs $(OBJS)
  43         $(CXX) $(CXXFLAGS) -o $@ $(OBJS) $(LDFLAGS)
  44
  45 $(OBJDIR)/%.o: %.cpp
  46         $(CXX) $< $(CXXFLAGS) -c -o $@
  47
  48 $(OBJDIR)/%.o: %.cu
  49         $(NVCC) $< $(NVCCFLAGS) -c -o $@

default -> executable -> objs -> objdir/main.o & objdir/saxpy.o
objdir/main.o -> main.cpp & main.cu
main.cpp -> g++ -m64 main.cpp -O3 -Wall -c -o main.o
objdir/saxpy.o -> saxpy.cpp & saxpy.cu
saxpy.cu -> nvcc saxpy.cu -O3 -m64 -c -o saxpy.o
dirs objs/main.o objs/saxpy.o -> g++ -m64 -o cudaSaxpy objs/main.o objs/saxpy.o -L/data/cuda/cuda-11.1/cuda/lib64 -lcudart

分离编译

在15418的assignment2的cuda render中，由于使用了动态并行，所以必须使用分离编译。这和cuda文档中给出的使用分离编译的原因不同（一般的分离编译是因为cu分散到了不同的文件中且存在变量的引用）
关于makefile中特殊符号的定义:
https://stackoverflow.com/questions/3220277/what-do-the-makefile-symbols-and-mean

makefile

这里主要需要注意的是(可能与其他的分离编译情况有区别)

1. 最终的host编译器需要加上-lcudadevrt 和 device-linker的输出文件(link.o)
2. 在编译cu时添加分离编译的选项，--device-c
3. 将可重定位的device代码使用device-linker进行链接
   
   code : https://github.com/ijpq/15418/commit/6b4d0617f072ad5239d8fda9e72cbcde01b08508

0x03 例子3 - 分离编译

a.h

void warperFoo();

a.cu

//---------- a.cu ----------
#include <cuda.h>
#include <cuda_runtime.h>
#include <stdio.h>
#include "b.h"
#include "a.h"


__global__ void foo (void) {
  printf("calling from kernel foo: %d\n", threadIdx.x);
  bar();
}

void warperFoo() {
    printf("calling from warperFoo\n");
    dim3 gdim(1);
    dim3 bdim(4);
    foo<<<gdim, bdim>>>();
    cudaDeviceSynchronize();
}

b.h

//---------- b.h ----------
#define N 8

extern __device__ int g[N];

extern __device__ void bar(void);

b.cu

#include "b.h"
#include <stdio.h>

__device__ int g[N];

__device__ void bar (void)
{
    printf("calling from kernel bar: %d \n", threadIdx.x);
    g[threadIdx.x]++;

}

main.cpp

#include <iostream>
#include "a.h"

using namespace std;


int main() {
    warperFoo();
}

makefile

.PHONY: clean
all:
        nvcc --gpu-architecture=sm_75 -dc a.cu b.cu
        nvcc --gpu-architecture=sm_75 -dlink a.o b.o -o link.o
        g++ -m64 -Wall a.o b.o link.o main.cpp -lcudart -L/data/cuda/cuda-10.2/cuda/lib64/
        ./a.out

clean:
        rm -rf *.o a.out

output

nvcc --gpu-architecture=sm_75 -dc a.cu b.cu
nvcc --gpu-architecture=sm_75 -dlink a.o b.o -o link.o
g++ -m64 -Wall a.o b.o link.o main.cpp -lcudart -L/data/cuda/cuda-10.2/cuda/lib64/
./a.out
calling from warperFoo
calling from kernel foo: 0
calling from kernel foo: 1
calling from kernel foo: 2
calling from kernel foo: 3
calling from kernel bar: 0
calling from kernel bar: 1
calling from kernel bar: 2
calling from kernel bar: 3

1. 这个例子中，将两个cu分离编译，再device link. 最后使用host的g++ 编译链接
2. main.cpp中调用kernel函数的wraper可以直接声明而不包含a.h，如下

#include <iostream>
// #include "a.h"

using namespace std;


void warperFoo();
int main() {
    warperFoo();
}

3. host端也可以先编译.cpp，再一起使用host linker将device的重定位和host的重定位链接起来，如下

.PHONY: clean
all:
        nvcc --gpu-architecture=sm_75 -dc a.cu b.cu
        nvcc --gpu-architecture=sm_75 -dlink a.o b.o -o link.o
        g++ -m64 -Wall -c main.cpp -o main.o
        g++ -m64 -Wall a.o b.o link.o main.o -lcudart -L/data/cuda/cuda-10.2/cuda/lib64/
        ./a.out

clean:
        rm -rf *.o a.out

0x02 NVCC编译选项

cuda_arch/cuda_code

--gpu-architecture指定virtual arch，一般是一个单独的值
--gpu-code指定real arch，一般是一系列值。nvcc对这个选项中的每一个real arch，执行ptx生成bin文件的过程.

example1

--gpu-code 也可以指定virtual arch. 那么此时，由ptx生成bin文件的过程将不会进行. 只会保留一个ptx文件。运行时，driver发现没有生成好的bin文件时，会将这个ptx文件作为输入，生成一遍bin.

example2

可以不指定--gpu-code的值，只有在这个情况下，--gpu-arch可以是一个real arch.
那么gpu code值是和--gpu-arch指定的最为相近的virtual arch值。
如果在不指定--gpu-code值的时候，--gpu-arch还是一个virtual值，那么--gpu-code也将会是一个virtual arch值，等效于example1中只生成ptx的效果

\