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