1; RUN: llc < %s -march=nvptx | FileCheck %s
2
3define 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 addrspace(1)* %in
8  store half %val, half addrspace(1) * %out
9  ret void
10}
11
12define 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 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
22define 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 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
32define 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 addrspace(1)* %in
37  %val32 = fpext half %val16 to float
38  store float %val32, float addrspace(1)* %out
39  ret void
40}
41
42define 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 addrspace(1)* %in
47  %val64 = fpext half %val16 to double
48  store double %val64, double addrspace(1)* %out
49  ret void
50}
51
52define 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 addrspace(1)* %in
57  %val16 = fptrunc float %val32 to half
58  store half %val16, half addrspace(1)* %out
59  ret void
60}
61
62define 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 addrspace(1)* %in
67  %val16 = fptrunc double %val32 to half
68  store half %val16, half addrspace(1)* %out
69  ret void
70}
71