13717251992

Committed 07 Mar 2025 09:01AM UTC coverage: 87.855% (-0.3%) from 88.138%

Build # 13717251992

Build Type

Pull #1984

github

Committed by

web-flow

Commit Message

Merge 5dca7bafb into 7aa6fb773

Pull Request Pull Request #1984: Add support for work_group_memory extension

Run Details

3159 of 3658 branches covered (86.36%)

Branch coverage included in aggregate %.

10 of 56 new or added lines in 2 files covered. (17.86%)

74 existing lines in 3 files now uncovered.

11981 of 13575 relevant lines covered (88.26%)

7122.61 hits per line

Source File
Press 'n' to go to next uncovered line, 'b' for previous

79.79

/libsyclinterface/source/dpctl_sycl_kernel_interface.cpp

//===--- dpctl_sycl_kernel_interface.cpp - Implements C API for sycl::kernel =//
//
//                      Data Parallel Control (dpctl)
//
// Copyright 2020-2025 Intel Corporation
//
// Licensed under the Apache License, Version 2.0 (the "License");
// you may not use this file except in compliance with the License.
// You may obtain a copy of the License at
//
//    http://www.apache.org/licenses/LICENSE-2.0
//
// Unless required by applicable law or agreed to in writing, software
// distributed under the License is distributed on an "AS IS" BASIS,
// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
// See the License for the specific language governing permissions and
// limitations under the License.
//
//===----------------------------------------------------------------------===//
///
/// \file
/// This file implements the functions declared in
/// dpctl_sycl_kernel_interface.h.
///
//===----------------------------------------------------------------------===//

#include "dpctl_sycl_kernel_interface.h"
#include "Config/dpctl_config.h"
#include "dpctl_error_handlers.h"
#include "dpctl_string_utils.hpp"
#include "dpctl_sycl_type_casters.hpp"
#include <cstdint>
#include <stddef.h>
#include <sycl/sycl.hpp> /* Sycl headers */

using namespace sycl;

namespace
{
static_assert(__SYCL_COMPILER_VERSION >= __SYCL_COMPILER_VERSION_REQUIRED,
              "The compiler does not meet minimum version requirement");

using namespace dpctl::syclinterface;
} // end of anonymous namespace

size_t DPCTLKernel_GetNumArgs(__dpctl_keep const DPCTLSyclKernelRef KRef)
{
    if (!KRef) {
        error_handler("Cannot get the number of arguments from "
                      "DPCTLSyclKernelRef as input is a nullptr.",
                      __FILE__, __func__, __LINE__);
        return -1;
    }

    auto sycl_kernel = unwrap<kernel>(KRef);
    auto num_args = sycl_kernel->get_info<info::kernel::num_args>();
    return static_cast<size_t>(num_args);
}

void DPCTLKernel_Delete(__dpctl_take DPCTLSyclKernelRef KRef)
{
    delete unwrap<kernel>(KRef);
}

__dpctl_give DPCTLSyclKernelRef
DPCTLKernel_Copy(__dpctl_keep const DPCTLSyclKernelRef KRef)
{
    auto Kernel = unwrap<kernel>(KRef);
    if (!Kernel) {
        error_handler("Cannot copy DPCTLSyclKernelRef as input is a nullptr",
                      __FILE__, __func__, __LINE__);
        return nullptr;
    }
    try {
        auto CopiedKernel = new kernel(*Kernel);
        return wrap<kernel>(CopiedKernel);
    } catch (std::exception const &e) {
        error_handler(e, __FILE__, __func__, __LINE__);
        return nullptr;
    }
}

