test/CodeGen/NVPTX/half.ll

   1 ; RUN: llc < %s -march=nvptx | FileCheck %s
   2
   3 define void @test_load_store(half addrspace(1)* %in, half addrspace(1)* %out) {
   4 ; CHECK-LABEL: @test_load_store
   5 ; CHECK: ld.global.u16 [[TMP:%rs[0-9]+]], [{{%r[0-9]+}}]
   6 ; CHECK: st.global.u16 [{{%r[0-9]+}}], [[TMP]]
   7   %val = load half, half addrspace(1)* %in
   8   store half %val, half addrspace(1) * %out
   9   ret void
  10 }
  11
  12 define void @test_bitcast_from_half(half addrspace(1)* %in, i16 addrspace(1)* %out) {
  13 ; CHECK-LABEL: @test_bitcast_from_half
  14 ; CHECK: ld.global.u16 [[TMP:%rs[0-9]+]], [{{%r[0-9]+}}]
  15 ; CHECK: st.global.u16 [{{%r[0-9]+}}], [[TMP]]
  16   %val = load half, half addrspace(1) * %in
  17   %val_int = bitcast half %val to i16
  18   store i16 %val_int, i16 addrspace(1)* %out
  19   ret void
  20 }
  21
  22 define void @test_bitcast_to_half(half addrspace(1)* %out, i16 addrspace(1)* %in) {
  23 ; CHECK-LABEL: @test_bitcast_to_half
  24 ; CHECK: ld.global.u16 [[TMP:%rs[0-9]+]], [{{%r[0-9]+}}]
  25 ; CHECK: st.global.u16 [{{%r[0-9]+}}], [[TMP]]
  26   %val = load i16, i16 addrspace(1)* %in
  27   %val_fp = bitcast i16 %val to half
  28   store half %val_fp, half addrspace(1)* %out
  29   ret void
  30 }
  31
  32 define void @test_extend32(half addrspace(1)* %in, float addrspace(1)* %out) {
  33 ; CHECK-LABEL: @test_extend32
  34 ; CHECK: cvt.f32.f16
  35
  36   %val16 = load half, half addrspace(1)* %in
  37   %val32 = fpext half %val16 to float
  38   store float %val32, float addrspace(1)* %out
  39   ret void
  40 }
  41
  42 define void @test_extend64(half addrspace(1)* %in, double addrspace(1)* %out) {
  43 ; CHECK-LABEL: @test_extend64
  44 ; CHECK: cvt.f64.f16
  45
  46   %val16 = load half, half addrspace(1)* %in
  47   %val64 = fpext half %val16 to double
  48   store double %val64, double addrspace(1)* %out
  49   ret void
  50 }
  51
  52 define void @test_trunc32(float addrspace(1)* %in, half addrspace(1)* %out) {
  53 ; CHECK-LABEL: test_trunc32
  54 ; CHECK: cvt.rn.f16.f32
  55
  56   %val32 = load float, float addrspace(1)* %in
  57   %val16 = fptrunc float %val32 to half
  58   store half %val16, half addrspace(1)* %out
  59   ret void
  60 }
  61
  62 define void @test_trunc64(double addrspace(1)* %in, half addrspace(1)* %out) {
  63 ; CHECK-LABEL: @test_trunc64
  64 ; CHECK: cvt.rn.f16.f64
  65
  66   %val32 = load double, double addrspace(1)* %in
  67   %val16 = fptrunc double %val32 to half
  68   store half %val16, half addrspace(1)* %out
  69   ret void
  70 }