NVPTX: Cannot select: = NVPTXISD::LoadParam<(load (s24), align 4)> #55764

kjetilkjeka · 2022-05-29T20:38:37Z

I'm currently working on improving nvptx support in rust. The rust compiler produces llvm-ir with struct return values bitcast into integer values. This is fine for most targets, but for the triple nvptx64-nvidia-cuda llc cannot select for these types. A description going deeper into the rust side of things can be found here I will try to keep the rest of this issue focused only on the llvm side.

The problem from the LLVM side of things is that well formed llvm-ir cannot be compiled into ptx. I'm not sure if it is coincidental that clang always uses the actual struct as the return type, or if it is because it is better to work around this bug than to fix it. I do think that it will be a ugly hack in rustc to special case this for the nvptx target, and is thus willing to put a little amount of work into helping out fixing this bug if it is a realistic option.

I do realize this is only tested on a rather old version of llvm. I looked at your issue tracker and do not believe it is fixed. I apologize if this is a duplicate or have already been fixed.

Simplified example code

I believe this is a LLVM backend bug, and therefore the most interesting source is perhaps a simplified .ll example. The following is generated with bugpoint:

; ModuleID = 'bugpoint-reduced-simplified.bc'
source_filename = "main_rs.99b55f54-cgu.0"
target datalayout = "e-i64:64-i128:128-v16:16-v32:32-n16:32:64"
target triple = "nvptx64-nvidia-cuda"

%Foo.0.1.2.3.4.5 = type { i8, i8, i8 }

declare dso_local i24 @_ZN7main_rs6device17h13d51938dd622c57E(i8) unnamed_addr #0

define void @kernel(i8* %input, %Foo.0.1.2.3.4.5* %output) unnamed_addr #0 {
start:
  %0 = call i24 @_ZN7main_rs6device17h13d51938dd622c57E(i8 undef) #2
  unreachable
}

; Function Attrs: argmemonly nofree nounwind willreturn
declare void @llvm.memcpy.p0i8.p0i8.i64(i8* noalias nocapture writeonly, i8* noalias nocapture readonly, i64, i1 immarg) #1

attributes #0 = { "target-cpu"="sm_30" }
attributes #1 = { argmemonly nofree nounwind willreturn }
attributes #2 = { nounwind }

!llvm.module.flags = !{!0}

!0 = !{i32 7, !"PIC Level", i32 2}

Backtrace

When attempting to compile this with llc-14 from the Ubuntu 22.04 apt repo, the following is returned:

LLVM ERROR: Cannot select: 0xbebd58: i32,ch,glue = NVPTXISD::LoadParam<(load (s24), align 4)> 0xbebcf0, Constant:i32<1>, Constant:i32<0>, 0xbebcf0:1
  0xbeba80: i32 = Constant<1>
  0xbeb8e0: i32 = Constant<0>
  0xbebcf0: ch,glue = NVPTXISD::CallArgEnd 0xbebc88, Constant:i32<1>, 0xbebc88:1
    0xbeba80: i32 = Constant<1>
    0xbebc88: ch,glue = NVPTXISD::LastCallArg 0xbebc20, Constant:i32<1>, Constant:i32<0>, 0xbebc20:1
      0xbeba80: i32 = Constant<1>
      0xbeb8e0: i32 = Constant<0>
      0xbebc20: ch,glue = NVPTXISD::CallArgBegin 0xbebbb8, 0xbebbb8:1
        0xbebbb8: ch,glue = NVPTXISD::CallVoid 0xbebb50, 0xbeb740, 0xbebb50:1
          0xbeb740: i64 = NVPTXISD::Wrapper TargetGlobalAddress:i64<i24 (i8)* @_ZN7main_rs6device17h13d51938dd622c57E> 0
            0xbeb670: i64 = TargetGlobalAddress<i24 (i8)* @_ZN7main_rs6device17h13d51938dd622c57E> 0
          0xbebb50: ch,glue = NVPTXISD::PrintCallUni 0xbebae8, Constant:i32<1>, 0xbebae8:1
            0xbeba80: i32 = Constant<1>
            0xbebae8: ch,glue = NVPTXISD::DeclareRet 0xbeba18, Constant:i32<1>, Constant:i32<32>, Constant:i32<0>, 0xbeba18:1
              0xbeba80: i32 = Constant<1>
              0xbeb948: i32 = Constant<32>
              0xbeb8e0: i32 = Constant<0>
              0xbeba18: ch,glue = NVPTXISD::StoreParam<(store (s32))> 0xbeb9b0, Constant:i32<0>, Constant:i32<0>, Constant:i32<0>, 0xbeb9b0:1
                0xbeb8e0: i32 = Constant<0>
                0xbeb8e0: i32 = Constant<0>
                0xbeb8e0: i32 = Constant<0>
                0xbeb9b0: ch,glue = NVPTXISD::DeclareScalarParam 0xbeb878, Constant:i32<0>, Constant:i32<32>, Constant:i32<0>, 0xbeb878:1
                  0xbeb8e0: i32 = Constant<0>
                  0xbeb948: i32 = Constant<32>
                  0xbeb8e0: i32 = Constant<0>
                  0xbeb878: ch,glue = callseq_start 0xbb7bb8, TargetConstant:i64<0>, TargetConstant:i64<0>


