[opaque pointer type] Add textual IR support for explicit type parameter to load...
[oota-llvm.git] / 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 }