llvm/test/CodeGen/NVPTX/bug26185.ll

   1 ; RUN: llc < %s -march=nvptx64 -mcpu=sm_35 -verify-machineinstrs | FileCheck %s
   2 ; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_35 | %ptxas-verify %}
   3
   4 ; Verify that we correctly emit code for i8 ldg/ldu. We do not expose 8-bit
   5 ; registers in the backend, so these loads need special handling.
   6
   7 target datalayout = "e-i64:64-v16:16-v32:32-n16:32:64"
   8 target triple = "nvptx64-unknown-unknown"
   9
  10 ; CHECK-LABEL: ex_zext
  11 define void @ex_zext(ptr noalias readonly %data, ptr %res) {
  12 entry:
  13 ; CHECK: ld.global.nc.u8
  14   %val = load i8, ptr %data
  15 ; CHECK: cvt.u32.u8
  16   %valext = zext i8 %val to i32
  17   store i32 %valext, ptr %res
  18   ret void
  19 }
  20
  21 ; CHECK-LABEL: ex_sext
  22 define void @ex_sext(ptr noalias readonly %data, ptr %res) {
  23 entry:
  24 ; CHECK: ld.global.nc.u8
  25   %val = load i8, ptr %data
  26 ; CHECK: cvt.s32.s8
  27   %valext = sext i8 %val to i32
  28   store i32 %valext, ptr %res
  29   ret void
  30 }
  31
  32 ; CHECK-LABEL: ex_zext_v2
  33 define void @ex_zext_v2(ptr noalias readonly %data, ptr %res) {
  34 entry:
  35 ; CHECK: ld.global.nc.v2.u8
  36   %val = load <2 x i8>, ptr %data
  37 ; CHECK: cvt.u32.u16
  38   %valext = zext <2 x i8> %val to <2 x i32>
  39   store <2 x i32> %valext, ptr %res
  40   ret void
  41 }
  42
  43 ; CHECK-LABEL: ex_sext_v2
  44 define void @ex_sext_v2(ptr noalias readonly %data, ptr %res) {
  45 entry:
  46 ; CHECK: ld.global.nc.v2.u8
  47   %val = load <2 x i8>, ptr %data
  48 ; CHECK: cvt.s32.s8
  49   %valext = sext <2 x i8> %val to <2 x i32>
  50   store <2 x i32> %valext, ptr %res
  51   ret void
  52 }
  53
  54 !nvvm.annotations = !{!0,!1,!2,!3}
  55 !0 = !{ptr @ex_zext, !"kernel", i32 1}
  56 !1 = !{ptr @ex_sext, !"kernel", i32 1}
  57 !2 = !{ptr @ex_zext_v2, !"kernel", i32 1}
  58 !3 = !{ptr @ex_sext_v2, !"kernel", i32 1}