test/CodeGen/NVPTX/bug26185.ll

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