In function: kernel
PLEASE submit a bug report to https://github.com/llvm/llvm-project/issues/ and include the crash backtrace.
Stack dump:
0.      Program arguments: llc-14 test.ll
1.      Running pass 'Function Pass Manager' on module 'test.ll'.
2.      Running pass 'NVPTX DAG->DAG Pattern Instruction Selection' on function '@kernel'
 #0 0x00007fce2350ed01 llvm::sys::PrintStackTrace(llvm::raw_ostream&, int) (/lib/x86_64-linux-gnu/libLLVM-14.so.1+0xe3fd01)
 #1 0x00007fce2350ca3e llvm::sys::RunSignalHandlers() (/lib/x86_64-linux-gnu/libLLVM-14.so.1+0xe3da3e)
 #2 0x00007fce2350f236 (/lib/x86_64-linux-gnu/libLLVM-14.so.1+0xe40236)
 #3 0x00007fce221b6520 (/lib/x86_64-linux-gnu/libc.so.6+0x42520)
 #4 0x00007fce2220aa7c __pthread_kill_implementation ./nptl/./nptl/pthread_kill.c:44:76
 #5 0x00007fce2220aa7c __pthread_kill_internal ./nptl/./nptl/pthread_kill.c:78:10
 #6 0x00007fce2220aa7c pthread_kill ./nptl/./nptl/pthread_kill.c:89:10
 #7 0x00007fce221b6476 gsignal ./signal/../sysdeps/posix/raise.c:27:6
 #8 0x00007fce2219c7f3 abort ./stdlib/./stdlib/abort.c:81:7
 #9 0x00007fce23449723 (/lib/x86_64-linux-gnu/libLLVM-14.so.1+0xd7a723)
