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 %foo, half 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]], half addrspace({{.}})* %bar, align 2 13} 14 15// CHECK-LABEL: @test_store_double(double %foo, half 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]], half addrspace({{.}})* %bar, align 2 21} 22 23// CHECK-LABEL: @test_load_float(float addrspace({{.}}){{.*}} %foo, half 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, half addrspace({{.}})* %bar 28// CHECK: [[FULL_VAL:%.*]] = fpext half [[HALF_VAL]] to float 29// CHECK: store float [[FULL_VAL]], float addrspace({{.}})* %foo 30} 31 32// CHECK-LABEL: @test_load_double(double addrspace({{.}}){{.*}} %foo, half 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, half addrspace({{.}})* %bar 37// CHECK: [[FULL_VAL:%.*]] = fpext half [[HALF_VAL]] to double 38// CHECK: store double [[FULL_VAL]], double addrspace({{.}})* %foo 39} 40