How to generate nvidia cuda bin (cubin) from MLIR?

[pal-stdr]

Hello,

I was trying to make the GPU compilation pipeline work for this example from llvm 17.0.6. But could not make it work.

And what I want for my RTX 3050 GPU (8.6 compute capability):

1. compile the .mlir file to gpu bin.
2. Then want to run the bin independently. (without using mlir-cpu-runner)
3. Want to profile it using Nvidia Nsight

Then I tried the solution by @bondhugula. Couldn’t make it work either.

Then started discussion in llvm discord with @fabianmc . Thanks to him , I have advanced a lot. But still couldn’t make it work. There you will find all the commands in chronological order that I have tried to make it work. I donot want to share it here, because it will take too long space.

But here I am going to share the final pieces that I have tried according to @fabianmc

./installation/bin/mlir-opt mlir/test/Integration/GPU/CUDA/gpu-to-cubin.mlir | ./installation/bin/mlir-opt -gpu-kernel-outlining | ./installation/bin/mlir-opt -pass-pipeline='builtin.module(gpu.module(strip-debuginfo,convert-gpu-to-nvvm,gpu-to-cubin))' | ./installation/bin/mlir-opt -gpu-to-llvm | ./installation/bin/mlir-translate --mlir-to-llvmir | ./installation/bin/llc -march=nvptx64 -mcpu=sm_86 - -o gpu.ptx

And then tried to compile the gpu.ptx file using nvcc

nvcc -rdc=true --cubin -arch=sm_86 gpu.ptx -o gpubin

But when I tried to run it

bash: ./gpubin: cannot execute binary file: Exec format error

The last response I’ve got from @fabianmc in discord is

the code you get here is not PTX, but rather llvm IR with the GPU module embedded as a string.

Now questions are

1. What am I missing in command pipeline? How can I make it work? Help much needed
2. Is there any official doc/tutorial? @herhut proposed an official doc for this years ago. . If I can find the way to properly organize the compilation pipeline, I am happy to write that doc as a contribution .
3. Advanced sorry for asking this very very noob question. I am using llvm-17.0.6. But got the idea from @mehdi_amini that main branch have all the latest cuda development for MLIR. But I am confused, the last commit in main branch is part of llvm-init-18? Or it is the latest for llvm-17? Logically it should be part of llvm 18. The reason for asking is to have a stable build, so that if I update it, the code doesn’t break.

Thanks in advance!!

[mehdi_amini]

@pal-stdr

And then tried to compile the gpu.ptx file using nvcc

nvcc -rdc=true --cubin -arch=sm_86 gpu.ptx -o gpubin

But when I tried to run it

bash: ./gpubin: cannot execute binary file: Exec format error

I’m a bit confused here: you’re only passing a gpu.ptx input to nvcc but there is no host code here, this can’t possibly generate a binary that you can run on the host, what am I missing?

The main branch is the current development branch for LLVM 18.
LLVM 17 is almost 5 months old now, so you’re missing on the last 5 months of developments in this area. It’s not much but at the same time many “solution” we may provide here won’t apply!

[jhuber6]

Like @mehdi_amini said, you need some form of host code to actually launch any kernels on the device. However, if you just want to call some arbitrary GPU function you could link your created binary in with some other kernel. e.g.

__device__ void your_func();

__global__ void kernel() { your_func(); }

int main() {
  kernel<<<1, 1, 1>>>();
  cudaDeviceSynchronize();
}

Forcing a GPU binary into the link job is a little difficult currently. I actually have a patch up for review that would make this slightly easier in clang at least [Clang][NVPTX] Allow passing arguments to the linker while standalone by jhuber6 · Pull Request #73030 · llvm/llvm-project · GitHub. That patch should allow you to do something like this.

clang++ -x cuda foo.cpp --offload-arch=native --offload-new-driver -fgpu-rdc -Xoffload-linker gpu.ptx

I’m assuming this is something close to what you’re asking for. If you’re talking about executing an NVIDIA GPU binary directly, I actually have some infrastructure for that in the libc project that simply executes a main function inside of a GPU executable. I use that to run the test suite.

[pal-stdr]

Thanks to you both for the reply.

@jhuber6

Forcing a GPU binary into the link job is a little difficult currently. I actually have a patch up for review that would make this slightly easier in clang at least [Clang][NVPTX] Allow passing arguments to the linker while standalone by jhuber6 · Pull Request #73030 · llvm/llvm-project · GitHub. That patch should allow you to do something like this.

clang++ -x cuda foo.cpp --offload-arch=native --offload-new-driver -fgpu-rdc -Xoffload-linker gpu.

@jhuber6 Here in the command, you have passed foo.cpp with clang++. But I don’t have any .cpp file to compile. I just want to compile .mlir.

@mehdi_amini

I’m a bit confused here: you’re only passing a gpu.ptx input to nvcc but there is no host code here, this can’t possibly generate a binary that you can run on the host, what am I missing?

@mehdi_amini & @jhuber6 I had the similar doubt about the host code. For your convenience, I am sharing the .mlir code here. I thought the @main() here is the host code. Because it is calling the kernel.

module attributes {gpu.container_module} {
    gpu.module @kernels {
        gpu.func @hello() kernel {
            %0 = gpu.thread_id x
            %csti8 = arith.constant 2 : i8
            %cstf32 = arith.constant 3.0 : f32
            gpu.printf "Hello from %lld, %d, %f\n" %0, %csti8, %cstf32  : index, i8, f32
            gpu.return
        }
    }

    func.func @main() {
        %c2 = arith.constant 2 : index
        %c1 = arith.constant 1 : index
        gpu.launch_func @kernels::@hello
            blocks in (%c1, %c1, %c1)
            threads in (%c2, %c1, %c1)
        return
    }
}

