svenvh created this revision.
svenvh added reviewers: Anastasia, airlied, azabaznov.
svenvh added a project: clang.
Herald added subscribers: ldrumm, yaxunl.
svenvh requested review of this revision.
Herald added a subscriber: cfe-commits.
Currently, -fdeclare-opencl-builtins always adds the generic address
space overloads of e.g. the vload builtin functions in OpenCL 3.0
mode, even when the generic address space feature is disabled.
Guard the generic address space overloads by the
`__opencl_c_generic_address_space` feature instead of by OpenCL
version.
Add a new field `RequireDisabledExtension` to the `Builtin` class so
that we can make certain builtins available only when an extension is
disabled. Thus, we can provide generic address space overloads OR
private/global/local address space overloads depending on the generic
address space feature availability.
Repository:
rG LLVM Github Monorepo
https://reviews.llvm.org/D107769
Files:
clang/lib/Sema/OpenCLBuiltins.td
clang/lib/Sema/SemaLookup.cpp
clang/test/CodeGenOpenCL/fdeclare-opencl-builtins.cl
clang/test/SemaOpenCL/fdeclare-opencl-builtins.cl
clang/utils/TableGen/ClangOpenCLBuiltinEmitter.cpp
Index: clang/utils/TableGen/ClangOpenCLBuiltinEmitter.cpp
===================================================================
--- clang/utils/TableGen/ClangOpenCLBuiltinEmitter.cpp
+++ clang/utils/TableGen/ClangOpenCLBuiltinEmitter.cpp
@@ -425,6 +425,8 @@
const bool IsConv : 1;
// OpenCL extension(s) required for this overload.
const unsigned short Extension;
+ // OpenCL extension(s) required to be disabled for this overload.
+ const unsigned short DisabledExtension;
// OpenCL versions in which this overload is available.
const unsigned short Versions;
};
@@ -611,6 +613,8 @@
for (const auto &Overload : SLM.second.Signatures) {
StringRef ExtName = Overload.first->getValueAsDef("Extension")->getName();
+ StringRef DisabledExtName =
+ Overload.first->getValueAsDef("RequireDisabledExtension")->getName();
unsigned int MinVersion =
Overload.first->getValueAsDef("MinVersion")->getValueAsInt("ID");
unsigned int MaxVersion =
@@ -622,6 +626,7 @@
<< (Overload.first->getValueAsBit("IsConst")) << ", "
<< (Overload.first->getValueAsBit("IsConv")) << ", "
<< FunctionExtensionIndex[ExtName] << ", "
+ << FunctionExtensionIndex[DisabledExtName] << ", "
<< EncodeVersions(MinVersion, MaxVersion) << " },\n";
Index++;
}
@@ -648,7 +653,9 @@
Rec->getValueAsDef("MaxVersion")->getValueAsInt("ID") ==
Rec2->getValueAsDef("MaxVersion")->getValueAsInt("ID") &&
Rec->getValueAsDef("Extension")->getName() ==
- Rec2->getValueAsDef("Extension")->getName()) {
+ Rec2->getValueAsDef("Extension")->getName() &&
+ Rec->getValueAsDef("RequireDisabledExtension")->getName() ==
+ Rec2->getValueAsDef("RequireDisabledExtension")->getName()) {
return true;
}
}
@@ -1085,11 +1092,27 @@
OpenCLBuiltinFileEmitterBase::emitExtensionGuard(const Record *Builtin) {
StringRef Extensions =
Builtin->getValueAsDef("Extension")->getValueAsString("ExtName");
- if (Extensions.empty())
- return "";
+ StringRef DisabledExtensions =
+ Builtin->getValueAsDef("RequireDisabledExtension")
+ ->getValueAsString("ExtName");
+
+ assert((Extensions.empty() || DisabledExtensions.empty()) &&
+ "enabling and disabling extensions simultaneously not supported yet!");
+
+ bool RequireDisabled = false;
+ if (Extensions.empty()) {
+ if (DisabledExtensions.empty())
+ return "";
+
+ Extensions = DisabledExtensions;
+ RequireDisabled = true;
+ }
OS << "#if";
+ // At this point, Extensions contains a space-separated list of either
+ // the required extensions or the required-to-be-disabled extensions.
+ // RequireDisabled is true if those extensions need to be disabled.
SmallVector<StringRef, 2> ExtVec;
Extensions.split(ExtVec, " ");
bool isFirst = true;
@@ -1097,7 +1120,11 @@
if (!isFirst) {
OS << " &&";
}
- OS << " defined(" << Ext << ")";
+ OS << " ";
+ if (RequireDisabled) {
+ OS << "!";
+ }
+ OS << "defined(" << Ext << ")";
isFirst = false;
}
OS << "\n";
Index: clang/test/SemaOpenCL/fdeclare-opencl-builtins.cl
===================================================================
--- clang/test/SemaOpenCL/fdeclare-opencl-builtins.cl
+++ clang/test/SemaOpenCL/fdeclare-opencl-builtins.cl
@@ -63,6 +63,7 @@
// Enable extensions that are enabled in opencl-c-base.h.
#if (defined(__OPENCL_CPP_VERSION__) || __OPENCL_C_VERSION__ >= 200)
+#define __opencl_c_generic_address_space 1
#define cl_khr_subgroup_extended_types 1
#define cl_khr_subgroup_ballot 1
#define cl_khr_subgroup_non_uniform_arithmetic 1
Index: clang/test/CodeGenOpenCL/fdeclare-opencl-builtins.cl
===================================================================
--- clang/test/CodeGenOpenCL/fdeclare-opencl-builtins.cl
+++ clang/test/CodeGenOpenCL/fdeclare-opencl-builtins.cl
@@ -1,5 +1,11 @@
-// RUN: %clang_cc1 -emit-llvm -o - -O0 -triple spir-unknown-unknown -cl-std=CL1.2 -finclude-default-header %s | FileCheck %s
-// RUN: %clang_cc1 -emit-llvm -o - -O0 -triple spir-unknown-unknown -cl-std=CL1.2 -fdeclare-opencl-builtins -finclude-default-header %s | FileCheck %s
+// RUN: %clang_cc1 -emit-llvm -o - -O0 -triple spir-unknown-unknown -cl-std=CL1.2 -finclude-default-header %s \
+// RUN: | FileCheck %s --check-prefixes CHECK,CHECK-NOGAS
+// RUN: %clang_cc1 -emit-llvm -o - -O0 -triple spir-unknown-unknown -cl-std=CL1.2 -fdeclare-opencl-builtins -finclude-default-header %s \
+// RUN: | FileCheck %s --check-prefixes CHECK,CHECK-NOGAS
+// RUN: %clang_cc1 -emit-llvm -o - -O0 -triple spir-unknown-unknown -cl-std=CL3.0 -fdeclare-opencl-builtins -finclude-default-header %s \
+// RUN: | FileCheck %s --check-prefixes CHECK,CHECK-GAS
+// RUN: %clang_cc1 -emit-llvm -o - -O0 -triple spir-unknown-unknown -cl-std=CL3.0 -fdeclare-opencl-builtins -finclude-default-header -cl-ext=-__opencl_c_generic_address_space,-__opencl_c_pipes %s \
+// RUN: | FileCheck %s --check-prefixes CHECK,CHECK-NOGAS
// Test that mix is correctly defined.
// CHECK-LABEL: @test_float
@@ -32,6 +38,15 @@
size_t lid = get_local_id(0);
}
+// Test that the correct builtin is called depending on the generic address
+// space feature availability.
+// CHECK-LABEL: @test_generic_optionality
+// CHECK-GAS: call spir_func float @_Z5fractfPU3AS4f
+// CHECK-NOGAS: call spir_func float @_Z5fractfPf
+void test_generic_optionality(float a, float *b) {
+ float res = fract(a, b);
+}
+
// CHECK: attributes [[ATTR_CONST]] =
// CHECK-SAME: readnone
// CHECK: attributes [[ATTR_PURE]] =
Index: clang/lib/Sema/SemaLookup.cpp
===================================================================
--- clang/lib/Sema/SemaLookup.cpp
+++ clang/lib/Sema/SemaLookup.cpp
@@ -809,6 +809,19 @@
ASTContext &Context = S.Context;
+ auto AreAllExtensionsDefined = [&S](StringRef Extensions) {
+ if (!Extensions.empty()) {
+ SmallVector<StringRef, 2> ExtVec;
+ Extensions.split(ExtVec, " ");
+ for (StringRef Ext : ExtVec) {
+ if (!S.getPreprocessor().isMacroDefined(Ext)) {
+ return false;
+ }
+ }
+ }
+ return true;
+ };
+
for (unsigned SignatureIndex = 0; SignatureIndex < Len; SignatureIndex++) {
const OpenCLBuiltinStruct &OpenCLBuiltin =
BuiltinTable[FctIndex + SignatureIndex];
@@ -823,19 +836,16 @@
// not defined. This indicates that the extension is not supported by the
// target, so the builtin function should not be available.
StringRef Extensions = FunctionExtensionTable[OpenCLBuiltin.Extension];
- if (!Extensions.empty()) {
- SmallVector<StringRef, 2> ExtVec;
- Extensions.split(ExtVec, " ");
- bool AllExtensionsDefined = true;
- for (StringRef Ext : ExtVec) {
- if (!S.getPreprocessor().isMacroDefined(Ext)) {
- AllExtensionsDefined = false;
- break;
- }
- }
- if (!AllExtensionsDefined)
- continue;
- }
+ if (!AreAllExtensionsDefined(Extensions))
+ continue;
+
+ // Ignore this builtin function if it carries extension macros that all
+ // have to be undefined, but all of them are actually defined.
+ StringRef DisabledExtensions =
+ FunctionExtensionTable[OpenCLBuiltin.DisabledExtension];
+ if (!DisabledExtensions.empty() &&
+ AreAllExtensionsDefined(DisabledExtensions))
+ continue;
SmallVector<QualType, 1> RetTypes;
SmallVector<SmallVector<QualType, 1>, 5> ArgTypes;
Index: clang/lib/Sema/OpenCLBuiltins.td
===================================================================
--- clang/lib/Sema/OpenCLBuiltins.td
+++ clang/lib/Sema/OpenCLBuiltins.td
@@ -83,6 +83,8 @@
def FuncExtKhrMipmapImageWrites : FunctionExtension<"cl_khr_mipmap_image_writes">;
def FuncExtKhrGlMsaaSharing : FunctionExtension<"cl_khr_gl_msaa_sharing">;
+def FuncExtOpenCLCGenericAddressSpace : FunctionExtension<"__opencl_c_generic_address_space">;
+
// Not a real extension, but a workaround to add C++ for OpenCL specific builtins.
def FuncExtOpenCLCxx : FunctionExtension<"__cplusplus">;
@@ -274,8 +276,10 @@
bit IsConst = _Attributes[1];
// Function attribute __attribute__((convergent))
bit IsConv = _Attributes[2];
- // OpenCL extensions to which the function belongs.
+ // OpenCL extensions that all need to be enabled for this builtin.
FunctionExtension Extension = FuncExtNone;
+ // OpenCL extensions that all need to be disabled for this builtin.
+ FunctionExtension RequireDisabledExtension = FuncExtNone;
// Version of OpenCL from which the function is available (e.g.: CL10).
// MinVersion is inclusive.
Version MinVersion = CL10;
@@ -563,10 +567,10 @@
}
}
-let MaxVersion = CL20 in {
+let RequireDisabledExtension = FuncExtOpenCLCGenericAddressSpace in {
defm : MathWithPointer<[GlobalAS, LocalAS, PrivateAS]>;
}
-let MinVersion = CL20 in {
+let Extension = FuncExtOpenCLCGenericAddressSpace in {
defm : MathWithPointer<[GenericAS]>;
}
@@ -821,10 +825,10 @@
}
}
-let MaxVersion = CL20 in {
+let RequireDisabledExtension = FuncExtOpenCLCGenericAddressSpace in {
defm : VloadVstore<[GlobalAS, LocalAS, PrivateAS], 1>;
}
-let MinVersion = CL20 in {
+let Extension = FuncExtOpenCLCGenericAddressSpace in {
defm : VloadVstore<[GenericAS], 1>;
}
// vload with constant address space is available regardless of version.
@@ -856,10 +860,10 @@
}
}
-let MaxVersion = CL20 in {
+let RequireDisabledExtension = FuncExtOpenCLCGenericAddressSpace in {
defm : VloadVstoreHalf<[GlobalAS, LocalAS, PrivateAS], 1>;
}
-let MinVersion = CL20 in {
+let Extension = FuncExtOpenCLCGenericAddressSpace in {
defm : VloadVstoreHalf<[GenericAS], 1>;
}
// vload with constant address space is available regardless of version.
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
