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. 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 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 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 RetTypes; SmallVector, 5> ArgTypes; 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/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/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 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";