1 ; RUN: llc < %s -march=nvptx -mcpu=sm_35 | FileCheck %s 2 3 ; Verify that we correctly emit code for extending ldg/ldu. We do not expose 4 ; extending variants in the backend, but the ldg/ldu selection code may pick 5 ; extending loads as candidates. We do want to support this, so make sure we 6 ; emit the necessary cvt.* instructions to implement the extension and let ptxas 7 ; emit the real extending loads. 8 9 target datalayout = "e-i64:64-v16:16-v32:32-n16:32:64" 10 target triple = "nvptx64-nvidia-cuda" 11 12 ; CHECK-LABEL: spam 13 define ptx_kernel void @spam(i8 addrspace(1)* noalias nocapture readonly %arg, i8 addrspace(1)* noalias nocapture %arg1, i64 %arg2, i64 %arg3) #0 { 14 bb: 15 %tmp = bitcast i8 addrspace(1)* %arg to i16 addrspace(1)* 16 %tmp4 = bitcast i8 addrspace(1)* %arg1 to i64 addrspace(1)* 17 %tmp5 = add nsw i64 %arg3, 8 18 %tmp6 = getelementptr i16, i16 addrspace(1)* %tmp, i64 %tmp5 19 ; CHECK: ld.global.nc.u16 20 %tmp7 = load i16, i16 addrspace(1)* %tmp6, align 2 21 ; CHECK: cvt.s32.s16 22 %tmp8 = sext i16 %tmp7 to i64 23 %tmp9 = mul nsw i64 %tmp8, %tmp8 24 %tmp10 = load i64, i64 addrspace(1)* %tmp4, align 8 25 %tmp11 = add nsw i64 %tmp9, %tmp10 26 store i64 %tmp11, i64 addrspace(1)* %tmp4, align 8 27 ret void 28 } 29 30 attributes #0 = { norecurse nounwind "polly.skip.fn" } 31 32 !nvvm.annotations = !{!0} 33 34 !0 = !{void (i8 addrspace(1)*, i8 addrspace(1)*, i64, i64)* @spam, !"maxntidx", i64 1, !"maxntidy", i64 1, !"maxntidz", i64 1} 35
