14d8e42eaSAlex MacLean; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 5 2*b279f6b0SFangrui Song; RUN: llc < %s -mtriple=nvptx64 | FileCheck %s 34d8e42eaSAlex MacLean 44d8e42eaSAlex MacLeantarget triple = "nvptx64-nvidia-cuda" 54d8e42eaSAlex MacLean 64d8e42eaSAlex MacLeandefine i32 @test_addr_mode_i64(ptr %x) { 74d8e42eaSAlex MacLean; CHECK-LABEL: test_addr_mode_i64( 84d8e42eaSAlex MacLean; CHECK: { 94d8e42eaSAlex MacLean; CHECK-NEXT: .reg .b32 %r<2>; 104d8e42eaSAlex MacLean; CHECK-NEXT: .reg .b64 %rd<2>; 114d8e42eaSAlex MacLean; CHECK-EMPTY: 124d8e42eaSAlex MacLean; CHECK-NEXT: // %bb.0: 134d8e42eaSAlex MacLean; CHECK-NEXT: ld.param.u64 %rd1, [test_addr_mode_i64_param_0]; 144d8e42eaSAlex MacLean; CHECK-NEXT: ld.u32 %r1, [%rd1+-4]; 150f0a96b8SYoungsuk Kim; CHECK-NEXT: st.param.b32 [func_retval0], %r1; 164d8e42eaSAlex MacLean; CHECK-NEXT: ret; 174d8e42eaSAlex MacLean %addr = getelementptr i32, ptr %x, i64 -1 184d8e42eaSAlex MacLean %res = load i32, ptr %addr 194d8e42eaSAlex MacLean ret i32 %res 204d8e42eaSAlex MacLean} 214d8e42eaSAlex MacLean 224d8e42eaSAlex MacLeandefine i32 @test_addr_mode_i32(ptr %x) { 234d8e42eaSAlex MacLean; CHECK-LABEL: test_addr_mode_i32( 244d8e42eaSAlex MacLean; CHECK: { 254d8e42eaSAlex MacLean; CHECK-NEXT: .reg .b32 %r<2>; 264d8e42eaSAlex MacLean; CHECK-NEXT: .reg .b64 %rd<2>; 274d8e42eaSAlex MacLean; CHECK-EMPTY: 284d8e42eaSAlex MacLean; CHECK-NEXT: // %bb.0: 294d8e42eaSAlex MacLean; CHECK-NEXT: ld.param.u64 %rd1, [test_addr_mode_i32_param_0]; 304d8e42eaSAlex MacLean; CHECK-NEXT: ld.u32 %r1, [%rd1+-4]; 310f0a96b8SYoungsuk Kim; CHECK-NEXT: st.param.b32 [func_retval0], %r1; 324d8e42eaSAlex MacLean; CHECK-NEXT: ret; 334d8e42eaSAlex MacLean %addr = getelementptr i32, ptr %x, i32 -1 344d8e42eaSAlex MacLean %res = load i32, ptr %addr 354d8e42eaSAlex MacLean ret i32 %res 364d8e42eaSAlex MacLean} 374d8e42eaSAlex MacLean 384d8e42eaSAlex MacLeandefine i32 @test_addr_mode_i16(ptr %x) { 394d8e42eaSAlex MacLean; CHECK-LABEL: test_addr_mode_i16( 404d8e42eaSAlex MacLean; CHECK: { 414d8e42eaSAlex MacLean; CHECK-NEXT: .reg .b32 %r<2>; 424d8e42eaSAlex MacLean; CHECK-NEXT: .reg .b64 %rd<2>; 434d8e42eaSAlex MacLean; CHECK-EMPTY: 444d8e42eaSAlex MacLean; CHECK-NEXT: // %bb.0: 454d8e42eaSAlex MacLean; CHECK-NEXT: ld.param.u64 %rd1, [test_addr_mode_i16_param_0]; 464d8e42eaSAlex MacLean; CHECK-NEXT: ld.u32 %r1, [%rd1+-4]; 470f0a96b8SYoungsuk Kim; CHECK-NEXT: st.param.b32 [func_retval0], %r1; 484d8e42eaSAlex MacLean; CHECK-NEXT: ret; 494d8e42eaSAlex MacLean %addr = getelementptr i32, ptr %x, i16 -1 504d8e42eaSAlex MacLean %res = load i32, ptr %addr 514d8e42eaSAlex MacLean ret i32 %res 524d8e42eaSAlex MacLean} 534d8e42eaSAlex MacLean 544d8e42eaSAlex MacLeandefine i32 @test_addr_mode_i8(ptr %x) { 554d8e42eaSAlex MacLean; CHECK-LABEL: test_addr_mode_i8( 564d8e42eaSAlex MacLean; CHECK: { 574d8e42eaSAlex MacLean; CHECK-NEXT: .reg .b32 %r<2>; 584d8e42eaSAlex MacLean; CHECK-NEXT: .reg .b64 %rd<2>; 594d8e42eaSAlex MacLean; CHECK-EMPTY: 604d8e42eaSAlex MacLean; CHECK-NEXT: // %bb.0: 614d8e42eaSAlex MacLean; CHECK-NEXT: ld.param.u64 %rd1, [test_addr_mode_i8_param_0]; 624d8e42eaSAlex MacLean; CHECK-NEXT: ld.u32 %r1, [%rd1+-4]; 630f0a96b8SYoungsuk Kim; CHECK-NEXT: st.param.b32 [func_retval0], %r1; 644d8e42eaSAlex MacLean; CHECK-NEXT: ret; 654d8e42eaSAlex MacLean %addr = getelementptr i32, ptr %x, i8 -1 664d8e42eaSAlex MacLean %res = load i32, ptr %addr 674d8e42eaSAlex MacLean ret i32 %res 684d8e42eaSAlex MacLean} 694d8e42eaSAlex MacLean 704d8e42eaSAlex MacLeandefine i32 @test_addr_mode_i64_large(ptr %x) { 714d8e42eaSAlex MacLean; CHECK-LABEL: test_addr_mode_i64_large( 724d8e42eaSAlex MacLean; CHECK: { 734d8e42eaSAlex MacLean; CHECK-NEXT: .reg .b32 %r<2>; 744d8e42eaSAlex MacLean; CHECK-NEXT: .reg .b64 %rd<3>; 754d8e42eaSAlex MacLean; CHECK-EMPTY: 764d8e42eaSAlex MacLean; CHECK-NEXT: // %bb.0: 774d8e42eaSAlex MacLean; CHECK-NEXT: ld.param.u64 %rd1, [test_addr_mode_i64_large_param_0]; 784d8e42eaSAlex MacLean; CHECK-NEXT: add.s64 %rd2, %rd1, 17179869172; 794d8e42eaSAlex MacLean; CHECK-NEXT: ld.u32 %r1, [%rd2]; 800f0a96b8SYoungsuk Kim; CHECK-NEXT: st.param.b32 [func_retval0], %r1; 814d8e42eaSAlex MacLean; CHECK-NEXT: ret; 824d8e42eaSAlex MacLean %addr = getelementptr i32, ptr %x, i64 4294967293 834d8e42eaSAlex MacLean %res = load i32, ptr %addr 844d8e42eaSAlex MacLean ret i32 %res 854d8e42eaSAlex MacLean} 86