#10 0x00007fce23d2fe61 (/lib/x86_64-linux-gnu/libLLVM-14.so.1+0x1660e61)
#11 0x00007fce23d2f308 llvm::SelectionDAGISel::SelectCodeCommon(llvm::SDNode*, unsigned char const*, unsigned int) (/lib/x86_64-linux-gnu/libLLVM-14.so.1+0x1660308)
#12 0x00007fce23d27eff llvm::SelectionDAGISel::DoInstructionSelection() (/lib/x86_64-linux-gnu/libLLVM-14.so.1+0x1658eff)
#13 0x00007fce23d275c4 llvm::SelectionDAGISel::CodeGenAndEmitDAG() (/lib/x86_64-linux-gnu/libLLVM-14.so.1+0x16585c4)
#14 0x00007fce23d269f8 llvm::SelectionDAGISel::SelectAllBasicBlocks(llvm::Function const&) (/lib/x86_64-linux-gnu/libLLVM-14.so.1+0x16579f8)
#15 0x00007fce23d24587 llvm::SelectionDAGISel::runOnMachineFunction(llvm::MachineFunction&) (/lib/x86_64-linux-gnu/libLLVM-14.so.1+0x1655587)
#16 0x00007fce238952be llvm::MachineFunctionPass::runOnFunction(llvm::Function&) (/lib/x86_64-linux-gnu/libLLVM-14.so.1+0x11c62be)
#17 0x00007fce236493c0 llvm::FPPassManager::runOnFunction(llvm::Function&) (/lib/x86_64-linux-gnu/libLLVM-14.so.1+0xf7a3c0)
#18 0x00007fce236509b3 llvm::FPPassManager::runOnModule(llvm::Module&) (/lib/x86_64-linux-gnu/libLLVM-14.so.1+0xf819b3)
#19 0x00007fce23649f66 llvm::legacy::PassManagerImpl::run(llvm::Module&) (/lib/x86_64-linux-gnu/libLLVM-14.so.1+0xf7af66)
#20 0x0000000000410223 main (/usr/lib/llvm-14/bin/llc+0x410223)
#21 0x00007fce2219dd90 __libc_start_call_main ./csu/../sysdeps/nptl/libc_start_call_main.h:58:16
#22 0x00007fce2219de40 call_init ./csu/../csu/libc-start.c:128:20
#23 0x00007fce2219de40 __libc_start_main ./csu/../csu/libc-start.c:379:5
#24 0x000000000040b065 _start (/usr/lib/llvm-14/bin/llc+0x40b065)
Aborted (core dumped)

Rust

#[repr(C)]
#[derive(Clone, Copy)]
pub struct Foo {
    a: u8,
    b: u8,
    c: u8,
}

#[inline(never)]
fn device(v: Foo) -> Foo {
    v
}

LLVM-IR

%Foo = type { i8, i8, i8 }

; main_rs::device
; Function Attrs: noinline nounwind
define internal i24 @_ZN7main_rs6device17h534b9a1ae91055fbE(i24 %0) unnamed_addr #1 {
start:
  %1 = alloca %Foo, align 1
  %2 = alloca i24, align 4
  %v = alloca %Foo, align 1
  store i24 %0, i24* %2, align 4
  %3 = bitcast %Foo* %v to i8*
  %4 = bitcast i24* %2 to i8*
  call void @llvm.memcpy.p0i8.p0i8.i64(i8* align 1 %3, i8* align 4 %4, i64 3, i1 false)
  %5 = bitcast %Foo* %1 to i8*
  %6 = bitcast %Foo* %v to i8*
  call void @llvm.memcpy.p0i8.p0i8.i64(i8* align 1 %5, i8* align 1 %6, i64 3, i1 false)
  %7 = bitcast %Foo* %1 to i24*
  %8 = load i24, i24* %7, align 1
  ret i24 %8
}

Copying between `Foo` and `i24`/`i32`

Rustc generates llvm-ir which uses the type Foo as the llvm-ir type %Foo = type { i8, i8, i8 } and as the llvm-ir type i24. Based on my observations, it seems like it is only bitcast into i24 when returning from or passing into functions.

If I understand correctly the reason this works is that bitcast is only being done on the pointers and not the aggregate type itself. The Foo and i24 is being converted between by doing a memcpy instead. I see that the pointer bitcast would still be compatible if we promoted i24 to i32.

My questions is: Are there any problems associated with only copying 3 bytes into the i32 that will be returned and leaving 8bits as undefined?

What (exactly) will need to be done.

I had a look in llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp but I couldn't figure out exactly how i24 is being lowered to ptx in general. My impression, mostly from looking at the methods and reading some code, is that it happens for all instructions individually. If this is correct then I assume some code needs to be appended to all the relevant methods (like LowerLOAD, LowerSTORE, LowerReturn, etc) and promote types to next level of power-of-two.

Is this understanding correct? Are you (@Artem-B) actively working on it or would it be beneficial if I took a stab at it?