size_t DPCTLKernel_GetWorkGroupSize(__dpctl_keep const DPCTLSyclKernelRef KRef)
{
    if (!KRef) {
        error_handler("Input DPCTKSyclKernelRef is nullptr.", __FILE__,
                      __func__, __LINE__);
        return 0;
    }

    auto sycl_kern = unwrap<kernel>(KRef);
    auto devs = sycl_kern->get_kernel_bundle().get_devices();
    if (devs.empty()) {
        error_handler("Input DPCTKSyclKernelRef has no associated device.",
                      __FILE__, __func__, __LINE__);
        return 0;
    }
    auto v = sycl_kern->get_info<info::kernel_device_specific::work_group_size>(
        devs[0]);
    return static_cast<size_t>(v);
}

size_t DPCTLKernel_GetPreferredWorkGroupSizeMultiple(
    __dpctl_keep const DPCTLSyclKernelRef KRef)
{
    if (!KRef) {
        error_handler("Input DPCTKSyclKernelRef is nullptr.", __FILE__,
                      __func__, __LINE__);
        return 0;
    }

    auto sycl_kern = unwrap<kernel>(KRef);
    auto devs = sycl_kern->get_kernel_bundle().get_devices();
    if (devs.empty()) {
        error_handler("Input DPCTKSyclKernelRef has no associated device.",
                      __FILE__, __func__, __LINE__);
        return 0;
    }
    auto v = sycl_kern->get_info<
        info::kernel_device_specific::preferred_work_group_size_multiple>(
        devs[0]);
    return static_cast<size_t>(v);
}

size_t DPCTLKernel_GetPrivateMemSize(__dpctl_keep const DPCTLSyclKernelRef KRef)
{
    if (!KRef) {
        error_handler("Input DPCTKSyclKernelRef is nullptr.", __FILE__,
                      __func__, __LINE__);
        return 0;
    }

    auto sycl_kern = unwrap<kernel>(KRef);
    auto devs = sycl_kern->get_kernel_bundle().get_devices();
    if (devs.empty()) {
        error_handler("Input DPCTKSyclKernelRef has no associated device.",
                      __FILE__, __func__, __LINE__);
        return 0;
    }
    auto v =
        sycl_kern->get_info<info::kernel_device_specific::private_mem_size>(
            devs[0]);
    return static_cast<size_t>(v);
}

uint32_t
DPCTLKernel_GetMaxNumSubGroups(__dpctl_keep const DPCTLSyclKernelRef KRef)
{
    if (!KRef) {
        error_handler("Input DPCTKSyclKernelRef is nullptr.", __FILE__,
                      __func__, __LINE__);
        return 0;
    }

    auto sycl_kern = unwrap<kernel>(KRef);
    auto devs = sycl_kern->get_kernel_bundle().get_devices();
    if (devs.empty()) {
        error_handler("Input DPCTKSyclKernelRef has no associated device.",
                      __FILE__, __func__, __LINE__);
        return 0;
    }
    auto v =
        sycl_kern->get_info<info::kernel_device_specific::max_num_sub_groups>(
            devs[0]);
    return static_cast<uint32_t>(v);
}

uint32_t
DPCTLKernel_GetMaxSubGroupSize(__dpctl_keep const DPCTLSyclKernelRef KRef)
{
    if (!KRef) {
        error_handler("Input DPCTKSyclKernelRef is nullptr.", __FILE__,
                      __func__, __LINE__);
        return 0;
    }

    auto sycl_kern = unwrap<kernel>(KRef);
    auto devs = sycl_kern->get_kernel_bundle().get_devices();
    if (devs.empty()) {
        error_handler("Input DPCTKSyclKernelRef has no associated device.",
                      __FILE__, __func__, __LINE__);
        return 0;
    }
    auto v =
        sycl_kern->get_info<info::kernel_device_specific::max_sub_group_size>(
            devs[0]);
    return v;
}

uint32_t
DPCTLKernel_GetCompileNumSubGroups(__dpctl_keep const DPCTLSyclKernelRef KRef)
{
    if (!KRef) {
        error_handler("Input DPCTKSyclKernelRef is nullptr.", __FILE__,
                      __func__, __LINE__);
        return 0;
    }

    auto sycl_kern = unwrap<kernel>(KRef);
    auto devs = sycl_kern->get_kernel_bundle().get_devices();
    if (devs.empty()) {
        error_handler("Input DPCTKSyclKernelRef has no associated device.",
                      __FILE__, __func__, __LINE__);
        return 0;
    }
    auto v =
        sycl_kern
            ->get_info<info::kernel_device_specific::compile_num_sub_groups>(
                devs[0]);
    return static_cast<uint32_t>(v);
}

