llvm-project

1
// RUN: %clang_cc1 %s -cl-std=cl2.0 -emit-llvm -o - -triple spir-unknown-unknown | FileCheck %s
2
// RUN: %clang_cc1 %s -cl-std=cl1.2 -emit-llvm -o - -triple spir-unknown-unknown | FileCheck %s
3
// RUN: %clang_cc1 %s -cl-std=cl1.1 -emit-llvm -o - -triple spir-unknown-unknown | FileCheck %s
4

5
#pragma OPENCL EXTENSION cl_khr_fp64:enable
6

7
// CHECK-LABEL: @test_store_float(float noundef %foo, ptr addrspace({{.}}){{.*}} %bar)
8
__kernel void test_store_float(float foo, __global half* bar)
9
{
10
	__builtin_store_halff(foo, bar);
11
// CHECK: [[HALF_VAL:%.*]] = fptrunc float %foo to half
12
// CHECK: store half [[HALF_VAL]], ptr addrspace({{.}}) %bar, align 2
13
}
14

15
// CHECK-LABEL: @test_store_double(double noundef %foo, ptr addrspace({{.}}){{.*}} %bar)
16
__kernel void test_store_double(double foo, __global half* bar)
17
{
18
	__builtin_store_half(foo, bar);
19
// CHECK: [[HALF_VAL:%.*]] = fptrunc double %foo to half
20
// CHECK: store half [[HALF_VAL]], ptr addrspace({{.}}) %bar, align 2
21
}
22

23
// CHECK-LABEL: @test_load_float(ptr addrspace({{.}}){{.*}} %foo, ptr addrspace({{.}}){{.*}} %bar)
24
__kernel void test_load_float(__global float* foo, __global half* bar)
25
{
26
	foo[0] = __builtin_load_halff(bar);
27
// CHECK: [[HALF_VAL:%.*]] = load half, ptr addrspace({{.}}) %bar
28
// CHECK: [[FULL_VAL:%.*]] = fpext half [[HALF_VAL]] to float
29
// CHECK: store float [[FULL_VAL]], ptr addrspace({{.}}) %foo
30
}
31

32
// CHECK-LABEL: @test_load_double(ptr addrspace({{.}}){{.*}} %foo, ptr addrspace({{.}}){{.*}} %bar)
33
__kernel void test_load_double(__global double* foo, __global half* bar)
34
{
35
	foo[0] = __builtin_load_half(bar);
36
// CHECK: [[HALF_VAL:%.*]] = load half, ptr addrspace({{.}}) %bar
37
// CHECK: [[FULL_VAL:%.*]] = fpext half [[HALF_VAL]] to double
38
// CHECK: store double [[FULL_VAL]], ptr addrspace({{.}}) %foo
39
}
40