So if there are no host code here, how can I compile those example codes that live in this directory llvm-project/mlir/test/Integration/GPU/CUDA/?

Thanks in advance!

[grypp]

Does following pipeline work for you?

# compile cubin, and embed
mlir-opt printf.mlir -test-lower-to-nvvm="cubin-chip=sm_86 opt-level=3" -o print-cubin.mlir
# translate mlir to llvm
mlir-translate -mlir-to-llvmir print-cubin.mlir -o exe.ll
# compile ll, and link
clang++  exe.ll -o exe  -lmlir_c_runner_utils -lmlir_cuda_runtime
#run 
$ ./exe
Hello from 0, 2, 3.000000
Hello from 1, 2, 3.000000

[pal-stdr]

@grypp Thanks a lot.
It worked!!!

Some small issues

1. I have got this following warning telling me, clang++ doesn’t understand the flag --cuda-gpu-arch=sm_86. But its very important for me to generate binary for particular GPU arch.

clang++: warning: argument unused during compilation: '--cuda-gpu-arch=sm_86' [-Wunused-command-line-argument]
warning: overriding the module target triple with x86_64-unknown-linux-gnu [-Woverride-module]
1 warning generated.

2. I was also trying to chain all the commands using pipe | Like following. But couldn’t do it. It was giving me error messages

./installation/bin/mlir-opt mlir/test/Integration/GPU/CUDA/printf.mlir | ./installation/bin/mlir-opt -test-lower-to-nvvm -o print-cubin.mlir

# Error Message
Translations to perform option: must be specified at least once!
Translations to perform option: must be specified at least once!
Translations to perform option: must be specified at least once!
.....

Any idea?

Thanks again for your valuable help @grypp

[grypp]

I fixed the pipeline in my message above. You need to pass cubin-chip=sm_86 to test-lower-to-nvvm pipeline (no need for clang) as it actually compiles the gpu side of the code.

[pal-stdr]

@grypp

"cubin-chip=sm_86 opt-level=3"

Is giving me

<Pass-Options-Parser>: no such option opt-level

But it is working without the opt-level=3. So any clue, where I can pass the opt-level=3? Or is it because my build is llvm 17.0.6. And those pass options are defined in last updates in main branch (i.e. for llvm-18)?

And If I use only cubin-chip=sm_86, then at clang++ compilation step, it’s still giving me

warning: overriding the module target triple with x86_64-unknown-linux-gnu [-Woverride-module]
1 warning generated.

[fabianmc]

@pal-stdr

Or is it because my build is llvm 17.0.6.

All up to date features are not in llvm/17.x. In particular opt-level is not there.
Actually, most of new developments on GPU compilation are not there.

[pal-stdr]

I have one last doubt as newbie
It’s about this this warning when I’m compiling the final cubin (i.e. exe here) with clang++

warning: overriding the module target triple with x86_64-unknown-linux-gnu [-Woverride-module]

Note: I am now using the latest llvm-18 with commit hash ecf881838045985f381003cc27569c73a207d0cc

1. So first I am generating print-cubin.mlir file with following command. And this contains the target GPU arch sm_86

/installation/bin/mlir-opt mlir/test/Integration/GPU/CUDA/printf.mlir | ./installation/bin/mlir-opt -test-lower-to-nvvm='cubin-chip=sm_86 opt-level=3' -o print-cubin.mlir

2. Then I am generating the .ll file (i.e. exe.ll here), which also contains the target nvidia chip information

@fabianmc

exe.ll is not a module with GPU code, instead it is a host module with a string containing the already compiled cubin.

/installation/bin/mlir-translate -mlir-to-llvmir print-cubin.mlir -o exe.ll

3. Finally I am compiling the exe.ll file with clang++ which actually generates the module target triple x86_64-unknown-linux-gnu warning. [Doesn’t need to point the libs in linker path $LD_LIBRARY_PATH]

./installation/bin/clang++ exe.ll -o exe -L$LOCAL_EXTERNAL_SSD_COMPILER_PROJ_PATH/llvm-18-src-build/installation/lib -lmlir_c_runner_utils -lmlir_cuda_runtime

So, the 3rd step is overriding the module target triple. What does it even mean? As far as I know about the target triple

Target triple refers to a string in LLVM IR, that describes the target architecture, vendor, and operating system for which a program or a compiler is being built.

Then I have searched inside the cuda binary with both sm_86 and x86_64-unknown-linux-gnu strings using following command. sm_86 is found, but the other one is not found.

hexdump -C exe | grep -i 'sm_86'
# Result
# 00002d30  2e 31 0a 2e 74 61 72 67  65 74 20 73 6d 5f 38 36  |.1..target sm_86|

Then why clang++ giving me the warning?

[fabianmc]

@pal-stdr

Finally I am compiling the exe.ll file with clang++ which actually generates the module target triple x86_64-unknown-linux-gnu warning. [Doesn’t need to point the libs in linker path $LD_LIBRARY_PATH]

That warning happens because mlir-translate is not setting the target triple, so when clang compiles exe.ll it sets the target triple to the native one (effectively overriding the triple).

One more thing, exe.ll is not a module with GPU code, instead it is a host module with a string containing the already compiled cubin.

[pal-stdr]

Thanks a lot @fabianmc

@fabianmc

One more thing, exe.ll is not a module with GPU code, instead it is a host module with a string containing the already compiled cubin.

So the target triple is actually for the host (i.e. cpu) part?
If yes, then how can I add this with mlir-translate?

[fabianmc]

@pal-stdr

If yes, then how can I add this with mlir-translate?

You can’t with mlir-translate, llc or clang sets it for you. It’s just a warning, you shouldn’t worry.