This revision was automatically updated to reflect the committed changes.
Closed by commit rGb250cca11d59: [OpenCL] Do not guard vload/store_half
builtins (authored by svenvh).
Repository:
rG LLVM Github Monorepo
CHANGES SINCE LAST ACTION
https://reviews.llvm.org/D125401/new/
https://reviews.llvm.org/D125401
Files:
clang/lib/Headers/opencl-c-base.h
clang/lib/Sema/OpenCLBuiltins.td
clang/test/SemaOpenCL/half.cl
Index: clang/test/SemaOpenCL/half.cl
===================================================================
--- clang/test/SemaOpenCL/half.cl
+++ clang/test/SemaOpenCL/half.cl
@@ -1,5 +1,5 @@
// RUN: %clang_cc1 %s -verify -pedantic -fsyntax-only -Wno-unused-value -triple spir-unknown-unknown
-// RUN: %clang_cc1 %s -verify -pedantic -fsyntax-only -Wno-unused-value -triple spir-unknown-unknown -fdeclare-opencl-builtins -finclude-default-header
+// RUN: %clang_cc1 %s -verify -pedantic -fsyntax-only -Wno-unused-value -triple spir-unknown-unknown -fdeclare-opencl-builtins -finclude-default-header -DHAVE_BUILTINS
constant float f = 1.0h; // expected-error{{half precision constant requires cl_khr_fp16}}
@@ -22,6 +22,11 @@
half *allowed2 = &*p;
half *allowed3 = p + 1;
+#ifdef HAVE_BUILTINS
+ (void)ilogb(*p); // expected-error{{loading directly from pointer to type '__private half' requires cl_khr_fp16. Use vector data load builtin functions instead}}
+ vstore_half(42.0f, 0, p);
+#endif
+
return h;
}
@@ -49,6 +54,11 @@
half *allowed2 = &*p;
half *allowed3 = p + 1;
+#ifdef HAVE_BUILTINS
+ (void)ilogb(*p);
+ vstore_half(42.0f, 0, p);
+#endif
+
return h;
}
Index: clang/lib/Sema/OpenCLBuiltins.td
===================================================================
--- clang/lib/Sema/OpenCLBuiltins.td
+++ clang/lib/Sema/OpenCLBuiltins.td
@@ -352,9 +352,22 @@
let Extension = Fp64TypeExt in {
def Double : Type<"double", QualType<"Context.DoubleTy">>;
}
+
+// The half type for builtins that require the cl_khr_fp16 extension.
let Extension = Fp16TypeExt in {
def Half : Type<"half", QualType<"Context.HalfTy">>;
}
+
+// Without the cl_khr_fp16 extension, the half type can only be used to declare
+// a pointer. Define const and non-const pointer types in all address spaces.
+// Use the "__half" alias to allow the TableGen emitter to distinguish the
+// (extensionless) pointee type of these pointer-to-half types from the "half"
+// type defined above that already carries the cl_khr_fp16 extension.
+foreach AS = [PrivateAS, GlobalAS, ConstantAS, LocalAS, GenericAS] in {
+ def "HalfPtr" # AS : PointerType<Type<"__half", QualType<"Context.HalfTy">>, AS>;
+ def "HalfPtrConst" # AS : PointerType<ConstType<Type<"__half", QualType<"Context.HalfTy">>>, AS>;
+}
+
def Size : Type<"size_t", QualType<"Context.getSizeType()">>;
def PtrDiff : Type<"ptrdiff_t", QualType<"Context.getPointerDiffType()">>;
def IntPtr : Type<"intptr_t", QualType<"Context.getIntPtrType()">>;
@@ -877,22 +890,22 @@
multiclass VloadVstoreHalf<list<AddressSpace> addrspaces, bit defStores> {
foreach AS = addrspaces in {
- def : Builtin<"vload_half", [Float, Size, PointerType<ConstType<Half>, AS>], Attr.Pure>;
+ def : Builtin<"vload_half", [Float, Size, !cast<Type>("HalfPtrConst" # AS)], Attr.Pure>;
foreach VSize = [2, 3, 4, 8, 16] in {
foreach name = ["vload_half" # VSize, "vloada_half" # VSize] in {
- def : Builtin<name, [VectorType<Float, VSize>, Size, PointerType<ConstType<Half>, AS>], Attr.Pure>;
+ def : Builtin<name, [VectorType<Float, VSize>, Size, !cast<Type>("HalfPtrConst" # AS)], Attr.Pure>;
}
}
if defStores then {
foreach rnd = ["", "_rte", "_rtz", "_rtp", "_rtn"] in {
foreach name = ["vstore_half" # rnd] in {
- def : Builtin<name, [Void, Float, Size, PointerType<Half, AS>]>;
- def : Builtin<name, [Void, Double, Size, PointerType<Half, AS>]>;
+ def : Builtin<name, [Void, Float, Size, !cast<Type>("HalfPtr" # AS)]>;
+ def : Builtin<name, [Void, Double, Size, !cast<Type>("HalfPtr" # AS)]>;
}
foreach VSize = [2, 3, 4, 8, 16] in {
foreach name = ["vstore_half" # VSize # rnd, "vstorea_half" # VSize # rnd] in {
- def : Builtin<name, [Void, VectorType<Float, VSize>, Size, PointerType<Half, AS>]>;
- def : Builtin<name, [Void, VectorType<Double, VSize>, Size, PointerType<Half, AS>]>;
+ def : Builtin<name, [Void, VectorType<Float, VSize>, Size, !cast<Type>("HalfPtr" # AS)]>;
+ def : Builtin<name, [Void, VectorType<Double, VSize>, Size, !cast<Type>("HalfPtr" # AS)]>;
}
}
}
Index: clang/lib/Headers/opencl-c-base.h
===================================================================
--- clang/lib/Headers/opencl-c-base.h
+++ clang/lib/Headers/opencl-c-base.h
@@ -202,6 +202,9 @@
typedef double double16 __attribute__((ext_vector_type(16)));
#endif
+// An internal alias for half, for use by OpenCLBuiltins.td.
+#define __half half
+
#if defined(__OPENCL_CPP_VERSION__)
#define NULL nullptr
#elif defined(__OPENCL_C_VERSION__)
_______________________________________________
cfe-commits mailing list
[email protected]
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