Artem-B · 2022-06-09T17:52:56Z

Are you (@Artem-B) actively working on it or would it be beneficial if I took a stab at it?

It's not at the top of my todo list and I don't know when I'll get to it, so feel free to poke at it.

how i24 is being lowered to ptx in general.

It's not, which is the reason for the crash you've reported.
Considering that PTX does not have 24-bit load/store variant, if you insist on operating on i24 types, it would have to be split into smaller loads/stores. E.g. a triple of i8 ops, or a pair of i8/i16 or i16/i8, depending on alignment and some glue code to split/merge those bits to/from i24. That would be relatively slow.

Performance-wise it would probably be better not to use i24 and, if you do need to pack into integers, pack into power-of-2 sized ones that can be lowered into efficient code.

I'm not quite sure what drives this hand-packing of aggregates on IR level. Is that part of the Rust ABI? Normally it's target-dependent. E.g. on x64 small aggregates are passed in registers, power-of-2 sized, at that. Something like that may make sense for the rust-generated code for NVPTX.

kjetilkjeka · 2022-06-16T18:23:51Z

I had the chance to look a bit more into this today and due to being unfamiliar with the LLVM structure I hope to get my perceived understanding confirmed and maybe also some pointers in the right direction.

It seems to me like the Legalizer exists partially for turning non-supported llvm types into supported ones already on the LLVM-IR level (before lowering). It also seems like there are no implementation of class NVPTXLegalizerInfo : public LegalizerInfo which is the reason the Cannot select: occurs for nvptx64-nvidia-cuda, where other architectures with a Legalizer (like x86_64) will have the i24 turned into something else. Is this understanding correct?

Is the solution to the problem to add the Legalizer for NVPTX target? Are there any problems by doing this? A chance of breaking existing ptx code or similar?

The implementation for Legalizer seems quite different for different targets. X86 seems like something I expected, AMDGPU seems surprisingly long and detail oriented, while RISC-V gets away with a single call getLegacyLegalizerInfo().computeTables();. At the same time it is mentioned in the header that these files should be generated by TableGen and the word Legacy occurs a lot when reading the code of existing Legalizer. What is the best approach to finding out exactly what the NVPTXLegalizerInfo should look like?

Artem-B · 2022-06-16T18:41:32Z

Legalizer does sound like the right place to handle this.

A lot of things in LLVM grow 'organically' so the new features/mechanisms are not always automatically propagated to all back-ends. In general, small reasonable changes (as in both understandable and sensible for the issue at hand) are the way to go. As long as they move LLVM in the right direction (incremental steps are fine), have tests, and don't break anyone upstreaming them should be pretty straightforward.

kjetilkjeka · 2022-06-17T17:21:31Z

I'm struggling a bit with making progress. It seems like in the NVPTXTargetLowering constructor it is possible to register combination of operations and registers as Legal, Custom, Promote, etc to get the desired behavior using targetLoweringBase::setOperationAction. The problem is that the function only accepts MVT while, if I understand correctly, types like i24 is an EVT.

Do you have a pointer in the right direction of where the corresponding EVT behavior can be found?

Artem-B · 2022-06-17T21:18:04Z

I don't really know. If the standard lowering machinery can't be used with odd-sized integers, I'd try checking what happens if the same code gets compiled by a back-end that can deal with it. Track what they do differently and that may provide enough of a hint on how to do it in NVPTX.

It may be easier to consider changing the IR generator so it does not use the types NVPTX can't handle.

nikic · 2022-06-24T14:44:57Z

Simpler reproducer:

target datalayout = "e-i64:64-i128:128-v16:16-v32:32-n16:32:64"
target triple = "nvptx64-nvidia-cuda"

define i24 @kernel() {
  ret i24 0
}

LegalizeInfo is a GlobalISel concept that should not be relevant for NVPTX.

From a cursory look, this seems to be a simple bug in LowerReturn:

llvm-project/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp

Line 2644 in 8c278a2

RetTy->isIntegerTy() && DL.getTypeAllocSizeInBits(RetTy) < 32;

