clang編譯CUDA程序，為何會生成多個LLVM bitcode文件?

01-14

請教一下，
這樣一個命令：
clang++ -g -emit-llvm axpy.cu -c

結果生成了兩個.bc文件bitcode文件，分別叫做axpy.bc和axpy-cuda-nvptx64-nvidia-cuda-sm_20.bc
大概看起來，第二個是device bitcode文件包含了CUDA kernel；第一個host bitcode則是所有東西，比如main，但也同樣含有cuda kernel；不過這裡的cuda kernel只是一個框架。

在後續的工作當中，如果我這樣編譯，
llc axpy.bc -o axpy.s
nvcc axpy.s -o axpy
整個程序可以可以運行不報錯，但是結果是不正確的。

但是host.bc在這過程中是怎麼參與進去的呢。。。？
整個原理是什麼？

附錄：
axpy.cu
#include &
__global__ void axpy(float a, float* x, float* y)

{
y[threadIdx.x] = a * x[threadIdx.x];
}

int main(int argc, char* argv[]) {...
...
// Copy data
// Launch the kernel. axpy&<&<&<1, kDataLen&>&>&>(a, device_x, device_y);
// Copy data...
...
return 0;
}

我在大約2015年的時候在Clang剛支持CUDA的時候做過研究，我知道你所遇到的問題，因為我當時也遇到過，但是我現在手上沒有CUDA的環境了，所以我只能憑藉著回憶來回答你的問題，但是時間有點久了，可能會有遺漏與不對的地方。

Clang會產生兩個.ll，那是由Clang編譯CUDA文件的編譯模型決定的。對於每一個.cu文件，Clang會編譯兩次，一次編譯代碼的device部分，一次編譯代碼的host部分，所以會dump出來兩個。其實你只要指定了正確的CUDA target選項，用Clang編譯其中那一個主的.ll 文件，Clang就會自動去編譯另外那一個nvidia-cuda-sm的.ll文件（不是用llc編譯。但是你用llc手動去編譯那兩個應該也是可以的，但是我沒有去嘗試過）。你直接這樣只編譯主的.ll的方式是可以編譯運行，但是你的結果會是不對的。因為Clang對於CUDA Kernerl的東西是放在一個特殊的ELF SECTION裡面的，所以你缺失了那個SECTION，也能運行，但是不對。

由於我沒有CUDA環境了（或者你可以給我一套CUDA的環境來讓我搞搞，然後來確認一下），所以我不能重現來仔細確認了，但是我能告訴你我當時是怎麼看的。我其實就很簡單，直接 -v，你就會發現過程了，然後查看產生的.o，注意觀察.section的東西，然後就能發現其中的奧秘了，一點都不複雜。