This revision was automatically updated to reflect the committed changes.
Closed by commit rL312742: [OpenCL] Add half load and store builtins (authored
by jvesely).
Changed prior to commit:
https://reviews.llvm.org/D37231?vs=113624&id=114240#toc
Repository:
rL LLVM
https://reviews.llvm.org/D37231
Files:
cfe/trunk/include/clang/Basic/Builtins.def
cfe/trunk/include/clang/Basic/Builtins.h
cfe/trunk/lib/Basic/Builtins.cpp
cfe/trunk/lib/CodeGen/CGBuiltin.cpp
cfe/trunk/test/CodeGenOpenCL/no-half.cl
Index: cfe/trunk/lib/CodeGen/CGBuiltin.cpp
===================================================================
--- cfe/trunk/lib/CodeGen/CGBuiltin.cpp
+++ cfe/trunk/lib/CodeGen/CGBuiltin.cpp
@@ -2768,6 +2768,24 @@
Name),
{NDRange, Block}));
}
+
+ case Builtin::BI__builtin_store_half:
+ case Builtin::BI__builtin_store_halff: {
+ Value *Val = EmitScalarExpr(E->getArg(0));
+ Address Address = EmitPointerWithAlignment(E->getArg(1));
+ Value *HalfVal = Builder.CreateFPTrunc(Val, Builder.getHalfTy());
+ return RValue::get(Builder.CreateStore(HalfVal, Address));
+ }
+ case Builtin::BI__builtin_load_half: {
+ Address Address = EmitPointerWithAlignment(E->getArg(0));
+ Value *HalfVal = Builder.CreateLoad(Address);
+ return RValue::get(Builder.CreateFPExt(HalfVal, Builder.getDoubleTy()));
+ }
+ case Builtin::BI__builtin_load_halff: {
+ Address Address = EmitPointerWithAlignment(E->getArg(0));
+ Value *HalfVal = Builder.CreateLoad(Address);
+ return RValue::get(Builder.CreateFPExt(HalfVal, Builder.getFloatTy()));
+ }
case Builtin::BIprintf:
if (getTarget().getTriple().isNVPTX())
return EmitNVPTXDevicePrintfCallExpr(E, ReturnValue);
Index: cfe/trunk/lib/Basic/Builtins.cpp
===================================================================
--- cfe/trunk/lib/Basic/Builtins.cpp
+++ cfe/trunk/lib/Basic/Builtins.cpp
@@ -69,9 +69,14 @@
bool MSModeUnsupported =
!LangOpts.MicrosoftExt && (BuiltinInfo.Langs & MS_LANG);
bool ObjCUnsupported = !LangOpts.ObjC1 && BuiltinInfo.Langs == OBJC_LANG;
- bool OclCUnsupported = LangOpts.OpenCLVersion != 200 &&
- BuiltinInfo.Langs == OCLC20_LANG;
+ bool OclC1Unsupported = (LangOpts.OpenCLVersion / 100) != 1 &&
+ (BuiltinInfo.Langs & ALL_OCLC_LANGUAGES ) == OCLC1X_LANG;
+ bool OclC2Unsupported = LangOpts.OpenCLVersion != 200 &&
+ (BuiltinInfo.Langs & ALL_OCLC_LANGUAGES) == OCLC20_LANG;
+ bool OclCUnsupported = !LangOpts.OpenCL &&
+ (BuiltinInfo.Langs & ALL_OCLC_LANGUAGES);
return !BuiltinsUnsupported && !MathBuiltinsUnsupported && !OclCUnsupported &&
+ !OclC1Unsupported && !OclC2Unsupported &&
!GnuModeUnsupported && !MSModeUnsupported && !ObjCUnsupported;
}
Index: cfe/trunk/include/clang/Basic/Builtins.h
===================================================================
--- cfe/trunk/include/clang/Basic/Builtins.h
+++ cfe/trunk/include/clang/Basic/Builtins.h
@@ -36,10 +36,12 @@
CXX_LANG = 0x4, // builtin for cplusplus only.
OBJC_LANG = 0x8, // builtin for objective-c and objective-c++
MS_LANG = 0x10, // builtin requires MS mode.
- OCLC20_LANG = 0x20, // builtin for OpenCL C only.
+ OCLC20_LANG = 0x20, // builtin for OpenCL C 2.0 only.
+ OCLC1X_LANG = 0x40, // builtin for OpenCL C 1.x only.
ALL_LANGUAGES = C_LANG | CXX_LANG | OBJC_LANG, // builtin for all languages.
ALL_GNU_LANGUAGES = ALL_LANGUAGES | GNU_LANG, // builtin requires GNU mode.
- ALL_MS_LANGUAGES = ALL_LANGUAGES | MS_LANG // builtin requires MS mode.
+ ALL_MS_LANGUAGES = ALL_LANGUAGES | MS_LANG, // builtin requires MS mode.
+ ALL_OCLC_LANGUAGES = OCLC1X_LANG | OCLC20_LANG // builtin for OCLC languages.
};
namespace Builtin {
Index: cfe/trunk/include/clang/Basic/Builtins.def
===================================================================
--- cfe/trunk/include/clang/Basic/Builtins.def
+++ cfe/trunk/include/clang/Basic/Builtins.def
@@ -1424,6 +1424,12 @@
LANGBUILTIN(to_local, "v*v*", "tn", OCLC20_LANG)
LANGBUILTIN(to_private, "v*v*", "tn", OCLC20_LANG)
+// OpenCL half load/store builtin
+LANGBUILTIN(__builtin_store_half, "vdh*", "n", ALL_OCLC_LANGUAGES)
+LANGBUILTIN(__builtin_store_halff, "vfh*", "n", ALL_OCLC_LANGUAGES)
+LANGBUILTIN(__builtin_load_half, "dhC*", "nc", ALL_OCLC_LANGUAGES)
+LANGBUILTIN(__builtin_load_halff, "fhC*", "nc", ALL_OCLC_LANGUAGES)
+
// Builtins for os_log/os_trace
BUILTIN(__builtin_os_log_format_buffer_size, "zcC*.", "p:0:nut")
BUILTIN(__builtin_os_log_format, "v*v*cC*.", "p:0:nt")
Index: cfe/trunk/test/CodeGenOpenCL/no-half.cl
===================================================================
--- cfe/trunk/test/CodeGenOpenCL/no-half.cl
+++ cfe/trunk/test/CodeGenOpenCL/no-half.cl
@@ -0,0 +1,39 @@
+// 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
+}
_______________________________________________
cfe-commits mailing list
[email protected]
http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