I think this should check getTypeSizeInBits() rather than getAllocSizeInBits(). For i24 the type size is 24, but the alloc size is 32.

Then an i24 return value will be handled by zero-extending to i32.

kjetilkjeka · 2022-06-24T15:43:29Z

Thanks for looking into this @nikic

I don't think it is entirely that simple. rustc promotes everything with size less than a pointer size to immediates. This means that the same must also be done for i40, i48 and i56. I also assume that a solution should preferably also work for non-divisible-by-eight bit widths?

I just found about the GlobalISel and that it is not used for NVPTX and that it is instead the NVPTXISelLowering I need to look at. I think it is correct that I need to do something similar to what you link from LowerReturn but I also think that it needs to be done for LowerCall and maybe also LowerFormalArguments.

I believe I can learn something from looking at ComputePTXValueVTs as it turns i128 into (i64, i64) for the mentioned functions. When attempting something similar I break a lot of tests. I assume this is due to transforming between types of different size and that I also need to do some bit extension of types. It's not possible to do this inside ComputePTXValueVTs as I do not have access to the DAG. I'm then not sure how it should interact with the ComputePTXValueVTs but my working hypothesis right now is that it needs to run directly after?

The status is more or less that I have managed to get the simple examples to compile correctly. But I am breaking tests and are trying to learn what I can from these. I'm slowly learning about how LLVM and NVPTX lowering works and I'm very happy for ideas and pointers like the one above.

kjetilkjeka · 2022-07-05T17:25:13Z

I made some progress today and believe I'm almost done!

The type promotion seems to work and can be viewed here
I've added some tests here

I've not yet ran the formating tools and friends (and will of course do that before submitting a patch) but if there are any problems to the general structure please let me know so I can fix them.

The final problem is that b48 and b56 is not turned into b64 in function signatures and I've started working on that here. But I'm unfortunately missing this promotion a few more places as my test output from the i48 test is.

; CHECK: .visible .func (.param .b64 func_retval0) callee
         ^
