diff --git a/sycl/doc/extensions/proposed/sycl_ext_oneapi_get_kernel_info.asciidoc b/sycl/doc/extensions/supported/sycl_ext_oneapi_get_kernel_info.asciidoc similarity index 94% rename from sycl/doc/extensions/proposed/sycl_ext_oneapi_get_kernel_info.asciidoc rename to sycl/doc/extensions/supported/sycl_ext_oneapi_get_kernel_info.asciidoc index e567a5297a6b3..33c9e615eeeed 100644 --- a/sycl/doc/extensions/proposed/sycl_ext_oneapi_get_kernel_info.asciidoc +++ b/sycl/doc/extensions/supported/sycl_ext_oneapi_get_kernel_info.asciidoc @@ -43,11 +43,7 @@ SYCL specification refer to that revision. == Status -This is a proposed extension specification, intended to gather community -feedback. Interfaces defined in this specification may not be implemented yet -or may be in a preliminary state. The specification itself may also change in -incompatible ways before it is finalized. *Shipping software products should -not rely on APIs defined in this specification.* +This extension is implemented and fully supported by DPC++. == Overview diff --git a/sycl/include/sycl/ext/oneapi/get_kernel_info.hpp b/sycl/include/sycl/ext/oneapi/get_kernel_info.hpp new file mode 100644 index 0000000000000..bd435426dfcd1 --- /dev/null +++ b/sycl/include/sycl/ext/oneapi/get_kernel_info.hpp @@ -0,0 +1,48 @@ +//==----- get_kernel_info.hpp --- SYCL get_kernel_info 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 + +namespace sycl { +inline namespace _V1 { +namespace ext::oneapi { + +template +typename sycl::detail::is_kernel_info_desc::return_type +get_kernel_info(const context &Ctx) { + auto Bundle = + sycl::get_kernel_bundle(Ctx); + return Bundle.template get_kernel().template get_info(); +} + +template +typename sycl::detail::is_kernel_device_specific_info_desc::return_type +get_kernel_info(const context &Ctx, const device &Dev) { + auto Bundle = + sycl::get_kernel_bundle(Ctx); + return Bundle.template get_kernel().template get_info(Dev); +} + +template +typename sycl::detail::is_kernel_device_specific_info_desc::return_type +get_kernel_info(const queue &Q) { + auto Bundle = + sycl::get_kernel_bundle( + Q.get_context()); + return Bundle.template get_kernel().template get_info( + Q.get_device()); +} + +} // namespace ext::oneapi +} // namespace _V1 +} // namespace sycl diff --git a/sycl/include/sycl/sycl.hpp b/sycl/include/sycl/sycl.hpp index ab44642639b0d..868fec035eed7 100644 --- a/sycl/include/sycl/sycl.hpp +++ b/sycl/include/sycl/sycl.hpp @@ -103,6 +103,7 @@ #include #include #include +#include #include #include #include diff --git a/sycl/source/feature_test.hpp.in b/sycl/source/feature_test.hpp.in index a61e504eb5e4c..8f4fb05752efc 100644 --- a/sycl/source/feature_test.hpp.in +++ b/sycl/source/feature_test.hpp.in @@ -108,6 +108,7 @@ inline namespace _V1 { #define SYCL_EXT_ONEAPI_RAW_KERNEL_ARG 1 #define SYCL_EXT_ONEAPI_PROFILING_TAG 1 #define SYCL_EXT_ONEAPI_ENQUEUE_NATIVE_COMMAND 1 +#define SYCL_EXT_ONEAPI_GET_KERNEL_INFO 1 // In progress yet #define SYCL_EXT_ONEAPI_ATOMIC16 0 diff --git a/sycl/test-e2e/Basic/kernel_info.cpp b/sycl/test-e2e/Basic/kernel_info.cpp index ffe8867ed2a8d..864db59dc43b5 100644 --- a/sycl/test-e2e/Basic/kernel_info.cpp +++ b/sycl/test-e2e/Basic/kernel_info.cpp @@ -1,12 +1,6 @@ // RUN: %{build} -o %t.out // RUN: %{run} %t.out // -// Fail is flaky for level_zero, enable when fixed. -// UNSUPPORTED: level_zero -// -// Consistently fails with opencl gpu, enable when fixed. -// XFAIL: opencl && gpu - //==--- kernel_info.cpp - SYCL kernel info test ----------------------------==// // // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. @@ -17,17 +11,36 @@ #include #include +#include using namespace sycl; +namespace syclex = sycl::ext::oneapi; + +auto checkExceptionIsThrown = [](auto &getInfoFunc, + const std::string &refErrMsg, + std::error_code refErrc) { + std::string errMsg = ""; + std::error_code errc; + bool exceptionWasThrown = false; + try { + std::ignore = getInfoFunc(); + } catch (exception &e) { + errMsg = e.what(); + errc = e.code(); + exceptionWasThrown = true; + } + assert(exceptionWasThrown); + assert(errMsg == refErrMsg); + assert(errc == refErrc); +}; int main() { queue q; - + auto ctx = q.get_context(); buffer buf(range<1>(1)); - auto KernelID = sycl::get_kernel_id(); - auto KB = - get_kernel_bundle(q.get_context(), {KernelID}); - kernel krn = KB.get_kernel(KernelID); + auto kernelID = sycl::get_kernel_id(); + auto kb = get_kernel_bundle(ctx, {kernelID}); + kernel krn = kb.get_kernel(kernelID); q.submit([&](handler &cgh) { auto acc = buf.get_access(cgh); @@ -37,30 +50,34 @@ int main() { const std::string krnName = krn.get_info(); assert(!krnName.empty()); - std::string ErrMsg = ""; - std::error_code Errc; - bool ExceptionWasThrown = false; - try { - const cl_uint krnArgCount = krn.get_info(); - } catch (exception &e) { - ErrMsg = e.what(); - Errc = e.code(); - ExceptionWasThrown = true; - } - assert(ExceptionWasThrown && "Invalid using of \"info::kernel::num_args\" " - "query should throw an exception."); - assert(ErrMsg == - "info::kernel::num_args descriptor may only be used to query a kernel " - "that resides in a kernel bundle constructed using a backend specific" - "interoperability function or to query a device built-in kernel"); - assert(Errc == errc::invalid); + auto refErrMsg = + "info::kernel::num_args descriptor may only be used to query a kernel " + "that resides in a kernel bundle constructed using a backend specific" + "interoperability function or to query a device built-in kernel"; + auto refErrc = errc::invalid; + auto getInfoNumArgsFunc = [&]() -> cl_uint { + return krn.get_info(); + }; + checkExceptionIsThrown(getInfoNumArgsFunc, refErrMsg, refErrc); + auto getInfoNumArgsFuncExt = [&]() { + return syclex::get_kernel_info(ctx); + }; + checkExceptionIsThrown(getInfoNumArgsFuncExt, refErrMsg, refErrc); const context krnCtx = krn.get_info(); assert(krnCtx == q.get_context()); const cl_uint krnRefCount = krn.get_info(); assert(krnRefCount > 0); - const std::string krnAttr = krn.get_info(); - assert(krnAttr.empty()); + + // Use ext_oneapi_get_kernel_info extension and check that answers match. + const context krnCtxExt = + syclex::get_kernel_info(ctx); + assert(krnCtxExt == krnCtx); + // Reference count might be different because we have to retain the kernel + // handle first to fetch the info. So just check that it is not 0. + const cl_uint krnRefCountExt = + syclex::get_kernel_info(ctx); + assert(krnRefCountExt > 0); device dev = q.get_device(); const size_t wgSize = @@ -82,34 +99,70 @@ int main() { krn.get_info(dev); assert(compileNumSg <= maxNumSg); - { - std::error_code Errc; - std::string ErrMsg = ""; - bool IsExceptionThrown = false; - try { - krn.get_info(dev); - auto BuiltInIds = dev.get_info(); - bool isBuiltInKernel = std::find(BuiltInIds.begin(), BuiltInIds.end(), - KernelID) != BuiltInIds.end(); - bool isCustomDevice = dev.get_info() == - sycl::info::device_type::custom; - assert((isCustomDevice || isBuiltInKernel) && - "info::kernel_device_specific::global_work_size descriptor can " - "only be used with custom device " - "or built-in kernel."); + // Use ext_oneapi_get_kernel_info extension and check that answers match. + const size_t wgSizeExt = syclex::get_kernel_info< + SingleTask, info::kernel_device_specific::work_group_size>(ctx, dev); + assert(wgSizeExt == wgSize); + const size_t prefWGSizeMultExt = syclex::get_kernel_info< + SingleTask, + info::kernel_device_specific::preferred_work_group_size_multiple>(ctx, + dev); + assert(prefWGSizeMultExt == prefWGSizeMult); + const cl_uint maxSgSizeExt = syclex::get_kernel_info< + SingleTask, info::kernel_device_specific::max_sub_group_size>(ctx, dev); + assert(maxSgSizeExt == maxSgSize); + const cl_uint compileSgSizeExt = syclex::get_kernel_info< + SingleTask, info::kernel_device_specific::compile_sub_group_size>(ctx, + dev); + assert(compileSgSizeExt == compileSgSize); + const cl_uint maxNumSgExt = syclex::get_kernel_info< + SingleTask, info::kernel_device_specific::max_num_sub_groups>(ctx, dev); + assert(maxNumSgExt == maxNumSg); + const cl_uint compileNumSgExt = syclex::get_kernel_info< + SingleTask, info::kernel_device_specific::compile_num_sub_groups>(ctx, + dev); + assert(compileNumSgExt == compileNumSg); - } catch (sycl::exception &e) { - IsExceptionThrown = true; - Errc = e.code(); - ErrMsg = e.what(); - } - assert(IsExceptionThrown && - "Invalid using of info::kernel_device_specific::global_work_size " - "query should throw an exception."); - assert(Errc == errc::invalid); - assert(ErrMsg == - "info::kernel_device_specific::global_work_size descriptor may only " - "be used if the device type is device_type::custom or if the " - "kernel is a built-in kernel."); - } + // Use ext_oneapi_get_kernel_info extension with queue parameter and check the + // result. + const size_t wgSizeExtQ = + syclex::get_kernel_info(q); + assert(wgSizeExtQ == wgSize); + const size_t prefWGSizeMultExtQ = syclex::get_kernel_info< + SingleTask, + info::kernel_device_specific::preferred_work_group_size_multiple>(q); + assert(prefWGSizeMultExtQ == prefWGSizeMult); + const cl_uint maxSgSizeExtQ = syclex::get_kernel_info< + SingleTask, info::kernel_device_specific::max_sub_group_size>(q); + assert(maxSgSizeExtQ == maxSgSize); + const cl_uint compileSgSizeExtQ = syclex::get_kernel_info< + SingleTask, info::kernel_device_specific::compile_sub_group_size>(q); + assert(compileSgSizeExtQ == compileSgSize); + const cl_uint maxNumSgExtQ = syclex::get_kernel_info< + SingleTask, info::kernel_device_specific::max_num_sub_groups>(q); + assert(maxNumSgExtQ == maxNumSg); + const cl_uint compileNumSgExtQ = syclex::get_kernel_info< + SingleTask, info::kernel_device_specific::compile_num_sub_groups>(q); + assert(compileNumSgExtQ == compileNumSg); + + refErrMsg = + "info::kernel_device_specific::global_work_size descriptor may only " + "be used if the device type is device_type::custom or if the " + "kernel is a built-in kernel."; + auto getInfoGWSFunc = [&]() { + return krn.get_info( + dev); + }; + checkExceptionIsThrown(getInfoGWSFunc, refErrMsg, refErrc); + auto getInfoGWSFuncExt = [&]() { + return syclex::get_kernel_info< + SingleTask, info::kernel_device_specific::global_work_size>(ctx, dev); + }; + checkExceptionIsThrown(getInfoGWSFuncExt, refErrMsg, refErrc); + auto getInfoGWSFuncExtQ = [&]() { + return syclex::get_kernel_info< + SingleTask, info::kernel_device_specific::global_work_size>(q); + }; + checkExceptionIsThrown(getInfoGWSFuncExtQ, refErrMsg, refErrc); } diff --git a/sycl/test-e2e/Basic/kernel_info_attr.cpp b/sycl/test-e2e/Basic/kernel_info_attr.cpp new file mode 100644 index 0000000000000..9ceb6b7a4005e --- /dev/null +++ b/sycl/test-e2e/Basic/kernel_info_attr.cpp @@ -0,0 +1,45 @@ +// RUN: %{build} -o %t.out +// RUN: %{run} %t.out +// +// Fail is flaky for level_zero, enable when fixed. +// UNSUPPORTED: level_zero +// +// Consistently fails with opencl gpu, enable when fixed. +// XFAIL: opencl && gpu +// XFAIL-TRACKER: GSD-8971 + +//==--- kernel_info_attr.cpp - SYCL info::kernel::attributes test ---==// +// +// 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 +// +//===---------------------------------------------------------------===// + +#include +#include +#include + +using namespace sycl; +namespace syclex = sycl::ext::oneapi; + +int main() { + queue q; + auto ctx = q.get_context(); + buffer buf(range<1>(1)); + auto KernelID = sycl::get_kernel_id(); + auto KB = get_kernel_bundle(ctx, {KernelID}); + kernel krn = KB.get_kernel(KernelID); + + q.submit([&](handler &cgh) { + auto acc = buf.get_access(cgh); + cgh.single_task(krn, [=]() { acc[0] = acc[0] + 1; }); + }); + + const std::string krnAttr = krn.get_info(); + assert(krnAttr.empty()); + const std::string krnAttrExt = + syclex::get_kernel_info(ctx); + assert(krnAttr == krnAttrExt); + return 0; +}