50 lines · plain
1; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 52; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_35 -verify-machineinstrs | FileCheck %s3; RUN: %if ptxas %{ llc < %s -mtriple=nvptx64 -mcpu=sm_35 | %ptxas-verify %}4 5; Check load from constant global variables. These loads should be6; ld.global.nc (aka ldg).7 8@gv_float = external constant float9@gv_float2 = external constant <2 x float>10@gv_float4 = external constant <4 x float>11 12define float @test_gv_float() {13; CHECK-LABEL: test_gv_float(14; CHECK: {15; CHECK-NEXT: .reg .b32 %r<2>;16; CHECK-EMPTY:17; CHECK-NEXT: // %bb.0:18; CHECK-NEXT: ld.global.nc.b32 %r1, [gv_float];19; CHECK-NEXT: st.param.b32 [func_retval0], %r1;20; CHECK-NEXT: ret;21 %v = load float, ptr @gv_float22 ret float %v23}24 25define <2 x float> @test_gv_float2() {26; CHECK-LABEL: test_gv_float2(27; CHECK: {28; CHECK-NEXT: .reg .b32 %r<3>;29; CHECK-EMPTY:30; CHECK-NEXT: // %bb.0:31; CHECK-NEXT: ld.global.nc.v2.b32 {%r1, %r2}, [gv_float2];32; CHECK-NEXT: st.param.v2.b32 [func_retval0], {%r1, %r2};33; CHECK-NEXT: ret;34 %v = load <2 x float>, ptr @gv_float235 ret <2 x float> %v36}37 38define <4 x float> @test_gv_float4() {39; CHECK-LABEL: test_gv_float4(40; CHECK: {41; CHECK-NEXT: .reg .b32 %r<5>;42; CHECK-EMPTY:43; CHECK-NEXT: // %bb.0:44; CHECK-NEXT: ld.global.nc.v4.b32 {%r1, %r2, %r3, %r4}, [gv_float4];45; CHECK-NEXT: st.param.v4.b32 [func_retval0], {%r1, %r2, %r3, %r4};46; CHECK-NEXT: ret;47 %v = load <4 x float>, ptr @gv_float448 ret <4 x float> %v49}50