OpenCL kernel (.cl) to RiscV assembly : "LLVM ERROR: Unsupported calling convention"

[Alexix24]

Hi,

I'm trying to build an OpenCL kernel (.cl) for a RiscV target using clang/llvm, but I'm facing a core dumped error when using llc.

• My OpenCL kernel (addVectors.cl):

__kernel void addVectors(__global const float *a, __global const float *b, __global float *c) {
    int gid = get_global_id(0);
    c[gid] = a[gid] + b[gid];
}

• .cl to .ll with clang (seems to work fine):

clang -c -emit-llvm -target riscv64 -o addVectors.ll addVectors.cl

• .ll to .s (riscv assembly) using llc (doesn't work):

llc addVectors.ll -o addVectors.s

LLVM error:

LLVM ERROR: Unsupported calling convention
PLEASE submit a bug report to https://github.com/llvm/llvm-project/issues/ and include the crash backtrace.
Stack dump:
0.	Program arguments: llc addVectors.ll -o addVectors.s
1.	Running pass 'Function Pass Manager' on module 'addVectors.ll'.
2.	Running pass 'RISCV DAG->DAG Pattern Instruction Selection' on function '@addVectors'
 #0 0x00007fe529c3fd01 llvm::sys::PrintStackTrace(llvm::raw_ostream&, int) (/lib/x86_64-linux-gnu/libLLVM-14.so.1+0xe3fd01)
 #1 0x00007fe529c3da3e llvm::sys::RunSignalHandlers() (/lib/x86_64-linux-gnu/libLLVM-14.so.1+0xe3da3e)
 #2 0x00007fe529c40236 (/lib/x86_64-linux-gnu/libLLVM-14.so.1+0xe40236)
 #3 0x00007fe528642520 (/lib/x86_64-linux-gnu/libc.so.6+0x42520)
 #4 0x00007fe5286969fc __pthread_kill_implementation ./nptl/./nptl/pthread_kill.c:44:76
 #5 0x00007fe5286969fc __pthread_kill_internal ./nptl/./nptl/pthread_kill.c:78:10
 #6 0x00007fe5286969fc pthread_kill ./nptl/./nptl/pthread_kill.c:89:10
 #7 0x00007fe528642476 gsignal ./signal/../sysdeps/posix/raise.c:27:6
 #8 0x00007fe5286287f3 abort ./stdlib/./stdlib/abort.c:81:7
 #9 0x00007fe529b7a723 (/lib/x86_64-linux-gnu/libLLVM-14.so.1+0xd7a723)
#10 0x00007fe529b7a556 (/lib/x86_64-linux-gnu/libLLVM-14.so.1+0xd7a556)
#11 0x00007fe52c1bee62 (/lib/x86_64-linux-gnu/libLLVM-14.so.1+0x33bee62)
#12 0x00007fe52a40475d llvm::SelectionDAGISel::LowerArguments(llvm::Function const&) (/lib/x86_64-linux-gnu/libLLVM-14.so.1+0x160475d)
#13 0x00007fe52a456516 llvm::SelectionDAGISel::SelectAllBasicBlocks(llvm::Function const&) (/lib/x86_64-linux-gnu/libLLVM-14.so.1+0x1656516)
#14 0x00007fe52a455557 llvm::SelectionDAGISel::runOnMachineFunction(llvm::MachineFunction&) (/lib/x86_64-linux-gnu/libLLVM-14.so.1+0x1655557)
#15 0x00007fe529fc628e llvm::MachineFunctionPass::runOnFunction(llvm::Function&) (/lib/x86_64-linux-gnu/libLLVM-14.so.1+0x11c628e)
#16 0x00007fe529d7a390 llvm::FPPassManager::runOnFunction(llvm::Function&) (/lib/x86_64-linux-gnu/libLLVM-14.so.1+0xf7a390)
#17 0x00007fe529d81983 llvm::FPPassManager::runOnModule(llvm::Module&) (/lib/x86_64-linux-gnu/libLLVM-14.so.1+0xf81983)
#18 0x00007fe529d7af36 llvm::legacy::PassManagerImpl::run(llvm::Module&) (/lib/x86_64-linux-gnu/libLLVM-14.so.1+0xf7af36)
#19 0x0000000000410223 main (/usr/lib/llvm-14/bin/llc+0x410223)
#20 0x00007fe528629d90 __libc_start_call_main ./csu/../sysdeps/nptl/libc_start_call_main.h:58:16
#21 0x00007fe528629e40 call_init ./csu/../csu/libc-start.c:128:20
#22 0x00007fe528629e40 __libc_start_main ./csu/../csu/libc-start.c:379:5
#23 0x000000000040b065 _start (/usr/lib/llvm-14/bin/llc+0x40b065)
Aborted (core dumped)

Thanks !

[dtcxzyw]

The calling convention SPIR_KERNEL is not supported yet.

[Alexix24]

So I can't build the kernel for a RiscV target using LLVM ? I thought opencl implementations used llvm to compile kernels.

[llvmbot]

@llvm/issue-subscribers-backend-risc-v

Author: None (Alexix24)

Hi,

I'm trying to build an OpenCL kernel (.cl) for a RiscV target using clang/llvm, but I'm facing a core dumped error when using llc.

• My OpenCL kernel (addVectors.cl):

__kernel void addVectors(__global const float *a, __global const float *b, __global float *c) {
    int gid = get_global_id(0);
    c[gid] = a[gid] + b[gid];
}

• .cl to .ll with clang (seems to work fine):

clang -c -emit-llvm -target riscv64 -o addVectors.ll addVectors.cl

• .ll to .s (riscv assembly) using llc (doesn't work):

llc addVectors.ll -o addVectors.s

LLVM error:

LLVM ERROR: Unsupported calling convention
PLEASE submit a bug report to https://github.com/llvm/llvm-project/issues/ and include the crash backtrace.
Stack dump:
0.	Program arguments: llc addVectors.ll -o addVectors.s
1.	Running pass 'Function Pass Manager' on module 'addVectors.ll'.
2.	Running pass 'RISCV DAG-&gt;DAG Pattern Instruction Selection' on function '@<!-- -->addVectors'
 #<!-- -->0 0x00007fe529c3fd01 llvm::sys::PrintStackTrace(llvm::raw_ostream&amp;, int) (/lib/x86_64-linux-gnu/libLLVM-14.so.1+0xe3fd01)
 #<!-- -->1 0x00007fe529c3da3e llvm::sys::RunSignalHandlers() (/lib/x86_64-linux-gnu/libLLVM-14.so.1+0xe3da3e)
 #<!-- -->2 0x00007fe529c40236 (/lib/x86_64-linux-gnu/libLLVM-14.so.1+0xe40236)
 #<!-- -->3 0x00007fe528642520 (/lib/x86_64-linux-gnu/libc.so.6+0x42520)
 #<!-- -->4 0x00007fe5286969fc __pthread_kill_implementation ./nptl/./nptl/pthread_kill.c:44:76
 #<!-- -->5 0x00007fe5286969fc __pthread_kill_internal ./nptl/./nptl/pthread_kill.c:78:10
 #<!-- -->6 0x00007fe5286969fc pthread_kill ./nptl/./nptl/pthread_kill.c:89:10
 #<!-- -->7 0x00007fe528642476 gsignal ./signal/../sysdeps/posix/raise.c:27:6
 #<!-- -->8 0x00007fe5286287f3 abort ./stdlib/./stdlib/abort.c:81:7
 #<!-- -->9 0x00007fe529b7a723 (/lib/x86_64-linux-gnu/libLLVM-14.so.1+0xd7a723)
#<!-- -->10 0x00007fe529b7a556 (/lib/x86_64-linux-gnu/libLLVM-14.so.1+0xd7a556)
#<!-- -->11 0x00007fe52c1bee62 (/lib/x86_64-linux-gnu/libLLVM-14.so.1+0x33bee62)
#<!-- -->12 0x00007fe52a40475d llvm::SelectionDAGISel::LowerArguments(llvm::Function const&amp;) (/lib/x86_64-linux-gnu/libLLVM-14.so.1+0x160475d)
#<!-- -->13 0x00007fe52a456516 llvm::SelectionDAGISel::SelectAllBasicBlocks(llvm::Function const&amp;) (/lib/x86_64-linux-gnu/libLLVM-14.so.1+0x1656516)
#<!-- -->14 0x00007fe52a455557 llvm::SelectionDAGISel::runOnMachineFunction(llvm::MachineFunction&amp;) (/lib/x86_64-linux-gnu/libLLVM-14.so.1+0x1655557)
#<!-- -->15 0x00007fe529fc628e llvm::MachineFunctionPass::runOnFunction(llvm::Function&amp;) (/lib/x86_64-linux-gnu/libLLVM-14.so.1+0x11c628e)
#<!-- -->16 0x00007fe529d7a390 llvm::FPPassManager::runOnFunction(llvm::Function&amp;) (/lib/x86_64-linux-gnu/libLLVM-14.so.1+0xf7a390)
#<!-- -->17 0x00007fe529d81983 llvm::FPPassManager::runOnModule(llvm::Module&amp;) (/lib/x86_64-linux-gnu/libLLVM-14.so.1+0xf81983)
#<!-- -->18 0x00007fe529d7af36 llvm::legacy::PassManagerImpl::run(llvm::Module&amp;) (/lib/x86_64-linux-gnu/libLLVM-14.so.1+0xf7af36)
#<!-- -->19 0x0000000000410223 main (/usr/lib/llvm-14/bin/llc+0x410223)
#<!-- -->20 0x00007fe528629d90 __libc_start_call_main ./csu/../sysdeps/nptl/libc_start_call_main.h:58:16
#<!-- -->21 0x00007fe528629e40 call_init ./csu/../csu/libc-start.c:128:20
#<!-- -->22 0x00007fe528629e40 __libc_start_main ./csu/../csu/libc-start.c:379:5
#<!-- -->23 0x000000000040b065 _start (/usr/lib/llvm-14/bin/llc+0x40b065)
Aborted (core dumped)

Thanks !

[dtcxzyw]

So I can't build the kernel for a RiscV target using LLVM ? I thought opencl implementations used llvm to compile kernels.

It is not the correct way for OpenCL runtime to generate device code. See also https://github.com/THU-DSP-LAB/llvm-project for the LLVM-based toolchain of an open-source RISC-V GPU.

BTW, the fix #69282 makes sense to me.

[wangpc-pp]

So I can't build the kernel for a RiscV target using LLVM ? I thought opencl implementations used llvm to compile kernels.

It is not the correct way for OpenCL runtime to generate device code. See also https://github.com/THU-DSP-LAB/llvm-project for the LLVM-based toolchain of an open-source RISC-V GPU.

OpenCL can run on CPU devices, so compiling kernel functions to CPU code (host?) is valid I think.

[Alexix24]

So I can't build the kernel for a RiscV target using LLVM ? I thought opencl implementations used llvm to compile kernels.

It is not the correct way for OpenCL runtime to generate device code. See also https://github.com/THU-DSP-LAB/llvm-project for the LLVM-based toolchain of an open-source RISC-V GPU.

OpenCL can run on CPU devices, so compiling kernel functions to CPU code (host?) is valid I think.

Sorry I didn't explain the context enough. I'm trying to develop my own OpenCL driver/implementation for the following architecture: a CPU (host) and a FPGA running RiscV cores (device). That's why I'm trying to build OpenCL kernels for a RiscV target, in order to run the kernel on the FPGA. But maybe that's not the right way, I'm still learning OpenCL. If you have any tips about it, it would be very useful (but I know it's not an OpenCL forum).

I just tried the fix #69282 and it works to build .cl into RiscV assembler, thanks a lot for your help !