| 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 | |
|---|