clang/test/CodeGenOpenCL/no-half.cl

   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 }