1*9880d681SAndroid Build Coastguard Worker; RUN: llc < %s -march=nvptx -mcpu=sm_20 -relocation-model=static | FileCheck %s --check-prefix CHK 2*9880d681SAndroid Build Coastguard Worker 3*9880d681SAndroid Build Coastguard Worker%MyStruct = type { i32, i32, float } 4*9880d681SAndroid Build Coastguard Worker@Gbl = internal addrspace(3) global [1024 x %MyStruct] zeroinitializer 5*9880d681SAndroid Build Coastguard Worker 6*9880d681SAndroid Build Coastguard Worker; CHK-LABEL: foo 7*9880d681SAndroid Build Coastguard Workerdefine void @foo(float %f) { 8*9880d681SAndroid Build Coastguard Workerentry: 9*9880d681SAndroid Build Coastguard Worker ; CHK: ld.shared.f32 %{{[a-zA-Z0-9]+}}, [Gbl+8]; 10*9880d681SAndroid Build Coastguard Worker %0 = load float, float addrspace(3)* getelementptr inbounds ([1024 x %MyStruct], [1024 x %MyStruct] addrspace(3)* @Gbl, i32 0, i32 0, i32 2) 11*9880d681SAndroid Build Coastguard Worker %add = fadd float %0, %f 12*9880d681SAndroid Build Coastguard Worker ; CHK: st.shared.f32 [Gbl+8], %{{[a-zA-Z0-9]+}}; 13*9880d681SAndroid Build Coastguard Worker store float %add, float addrspace(3)* getelementptr inbounds ([1024 x %MyStruct], [1024 x %MyStruct] addrspace(3)* @Gbl, i32 0, i32 0, i32 2) 14*9880d681SAndroid Build Coastguard Worker ret void 15*9880d681SAndroid Build Coastguard Worker} 16