<stdin>:1:1: note: scanning from here
//
^
<stdin>:11:1: note: possible intended match here
.visible .func (.param .b48 func_retval0) callee(
^

Input file: <stdin>
Check file: <path>/llvm-project/llvm/test/CodeGen/NVPTX/i48-param.ll

-dump-input=help explains the following input dump.

Input was:
<<<<<<
           1: // 
check:6'0     X~~ error: no match found
           2: // Generated by LLVM NVPTX Back-End 
check:6'0     ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
           3: // 
check:6'0     ~~~
           4:  
check:6'0     ~
           5: .version 3.2 
check:6'0     ~~~~~~~~~~~~~
           6: .target sm_20 
check:6'0     ~~~~~~~~~~~~~~
           7: .address_size 32 
check:6'0     ~~~~~~~~~~~~~~~~~
           8:  
check:6'0     ~
           9:  // .globl callee // -- Begin function callee 
check:6'0     ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
          10:  // @callee 
check:6'0     ~~~~~~~~~~~~
          11: .visible .func (.param .b48 func_retval0) callee( 
check:6'0     ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
check:6'1     ?                                                  possible intended match
          12:  .param .b48 callee_param_0 
check:6'0     ~~~~~~~~~~~~~~~~~~~~~~~~~~~~
          13: ) 
check:6'0     ~~
          14: { 
check:6'0     ~~
          15:  .reg .b64 %rd<5>; 
check:6'0     ~~~~~~~~~~~~~~~~~~~
          16:  
check:6'0     ~

kjetilkjeka · 2022-07-07T14:29:52Z

I submitted a patch on phabricator today. https://reviews.llvm.org/D129291

…hen passing Today llc will crash when attempting to use non-power-of-two integer types as function arguments or returns. This patch enables passing non standard integer values in functions by promoting them before store and truncating after load. The main motivation of implementing this change is that rust casts small structs (less than pointer size) into an integer of the same size. As an example, if a struct contains three u8 then it will be passed as an i24. This patch is a step towards enabling rust compilation to ptx while retaining the target independent optimizations. More context can be found in #55764 Differential Revision: https://reviews.llvm.org/D129291

kjetilkjeka · 2022-08-08T09:50:19Z

Closing as the fix have been merged

…hen passing Today llc will crash when attempting to use non-power-of-two integer types as function arguments or returns. This patch enables passing non standard integer values in functions by promoting them before store and truncating after load. The main motivation of implementing this change is that rust casts small structs (less than pointer size) into an integer of the same size. As an example, if a struct contains three u8 then it will be passed as an i24. This patch is a step towards enabling rust compilation to ptx while retaining the target independent optimizations. More context can be found in llvm/llvm-project#55764 Differential Revision: https://reviews.llvm.org/D129291

github-actions bot added the new issue label May 29, 2022

kjetilkjeka mentioned this issue May 29, 2022

NVPTX: "LLVM ERROR: Cannot select" when returning struct with 3byte size from "device function" rust-lang/rust#97174

Closed

EugeneZelenko removed the new issue label May 29, 2022

Artem-B self-assigned this Jun 1, 2022

fhahn added llvm:crash backend:NVPTX labels Jul 7, 2022

kjetilkjeka closed this as completed Aug 8, 2022

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

NVPTX: Cannot select: = NVPTXISD::LoadParam<(load (s24), align 4)> #55764

NVPTX: Cannot select: = NVPTXISD::LoadParam<(load (s24), align 4)> #55764

kjetilkjeka commented May 29, 2022 •

edited by VoltrexKeyva

Loading

Artem-B commented Jun 1, 2022

kjetilkjeka commented Jun 2, 2022 •

edited

Loading

Artem-B commented Jun 2, 2022

kjetilkjeka commented Jun 2, 2022 •

edited

Loading

kjetilkjeka commented Jun 9, 2022 •

edited

Loading

Artem-B commented Jun 9, 2022 •

edited

Loading

kjetilkjeka commented Jun 16, 2022

Artem-B commented Jun 16, 2022

kjetilkjeka commented Jun 17, 2022

Artem-B commented Jun 17, 2022

nikic commented Jun 24, 2022

kjetilkjeka commented Jun 24, 2022

kjetilkjeka commented Jul 5, 2022

kjetilkjeka commented Jul 7, 2022

kjetilkjeka commented Aug 8, 2022

NVPTX: Cannot select: = NVPTXISD::LoadParam<(load (s24), align 4)> #55764

NVPTX: Cannot select: = NVPTXISD::LoadParam<(load (s24), align 4)> #55764

Comments

kjetilkjeka commented May 29, 2022 • edited by VoltrexKeyva Loading

Simplified example code

Backtrace

Meta

Artem-B commented Jun 1, 2022

kjetilkjeka commented Jun 2, 2022 • edited Loading

Artem-B commented Jun 2, 2022

kjetilkjeka commented Jun 2, 2022 • edited Loading

kjetilkjeka commented Jun 9, 2022 • edited Loading

Rust

LLVM-IR

Copying between Foo and i24/i32

What (exactly) will need to be done.

Artem-B commented Jun 9, 2022 • edited Loading

kjetilkjeka commented Jun 16, 2022

Artem-B commented Jun 16, 2022

kjetilkjeka commented Jun 17, 2022

Artem-B commented Jun 17, 2022

nikic commented Jun 24, 2022

kjetilkjeka commented Jun 24, 2022

kjetilkjeka commented Jul 5, 2022

kjetilkjeka commented Jul 7, 2022

kjetilkjeka commented Aug 8, 2022

kjetilkjeka commented May 29, 2022 •

edited by VoltrexKeyva

Loading

kjetilkjeka commented Jun 2, 2022 •

edited

Loading

kjetilkjeka commented Jun 2, 2022 •

edited

Loading

kjetilkjeka commented Jun 9, 2022 •

edited

Loading

Copying between `Foo` and `i24`/`i32`

Artem-B commented Jun 9, 2022 •

edited

Loading