From 1c4379c9fca063a8bfd136f8ce5085a2a1e4c062 Mon Sep 17 00:00:00 2001 From: "Chen, Brox" Date: Thu, 29 Sep 2022 11:52:02 -0700 Subject: [PATCH 01/26] added sycl kernel arg annotation support --- .../sycl-kernel-arg-annotations.ll | 38 +++++++++++++ .../sycl-post-link/CompileTimeProperties.def | 15 ++++++ .../CompileTimePropertiesPass.cpp | 54 +++++++++++++++++++ 3 files changed, 107 insertions(+) create mode 100644 llvm/test/tools/sycl-post-link/sycl-kernel-arg-annotations.ll diff --git a/llvm/test/tools/sycl-post-link/sycl-kernel-arg-annotations.ll b/llvm/test/tools/sycl-post-link/sycl-kernel-arg-annotations.ll new file mode 100644 index 0000000000000..98f9eba6d247a --- /dev/null +++ b/llvm/test/tools/sycl-post-link/sycl-kernel-arg-annotations.ll @@ -0,0 +1,38 @@ +; RUN: sycl-post-link --device-globals --ir-output-only -S %s -o %t.ll +; RUN: FileCheck %s -input-file=%t.ll +; +; TODO: Remove --device-globals once other features start using compile-time +; properties. +; +; Tests the translation of "sycl-kernel-arg-attribute" to "spirv.ParameterDecorations" metadata + +target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-n8:16:32:64" +target triple = "spir64_fpga-unknown-unknown" + +$singleArg = comdat any + +; Function Attrs: convergent mustprogress norecurse +define weak_odr dso_local spir_kernel void @singleArg(i32 addrspace(4)* noundef align 4 "sycl-alignment"="4" "sycl-awidth"="32" "sycl-buffer-location"="10" "sycl-conduit" "sycl-dwidth"="64" "sycl-latency"="1" "sycl-maxburst"="3" "sycl-read-write-mode"="2" "sycl-register-map" "sycl-stable" "sycl-strict" "sycl-wait-request"="5" %_arg_p) #0 comdat !kernel_arg_buffer_location !1587 +; CHECK-DAG: !spirv.ParameterDecorations ![[PARMDECOR:[0-9]+]] +{ + +entry: + ret void +} + +!1587 = !{i32 -1} +; CHECK-DAG: ![[PARMDECOR]] = !{![[ARG1:[0-9]+]]} +; CHECK-DAG: ![[ARG1]] = !{![[ALIGN:[0-9]+]], ![[AWIDTH:[0-9]+]], ![[BL:[0-9]+]], ![[CONDUIT:[0-9]+]], ![[DWIDTH:[0-9]+]], ![[LATENCY:[0-9]+]], ![[MAXBURST:[0-9]+]], ![[RWMODE:[0-9]+]], ![[REGMAP:[0-9]+]], ![[STABLE:[0-9]+]], ![[STRICT:[0-9]+]], ![[WAITREQ:[0-9]+]]} + +; CHECK: ![[ALIGN]] = !{i32 6182, i32 4} +; CHECK: ![[AWIDTH]] = !{i32 6177, i32 32} +; CHECK: ![[BL]] = !{i32 5921, i32 10} +; CHECK: ![[CONDUIT]] = !{i32 6175, i32 1} +; CHECK: ![[DWIDTH]] = !{i32 6178, i32 64} +; CHECK: ![[LATENCY]] = !{i32 6179, i32 1} +; CHECK: ![[MAXBURST]] = !{i32 6181, i32 3} +; CHECK: ![[RWMODE]] = !{i32 6180, i32 2} +; CHECK: ![[REGMAP]] = !{i32 6176, i32 1} +; CHECK: ![[STABLE]] = !{i32 6184, i32 1} +; CHECK: ![[STRICT]] = !{i32 19, i32 1} +; CHECK: ![[WAITREQ]] = !{i32 6183, i32 5} diff --git a/llvm/tools/sycl-post-link/CompileTimeProperties.def b/llvm/tools/sycl-post-link/CompileTimeProperties.def index 24628840dc360..139b9b187911f 100644 --- a/llvm/tools/sycl-post-link/CompileTimeProperties.def +++ b/llvm/tools/sycl-post-link/CompileTimeProperties.def @@ -16,3 +16,18 @@ // https://github.com/intel/llvm/blob/sycl/sycl/doc/extensions/DeviceGlobal/SPV_INTEL_global_variable_decorations.asciidoc#decoration SYCL_COMPILE_TIME_PROPERTY("sycl-init-mode", 6148, DecorValueTy::uint32) SYCL_COMPILE_TIME_PROPERTY("sycl-implement-in-csr", 6149, DecorValueTy::boolean) + +// The corresponding SPIR-V OpCodes for kernel pointer/arg properties +// are documented in the +SYCL_COMPILE_TIME_PROPERTY("sycl-buffer-location", 5921, DecorValueTy::uint32) +SYCL_COMPILE_TIME_PROPERTY("sycl-conduit", 6175, DecorValueTy::boolean) +SYCL_COMPILE_TIME_PROPERTY("sycl-register-map", 6176, DecorValueTy::boolean) +SYCL_COMPILE_TIME_PROPERTY("sycl-awidth", 6177, DecorValueTy::uint32) +SYCL_COMPILE_TIME_PROPERTY("sycl-dwidth", 6178, DecorValueTy::uint32) +SYCL_COMPILE_TIME_PROPERTY("sycl-latency", 6179, DecorValueTy::uint32) +SYCL_COMPILE_TIME_PROPERTY("sycl-read-write-mode", 6180, DecorValueTy::uint32) +SYCL_COMPILE_TIME_PROPERTY("sycl-maxburst", 6181, DecorValueTy::uint32) +SYCL_COMPILE_TIME_PROPERTY("sycl-alignment", 6182, DecorValueTy::uint32) +SYCL_COMPILE_TIME_PROPERTY("sycl-wait-request", 6183, DecorValueTy::uint32) +SYCL_COMPILE_TIME_PROPERTY("sycl-stable", 6184, DecorValueTy::boolean) +SYCL_COMPILE_TIME_PROPERTY("sycl-strict", 19, DecorValueTy::boolean) diff --git a/llvm/tools/sycl-post-link/CompileTimePropertiesPass.cpp b/llvm/tools/sycl-post-link/CompileTimePropertiesPass.cpp index 10d73d5ed657c..58c8c337062ee 100644 --- a/llvm/tools/sycl-post-link/CompileTimePropertiesPass.cpp +++ b/llvm/tools/sycl-post-link/CompileTimePropertiesPass.cpp @@ -26,6 +26,7 @@ namespace { constexpr StringRef SYCL_HOST_ACCESS_ATTR = "sycl-host-access"; constexpr StringRef SPIRV_DECOR_MD_KIND = "spirv.Decorations"; +constexpr StringRef SPIRV_PARAM_DECOR_MD_KIND = "spirv.ParameterDecorations"; // The corresponding SPIR-V OpCode for the host_access property is documented // in the SPV_INTEL_global_variable_decorations design document: // https://github.com/intel/llvm/blob/sycl/sycl/doc/extensions/DeviceGlobal/SPV_INTEL_global_variable_decorations.asciidoc#decoration @@ -102,12 +103,43 @@ Optional getGlobalVariableString(const Value *StringV) { return {}; } +/// Tries to generate a SPIR-V decorate metadata node from an attribute +// of kernel arguments. If the attribute is unknown \c nullptr will be returned. +/// +/// @param Ctx [in] the LLVM context. +/// @param Attr [in] the LLVM attribute to generate metadata for. +/// +/// @returns a pointer to a new metadata node if \c Attr is an attribute with a +/// known corresponding SPIR-V decorate and the arguments are valid. +/// Otherwise \c nullptr is returned. +static MDNode *kernelArgAttributeToDecorateMetadata(LLVMContext &Ctx, + const Attribute &Attr) { + // Currently, only string attributes are supported + if (!Attr.isStringAttribute()) + return nullptr; + auto DecorIt = SpirvDecorMap.find(Attr.getKindAsString()); + if (DecorIt == SpirvDecorMap.end()) + return nullptr; + auto Decor = DecorIt->second; + auto DecorCode = Decor.Code; + switch (Decor.Type) { + case DecorValueTy::uint32: + return buildSpirvDecorMetadata(Ctx, DecorCode, + getAttributeAsInteger(Attr)); + case DecorValueTy::boolean: + return buildSpirvDecorMetadata(Ctx, DecorCode, hasProperty(Attr)); + default: + llvm_unreachable("Unhandled decorator type."); + } +} + } // anonymous namespace PreservedAnalyses CompileTimePropertiesPass::run(Module &M, ModuleAnalysisManager &MAM) { LLVMContext &Ctx = M.getContext(); unsigned MDKindID = Ctx.getMDKindID(SPIRV_DECOR_MD_KIND); + unsigned MDParamKindID = Ctx.getMDKindID(SPIRV_PARAM_DECOR_MD_KIND); bool CompileTimePropertiesMet = false; // Let's process all the globals @@ -154,6 +186,28 @@ PreservedAnalyses CompileTimePropertiesPass::run(Module &M, } } + // Process all properties on kernels arugments + for (Function &F : M) { + // Only consider kernels. + if (F.getCallingConv() != CallingConv::SPIR_KERNEL) + continue; + + SmallVector MDOps; + for (unsigned i = 0 ; i < F.arg_size(); i++) { + SmallVector MDArgOps; + for (auto &Attribute : F.getAttributes().getParamAttrs(i)) { + if (MDNode *SPIRVMetadata = kernelArgAttributeToDecorateMetadata(Ctx, Attribute)) + MDArgOps.push_back(SPIRVMetadata); + } + MDOps.push_back(MDNode::get(Ctx, MDArgOps)); + } + // Add the generated metadata to the kernel function. + if (!MDOps.empty()) { + F.addMetadata(MDParamKindID, *MDNode::get(Ctx, MDOps)); + CompileTimePropertiesMet = true; + } + } + // Check pointer annotations. SmallVector RemovableAnnots; for (Function &F : M) From 21d5629d403f54c5741f87c62ff7d2a512e5de89 Mon Sep 17 00:00:00 2001 From: "Chen, Brox" Date: Thu, 29 Sep 2022 12:01:48 -0700 Subject: [PATCH 02/26] run clang-format --- .../sycl-post-link/CompileTimeProperties.def | 25 +++++----- .../CompileTimePropertiesPass.cpp | 46 ++++++++++--------- 2 files changed, 37 insertions(+), 34 deletions(-) diff --git a/llvm/tools/sycl-post-link/CompileTimeProperties.def b/llvm/tools/sycl-post-link/CompileTimeProperties.def index 139b9b187911f..28a827a32689a 100644 --- a/llvm/tools/sycl-post-link/CompileTimeProperties.def +++ b/llvm/tools/sycl-post-link/CompileTimeProperties.def @@ -7,7 +7,8 @@ \*===----------------------------------------------------------------------===*/ #ifndef SYCL_COMPILE_TIME_PROPERTY -#error "SYCL_COMPILE_TIME_PROPERTY(PropertyName, Decoration, ValueType) is not defined." +#error \ + "SYCL_COMPILE_TIME_PROPERTY(PropertyName, Decoration, ValueType) is not defined." #endif // The corresponding SPIR-V OpCodes for the sycl-init-mode and @@ -18,16 +19,16 @@ SYCL_COMPILE_TIME_PROPERTY("sycl-init-mode", 6148, DecorValueTy::uint32) SYCL_COMPILE_TIME_PROPERTY("sycl-implement-in-csr", 6149, DecorValueTy::boolean) // The corresponding SPIR-V OpCodes for kernel pointer/arg properties -// are documented in the +// docs are TBD SYCL_COMPILE_TIME_PROPERTY("sycl-buffer-location", 5921, DecorValueTy::uint32) -SYCL_COMPILE_TIME_PROPERTY("sycl-conduit", 6175, DecorValueTy::boolean) -SYCL_COMPILE_TIME_PROPERTY("sycl-register-map", 6176, DecorValueTy::boolean) -SYCL_COMPILE_TIME_PROPERTY("sycl-awidth", 6177, DecorValueTy::uint32) -SYCL_COMPILE_TIME_PROPERTY("sycl-dwidth", 6178, DecorValueTy::uint32) -SYCL_COMPILE_TIME_PROPERTY("sycl-latency", 6179, DecorValueTy::uint32) +SYCL_COMPILE_TIME_PROPERTY("sycl-conduit", 6175, DecorValueTy::boolean) +SYCL_COMPILE_TIME_PROPERTY("sycl-register-map", 6176, DecorValueTy::boolean) +SYCL_COMPILE_TIME_PROPERTY("sycl-awidth", 6177, DecorValueTy::uint32) +SYCL_COMPILE_TIME_PROPERTY("sycl-dwidth", 6178, DecorValueTy::uint32) +SYCL_COMPILE_TIME_PROPERTY("sycl-latency", 6179, DecorValueTy::uint32) SYCL_COMPILE_TIME_PROPERTY("sycl-read-write-mode", 6180, DecorValueTy::uint32) -SYCL_COMPILE_TIME_PROPERTY("sycl-maxburst", 6181, DecorValueTy::uint32) -SYCL_COMPILE_TIME_PROPERTY("sycl-alignment", 6182, DecorValueTy::uint32) -SYCL_COMPILE_TIME_PROPERTY("sycl-wait-request", 6183, DecorValueTy::uint32) -SYCL_COMPILE_TIME_PROPERTY("sycl-stable", 6184, DecorValueTy::boolean) -SYCL_COMPILE_TIME_PROPERTY("sycl-strict", 19, DecorValueTy::boolean) +SYCL_COMPILE_TIME_PROPERTY("sycl-maxburst", 6181, DecorValueTy::uint32) +SYCL_COMPILE_TIME_PROPERTY("sycl-alignment", 6182, DecorValueTy::uint32) +SYCL_COMPILE_TIME_PROPERTY("sycl-wait-request", 6183, DecorValueTy::uint32) +SYCL_COMPILE_TIME_PROPERTY("sycl-stable", 6184, DecorValueTy::boolean) +SYCL_COMPILE_TIME_PROPERTY("sycl-strict", 19, DecorValueTy::boolean) diff --git a/llvm/tools/sycl-post-link/CompileTimePropertiesPass.cpp b/llvm/tools/sycl-post-link/CompileTimePropertiesPass.cpp index 58c8c337062ee..ae5ce546f3ee5 100644 --- a/llvm/tools/sycl-post-link/CompileTimePropertiesPass.cpp +++ b/llvm/tools/sycl-post-link/CompileTimePropertiesPass.cpp @@ -104,7 +104,8 @@ Optional getGlobalVariableString(const Value *StringV) { } /// Tries to generate a SPIR-V decorate metadata node from an attribute -// of kernel arguments. If the attribute is unknown \c nullptr will be returned. +// of kernel arguments. If the attribute is unknown \c nullptr will be +// returned. /// /// @param Ctx [in] the LLVM context. /// @param Attr [in] the LLVM attribute to generate metadata for. @@ -113,7 +114,7 @@ Optional getGlobalVariableString(const Value *StringV) { /// known corresponding SPIR-V decorate and the arguments are valid. /// Otherwise \c nullptr is returned. static MDNode *kernelArgAttributeToDecorateMetadata(LLVMContext &Ctx, - const Attribute &Attr) { + const Attribute &Attr) { // Currently, only string attributes are supported if (!Attr.isStringAttribute()) return nullptr; @@ -123,13 +124,13 @@ static MDNode *kernelArgAttributeToDecorateMetadata(LLVMContext &Ctx, auto Decor = DecorIt->second; auto DecorCode = Decor.Code; switch (Decor.Type) { - case DecorValueTy::uint32: - return buildSpirvDecorMetadata(Ctx, DecorCode, - getAttributeAsInteger(Attr)); - case DecorValueTy::boolean: - return buildSpirvDecorMetadata(Ctx, DecorCode, hasProperty(Attr)); - default: - llvm_unreachable("Unhandled decorator type."); + case DecorValueTy::uint32: + return buildSpirvDecorMetadata(Ctx, DecorCode, + getAttributeAsInteger(Attr)); + case DecorValueTy::boolean: + return buildSpirvDecorMetadata(Ctx, DecorCode, hasProperty(Attr)); + default: + llvm_unreachable("Unhandled decorator type."); } } @@ -193,19 +194,20 @@ PreservedAnalyses CompileTimePropertiesPass::run(Module &M, continue; SmallVector MDOps; - for (unsigned i = 0 ; i < F.arg_size(); i++) { - SmallVector MDArgOps; - for (auto &Attribute : F.getAttributes().getParamAttrs(i)) { - if (MDNode *SPIRVMetadata = kernelArgAttributeToDecorateMetadata(Ctx, Attribute)) - MDArgOps.push_back(SPIRVMetadata); - } - MDOps.push_back(MDNode::get(Ctx, MDArgOps)); - } - // Add the generated metadata to the kernel function. - if (!MDOps.empty()) { - F.addMetadata(MDParamKindID, *MDNode::get(Ctx, MDOps)); - CompileTimePropertiesMet = true; - } + for (unsigned i = 0; i < F.arg_size(); i++) { + SmallVector MDArgOps; + for (auto &Attribute : F.getAttributes().getParamAttrs(i)) { + if (MDNode *SPIRVMetadata = + kernelArgAttributeToDecorateMetadata(Ctx, Attribute)) + MDArgOps.push_back(SPIRVMetadata); + } + MDOps.push_back(MDNode::get(Ctx, MDArgOps)); + } + // Add the generated metadata to the kernel function. + if (!MDOps.empty()) { + F.addMetadata(MDParamKindID, *MDNode::get(Ctx, MDOps)); + CompileTimePropertiesMet = true; + } } // Check pointer annotations. From 7a50687cbd657de65c09fa0a479275906964fa64 Mon Sep 17 00:00:00 2001 From: "Chen, Brox" Date: Thu, 27 Oct 2022 09:57:47 -0700 Subject: [PATCH 03/26] resolve merge error --- llvm/tools/sycl-post-link/CompileTimePropertiesPass.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/llvm/tools/sycl-post-link/CompileTimePropertiesPass.cpp b/llvm/tools/sycl-post-link/CompileTimePropertiesPass.cpp index cfb2b6f48bdcb..baa044f1e0d8e 100644 --- a/llvm/tools/sycl-post-link/CompileTimePropertiesPass.cpp +++ b/llvm/tools/sycl-post-link/CompileTimePropertiesPass.cpp @@ -212,7 +212,6 @@ attributeToExecModeMetadata(Module &M, const Attribute &Attr) { return None; } ->>>>>>> f74664a57c1caa448b8ed3014dd51c4277a8a14c } // anonymous namespace PreservedAnalyses CompileTimePropertiesPass::run(Module &M, @@ -282,6 +281,7 @@ PreservedAnalyses CompileTimePropertiesPass::run(Module &M, { // Process all properties on kernels. + SmallVector MDOps; SmallVector, 8> NamedMDOps; for (const Attribute &Attribute : F.getAttributes().getFnAttrs()) { if (MDNode *SPIRVMetadata = attributeToDecorateMetadata(Ctx, Attribute)) From 9fe0a3da7792e64e20c536ec5bc1dc53dae99527 Mon Sep 17 00:00:00 2001 From: "Chen, Brox" Date: Thu, 27 Oct 2022 10:55:10 -0700 Subject: [PATCH 04/26] run clang-format --- .../CompileTimePropertiesPass.cpp | 106 +++++++++--------- 1 file changed, 54 insertions(+), 52 deletions(-) diff --git a/llvm/tools/sycl-post-link/CompileTimePropertiesPass.cpp b/llvm/tools/sycl-post-link/CompileTimePropertiesPass.cpp index baa044f1e0d8e..946e3feb62a4e 100644 --- a/llvm/tools/sycl-post-link/CompileTimePropertiesPass.cpp +++ b/llvm/tools/sycl-post-link/CompileTimePropertiesPass.cpp @@ -131,13 +131,13 @@ MDNode *attributeToDecorateMetadata(LLVMContext &Ctx, const Attribute &Attr) { Decor DecorFound = DecorIt->second; uint32_t DecorCode = DecorFound.Code; switch (DecorFound.Type) { - case DecorValueTy::uint32: - return buildSpirvDecorMetadata(Ctx, DecorCode, - getAttributeAsInteger(Attr)); - case DecorValueTy::boolean: - return buildSpirvDecorMetadata(Ctx, DecorCode, hasProperty(Attr)); - default: - llvm_unreachable("Unhandled decorator type."); + case DecorValueTy::uint32: + return buildSpirvDecorMetadata(Ctx, DecorCode, + getAttributeAsInteger(Attr)); + case DecorValueTy::boolean: + return buildSpirvDecorMetadata(Ctx, DecorCode, hasProperty(Attr)); + default: + llvm_unreachable("Unhandled decorator type."); } } @@ -184,8 +184,8 @@ attributeToExecModeMetadata(Module &M, const Attribute &Attr) { // Get the integers from the strings. SmallVector MDVals; for (StringRef ValStr : ValStrs) - MDVals.push_back(ConstantAsMetadata::get(Constant::getIntegerValue( - SizeTTy, APInt(SizeTBitSize, ValStr, 10)))); + MDVals.push_back(ConstantAsMetadata::get( + Constant::getIntegerValue(SizeTTy, APInt(SizeTBitSize, ValStr, 10)))); // The SPIR-V translator expects 3 values, so we pad the remaining // dimensions with 1. @@ -261,49 +261,51 @@ PreservedAnalyses CompileTimePropertiesPass::run(Module &M, if (F.getCallingConv() != CallingConv::SPIR_KERNEL) continue; - { - // Process all properties on kernels arugments - SmallVector MDOps; - for (unsigned i = 0; i < F.arg_size(); i++) { - SmallVector MDArgOps; - for (auto &Attribute : F.getAttributes().getParamAttrs(i)) { - if (MDNode *SPIRVMetadata = attributeToDecorateMetadata(Ctx, Attribute)) - MDArgOps.push_back(SPIRVMetadata); - } - MDOps.push_back(MDNode::get(Ctx, MDArgOps)); - } - // Add the generated metadata to the kernel function. - if (!MDOps.empty()) { - F.addMetadata(MDParamKindID, *MDNode::get(Ctx, MDOps)); - CompileTimePropertiesMet = true; - } - } - - { - // Process all properties on kernels. - SmallVector MDOps; - SmallVector, 8> NamedMDOps; - for (const Attribute &Attribute : F.getAttributes().getFnAttrs()) { - if (MDNode *SPIRVMetadata = attributeToDecorateMetadata(Ctx, Attribute)) - MDOps.push_back(SPIRVMetadata); - else if (auto NamedMetadata = attributeToExecModeMetadata(M, Attribute)) - NamedMDOps.push_back(*NamedMetadata); - } - - // Add the generated metadata to the kernel function. - if (!MDOps.empty()) { - F.addMetadata(MDKindID, *MDNode::get(Ctx, MDOps)); - CompileTimePropertiesMet = true; - } - - // Add the new named metadata to the kernel function. - for (std::pair NamedMD : NamedMDOps) { - // If multiple sources defined this metadata, prioritize the existing one. - if (F.hasMetadata(NamedMD.first)) - continue; - F.addMetadata(NamedMD.first, *NamedMD.second); - } - } + { + // Process all properties on kernels arugments + SmallVector MDOps; + for (unsigned i = 0; i < F.arg_size(); i++) { + SmallVector MDArgOps; + for (auto &Attribute : F.getAttributes().getParamAttrs(i)) { + if (MDNode *SPIRVMetadata = + attributeToDecorateMetadata(Ctx, Attribute)) + MDArgOps.push_back(SPIRVMetadata); + } + MDOps.push_back(MDNode::get(Ctx, MDArgOps)); + } + // Add the generated metadata to the kernel function. + if (!MDOps.empty()) { + F.addMetadata(MDParamKindID, *MDNode::get(Ctx, MDOps)); + CompileTimePropertiesMet = true; + } + } + + { + // Process all properties on kernels. + SmallVector MDOps; + SmallVector, 8> NamedMDOps; + for (const Attribute &Attribute : F.getAttributes().getFnAttrs()) { + if (MDNode *SPIRVMetadata = attributeToDecorateMetadata(Ctx, Attribute)) + MDOps.push_back(SPIRVMetadata); + else if (auto NamedMetadata = attributeToExecModeMetadata(M, Attribute)) + NamedMDOps.push_back(*NamedMetadata); + } + + // Add the generated metadata to the kernel function. + if (!MDOps.empty()) { + F.addMetadata(MDKindID, *MDNode::get(Ctx, MDOps)); + CompileTimePropertiesMet = true; + } + + // Add the new named metadata to the kernel function. + for (std::pair NamedMD : NamedMDOps) { + // If multiple sources defined this metadata, prioritize the existing + // one. + if (F.hasMetadata(NamedMD.first)) + continue; + F.addMetadata(NamedMD.first, *NamedMD.second); + } + } } // Check pointer annotations. From f53534c3d35434e4d5245715ca779ab82295f28b Mon Sep 17 00:00:00 2001 From: "Wang, Di5" Date: Thu, 29 Sep 2022 19:28:40 -0700 Subject: [PATCH 05/26] combine changes on spirv-header & spirv-translator --- llvm-spirv/CMakeLists.txt | 4 ++-- llvm-spirv/lib/SPIRV/SPIRVWriter.cpp | 24 +++++++++++++++++++ llvm-spirv/lib/SPIRV/libSPIRV/SPIRVEnum.h | 20 ++++++++++++++++ .../lib/SPIRV/libSPIRV/SPIRVNameMapEnum.h | 11 +++++++++ 4 files changed, 57 insertions(+), 2 deletions(-) diff --git a/llvm-spirv/CMakeLists.txt b/llvm-spirv/CMakeLists.txt index d2d668f3495ed..34403380634ab 100644 --- a/llvm-spirv/CMakeLists.txt +++ b/llvm-spirv/CMakeLists.txt @@ -38,8 +38,8 @@ if(NOT DEFINED LLVM_TOOL_SPIRV_HEADERS_BUILD AND # Strip the potential trailing newline from tag string(STRIP "${SPIRV_HEADERS_TAG}" SPIRV_HEADERS_TAG) FetchContent_Declare(spirv-headers - GIT_REPOSITORY https://github.com/KhronosGroup/SPIRV-Headers.git - GIT_TAG ${SPIRV_HEADERS_TAG} + GIT_REPOSITORY https://github.com/broxigarchen/SPIRV-Headers.git + GIT_TAG kernelArgProperties SOURCE_DIR ${LLVM_EXTERNAL_SPIRV_HEADERS_SOURCE_DIR} ) FetchContent_MakeAvailable(spirv-headers) diff --git a/llvm-spirv/lib/SPIRV/SPIRVWriter.cpp b/llvm-spirv/lib/SPIRV/SPIRVWriter.cpp index 48310473ea8f4..24ce59be36bc1 100644 --- a/llvm-spirv/lib/SPIRV/SPIRVWriter.cpp +++ b/llvm-spirv/lib/SPIRV/SPIRVWriter.cpp @@ -2349,6 +2349,30 @@ static void transMetadataDecorations(Metadata *MD, SPIRVEntry *Target) { SPIRVWord); TWO_INT_DECORATION_CASE(MathOpDSPModeINTEL, spv::internal, SPIRVWord, SPIRVWord); + case DecorationConduitKernelArgumentINTEL: + case DecorationRegisterMapKernelArgumentINTEL: + case DecorationStableKernelArgumentINTEL: + case DecorationRestrict: { + Target->addDecorate(new SPIRVDecorate(DecoKind, Target)); + break; + } + case DecorationBufferLocationINTEL: + case DecorationMMHostInterfaceReadWriteModeINTEL: + case DecorationMMHostInterfaceAddressWidthINTEL: + case DecorationMMHostInterfaceDataWidthINTEL: + case DecorationMMHostInterfaceLatencyINTEL: + case DecorationMMHostInterfaceMaxBurstINTEL: + case DecorationMMHostInterfaceAlignmentINTEL: + case DecorationMMHostInterfaceWaitRequestINTEL: { + ErrLog.checkError(NumOperands == 2, SPIRVEC_InvalidLlvmModule, + "MMHost Kernel Argument Annotation requires exactly 2 " + "extra operands"); + auto *DecoValEO1 = + mdconst::dyn_extract(DecoMD->getOperand(1)); + Target->addDecorate( + new SPIRVDecorate(DecoKind, Target, DecoValEO1->getZExtValue())); + break; + } case DecorationStallEnableINTEL: { Target->addDecorate(new SPIRVDecorateStallEnableINTEL(Target)); break; diff --git a/llvm-spirv/lib/SPIRV/libSPIRV/SPIRVEnum.h b/llvm-spirv/lib/SPIRV/libSPIRV/SPIRVEnum.h index 46024d2fefbc9..144952b0ff432 100644 --- a/llvm-spirv/lib/SPIRV/libSPIRV/SPIRVEnum.h +++ b/llvm-spirv/lib/SPIRV/libSPIRV/SPIRVEnum.h @@ -456,6 +456,26 @@ template <> inline void SPIRVMap::init() { {internal::CapabilityGlobalVariableDecorationsINTEL}); ADD_VEC_INIT(internal::DecorationArgumentAttributeINTEL, {CapabilityFunctionPointersINTEL}); + ADD_VEC_INIT(DecorationConduitKernelArgumentINTEL, + {CapabilityFPGAKernelArgumentPropertiesINTEL}); + ADD_VEC_INIT(DecorationRegisterMapKernelArgumentINTEL, + {CapabilityFPGAKernelArgumentPropertiesINTEL}); + ADD_VEC_INIT(DecorationMMHostInterfaceAddressWidthINTEL, + {CapabilityFPGAKernelArgumentPropertiesINTEL}); + ADD_VEC_INIT(DecorationMMHostInterfaceDataWidthINTEL, + {CapabilityFPGAKernelArgumentPropertiesINTEL}); + ADD_VEC_INIT(DecorationMMHostInterfaceLatencyINTEL, + {CapabilityFPGAKernelArgumentPropertiesINTEL}); + ADD_VEC_INIT(DecorationMMHostInterfaceReadWriteModeINTEL, + {CapabilityFPGAKernelArgumentPropertiesINTEL}); + ADD_VEC_INIT(DecorationMMHostInterfaceMaxBurstINTEL, + {CapabilityFPGAKernelArgumentPropertiesINTEL}); + ADD_VEC_INIT(DecorationMMHostInterfaceAlignmentINTEL, + {CapabilityFPGAKernelArgumentPropertiesINTEL}); + ADD_VEC_INIT(DecorationMMHostInterfaceWaitRequestINTEL, + {CapabilityFPGAKernelArgumentPropertiesINTEL}); + ADD_VEC_INIT(DecorationStableKernelArgumentINTEL, + {CapabilityFPGAKernelArgumentPropertiesINTEL}); } template <> inline void SPIRVMap::init() { diff --git a/llvm-spirv/lib/SPIRV/libSPIRV/SPIRVNameMapEnum.h b/llvm-spirv/lib/SPIRV/libSPIRV/SPIRVNameMapEnum.h index 4520a5a4602a9..6ff309e99a338 100644 --- a/llvm-spirv/lib/SPIRV/libSPIRV/SPIRVNameMapEnum.h +++ b/llvm-spirv/lib/SPIRV/libSPIRV/SPIRVNameMapEnum.h @@ -179,6 +179,16 @@ template <> inline void SPIRVMap::init() { add(DecorationMediaBlockIOINTEL, "MediaBlockIOINTEL"); add(DecorationAliasScopeINTEL, "AliasScopeINTEL"); add(DecorationNoAliasINTEL, "NoAliasINTEL"); + add(DecorationConduitKernelArgumentINTEL, "ConduitKernelArgumentINTEL"); + add(DecorationRegisterMapKernelArgumentINTEL, "RegisterMapKernelArgumentINTEL"); + add(DecorationMMHostInterfaceAddressWidthINTEL, "MMHostInterfaceAddressWidthINTEL"); + add(DecorationMMHostInterfaceDataWidthINTEL, "MMHostInterfaceDataWidthINTEL"); + add(DecorationMMHostInterfaceLatencyINTEL, "MMHostInterfaceLatencyINTEL"); + add(DecorationMMHostInterfaceReadWriteModeINTEL, "MMHostInterfaceReadWriteModeINTEL"); + add(DecorationMMHostInterfaceMaxBurstINTEL, "MMHostInterfaceMaxBurstINTEL"); + add(DecorationMMHostInterfaceAlignmentINTEL, "MMHostInterfaceAlignmentINTEL"); + add(DecorationMMHostInterfaceWaitRequestINTEL, "MMHostInterfaceWaitRequestINTEL"); + add(DecorationStableKernelArgumentINTEL, "StableKernelArgumentINTEL"); // From spirv_internal.hpp add(internal::DecorationFuncParamKindINTEL, "FuncParamKindINTEL"); @@ -599,6 +609,7 @@ template <> inline void SPIRVMap::init() { add(CapabilityDebugInfoModuleINTEL, "DebugInfoModuleINTEL"); add(CapabilitySplitBarrierINTEL, "SplitBarrierINTEL"); add(CapabilityGroupUniformArithmeticKHR, "GroupUniformArithmeticKHR"); + add(CapabilityFPGAKernelArgumentPropertiesINTEL, "FPGAKernelArgumentPropertiesINTEL"); // From spirv_internal.hpp add(internal::CapabilityFPGADSPControlINTEL, "FPGADSPControlINTEL"); From 27a8714a992ec0f5a031ea19146c5f1886202928 Mon Sep 17 00:00:00 2001 From: "Wang, Di5" Date: Fri, 28 Oct 2022 11:12:50 -0700 Subject: [PATCH 06/26] add annotated_arg --- .../oneapi/annotated_arg/annotated_arg.hpp | 291 ++++++++++++++++ .../ext/oneapi/annotated_arg/properties.hpp | 314 ++++++++++++++++++ .../sycl/ext/oneapi/properties/property.hpp | 12 +- sycl/include/sycl/sycl.hpp | 2 + .../annotated_arg_properties.cpp | 78 +++++ 5 files changed, 696 insertions(+), 1 deletion(-) create mode 100644 sycl/include/sycl/ext/oneapi/annotated_arg/annotated_arg.hpp create mode 100644 sycl/include/sycl/ext/oneapi/annotated_arg/properties.hpp create mode 100644 sycl/test/extensions/annotated_arg/annotated_arg_properties.cpp diff --git a/sycl/include/sycl/ext/oneapi/annotated_arg/annotated_arg.hpp b/sycl/include/sycl/ext/oneapi/annotated_arg/annotated_arg.hpp new file mode 100644 index 0000000000000..6b6ad55141b87 --- /dev/null +++ b/sycl/include/sycl/ext/oneapi/annotated_arg/annotated_arg.hpp @@ -0,0 +1,291 @@ +//==----------- annotated_arg.hpp - SYCL annotated_arg extension -----------==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// + +#pragma once + +#include +#include + +#include +#include +#include +#include + +#ifdef __SYCL_DEVICE_ONLY__ +#define __SYCL_HOST_NOT_SUPPORTED(Op) +#else +#define __SYCL_HOST_NOT_SUPPORTED(Op) \ + throw sycl::exception( \ + sycl::make_error_code(sycl::errc::feature_not_supported), \ + Op " is not supported on host device."); +#endif + +#ifdef __SYCL_DEVICE_ONLY__ +#define __OPENCL_GLOBAL_AS__ __attribute__((opencl_global)) +#ifdef __ENABLE_USM_ADDR_SPACE__ +#define __OPENCL_GLOBAL_DEVICE_AS__ __attribute__((opencl_global_device)) +#define __OPENCL_GLOBAL_HOST_AS__ __attribute__((opencl_global_host)) +#else +#define __OPENCL_GLOBAL_DEVICE_AS__ __attribute__((opencl_global)) +#define __OPENCL_GLOBAL_HOST_AS__ __attribute__((opencl_global)) +#endif // __ENABLE_USM_ADDR_SPACE__ +#define __OPENCL_LOCAL_AS__ __attribute__((opencl_local)) +#define __OPENCL_CONSTANT_AS__ __attribute__((opencl_constant)) +#define __OPENCL_PRIVATE_AS__ __attribute__((opencl_private)) +#else +#define __OPENCL_GLOBAL_AS__ +#define __OPENCL_GLOBAL_DEVICE_AS__ +#define __OPENCL_GLOBAL_HOST_AS__ +#define __OPENCL_LOCAL_AS__ +#define __OPENCL_CONSTANT_AS__ +#define __OPENCL_PRIVATE_AS__ +#endif + +namespace sycl { +__SYCL_INLINE_VER_NAMESPACE(_V1) { +namespace ext { +namespace oneapi { +namespace experimental { + + +template +class annotated_arg { + // This should always fail when instantiating the unspecialized version. + static_assert(is_property_list::value, + "Property list is invalid."); +}; + +// Partial specialization for pointer type +template +class __SYCL_SPECIAL_CLASS annotated_arg, typename std::enable_if::value>::type> { + using property_list_t = detail::properties_t; + using UnderlyingT = typename std::remove_pointer::type; + __OPENCL_GLOBAL_AS__ UnderlyingT *ptr; + + #ifdef __SYCL_DEVICE_ONLY__ + void __init( + [[__sycl_detail__::add_ir_attributes_kernel_parameter( + detail::PropertyMetaInfo::name..., + detail::PropertyMetaInfo::value... + )]] + __OPENCL_GLOBAL_AS__ UnderlyingT* _ptr) { + ptr = _ptr; + } + #endif + +public: + static_assert(std::is_trivially_destructible::value, + "Type T must be trivially destructible."); + static_assert(is_property_list::value, + "Property list is invalid."); + + annotated_arg() = default; + // annotated_arg(const annotated_arg&) = default; + annotated_arg(UnderlyingT *_ptr) : ptr((__OPENCL_GLOBAL_AS__ UnderlyingT*)_ptr) {}; + + operator T&() { + __SYCL_HOST_NOT_SUPPORTED("Implicit conversion of annotated_arg to T") + return ptr; + } + operator const T&() const { + __SYCL_HOST_NOT_SUPPORTED("Implicit conversion of annotated_arg to T") + return ptr; + } + + // template::value>> + // std::remove_pointer_t operator [](std::ptrdiff_t idx) { + // __SYCL_HOST_NOT_SUPPORTED("operator[] on an annotated_arg") + // return ptr[idx]; + // } + + // auto operator [](std::ptrdiff_t idx) { + // __SYCL_HOST_NOT_SUPPORTED("operator[] on an annotated_arg") + // return ptr[idx]; + // } + + // inline T& get() { + // __SYCL_HOST_NOT_SUPPORTED("get()") + // return ptr; + // } + // inline const T& get() const { + // __SYCL_HOST_NOT_SUPPORTED("get()") + // return ptr; + // } + + inline T get() const { + __SYCL_HOST_NOT_SUPPORTED("get()") + return ptr; + } + + template static constexpr bool has_property() { + return property_list_t::template has_property(); + } + + template static constexpr auto get_property() { + return property_list_t::template get_property(); + } +}; + +// Partial specialization for non-pointer type +template +class __SYCL_SPECIAL_CLASS annotated_arg , typename std::enable_if::value>::type> { + using property_list_t = detail::properties_t; + + T obj; + + #ifdef __SYCL_DEVICE_ONLY__ + void __init( + [[__sycl_detail__::add_ir_attributes_kernel_parameter( + detail::PropertyMetaInfo::name..., + detail::PropertyMetaInfo::value... + )]] + T _obj) { + obj = _obj; + } + #endif + +public: + // T should be trivially copy constructible to be device copyable + static_assert(std::is_trivially_copy_constructible::value, + "Type T must be trivially copy constructable."); + static_assert(std::is_trivially_destructible::value, + "Type T must be trivially destructible."); + static_assert(is_property_list::value, + "Property list is invalid."); + + annotated_arg() = default; + annotated_arg(const annotated_arg&) = default; + annotated_arg(const T &_obj) : obj(_obj) {}; + + operator T&() { + __SYCL_HOST_NOT_SUPPORTED("Implicit conversion of annotated_arg to T") + return obj; + } + operator const T&() const { + __SYCL_HOST_NOT_SUPPORTED("Implicit conversion of annotated_arg to T") + return obj; + } + + inline T& get() { + __SYCL_HOST_NOT_SUPPORTED("get()") + return obj; + } + inline const T& get() const { + __SYCL_HOST_NOT_SUPPORTED("get()") + return obj; + } + + template static constexpr bool has_property() { + return property_list_t::template has_property(); + } + + template static constexpr auto get_property() { + return property_list_t::template get_property(); + } +}; + + +/* +template +class annotated_arg { + // This should always fail when instantiating the unspecialized version. + static_assert(is_property_list::value, + "Property list is invalid."); +}; + +// Partial specialization to make PropertyListT visible as a parameter pack +// of properties. +template +class __SYCL_SPECIAL_CLASS annotated_arg> { + using property_list_t = detail::properties_t; + // using CondT = std::conditional::value, __OPENCL_GLOBAL_AS__ UnderlyingT, T>::type; + + #ifdef __SYCL_DEVICE_ONLY__ + void __init( + [[__sycl_detail__::add_ir_attributes_kernel_parameter( + detail::PropertyMetaInfo::name..., + detail::PropertyMetaInfo::value... + )]] + __OPENCL_GLOBAL_AS__ T* _obj) { + obj = _obj; + } + #endif + +public: + // T should be trivially copy constructible to be device copyable + static_assert(std::is_trivially_copy_constructible::value, + "Type T must be trivially copy constructable."); + static_assert(std::is_trivially_destructible::value, + "Type T must be trivially destructible."); + static_assert(is_property_list::value, + "Property list is invalid."); + + // Check compability of each property values in the property list + // static_assert(check_property_list::value, + // "property list contains invalid property."); + + annotated_arg() = default; + annotated_arg(const annotated_arg&) = default; + // annotated_arg(const __OPENCL_GLOBAL_AS__ T &_obj) : obj(_obj) {}; + annotated_arg(T *_obj) : obj((__OPENCL_GLOBAL_AS__ T*)_obj) {}; + + operator T&() { + __SYCL_HOST_NOT_SUPPORTED("Implicit conversion of annotated_arg to T") + return obj; + } + operator const T&() const { + __SYCL_HOST_NOT_SUPPORTED("Implicit conversion of annotated_arg to T") + return obj; + } + + // template::value>> + // std::remove_pointer_t operator [](std::ptrdiff_t idx) { + // __SYCL_HOST_NOT_SUPPORTED("operator[] on an annotated_arg") + // return obj[idx]; + // } + + // auto operator [](std::ptrdiff_t idx) { + // __SYCL_HOST_NOT_SUPPORTED("operator[] on an annotated_arg") + // return obj[idx]; + // } + + // inline T& get() { + // __SYCL_HOST_NOT_SUPPORTED("get()") + // return obj; + // } + // inline const T& get() const { + // __SYCL_HOST_NOT_SUPPORTED("get()") + // return obj; + // } + + inline T* get() const { + __SYCL_HOST_NOT_SUPPORTED("get()") + return obj; + } + // inline const T* get() const { + // __SYCL_HOST_NOT_SUPPORTED("get()") + // return obj; + // } + + template static constexpr bool has_property() { + return property_list_t::template has_property(); + } + + template static constexpr auto get_property() { + return property_list_t::template get_property(); + } +}; +*/ + +} // namespace experimental +} // namespace oneapi +} // namespace ext +} // __SYCL_INLINE_VER_NAMESPACE(_V1) +} // namespace sycl + +#undef __SYCL_HOST_NOT_SUPPORTED diff --git a/sycl/include/sycl/ext/oneapi/annotated_arg/properties.hpp b/sycl/include/sycl/ext/oneapi/annotated_arg/properties.hpp new file mode 100644 index 0000000000000..bdcabd06d5367 --- /dev/null +++ b/sycl/include/sycl/ext/oneapi/annotated_arg/properties.hpp @@ -0,0 +1,314 @@ +//==----- properties.hpp - SYCL properties associated with annotated_arg ---==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// + +#pragma once + +#include +#include + +namespace sycl { +__SYCL_INLINE_VER_NAMESPACE(_V1) { +namespace ext { +namespace oneapi { +namespace experimental { + +template class annotated_arg; + +struct register_map_key { + using value_t = property_value; +}; + +struct conduit_key { + using value_t = property_value; +}; + +struct stable_key { + using value_t = property_value; +}; + +struct buffer_location_key { + template + using value_t = property_value>; +}; + +struct awidth_key { + template + using value_t = property_value>; +}; + +struct dwidth_key { + template + using value_t = property_value>; +}; + +struct latency_key { + template + using value_t = property_value>; +}; + +enum class read_write_mode_enum : std::uint16_t { + read_write, + read, + write +}; + +struct read_write_mode_key { + template + using value_t = property_value< + read_write_mode_key, std::integral_constant>; +}; + +struct maxburst_key { + template + using value_t = property_value>; +}; + +struct wait_request_key { + template + using value_t = property_value>; +}; + + +#if __cplusplus >= 201703L // inline variables +// non-mmhost properties +inline constexpr register_map_key::value_t register_map; +inline constexpr conduit_key::value_t conduit; +inline constexpr stable_key::value_t stable; + +// mmhost properties +template +inline constexpr buffer_location_key::value_t buffer_location; +template +inline constexpr awidth_key::value_t awidth; +template +inline constexpr dwidth_key::value_t dwidth; +template +inline constexpr latency_key::value_t latency; +template +inline constexpr maxburst_key::value_t maxburst; +template +inline constexpr wait_request_key::value_t wait_request; + +template +inline constexpr read_write_mode_key::value_t read_write_mode; +inline constexpr read_write_mode_key::value_t read_only; +inline constexpr read_write_mode_key::value_t write_only; +inline constexpr read_write_mode_key::value_t read_write; + + +#endif // __cplusplus >= 201703L + +template <> struct is_property_key : std::true_type {}; +template <> struct is_property_key : std::true_type {}; +template <> struct is_property_key : std::true_type {}; + +template <> struct is_property_key : std::true_type {}; +template <> struct is_property_key : std::true_type {}; +template <> struct is_property_key : std::true_type {}; +template <> struct is_property_key : std::true_type {}; +template <> struct is_property_key : std::true_type {}; +template <> struct is_property_key : std::true_type {}; +template <> struct is_property_key : std::true_type {}; + + +template +struct is_property_key_of> + : std::true_type {}; + +template +struct is_property_key_of> + : std::true_type {}; + +template +struct is_property_key_of> + : std::true_type {}; + +template +struct is_property_key_of> + : std::true_type {}; + +template +struct is_property_key_of> + : std::true_type {}; + +template +struct is_property_key_of> + : std::true_type {}; + +template +struct is_property_key_of> + : std::true_type {}; + +template +struct is_property_key_of> + : std::true_type {}; + +template +struct is_property_key_of> + : std::true_type {}; + +template +struct is_property_key_of> + : std::true_type {}; + + +namespace detail { +template <> struct PropertyToKind { + static constexpr PropKind Kind = PropKind::RegisterMap; +}; +template <> struct PropertyToKind { + static constexpr PropKind Kind = PropKind::Conduit; +}; +template <> struct PropertyToKind { + static constexpr PropKind Kind = PropKind::Stable; +}; +template <> struct PropertyToKind { + static constexpr PropKind Kind = PropKind::BufferLocation; +}; +template <> struct PropertyToKind { + static constexpr PropKind Kind = PropKind::AddrWidth; +}; +template <> struct PropertyToKind { + static constexpr PropKind Kind = PropKind::DataWidth; +}; +template <> struct PropertyToKind { + static constexpr PropKind Kind = PropKind::Latency; +}; +template <> struct PropertyToKind { + static constexpr PropKind Kind = PropKind::RWMode; +}; +template <> struct PropertyToKind { + static constexpr PropKind Kind = PropKind::MaxBurst; +}; +template <> struct PropertyToKind { + static constexpr PropKind Kind = PropKind::WaitRequest; +}; + +template <> struct IsCompileTimeProperty : std::true_type {}; +template <> struct IsCompileTimeProperty : std::true_type {}; +template <> struct IsCompileTimeProperty : std::true_type {}; + +template <> struct IsCompileTimeProperty : std::true_type {}; +template <> struct IsCompileTimeProperty : std::true_type {}; +template <> struct IsCompileTimeProperty : std::true_type {}; +template <> struct IsCompileTimeProperty : std::true_type {}; +template <> struct IsCompileTimeProperty : std::true_type {}; +template <> struct IsCompileTimeProperty : std::true_type {}; +template <> struct IsCompileTimeProperty : std::true_type {}; + +template <> +struct PropertyMetaInfo { + static constexpr const char *name = "sycl-register-map"; + static constexpr std::nullptr_t value = nullptr; +}; +template <> +struct PropertyMetaInfo { + static constexpr const char *name = "sycl-conduit"; + static constexpr std::nullptr_t value = nullptr; +}; +template <> +struct PropertyMetaInfo { + static constexpr const char *name = "sycl-stable"; + static constexpr std::nullptr_t value = nullptr; +}; + +template +struct PropertyMetaInfo> { + static constexpr const char *name = "sycl-buffer-location"; + static constexpr int value = N; +}; +template +struct PropertyMetaInfo> { + static constexpr const char *name = "sycl-awidth"; + static constexpr int value = W; +}; +template +struct PropertyMetaInfo> { + static constexpr const char *name = "sycl-dwidth"; + static constexpr int value = W; +}; +template +struct PropertyMetaInfo> { + static constexpr const char *name = "sycl-latency"; + static constexpr int value = N; +}; +template +struct PropertyMetaInfo> { + static constexpr const char *name = "sycl-maxburst"; + static constexpr int value = N; +}; +template +struct PropertyMetaInfo> { + static constexpr const char *name = "sycl-wait-request"; + static constexpr int value = Enable; +}; +template +struct PropertyMetaInfo> { + static constexpr const char *name = "sycl-read-write-mode"; + static constexpr read_write_mode_enum value = Mode; +}; + +} // namespace detail + + +//===---------------------Properties Verification----------------------===// + +// 'buffer_location' and mmhost properties are pointers-only +template +struct is_property_compatible : std::false_type {}; + +template +struct is_property_compatible> + : std::bool_constant::value> {}; + +template +struct is_property_compatible> + : std::bool_constant::value> {}; + +template +struct is_property_compatible> + : std::bool_constant::value> {}; + +template +struct is_property_compatible> + : std::bool_constant::value> {}; + +template +struct is_property_compatible> + : std::bool_constant::value> {}; + +template +struct is_property_compatible> + : std::bool_constant::value> {}; + +template +struct is_property_compatible> + : std::bool_constant::value> {}; + +// 'register_map', 'conduit', 'stable' are common properties for pointers +// and non pointers; +template +struct is_property_compatible : std::true_type {}; +template +struct is_property_compatible : std::true_type {}; +template +struct is_property_compatible : std::true_type {}; + +template +struct check_property_list : std::true_type {}; +template +struct check_property_list + : std::conditional_t::value, + check_property_list, + std::false_type> {}; + +} // namespace experimental +} // namespace oneapi +} // namespace ext +} // __SYCL_INLINE_VER_NAMESPACE(_V1) +} // namespace sycl diff --git a/sycl/include/sycl/ext/oneapi/properties/property.hpp b/sycl/include/sycl/ext/oneapi/properties/property.hpp index 08aa8a5f76dd4..840c382f3c15f 100644 --- a/sycl/include/sycl/ext/oneapi/properties/property.hpp +++ b/sycl/include/sycl/ext/oneapi/properties/property.hpp @@ -173,8 +173,18 @@ enum PropKind : uint32_t { WorkGroupSizeHint = 7, SubGroupSize = 8, DeviceHas = 9, + RegisterMap = 10, + Conduit = 11, + Stable = 12, + BufferLocation = 13, + AddrWidth = 14, + DataWidth = 15, + Latency = 16, + RWMode = 17, + MaxBurst = 18, + WaitRequest = 19, // PropKindSize must always be the last value. - PropKindSize = 10, + PropKindSize = 20, }; // This trait must be specialized for all properties and must have a unique diff --git a/sycl/include/sycl/sycl.hpp b/sycl/include/sycl/sycl.hpp index 70a29d96d5097..bb6dd463cddbf 100644 --- a/sycl/include/sycl/sycl.hpp +++ b/sycl/include/sycl/sycl.hpp @@ -61,6 +61,8 @@ #endif #include #include +#include +#include #include #include #include diff --git a/sycl/test/extensions/annotated_arg/annotated_arg_properties.cpp b/sycl/test/extensions/annotated_arg/annotated_arg_properties.cpp new file mode 100644 index 0000000000000..c579f471eab0a --- /dev/null +++ b/sycl/test/extensions/annotated_arg/annotated_arg_properties.cpp @@ -0,0 +1,78 @@ +// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple -fsyntax-only -Xclang -verify -Xclang -verify-ignore-unexpected=note %s +// expected-no-diagnostics + +#include + +using namespace sycl::ext::oneapi::experimental; + +static annotated_arg AnnotatedArg1; +static annotated_arg + AnnotatedArg2; +static annotated_arg + AnnotatedArg3; +static annotated_arg, read_only, + stable, conduit))> + AnnotatedArg4; + +// Checks is_property_key_of and is_property_value_of for T. +template void checkIsPropertyOf() { + static_assert(is_property_key_of::value); + static_assert(is_property_key_of::value); + static_assert(is_property_key_of::value); + + static_assert(is_property_key_of::value); + static_assert(is_property_key_of::value); + static_assert(is_property_key_of::value); + static_assert(is_property_key_of::value); + static_assert(is_property_key_of::value); + static_assert(is_property_key_of::value); + static_assert(is_property_key_of::value); + + static_assert(is_property_value_of::value); + static_assert(is_property_value_of::value); + static_assert(is_property_value_of::value); + + static_assert(is_property_value_of), T>::value); + static_assert(is_property_value_of), T>::value); + static_assert(is_property_value_of), T>::value); + static_assert(is_property_value_of), T>::value); + static_assert(is_property_value_of::value); + static_assert(is_property_value_of), T>::value); + static_assert(is_property_value_of), T>::value); +} + +int main() { + static_assert(is_property_key::value); + static_assert(is_property_key::value); + + checkIsPropertyOf(); + static_assert(!AnnotatedArg1.has_property()); + static_assert(!AnnotatedArg1.has_property()); + + checkIsPropertyOf(); + static_assert(AnnotatedArg2.has_property()); + static_assert(!AnnotatedArg2.has_property()); + static_assert(!AnnotatedArg2.has_property()); + static_assert(AnnotatedArg2.get_property() == register_map); + + checkIsPropertyOf(); + static_assert(AnnotatedArg3.has_property()); + static_assert(AnnotatedArg3.has_property()); + static_assert(!AnnotatedArg3.has_property()); + static_assert(!AnnotatedArg3.has_property()); + static_assert(AnnotatedArg3.get_property() == register_map); + static_assert(AnnotatedArg3.get_property() == conduit); + + checkIsPropertyOf(); + static_assert(!AnnotatedArg4.has_property()); + static_assert(AnnotatedArg4.has_property()); + static_assert(AnnotatedArg4.has_property()); + static_assert(AnnotatedArg4.has_property()); + static_assert(AnnotatedArg4.has_property()); + static_assert(AnnotatedArg4.get_property() == conduit); + static_assert(AnnotatedArg4.get_property() == stable); + static_assert(AnnotatedArg4.get_property() == buffer_location<1>); + static_assert(AnnotatedArg4.get_property() == read_only); + + return 0; +} From a657d7080b3558efd2b769f8f8f3416715030d1f Mon Sep 17 00:00:00 2001 From: "Wang, Di5" Date: Mon, 21 Nov 2022 06:42:01 -0800 Subject: [PATCH 07/26] implement based on annotated_arg spec --- clang/lib/Driver/ToolChains/Clang.cpp | 2 +- .../oneapi/annotated_arg/annotated_arg.hpp | 132 ++++++++++++------ 2 files changed, 94 insertions(+), 40 deletions(-) diff --git a/clang/lib/Driver/ToolChains/Clang.cpp b/clang/lib/Driver/ToolChains/Clang.cpp index 0920e44f41cb2..140eafe72564f 100644 --- a/clang/lib/Driver/ToolChains/Clang.cpp +++ b/clang/lib/Driver/ToolChains/Clang.cpp @@ -9531,13 +9531,13 @@ void SPIRVTranslator::ConstructJob(Compilation &C, const JobAction &JA, ",+SPV_INTEL_variable_length_array,+SPV_INTEL_fp_fast_math_mode" ",+SPV_INTEL_long_constant_composite" ",+SPV_INTEL_arithmetic_fence" + ",+SPV_INTEL_fpga_buffer_location" ",+SPV_INTEL_global_variable_decorations"; ExtArg = ExtArg + DefaultExtArg + INTELExtArg; if (!C.getDriver().isFPGAEmulationMode()) // Enable several extensions on FPGA H/W exclusively ExtArg += ",+SPV_INTEL_usm_storage_classes,+SPV_INTEL_runtime_aligned" ",+SPV_INTEL_fpga_cluster_attributes,+SPV_INTEL_loop_fuse" - ",+SPV_INTEL_fpga_buffer_location" ",+SPV_INTEL_fpga_invocation_pipelining_attributes" ",+SPV_INTEL_fpga_dsp_control,+SPV_INTEL_fpga_memory_accesses" ",+SPV_INTEL_fpga_memory_attributes"; diff --git a/sycl/include/sycl/ext/oneapi/annotated_arg/annotated_arg.hpp b/sycl/include/sycl/ext/oneapi/annotated_arg/annotated_arg.hpp index 6b6ad55141b87..7f8ae15eaf6a4 100644 --- a/sycl/include/sycl/ext/oneapi/annotated_arg/annotated_arg.hpp +++ b/sycl/include/sycl/ext/oneapi/annotated_arg/annotated_arg.hpp @@ -52,6 +52,24 @@ namespace ext { namespace oneapi { namespace experimental { +namespace detail { +// Type-trait for checking if a type defines `operator->`. +template +struct HasParenthesisOperator : std::false_type {}; +template +struct HasParenthesisOperator< + T, sycl::detail::void_t().operator()())>> + : std::true_type {}; + +template +struct HasSubscriptOperator : std::false_type {}; + +template +struct HasSubscriptOperator< + T, sycl::detail::void_t().operator[]())>> + : std::true_type {}; + +} // namespace detail template class annotated_arg { @@ -62,10 +80,10 @@ class annotated_arg { // Partial specialization for pointer type template -class __SYCL_SPECIAL_CLASS annotated_arg, typename std::enable_if::value>::type> { +class __SYCL_SPECIAL_CLASS __SYCL_TYPE(annotated_arg) annotated_arg, typename std::enable_if::value>::type> { using property_list_t = detail::properties_t; using UnderlyingT = typename std::remove_pointer::type; - __OPENCL_GLOBAL_AS__ UnderlyingT *ptr; + __OPENCL_GLOBAL_AS__ UnderlyingT *g_ptr; #ifdef __SYCL_DEVICE_ONLY__ void __init( @@ -73,8 +91,8 @@ class __SYCL_SPECIAL_CLASS annotated_arg, type detail::PropertyMetaInfo::name..., detail::PropertyMetaInfo::value... )]] - __OPENCL_GLOBAL_AS__ UnderlyingT* _ptr) { - ptr = _ptr; + __OPENCL_GLOBAL_AS__ UnderlyingT* _g_ptr) { + g_ptr = _g_ptr; } #endif @@ -84,17 +102,17 @@ class __SYCL_SPECIAL_CLASS annotated_arg, type static_assert(is_property_list::value, "Property list is invalid."); - annotated_arg() = default; - // annotated_arg(const annotated_arg&) = default; - annotated_arg(UnderlyingT *_ptr) : ptr((__OPENCL_GLOBAL_AS__ UnderlyingT*)_ptr) {}; + annotated_arg() noexcept = default; + annotated_arg(T _ptr) + : g_ptr((__OPENCL_GLOBAL_AS__ UnderlyingT*)_ptr) {}; - operator T&() { + operator T&() noexcept { __SYCL_HOST_NOT_SUPPORTED("Implicit conversion of annotated_arg to T") - return ptr; + return g_ptr; } - operator const T&() const { + operator const T&() const noexcept { __SYCL_HOST_NOT_SUPPORTED("Implicit conversion of annotated_arg to T") - return ptr; + return g_ptr; } // template::value>> @@ -103,23 +121,9 @@ class __SYCL_SPECIAL_CLASS annotated_arg, type // return ptr[idx]; // } - // auto operator [](std::ptrdiff_t idx) { - // __SYCL_HOST_NOT_SUPPORTED("operator[] on an annotated_arg") - // return ptr[idx]; - // } - - // inline T& get() { - // __SYCL_HOST_NOT_SUPPORTED("get()") - // return ptr; - // } - // inline const T& get() const { - // __SYCL_HOST_NOT_SUPPORTED("get()") - // return ptr; - // } - - inline T get() const { - __SYCL_HOST_NOT_SUPPORTED("get()") - return ptr; + auto operator [](std::ptrdiff_t idx) { + __SYCL_HOST_NOT_SUPPORTED("operator[] on an annotated_arg") + return g_ptr[idx]; } template static constexpr bool has_property() { @@ -133,7 +137,7 @@ class __SYCL_SPECIAL_CLASS annotated_arg, type // Partial specialization for non-pointer type template -class __SYCL_SPECIAL_CLASS annotated_arg , typename std::enable_if::value>::type> { +class __SYCL_SPECIAL_CLASS __SYCL_TYPE(annotated_arg) annotated_arg , typename std::enable_if::value>::type> { using property_list_t = detail::properties_t; T obj; @@ -158,9 +162,51 @@ class __SYCL_SPECIAL_CLASS annotated_arg , typ static_assert(is_property_list::value, "Property list is invalid."); - annotated_arg() = default; + annotated_arg() noexcept = default; annotated_arg(const annotated_arg&) = default; - annotated_arg(const T &_obj) : obj(_obj) {}; + annotated_arg& operator=(annotated_arg&) = default; + + // annotated_arg(const T& _obj, const property_list_t &PropList = properties{}) noexcept : obj(_obj) {} + // template + template + annotated_arg(const T& _obj, PropertyValueTs... props) noexcept : obj(_obj) { + static_assert( + std::is_same< + // std::tuple, + // detail::MergeProperties< + // std::tuple, + // std::tuple + // >::type + property_list_t, + detail::merged_properties_t< property_list_t, detail::properties_t > >::value, + "The property list must contain all properties of the input of the constructor" + ); + } + + // // Constructs an annotated_arg object from another annotated_arg object. + // // The property set PropertyListT contains all properties of the input annotated_arg object. + // // If there are duplicate properties present in the property list of the input annotated_arg object, + // // the values of the duplicate properties must be the same. + // template ::value>::type> + // annotated_arg(const annotated_arg &other) noexcept : obj(other.obj) { + // static_assert( + // std::is_same< + // property_list_t, + // detail::MergeProperties::type>::value, + // "The property list must contain all properties of the input of the copy constructor"); + // } + + // template ::value>::type> + // explicit annotated_arg(const annotated_arg &other, + // properties proplist) noexcept { + // static_assert( + // std::is_same< + // property_list_t, + // detail::MergeProperties::type>::value, + // "The property list must contain all properties of the input of the copy constructor"); + // this->obj = other.obj; + // } operator T&() { __SYCL_HOST_NOT_SUPPORTED("Implicit conversion of annotated_arg to T") @@ -171,14 +217,22 @@ class __SYCL_SPECIAL_CLASS annotated_arg , typ return obj; } - inline T& get() { - __SYCL_HOST_NOT_SUPPORTED("get()") - return obj; - } - inline const T& get() const { - __SYCL_HOST_NOT_SUPPORTED("get()") - return obj; - } + // template + // template + // std::enable_if_t::value> + // &operator()(Args... args) noexcept { + // __SYCL_HOST_NOT_SUPPORTED("operator() on an annotated_arg") + // return obj.operator(args); + // } + + // inline T& get() { + // __SYCL_HOST_NOT_SUPPORTED("get()") + // return obj; + // } + // inline const T& get() const { + // __SYCL_HOST_NOT_SUPPORTED("get()") + // return obj; + // } template static constexpr bool has_property() { return property_list_t::template has_property(); From addd1c2f889d606cab58894fb0183dae1fb9b466 Mon Sep 17 00:00:00 2001 From: "Wang, Di5" Date: Tue, 22 Nov 2022 19:35:41 -0800 Subject: [PATCH 08/26] add ctors for case when T is a pointer type --- .../oneapi/annotated_arg/annotated_arg.hpp | 97 +++++++++++++------ .../ext/oneapi/annotated_arg/properties.hpp | 8 +- 2 files changed, 73 insertions(+), 32 deletions(-) diff --git a/sycl/include/sycl/ext/oneapi/annotated_arg/annotated_arg.hpp b/sycl/include/sycl/ext/oneapi/annotated_arg/annotated_arg.hpp index 7f8ae15eaf6a4..a52d25a6476b8 100644 --- a/sycl/include/sycl/ext/oneapi/annotated_arg/annotated_arg.hpp +++ b/sycl/include/sycl/ext/oneapi/annotated_arg/annotated_arg.hpp @@ -103,13 +103,52 @@ class __SYCL_SPECIAL_CLASS __SYCL_TYPE(annotated_arg) annotated_arg + annotated_arg(const T& _ptr, PropertyValueTs... props) noexcept : g_ptr((__OPENCL_GLOBAL_AS__ UnderlyingT*)_ptr) { + static_assert( + std::is_same< + property_list_t, + detail::merged_properties_t>>::value, + "The property list must contain all properties of the input of the constructor" + ); + } + + // Constructs an annotated_arg object from another annotated_arg object. + // The property set PropertyListT contains all properties of the input annotated_arg object. + // If there are duplicate properties present in the property list of the input annotated_arg object, + // the values of the duplicate properties must be the same. + template ::value>::type> + explicit annotated_arg(const annotated_arg &other) noexcept : g_ptr(other.g_ptr) { + static_assert( + std::is_same< + property_list_t, + detail::merged_properties_t>::value, + "The property list must contain all properties of the input of the copy constructor"); + } + + template ::value>::type> + explicit annotated_arg(const annotated_arg &other, + properties proplist) noexcept : g_ptr(other.g_ptr) { + static_assert( + std::is_same< + property_list_t, + detail::merged_properties_t>::value, + "The property list must contain all properties of the input of the copy constructor"); + } operator T&() noexcept { __SYCL_HOST_NOT_SUPPORTED("Implicit conversion of annotated_arg to T") - return g_ptr; + // return (T&) g_ptr; + return g_ptr; } + operator const T&() const noexcept { __SYCL_HOST_NOT_SUPPORTED("Implicit conversion of annotated_arg to T") return g_ptr; @@ -166,8 +205,8 @@ class __SYCL_SPECIAL_CLASS __SYCL_TYPE(annotated_arg) annotated_arg + annotated_arg(const T& _obj, const property_list_t &PropList = properties{}) noexcept : obj(_obj) {} + template annotated_arg(const T& _obj, PropertyValueTs... props) noexcept : obj(_obj) { static_assert( @@ -178,35 +217,35 @@ class __SYCL_SPECIAL_CLASS __SYCL_TYPE(annotated_arg) annotated_arg // >::type property_list_t, - detail::merged_properties_t< property_list_t, detail::properties_t > >::value, + detail::merged_properties_t>>::value, "The property list must contain all properties of the input of the constructor" ); } - // // Constructs an annotated_arg object from another annotated_arg object. - // // The property set PropertyListT contains all properties of the input annotated_arg object. - // // If there are duplicate properties present in the property list of the input annotated_arg object, - // // the values of the duplicate properties must be the same. - // template ::value>::type> - // annotated_arg(const annotated_arg &other) noexcept : obj(other.obj) { - // static_assert( - // std::is_same< - // property_list_t, - // detail::MergeProperties::type>::value, - // "The property list must contain all properties of the input of the copy constructor"); - // } + // Constructs an annotated_arg object from another annotated_arg object. + // The property set PropertyListT contains all properties of the input annotated_arg object. + // If there are duplicate properties present in the property list of the input annotated_arg object, + // the values of the duplicate properties must be the same. + template ::value>::type> + explicit annotated_arg(const annotated_arg &other) noexcept : obj(other.obj) { + static_assert( + std::is_same< + property_list_t, + detail::merged_properties_t>::value, + "The property list must contain all properties of the input of the copy constructor"); + } - // template ::value>::type> - // explicit annotated_arg(const annotated_arg &other, - // properties proplist) noexcept { - // static_assert( - // std::is_same< - // property_list_t, - // detail::MergeProperties::type>::value, - // "The property list must contain all properties of the input of the copy constructor"); - // this->obj = other.obj; - // } + template ::value>::type> + explicit annotated_arg(const annotated_arg &other, + properties proplist) noexcept { + static_assert( + std::is_same< + property_list_t, + detail::merged_properties_t>::value, + "The property list must contain all properties of the input of the copy constructor"); + this->obj = other.obj; + } operator T&() { __SYCL_HOST_NOT_SUPPORTED("Implicit conversion of annotated_arg to T") diff --git a/sycl/include/sycl/ext/oneapi/annotated_arg/properties.hpp b/sycl/include/sycl/ext/oneapi/annotated_arg/properties.hpp index bdcabd06d5367..e2a734e51632f 100644 --- a/sycl/include/sycl/ext/oneapi/annotated_arg/properties.hpp +++ b/sycl/include/sycl/ext/oneapi/annotated_arg/properties.hpp @@ -93,12 +93,14 @@ template inline constexpr maxburst_key::value_t maxburst; template inline constexpr wait_request_key::value_t wait_request; +inline constexpr wait_request_key::value_t<1> wait_request_requested; +inline constexpr wait_request_key::value_t<0> wait_request_not_requested; template inline constexpr read_write_mode_key::value_t read_write_mode; -inline constexpr read_write_mode_key::value_t read_only; -inline constexpr read_write_mode_key::value_t write_only; -inline constexpr read_write_mode_key::value_t read_write; +inline constexpr read_write_mode_key::value_t read_write_mode_read; +inline constexpr read_write_mode_key::value_t read_write_mode_write; +inline constexpr read_write_mode_key::value_t read_write_mode_readwrite; #endif // __cplusplus >= 201703L From 9360f9f9f08269654e83b5408912229e3e4c2580 Mon Sep 17 00:00:00 2001 From: "Wang, Di5" Date: Mon, 28 Nov 2022 12:25:17 -0800 Subject: [PATCH 09/26] add deduction guide --- .../oneapi/annotated_arg/annotated_arg.hpp | 57 ++++++++++++++++--- 1 file changed, 49 insertions(+), 8 deletions(-) diff --git a/sycl/include/sycl/ext/oneapi/annotated_arg/annotated_arg.hpp b/sycl/include/sycl/ext/oneapi/annotated_arg/annotated_arg.hpp index a52d25a6476b8..2fa55662d8e94 100644 --- a/sycl/include/sycl/ext/oneapi/annotated_arg/annotated_arg.hpp +++ b/sycl/include/sycl/ext/oneapi/annotated_arg/annotated_arg.hpp @@ -71,6 +71,24 @@ struct HasSubscriptOperator< } // namespace detail +// template +// annotated_arg(T, Args... args) -> annotated_arg; + +// template +// annotated_arg(annotated_arg, ArgT newp) -> annotated_arg, Enable>; + +template ::value>::type> +annotated_arg(T, Args... args) -> annotated_arg; + +template ::value>::type> +annotated_arg(T, Args... args) -> annotated_arg; + +// template +// annotated_arg(T, Args... args) -> annotated_arg::value>::type>; + +template +annotated_arg(annotated_arg, ArgT newp) -> annotated_arg, Enable>; + template class annotated_arg { // This should always fail when instantiating the unspecialized version. @@ -85,6 +103,13 @@ class __SYCL_SPECIAL_CLASS __SYCL_TYPE(annotated_arg) annotated_arg::type; __OPENCL_GLOBAL_AS__ UnderlyingT *g_ptr; + using base_t = annotated_arg::value>::type>; + + using IsPtr = typename std::enable_if::value>::type; + + template + friend class annotated_arg; + #ifdef __SYCL_DEVICE_ONLY__ void __init( [[__sycl_detail__::add_ir_attributes_kernel_parameter( @@ -106,17 +131,23 @@ class __SYCL_SPECIAL_CLASS __SYCL_TYPE(annotated_arg) annotated_arg - annotated_arg(const T& _ptr, PropertyValueTs... props) noexcept : g_ptr((__OPENCL_GLOBAL_AS__ UnderlyingT*)_ptr) { + explicit annotated_arg(const T& _ptr, PropertyValueTs... props) noexcept : g_ptr((__OPENCL_GLOBAL_AS__ UnderlyingT*)_ptr) { static_assert( std::is_same< property_list_t, detail::merged_properties_t>>::value, "The property list must contain all properties of the input of the constructor" ); + // static_assert( + // std::is_same< + // property_list_t, + // typename detail::merged_properties::type>::value, + // "The property list must contain all properties of the input of the constructor" + // ); } // Constructs an annotated_arg object from another annotated_arg object. @@ -132,10 +163,20 @@ class __SYCL_SPECIAL_CLASS __SYCL_TYPE(annotated_arg) annotated_arg::value>::type> - explicit annotated_arg(const annotated_arg &other, - properties proplist) noexcept : g_ptr(other.g_ptr) { + // template ::value>::type> + // explicit annotated_arg(const annotated_arg &other, + // properties proplist) noexcept : g_ptr(other.g_ptr) { + // static_assert( + // std::is_same< + // property_list_t, + // detail::merged_properties_t>::value, + // "The property list must contain all properties of the input of the copy constructor"); + // } + + template + explicit annotated_arg(const annotated_arg &other, + const PropertyListV& proplist) noexcept : g_ptr(other.g_ptr) { static_assert( std::is_same< property_list_t, @@ -143,13 +184,13 @@ class __SYCL_SPECIAL_CLASS __SYCL_TYPE(annotated_arg) annotated_arg Date: Wed, 30 Nov 2022 06:02:09 -0800 Subject: [PATCH 10/26] finish pointer's implementation of annnotated_arg --- .../oneapi/annotated_arg/annotated_arg.hpp | 212 ++++-------------- .../ext/oneapi/annotated_arg/properties.hpp | 42 ++-- 2 files changed, 63 insertions(+), 191 deletions(-) diff --git a/sycl/include/sycl/ext/oneapi/annotated_arg/annotated_arg.hpp b/sycl/include/sycl/ext/oneapi/annotated_arg/annotated_arg.hpp index 2fa55662d8e94..f70f0835715d1 100644 --- a/sycl/include/sycl/ext/oneapi/annotated_arg/annotated_arg.hpp +++ b/sycl/include/sycl/ext/oneapi/annotated_arg/annotated_arg.hpp @@ -71,25 +71,18 @@ struct HasSubscriptOperator< } // namespace detail -// template -// annotated_arg(T, Args... args) -> annotated_arg; -// template -// annotated_arg(annotated_arg, ArgT newp) -> annotated_arg, Enable>; +template +annotated_arg(T, Args... args) -> annotated_arg, std::is_pointer::value>; -template ::value>::type> -annotated_arg(T, Args... args) -> annotated_arg; +// template +// annotated_arg(annotated_arg) -> annotated_arg; -template ::value>::type> -annotated_arg(T, Args... args) -> annotated_arg; +template +annotated_arg(annotated_arg, ArgT newp) -> annotated_arg, IsPtr>; -// template -// annotated_arg(T, Args... args) -> annotated_arg::value>::type>; -template -annotated_arg(annotated_arg, ArgT newp) -> annotated_arg, Enable>; - -template +template ::value> class annotated_arg { // This should always fail when instantiating the unspecialized version. static_assert(is_property_list::value, @@ -98,16 +91,12 @@ class annotated_arg { // Partial specialization for pointer type template -class __SYCL_SPECIAL_CLASS __SYCL_TYPE(annotated_arg) annotated_arg, typename std::enable_if::value>::type> { +class __SYCL_SPECIAL_CLASS __SYCL_TYPE(annotated_arg) annotated_arg, true> { using property_list_t = detail::properties_t; using UnderlyingT = typename std::remove_pointer::type; - __OPENCL_GLOBAL_AS__ UnderlyingT *g_ptr; - - using base_t = annotated_arg::value>::type>; - - using IsPtr = typename std::enable_if::value>::type; + __OPENCL_GLOBAL_AS__ UnderlyingT *obj; - template + template friend class annotated_arg; #ifdef __SYCL_DEVICE_ONLY__ @@ -116,8 +105,8 @@ class __SYCL_SPECIAL_CLASS __SYCL_TYPE(annotated_arg) annotated_arg::name..., detail::PropertyMetaInfo::value... )]] - __OPENCL_GLOBAL_AS__ UnderlyingT* _g_ptr) { - g_ptr = _g_ptr; + __OPENCL_GLOBAL_AS__ UnderlyingT* _obj) { + obj = _obj; } #endif @@ -132,78 +121,62 @@ class __SYCL_SPECIAL_CLASS __SYCL_TYPE(annotated_arg) annotated_arg - explicit annotated_arg(const T& _ptr, PropertyValueTs... props) noexcept : g_ptr((__OPENCL_GLOBAL_AS__ UnderlyingT*)_ptr) { + explicit annotated_arg(const T& _ptr, PropertyValueTs... props) noexcept : obj((__OPENCL_GLOBAL_AS__ UnderlyingT*)_ptr) { static_assert( std::is_same< property_list_t, detail::merged_properties_t>>::value, "The property list must contain all properties of the input of the constructor" ); - // static_assert( - // std::is_same< - // property_list_t, - // typename detail::merged_properties::type>::value, - // "The property list must contain all properties of the input of the constructor" - // ); } // Constructs an annotated_arg object from another annotated_arg object. // The property set PropertyListT contains all properties of the input annotated_arg object. // If there are duplicate properties present in the property list of the input annotated_arg object, // the values of the duplicate properties must be the same. - template ::value>::type> - explicit annotated_arg(const annotated_arg &other) noexcept : g_ptr(other.g_ptr) { + template + explicit annotated_arg(const annotated_arg &other) noexcept : obj(other.obj) { + + static_assert(std::is_convertible::value, + "The underlying data type of the input annotated_arg is not compatible"); + static_assert( std::is_same< property_list_t, detail::merged_properties_t>::value, - "The property list must contain all properties of the input of the copy constructor"); + "The constructed annotated_arg type must contain all the properties of the input annotated_arg"); } - // template ::value>::type> - // explicit annotated_arg(const annotated_arg &other, - // properties proplist) noexcept : g_ptr(other.g_ptr) { - // static_assert( - // std::is_same< - // property_list_t, - // detail::merged_properties_t>::value, - // "The property list must contain all properties of the input of the copy constructor"); - // } + template + explicit annotated_arg(const annotated_arg &other, + const PropertyListV& proplist) noexcept : obj(other.obj) { - template - explicit annotated_arg(const annotated_arg &other, - const PropertyListV& proplist) noexcept : g_ptr(other.g_ptr) { - static_assert( + static_assert(std::is_convertible::value, + "The underlying data type of the input annotated_arg is not compatible"); + + static_assert( std::is_same< property_list_t, detail::merged_properties_t>::value, - "The property list must contain all properties of the input of the copy constructor"); + "The property list of constructed annotated_arg type must be the union of the input property lists"); } - explicit operator T&() noexcept { + explicit operator T() noexcept { __SYCL_HOST_NOT_SUPPORTED("Implicit conversion of annotated_arg to T") - // return (T&) g_ptr; - return g_ptr; + return obj; } - explicit operator const T&() const noexcept { + explicit operator const T() const noexcept { __SYCL_HOST_NOT_SUPPORTED("Implicit conversion of annotated_arg to T") - return g_ptr; + return obj; } - // template::value>> - // std::remove_pointer_t operator [](std::ptrdiff_t idx) { - // __SYCL_HOST_NOT_SUPPORTED("operator[] on an annotated_arg") - // return ptr[idx]; - // } - - auto operator [](std::ptrdiff_t idx) { + UnderlyingT& operator [](std::ptrdiff_t idx) const noexcept { __SYCL_HOST_NOT_SUPPORTED("operator[] on an annotated_arg") - return g_ptr[idx]; + return obj[idx]; } template static constexpr bool has_property() { @@ -215,11 +188,15 @@ class __SYCL_SPECIAL_CLASS __SYCL_TYPE(annotated_arg) annotated_arg -class __SYCL_SPECIAL_CLASS __SYCL_TYPE(annotated_arg) annotated_arg , typename std::enable_if::value>::type> { +class __SYCL_SPECIAL_CLASS __SYCL_TYPE(annotated_arg) annotated_arg , false> { using property_list_t = detail::properties_t; + template + friend class annotated_arg; + T obj; #ifdef __SYCL_DEVICE_ONLY__ @@ -252,11 +229,6 @@ class __SYCL_SPECIAL_CLASS __SYCL_TYPE(annotated_arg) annotated_arg , - // detail::MergeProperties< - // std::tuple, - // std::tuple - // >::type property_list_t, detail::merged_properties_t>>::value, "The property list must contain all properties of the input of the constructor" @@ -267,7 +239,8 @@ class __SYCL_SPECIAL_CLASS __SYCL_TYPE(annotated_arg) annotated_arg ::value>::type> + // template ::value>::type> + template explicit annotated_arg(const annotated_arg &other) noexcept : obj(other.obj) { static_assert( std::is_same< @@ -305,15 +278,6 @@ class __SYCL_SPECIAL_CLASS __SYCL_TYPE(annotated_arg) annotated_arg static constexpr bool has_property() { return property_list_t::template has_property(); } @@ -324,98 +288,6 @@ class __SYCL_SPECIAL_CLASS __SYCL_TYPE(annotated_arg) annotated_arg -class annotated_arg { - // This should always fail when instantiating the unspecialized version. - static_assert(is_property_list::value, - "Property list is invalid."); -}; - -// Partial specialization to make PropertyListT visible as a parameter pack -// of properties. -template -class __SYCL_SPECIAL_CLASS annotated_arg> { - using property_list_t = detail::properties_t; - // using CondT = std::conditional::value, __OPENCL_GLOBAL_AS__ UnderlyingT, T>::type; - - #ifdef __SYCL_DEVICE_ONLY__ - void __init( - [[__sycl_detail__::add_ir_attributes_kernel_parameter( - detail::PropertyMetaInfo::name..., - detail::PropertyMetaInfo::value... - )]] - __OPENCL_GLOBAL_AS__ T* _obj) { - obj = _obj; - } - #endif - -public: - // T should be trivially copy constructible to be device copyable - static_assert(std::is_trivially_copy_constructible::value, - "Type T must be trivially copy constructable."); - static_assert(std::is_trivially_destructible::value, - "Type T must be trivially destructible."); - static_assert(is_property_list::value, - "Property list is invalid."); - - // Check compability of each property values in the property list - // static_assert(check_property_list::value, - // "property list contains invalid property."); - - annotated_arg() = default; - annotated_arg(const annotated_arg&) = default; - // annotated_arg(const __OPENCL_GLOBAL_AS__ T &_obj) : obj(_obj) {}; - annotated_arg(T *_obj) : obj((__OPENCL_GLOBAL_AS__ T*)_obj) {}; - - operator T&() { - __SYCL_HOST_NOT_SUPPORTED("Implicit conversion of annotated_arg to T") - return obj; - } - operator const T&() const { - __SYCL_HOST_NOT_SUPPORTED("Implicit conversion of annotated_arg to T") - return obj; - } - - // template::value>> - // std::remove_pointer_t operator [](std::ptrdiff_t idx) { - // __SYCL_HOST_NOT_SUPPORTED("operator[] on an annotated_arg") - // return obj[idx]; - // } - - // auto operator [](std::ptrdiff_t idx) { - // __SYCL_HOST_NOT_SUPPORTED("operator[] on an annotated_arg") - // return obj[idx]; - // } - - // inline T& get() { - // __SYCL_HOST_NOT_SUPPORTED("get()") - // return obj; - // } - // inline const T& get() const { - // __SYCL_HOST_NOT_SUPPORTED("get()") - // return obj; - // } - - inline T* get() const { - __SYCL_HOST_NOT_SUPPORTED("get()") - return obj; - } - // inline const T* get() const { - // __SYCL_HOST_NOT_SUPPORTED("get()") - // return obj; - // } - - template static constexpr bool has_property() { - return property_list_t::template has_property(); - } - - template static constexpr auto get_property() { - return property_list_t::template get_property(); - } -}; -*/ - } // namespace experimental } // namespace oneapi } // namespace ext diff --git a/sycl/include/sycl/ext/oneapi/annotated_arg/properties.hpp b/sycl/include/sycl/ext/oneapi/annotated_arg/properties.hpp index e2a734e51632f..5e2870f2454f0 100644 --- a/sycl/include/sycl/ext/oneapi/annotated_arg/properties.hpp +++ b/sycl/include/sycl/ext/oneapi/annotated_arg/properties.hpp @@ -17,7 +17,7 @@ namespace ext { namespace oneapi { namespace experimental { -template class annotated_arg; +template class annotated_arg; struct register_map_key { using value_t = property_value; @@ -118,44 +118,44 @@ template <> struct is_property_key : std::true_type {}; template <> struct is_property_key : std::true_type {}; -template -struct is_property_key_of> +template +struct is_property_key_of> : std::true_type {}; -template -struct is_property_key_of> +template +struct is_property_key_of> : std::true_type {}; -template -struct is_property_key_of> +template +struct is_property_key_of> : std::true_type {}; -template -struct is_property_key_of> +template +struct is_property_key_of> : std::true_type {}; -template -struct is_property_key_of> +template +struct is_property_key_of> : std::true_type {}; -template -struct is_property_key_of> +template +struct is_property_key_of> : std::true_type {}; -template -struct is_property_key_of> +template +struct is_property_key_of> : std::true_type {}; -template -struct is_property_key_of> +template +struct is_property_key_of> : std::true_type {}; -template -struct is_property_key_of> +template +struct is_property_key_of> : std::true_type {}; -template -struct is_property_key_of> +template +struct is_property_key_of> : std::true_type {}; From 5ab61cf6304285f152068ea7f6ee02abbd27aeec Mon Sep 17 00:00:00 2001 From: "Wang, Di5" Date: Wed, 30 Nov 2022 13:47:44 -0800 Subject: [PATCH 11/26] finish non-pointer's implementation of annnotated_arg --- .../oneapi/annotated_arg/annotated_arg.hpp | 59 ++++++++++--------- .../ext/oneapi/annotated_arg/properties.hpp | 28 +++++---- 2 files changed, 45 insertions(+), 42 deletions(-) diff --git a/sycl/include/sycl/ext/oneapi/annotated_arg/annotated_arg.hpp b/sycl/include/sycl/ext/oneapi/annotated_arg/annotated_arg.hpp index f70f0835715d1..2f837c993c3a6 100644 --- a/sycl/include/sycl/ext/oneapi/annotated_arg/annotated_arg.hpp +++ b/sycl/include/sycl/ext/oneapi/annotated_arg/annotated_arg.hpp @@ -75,13 +75,12 @@ struct HasSubscriptOperator< template annotated_arg(T, Args... args) -> annotated_arg, std::is_pointer::value>; -// template -// annotated_arg(annotated_arg) -> annotated_arg; +// template +// annotated_arg(T, properties>) -> annotated_arg, std::is_pointer::value>; template annotated_arg(annotated_arg, ArgT newp) -> annotated_arg, IsPtr>; - template ::value> class annotated_arg { // This should always fail when instantiating the unspecialized version. @@ -113,15 +112,17 @@ class __SYCL_SPECIAL_CLASS __SYCL_TYPE(annotated_arg) annotated_arg::value, "Type T must be trivially destructible."); - static_assert(is_property_list::value, - "Property list is invalid."); + // static_assert(is_property_list::value, + // "Property list is invalid."); + static_assert(check_property_list::value, + "The property list contains invalid property."); annotated_arg() noexcept = default; annotated_arg(const annotated_arg&) = default; annotated_arg& operator=(annotated_arg&) = default; - explicit annotated_arg(const T& _ptr, const property_list_t &PropList = properties{}) noexcept - : obj((__OPENCL_GLOBAL_AS__ UnderlyingT*)_ptr) {} + // explicit annotated_arg(const T& _ptr, const property_list_t &PropList = properties{}) noexcept + // : obj((__OPENCL_GLOBAL_AS__ UnderlyingT*)_ptr) {} template explicit annotated_arg(const T& _ptr, PropertyValueTs... props) noexcept : obj((__OPENCL_GLOBAL_AS__ UnderlyingT*)_ptr) { @@ -139,7 +140,6 @@ class __SYCL_SPECIAL_CLASS __SYCL_TYPE(annotated_arg) annotated_arg explicit annotated_arg(const annotated_arg &other) noexcept : obj(other.obj) { - static_assert(std::is_convertible::value, "The underlying data type of the input annotated_arg is not compatible"); @@ -153,7 +153,6 @@ class __SYCL_SPECIAL_CLASS __SYCL_TYPE(annotated_arg) annotated_arg explicit annotated_arg(const annotated_arg &other, const PropertyListV& proplist) noexcept : obj(other.obj) { - static_assert(std::is_convertible::value, "The underlying data type of the input annotated_arg is not compatible"); @@ -164,18 +163,15 @@ class __SYCL_SPECIAL_CLASS __SYCL_TYPE(annotated_arg) annotated_arg::value, - "Type T must be trivially copy constructable."); + static_assert(std::is_trivially_copyable::value, + "Type T must be trivially copyable."); + // static_assert(std::is_trivially_copy_constructible::value, + // "Type T must be trivially copy constructable."); static_assert(std::is_trivially_destructible::value, "Type T must be trivially destructible."); static_assert(is_property_list::value, "Property list is invalid."); + static_assert(check_property_list::value, + "The property list contains invalid property."); annotated_arg() noexcept = default; annotated_arg(const annotated_arg&) = default; annotated_arg& operator=(annotated_arg&) = default; - annotated_arg(const T& _obj, const property_list_t &PropList = properties{}) noexcept : obj(_obj) {} + explicit annotated_arg(const T& _obj, const property_list_t &PropList = properties{}) noexcept : obj(_obj) {} template - annotated_arg(const T& _obj, PropertyValueTs... props) noexcept : obj(_obj) { + explicit annotated_arg(const T& _obj, PropertyValueTs... props) noexcept : obj(_obj) { static_assert( std::is_same< property_list_t, @@ -239,34 +239,35 @@ class __SYCL_SPECIAL_CLASS __SYCL_TYPE(annotated_arg) annotated_arg ::value>::type> template explicit annotated_arg(const annotated_arg &other) noexcept : obj(other.obj) { + static_assert(std::is_convertible::value, + "The underlying data type of the input annotated_arg is not compatible"); + static_assert( std::is_same< property_list_t, detail::merged_properties_t>::value, - "The property list must contain all properties of the input of the copy constructor"); + "The constructed annotated_arg type must contain all the properties of the input annotated_arg"); } - template ::value>::type> + template explicit annotated_arg(const annotated_arg &other, - properties proplist) noexcept { + const PropertyListV& proplist) noexcept : obj(other.obj) { + static_assert(std::is_convertible::value, + "The underlying data type of the input annotated_arg is not compatible"); + static_assert( std::is_same< property_list_t, detail::merged_properties_t>::value, - "The property list must contain all properties of the input of the copy constructor"); - this->obj = other.obj; + "The property list of constructed annotated_arg type must be the union of the input property lists"); } - operator T&() { - __SYCL_HOST_NOT_SUPPORTED("Implicit conversion of annotated_arg to T") + operator T() noexcept { return obj; } - operator const T&() const { - __SYCL_HOST_NOT_SUPPORTED("Implicit conversion of annotated_arg to T") + operator const T() const noexcept { return obj; } diff --git a/sycl/include/sycl/ext/oneapi/annotated_arg/properties.hpp b/sycl/include/sycl/ext/oneapi/annotated_arg/properties.hpp index 5e2870f2454f0..b508549df5a42 100644 --- a/sycl/include/sycl/ext/oneapi/annotated_arg/properties.hpp +++ b/sycl/include/sycl/ext/oneapi/annotated_arg/properties.hpp @@ -262,52 +262,54 @@ struct PropertyMetaInfo> { // 'buffer_location' and mmhost properties are pointers-only template -struct is_property_compatible : std::false_type {}; +struct is_valid_property : std::false_type {}; template -struct is_property_compatible> +struct is_valid_property> : std::bool_constant::value> {}; template -struct is_property_compatible> +struct is_valid_property> : std::bool_constant::value> {}; template -struct is_property_compatible> +struct is_valid_property> : std::bool_constant::value> {}; template -struct is_property_compatible> +struct is_valid_property> : std::bool_constant::value> {}; template -struct is_property_compatible> +struct is_valid_property> : std::bool_constant::value> {}; template -struct is_property_compatible> +struct is_valid_property> : std::bool_constant::value> {}; template -struct is_property_compatible> +struct is_valid_property> : std::bool_constant::value> {}; // 'register_map', 'conduit', 'stable' are common properties for pointers // and non pointers; template -struct is_property_compatible : std::true_type {}; +struct is_valid_property : std::true_type {}; template -struct is_property_compatible : std::true_type {}; +struct is_valid_property : std::true_type {}; template -struct is_property_compatible : std::true_type {}; +struct is_valid_property : std::true_type {}; template struct check_property_list : std::true_type {}; template struct check_property_list - : std::conditional_t::value, + : std::conditional_t::value, check_property_list, - std::false_type> {}; + std::false_type> { + static_assert(is_valid_property::value, "Property is invalid for the given type."); +}; } // namespace experimental } // namespace oneapi From 33a2557c76af6423e1dd4c0ae9f42aa96944434a Mon Sep 17 00:00:00 2001 From: "Wang, Di5" Date: Thu, 1 Dec 2022 02:35:47 -0800 Subject: [PATCH 12/26] small fixes --- .../oneapi/annotated_arg/annotated_arg.hpp | 21 +++++++++++-------- 1 file changed, 12 insertions(+), 9 deletions(-) diff --git a/sycl/include/sycl/ext/oneapi/annotated_arg/annotated_arg.hpp b/sycl/include/sycl/ext/oneapi/annotated_arg/annotated_arg.hpp index 2f837c993c3a6..7f4563b34c470 100644 --- a/sycl/include/sycl/ext/oneapi/annotated_arg/annotated_arg.hpp +++ b/sycl/include/sycl/ext/oneapi/annotated_arg/annotated_arg.hpp @@ -25,6 +25,7 @@ Op " is not supported on host device."); #endif + #ifdef __SYCL_DEVICE_ONLY__ #define __OPENCL_GLOBAL_AS__ __attribute__((opencl_global)) #ifdef __ENABLE_USM_ADDR_SPACE__ @@ -75,11 +76,11 @@ struct HasSubscriptOperator< template annotated_arg(T, Args... args) -> annotated_arg, std::is_pointer::value>; -// template -// annotated_arg(T, properties>) -> annotated_arg, std::is_pointer::value>; +template +annotated_arg(T, properties>) -> annotated_arg, std::is_pointer::value>; -template -annotated_arg(annotated_arg, ArgT newp) -> annotated_arg, IsPtr>; +template +annotated_arg(annotated_arg, properties>) -> annotated_arg>, IsPtr>; template ::value> class annotated_arg { @@ -116,16 +117,18 @@ class __SYCL_SPECIAL_CLASS __SYCL_TYPE(annotated_arg) annotated_arg::value, "The property list contains invalid property."); + static_assert(detail::SortedAllUnique>::value, + "Duplicate properties in property list."); annotated_arg() noexcept = default; annotated_arg(const annotated_arg&) = default; annotated_arg& operator=(annotated_arg&) = default; - // explicit annotated_arg(const T& _ptr, const property_list_t &PropList = properties{}) noexcept - // : obj((__OPENCL_GLOBAL_AS__ UnderlyingT*)_ptr) {} + annotated_arg(const T& _ptr, const property_list_t &PropList = properties{}) noexcept + : obj((__OPENCL_GLOBAL_AS__ UnderlyingT*)_ptr) {} template - explicit annotated_arg(const T& _ptr, PropertyValueTs... props) noexcept : obj((__OPENCL_GLOBAL_AS__ UnderlyingT*)_ptr) { + annotated_arg(const T& _ptr, PropertyValueTs... props) noexcept : obj((__OPENCL_GLOBAL_AS__ UnderlyingT*)_ptr) { static_assert( std::is_same< property_list_t, @@ -223,10 +226,10 @@ class __SYCL_SPECIAL_CLASS __SYCL_TYPE(annotated_arg) annotated_arg - explicit annotated_arg(const T& _obj, PropertyValueTs... props) noexcept : obj(_obj) { + annotated_arg(const T& _obj, PropertyValueTs... props) noexcept : obj(_obj) { static_assert( std::is_same< property_list_t, From acc0beea0ca0db4b16de20bc5e49a487daff7946 Mon Sep 17 00:00:00 2001 From: "Wang, Di5" Date: Thu, 1 Dec 2022 06:49:19 -0800 Subject: [PATCH 13/26] add operator[] --- .../oneapi/annotated_arg/annotated_arg.hpp | 32 ++++++++++--------- 1 file changed, 17 insertions(+), 15 deletions(-) diff --git a/sycl/include/sycl/ext/oneapi/annotated_arg/annotated_arg.hpp b/sycl/include/sycl/ext/oneapi/annotated_arg/annotated_arg.hpp index 7f4563b34c470..020c7f652d2a8 100644 --- a/sycl/include/sycl/ext/oneapi/annotated_arg/annotated_arg.hpp +++ b/sycl/include/sycl/ext/oneapi/annotated_arg/annotated_arg.hpp @@ -54,20 +54,14 @@ namespace oneapi { namespace experimental { namespace detail { -// Type-trait for checking if a type defines `operator->`. -template -struct HasParenthesisOperator : std::false_type {}; -template -struct HasParenthesisOperator< - T, sycl::detail::void_t().operator()())>> - : std::true_type {}; +// Type-trait for checking if a type defines `operator[]`. template struct HasSubscriptOperator : std::false_type {}; template struct HasSubscriptOperator< - T, sycl::detail::void_t().operator[]())>> + T, typename std::enable_if_t().operator[](0))>::value>> : std::true_type {}; } // namespace detail @@ -274,13 +268,21 @@ class __SYCL_SPECIAL_CLASS __SYCL_TYPE(annotated_arg) annotated_arg - // template - // std::enable_if_t::value> - // &operator()(Args... args) noexcept { - // __SYCL_HOST_NOT_SUPPORTED("operator() on an annotated_arg") - // return obj.operator(args); - // } + template + std::enable_if_t< + detail::HasSubscriptOperator::value, + const decltype(std::declval().operator[](0))> + &operator[](std::ptrdiff_t idx) const noexcept { + return obj.operator[](idx); + } + + template + std::enable_if_t< + detail::HasSubscriptOperator::value, + decltype(std::declval().operator[](0))> + &operator[](std::ptrdiff_t idx) noexcept { + return obj.operator[](idx); + } template static constexpr bool has_property() { return property_list_t::template has_property(); From 55f300e4db91a3d80cfaf80ed8f69e7e5eed4344 Mon Sep 17 00:00:00 2001 From: "Wang, Di5" Date: Thu, 1 Dec 2022 16:15:30 -0800 Subject: [PATCH 14/26] revert SPIRV changes --- clang/lib/Driver/ToolChains/Clang.cpp | 2 +- llvm-spirv/CMakeLists.txt | 4 +- llvm-spirv/lib/SPIRV/SPIRVWriter.cpp | 24 ------ llvm-spirv/lib/SPIRV/libSPIRV/SPIRVEnum.h | 20 ----- .../lib/SPIRV/libSPIRV/SPIRVNameMapEnum.h | 11 --- .../sycl-kernel-arg-annotations.ll | 38 --------- .../sycl-post-link/CompileTimeProperties.def | 18 +--- .../CompileTimePropertiesPass.cpp | 84 +++++++------------ 8 files changed, 34 insertions(+), 167 deletions(-) delete mode 100644 llvm/test/tools/sycl-post-link/sycl-kernel-arg-annotations.ll diff --git a/clang/lib/Driver/ToolChains/Clang.cpp b/clang/lib/Driver/ToolChains/Clang.cpp index 64f021b4de5af..e53e195a1c673 100644 --- a/clang/lib/Driver/ToolChains/Clang.cpp +++ b/clang/lib/Driver/ToolChains/Clang.cpp @@ -9534,13 +9534,13 @@ void SPIRVTranslator::ConstructJob(Compilation &C, const JobAction &JA, ",+SPV_INTEL_variable_length_array,+SPV_INTEL_fp_fast_math_mode" ",+SPV_INTEL_long_constant_composite" ",+SPV_INTEL_arithmetic_fence" - ",+SPV_INTEL_fpga_buffer_location" ",+SPV_INTEL_global_variable_decorations"; ExtArg = ExtArg + DefaultExtArg + INTELExtArg; if (!C.getDriver().isFPGAEmulationMode()) // Enable several extensions on FPGA H/W exclusively ExtArg += ",+SPV_INTEL_usm_storage_classes,+SPV_INTEL_runtime_aligned" ",+SPV_INTEL_fpga_cluster_attributes,+SPV_INTEL_loop_fuse" + ",+SPV_INTEL_fpga_buffer_location" ",+SPV_INTEL_fpga_invocation_pipelining_attributes" ",+SPV_INTEL_fpga_dsp_control,+SPV_INTEL_fpga_memory_accesses" ",+SPV_INTEL_fpga_memory_attributes"; diff --git a/llvm-spirv/CMakeLists.txt b/llvm-spirv/CMakeLists.txt index 34403380634ab..d2d668f3495ed 100644 --- a/llvm-spirv/CMakeLists.txt +++ b/llvm-spirv/CMakeLists.txt @@ -38,8 +38,8 @@ if(NOT DEFINED LLVM_TOOL_SPIRV_HEADERS_BUILD AND # Strip the potential trailing newline from tag string(STRIP "${SPIRV_HEADERS_TAG}" SPIRV_HEADERS_TAG) FetchContent_Declare(spirv-headers - GIT_REPOSITORY https://github.com/broxigarchen/SPIRV-Headers.git - GIT_TAG kernelArgProperties + GIT_REPOSITORY https://github.com/KhronosGroup/SPIRV-Headers.git + GIT_TAG ${SPIRV_HEADERS_TAG} SOURCE_DIR ${LLVM_EXTERNAL_SPIRV_HEADERS_SOURCE_DIR} ) FetchContent_MakeAvailable(spirv-headers) diff --git a/llvm-spirv/lib/SPIRV/SPIRVWriter.cpp b/llvm-spirv/lib/SPIRV/SPIRVWriter.cpp index a6f59ca525f56..3277d1ed41b04 100644 --- a/llvm-spirv/lib/SPIRV/SPIRVWriter.cpp +++ b/llvm-spirv/lib/SPIRV/SPIRVWriter.cpp @@ -2393,30 +2393,6 @@ static void transMetadataDecorations(Metadata *MD, SPIRVEntry *Target) { SPIRVWord); TWO_INT_DECORATION_CASE(MathOpDSPModeINTEL, spv::internal, SPIRVWord, SPIRVWord); - case DecorationConduitKernelArgumentINTEL: - case DecorationRegisterMapKernelArgumentINTEL: - case DecorationStableKernelArgumentINTEL: - case DecorationRestrict: { - Target->addDecorate(new SPIRVDecorate(DecoKind, Target)); - break; - } - case DecorationBufferLocationINTEL: - case DecorationMMHostInterfaceReadWriteModeINTEL: - case DecorationMMHostInterfaceAddressWidthINTEL: - case DecorationMMHostInterfaceDataWidthINTEL: - case DecorationMMHostInterfaceLatencyINTEL: - case DecorationMMHostInterfaceMaxBurstINTEL: - case DecorationMMHostInterfaceAlignmentINTEL: - case DecorationMMHostInterfaceWaitRequestINTEL: { - ErrLog.checkError(NumOperands == 2, SPIRVEC_InvalidLlvmModule, - "MMHost Kernel Argument Annotation requires exactly 2 " - "extra operands"); - auto *DecoValEO1 = - mdconst::dyn_extract(DecoMD->getOperand(1)); - Target->addDecorate( - new SPIRVDecorate(DecoKind, Target, DecoValEO1->getZExtValue())); - break; - } case DecorationStallEnableINTEL: { Target->addDecorate(new SPIRVDecorateStallEnableINTEL(Target)); break; diff --git a/llvm-spirv/lib/SPIRV/libSPIRV/SPIRVEnum.h b/llvm-spirv/lib/SPIRV/libSPIRV/SPIRVEnum.h index 144952b0ff432..46024d2fefbc9 100644 --- a/llvm-spirv/lib/SPIRV/libSPIRV/SPIRVEnum.h +++ b/llvm-spirv/lib/SPIRV/libSPIRV/SPIRVEnum.h @@ -456,26 +456,6 @@ template <> inline void SPIRVMap::init() { {internal::CapabilityGlobalVariableDecorationsINTEL}); ADD_VEC_INIT(internal::DecorationArgumentAttributeINTEL, {CapabilityFunctionPointersINTEL}); - ADD_VEC_INIT(DecorationConduitKernelArgumentINTEL, - {CapabilityFPGAKernelArgumentPropertiesINTEL}); - ADD_VEC_INIT(DecorationRegisterMapKernelArgumentINTEL, - {CapabilityFPGAKernelArgumentPropertiesINTEL}); - ADD_VEC_INIT(DecorationMMHostInterfaceAddressWidthINTEL, - {CapabilityFPGAKernelArgumentPropertiesINTEL}); - ADD_VEC_INIT(DecorationMMHostInterfaceDataWidthINTEL, - {CapabilityFPGAKernelArgumentPropertiesINTEL}); - ADD_VEC_INIT(DecorationMMHostInterfaceLatencyINTEL, - {CapabilityFPGAKernelArgumentPropertiesINTEL}); - ADD_VEC_INIT(DecorationMMHostInterfaceReadWriteModeINTEL, - {CapabilityFPGAKernelArgumentPropertiesINTEL}); - ADD_VEC_INIT(DecorationMMHostInterfaceMaxBurstINTEL, - {CapabilityFPGAKernelArgumentPropertiesINTEL}); - ADD_VEC_INIT(DecorationMMHostInterfaceAlignmentINTEL, - {CapabilityFPGAKernelArgumentPropertiesINTEL}); - ADD_VEC_INIT(DecorationMMHostInterfaceWaitRequestINTEL, - {CapabilityFPGAKernelArgumentPropertiesINTEL}); - ADD_VEC_INIT(DecorationStableKernelArgumentINTEL, - {CapabilityFPGAKernelArgumentPropertiesINTEL}); } template <> inline void SPIRVMap::init() { diff --git a/llvm-spirv/lib/SPIRV/libSPIRV/SPIRVNameMapEnum.h b/llvm-spirv/lib/SPIRV/libSPIRV/SPIRVNameMapEnum.h index a1f0f130b987f..d098ff0d8b244 100644 --- a/llvm-spirv/lib/SPIRV/libSPIRV/SPIRVNameMapEnum.h +++ b/llvm-spirv/lib/SPIRV/libSPIRV/SPIRVNameMapEnum.h @@ -179,16 +179,6 @@ template <> inline void SPIRVMap::init() { add(DecorationMediaBlockIOINTEL, "MediaBlockIOINTEL"); add(DecorationAliasScopeINTEL, "AliasScopeINTEL"); add(DecorationNoAliasINTEL, "NoAliasINTEL"); - add(DecorationConduitKernelArgumentINTEL, "ConduitKernelArgumentINTEL"); - add(DecorationRegisterMapKernelArgumentINTEL, "RegisterMapKernelArgumentINTEL"); - add(DecorationMMHostInterfaceAddressWidthINTEL, "MMHostInterfaceAddressWidthINTEL"); - add(DecorationMMHostInterfaceDataWidthINTEL, "MMHostInterfaceDataWidthINTEL"); - add(DecorationMMHostInterfaceLatencyINTEL, "MMHostInterfaceLatencyINTEL"); - add(DecorationMMHostInterfaceReadWriteModeINTEL, "MMHostInterfaceReadWriteModeINTEL"); - add(DecorationMMHostInterfaceMaxBurstINTEL, "MMHostInterfaceMaxBurstINTEL"); - add(DecorationMMHostInterfaceAlignmentINTEL, "MMHostInterfaceAlignmentINTEL"); - add(DecorationMMHostInterfaceWaitRequestINTEL, "MMHostInterfaceWaitRequestINTEL"); - add(DecorationStableKernelArgumentINTEL, "StableKernelArgumentINTEL"); // From spirv_internal.hpp add(internal::DecorationFuncParamKindINTEL, "FuncParamKindINTEL"); @@ -609,7 +599,6 @@ template <> inline void SPIRVMap::init() { add(CapabilityDebugInfoModuleINTEL, "DebugInfoModuleINTEL"); add(CapabilitySplitBarrierINTEL, "SplitBarrierINTEL"); add(CapabilityGroupUniformArithmeticKHR, "GroupUniformArithmeticKHR"); - add(CapabilityFPGAKernelArgumentPropertiesINTEL, "FPGAKernelArgumentPropertiesINTEL"); // From spirv_internal.hpp add(internal::CapabilityFPGADSPControlINTEL, "FPGADSPControlINTEL"); diff --git a/llvm/test/tools/sycl-post-link/sycl-kernel-arg-annotations.ll b/llvm/test/tools/sycl-post-link/sycl-kernel-arg-annotations.ll deleted file mode 100644 index 98f9eba6d247a..0000000000000 --- a/llvm/test/tools/sycl-post-link/sycl-kernel-arg-annotations.ll +++ /dev/null @@ -1,38 +0,0 @@ -; RUN: sycl-post-link --device-globals --ir-output-only -S %s -o %t.ll -; RUN: FileCheck %s -input-file=%t.ll -; -; TODO: Remove --device-globals once other features start using compile-time -; properties. -; -; Tests the translation of "sycl-kernel-arg-attribute" to "spirv.ParameterDecorations" metadata - -target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-n8:16:32:64" -target triple = "spir64_fpga-unknown-unknown" - -$singleArg = comdat any - -; Function Attrs: convergent mustprogress norecurse -define weak_odr dso_local spir_kernel void @singleArg(i32 addrspace(4)* noundef align 4 "sycl-alignment"="4" "sycl-awidth"="32" "sycl-buffer-location"="10" "sycl-conduit" "sycl-dwidth"="64" "sycl-latency"="1" "sycl-maxburst"="3" "sycl-read-write-mode"="2" "sycl-register-map" "sycl-stable" "sycl-strict" "sycl-wait-request"="5" %_arg_p) #0 comdat !kernel_arg_buffer_location !1587 -; CHECK-DAG: !spirv.ParameterDecorations ![[PARMDECOR:[0-9]+]] -{ - -entry: - ret void -} - -!1587 = !{i32 -1} -; CHECK-DAG: ![[PARMDECOR]] = !{![[ARG1:[0-9]+]]} -; CHECK-DAG: ![[ARG1]] = !{![[ALIGN:[0-9]+]], ![[AWIDTH:[0-9]+]], ![[BL:[0-9]+]], ![[CONDUIT:[0-9]+]], ![[DWIDTH:[0-9]+]], ![[LATENCY:[0-9]+]], ![[MAXBURST:[0-9]+]], ![[RWMODE:[0-9]+]], ![[REGMAP:[0-9]+]], ![[STABLE:[0-9]+]], ![[STRICT:[0-9]+]], ![[WAITREQ:[0-9]+]]} - -; CHECK: ![[ALIGN]] = !{i32 6182, i32 4} -; CHECK: ![[AWIDTH]] = !{i32 6177, i32 32} -; CHECK: ![[BL]] = !{i32 5921, i32 10} -; CHECK: ![[CONDUIT]] = !{i32 6175, i32 1} -; CHECK: ![[DWIDTH]] = !{i32 6178, i32 64} -; CHECK: ![[LATENCY]] = !{i32 6179, i32 1} -; CHECK: ![[MAXBURST]] = !{i32 6181, i32 3} -; CHECK: ![[RWMODE]] = !{i32 6180, i32 2} -; CHECK: ![[REGMAP]] = !{i32 6176, i32 1} -; CHECK: ![[STABLE]] = !{i32 6184, i32 1} -; CHECK: ![[STRICT]] = !{i32 19, i32 1} -; CHECK: ![[WAITREQ]] = !{i32 6183, i32 5} diff --git a/llvm/tools/sycl-post-link/CompileTimeProperties.def b/llvm/tools/sycl-post-link/CompileTimeProperties.def index 28a827a32689a..24628840dc360 100644 --- a/llvm/tools/sycl-post-link/CompileTimeProperties.def +++ b/llvm/tools/sycl-post-link/CompileTimeProperties.def @@ -7,8 +7,7 @@ \*===----------------------------------------------------------------------===*/ #ifndef SYCL_COMPILE_TIME_PROPERTY -#error \ - "SYCL_COMPILE_TIME_PROPERTY(PropertyName, Decoration, ValueType) is not defined." +#error "SYCL_COMPILE_TIME_PROPERTY(PropertyName, Decoration, ValueType) is not defined." #endif // The corresponding SPIR-V OpCodes for the sycl-init-mode and @@ -17,18 +16,3 @@ // https://github.com/intel/llvm/blob/sycl/sycl/doc/extensions/DeviceGlobal/SPV_INTEL_global_variable_decorations.asciidoc#decoration SYCL_COMPILE_TIME_PROPERTY("sycl-init-mode", 6148, DecorValueTy::uint32) SYCL_COMPILE_TIME_PROPERTY("sycl-implement-in-csr", 6149, DecorValueTy::boolean) - -// The corresponding SPIR-V OpCodes for kernel pointer/arg properties -// docs are TBD -SYCL_COMPILE_TIME_PROPERTY("sycl-buffer-location", 5921, DecorValueTy::uint32) -SYCL_COMPILE_TIME_PROPERTY("sycl-conduit", 6175, DecorValueTy::boolean) -SYCL_COMPILE_TIME_PROPERTY("sycl-register-map", 6176, DecorValueTy::boolean) -SYCL_COMPILE_TIME_PROPERTY("sycl-awidth", 6177, DecorValueTy::uint32) -SYCL_COMPILE_TIME_PROPERTY("sycl-dwidth", 6178, DecorValueTy::uint32) -SYCL_COMPILE_TIME_PROPERTY("sycl-latency", 6179, DecorValueTy::uint32) -SYCL_COMPILE_TIME_PROPERTY("sycl-read-write-mode", 6180, DecorValueTy::uint32) -SYCL_COMPILE_TIME_PROPERTY("sycl-maxburst", 6181, DecorValueTy::uint32) -SYCL_COMPILE_TIME_PROPERTY("sycl-alignment", 6182, DecorValueTy::uint32) -SYCL_COMPILE_TIME_PROPERTY("sycl-wait-request", 6183, DecorValueTy::uint32) -SYCL_COMPILE_TIME_PROPERTY("sycl-stable", 6184, DecorValueTy::boolean) -SYCL_COMPILE_TIME_PROPERTY("sycl-strict", 19, DecorValueTy::boolean) diff --git a/llvm/tools/sycl-post-link/CompileTimePropertiesPass.cpp b/llvm/tools/sycl-post-link/CompileTimePropertiesPass.cpp index 5f7c44c18c5f9..8afdfd899f320 100644 --- a/llvm/tools/sycl-post-link/CompileTimePropertiesPass.cpp +++ b/llvm/tools/sycl-post-link/CompileTimePropertiesPass.cpp @@ -26,7 +26,6 @@ namespace { constexpr StringRef SYCL_HOST_ACCESS_ATTR = "sycl-host-access"; constexpr StringRef SPIRV_DECOR_MD_KIND = "spirv.Decorations"; -constexpr StringRef SPIRV_PARAM_DECOR_MD_KIND = "spirv.ParameterDecorations"; // The corresponding SPIR-V OpCode for the host_access property is documented // in the SPV_INTEL_global_variable_decorations design document: // https://github.com/intel/llvm/blob/sycl/sycl/doc/extensions/DeviceGlobal/SPV_INTEL_global_variable_decorations.asciidoc#decoration @@ -113,8 +112,8 @@ Optional getGlobalVariableString(const Value *StringV) { } /// Tries to generate a SPIR-V decorate metadata node from an attribute. If -///// the attribute is unknown \c nullptr will be returned. -///// +/// the attribute is unknown \c nullptr will be returned. +/// /// @param Ctx [in] the LLVM context. /// @param Attr [in] the LLVM attribute to generate metadata for. /// @@ -131,13 +130,13 @@ MDNode *attributeToDecorateMetadata(LLVMContext &Ctx, const Attribute &Attr) { Decor DecorFound = DecorIt->second; uint32_t DecorCode = DecorFound.Code; switch (DecorFound.Type) { - case DecorValueTy::uint32: - return buildSpirvDecorMetadata(Ctx, DecorCode, - getAttributeAsInteger(Attr)); - case DecorValueTy::boolean: - return buildSpirvDecorMetadata(Ctx, DecorCode, hasProperty(Attr)); - default: - llvm_unreachable("Unhandled decorator type."); + case DecorValueTy::uint32: + return buildSpirvDecorMetadata(Ctx, DecorCode, + getAttributeAsInteger(Attr)); + case DecorValueTy::boolean: + return buildSpirvDecorMetadata(Ctx, DecorCode, hasProperty(Attr)); + default: + llvm_unreachable("Unhandled decorator type."); } } @@ -184,8 +183,8 @@ attributeToExecModeMetadata(Module &M, const Attribute &Attr) { // Get the integers from the strings. SmallVector MDVals; for (StringRef ValStr : ValStrs) - MDVals.push_back(ConstantAsMetadata::get( - Constant::getIntegerValue(SizeTTy, APInt(SizeTBitSize, ValStr, 10)))); + MDVals.push_back(ConstantAsMetadata::get(Constant::getIntegerValue( + SizeTTy, APInt(SizeTBitSize, ValStr, 10)))); const char *MDName = (AttrKindStr == "sycl-work-group-size") ? "reqd_work_group_size" @@ -212,7 +211,6 @@ PreservedAnalyses CompileTimePropertiesPass::run(Module &M, ModuleAnalysisManager &MAM) { LLVMContext &Ctx = M.getContext(); unsigned MDKindID = Ctx.getMDKindID(SPIRV_DECOR_MD_KIND); - unsigned MDParamKindID = Ctx.getMDKindID(SPIRV_PARAM_DECOR_MD_KIND); bool CompileTimePropertiesMet = false; // Let's process all the globals @@ -250,55 +248,33 @@ PreservedAnalyses CompileTimePropertiesPass::run(Module &M, } } + // Process all properties on kernels. for (Function &F : M) { // Only consider kernels. if (F.getCallingConv() != CallingConv::SPIR_KERNEL) continue; - { - // Process all properties on kernels arugments - SmallVector MDOps; - for (unsigned i = 0; i < F.arg_size(); i++) { - SmallVector MDArgOps; - for (auto &Attribute : F.getAttributes().getParamAttrs(i)) { - if (MDNode *SPIRVMetadata = - attributeToDecorateMetadata(Ctx, Attribute)) - MDArgOps.push_back(SPIRVMetadata); - } - MDOps.push_back(MDNode::get(Ctx, MDArgOps)); - } - // Add the generated metadata to the kernel function. - if (!MDOps.empty()) { - F.addMetadata(MDParamKindID, *MDNode::get(Ctx, MDOps)); - CompileTimePropertiesMet = true; - } + SmallVector MDOps; + SmallVector, 8> NamedMDOps; + for (const Attribute &Attribute : F.getAttributes().getFnAttrs()) { + if (MDNode *SPIRVMetadata = attributeToDecorateMetadata(Ctx, Attribute)) + MDOps.push_back(SPIRVMetadata); + else if (auto NamedMetadata = attributeToExecModeMetadata(M, Attribute)) + NamedMDOps.push_back(*NamedMetadata); } - { - // Process all properties on kernels. - SmallVector MDOps; - SmallVector, 8> NamedMDOps; - for (const Attribute &Attribute : F.getAttributes().getFnAttrs()) { - if (MDNode *SPIRVMetadata = attributeToDecorateMetadata(Ctx, Attribute)) - MDOps.push_back(SPIRVMetadata); - else if (auto NamedMetadata = attributeToExecModeMetadata(M, Attribute)) - NamedMDOps.push_back(*NamedMetadata); - } - - // Add the generated metadata to the kernel function. - if (!MDOps.empty()) { - F.addMetadata(MDKindID, *MDNode::get(Ctx, MDOps)); - CompileTimePropertiesMet = true; - } + // Add the generated metadata to the kernel function. + if (!MDOps.empty()) { + F.addMetadata(MDKindID, *MDNode::get(Ctx, MDOps)); + CompileTimePropertiesMet = true; + } - // Add the new named metadata to the kernel function. - for (std::pair NamedMD : NamedMDOps) { - // If multiple sources defined this metadata, prioritize the existing - // one. - if (F.hasMetadata(NamedMD.first)) - continue; - F.addMetadata(NamedMD.first, *NamedMD.second); - } + // Add the new named metadata to the kernel function. + for (std::pair NamedMD : NamedMDOps) { + // If multiple sources defined this metadata, prioritize the existing one. + if (F.hasMetadata(NamedMD.first)) + continue; + F.addMetadata(NamedMD.first, *NamedMD.second); } } From d7c82c9365d8d03be018bf7b587622a86738009c Mon Sep 17 00:00:00 2001 From: "Wang, Di5" Date: Thu, 1 Dec 2022 16:54:51 -0800 Subject: [PATCH 15/26] add extension tests --- .../annotated_arg_for_non_ptr.cpp | 151 ++++++++++++++++ .../annotated_arg/annotated_arg_for_ptr.cpp | 161 ++++++++++++++++++ 2 files changed, 312 insertions(+) create mode 100644 sycl/test/extensions/annotated_arg/annotated_arg_for_non_ptr.cpp create mode 100644 sycl/test/extensions/annotated_arg/annotated_arg_for_ptr.cpp diff --git a/sycl/test/extensions/annotated_arg/annotated_arg_for_non_ptr.cpp b/sycl/test/extensions/annotated_arg/annotated_arg_for_non_ptr.cpp new file mode 100644 index 0000000000000..9373f78fb3b29 --- /dev/null +++ b/sycl/test/extensions/annotated_arg/annotated_arg_for_non_ptr.cpp @@ -0,0 +1,151 @@ +// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple -fsyntax-only -Xclang -verify -Xclang -verify-ignore-unexpected=note %s +// expected-no-diagnostics + +#include "sycl/sycl.hpp" +#include + +#include + +// clang-format on + +using namespace sycl; +using namespace ext::oneapi::experimental; + +using annotated_arg_t1 = annotated_arg; + +using annotated_arg_t3 = annotated_arg; + + +struct B {}; + +struct A : public B { + int x; + A() {} + A(int x_) : x(x_) {} + + const int& operator[](std::ptrdiff_t idx) const { + return x; + } +}; + +struct MyIP { + int *a; + + annotated_arg_t1 b; + + MyIP(int *a_, const A& b_) : a(a_), b(b_) {} + + void operator()() const { + const A& tmp = b; + + for (int i = 0; i < tmp.x; i++) { + *a += 1; + } + *a += b[0]; + } +}; + +template +T foo() { + A obj(5); + return annotated_arg(obj, conduit); +} + +void TestVectorAddWithAnnotatedMMHosts() { +#ifdef FPGA_EMULATOR + using testconfig_selector = sycl::ext::intel::fpga_emulator_selector; +#elif FPGA_SIMULATOR + using testconfig_selector = sycl::ext::intel::fpga_simulator_selector; +#else + using testconfig_selector = sycl::ext::intel::fpga_selector; +#endif + + // Create the SYCL device queue + queue q(testconfig_selector{}); + auto raw = malloc_shared(1, q); + + + A obj(0); + // default ctor + annotated_arg_t3 a1(obj); + // copy ctor + auto a2(a1); + auto a3(foo()); + // // assign ctor + auto a4 = a3; + + // Construct from A instance + auto tmp11 = annotated_arg(obj); // empty property list + + // Construct from A instance and a property list + // auto tmp12 = annotated_arg(obj, properties{conduit}); + auto tmp12 = annotated_arg(obj, conduit); + static_assert(std::is_same::value, "deduction guide failed 1"); + + // Construct from A instance and variadic properties + auto tmp13 = annotated_arg(obj, conduit, stable); // deduction guide + static_assert(std::is_same::value, "deduction guide failed 2"); + + // property list contains invalid property + // auto tmp14 = annotated_arg(obj, awidth<32>); // ERR + + // Construct from another annotated_arg + // templated copy constructor + annotated_arg arg11(tmp11); + annotated_arg arg14(tmp11); // convertible type + auto arg12 = annotated_arg(tmp11); + + // default copy constructor + auto arg13 = annotated_arg(tmp12); + static_assert(std::is_same::value, "deduction guide failed 3"); + + // Construct from another annotated_arg and a property list + // annotated_arg arg21(tmp11, properties{stable}); // ERR: the type properties should be the union of the inputs + annotated_arg arg22(tmp12, properties{stable}); + auto arg23 = annotated_arg(tmp12, properties{stable}); // deduction guide + static_assert(std::is_same::value, "deduction guide failed 4"); + static_assert(std::is_same::value, "deduction guide failed 5"); + annotated_arg arg24(tmp12, properties{stable}); // convertible type + + // Property merge + auto arg31 = annotated_arg_t3(obj, conduit); // OK + auto arg32 = annotated_arg(arg31, properties{stable}); // OK + auto arg33 = annotated_arg(arg32, properties{conduit, stable}); // OK + // auto arg34 = annotated_arg(arg32, properties{conduit, latency<22>}); // ERR: invalid property + static_assert(std::is_same::value, "deduction guide failed 6"); + static_assert(std::is_same::value, "deduction guide failed 7"); + // auto arg35 = annotated_arg(arg32, properties{conduit, dwidth<22>}); // ERR: two input property lists are conflict + // annotated_arg arg36(arg31, properties{latency<32>, stable}); // ERR: input property list is conflict with the declared type + + // Implicit Conversion + const A& x13 = arg32; // OK + // A& x11 = arg32; // ERR: non-const lvalue reference to type 'A' cannot bind to a value of unrelated type + + + // has/get property + static_assert(annotated_arg_t1::has_property(), "has property 1"); + static_assert(annotated_arg_t1::get_property() == conduit, "get property 1"); + static_assert(annotated_arg_t1::has_property() == false, "has property 2"); + + static_assert(annotated_arg_t3::has_property() == false, "has property 3"); + // auto stable_prop = annotated_arg_t3::get_property(); // ERR: can't get non-existing property + + + *raw = 0; + q.submit([&](handler &h) { + h.single_task(MyIP{raw, A(5)}); + }).wait(); + + std::cout << raw[0] << std::endl; + free(raw, q); +} + +int main() { + TestVectorAddWithAnnotatedMMHosts(); + return 0; +} diff --git a/sycl/test/extensions/annotated_arg/annotated_arg_for_ptr.cpp b/sycl/test/extensions/annotated_arg/annotated_arg_for_ptr.cpp new file mode 100644 index 0000000000000..f445fd5cee017 --- /dev/null +++ b/sycl/test/extensions/annotated_arg/annotated_arg_for_ptr.cpp @@ -0,0 +1,161 @@ +// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple -fsyntax-only -Xclang -verify -Xclang -verify-ignore-unexpected=note %s +// expected-no-diagnostics + +#include "sycl/sycl.hpp" +#include + +#include + +// clang-format on + +using namespace sycl; +using namespace ext::oneapi::experimental; + +using annotated_arg_t1 = annotated_arg, + dwidth<32> + ))>; + +using annotated_arg_t2 = annotated_arg; + +using annotated_arg_t3 = annotated_arg + ))>; + +struct MyIP { + annotated_arg, + dwidth<32> + ))> a; + + int b; + + MyIP(int *a_, int b_) : a(a_), b(b_) {} + + void operator()() const { + int *p = a; + + for (int i = 0; i < b; i++) { + p[i] = i; + a[i] += 1; + } + } +}; + +template +T foo() { + auto raw = new int; + return annotated_arg(raw, awidth<32>); +} + +void TestVectorAddWithAnnotatedMMHosts() { +#ifdef FPGA_EMULATOR + using testconfig_selector = sycl::ext::intel::fpga_emulator_selector; +#elif FPGA_SIMULATOR + using testconfig_selector = sycl::ext::intel::fpga_simulator_selector; +#else + using testconfig_selector = sycl::ext::intel::fpga_selector; +#endif + + // Create the SYCL device queue + queue q(testconfig_selector{}); + auto raw = malloc_shared(5, q); + for (int i = 0; i < 5; i++) { + raw[i] = 0; + } + + + // default ctor + annotated_arg_t3 a1; + // copy ctor + auto a2(a1); + auto a3(foo()); + // // assign ctor + auto a4 = a1; + + // Construct from raw pointers + auto tmp11 = annotated_arg(raw); // empty property list + // Construct from raw pointers and a property list + auto tmp12 = annotated_arg})>(raw, properties{awidth<32>}); + auto tmp14 = annotated_arg(raw, properties{awidth<32>}); // deduction guide + static_assert(std::is_same::value, "deduction guide failed 1"); + // Construct from raw pointers and variadic properties + auto tmp13 = annotated_arg(raw, awidth<32>, dwidth<32>); // deduction guide + static_assert(std::is_same::value, "deduction guide failed 2"); + auto tmp15 = annotated_arg(raw, awidth<32>); + static_assert(std::is_same::value, "deduction guide failed 1"); + + // Property list can't have duplicated properties + // auto tmp16 = annotated_arg(raw, awidth<32>, awidth<32>); // ERR + // auto tmp17 = annotated_arg(raw, awidth<32>, awidth<22>); // ERR + + // auto tmp18 = annotated_arg(raw, properties{awidth<32>, dwidth<32>, awidth<32>}); // ERR: Duplicate properties in property list + // auto tmp19 = annotated_arg(raw, properties{awidth<32>, awidth<22>}); // ERR + + + // Construct from another annotated_arg + // templated copy constructor + annotated_arg, dwidth<32>})> arg11(tmp11); + auto arg12 = annotated_arg, dwidth<32>})>(tmp11); + + // default copy constructor + auto arg13 = annotated_arg(tmp12); + static_assert(std::is_same::value, "deduction guide failed 3"); + + // Construct from another annotated_arg and a property list + // annotated_arg, dwidth<32>})> arg21(tmp11, properties{dwidth<32>}); // ERR: the type properties should be the union of the inputs + annotated_arg, dwidth<32>})> arg22(tmp12, properties{dwidth<32>}); + auto arg23 = annotated_arg(tmp12, properties{dwidth<32>}); // deduction guide + static_assert(std::is_same::value, "deduction guide failed 4"); + static_assert(std::is_same::value, "deduction guide failed 5"); + + // Construct from inconvertible type + // annotated_arg tmp21; + // annotated_arg})> arg24(tmp21, properties{dwidth<32>}); // ERR + + // Property merge + auto arg31 = annotated_arg_t3(raw, awidth<32>); // OK + auto arg32 = annotated_arg(arg31, properties{dwidth<32>}); // OK + auto arg33 = annotated_arg(arg32, properties{awidth<32>, dwidth<32>}); // OK + auto arg34 = annotated_arg(arg32, properties{awidth<32>, latency<22>}); // OK + static_assert(std::is_same::value, "deduction guide failed 6"); + static_assert(std::is_same::value, "deduction guide failed 7"); + // auto arg34 = annotated_arg(arg32, properties{awidth<32>, dwidth<22>}); // ERR: two input property lists are conflict + // annotated_arg, dwidth<32>})> arg35(arg31, properties{latency<32>, dwidth<32>}); // ERR: input property list is conflict with the declared type + + + // Implicit Conversion + int* x11 = arg13; + const int* x13 = arg32; + + // operator() + + + // has/get property + static_assert(annotated_arg_t1::has_property(), "has property 1"); + static_assert(annotated_arg_t1::get_property() == awidth<32>, "get property 1"); + static_assert(annotated_arg_t1::has_property() == false, "has property 2"); + + static_assert(annotated_arg_t3::has_property() == false, "has property 3"); + // auto dwidth_prop = annotated_arg_t3::get_property(); // ERR + + + q.submit([&](handler &h) { + h.single_task(MyIP{raw, 5}); + }).wait(); + + + for (int i = 0; i < 5; i++) { + std::cout << raw[i] << std::endl; + } + + free(raw, q); +} + +int main() { + TestVectorAddWithAnnotatedMMHosts(); + return 0; +} From 863b3552e58622170ca0ea66c15acc63c56ce8de Mon Sep 17 00:00:00 2001 From: "Wang, Di5" Date: Thu, 1 Dec 2022 19:42:08 -0800 Subject: [PATCH 16/26] reformat --- .../oneapi/annotated_arg/annotated_arg.hpp | 220 +++++++++--------- .../ext/oneapi/annotated_arg/properties.hpp | 145 ++++++------ sycl/include/sycl/sycl.hpp | 4 +- .../annotated_arg_for_non_ptr.cpp | 114 ++++----- .../annotated_arg/annotated_arg_for_ptr.cpp | 124 +++++----- .../annotated_arg_properties.cpp | 12 +- 6 files changed, 316 insertions(+), 303 deletions(-) diff --git a/sycl/include/sycl/ext/oneapi/annotated_arg/annotated_arg.hpp b/sycl/include/sycl/ext/oneapi/annotated_arg/annotated_arg.hpp index 020c7f652d2a8..37afd732002c7 100644 --- a/sycl/include/sycl/ext/oneapi/annotated_arg/annotated_arg.hpp +++ b/sycl/include/sycl/ext/oneapi/annotated_arg/annotated_arg.hpp @@ -25,7 +25,6 @@ Op " is not supported on host device."); #endif - #ifdef __SYCL_DEVICE_ONLY__ #define __OPENCL_GLOBAL_AS__ __attribute__((opencl_global)) #ifdef __ENABLE_USM_ADDR_SPACE__ @@ -61,22 +60,30 @@ struct HasSubscriptOperator : std::false_type {}; template struct HasSubscriptOperator< - T, typename std::enable_if_t().operator[](0))>::value>> - : std::true_type {}; + T, typename std::enable_if_t< + !std::is_void().operator[](0))>::value>> + : std::true_type{}; } // namespace detail - template -annotated_arg(T, Args... args) -> annotated_arg, std::is_pointer::value>; +annotated_arg(T, Args... args) + -> annotated_arg, + std::is_pointer::value>; template -annotated_arg(T, properties>) -> annotated_arg, std::is_pointer::value>; +annotated_arg(T, properties>) + -> annotated_arg, + std::is_pointer::value>; template -annotated_arg(annotated_arg, properties>) -> annotated_arg>, IsPtr>; +annotated_arg(annotated_arg, properties>) + -> annotated_arg< + T, detail::merged_properties_t>, + IsPtr>; -template ::value> +template ::value> class annotated_arg { // This should always fail when instantiating the unspecialized version. static_assert(is_property_list::value, @@ -85,90 +92,94 @@ class annotated_arg { // Partial specialization for pointer type template -class __SYCL_SPECIAL_CLASS __SYCL_TYPE(annotated_arg) annotated_arg, true> { +class __SYCL_SPECIAL_CLASS __SYCL_TYPE(annotated_arg) + annotated_arg, true> { using property_list_t = detail::properties_t; using UnderlyingT = typename std::remove_pointer::type; __OPENCL_GLOBAL_AS__ UnderlyingT *obj; - template + template friend class annotated_arg; - #ifdef __SYCL_DEVICE_ONLY__ - void __init( - [[__sycl_detail__::add_ir_attributes_kernel_parameter( - detail::PropertyMetaInfo::name..., - detail::PropertyMetaInfo::value... - )]] - __OPENCL_GLOBAL_AS__ UnderlyingT* _obj) { - obj = _obj; - } - #endif +#ifdef __SYCL_DEVICE_ONLY__ + void __init([[__sycl_detail__::add_ir_attributes_kernel_parameter( + detail::PropertyMetaInfo::name..., + detail::PropertyMetaInfo::value...)]] __OPENCL_GLOBAL_AS__ + UnderlyingT *_obj) { + obj = _obj; + } +#endif public: static_assert(std::is_trivially_destructible::value, "Type T must be trivially destructible."); // static_assert(is_property_list::value, - // "Property list is invalid."); + // "Property list is invalid."); static_assert(check_property_list::value, "The property list contains invalid property."); static_assert(detail::SortedAllUnique>::value, "Duplicate properties in property list."); annotated_arg() noexcept = default; - annotated_arg(const annotated_arg&) = default; - annotated_arg& operator=(annotated_arg&) = default; + annotated_arg(const annotated_arg &) = default; + annotated_arg &operator=(annotated_arg &) = default; - annotated_arg(const T& _ptr, const property_list_t &PropList = properties{}) noexcept - : obj((__OPENCL_GLOBAL_AS__ UnderlyingT*)_ptr) {} + annotated_arg(const T &_ptr, + const property_list_t &PropList = properties{}) noexcept + : obj((__OPENCL_GLOBAL_AS__ UnderlyingT *)_ptr) {} - template - annotated_arg(const T& _ptr, PropertyValueTs... props) noexcept : obj((__OPENCL_GLOBAL_AS__ UnderlyingT*)_ptr) { + template + annotated_arg(const T &_ptr, PropertyValueTs... props) noexcept + : obj((__OPENCL_GLOBAL_AS__ UnderlyingT *)_ptr) { static_assert( - std::is_same< - property_list_t, - detail::merged_properties_t>>::value, - "The property list must contain all properties of the input of the constructor" - ); + std::is_same>>::value, + "The property list must contain all properties of the input of the " + "constructor"); } - + // Constructs an annotated_arg object from another annotated_arg object. - // The property set PropertyListT contains all properties of the input annotated_arg object. - // If there are duplicate properties present in the property list of the input annotated_arg object, - // the values of the duplicate properties must be the same. + // The property set PropertyListT contains all properties of the input + // annotated_arg object. If there are duplicate properties present in the + // property list of the input annotated_arg object, the values of the + // duplicate properties must be the same. template - explicit annotated_arg(const annotated_arg &other) noexcept : obj(other.obj) { - static_assert(std::is_convertible::value, - "The underlying data type of the input annotated_arg is not compatible"); + explicit annotated_arg(const annotated_arg &other) noexcept + : obj(other.obj) { + static_assert(std::is_convertible::value, + "The underlying data type of the input annotated_arg is not " + "compatible"); static_assert( std::is_same< property_list_t, detail::merged_properties_t>::value, - "The constructed annotated_arg type must contain all the properties of the input annotated_arg"); + "The constructed annotated_arg type must contain all the properties of " + "the input annotated_arg"); } template explicit annotated_arg(const annotated_arg &other, - const PropertyListV& proplist) noexcept : obj(other.obj) { - static_assert(std::is_convertible::value, - "The underlying data type of the input annotated_arg is not compatible"); + const PropertyListV &proplist) noexcept + : obj(other.obj) { + static_assert(std::is_convertible::value, + "The underlying data type of the input annotated_arg is not " + "compatible"); static_assert( - std::is_same< - property_list_t, - detail::merged_properties_t>::value, - "The property list of constructed annotated_arg type must be the union of the input property lists"); + std::is_same>::value, + "The property list of constructed annotated_arg type must be the union " + "of the input property lists"); } - operator T() noexcept { - return obj; - } + operator T() noexcept { return obj; } - operator const T() const noexcept { - return obj; - } + operator const T() const noexcept { return obj; } - UnderlyingT& operator [](std::ptrdiff_t idx) const noexcept { + UnderlyingT &operator[](std::ptrdiff_t idx) const noexcept { return obj[idx]; } @@ -181,27 +192,24 @@ class __SYCL_SPECIAL_CLASS __SYCL_TYPE(annotated_arg) annotated_arg -class __SYCL_SPECIAL_CLASS __SYCL_TYPE(annotated_arg) annotated_arg , false> { +class __SYCL_SPECIAL_CLASS __SYCL_TYPE(annotated_arg) + annotated_arg, false> { using property_list_t = detail::properties_t; - template + template friend class annotated_arg; T obj; - #ifdef __SYCL_DEVICE_ONLY__ - void __init( - [[__sycl_detail__::add_ir_attributes_kernel_parameter( - detail::PropertyMetaInfo::name..., - detail::PropertyMetaInfo::value... - )]] - T _obj) { - obj = _obj; - } - #endif +#ifdef __SYCL_DEVICE_ONLY__ + void __init([[__sycl_detail__::add_ir_attributes_kernel_parameter( + detail::PropertyMetaInfo::name..., + detail::PropertyMetaInfo::value...)]] T _obj) { + obj = _obj; + } +#endif public: // T should be trivially copy constructible to be device copyable @@ -217,70 +225,73 @@ class __SYCL_SPECIAL_CLASS __SYCL_TYPE(annotated_arg) annotated_arg - annotated_arg(const T& _obj, PropertyValueTs... props) noexcept : obj(_obj) { + template + annotated_arg(const T &_obj, PropertyValueTs... props) noexcept : obj(_obj) { static_assert( - std::is_same< - property_list_t, - detail::merged_properties_t>>::value, - "The property list must contain all properties of the input of the constructor" - ); + std::is_same>>::value, + "The property list must contain all properties of the input of the " + "constructor"); } // Constructs an annotated_arg object from another annotated_arg object. - // The property set PropertyListT contains all properties of the input annotated_arg object. - // If there are duplicate properties present in the property list of the input annotated_arg object, - // the values of the duplicate properties must be the same. + // The property set PropertyListT contains all properties of the input + // annotated_arg object. If there are duplicate properties present in the + // property list of the input annotated_arg object, the values of the + // duplicate properties must be the same. template - explicit annotated_arg(const annotated_arg &other) noexcept : obj(other.obj) { + explicit annotated_arg(const annotated_arg &other) noexcept + : obj(other.obj) { static_assert(std::is_convertible::value, - "The underlying data type of the input annotated_arg is not compatible"); + "The underlying data type of the input annotated_arg is not " + "compatible"); static_assert( std::is_same< property_list_t, detail::merged_properties_t>::value, - "The constructed annotated_arg type must contain all the properties of the input annotated_arg"); + "The constructed annotated_arg type must contain all the properties of " + "the input annotated_arg"); } template explicit annotated_arg(const annotated_arg &other, - const PropertyListV& proplist) noexcept : obj(other.obj) { + const PropertyListV &proplist) noexcept + : obj(other.obj) { static_assert(std::is_convertible::value, - "The underlying data type of the input annotated_arg is not compatible"); + "The underlying data type of the input annotated_arg is not " + "compatible"); - static_assert( - std::is_same< - property_list_t, - detail::merged_properties_t>::value, - "The property list of constructed annotated_arg type must be the union of the input property lists"); + static_assert( + std::is_same>::value, + "The property list of constructed annotated_arg type must be the union " + "of the input property lists"); } - operator T() noexcept { - return obj; - } - operator const T() const noexcept { - return obj; - } + operator T() noexcept { return obj; } + operator const T() const noexcept { return obj; } template - std::enable_if_t< - detail::HasSubscriptOperator::value, - const decltype(std::declval().operator[](0))> - &operator[](std::ptrdiff_t idx) const noexcept { + std::enable_if_t::value, + const decltype(std::declval().operator[](0))> & + operator[](std::ptrdiff_t idx) const noexcept { return obj.operator[](idx); } template - std::enable_if_t< - detail::HasSubscriptOperator::value, - decltype(std::declval().operator[](0))> - &operator[](std::ptrdiff_t idx) noexcept { + std::enable_if_t::value, + decltype(std::declval().operator[](0))> & + operator[](std::ptrdiff_t idx) noexcept { return obj.operator[](idx); } @@ -293,7 +304,6 @@ class __SYCL_SPECIAL_CLASS __SYCL_TYPE(annotated_arg) annotated_arg - using value_t = property_value>; + template + using value_t = + property_value>; }; struct awidth_key { - template + template using value_t = property_value>; }; struct dwidth_key { - template + template using value_t = property_value>; }; struct latency_key { - template + template using value_t = property_value>; }; -enum class read_write_mode_enum : std::uint16_t { - read_write, - read, - write -}; +enum class read_write_mode_enum : std::uint16_t { read_write, read, write }; struct read_write_mode_key { - template - using value_t = property_value< - read_write_mode_key, std::integral_constant>; + template + using value_t = + property_value>; }; struct maxburst_key { - template + template using value_t = property_value>; }; struct wait_request_key { - template - using value_t = property_value>; + template + using value_t = + property_value>; }; - #if __cplusplus >= 201703L // inline variables // non-mmhost properties inline constexpr register_map_key::value_t register_map; inline constexpr conduit_key::value_t conduit; inline constexpr stable_key::value_t stable; -// mmhost properties -template +// mmhost properties +template inline constexpr buffer_location_key::value_t buffer_location; -template -inline constexpr awidth_key::value_t awidth; -template -inline constexpr dwidth_key::value_t dwidth; -template -inline constexpr latency_key::value_t latency; -template -inline constexpr maxburst_key::value_t maxburst; -template +template inline constexpr awidth_key::value_t awidth; +template inline constexpr dwidth_key::value_t dwidth; +template inline constexpr latency_key::value_t latency; +template inline constexpr maxburst_key::value_t maxburst; +template inline constexpr wait_request_key::value_t wait_request; inline constexpr wait_request_key::value_t<1> wait_request_requested; inline constexpr wait_request_key::value_t<0> wait_request_not_requested; -template +template inline constexpr read_write_mode_key::value_t read_write_mode; -inline constexpr read_write_mode_key::value_t read_write_mode_read; -inline constexpr read_write_mode_key::value_t read_write_mode_write; -inline constexpr read_write_mode_key::value_t read_write_mode_readwrite; - +inline constexpr read_write_mode_key::value_t + read_write_mode_read; +inline constexpr read_write_mode_key::value_t + read_write_mode_write; +inline constexpr read_write_mode_key::value_t + read_write_mode_readwrite; #endif // __cplusplus >= 201703L @@ -117,9 +113,9 @@ template <> struct is_property_key : std::true_type {}; template <> struct is_property_key : std::true_type {}; template <> struct is_property_key : std::true_type {}; - template -struct is_property_key_of> +struct is_property_key_of> : std::true_type {}; template @@ -131,7 +127,8 @@ struct is_property_key_of> : std::true_type {}; template -struct is_property_key_of> +struct is_property_key_of> : std::true_type {}; template @@ -147,7 +144,8 @@ struct is_property_key_of> : std::true_type {}; template -struct is_property_key_of> +struct is_property_key_of> : std::true_type {}; template @@ -155,10 +153,10 @@ struct is_property_key_of> : std::true_type {}; template -struct is_property_key_of> +struct is_property_key_of> : std::true_type {}; - namespace detail { template <> struct PropertyToKind { static constexpr PropKind Kind = PropKind::RegisterMap; @@ -195,52 +193,46 @@ template <> struct IsCompileTimeProperty : std::true_type {}; template <> struct IsCompileTimeProperty : std::true_type {}; template <> struct IsCompileTimeProperty : std::true_type {}; -template <> struct IsCompileTimeProperty : std::true_type {}; +template <> +struct IsCompileTimeProperty : std::true_type {}; template <> struct IsCompileTimeProperty : std::true_type {}; template <> struct IsCompileTimeProperty : std::true_type {}; -template <> struct IsCompileTimeProperty : std::true_type {}; +template <> +struct IsCompileTimeProperty : std::true_type {}; template <> struct IsCompileTimeProperty : std::true_type {}; template <> struct IsCompileTimeProperty : std::true_type {}; template <> struct IsCompileTimeProperty : std::true_type {}; -template <> -struct PropertyMetaInfo { +template <> struct PropertyMetaInfo { static constexpr const char *name = "sycl-register-map"; static constexpr std::nullptr_t value = nullptr; }; -template <> -struct PropertyMetaInfo { +template <> struct PropertyMetaInfo { static constexpr const char *name = "sycl-conduit"; static constexpr std::nullptr_t value = nullptr; }; -template <> -struct PropertyMetaInfo { +template <> struct PropertyMetaInfo { static constexpr const char *name = "sycl-stable"; static constexpr std::nullptr_t value = nullptr; }; -template -struct PropertyMetaInfo> { +template struct PropertyMetaInfo> { static constexpr const char *name = "sycl-buffer-location"; static constexpr int value = N; }; -template -struct PropertyMetaInfo> { +template struct PropertyMetaInfo> { static constexpr const char *name = "sycl-awidth"; static constexpr int value = W; }; -template -struct PropertyMetaInfo> { +template struct PropertyMetaInfo> { static constexpr const char *name = "sycl-dwidth"; static constexpr int value = W; }; -template -struct PropertyMetaInfo> { +template struct PropertyMetaInfo> { static constexpr const char *name = "sycl-latency"; static constexpr int value = N; }; -template -struct PropertyMetaInfo> { +template struct PropertyMetaInfo> { static constexpr const char *name = "sycl-maxburst"; static constexpr int value = N; }; @@ -257,58 +249,57 @@ struct PropertyMetaInfo> { } // namespace detail - //===---------------------Properties Verification----------------------===// // 'buffer_location' and mmhost properties are pointers-only -template +template struct is_valid_property : std::false_type {}; -template +template struct is_valid_property> : std::bool_constant::value> {}; -template +template struct is_valid_property> : std::bool_constant::value> {}; -template +template struct is_valid_property> : std::bool_constant::value> {}; -template +template struct is_valid_property> : std::bool_constant::value> {}; -template +template struct is_valid_property> : std::bool_constant::value> {}; -template +template struct is_valid_property> : std::bool_constant::value> {}; -template +template struct is_valid_property> : std::bool_constant::value> {}; -// 'register_map', 'conduit', 'stable' are common properties for pointers -// and non pointers; -template +// 'register_map', 'conduit', 'stable' are common properties for pointers +// and non pointers; +template struct is_valid_property : std::true_type {}; -template +template struct is_valid_property : std::true_type {}; -template +template struct is_valid_property : std::true_type {}; -template +template struct check_property_list : std::true_type {}; -template -struct check_property_list +template +struct check_property_list : std::conditional_t::value, - check_property_list, - std::false_type> { - static_assert(is_valid_property::value, "Property is invalid for the given type."); + check_property_list, std::false_type> { + static_assert(is_valid_property::value, + "Property is invalid for the given type."); }; } // namespace experimental diff --git a/sycl/include/sycl/sycl.hpp b/sycl/include/sycl/sycl.hpp index a6d92fa81f0a9..2f9f46a54b81b 100644 --- a/sycl/include/sycl/sycl.hpp +++ b/sycl/include/sycl/sycl.hpp @@ -60,10 +60,10 @@ #include #endif #include -#include -#include #include #include +#include +#include #include #include #include diff --git a/sycl/test/extensions/annotated_arg/annotated_arg_for_non_ptr.cpp b/sycl/test/extensions/annotated_arg/annotated_arg_for_non_ptr.cpp index 9373f78fb3b29..71bdb6171fda3 100644 --- a/sycl/test/extensions/annotated_arg/annotated_arg_for_non_ptr.cpp +++ b/sycl/test/extensions/annotated_arg/annotated_arg_for_non_ptr.cpp @@ -1,7 +1,7 @@ // RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple -fsyntax-only -Xclang -verify -Xclang -verify-ignore-unexpected=note %s // expected-no-diagnostics -#include "sycl/sycl.hpp" +#include "sycl/sycl.hpp" #include #include @@ -11,15 +11,10 @@ using namespace sycl; using namespace ext::oneapi::experimental; -using annotated_arg_t1 = annotated_arg; - -using annotated_arg_t3 = annotated_arg; +using annotated_arg_t1 = + annotated_arg; +using annotated_arg_t3 = annotated_arg; struct B {}; @@ -28,9 +23,7 @@ struct A : public B { A() {} A(int x_) : x(x_) {} - const int& operator[](std::ptrdiff_t idx) const { - return x; - } + const int &operator[](std::ptrdiff_t idx) const { return x; } }; struct MyIP { @@ -38,10 +31,10 @@ struct MyIP { annotated_arg_t1 b; - MyIP(int *a_, const A& b_) : a(a_), b(b_) {} + MyIP(int *a_, const A &b_) : a(a_), b(b_) {} void operator()() const { - const A& tmp = b; + const A &tmp = b; for (int i = 0; i < tmp.x; i++) { *a += 1; @@ -50,8 +43,7 @@ struct MyIP { } }; -template -T foo() { +template T foo() { A obj(5); return annotated_arg(obj, conduit); } @@ -66,10 +58,9 @@ void TestVectorAddWithAnnotatedMMHosts() { #endif // Create the SYCL device queue - queue q(testconfig_selector{}); + queue q(testconfig_selector{}); auto raw = malloc_shared(1, q); - A obj(0); // default ctor annotated_arg_t3 a1(obj); @@ -80,66 +71,83 @@ void TestVectorAddWithAnnotatedMMHosts() { auto a4 = a3; // Construct from A instance - auto tmp11 = annotated_arg(obj); // empty property list + auto tmp11 = annotated_arg(obj); // empty property list // Construct from A instance and a property list // auto tmp12 = annotated_arg(obj, properties{conduit}); auto tmp12 = annotated_arg(obj, conduit); - static_assert(std::is_same::value, "deduction guide failed 1"); + static_assert(std::is_same::value, + "deduction guide failed 1"); // Construct from A instance and variadic properties - auto tmp13 = annotated_arg(obj, conduit, stable); // deduction guide - static_assert(std::is_same::value, "deduction guide failed 2"); + auto tmp13 = annotated_arg(obj, conduit, stable); // deduction guide + static_assert(std::is_same::value, + "deduction guide failed 2"); // property list contains invalid property // auto tmp14 = annotated_arg(obj, awidth<32>); // ERR // Construct from another annotated_arg // templated copy constructor - annotated_arg arg11(tmp11); - annotated_arg arg14(tmp11); // convertible type + annotated_arg arg11(tmp11); + annotated_arg arg14( + tmp11); // convertible type auto arg12 = annotated_arg(tmp11); - + // default copy constructor - auto arg13 = annotated_arg(tmp12); - static_assert(std::is_same::value, "deduction guide failed 3"); + auto arg13 = annotated_arg(tmp12); + static_assert(std::is_same::value, + "deduction guide failed 3"); // Construct from another annotated_arg and a property list - // annotated_arg arg21(tmp11, properties{stable}); // ERR: the type properties should be the union of the inputs - annotated_arg arg22(tmp12, properties{stable}); + // annotated_arg arg21(tmp11, + // properties{stable}); // ERR: the type properties should be the union of + // the inputs + annotated_arg arg22( + tmp12, properties{stable}); auto arg23 = annotated_arg(tmp12, properties{stable}); // deduction guide - static_assert(std::is_same::value, "deduction guide failed 4"); - static_assert(std::is_same::value, "deduction guide failed 5"); - annotated_arg arg24(tmp12, properties{stable}); // convertible type + static_assert(std::is_same::value, + "deduction guide failed 4"); + static_assert(std::is_same::value, + "deduction guide failed 5"); + annotated_arg arg24( + tmp12, properties{stable}); // convertible type // Property merge - auto arg31 = annotated_arg_t3(obj, conduit); // OK - auto arg32 = annotated_arg(arg31, properties{stable}); // OK - auto arg33 = annotated_arg(arg32, properties{conduit, stable}); // OK - // auto arg34 = annotated_arg(arg32, properties{conduit, latency<22>}); // ERR: invalid property - static_assert(std::is_same::value, "deduction guide failed 6"); - static_assert(std::is_same::value, "deduction guide failed 7"); - // auto arg35 = annotated_arg(arg32, properties{conduit, dwidth<22>}); // ERR: two input property lists are conflict - // annotated_arg arg36(arg31, properties{latency<32>, stable}); // ERR: input property list is conflict with the declared type + auto arg31 = annotated_arg_t3(obj, conduit); // OK + auto arg32 = annotated_arg(arg31, properties{stable}); // OK + auto arg33 = annotated_arg(arg32, properties{conduit, stable}); // OK + // auto arg34 = annotated_arg(arg32, properties{conduit, latency<22>}); // + // ERR: invalid property + static_assert(std::is_same::value, + "deduction guide failed 6"); + static_assert(std::is_same::value, + "deduction guide failed 7"); + // auto arg35 = annotated_arg(arg32, properties{conduit, dwidth<22>}); // + // ERR: two input property lists are conflict annotated_arg arg36(arg31, properties{latency<32>, + // stable}); // ERR: input property list is conflict with the declared type // Implicit Conversion - const A& x13 = arg32; // OK - // A& x11 = arg32; // ERR: non-const lvalue reference to type 'A' cannot bind to a value of unrelated type - + const A &x13 = arg32; // OK + // A& x11 = arg32; // ERR: non-const lvalue reference to type 'A' cannot + // bind to a value of unrelated type // has/get property - static_assert(annotated_arg_t1::has_property(), "has property 1"); - static_assert(annotated_arg_t1::get_property() == conduit, "get property 1"); - static_assert(annotated_arg_t1::has_property() == false, "has property 2"); - - static_assert(annotated_arg_t3::has_property() == false, "has property 3"); - // auto stable_prop = annotated_arg_t3::get_property(); // ERR: can't get non-existing property - + static_assert(annotated_arg_t1::has_property(), + "has property 1"); + static_assert(annotated_arg_t1::get_property() == conduit, + "get property 1"); + static_assert(annotated_arg_t1::has_property() == false, + "has property 2"); + + static_assert(annotated_arg_t3::has_property() == false, + "has property 3"); + // auto stable_prop = annotated_arg_t3::get_property(); // ERR: + // can't get non-existing property *raw = 0; - q.submit([&](handler &h) { - h.single_task(MyIP{raw, A(5)}); - }).wait(); + q.submit([&](handler &h) { h.single_task(MyIP{raw, A(5)}); }).wait(); std::cout << raw[0] << std::endl; free(raw, q); diff --git a/sycl/test/extensions/annotated_arg/annotated_arg_for_ptr.cpp b/sycl/test/extensions/annotated_arg/annotated_arg_for_ptr.cpp index f445fd5cee017..df9260c6998d9 100644 --- a/sycl/test/extensions/annotated_arg/annotated_arg_for_ptr.cpp +++ b/sycl/test/extensions/annotated_arg/annotated_arg_for_ptr.cpp @@ -1,7 +1,7 @@ // RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple -fsyntax-only -Xclang -verify -Xclang -verify-ignore-unexpected=note %s // expected-no-diagnostics -#include "sycl/sycl.hpp" +#include "sycl/sycl.hpp" #include #include @@ -11,28 +11,19 @@ using namespace sycl; using namespace ext::oneapi::experimental; -using annotated_arg_t1 = annotated_arg, - dwidth<32> - ))>; +using annotated_arg_t1 = + annotated_arg, dwidth<32>))>; -using annotated_arg_t2 = annotated_arg; +using annotated_arg_t2 = + annotated_arg; -using annotated_arg_t3 = annotated_arg - ))>; +using annotated_arg_t3 = annotated_arg))>; struct MyIP { - annotated_arg, - dwidth<32> - ))> a; + annotated_arg, dwidth<32>))> a; int b; - + MyIP(int *a_, int b_) : a(a_), b(b_) {} void operator()() const { @@ -45,8 +36,7 @@ struct MyIP { } }; -template -T foo() { +template T foo() { auto raw = new int; return annotated_arg(raw, awidth<32>); } @@ -61,13 +51,12 @@ void TestVectorAddWithAnnotatedMMHosts() { #endif // Create the SYCL device queue - queue q(testconfig_selector{}); + queue q(testconfig_selector{}); auto raw = malloc_shared(5, q); for (int i = 0; i < 5; i++) { raw[i] = 0; } - // default ctor annotated_arg_t3 a1; // copy ctor @@ -77,76 +66,91 @@ void TestVectorAddWithAnnotatedMMHosts() { auto a4 = a1; // Construct from raw pointers - auto tmp11 = annotated_arg(raw); // empty property list + auto tmp11 = annotated_arg(raw); // empty property list // Construct from raw pointers and a property list - auto tmp12 = annotated_arg})>(raw, properties{awidth<32>}); - auto tmp14 = annotated_arg(raw, properties{awidth<32>}); // deduction guide - static_assert(std::is_same::value, "deduction guide failed 1"); + auto tmp12 = annotated_arg})>( + raw, properties{awidth<32>}); + auto tmp14 = annotated_arg(raw, properties{awidth<32>}); // deduction guide + static_assert(std::is_same::value, + "deduction guide failed 1"); // Construct from raw pointers and variadic properties - auto tmp13 = annotated_arg(raw, awidth<32>, dwidth<32>); // deduction guide - static_assert(std::is_same::value, "deduction guide failed 2"); + auto tmp13 = annotated_arg(raw, awidth<32>, dwidth<32>); // deduction guide + static_assert(std::is_same::value, + "deduction guide failed 2"); auto tmp15 = annotated_arg(raw, awidth<32>); - static_assert(std::is_same::value, "deduction guide failed 1"); + static_assert(std::is_same::value, + "deduction guide failed 1"); // Property list can't have duplicated properties // auto tmp16 = annotated_arg(raw, awidth<32>, awidth<32>); // ERR // auto tmp17 = annotated_arg(raw, awidth<32>, awidth<22>); // ERR - // auto tmp18 = annotated_arg(raw, properties{awidth<32>, dwidth<32>, awidth<32>}); // ERR: Duplicate properties in property list - // auto tmp19 = annotated_arg(raw, properties{awidth<32>, awidth<22>}); // ERR - + // auto tmp18 = annotated_arg(raw, properties{awidth<32>, dwidth<32>, + // awidth<32>}); // ERR: Duplicate properties in property list auto tmp19 = + // annotated_arg(raw, properties{awidth<32>, awidth<22>}); // ERR // Construct from another annotated_arg // templated copy constructor - annotated_arg, dwidth<32>})> arg11(tmp11); - auto arg12 = annotated_arg, dwidth<32>})>(tmp11); - + annotated_arg, dwidth<32>})> arg11( + tmp11); + auto arg12 = + annotated_arg, dwidth<32>})>(tmp11); + // default copy constructor - auto arg13 = annotated_arg(tmp12); - static_assert(std::is_same::value, "deduction guide failed 3"); + auto arg13 = annotated_arg(tmp12); + static_assert(std::is_same::value, + "deduction guide failed 3"); // Construct from another annotated_arg and a property list - // annotated_arg, dwidth<32>})> arg21(tmp11, properties{dwidth<32>}); // ERR: the type properties should be the union of the inputs - annotated_arg, dwidth<32>})> arg22(tmp12, properties{dwidth<32>}); + // annotated_arg, dwidth<32>})> + // arg21(tmp11, properties{dwidth<32>}); // ERR: the type properties should + // be the union of the inputs + annotated_arg, dwidth<32>})> arg22( + tmp12, properties{dwidth<32>}); auto arg23 = annotated_arg(tmp12, properties{dwidth<32>}); // deduction guide - static_assert(std::is_same::value, "deduction guide failed 4"); - static_assert(std::is_same::value, "deduction guide failed 5"); + static_assert(std::is_same::value, + "deduction guide failed 4"); + static_assert(std::is_same::value, + "deduction guide failed 5"); // Construct from inconvertible type // annotated_arg tmp21; - // annotated_arg})> arg24(tmp21, properties{dwidth<32>}); // ERR + // annotated_arg})> arg24(tmp21, + // properties{dwidth<32>}); // ERR // Property merge - auto arg31 = annotated_arg_t3(raw, awidth<32>); // OK - auto arg32 = annotated_arg(arg31, properties{dwidth<32>}); // OK + auto arg31 = annotated_arg_t3(raw, awidth<32>); // OK + auto arg32 = annotated_arg(arg31, properties{dwidth<32>}); // OK auto arg33 = annotated_arg(arg32, properties{awidth<32>, dwidth<32>}); // OK - auto arg34 = annotated_arg(arg32, properties{awidth<32>, latency<22>}); // OK - static_assert(std::is_same::value, "deduction guide failed 6"); - static_assert(std::is_same::value, "deduction guide failed 7"); - // auto arg34 = annotated_arg(arg32, properties{awidth<32>, dwidth<22>}); // ERR: two input property lists are conflict - // annotated_arg, dwidth<32>})> arg35(arg31, properties{latency<32>, dwidth<32>}); // ERR: input property list is conflict with the declared type - + auto arg34 = annotated_arg(arg32, properties{awidth<32>, latency<22>}); // OK + static_assert(std::is_same::value, + "deduction guide failed 6"); + static_assert(std::is_same::value, + "deduction guide failed 7"); + // auto arg34 = annotated_arg(arg32, properties{awidth<32>, dwidth<22>}); // + // ERR: two input property lists are conflict annotated_arg, dwidth<32>})> arg35(arg31, + // properties{latency<32>, dwidth<32>}); // ERR: input property list is + // conflict with the declared type // Implicit Conversion - int* x11 = arg13; - const int* x13 = arg32; + int *x11 = arg13; + const int *x13 = arg32; // operator() - // has/get property static_assert(annotated_arg_t1::has_property(), "has property 1"); - static_assert(annotated_arg_t1::get_property() == awidth<32>, "get property 1"); - static_assert(annotated_arg_t1::has_property() == false, "has property 2"); + static_assert(annotated_arg_t1::get_property() == awidth<32>, + "get property 1"); + static_assert(annotated_arg_t1::has_property() == false, + "has property 2"); - static_assert(annotated_arg_t3::has_property() == false, "has property 3"); + static_assert(annotated_arg_t3::has_property() == false, + "has property 3"); // auto dwidth_prop = annotated_arg_t3::get_property(); // ERR - - q.submit([&](handler &h) { - h.single_task(MyIP{raw, 5}); - }).wait(); - + q.submit([&](handler &h) { h.single_task(MyIP{raw, 5}); }).wait(); for (int i = 0; i < 5; i++) { std::cout << raw[i] << std::endl; diff --git a/sycl/test/extensions/annotated_arg/annotated_arg_properties.cpp b/sycl/test/extensions/annotated_arg/annotated_arg_properties.cpp index c579f471eab0a..e86e95f6b37ec 100644 --- a/sycl/test/extensions/annotated_arg/annotated_arg_properties.cpp +++ b/sycl/test/extensions/annotated_arg/annotated_arg_properties.cpp @@ -6,12 +6,11 @@ using namespace sycl::ext::oneapi::experimental; static annotated_arg AnnotatedArg1; -static annotated_arg - AnnotatedArg2; -static annotated_arg +static annotated_arg AnnotatedArg2; +static annotated_arg AnnotatedArg3; -static annotated_arg, read_only, - stable, conduit))> +static annotated_arg, read_only, + stable, conduit))> AnnotatedArg4; // Checks is_property_key_of and is_property_value_of for T. @@ -71,7 +70,8 @@ int main() { static_assert(AnnotatedArg4.has_property()); static_assert(AnnotatedArg4.get_property() == conduit); static_assert(AnnotatedArg4.get_property() == stable); - static_assert(AnnotatedArg4.get_property() == buffer_location<1>); + static_assert(AnnotatedArg4.get_property() == + buffer_location<1>); static_assert(AnnotatedArg4.get_property() == read_only); return 0; From 78f0fd15978d865cd47f4eba479ed41ae477031f Mon Sep 17 00:00:00 2001 From: "Wang, Di5" Date: Mon, 5 Dec 2022 07:43:18 -0800 Subject: [PATCH 17/26] code refactor --- .../oneapi/annotated_arg/annotated_arg.hpp | 68 ++++++++----------- .../ext/oneapi/annotated_arg/properties.hpp | 54 +++++++-------- .../annotated_arg_for_non_ptr.cpp | 10 +-- .../annotated_arg_properties.cpp | 34 +++++++--- 4 files changed, 84 insertions(+), 82 deletions(-) diff --git a/sycl/include/sycl/ext/oneapi/annotated_arg/annotated_arg.hpp b/sycl/include/sycl/ext/oneapi/annotated_arg/annotated_arg.hpp index 37afd732002c7..e6731926fa9e1 100644 --- a/sycl/include/sycl/ext/oneapi/annotated_arg/annotated_arg.hpp +++ b/sycl/include/sycl/ext/oneapi/annotated_arg/annotated_arg.hpp @@ -66,24 +66,21 @@ struct HasSubscriptOperator< } // namespace detail +// Deduction guide template annotated_arg(T, Args... args) - -> annotated_arg, - std::is_pointer::value>; + -> annotated_arg>; template annotated_arg(T, properties>) - -> annotated_arg, - std::is_pointer::value>; + -> annotated_arg>; -template -annotated_arg(annotated_arg, properties>) +template +annotated_arg(annotated_arg, properties>) -> annotated_arg< - T, detail::merged_properties_t>, - IsPtr>; + T, detail::merged_properties_t>>; -template ::value> +template class annotated_arg { // This should always fail when instantiating the unspecialized version. static_assert(is_property_list::value, @@ -92,20 +89,18 @@ class annotated_arg { // Partial specialization for pointer type template -class __SYCL_SPECIAL_CLASS __SYCL_TYPE(annotated_arg) - annotated_arg, true> { +class __SYCL_SPECIAL_CLASS +__SYCL_TYPE(annotated_arg) annotated_arg> { using property_list_t = detail::properties_t; - using UnderlyingT = typename std::remove_pointer::type; - __OPENCL_GLOBAL_AS__ UnderlyingT *obj; + __OPENCL_GLOBAL_AS__ T *obj; - template - friend class annotated_arg; + template friend class annotated_arg; #ifdef __SYCL_DEVICE_ONLY__ void __init([[__sycl_detail__::add_ir_attributes_kernel_parameter( detail::PropertyMetaInfo::name..., - detail::PropertyMetaInfo::value...)]] __OPENCL_GLOBAL_AS__ - UnderlyingT *_obj) { + detail::PropertyMetaInfo::value...)]] __OPENCL_GLOBAL_AS__ T + *_obj) { obj = _obj; } #endif @@ -113,10 +108,8 @@ class __SYCL_SPECIAL_CLASS __SYCL_TYPE(annotated_arg) public: static_assert(std::is_trivially_destructible::value, "Type T must be trivially destructible."); - // static_assert(is_property_list::value, - // "Property list is invalid."); - static_assert(check_property_list::value, - "The property list contains invalid property."); + static_assert(is_property_list::value, + "Property list is invalid."); static_assert(detail::SortedAllUnique>::value, "Duplicate properties in property list."); @@ -124,13 +117,13 @@ class __SYCL_SPECIAL_CLASS __SYCL_TYPE(annotated_arg) annotated_arg(const annotated_arg &) = default; annotated_arg &operator=(annotated_arg &) = default; - annotated_arg(const T &_ptr, + annotated_arg(T *_ptr, const property_list_t &PropList = properties{}) noexcept - : obj((__OPENCL_GLOBAL_AS__ UnderlyingT *)_ptr) {} + : obj((__OPENCL_GLOBAL_AS__ T *)_ptr) {} template - annotated_arg(const T &_ptr, PropertyValueTs... props) noexcept - : obj((__OPENCL_GLOBAL_AS__ UnderlyingT *)_ptr) { + annotated_arg(T *_ptr, PropertyValueTs... props) noexcept + : obj((__OPENCL_GLOBAL_AS__ T *)_ptr) { static_assert( std::is_same explicit annotated_arg(const annotated_arg &other) noexcept : obj(other.obj) { - static_assert(std::is_convertible::value, + static_assert(std::is_convertible::value, "The underlying data type of the input annotated_arg is not " "compatible"); @@ -164,7 +157,7 @@ class __SYCL_SPECIAL_CLASS __SYCL_TYPE(annotated_arg) explicit annotated_arg(const annotated_arg &other, const PropertyListV &proplist) noexcept : obj(other.obj) { - static_assert(std::is_convertible::value, + static_assert(std::is_convertible::value, "The underlying data type of the input annotated_arg is not " "compatible"); @@ -175,13 +168,11 @@ class __SYCL_SPECIAL_CLASS __SYCL_TYPE(annotated_arg) "of the input property lists"); } - operator T() noexcept { return obj; } + operator T *() const noexcept { return obj; } - operator const T() const noexcept { return obj; } + operator const T *() const noexcept { return obj; } - UnderlyingT &operator[](std::ptrdiff_t idx) const noexcept { - return obj[idx]; - } + T &operator[](std::ptrdiff_t idx) const noexcept { return obj[idx]; } template static constexpr bool has_property() { return property_list_t::template has_property(); @@ -194,12 +185,11 @@ class __SYCL_SPECIAL_CLASS __SYCL_TYPE(annotated_arg) // Partial specialization for non-pointer type template -class __SYCL_SPECIAL_CLASS __SYCL_TYPE(annotated_arg) - annotated_arg, false> { +class __SYCL_SPECIAL_CLASS +__SYCL_TYPE(annotated_arg) annotated_arg> { using property_list_t = detail::properties_t; - template - friend class annotated_arg; + template friend class annotated_arg; T obj; @@ -212,11 +202,9 @@ class __SYCL_SPECIAL_CLASS __SYCL_TYPE(annotated_arg) #endif public: - // T should be trivially copy constructible to be device copyable + // T should be trivially copyable to be device-copyable static_assert(std::is_trivially_copyable::value, "Type T must be trivially copyable."); - // static_assert(std::is_trivially_copy_constructible::value, - // "Type T must be trivially copy constructable."); static_assert(std::is_trivially_destructible::value, "Type T must be trivially destructible."); static_assert(is_property_list::value, diff --git a/sycl/include/sycl/ext/oneapi/annotated_arg/properties.hpp b/sycl/include/sycl/ext/oneapi/annotated_arg/properties.hpp index 5a218a12b5bee..d7d16b3074039 100644 --- a/sycl/include/sycl/ext/oneapi/annotated_arg/properties.hpp +++ b/sycl/include/sycl/ext/oneapi/annotated_arg/properties.hpp @@ -1,4 +1,4 @@ -//==----- properties.hpp - SYCL properties associated with annotated_arg ---==// +// properties.hpp - SYCL properties associated with annotated_arg/ptr // // // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. // See https://llvm.org/LICENSE.txt for license information. @@ -17,8 +17,11 @@ namespace ext { namespace oneapi { namespace experimental { -template class annotated_arg; +template class annotated_arg; +//===----------------------------------------------------------------------===// +// Common properties of annotated_arg/annotated_ptr +//===----------------------------------------------------------------------===// struct register_map_key { using value_t = property_value; }; @@ -113,48 +116,44 @@ template <> struct is_property_key : std::true_type {}; template <> struct is_property_key : std::true_type {}; template <> struct is_property_key : std::true_type {}; -template -struct is_property_key_of> +template +struct is_property_key_of> : std::true_type {}; -template -struct is_property_key_of> +template +struct is_property_key_of> : std::true_type {}; -template -struct is_property_key_of> +template +struct is_property_key_of> : std::true_type {}; -template -struct is_property_key_of> +template +struct is_property_key_of> : std::true_type {}; -template -struct is_property_key_of> +template +struct is_property_key_of> : std::true_type {}; -template -struct is_property_key_of> +template +struct is_property_key_of> : std::true_type {}; -template -struct is_property_key_of> +template +struct is_property_key_of> : std::true_type {}; -template -struct is_property_key_of> +template +struct is_property_key_of> : std::true_type {}; -template -struct is_property_key_of> +template +struct is_property_key_of> : std::true_type {}; -template -struct is_property_key_of> +template +struct is_property_key_of> : std::true_type {}; namespace detail { @@ -249,8 +248,6 @@ struct PropertyMetaInfo> { } // namespace detail -//===---------------------Properties Verification----------------------===// - // 'buffer_location' and mmhost properties are pointers-only template struct is_valid_property : std::false_type {}; @@ -294,6 +291,7 @@ struct is_valid_property : std::true_type {}; template struct check_property_list : std::true_type {}; + template struct check_property_list : std::conditional_t::value, diff --git a/sycl/test/extensions/annotated_arg/annotated_arg_for_non_ptr.cpp b/sycl/test/extensions/annotated_arg/annotated_arg_for_non_ptr.cpp index 71bdb6171fda3..430d239801a76 100644 --- a/sycl/test/extensions/annotated_arg/annotated_arg_for_non_ptr.cpp +++ b/sycl/test/extensions/annotated_arg/annotated_arg_for_non_ptr.cpp @@ -11,11 +11,6 @@ using namespace sycl; using namespace ext::oneapi::experimental; -using annotated_arg_t1 = - annotated_arg; - -using annotated_arg_t3 = annotated_arg; - struct B {}; struct A : public B { @@ -26,6 +21,11 @@ struct A : public B { const int &operator[](std::ptrdiff_t idx) const { return x; } }; +using annotated_arg_t1 = + annotated_arg; + +using annotated_arg_t3 = annotated_arg; + struct MyIP { int *a; diff --git a/sycl/test/extensions/annotated_arg/annotated_arg_properties.cpp b/sycl/test/extensions/annotated_arg/annotated_arg_properties.cpp index e86e95f6b37ec..e622e67652808 100644 --- a/sycl/test/extensions/annotated_arg/annotated_arg_properties.cpp +++ b/sycl/test/extensions/annotated_arg/annotated_arg_properties.cpp @@ -7,12 +7,15 @@ using namespace sycl::ext::oneapi::experimental; static annotated_arg AnnotatedArg1; static annotated_arg AnnotatedArg2; -static annotated_arg +static annotated_arg AnnotatedArg3; -static annotated_arg, read_only, - stable, conduit))> +static annotated_arg, + read_write_mode_read, stable, + conduit))> AnnotatedArg4; +struct A {}; + // Checks is_property_key_of and is_property_value_of for T. template void checkIsPropertyOf() { static_assert(is_property_key_of::value); @@ -35,9 +38,19 @@ template void checkIsPropertyOf() { static_assert(is_property_value_of), T>::value); static_assert(is_property_value_of), T>::value); static_assert(is_property_value_of), T>::value); - static_assert(is_property_value_of::value); + static_assert(is_property_value_of::value); static_assert(is_property_value_of), T>::value); - static_assert(is_property_value_of), T>::value); + static_assert( + is_property_value_of::value); +} + +// Checks is_property_key_of and is_property_value_of are false for non-pointer +// type T. +template void checkIsValidPropertyOfNonPtr() { + static_assert( + is_valid_property::value == + false); + static_assert(is_valid_property)>::value == false); } int main() { @@ -55,11 +68,11 @@ int main() { static_assert(AnnotatedArg2.get_property() == register_map); checkIsPropertyOf(); - static_assert(AnnotatedArg3.has_property()); + static_assert(!AnnotatedArg3.has_property()); static_assert(AnnotatedArg3.has_property()); - static_assert(!AnnotatedArg3.has_property()); + static_assert(AnnotatedArg3.has_property()); static_assert(!AnnotatedArg3.has_property()); - static_assert(AnnotatedArg3.get_property() == register_map); + static_assert(AnnotatedArg3.get_property() == stable); static_assert(AnnotatedArg3.get_property() == conduit); checkIsPropertyOf(); @@ -72,7 +85,10 @@ int main() { static_assert(AnnotatedArg4.get_property() == stable); static_assert(AnnotatedArg4.get_property() == buffer_location<1>); - static_assert(AnnotatedArg4.get_property() == read_only); + static_assert(AnnotatedArg4.get_property() == + read_write_mode_read); + // Check if a property is valid for a given type + checkIsValidPropertyOfNonPtr(); return 0; } From 8155e5b2f8912677432591b16ae641872671681f Mon Sep 17 00:00:00 2001 From: "Wang, Di5" Date: Mon, 5 Dec 2022 11:07:44 -0800 Subject: [PATCH 18/26] fix warning --- .../annotated_arg/annotated_arg_for_non_ptr.cpp | 10 +--------- .../extensions/annotated_arg/annotated_arg_for_ptr.cpp | 10 +--------- 2 files changed, 2 insertions(+), 18 deletions(-) diff --git a/sycl/test/extensions/annotated_arg/annotated_arg_for_non_ptr.cpp b/sycl/test/extensions/annotated_arg/annotated_arg_for_non_ptr.cpp index 430d239801a76..cb35aae94c933 100644 --- a/sycl/test/extensions/annotated_arg/annotated_arg_for_non_ptr.cpp +++ b/sycl/test/extensions/annotated_arg/annotated_arg_for_non_ptr.cpp @@ -49,16 +49,8 @@ template T foo() { } void TestVectorAddWithAnnotatedMMHosts() { -#ifdef FPGA_EMULATOR - using testconfig_selector = sycl::ext::intel::fpga_emulator_selector; -#elif FPGA_SIMULATOR - using testconfig_selector = sycl::ext::intel::fpga_simulator_selector; -#else - using testconfig_selector = sycl::ext::intel::fpga_selector; -#endif - // Create the SYCL device queue - queue q(testconfig_selector{}); + queue q(sycl::ext::intel::fpga_selector_v); auto raw = malloc_shared(1, q); A obj(0); diff --git a/sycl/test/extensions/annotated_arg/annotated_arg_for_ptr.cpp b/sycl/test/extensions/annotated_arg/annotated_arg_for_ptr.cpp index df9260c6998d9..20a79563ef61e 100644 --- a/sycl/test/extensions/annotated_arg/annotated_arg_for_ptr.cpp +++ b/sycl/test/extensions/annotated_arg/annotated_arg_for_ptr.cpp @@ -42,16 +42,8 @@ template T foo() { } void TestVectorAddWithAnnotatedMMHosts() { -#ifdef FPGA_EMULATOR - using testconfig_selector = sycl::ext::intel::fpga_emulator_selector; -#elif FPGA_SIMULATOR - using testconfig_selector = sycl::ext::intel::fpga_simulator_selector; -#else - using testconfig_selector = sycl::ext::intel::fpga_selector; -#endif - // Create the SYCL device queue - queue q(testconfig_selector{}); + queue q(sycl::ext::intel::fpga_selector_v); auto raw = malloc_shared(5, q); for (int i = 0; i < 5; i++) { raw[i] = 0; From b0a81196dd7eb4cfc5b56dbf2f7a2b01fccc3df3 Mon Sep 17 00:00:00 2001 From: "Wang, Di5" Date: Tue, 6 Dec 2022 08:16:17 -0800 Subject: [PATCH 19/26] resolve comments --- .../oneapi/annotated_arg/annotated_arg.hpp | 31 ++++++++++--------- 1 file changed, 17 insertions(+), 14 deletions(-) diff --git a/sycl/include/sycl/ext/oneapi/annotated_arg/annotated_arg.hpp b/sycl/include/sycl/ext/oneapi/annotated_arg/annotated_arg.hpp index e6731926fa9e1..7547eadf6a1b0 100644 --- a/sycl/include/sycl/ext/oneapi/annotated_arg/annotated_arg.hpp +++ b/sycl/include/sycl/ext/oneapi/annotated_arg/annotated_arg.hpp @@ -55,14 +55,10 @@ namespace experimental { namespace detail { // Type-trait for checking if a type defines `operator[]`. -template -struct HasSubscriptOperator : std::false_type {}; - template -struct HasSubscriptOperator< - T, typename std::enable_if_t< - !std::is_void().operator[](0))>::value>> - : std::true_type{}; +struct HasSubscriptOperator + : std::bool_constant< + !std::is_void().operator[](0))>::value>{}; } // namespace detail @@ -92,15 +88,22 @@ template class __SYCL_SPECIAL_CLASS __SYCL_TYPE(annotated_arg) annotated_arg> { using property_list_t = detail::properties_t; - __OPENCL_GLOBAL_AS__ T *obj; + +#ifdef __SYCL_DEVICE_ONLY__ + using global_pointer_t = typename sycl::detail::DecoratedType< + T, access::address_space::global_space>::type *; +#else + using global_pointer_t = T *; +#endif + + global_pointer_t obj; template friend class annotated_arg; #ifdef __SYCL_DEVICE_ONLY__ void __init([[__sycl_detail__::add_ir_attributes_kernel_parameter( detail::PropertyMetaInfo::name..., - detail::PropertyMetaInfo::value...)]] __OPENCL_GLOBAL_AS__ T - *_obj) { + detail::PropertyMetaInfo::value...)]] global_pointer_t _obj) { obj = _obj; } #endif @@ -110,8 +113,8 @@ __SYCL_TYPE(annotated_arg) annotated_arg> { "Type T must be trivially destructible."); static_assert(is_property_list::value, "Property list is invalid."); - static_assert(detail::SortedAllUnique>::value, - "Duplicate properties in property list."); + // static_assert(detail::SortedAllUnique>::value, + // "Duplicate properties in property list."); annotated_arg() noexcept = default; annotated_arg(const annotated_arg &) = default; @@ -119,11 +122,11 @@ __SYCL_TYPE(annotated_arg) annotated_arg> { annotated_arg(T *_ptr, const property_list_t &PropList = properties{}) noexcept - : obj((__OPENCL_GLOBAL_AS__ T *)_ptr) {} + : obj(sycl::detail::cast_AS(_ptr)) {} template annotated_arg(T *_ptr, PropertyValueTs... props) noexcept - : obj((__OPENCL_GLOBAL_AS__ T *)_ptr) { + : obj(sycl::detail::cast_AS(_ptr)) { static_assert( std::is_same Date: Tue, 6 Dec 2022 08:18:27 -0800 Subject: [PATCH 20/26] code cleanup --- .../oneapi/annotated_arg/annotated_arg.hpp | 21 ------------------- 1 file changed, 21 deletions(-) diff --git a/sycl/include/sycl/ext/oneapi/annotated_arg/annotated_arg.hpp b/sycl/include/sycl/ext/oneapi/annotated_arg/annotated_arg.hpp index 7547eadf6a1b0..39bc098233d08 100644 --- a/sycl/include/sycl/ext/oneapi/annotated_arg/annotated_arg.hpp +++ b/sycl/include/sycl/ext/oneapi/annotated_arg/annotated_arg.hpp @@ -25,27 +25,6 @@ Op " is not supported on host device."); #endif -#ifdef __SYCL_DEVICE_ONLY__ -#define __OPENCL_GLOBAL_AS__ __attribute__((opencl_global)) -#ifdef __ENABLE_USM_ADDR_SPACE__ -#define __OPENCL_GLOBAL_DEVICE_AS__ __attribute__((opencl_global_device)) -#define __OPENCL_GLOBAL_HOST_AS__ __attribute__((opencl_global_host)) -#else -#define __OPENCL_GLOBAL_DEVICE_AS__ __attribute__((opencl_global)) -#define __OPENCL_GLOBAL_HOST_AS__ __attribute__((opencl_global)) -#endif // __ENABLE_USM_ADDR_SPACE__ -#define __OPENCL_LOCAL_AS__ __attribute__((opencl_local)) -#define __OPENCL_CONSTANT_AS__ __attribute__((opencl_constant)) -#define __OPENCL_PRIVATE_AS__ __attribute__((opencl_private)) -#else -#define __OPENCL_GLOBAL_AS__ -#define __OPENCL_GLOBAL_DEVICE_AS__ -#define __OPENCL_GLOBAL_HOST_AS__ -#define __OPENCL_LOCAL_AS__ -#define __OPENCL_CONSTANT_AS__ -#define __OPENCL_PRIVATE_AS__ -#endif - namespace sycl { __SYCL_INLINE_VER_NAMESPACE(_V1) { namespace ext { From c5b90ce06e614a8d7561012ade04dfbb2e69f5ba Mon Sep 17 00:00:00 2001 From: "Wang, Di5" Date: Mon, 12 Dec 2022 02:10:20 -0800 Subject: [PATCH 21/26] sort properties in deduction guide --- .../oneapi/annotated_arg/annotated_arg.hpp | 82 +++++++++++-------- .../ext/oneapi/annotated_arg/properties.hpp | 3 - 2 files changed, 49 insertions(+), 36 deletions(-) diff --git a/sycl/include/sycl/ext/oneapi/annotated_arg/annotated_arg.hpp b/sycl/include/sycl/ext/oneapi/annotated_arg/annotated_arg.hpp index 39bc098233d08..80929e38db656 100644 --- a/sycl/include/sycl/ext/oneapi/annotated_arg/annotated_arg.hpp +++ b/sycl/include/sycl/ext/oneapi/annotated_arg/annotated_arg.hpp @@ -16,15 +16,6 @@ #include #include -#ifdef __SYCL_DEVICE_ONLY__ -#define __SYCL_HOST_NOT_SUPPORTED(Op) -#else -#define __SYCL_HOST_NOT_SUPPORTED(Op) \ - throw sycl::exception( \ - sycl::make_error_code(sycl::errc::feature_not_supported), \ - Op " is not supported on host device."); -#endif - namespace sycl { __SYCL_INLINE_VER_NAMESPACE(_V1) { namespace ext { @@ -41,14 +32,22 @@ struct HasSubscriptOperator } // namespace detail +namespace { +// Performs merge-sort on types with PropertyID. +template struct SortedProperties { + using split = typename detail::CreateTuplePairs::type; + using type = typename detail::MergeAll::type; +}; +} // namespace + // Deduction guide template -annotated_arg(T, Args... args) +annotated_arg(T, detail::properties_t) -> annotated_arg>; template -annotated_arg(T, properties>) - -> annotated_arg>; +annotated_arg(T, Args... args) + -> annotated_arg::type>>; template annotated_arg(annotated_arg, properties>) @@ -92,8 +91,6 @@ __SYCL_TYPE(annotated_arg) annotated_arg> { "Type T must be trivially destructible."); static_assert(is_property_list::value, "Property list is invalid."); - // static_assert(detail::SortedAllUnique>::value, - // "Duplicate properties in property list."); annotated_arg() noexcept = default; annotated_arg(const annotated_arg &) = default; @@ -103,23 +100,29 @@ __SYCL_TYPE(annotated_arg) annotated_arg> { const property_list_t &PropList = properties{}) noexcept : obj(sycl::detail::cast_AS(_ptr)) {} + // Constructs an annotated_arg object from a raw pointer and variadic + // properties. The new property set contains all properties of the input + // variadic properties. The same property in `Props...` and + // `PropertyValueTs...` must have the same property value. template - annotated_arg(T *_ptr, PropertyValueTs... props) noexcept + annotated_arg(T *_ptr, const PropertyValueTs &...props) noexcept : obj(sycl::detail::cast_AS(_ptr)) { + static_assert(detail::SortedAllUnique< + typename detail::Sorted::type>::value, + "Duplicate properties in the variadic properties."); static_assert( - std::is_same>>::value, + std::is_same< + property_list_t, + detail::merged_properties_t>::value, "The property list must contain all properties of the input of the " "constructor"); } // Constructs an annotated_arg object from another annotated_arg object. - // The property set PropertyListT contains all properties of the input - // annotated_arg object. If there are duplicate properties present in the - // property list of the input annotated_arg object, the values of the - // duplicate properties must be the same. + // The new property set contains all properties of the input + // annotated_arg object. The same property in `Props...` and `PropertyList2` + // must have the same property value. template explicit annotated_arg(const annotated_arg &other) noexcept : obj(other.obj) { @@ -135,6 +138,10 @@ __SYCL_TYPE(annotated_arg) annotated_arg> { "the input annotated_arg"); } + // Constructs an annotated_arg object from another annotated_arg object and a + // property list. The new property set is the union of property lists + // `PropertyListU` and `PropertyListV`. The same property in `PropertyListU` + // and `PropertyListV` must have the same property value. template explicit annotated_arg(const annotated_arg &other, const PropertyListV &proplist) noexcept @@ -202,22 +209,29 @@ __SYCL_TYPE(annotated_arg) annotated_arg> { const property_list_t &PropList = properties{}) noexcept : obj(_obj) {} + // Constructs an annotated_arg object from a raw pointer and variadic + // properties. The new property set contains all properties of the input + // variadic properties. The same property in `Props...` and + // `PropertyValueTs...` must have the same property value. template annotated_arg(const T &_obj, PropertyValueTs... props) noexcept : obj(_obj) { + static_assert(detail::SortedAllUnique< + typename detail::Sorted::type>::value, + "Duplicate properties in the variadic properties."); + static_assert( - std::is_same>>::value, + std::is_same< + property_list_t, + detail::merged_properties_t>::value, "The property list must contain all properties of the input of the " "constructor"); } // Constructs an annotated_arg object from another annotated_arg object. - // The property set PropertyListT contains all properties of the input - // annotated_arg object. If there are duplicate properties present in the - // property list of the input annotated_arg object, the values of the - // duplicate properties must be the same. + // The new property set contains all properties of the input + // annotated_arg object. The same property in `Props...` and `PropertyList2` + // must have the same property value. template explicit annotated_arg(const annotated_arg &other) noexcept : obj(other.obj) { @@ -233,6 +247,10 @@ __SYCL_TYPE(annotated_arg) annotated_arg> { "the input annotated_arg"); } + // Constructs an annotated_arg object from another annotated_arg object and a + // property list. The new property set is the union of property lists + // `PropertyListU` and `PropertyListV`. The same property in `PropertyListU` + // and `PropertyListV` must have the same property value. template explicit annotated_arg(const annotated_arg &other, const PropertyListV &proplist) noexcept @@ -279,5 +297,3 @@ __SYCL_TYPE(annotated_arg) annotated_arg> { } // namespace ext } // __SYCL_INLINE_VER_NAMESPACE(_V1) } // namespace sycl - -#undef __SYCL_HOST_NOT_SUPPORTED diff --git a/sycl/include/sycl/ext/oneapi/annotated_arg/properties.hpp b/sycl/include/sycl/ext/oneapi/annotated_arg/properties.hpp index d7d16b3074039..3d2584f658c8f 100644 --- a/sycl/include/sycl/ext/oneapi/annotated_arg/properties.hpp +++ b/sycl/include/sycl/ext/oneapi/annotated_arg/properties.hpp @@ -75,7 +75,6 @@ struct wait_request_key { property_value>; }; -#if __cplusplus >= 201703L // inline variables // non-mmhost properties inline constexpr register_map_key::value_t register_map; inline constexpr conduit_key::value_t conduit; @@ -102,8 +101,6 @@ inline constexpr read_write_mode_key::value_t inline constexpr read_write_mode_key::value_t read_write_mode_readwrite; -#endif // __cplusplus >= 201703L - template <> struct is_property_key : std::true_type {}; template <> struct is_property_key : std::true_type {}; template <> struct is_property_key : std::true_type {}; From 4558641787a098ada3ca0eda0a7d343722d3c73f Mon Sep 17 00:00:00 2001 From: "Wang, Di5" Date: Tue, 13 Dec 2022 06:21:16 -0800 Subject: [PATCH 22/26] resolve comments --- .../oneapi/annotated_arg/annotated_arg.hpp | 48 +++++++------------ .../annotated_arg_for_non_ptr.cpp | 30 ++++++++---- .../annotated_arg/annotated_arg_for_ptr.cpp | 24 +++++----- 3 files changed, 51 insertions(+), 51 deletions(-) diff --git a/sycl/include/sycl/ext/oneapi/annotated_arg/annotated_arg.hpp b/sycl/include/sycl/ext/oneapi/annotated_arg/annotated_arg.hpp index 80929e38db656..e62d7c5172ecb 100644 --- a/sycl/include/sycl/ext/oneapi/annotated_arg/annotated_arg.hpp +++ b/sycl/include/sycl/ext/oneapi/annotated_arg/annotated_arg.hpp @@ -30,24 +30,24 @@ struct HasSubscriptOperator : std::bool_constant< !std::is_void().operator[](0))>::value>{}; -} // namespace detail +// Deduce a `properties<>` type from given variadic properties +template struct DeducedProperties { + using type = decltype(properties{std::declval()...}); +}; -namespace { -// Performs merge-sort on types with PropertyID. -template struct SortedProperties { - using split = typename detail::CreateTuplePairs::type; - using type = typename detail::MergeAll::type; +// Partial specialization for deducing a `properties<>` type by forwarding the +// given `properties<>` type +template +struct DeducedProperties> { + using type = detail::properties_t; }; -} // namespace -// Deduction guide -template -annotated_arg(T, detail::properties_t) - -> annotated_arg>; +} // namespace detail +// Deduction guide template -annotated_arg(T, Args... args) - -> annotated_arg::type>>; +annotated_arg(T, Args...) + -> annotated_arg::type>; template annotated_arg(annotated_arg, properties>) @@ -87,8 +87,6 @@ __SYCL_TYPE(annotated_arg) annotated_arg> { #endif public: - static_assert(std::is_trivially_destructible::value, - "Type T must be trivially destructible."); static_assert(is_property_list::value, "Property list is invalid."); @@ -157,10 +155,9 @@ __SYCL_TYPE(annotated_arg) annotated_arg> { "of the input property lists"); } + operator T *() noexcept { return obj; } operator T *() const noexcept { return obj; } - operator const T *() const noexcept { return obj; } - T &operator[](std::ptrdiff_t idx) const noexcept { return obj[idx]; } template static constexpr bool has_property() { @@ -192,10 +189,8 @@ __SYCL_TYPE(annotated_arg) annotated_arg> { public: // T should be trivially copyable to be device-copyable - static_assert(std::is_trivially_copyable::value, - "Type T must be trivially copyable."); - static_assert(std::is_trivially_destructible::value, - "Type T must be trivially destructible."); + // static_assert(std::is_trivially_copyable::value, + // "Type T must be trivially copyable."); static_assert(is_property_list::value, "Property list is invalid."); static_assert(check_property_list::value, @@ -267,19 +262,12 @@ __SYCL_TYPE(annotated_arg) annotated_arg> { } operator T() noexcept { return obj; } - operator const T() const noexcept { return obj; } - - template - std::enable_if_t::value, - const decltype(std::declval().operator[](0))> & - operator[](std::ptrdiff_t idx) const noexcept { - return obj.operator[](idx); - } + operator T() const noexcept { return obj; } template std::enable_if_t::value, decltype(std::declval().operator[](0))> & - operator[](std::ptrdiff_t idx) noexcept { + operator[](std::ptrdiff_t idx) const noexcept { return obj.operator[](idx); } diff --git a/sycl/test/extensions/annotated_arg/annotated_arg_for_non_ptr.cpp b/sycl/test/extensions/annotated_arg/annotated_arg_for_non_ptr.cpp index cb35aae94c933..9d3d535bc408e 100644 --- a/sycl/test/extensions/annotated_arg/annotated_arg_for_non_ptr.cpp +++ b/sycl/test/extensions/annotated_arg/annotated_arg_for_non_ptr.cpp @@ -18,7 +18,8 @@ struct A : public B { A() {} A(int x_) : x(x_) {} - const int &operator[](std::ptrdiff_t idx) const { return x; } + int &operator[](std::ptrdiff_t idx) { return x; } + int &operator[](std::ptrdiff_t idx) const { return const_cast(x); } }; using annotated_arg_t1 = @@ -34,7 +35,10 @@ struct MyIP { MyIP(int *a_, const A &b_) : a(a_), b(b_) {} void operator()() const { + b[0] = 10; + const A &tmp = b; + A tmp2 = b; for (int i = 0; i < tmp.x; i++) { *a += 1; @@ -56,6 +60,7 @@ void TestVectorAddWithAnnotatedMMHosts() { A obj(0); // default ctor annotated_arg_t3 a1(obj); + // copy ctor auto a2(a1); auto a3(foo()); @@ -72,7 +77,7 @@ void TestVectorAddWithAnnotatedMMHosts() { "deduction guide failed 1"); // Construct from A instance and variadic properties - auto tmp13 = annotated_arg(obj, conduit, stable); // deduction guide + auto tmp13 = annotated_arg(obj, stable, conduit); // deduction guide static_assert(std::is_same::value, "deduction guide failed 2"); @@ -82,9 +87,9 @@ void TestVectorAddWithAnnotatedMMHosts() { // Construct from another annotated_arg // templated copy constructor annotated_arg arg11(tmp11); - annotated_arg arg14( + annotated_arg arg14( tmp11); // convertible type - auto arg12 = annotated_arg(tmp11); + auto arg12 = annotated_arg(tmp11); // default copy constructor auto arg13 = annotated_arg(tmp12); @@ -102,13 +107,13 @@ void TestVectorAddWithAnnotatedMMHosts() { "deduction guide failed 4"); static_assert(std::is_same::value, "deduction guide failed 5"); - annotated_arg arg24( + annotated_arg arg24( tmp12, properties{stable}); // convertible type // Property merge auto arg31 = annotated_arg_t3(obj, conduit); // OK auto arg32 = annotated_arg(arg31, properties{stable}); // OK - auto arg33 = annotated_arg(arg32, properties{conduit, stable}); // OK + auto arg33 = annotated_arg(arg32, properties{stable, conduit}); // OK // auto arg34 = annotated_arg(arg32, properties{conduit, latency<22>}); // // ERR: invalid property static_assert(std::is_same::value, @@ -116,15 +121,20 @@ void TestVectorAddWithAnnotatedMMHosts() { static_assert(std::is_same::value, "deduction guide failed 7"); // auto arg35 = annotated_arg(arg32, properties{conduit, dwidth<22>}); // - // ERR: two input property lists are conflict annotated_arg arg36(arg31, properties{latency<32>, - // stable}); // ERR: input property list is conflict with the declared type + // ERR: two input property lists are conflict + // annotated_arg + // arg36(arg31, properties{latency<32>, stable}); // ERR: input + // property list is conflict with the declared type // Implicit Conversion const A &x13 = arg32; // OK + A x14 = arg32; // OK // A& x11 = arg32; // ERR: non-const lvalue reference to type 'A' cannot // bind to a value of unrelated type + // operator[] + a1[0] = 5; + // has/get property static_assert(annotated_arg_t1::has_property(), "has property 1"); @@ -139,7 +149,7 @@ void TestVectorAddWithAnnotatedMMHosts() { // can't get non-existing property *raw = 0; - q.submit([&](handler &h) { h.single_task(MyIP{raw, A(5)}); }).wait(); + q.submit([&](handler &h) { h.single_task(MyIP{raw, a1}); }).wait(); std::cout << raw[0] << std::endl; free(raw, q); diff --git a/sycl/test/extensions/annotated_arg/annotated_arg_for_ptr.cpp b/sycl/test/extensions/annotated_arg/annotated_arg_for_ptr.cpp index 20a79563ef61e..12549bbbe4ebf 100644 --- a/sycl/test/extensions/annotated_arg/annotated_arg_for_ptr.cpp +++ b/sycl/test/extensions/annotated_arg/annotated_arg_for_ptr.cpp @@ -28,6 +28,7 @@ struct MyIP { void operator()() const { int *p = a; + const int *p2 = a; for (int i = 0; i < b; i++) { p[i] = i; @@ -45,9 +46,6 @@ void TestVectorAddWithAnnotatedMMHosts() { // Create the SYCL device queue queue q(sycl::ext::intel::fpga_selector_v); auto raw = malloc_shared(5, q); - for (int i = 0; i < 5; i++) { - raw[i] = 0; - } // default ctor annotated_arg_t3 a1; @@ -66,7 +64,7 @@ void TestVectorAddWithAnnotatedMMHosts() { static_assert(std::is_same::value, "deduction guide failed 1"); // Construct from raw pointers and variadic properties - auto tmp13 = annotated_arg(raw, awidth<32>, dwidth<32>); // deduction guide + auto tmp13 = annotated_arg(raw, dwidth<32>, awidth<32>); // deduction guide static_assert(std::is_same::value, "deduction guide failed 2"); auto tmp15 = annotated_arg(raw, awidth<32>); @@ -86,7 +84,7 @@ void TestVectorAddWithAnnotatedMMHosts() { annotated_arg, dwidth<32>})> arg11( tmp11); auto arg12 = - annotated_arg, dwidth<32>})>(tmp11); + annotated_arg, awidth<32>})>(tmp11); // default copy constructor auto arg13 = annotated_arg(tmp12); @@ -113,23 +111,27 @@ void TestVectorAddWithAnnotatedMMHosts() { // Property merge auto arg31 = annotated_arg_t3(raw, awidth<32>); // OK auto arg32 = annotated_arg(arg31, properties{dwidth<32>}); // OK - auto arg33 = annotated_arg(arg32, properties{awidth<32>, dwidth<32>}); // OK + auto arg33 = annotated_arg(arg32, properties{dwidth<32>, awidth<32>}); // OK auto arg34 = annotated_arg(arg32, properties{awidth<32>, latency<22>}); // OK static_assert(std::is_same::value, "deduction guide failed 6"); static_assert(std::is_same::value, "deduction guide failed 7"); // auto arg34 = annotated_arg(arg32, properties{awidth<32>, dwidth<22>}); // - // ERR: two input property lists are conflict annotated_arg, dwidth<32>})> arg35(arg31, - // properties{latency<32>, dwidth<32>}); // ERR: input property list is - // conflict with the declared type + // ERR: two input property lists are conflict + // annotated_arg, dwidth<32>})> + // arg35(arg31, properties{latency<32>, dwidth<32>}); // ERR: input + // property list is conflict with the declared type // Implicit Conversion int *x11 = arg13; const int *x13 = arg32; - // operator() + // operator[] + arg31[0] = 1; + for (int i = 1; i < 5; i++) { + arg31[i] = arg31[i - 1]; + } // has/get property static_assert(annotated_arg_t1::has_property(), "has property 1"); From 856221d37886a5ec9b07a62e03d7983f1fda4d3a Mon Sep 17 00:00:00 2001 From: "Wang, Di5" Date: Tue, 13 Dec 2022 06:51:27 -0800 Subject: [PATCH 23/26] use multi_pointer for decorated global pointer --- .../sycl/ext/oneapi/annotated_arg/annotated_arg.hpp | 7 +++---- 1 file changed, 3 insertions(+), 4 deletions(-) diff --git a/sycl/include/sycl/ext/oneapi/annotated_arg/annotated_arg.hpp b/sycl/include/sycl/ext/oneapi/annotated_arg/annotated_arg.hpp index e62d7c5172ecb..07330a8e7bca1 100644 --- a/sycl/include/sycl/ext/oneapi/annotated_arg/annotated_arg.hpp +++ b/sycl/include/sycl/ext/oneapi/annotated_arg/annotated_arg.hpp @@ -68,8 +68,7 @@ __SYCL_TYPE(annotated_arg) annotated_arg> { using property_list_t = detail::properties_t; #ifdef __SYCL_DEVICE_ONLY__ - using global_pointer_t = typename sycl::detail::DecoratedType< - T, access::address_space::global_space>::type *; + using global_pointer_t = decorated_global_ptr::pointer; #else using global_pointer_t = T *; #endif @@ -96,7 +95,7 @@ __SYCL_TYPE(annotated_arg) annotated_arg> { annotated_arg(T *_ptr, const property_list_t &PropList = properties{}) noexcept - : obj(sycl::detail::cast_AS(_ptr)) {} + : obj(global_pointer_t(_ptr)) {} // Constructs an annotated_arg object from a raw pointer and variadic // properties. The new property set contains all properties of the input @@ -104,7 +103,7 @@ __SYCL_TYPE(annotated_arg) annotated_arg> { // `PropertyValueTs...` must have the same property value. template annotated_arg(T *_ptr, const PropertyValueTs &...props) noexcept - : obj(sycl::detail::cast_AS(_ptr)) { + : obj(global_pointer_t(_ptr)) { static_assert(detail::SortedAllUnique< typename detail::Sorted::type>::value, "Duplicate properties in the variadic properties."); From 41fe464cd7bd1d5d43575ade2cd6a53ea4be4e1a Mon Sep 17 00:00:00 2001 From: "Wang, Di5" Date: Tue, 13 Dec 2022 07:31:41 -0800 Subject: [PATCH 24/26] fix LIT error --- sycl/include/sycl/ext/oneapi/annotated_arg/annotated_arg.hpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/sycl/include/sycl/ext/oneapi/annotated_arg/annotated_arg.hpp b/sycl/include/sycl/ext/oneapi/annotated_arg/annotated_arg.hpp index 07330a8e7bca1..5a0b065f85dd9 100644 --- a/sycl/include/sycl/ext/oneapi/annotated_arg/annotated_arg.hpp +++ b/sycl/include/sycl/ext/oneapi/annotated_arg/annotated_arg.hpp @@ -68,7 +68,7 @@ __SYCL_TYPE(annotated_arg) annotated_arg> { using property_list_t = detail::properties_t; #ifdef __SYCL_DEVICE_ONLY__ - using global_pointer_t = decorated_global_ptr::pointer; + using global_pointer_t = typename decorated_global_ptr::pointer; #else using global_pointer_t = T *; #endif From b3e4fd99c4988a21e28de3397c4385ee0be69a29 Mon Sep 17 00:00:00 2001 From: "Wang, Di5" Date: Wed, 14 Dec 2022 04:54:58 -0800 Subject: [PATCH 25/26] check is_device_copyable --- .../sycl/ext/oneapi/annotated_arg/annotated_arg.hpp | 11 +---------- .../sycl/ext/oneapi/annotated_arg/properties.hpp | 2 +- 2 files changed, 2 insertions(+), 11 deletions(-) diff --git a/sycl/include/sycl/ext/oneapi/annotated_arg/annotated_arg.hpp b/sycl/include/sycl/ext/oneapi/annotated_arg/annotated_arg.hpp index 5a0b065f85dd9..f58d1596248c8 100644 --- a/sycl/include/sycl/ext/oneapi/annotated_arg/annotated_arg.hpp +++ b/sycl/include/sycl/ext/oneapi/annotated_arg/annotated_arg.hpp @@ -104,9 +104,6 @@ __SYCL_TYPE(annotated_arg) annotated_arg> { template annotated_arg(T *_ptr, const PropertyValueTs &...props) noexcept : obj(global_pointer_t(_ptr)) { - static_assert(detail::SortedAllUnique< - typename detail::Sorted::type>::value, - "Duplicate properties in the variadic properties."); static_assert( std::is_same< property_list_t, @@ -187,9 +184,7 @@ __SYCL_TYPE(annotated_arg) annotated_arg> { #endif public: - // T should be trivially copyable to be device-copyable - // static_assert(std::is_trivially_copyable::value, - // "Type T must be trivially copyable."); + static_assert(is_device_copyable_v, "Type T must be device copyable."); static_assert(is_property_list::value, "Property list is invalid."); static_assert(check_property_list::value, @@ -209,10 +204,6 @@ __SYCL_TYPE(annotated_arg) annotated_arg> { // `PropertyValueTs...` must have the same property value. template annotated_arg(const T &_obj, PropertyValueTs... props) noexcept : obj(_obj) { - static_assert(detail::SortedAllUnique< - typename detail::Sorted::type>::value, - "Duplicate properties in the variadic properties."); - static_assert( std::is_same< property_list_t, diff --git a/sycl/include/sycl/ext/oneapi/annotated_arg/properties.hpp b/sycl/include/sycl/ext/oneapi/annotated_arg/properties.hpp index 3d2584f658c8f..eb21a01e6d229 100644 --- a/sycl/include/sycl/ext/oneapi/annotated_arg/properties.hpp +++ b/sycl/include/sycl/ext/oneapi/annotated_arg/properties.hpp @@ -1,4 +1,4 @@ -// properties.hpp - SYCL properties associated with annotated_arg/ptr // +//==-- properties.hpp - SYCL properties associated with annotated_arg/ptr --==// // // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. // See https://llvm.org/LICENSE.txt for license information. From 5567daaab08757d5c46da2f1d0d3bad8523daa98 Mon Sep 17 00:00:00 2001 From: "Wang, Di5" Date: Tue, 10 Jan 2023 13:08:04 -0800 Subject: [PATCH 26/26] run clang-format --- sycl/include/sycl/ext/oneapi/properties/property.hpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/sycl/include/sycl/ext/oneapi/properties/property.hpp b/sycl/include/sycl/ext/oneapi/properties/property.hpp index dbfc81e2b3563..08f6d41ecffbf 100644 --- a/sycl/include/sycl/ext/oneapi/properties/property.hpp +++ b/sycl/include/sycl/ext/oneapi/properties/property.hpp @@ -173,10 +173,10 @@ enum PropKind : uint32_t { WorkGroupSizeHint = 7, SubGroupSize = 8, DeviceHas = 9, - StreamingInterface = 10, // kernel attribute + StreamingInterface = 10, // kernel attribute RegisterMapInterface = 11, Pipelined = 12, - RegisterMap = 13, // kernel argument attribute + RegisterMap = 13, // kernel argument attribute Conduit = 14, Stable = 15, BufferLocation = 16,