1 // RUN
: %clang_cc1 -no-opaque-pointers %s -cl-std
=cl2.0 -emit-llvm -o - -triple spir-unknown-unknown | FileCheck %s
2 // RUN
: %clang_cc1 -no-opaque-pointers %s -cl-std
=cl1.2 -emit-llvm -o - -triple spir-unknown-unknown | FileCheck %s
3 // RUN
: %clang_cc1 -no-opaque-pointers %s -cl-std
=cl1.1 -emit-llvm -o - -triple spir-unknown-unknown | FileCheck %s
5 #pragma OPENCL EXTENSION cl_khr_fp64
:enable
7 // CHECK-LABEL
: @test_store_float
(float noundef %foo
, half addrspace
({{.
}}){{.
*}} %bar
)
8 __kernel void test_store_float
(float foo
, __global half
* bar
)
10 __builtin_store_halff
(foo, bar
);
11 // CHECK
: [[HALF_VAL
:%.
*]] = fptrunc float %foo to half
12 // CHECK
: store half
[[HALF_VAL]], half addrspace({{.}})* %bar, align 2
15 // CHECK-LABEL: @test_store_double(double noundef %foo, half addrspace({{.}}){{.*}} %bar)
16 __kernel void test_store_double(double foo, __global half* bar)
18 __builtin_store_half(foo, bar);
19 // CHECK: [[HALF_VAL:%.*]] = fptrunc double %foo to half
20 // CHECK: store half [[HALF_VAL]], half addrspace
({{.
}})* %bar
, align
2
23 // CHECK-LABEL
: @test_load_float
(float addrspace
({{.
}}){{.
*}} %foo
, half addrspace
({{.
}}){{.
*}} %bar
)
24 __kernel void test_load_float
(__global float
* foo
, __global half
* bar
)
26 foo
[0] = __builtin_load_halff(bar);
27 // CHECK: [[HALF_VAL:%.*]] = load half, half addrspace({{.}})* %bar
28 // CHECK: [[FULL_VAL:%.*]] = fpext half [[HALF_VAL]] to float
29 // CHECK: store float [[FULL_VAL]], float addrspace({{.}})* %foo
32 // CHECK-LABEL: @test_load_double(double addrspace({{.}}){{.*}} %foo, half addrspace({{.}}){{.*}} %bar)
33 __kernel void test_load_double(__global double* foo, __global half* bar)
35 foo[0] = __builtin_load_half
(bar);
36 // CHECK
: [[HALF_VAL
:%.
*]] = load half
, half addrspace
({{.
}})* %bar
37 // CHECK
: [[FULL_VAL
:%.
*]] = fpext half
[[HALF_VAL]] to double
38 // CHECK: store double [[FULL_VAL]], double addrspace({{.}})* %foo