Skip to content

Commit 789cb57

Browse files
[SYCL] Split throwOnKernelParameterMisuse to improve compile time
For the following code (note no submit, just bare minimum `cgh.single_task`): ``` int main() { int *p; sycl::detail::loop<2>([&](auto outer_idx) { sycl::detail::loop<200>([&](auto idx) { auto krn = [=]() { *p = 42; }; auto s = [&](sycl::handler &cgh) { cgh.single_task(krn); }; (void)sycl::detail::type_erased_cgfo_ty{s}; }); }); } ``` compiled as ` $ time clang++ -isystem ~/sycl/build/include a.cpp -c -o /dev/null -Wno-deprecated` (to simulate host-only compilation), improves 3.4s -> 3.0s.
1 parent 8e7ec0f commit 789cb57

File tree

1 file changed

+14
-7
lines changed

1 file changed

+14
-7
lines changed

sycl/include/sycl/handler.hpp

Lines changed: 14 additions & 7 deletions
Original file line numberDiff line numberDiff line change
@@ -3745,13 +3745,13 @@ class __SYCL_EXPORT handler {
37453745
/// According to section 4.7.6.11. of the SYCL specification, a local accessor
37463746
/// must not be used in a SYCL kernel function that is invoked via single_task
37473747
/// or via the simple form of parallel_for that takes a range parameter.
3748-
template <typename KernelName, typename KernelType>
3749-
void throwOnKernelParameterMisuse() const {
3750-
using NameT =
3751-
typename detail::get_kernel_name_t<KernelName, KernelType>::name;
3752-
for (unsigned I = 0; I < detail::getKernelNumParams<NameT>(); ++I) {
3753-
const detail::kernel_param_desc_t ParamDesc =
3754-
detail::getKernelParamDesc<NameT>(I);
3748+
//
3749+
// Exception handling generates lots of code, outline it out of template
3750+
// method to improve compilation times.
3751+
void throwOnKernelParameterMisuseHelper(
3752+
int N, detail::kernel_param_desc_t (*f)(int)) const {
3753+
for (int I = 0; I < N; ++I) {
3754+
detail::kernel_param_desc_t ParamDesc = (*f)(I);
37553755
const detail::kernel_param_kind_t &Kind = ParamDesc.kind;
37563756
const access::target AccTarget =
37573757
static_cast<access::target>(ParamDesc.info & AccessTargetMask);
@@ -3770,6 +3770,13 @@ class __SYCL_EXPORT handler {
37703770
"of parallel_for that takes a range parameter.");
37713771
}
37723772
}
3773+
template <typename KernelName, typename KernelType>
3774+
void throwOnKernelParameterMisuse() const {
3775+
using NameT =
3776+
typename detail::get_kernel_name_t<KernelName, KernelType>::name;
3777+
throwOnKernelParameterMisuseHelper(detail::getKernelNumParams<NameT>(),
3778+
&detail::getKernelParamDesc<NameT>);
3779+
}
37733780

37743781
template <typename T, int Dims, access::mode AccessMode,
37753782
access::target AccessTarget,

0 commit comments

Comments
 (0)