86 lines · plain
1; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 52; RUN: llc < %s -mtriple=nvptx64 | FileCheck %s3 4target triple = "nvptx64-nvidia-cuda"5 6define i32 @test_addr_mode_i64(ptr %x) {7; CHECK-LABEL: test_addr_mode_i64(8; CHECK: {9; CHECK-NEXT: .reg .b32 %r<2>;10; CHECK-NEXT: .reg .b64 %rd<2>;11; CHECK-EMPTY:12; CHECK-NEXT: // %bb.0:13; CHECK-NEXT: ld.param.b64 %rd1, [test_addr_mode_i64_param_0];14; CHECK-NEXT: ld.b32 %r1, [%rd1+-4];15; CHECK-NEXT: st.param.b32 [func_retval0], %r1;16; CHECK-NEXT: ret;17 %addr = getelementptr i32, ptr %x, i64 -118 %res = load i32, ptr %addr19 ret i32 %res20}21 22define i32 @test_addr_mode_i32(ptr %x) {23; CHECK-LABEL: test_addr_mode_i32(24; CHECK: {25; CHECK-NEXT: .reg .b32 %r<2>;26; CHECK-NEXT: .reg .b64 %rd<2>;27; CHECK-EMPTY:28; CHECK-NEXT: // %bb.0:29; CHECK-NEXT: ld.param.b64 %rd1, [test_addr_mode_i32_param_0];30; CHECK-NEXT: ld.b32 %r1, [%rd1+-4];31; CHECK-NEXT: st.param.b32 [func_retval0], %r1;32; CHECK-NEXT: ret;33 %addr = getelementptr i32, ptr %x, i32 -134 %res = load i32, ptr %addr35 ret i32 %res36}37 38define i32 @test_addr_mode_i16(ptr %x) {39; CHECK-LABEL: test_addr_mode_i16(40; CHECK: {41; CHECK-NEXT: .reg .b32 %r<2>;42; CHECK-NEXT: .reg .b64 %rd<2>;43; CHECK-EMPTY:44; CHECK-NEXT: // %bb.0:45; CHECK-NEXT: ld.param.b64 %rd1, [test_addr_mode_i16_param_0];46; CHECK-NEXT: ld.b32 %r1, [%rd1+-4];47; CHECK-NEXT: st.param.b32 [func_retval0], %r1;48; CHECK-NEXT: ret;49 %addr = getelementptr i32, ptr %x, i16 -150 %res = load i32, ptr %addr51 ret i32 %res52}53 54define i32 @test_addr_mode_i8(ptr %x) {55; CHECK-LABEL: test_addr_mode_i8(56; CHECK: {57; CHECK-NEXT: .reg .b32 %r<2>;58; CHECK-NEXT: .reg .b64 %rd<2>;59; CHECK-EMPTY:60; CHECK-NEXT: // %bb.0:61; CHECK-NEXT: ld.param.b64 %rd1, [test_addr_mode_i8_param_0];62; CHECK-NEXT: ld.b32 %r1, [%rd1+-4];63; CHECK-NEXT: st.param.b32 [func_retval0], %r1;64; CHECK-NEXT: ret;65 %addr = getelementptr i32, ptr %x, i8 -166 %res = load i32, ptr %addr67 ret i32 %res68}69 70define i32 @test_addr_mode_i64_large(ptr %x) {71; CHECK-LABEL: test_addr_mode_i64_large(72; CHECK: {73; CHECK-NEXT: .reg .b32 %r<2>;74; CHECK-NEXT: .reg .b64 %rd<3>;75; CHECK-EMPTY:76; CHECK-NEXT: // %bb.0:77; CHECK-NEXT: ld.param.b64 %rd1, [test_addr_mode_i64_large_param_0];78; CHECK-NEXT: add.s64 %rd2, %rd1, 17179869172;79; CHECK-NEXT: ld.b32 %r1, [%rd2];80; CHECK-NEXT: st.param.b32 [func_retval0], %r1;81; CHECK-NEXT: ret;82 %addr = getelementptr i32, ptr %x, i64 429496729383 %res = load i32, ptr %addr84 ret i32 %res85}86