diff --git a/clang/include/clang/AST/Type.h b/clang/include/clang/AST/Type.h index 95ecc5f0a38da..b4fbfc080034d 100644 --- a/clang/include/clang/AST/Type.h +++ b/clang/include/clang/AST/Type.h @@ -486,15 +486,12 @@ class Qualifiers { /// Returns true if the address space in these qualifiers is equal to or /// a superset of the address space in the argument qualifiers. bool isAddressSpaceSupersetOf(Qualifiers other) const { - - return - isAddressSpaceSupersetOf(getAddressSpace(), other.getAddressSpace()) || - (!hasAddressSpace() && - (other.getAddressSpace() == LangAS::sycl_private || - other.getAddressSpace() == LangAS::sycl_local || - other.getAddressSpace() == LangAS::sycl_global || - other.getAddressSpace() == LangAS::sycl_constant || - other.getAddressSpace() == LangAS::sycl_generic)); + return isAddressSpaceSupersetOf(getAddressSpace(), + other.getAddressSpace()) || + (!hasAddressSpace() && + (other.getAddressSpace() == LangAS::opencl_private || + other.getAddressSpace() == LangAS::opencl_local || + other.getAddressSpace() == LangAS::opencl_global)); } /// Determines if these qualifiers compatibly include another set. diff --git a/clang/include/clang/Basic/AddressSpaces.h b/clang/include/clang/Basic/AddressSpaces.h index 996b6f03aadf5..faf7f303aa2d6 100644 --- a/clang/include/clang/Basic/AddressSpaces.h +++ b/clang/include/clang/Basic/AddressSpaces.h @@ -42,14 +42,6 @@ enum class LangAS : unsigned { cuda_constant, cuda_shared, - sycl_global, - sycl_local, - sycl_constant, - sycl_private, - // Likely never used, but useful in the future to reserve the spot in the - // enum. - sycl_generic, - // Pointer size and extension address spaces. ptr32_sptr, ptr32_uptr, diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td index 86bb123137432..c1a12e9d4fab0 100644 --- a/clang/include/clang/Basic/DiagnosticSemaKinds.td +++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td @@ -10292,8 +10292,6 @@ def err_builtin_launder_invalid_arg : Error< "'__builtin_launder' is not allowed">; // SYCL-specific diagnostics -def err_sycl_attribute_address_space_invalid : Error< - "address space is outside the valid range of values">; def err_sycl_kernel_name_class_not_top_level : Error< "kernel name class and its template argument classes' declarations can only " "nest in a namespace: %0">; diff --git a/clang/include/clang/Sema/ParsedAttr.h b/clang/include/clang/Sema/ParsedAttr.h index fb06e7857994f..d9d8585970d99 100644 --- a/clang/include/clang/Sema/ParsedAttr.h +++ b/clang/include/clang/Sema/ParsedAttr.h @@ -534,24 +534,6 @@ class ParsedAttr final } } - /// If this is an OpenCL addr space attribute returns its SYCL representation - /// in LangAS, otherwise returns default addr space. - LangAS asSYCLLangAS() const { - switch (getKind()) { - case ParsedAttr::AT_OpenCLConstantAddressSpace: - return LangAS::sycl_constant; - case ParsedAttr::AT_OpenCLGlobalAddressSpace: - return LangAS::sycl_global; - case ParsedAttr::AT_OpenCLLocalAddressSpace: - return LangAS::sycl_local; - case ParsedAttr::AT_OpenCLPrivateAddressSpace: - return LangAS::sycl_private; - case ParsedAttr::AT_OpenCLGenericAddressSpace: - default: - return LangAS::Default; - } - } - AttributeCommonInfo::Kind getKind() const { return getParsedKind(); } }; diff --git a/clang/lib/AST/ASTContext.cpp b/clang/lib/AST/ASTContext.cpp index 869deeebb1ec4..645dcc162165e 100644 --- a/clang/lib/AST/ASTContext.cpp +++ b/clang/lib/AST/ASTContext.cpp @@ -836,11 +836,6 @@ static const LangASMap *getAddressSpaceMap(const TargetInfo &T, 5, // cuda_device 6, // cuda_constant 7, // cuda_shared - 1, // sycl_global - 3, // sycl_local - 2, // sycl_constant - 0, // sycl_private - 4, // sycl_generic 8, // ptr32_sptr 9, // ptr32_uptr 10 // ptr64 diff --git a/clang/lib/AST/TypePrinter.cpp b/clang/lib/AST/TypePrinter.cpp index 808afb09aebea..bf9318cf60c09 100644 --- a/clang/lib/AST/TypePrinter.cpp +++ b/clang/lib/AST/TypePrinter.cpp @@ -1792,16 +1792,12 @@ std::string Qualifiers::getAddrSpaceAsString(LangAS AS) { case LangAS::Default: return ""; case LangAS::opencl_global: - case LangAS::sycl_global: return "__global"; case LangAS::opencl_local: - case LangAS::sycl_local: return "__local"; case LangAS::opencl_private: - case LangAS::sycl_private: return "__private"; case LangAS::opencl_constant: - case LangAS::sycl_constant: return "__constant"; case LangAS::opencl_generic: return "__generic"; diff --git a/clang/lib/Basic/Targets/AMDGPU.cpp b/clang/lib/Basic/Targets/AMDGPU.cpp index 34a1a2375d8d8..6f4ad41739f8c 100644 --- a/clang/lib/Basic/Targets/AMDGPU.cpp +++ b/clang/lib/Basic/Targets/AMDGPU.cpp @@ -48,11 +48,6 @@ const LangASMap AMDGPUTargetInfo::AMDGPUDefIsGenMap = { Global, // cuda_device Constant, // cuda_constant Local, // cuda_shared - Global, // sycl_global - Local, // sycl_local - Constant, // sycl_constant - Private, // sycl_private - Generic, // sycl_generic Generic, // ptr32_sptr Generic, // ptr32_uptr Generic // ptr64 @@ -68,11 +63,6 @@ const LangASMap AMDGPUTargetInfo::AMDGPUDefIsPrivMap = { Global, // cuda_device Constant, // cuda_constant Local, // cuda_shared - Global, // sycl_global - Local, // sycl_local - Constant, // sycl_constant - Private, // sycl_private - Generic, // sycl_generic Generic, // ptr32_sptr Generic, // ptr32_uptr Generic // ptr64 diff --git a/clang/lib/Basic/Targets/NVPTX.h b/clang/lib/Basic/Targets/NVPTX.h index aa97741353da9..63780789c474e 100644 --- a/clang/lib/Basic/Targets/NVPTX.h +++ b/clang/lib/Basic/Targets/NVPTX.h @@ -33,12 +33,6 @@ static const unsigned NVPTXAddrSpaceMap[] = { 1, // cuda_device 4, // cuda_constant 3, // cuda_shared - 1, // sycl_global - 3, // sycl_local - 4, // sycl_constant - 0, // sycl_private - // FIXME: generic has to be added to the target - 0, // sycl_generic 0, // ptr32_sptr 0, // ptr32_uptr 0 // ptr64 diff --git a/clang/lib/Basic/Targets/SPIR.h b/clang/lib/Basic/Targets/SPIR.h index b24d0107d51a0..b250a72ad6c76 100644 --- a/clang/lib/Basic/Targets/SPIR.h +++ b/clang/lib/Basic/Targets/SPIR.h @@ -33,11 +33,6 @@ static const unsigned SPIRAddrSpaceMap[] = { 0, // cuda_device 0, // cuda_constant 0, // cuda_shared - 1, // sycl_global - 3, // sycl_local - 2, // sycl_constant - 0, // sycl_private - 4, // sycl_generic 0, // ptr32_sptr 0, // ptr32_uptr 0 // ptr64 @@ -53,11 +48,6 @@ static const unsigned SYCLAddrSpaceMap[] = { 0, // cuda_device 0, // cuda_constant 0, // cuda_shared - 1, // sycl_global - 3, // sycl_local - 2, // sycl_constant - 0, // sycl_private - 4, // sycl_generic 0, // ptr32_sptr 0, // ptr32_uptr 0 // ptr64 @@ -70,11 +60,9 @@ class LLVM_LIBRARY_VISIBILITY SPIRTargetInfo : public TargetInfo { TLSSupported = false; VLASupported = false; LongWidth = LongAlign = 64; - if (Triple.getEnvironment() == llvm::Triple::SYCLDevice) { - AddrSpaceMap = &SYCLAddrSpaceMap; - } else { - AddrSpaceMap = &SPIRAddrSpaceMap; - } + AddrSpaceMap = (Triple.getEnvironment() == llvm::Triple::SYCLDevice) + ? &SYCLAddrSpaceMap + : &SPIRAddrSpaceMap; UseAddrSpaceMapMangling = true; HasLegalHalfType = true; HasFloat16 = true; diff --git a/clang/lib/Basic/Targets/TCE.h b/clang/lib/Basic/Targets/TCE.h index f7e2bb99e9371..9cbf2a3688a2e 100644 --- a/clang/lib/Basic/Targets/TCE.h +++ b/clang/lib/Basic/Targets/TCE.h @@ -40,12 +40,6 @@ static const unsigned TCEOpenCLAddrSpaceMap[] = { 0, // cuda_device 0, // cuda_constant 0, // cuda_shared - 3, // sycl_global - 4, // sycl_local - 5, // sycl_constant - 0, // sycl_private - // FIXME: generic has to be added to the target - 0, // sycl_generic 0, // ptr32_sptr 0, // ptr32_uptr 0, // ptr64 diff --git a/clang/lib/Basic/Targets/X86.h b/clang/lib/Basic/Targets/X86.h index e05de294fb0f6..5b5e284e51419 100644 --- a/clang/lib/Basic/Targets/X86.h +++ b/clang/lib/Basic/Targets/X86.h @@ -32,11 +32,6 @@ static const unsigned X86AddrSpaceMap[] = { 0, // cuda_device 0, // cuda_constant 0, // cuda_shared - 0, // sycl_global - 0, // sycl_local - 0, // sycl_constant - 0, // sycl_private - 0, // sycl_generic 270, // ptr32_sptr 271, // ptr32_uptr 272 // ptr64 diff --git a/clang/lib/Sema/SemaType.cpp b/clang/lib/Sema/SemaType.cpp index 611c649338257..97ee00c93c70c 100644 --- a/clang/lib/Sema/SemaType.cpp +++ b/clang/lib/Sema/SemaType.cpp @@ -5963,35 +5963,14 @@ static bool BuildAddressSpaceIndex(Sema &S, LangAS &ASIdx, llvm::APSInt max(addrSpace.getBitWidth()); max = Qualifiers::MaxAddressSpace - (unsigned)LangAS::FirstTargetAddressSpace; - if (addrSpace > max) { S.Diag(AttrLoc, diag::err_attribute_address_space_too_high) << (unsigned)max.getZExtValue() << AddrSpace->getSourceRange(); return false; } - if (S.LangOpts.SYCLIsDevice && (addrSpace >= 4)) { - S.Diag(AttrLoc, diag::err_sycl_attribute_address_space_invalid) - << AddrSpace->getSourceRange(); - return false; - } - - ASIdx = getLangASFromTargetAS( - static_cast(addrSpace.getZExtValue())); - - if (S.LangOpts.SYCLIsDevice) { - ASIdx = - [](unsigned AS) { - switch (AS) { - case 0: return LangAS::sycl_private; - case 1: return LangAS::sycl_global; - case 2: return LangAS::sycl_constant; - case 3: return LangAS::sycl_local; - case 4: default: llvm_unreachable("Invalid SYCL AS"); - } - }(static_cast(ASIdx) - - static_cast(LangAS::FirstTargetAddressSpace)); - } + ASIdx = + getLangASFromTargetAS(static_cast(addrSpace.getZExtValue())); return true; } @@ -6117,8 +6096,7 @@ static void HandleAddressSpaceTypeAttribute(QualType &Type, Attr.setInvalid(); } else { // The keyword-based type attributes imply which address space to use. - ASIdx = S.getLangOpts().SYCLIsDevice ? - Attr.asSYCLLangAS() : Attr.asOpenCLLangAS(); + ASIdx = Attr.asOpenCLLangAS(); if (ASIdx == LangAS::Default) llvm_unreachable("Invalid address space"); diff --git a/clang/test/CodeGenSYCL/address-space-parameter-conversions.cpp b/clang/test/CodeGenSYCL/address-space-parameter-conversions.cpp index 4057603281239..7ef8f3176dd04 100644 --- a/clang/test/CodeGenSYCL/address-space-parameter-conversions.cpp +++ b/clang/test/CodeGenSYCL/address-space-parameter-conversions.cpp @@ -9,7 +9,7 @@ void foo(int * Data) {} // CHECK-DAG: define spir_func void @[[RAW_PTR:[a-zA-Z0-9_]+]](i32 addrspace(4)* % void foo2(int * Data) {} // CHECK-DAG: define spir_func void @[[RAW_PTR2:[a-zA-Z0-9_]+]](i32 addrspace(4)* % -void foo(__attribute__((address_space(3))) int * Data) {} +void foo(__attribute__((opencl_local)) int * Data) {} // CHECK-DAG: define spir_func void [[LOC_PTR:@[a-zA-Z0-9_]+]](i32 addrspace(3)* % template @@ -18,12 +18,11 @@ void tmpl(T t){} void usages() { // CHECK-DAG: [[GLOB:%[a-zA-Z0-9]+]] = alloca i32 addrspace(1)* - __attribute__((address_space(1))) int *GLOB; + __attribute__((opencl_global)) int *GLOB; // CHECK-DAG: [[LOC:%[a-zA-Z0-9]+]] = alloca i32 addrspace(3)* __attribute__((opencl_local)) int *LOC; // CHECK-DAG: [[NoAS:%[a-zA-Z0-9]+]] = alloca i32 addrspace(4)* int *NoAS; - // CHECK-DAG: [[PRIV:%[a-zA-Z0-9]+]] = alloca i32* __attribute__((opencl_private)) int *PRIV; @@ -94,57 +93,23 @@ void usages() { // CHECK-DAG: define linkonce_odr spir_func void [[GEN_TMPL]](i32 addrspace(4)* % void usages2() { - __attribute__((address_space(0))) int *PRIV_NUM; - // CHECK-DAG: [[PRIV_NUM:%[a-zA-Z0-9_]+]] = alloca i32* - __attribute__((address_space(0))) int *PRIV_NUM2; - // CHECK-DAG: [[PRIV_NUM2:%[a-zA-Z0-9_]+]] = alloca i32* __attribute__((opencl_private)) int *PRIV; // CHECK-DAG: [[PRIV:%[a-zA-Z0-9_]+]] = alloca i32* - __attribute__((address_space(1))) int *GLOB_NUM; - // CHECK-DAG: [[GLOB_NUM:%[a-zA-Z0-9_]+]] = alloca i32 addrspace(1)* __attribute__((opencl_global)) int *GLOB; // CHECK-DAG: [[GLOB:%[a-zA-Z0-9_]+]] = alloca i32 addrspace(1)* - __attribute__((address_space(2))) int *CONST_NUM; - // CHECK-DAG: [[CONST_NUM:%[a-zA-Z0-9_]+]] = alloca i32 addrspace(2)* __attribute__((opencl_constant)) int *CONST; // CHECK-DAG: [[CONST:%[a-zA-Z0-9_]+]] = alloca i32 addrspace(2)* - __attribute__((address_space(3))) int *LOCAL_NUM; - // CHECK-DAG: [[LOCAL_NUM:%[a-zA-Z0-9_]+]] = alloca i32 addrspace(3)* __attribute__((opencl_local)) int *LOCAL; // CHECK-DAG: [[LOCAL:%[a-zA-Z0-9_]+]] = alloca i32 addrspace(3)* - bar(*PRIV_NUM); - // CHECK-DAG: [[PRIV_NUM_LOAD:%[a-zA-Z0-9]+]] = load i32*, i32** [[PRIV_NUM]] - // CHECK-DAG: [[PRIV_NUM_ASCAST:%[a-zA-Z0-9]+]] = addrspacecast i32* [[PRIV_NUM_LOAD]] to i32 addrspace(4)* - // CHECK-DAG: call spir_func void @[[RAW_REF]](i32 addrspace(4)* dereferenceable(4) [[PRIV_NUM_ASCAST]]) - bar(*PRIV_NUM2); - // CHECK-DAG: [[PRIV_NUM2_LOAD:%[a-zA-Z0-9]+]] = load i32*, i32** [[PRIV_NUM2]] - // CHECK-DAG: [[PRIV_NUM2_ASCAST:%[a-zA-Z0-9]+]] = addrspacecast i32* [[PRIV_NUM2_LOAD]] to i32 addrspace(4)* - // CHECK-DAG: call spir_func void @[[RAW_REF]](i32 addrspace(4)* dereferenceable(4) [[PRIV_NUM2_ASCAST]]) bar(*PRIV); // CHECK-DAG: [[PRIV_LOAD:%[a-zA-Z0-9]+]] = load i32*, i32** [[PRIV]] // CHECK-DAG: [[PRIV_ASCAST:%[a-zA-Z0-9]+]] = addrspacecast i32* [[PRIV_LOAD]] to i32 addrspace(4)* // CHECK-DAG: call spir_func void @[[RAW_REF]](i32 addrspace(4)* dereferenceable(4) [[PRIV_ASCAST]]) - bar(*GLOB_NUM); - // CHECK-DAG: [[GLOB_NUM_LOAD:%[a-zA-Z0-9]+]] = load i32 addrspace(1)*, i32 addrspace(1)** [[GLOB_NUM]] - // CHECK-DAG: [[GLOB_NUM_CAST:%[a-zA-Z0-9]+]] = addrspacecast i32 addrspace(1)* [[GLOB_NUM_LOAD]] to i32 addrspace(4)* - // CHECK-DAG: call spir_func void @[[RAW_REF]](i32 addrspace(4)* dereferenceable(4) [[GLOB_NUM_CAST]]) bar(*GLOB); // CHECK-DAG: [[GLOB_LOAD:%[a-zA-Z0-9]+]] = load i32 addrspace(1)*, i32 addrspace(1)** [[GLOB]] // CHECK-DAG: [[GLOB_CAST:%[a-zA-Z0-9]+]] = addrspacecast i32 addrspace(1)* [[GLOB_LOAD]] to i32 addrspace(4)* // CHECK-DAG: call spir_func void @[[RAW_REF]](i32 addrspace(4)* dereferenceable(4) [[GLOB_CAST]]) - bar(*CONST_NUM); - // CHECK-DAG: [[CONST_NUM_LOAD:%[a-zA-Z0-9]+]] = load i32 addrspace(2)*, i32 addrspace(2)** [[CONST_NUM]] - // CHECK-DAG: [[CONST_NUM_CAST:%[a-zA-Z0-9]+]] = addrspacecast i32 addrspace(2)* [[CONST_NUM_LOAD]] to i32 addrspace(4)* - // CHECK-DAG: call spir_func void @[[RAW_REF]](i32 addrspace(4)* dereferenceable(4) [[CONST_NUM_CAST]]) - bar(*CONST); - // CHECK-DAG: [[CONST_LOAD:%[a-zA-Z0-9]+]] = load i32 addrspace(2)*, i32 addrspace(2)** [[CONST]] - // CHECK-DAG: [[CONST_CAST:%[a-zA-Z0-9]+]] = addrspacecast i32 addrspace(2)* [[CONST_LOAD]] to i32 addrspace(4)* - // CHECK-DAG: call spir_func void @[[RAW_REF]](i32 addrspace(4)* dereferenceable(4) [[CONST_CAST]]) - bar2(*LOCAL_NUM); - // CHECK-DAG: [[LOCAL_NUM_LOAD:%[a-zA-Z0-9]+]] = load i32 addrspace(3)*, i32 addrspace(3)** [[LOCAL_NUM]] - // CHECK-DAG: [[LOCAL_NUM_CAST:%[a-zA-Z0-9]+]] = addrspacecast i32 addrspace(3)* [[LOCAL_NUM_LOAD]] to i32 addrspace(4)* - // CHECK-DAG: call spir_func void @[[RAW_REF2]](i32 addrspace(4)* dereferenceable(4) [[LOCAL_NUM_CAST]]) bar2(*LOCAL); // CHECK-DAG: [[LOCAL_LOAD:%[a-zA-Z0-9]+]] = load i32 addrspace(3)*, i32 addrspace(3)** [[LOCAL]] // CHECK-DAG: [[LOCAL_CAST:%[a-zA-Z0-9]+]] = addrspacecast i32 addrspace(3)* [[LOCAL_LOAD]] to i32 addrspace(4)* diff --git a/clang/test/SemaOpenCLCXX/address-space-lambda.cl b/clang/test/SemaOpenCLCXX/address-space-lambda.cl index eeea71e6353f6..f94717965016e 100644 --- a/clang/test/SemaOpenCLCXX/address-space-lambda.cl +++ b/clang/test/SemaOpenCLCXX/address-space-lambda.cl @@ -31,8 +31,8 @@ __kernel void test_qual() { //CHECK: |-CXXMethodDecl {{.*}} constexpr operator() 'void () const __generic' auto priv2 = []() __generic {}; priv2(); - auto priv3 = []() __global {}; //expected-note{{candidate function not viable: 'this' object is in address space '__private', but method expects object in address space '__global'}} //expected-note{{conversion candidate of type 'void (*)()'}} - priv3(); //expected-error{{no matching function for call to object of type}} + auto priv3 = []() __global {}; //ex pected-note{{candidate function not viable: 'this' object is in address space '__private', but method expects object in address space '__global'}} //ex pected-note{{conversion candidate of type 'void (*)()'}} + priv3(); //ex pected-error{{no matching function for call to object of type}} __constant auto const1 = []() __private{}; //expected-note{{candidate function not viable: 'this' object is in address space '__constant', but method expects object in address space '__private'}} //expected-note{{conversion candidate of type 'void (*)()'}} const1(); //expected-error{{no matching function for call to object of type '__constant (lambda at}} diff --git a/clang/test/SemaSYCL/address-space-parameter-conversions.cpp b/clang/test/SemaSYCL/address-space-parameter-conversions.cpp index 4633964ac2235..bd7c6f3c22285 100644 --- a/clang/test/SemaSYCL/address-space-parameter-conversions.cpp +++ b/clang/test/SemaSYCL/address-space-parameter-conversions.cpp @@ -13,7 +13,7 @@ void tmpl(T *t){} void usages() { __attribute__((opencl_global)) int *GLOB; __attribute__((opencl_private)) int *PRIV; - __attribute__((address_space(3))) int *LOC; + __attribute__((opencl_local)) int *LOC; int *NoAS; bar(*GLOB); @@ -53,10 +53,6 @@ void usages() { // expected-error@+1{{address space is negative}} __attribute__((address_space(-1))) int *TooLow; - // expected-error@+1{{address space is outside the valid range of values}} - __attribute__((address_space(6))) int *TooHigh; - // expected-error@+1{{address space is outside the valid range of values}} - __attribute__((address_space(4))) int *TriedGeneric; // expected-error@+1{{unknown type name '__generic'}} __generic int *IsGeneric; diff --git a/clang/test/SemaTemplate/address_space-dependent.cpp b/clang/test/SemaTemplate/address_space-dependent.cpp index 6983a39e5e360..76cd338769eaa 100644 --- a/clang/test/SemaTemplate/address_space-dependent.cpp +++ b/clang/test/SemaTemplate/address_space-dependent.cpp @@ -43,7 +43,7 @@ void neg() { template void tooBig() { - __attribute__((address_space(I))) int *bounds; // expected-error {{address space is larger than the maximum supported (8388590)}} + __attribute__((address_space(I))) int *bounds; // expected-error {{address space is larger than the maximum supported (8388595)}} } template diff --git a/sycl/include/CL/__spirv/spirv_vars.hpp b/sycl/include/CL/__spirv/spirv_vars.hpp index dd0eaa0dce42e..ff5ede3166c92 100644 --- a/sycl/include/CL/__spirv/spirv_vars.hpp +++ b/sycl/include/CL/__spirv/spirv_vars.hpp @@ -10,43 +10,35 @@ #ifdef __SYCL_DEVICE_ONLY__ -typedef size_t size_t_vec __attribute__((ext_vector_type(3))); -extern "C" const __attribute__((opencl_constant)) size_t_vec __spirv_BuiltInGlobalSize; -extern "C" const __attribute__((opencl_constant)) size_t_vec __spirv_BuiltInGlobalInvocationId; -extern "C" const __attribute__((opencl_constant)) size_t_vec __spirv_BuiltInWorkgroupSize; -extern "C" const __attribute__((opencl_constant)) size_t_vec __spirv_BuiltInNumWorkgroups; -extern "C" const __attribute__((opencl_constant)) size_t_vec __spirv_BuiltInLocalInvocationId; -extern "C" const __attribute__((opencl_constant)) size_t_vec __spirv_BuiltInWorkgroupId; -extern "C" const __attribute__((opencl_constant)) size_t_vec __spirv_BuiltInGlobalOffset; - -#define DEFINE_INT_ID_TO_XYZ_CONVERTER(POSTFIX) \ - template static size_t get##POSTFIX(); \ - template <> size_t get##POSTFIX<0>() { return __spirv_BuiltIn##POSTFIX.x; } \ - template <> size_t get##POSTFIX<1>() { return __spirv_BuiltIn##POSTFIX.y; } \ - template <> size_t get##POSTFIX<2>() { return __spirv_BuiltIn##POSTFIX.z; } - -namespace __spirv { +#define __SPIRV_VAR_QUALIFIERS extern "C" const __attribute__((opencl_global)) -DEFINE_INT_ID_TO_XYZ_CONVERTER(GlobalSize); -DEFINE_INT_ID_TO_XYZ_CONVERTER(GlobalInvocationId) -DEFINE_INT_ID_TO_XYZ_CONVERTER(WorkgroupSize) -DEFINE_INT_ID_TO_XYZ_CONVERTER(NumWorkgroups) -DEFINE_INT_ID_TO_XYZ_CONVERTER(LocalInvocationId) -DEFINE_INT_ID_TO_XYZ_CONVERTER(WorkgroupId) -DEFINE_INT_ID_TO_XYZ_CONVERTER(GlobalOffset) +__SPIRV_VAR_QUALIFIERS uint32_t __spirv_BuiltInSubgroupSize; +__SPIRV_VAR_QUALIFIERS uint32_t __spirv_BuiltInSubgroupMaxSize; +__SPIRV_VAR_QUALIFIERS uint32_t __spirv_BuiltInNumSubgroups; +__SPIRV_VAR_QUALIFIERS uint32_t __spirv_BuiltInNumEnqueuedSubgroups; +__SPIRV_VAR_QUALIFIERS uint32_t __spirv_BuiltInSubgroupId; +__SPIRV_VAR_QUALIFIERS uint32_t __spirv_BuiltInSubgroupLocalInvocationId; -} // namespace __spirv +typedef size_t size_t_vec __attribute__((ext_vector_type(3))); +__SPIRV_VAR_QUALIFIERS size_t_vec __spirv_BuiltInGlobalSize; +__SPIRV_VAR_QUALIFIERS size_t_vec __spirv_BuiltInGlobalInvocationId; +__SPIRV_VAR_QUALIFIERS size_t_vec __spirv_BuiltInWorkgroupSize; +__SPIRV_VAR_QUALIFIERS size_t_vec __spirv_BuiltInNumWorkgroups; +__SPIRV_VAR_QUALIFIERS size_t_vec __spirv_BuiltInLocalInvocationId; +__SPIRV_VAR_QUALIFIERS size_t_vec __spirv_BuiltInWorkgroupId; +__SPIRV_VAR_QUALIFIERS size_t_vec __spirv_BuiltInGlobalOffset; -#undef DEFINE_INT_ID_TO_XYZ_CONVERTER +#undef __SPIRV_VAR_QUALIFIERS -extern "C" const __attribute__((opencl_constant)) uint32_t __spirv_BuiltInSubgroupSize; -extern "C" const __attribute__((opencl_constant)) uint32_t __spirv_BuiltInSubgroupMaxSize; -extern "C" const __attribute__((opencl_constant)) uint32_t __spirv_BuiltInNumSubgroups; -extern "C" const __attribute__((opencl_constant)) uint32_t __spirv_BuiltInNumEnqueuedSubgroups; -extern "C" const __attribute__((opencl_constant)) uint32_t __spirv_BuiltInSubgroupId; -extern "C" const __attribute__((opencl_constant)) uint32_t __spirv_BuiltInSubgroupLocalInvocationId; +namespace __spirv { -#define DEFINE_INIT_SIZES(POSTFIX) \ +// Helper function templates to initialize and get vector component from SPIR-V +// built-in variables +#define __SPIRV_DEFINE_INIT_AND_GET_HELPERS(POSTFIX) \ + template static size_t get##POSTFIX(); \ + template <> size_t get##POSTFIX<0>() { return __spirv_BuiltIn##POSTFIX.x; } \ + template <> size_t get##POSTFIX<1>() { return __spirv_BuiltIn##POSTFIX.y; } \ + template <> size_t get##POSTFIX<2>() { return __spirv_BuiltIn##POSTFIX.z; } \ \ template struct InitSizesST##POSTFIX; \ \ @@ -68,18 +60,16 @@ extern "C" const __attribute__((opencl_constant)) uint32_t __spirv_BuiltInSubgro return InitSizesST##POSTFIX::initSize(); \ } -namespace __spirv { +__SPIRV_DEFINE_INIT_AND_GET_HELPERS(GlobalSize); +__SPIRV_DEFINE_INIT_AND_GET_HELPERS(GlobalInvocationId) +__SPIRV_DEFINE_INIT_AND_GET_HELPERS(WorkgroupSize) +__SPIRV_DEFINE_INIT_AND_GET_HELPERS(NumWorkgroups) +__SPIRV_DEFINE_INIT_AND_GET_HELPERS(LocalInvocationId) +__SPIRV_DEFINE_INIT_AND_GET_HELPERS(WorkgroupId) +__SPIRV_DEFINE_INIT_AND_GET_HELPERS(GlobalOffset) -DEFINE_INIT_SIZES(GlobalSize); -DEFINE_INIT_SIZES(GlobalInvocationId) -DEFINE_INIT_SIZES(WorkgroupSize) -DEFINE_INIT_SIZES(NumWorkgroups) -DEFINE_INIT_SIZES(LocalInvocationId) -DEFINE_INIT_SIZES(WorkgroupId) -DEFINE_INIT_SIZES(GlobalOffset) +#undef __SPIRV_DEFINE_INIT_AND_GET_HELPERS } // namespace __spirv -#undef DEFINE_INIT_SIZES - #endif // __SYCL_DEVICE_ONLY__