40 lines · plain
1// RUN: %clang_cc1 %s -cl-std=cl2.0 -emit-llvm -o - -triple spir-unknown-unknown | FileCheck %s2// RUN: %clang_cc1 %s -cl-std=cl1.2 -emit-llvm -o - -triple spir-unknown-unknown | FileCheck %s3// RUN: %clang_cc1 %s -cl-std=cl1.1 -emit-llvm -o - -triple spir-unknown-unknown | FileCheck %s4 5#pragma OPENCL EXTENSION cl_khr_fp64:enable6 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 half12// CHECK: store half [[HALF_VAL]], ptr addrspace({{.}}) %bar, align 213}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 half20// CHECK: store half [[HALF_VAL]], ptr addrspace({{.}}) %bar, align 221}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({{.}}) %bar28// CHECK: [[FULL_VAL:%.*]] = fpext half [[HALF_VAL]] to float29// CHECK: store float [[FULL_VAL]], ptr addrspace({{.}}) %foo30}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({{.}}) %bar37// CHECK: [[FULL_VAL:%.*]] = fpext half [[HALF_VAL]] to double38// CHECK: store double [[FULL_VAL]], ptr addrspace({{.}}) %foo39}40