uint32_t
DPCTLKernel_GetCompileSubGroupSize(__dpctl_keep const DPCTLSyclKernelRef KRef)
{
    if (!KRef) {
        error_handler("Input DPCTKSyclKernelRef is nullptr.", __FILE__,
                      __func__, __LINE__);
        return 0;
    }

    auto sycl_kern = unwrap<kernel>(KRef);
    auto devs = sycl_kern->get_kernel_bundle().get_devices();
    if (devs.empty()) {
        error_handler("Input DPCTKSyclKernelRef has no associated device.",
                      __FILE__, __func__, __LINE__);
        return 0;
    }
    auto v =
        sycl_kern
            ->get_info<info::kernel_device_specific::compile_sub_group_size>(
                devs[0]);
    return static_cast<uint32_t>(v);
}

1	//===--- dpctl_sycl_kernel_interface.cpp - Implements C API for sycl::kernel =//
2	//
3	// Data Parallel Control (dpctl)
4	//
5	// Copyright 2020-2025 Intel Corporation
6	//
7	// Licensed under the Apache License, Version 2.0 (the "License");
8	// you may not use this file except in compliance with the License.
9	// You may obtain a copy of the License at
10	//
11	// http://www.apache.org/licenses/LICENSE-2.0
12	//
13	// Unless required by applicable law or agreed to in writing, software
14	// distributed under the License is distributed on an "AS IS" BASIS,
15	// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
16	// See the License for the specific language governing permissions and
17	// limitations under the License.
18	//
19	//===----------------------------------------------------------------------===//
20	///
21	/// \file
22	/// This file implements the functions declared in
23	/// dpctl_sycl_kernel_interface.h.
24	///
25	//===----------------------------------------------------------------------===//
26
27	#include "dpctl_sycl_kernel_interface.h"
28	#include "Config/dpctl_config.h"
29	#include "dpctl_error_handlers.h"
30	#include "dpctl_string_utils.hpp"
31	#include "dpctl_sycl_type_casters.hpp"
32	#include <cstdint>
33	#include <stddef.h>
34	#include <sycl/sycl.hpp> /* Sycl headers */
35
36	using namespace sycl;
37
38	namespace
39	{
40	static_assert(__SYCL_COMPILER_VERSION >= __SYCL_COMPILER_VERSION_REQUIRED,
41	"The compiler does not meet minimum version requirement");
42
43	using namespace dpctl::syclinterface;
44	} // end of anonymous namespace
45
46	size_t DPCTLKernel_GetNumArgs(__dpctl_keep const DPCTLSyclKernelRef KRef)
47	{	32✔
48	if (!KRef) {	32✔
49	error_handler("Cannot get the number of arguments from "	1✔
50	"DPCTLSyclKernelRef as input is a nullptr.",	1✔
51	__FILE__, __func__, __LINE__);	1✔
52	return -1;	1✔
53	}	1✔
54
55	auto sycl_kernel = unwrap<kernel>(KRef);	31✔
56	auto num_args = sycl_kernel->get_info<info::kernel::num_args>();	31✔
57	return static_cast<size_t>(num_args);	31✔
58	}	32✔
59
60	void DPCTLKernel_Delete(__dpctl_take DPCTLSyclKernelRef KRef)
61	{	96✔
62	delete unwrap<kernel>(KRef);	96✔
63	}	96✔
64
65	__dpctl_give DPCTLSyclKernelRef
66	DPCTLKernel_Copy(__dpctl_keep const DPCTLSyclKernelRef KRef)
67	{	10✔
68	auto Kernel = unwrap<kernel>(KRef);	10✔
69	if (!Kernel) {	10✔
70	error_handler("Cannot copy DPCTLSyclKernelRef as input is a nullptr",	1✔
71	__FILE__, __func__, __LINE__);	1✔
72	return nullptr;	1✔
73	}	1✔
74	try {	9✔
75	auto CopiedKernel = new kernel(*Kernel);	9✔
76	return wrap<kernel>(CopiedKernel);	9✔
77	} catch (std::exception const &e) {	9✔
78	error_handler(e, __FILE__, __func__, __LINE__);	×
79	return nullptr;	×
80	}	×
81	}	9✔
82
83	size_t DPCTLKernel_GetWorkGroupSize(__dpctl_keep const DPCTLSyclKernelRef KRef)
84	{	23✔
85	if (!KRef) {	23✔
86	error_handler("Input DPCTKSyclKernelRef is nullptr.", __FILE__,	1✔
87	__func__, __LINE__);	1✔
88	return 0;	1✔
89	}	1✔
90
91	auto sycl_kern = unwrap<kernel>(KRef);	22✔
92	auto devs = sycl_kern->get_kernel_bundle().get_devices();	22✔
93	if (devs.empty()) {	22!
94	error_handler("Input DPCTKSyclKernelRef has no associated device.",	×
95	__FILE__, __func__, __LINE__);	×
96	return 0;	×
97	}	×
98	auto v = sycl_kern->get_info<info::kernel_device_specific::work_group_size>(	22✔
99	devs[0]);	22✔
100	return static_cast<size_t>(v);	22✔
101	}	22✔
102
103	size_t DPCTLKernel_GetPreferredWorkGroupSizeMultiple(
104	__dpctl_keep const DPCTLSyclKernelRef KRef)
105	{	7✔
106	if (!KRef) {	7✔
107	error_handler("Input DPCTKSyclKernelRef is nullptr.", __FILE__,	1✔
108	__func__, __LINE__);	1✔
109	return 0;	1✔
110	}	1✔
111
112	auto sycl_kern = unwrap<kernel>(KRef);	6✔
113	auto devs = sycl_kern->get_kernel_bundle().get_devices();	6✔
114	if (devs.empty()) {	6!
115	error_handler("Input DPCTKSyclKernelRef has no associated device.",	×
116	__FILE__, __func__, __LINE__);	×
117	return 0;	×
118	}	×
119	auto v = sycl_kern->get_info<	6✔
120	info::kernel_device_specific::preferred_work_group_size_multiple>(	6✔
121	devs[0]);	6✔
122	return static_cast<size_t>(v);	6✔
123	}	6✔
124
125	size_t DPCTLKernel_GetPrivateMemSize(__dpctl_keep const DPCTLSyclKernelRef KRef)
126	{	7✔
127	if (!KRef) {	7✔
128	error_handler("Input DPCTKSyclKernelRef is nullptr.", __FILE__,	1✔
129	__func__, __LINE__);	1✔
130	return 0;	1✔
131	}	1✔
132
133	auto sycl_kern = unwrap<kernel>(KRef);	6✔
134	auto devs = sycl_kern->get_kernel_bundle().get_devices();	6✔
135	if (devs.empty()) {	6!
136	error_handler("Input DPCTKSyclKernelRef has no associated device.",	×
137	__FILE__, __func__, __LINE__);	×
138	return 0;	×
139	}	×
140	auto v =	6✔
141	sycl_kern->get_info<info::kernel_device_specific::private_mem_size>(	6✔
142	devs[0]);	6✔
143	return static_cast<size_t>(v);	6✔
144	}	6✔
145
146	uint32_t
147	DPCTLKernel_GetMaxNumSubGroups(__dpctl_keep const DPCTLSyclKernelRef KRef)
148	{	7✔
149	if (!KRef) {	7✔
150	error_handler("Input DPCTKSyclKernelRef is nullptr.", __FILE__,	1✔
151	__func__, __LINE__);	1✔
152	return 0;	1✔
153	}	1✔
154
155	auto sycl_kern = unwrap<kernel>(KRef);	6✔
156	auto devs = sycl_kern->get_kernel_bundle().get_devices();	6✔
157	if (devs.empty()) {	6!
158	error_handler("Input DPCTKSyclKernelRef has no associated device.",	×
159	__FILE__, __func__, __LINE__);	×
160	return 0;	×
161	}	×
162	auto v =	6✔
163	sycl_kern->get_info<info::kernel_device_specific::max_num_sub_groups>(	6✔
164	devs[0]);	6✔
165	return static_cast<uint32_t>(v);	6✔
166	}	6✔
167
168	uint32_t
169	DPCTLKernel_GetMaxSubGroupSize(__dpctl_keep const DPCTLSyclKernelRef KRef)
170	{	7✔
171	if (!KRef) {	7✔
172	error_handler("Input DPCTKSyclKernelRef is nullptr.", __FILE__,	1✔
173	__func__, __LINE__);	1✔
174	return 0;	1✔
175	}	1✔
176
177	auto sycl_kern = unwrap<kernel>(KRef);	6✔
178	auto devs = sycl_kern->get_kernel_bundle().get_devices();	6✔
179	if (devs.empty()) {	6!
180	error_handler("Input DPCTKSyclKernelRef has no associated device.",	×
181	__FILE__, __func__, __LINE__);	×
182	return 0;	×
183	}	×
184	auto v =	6✔
185	sycl_kern->get_info<info::kernel_device_specific::max_sub_group_size>(	6✔
186	devs[0]);	6✔
187	return v;	6✔
188	}	6✔
189
190	uint32_t
191	DPCTLKernel_GetCompileNumSubGroups(__dpctl_keep const DPCTLSyclKernelRef KRef)
192	{	7✔
193	if (!KRef) {	7✔
194	error_handler("Input DPCTKSyclKernelRef is nullptr.", __FILE__,	1✔
195	__func__, __LINE__);	1✔
196	return 0;	1✔
197	}	1✔
198
199	auto sycl_kern = unwrap<kernel>(KRef);	6✔
200	auto devs = sycl_kern->get_kernel_bundle().get_devices();	6✔
201	if (devs.empty()) {	6!
202	error_handler("Input DPCTKSyclKernelRef has no associated device.",	×
203	__FILE__, __func__, __LINE__);	×
204	return 0;	×
205	}	×
206	auto v =	6✔
207	sycl_kern	6✔
208	->get_info<info::kernel_device_specific::compile_num_sub_groups>(	6✔
209	devs[0]);	6✔
210	return static_cast<uint32_t>(v);	6✔
211	}	6✔
212
213	uint32_t
214	DPCTLKernel_GetCompileSubGroupSize(__dpctl_keep const DPCTLSyclKernelRef KRef)
215	{	7✔
216	if (!KRef) {	7✔
217	error_handler("Input DPCTKSyclKernelRef is nullptr.", __FILE__,	1✔
218	__func__, __LINE__);	1✔
219	return 0;	1✔
220	}	1✔
221
222	auto sycl_kern = unwrap<kernel>(KRef);	6✔
223	auto devs = sycl_kern->get_kernel_bundle().get_devices();	6✔
224	if (devs.empty()) {	6!
225	error_handler("Input DPCTKSyclKernelRef has no associated device.",	×
226	__FILE__, __func__, __LINE__);	×
227	return 0;	×
228	}	×
229	auto v =	6✔
230	sycl_kern	6✔
231	->get_info<info::kernel_device_specific::compile_sub_group_size>(	6✔
232	devs[0]);	6✔
233	return static_cast<uint32_t>(v);	6✔
234	}	6✔

IntelPython / dpctl / 13717251992

Source File Press 'n' to go to next uncovered line, 'b' for previous

Source File
Press 'n' to go to next uncovered line, 'b' for previous