f193bcc701
This reverts the following commits: 37ca7a795b277c20c02a218bf44052278c03344b 9aa6c72b92b6c89cc6d23b693257df9af7de2d15 705387c5074bcca36d626882462ebbc2bcc3bed4 8ca4b3ef19fe82d7ad6a6e1515317dcc01b41515 80dba72a669b5416e97a42fd2c2a7bc5a6d3f44a
40 lines
1.7 KiB
Common Lisp
40 lines
1.7 KiB
Common Lisp
// RUN: %clang_cc1 %s -cl-std=cl2.0 -emit-llvm -o - -triple spir-unknown-unknown | FileCheck %s
|
|
// RUN: %clang_cc1 %s -cl-std=cl1.2 -emit-llvm -o - -triple spir-unknown-unknown | FileCheck %s
|
|
// RUN: %clang_cc1 %s -cl-std=cl1.1 -emit-llvm -o - -triple spir-unknown-unknown | FileCheck %s
|
|
|
|
#pragma OPENCL EXTENSION cl_khr_fp64:enable
|
|
|
|
// CHECK-LABEL: @test_store_float(float %foo, half addrspace({{.}}){{.*}} %bar)
|
|
__kernel void test_store_float(float foo, __global half* bar)
|
|
{
|
|
__builtin_store_halff(foo, bar);
|
|
// CHECK: [[HALF_VAL:%.*]] = fptrunc float %foo to half
|
|
// CHECK: store half [[HALF_VAL]], half addrspace({{.}})* %bar, align 2
|
|
}
|
|
|
|
// CHECK-LABEL: @test_store_double(double %foo, half addrspace({{.}}){{.*}} %bar)
|
|
__kernel void test_store_double(double foo, __global half* bar)
|
|
{
|
|
__builtin_store_half(foo, bar);
|
|
// CHECK: [[HALF_VAL:%.*]] = fptrunc double %foo to half
|
|
// CHECK: store half [[HALF_VAL]], half addrspace({{.}})* %bar, align 2
|
|
}
|
|
|
|
// CHECK-LABEL: @test_load_float(float addrspace({{.}}){{.*}} %foo, half addrspace({{.}}){{.*}} %bar)
|
|
__kernel void test_load_float(__global float* foo, __global half* bar)
|
|
{
|
|
foo[0] = __builtin_load_halff(bar);
|
|
// CHECK: [[HALF_VAL:%.*]] = load half, half addrspace({{.}})* %bar
|
|
// CHECK: [[FULL_VAL:%.*]] = fpext half [[HALF_VAL]] to float
|
|
// CHECK: store float [[FULL_VAL]], float addrspace({{.}})* %foo
|
|
}
|
|
|
|
// CHECK-LABEL: @test_load_double(double addrspace({{.}}){{.*}} %foo, half addrspace({{.}}){{.*}} %bar)
|
|
__kernel void test_load_double(__global double* foo, __global half* bar)
|
|
{
|
|
foo[0] = __builtin_load_half(bar);
|
|
// CHECK: [[HALF_VAL:%.*]] = load half, half addrspace({{.}})* %bar
|
|
// CHECK: [[FULL_VAL:%.*]] = fpext half [[HALF_VAL]] to double
|
|
// CHECK: store double [[FULL_VAL]], double addrspace({{.}})* %foo
|
|
}
|