From 6535f35710daa50fd08cf657d4c815b0fcfc5f44 Mon Sep 17 00:00:00 2001 From: "Hu, Peisen" Date: Tue, 28 Oct 2025 06:40:56 -0700 Subject: [PATCH 1/5] [SYCL] Update enqueue_functions.hpp Signed-off-by: Hu, Peisen --- .../ext/oneapi/experimental/enqueue_functions.hpp | 11 +++++++++-- 1 file changed, 9 insertions(+), 2 deletions(-) diff --git a/sycl/include/sycl/ext/oneapi/experimental/enqueue_functions.hpp b/sycl/include/sycl/ext/oneapi/experimental/enqueue_functions.hpp index e393bd626d4d..a6754550a5de 100644 --- a/sycl/include/sycl/ext/oneapi/experimental/enqueue_functions.hpp +++ b/sycl/include/sycl/ext/oneapi/experimental/enqueue_functions.hpp @@ -251,8 +251,15 @@ template void nd_launch(handler &CGH, nd_range Range, const KernelType &KernelObj, ReductionsT &&...Reductions) { - CGH.parallel_for(Range, std::forward(Reductions)..., - KernelObj); + if constexpr (ext::oneapi::experimental::detail::HasKernelPropertiesGetMethod< + const KernelType &>::value) { + CGH.parallel_for( + Range, KernelObj.get(ext::oneapi::experimental::properties_tag{}), + std::forward(Reductions)..., KernelObj); + } else { + CGH.parallel_for( + Range, std::forward(Reductions)..., KernelObj); + } } template Date: Tue, 28 Oct 2025 07:55:09 -0700 Subject: [PATCH 2/5] [SYCL] Revise tset case Signed-off-by: Hu, Peisen --- .../properties_kernel_device_has.cpp | 102 ++++++++++++------ 1 file changed, 69 insertions(+), 33 deletions(-) diff --git a/sycl/test/check_device_code/extensions/properties/properties_kernel_device_has.cpp b/sycl/test/check_device_code/extensions/properties/properties_kernel_device_has.cpp index 3d1c528744af..018115a0ef2b 100644 --- a/sycl/test/check_device_code/extensions/properties/properties_kernel_device_has.cpp +++ b/sycl/test/check_device_code/extensions/properties/properties_kernel_device_has.cpp @@ -27,6 +27,37 @@ static constexpr auto device_has_all = device_has< aspect::usm_host_allocations, aspect::usm_shared_allocations, aspect::ext_intel_free_memory, aspect::ext_intel_device_id>; +struct TestKernelHasDevice { + void operator()() const {} + auto get(properties_tag) const { return properties{device_has_all}; } +}; + +struct TestKernelHasDevice_id1 { + void operator()(id<1>) const {} + auto get(properties_tag) const { return properties{device_has_all}; } +}; + +struct TestKernelHasDevice_id1_1 { + template void operator()(id<1>, T1 &) const {} + auto get(properties_tag) const { return properties{device_has_all}; } +}; + +struct TestKernelHasDevice_nd_item1 { + void operator()(nd_item<1>) const {} + auto get(properties_tag) const { return properties{device_has_all}; } +}; + +struct TestKernelHasDevice_nd_item1_1 { + template void operator()(nd_item<1>, T1 &) const {} + auto get(properties_tag) const { return properties{device_has_all}; } +}; + +struct TestKernelHasDevice_nd_item1_2 { + template + void operator()(nd_item<1>, T1 &, T2 &) const {} + auto get(properties_tag) const { return properties{device_has_all}; } +}; + int main() { queue Q; event Ev; @@ -40,18 +71,18 @@ int main() { auto Redu2 = reduction(nullptr, multiplies()); // CHECK-IR: spir_kernel void @{{.*}}WGSizeKernel0(){{.*}} #[[DHAttr1:[0-9]+]] - Q.single_task(Props, []() {}); + Q.single_task(TestKernelHasDevice{}); // CHECK-IR: spir_kernel void @{{.*}}WGSizeKernel1(){{.*}} #[[DHAttr1]] - Q.single_task(Ev, Props, []() {}); + Q.single_task(Ev, TestKernelHasDevice{}); // CHECK-IR: spir_kernel void @{{.*}}WGSizeKernel2(){{.*}} #[[DHAttr1]] - Q.single_task({Ev}, Props, []() {}); + Q.single_task({Ev}, TestKernelHasDevice{}); // CHECK-IR: spir_kernel void @{{.*}}WGSizeKernel3(){{.*}} #[[DHAttr2:[0-9]+]] - Q.parallel_for(R1, Props, [](id<1>) {}); + Q.parallel_for(R1, TestKernelHasDevice_id1{}); // CHECK-IR: spir_kernel void @{{.*}}WGSizeKernel4(){{.*}} #[[DHAttr2]] - Q.parallel_for(R1, Ev, Props, [](id<1>) {}); + Q.parallel_for(R1, Ev, TestKernelHasDevice_id1{}); // CHECK-IR: spir_kernel void @{{.*}}WGSizeKernel5(){{.*}} #[[DHAttr2]] - Q.parallel_for(R1, {Ev}, Props, [](id<1>) {}); + Q.parallel_for(R1, {Ev}, TestKernelHasDevice_id1{}); // CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}WGSizeKernel6{{.*}}{{.*}} #[[DHAttr2:[0-9]+]] Q.parallel_for(R1, Props, Redu1, [](id<1>, auto &) {}); @@ -70,57 +101,62 @@ int main() { Q.parallel_for(NDR1, {Ev}, Props, [](nd_item<1>) {}); // CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}WGSizeKernel12{{.*}}{{.*}} #[[DHAttr2]] - Q.parallel_for(NDR1, Props, Redu1, - [](nd_item<1>, auto &) {}); + nd_launch(Q, NDR1, TestKernelHasDevice_nd_item1_1{}, + Redu1); // CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}WGSizeKernel13{{.*}}{{.*}} #[[DHAttr2]] - Q.parallel_for(NDR1, Ev, Props, Redu1, - [](nd_item<1>, auto &) {}); + Q.submit([&](sycl::handler &CGH) { + CGH.depends_on(Ev); + nd_launch(CGH, NDR1, TestKernelHasDevice_nd_item1_1{}, + Redu1); + }); // CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}WGSizeKernel14{{.*}}{{.*}} #[[DHAttr2]] - Q.parallel_for(NDR1, {Ev}, Props, Redu1, - [](nd_item<1>, auto &) {}); + Q.submit([&](sycl::handler &CGH) { + CGH.depends_on({Ev}); + nd_launch(CGH, NDR1, TestKernelHasDevice_nd_item1_1{}, + Redu1); + }); // CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}WGSizeKernel15{{.*}}{{.*}} #[[DHAttr2]] - Q.parallel_for(NDR1, Props, Redu1, Redu2, - [](nd_item<1>, auto &, auto &) {}); + nd_launch(Q, NDR1, TestKernelHasDevice_nd_item1_2{}, + Redu1, Redu2); // CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}WGSizeKernel16{{.*}}{{.*}} #[[DHAttr2]] - Q.parallel_for(NDR1, Ev, Props, Redu1, Redu2, - [](nd_item<1>, auto &, auto &) {}); + Q.submit([&](sycl::handler &CGH) { + CGH.depends_on(Ev); + nd_launch(CGH, NDR1, TestKernelHasDevice_nd_item1_2{}, + Redu1, Redu2); + }); // CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}WGSizeKernel17{{.*}}{{.*}} #[[DHAttr2]] - Q.parallel_for(NDR1, {Ev}, Props, Redu1, Redu2, - [](nd_item<1>, auto &, auto &) {}); + Q.submit([&](sycl::handler &CGH) { + CGH.depends_on({Ev}); + nd_launch(CGH, NDR1, TestKernelHasDevice_nd_item1_2{}, + Redu1, Redu2); + }); // CHECK-IR: spir_kernel void @{{.*}}WGSizeKernel18(){{.*}} #[[DHAttr1]] Q.submit([&](handler &CGH) { - CGH.single_task(Props, []() {}); + CGH.single_task(TestKernelHasDevice{}); }); // CHECK-IR: spir_kernel void @{{.*}}WGSizeKernel19(){{.*}} #[[DHAttr2]] Q.submit([&](handler &CGH) { - CGH.parallel_for(R1, Props, [](id<1>) {}); + CGH.parallel_for(R1, TestKernelHasDevice_id1{}); }); // CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}WGSizeKernel20{{.*}}{{.*}} #[[DHAttr2]] Q.submit([&](handler &CGH) { - CGH.parallel_for(R1, Props, Redu1, - [](id<1>, auto &) {}); + CGH.parallel_for(R1, Props, Redu1, [](id<1>, auto &) { + }); // note: this one still doesn't work }); // CHECK-IR: spir_kernel void @{{.*}}WGSizeKernel21(){{.*}} #[[DHAttr2]] Q.submit([&](handler &CGH) { - CGH.parallel_for(NDR1, Props, [](nd_item<1>) {}); + CGH.parallel_for(NDR1, + TestKernelHasDevice_nd_item1{}); }); - // CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}WGSizeKernel22{{.*}}{{.*}} #[[DHAttr2]] - Q.submit([&](handler &CGH) { - CGH.parallel_for(NDR1, Props, Redu1, - [](nd_item<1>, auto &) {}); - }); + // DUPLICATE, REMOVED - // CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}WGSizeKernel23{{.*}}{{.*}} #[[DHAttr2]] - Q.submit([&](handler &CGH) { - CGH.parallel_for(NDR1, Props, Redu1, Redu2, - [](nd_item<1>, auto &, auto &) {}); - }); + // DUPLICATE, REMOVED // CHECK-IR: spir_kernel void @{{.*}}WGSizeKernel24(){{.*}} #[[DHAttr2]] Q.submit([&](handler &CGH) { From a31dbb50712c2075a7f69bddca74971580582f7a Mon Sep 17 00:00:00 2001 From: "Hu, Peisen" Date: Tue, 4 Nov 2025 05:18:26 -0800 Subject: [PATCH 3/5] [SYCL] Fix problem at lower level Signed-off-by: Hu, Peisen --- .../oneapi/experimental/enqueue_functions.hpp | 11 +--- sycl/include/sycl/handler.hpp | 66 +++++++++++++++---- .../properties_kernel_device_has.cpp | 63 ++++++++++++------ 3 files changed, 98 insertions(+), 42 deletions(-) diff --git a/sycl/include/sycl/ext/oneapi/experimental/enqueue_functions.hpp b/sycl/include/sycl/ext/oneapi/experimental/enqueue_functions.hpp index e8d1b2be337a..8c8488a99e35 100644 --- a/sycl/include/sycl/ext/oneapi/experimental/enqueue_functions.hpp +++ b/sycl/include/sycl/ext/oneapi/experimental/enqueue_functions.hpp @@ -263,15 +263,8 @@ template void nd_launch(handler &CGH, nd_range Range, const KernelType &KernelObj, ReductionsT &&...Reductions) { - if constexpr (ext::oneapi::experimental::detail::HasKernelPropertiesGetMethod< - const KernelType &>::value) { - CGH.parallel_for( - Range, KernelObj.get(ext::oneapi::experimental::properties_tag{}), - std::forward(Reductions)..., KernelObj); - } else { - CGH.parallel_for( - Range, std::forward(Reductions)..., KernelObj); - } + CGH.parallel_for(Range, std::forward(Reductions)..., + KernelObj); } template (); #endif + detail::reduction_parallel_for(*this, Range, Properties, std::forward(Rest)...); } @@ -2172,6 +2173,7 @@ class __SYCL_EXPORT handler { throwIfGraphAssociated(); #endif + detail::reduction_parallel_for(*this, Range, Properties, std::forward(Rest)...); } @@ -2192,6 +2194,7 @@ class __SYCL_EXPORT handler { throwIfGraphAssociated(); #endif + detail::reduction_parallel_for(*this, Range, Properties, std::forward(Rest)...); } @@ -2200,27 +2203,54 @@ class __SYCL_EXPORT handler { std::enable_if_t::value && (sizeof...(RestT) > 1)> parallel_for(range<1> Range, RestT &&...Rest) { - parallel_for(Range, - ext::oneapi::experimental::empty_properties_t{}, - std::forward(Rest)...); + const auto &KernelObj = (Rest, ...); + if constexpr (ext::oneapi::experimental::detail:: + HasKernelPropertiesGetMethod< + decltype(KernelObj)>::value) { + parallel_for( + Range, KernelObj.get(ext::oneapi::experimental::properties_tag{}), + std::forward(Rest)...); + } else { + parallel_for(Range, + ext::oneapi::experimental::empty_properties_t{}, + std::forward(Rest)...); + } } template std::enable_if_t::value && (sizeof...(RestT) > 1)> parallel_for(range<2> Range, RestT &&...Rest) { - parallel_for(Range, - ext::oneapi::experimental::empty_properties_t{}, - std::forward(Rest)...); + const auto &KernelObj = (Rest, ...); + if constexpr (ext::oneapi::experimental::detail:: + HasKernelPropertiesGetMethod< + decltype(KernelObj)>::value) { + parallel_for( + Range, KernelObj.get(ext::oneapi::experimental::properties_tag{}), + std::forward(Rest)...); + } else { + parallel_for(Range, + ext::oneapi::experimental::empty_properties_t{}, + std::forward(Rest)...); + } } template std::enable_if_t::value && (sizeof...(RestT) > 1)> parallel_for(range<3> Range, RestT &&...Rest) { - parallel_for(Range, - ext::oneapi::experimental::empty_properties_t{}, - std::forward(Rest)...); + const auto &KernelObj = (Rest, ...); + if constexpr (ext::oneapi::experimental::detail:: + HasKernelPropertiesGetMethod< + decltype(KernelObj)>::value) { + parallel_for( + Range, KernelObj.get(ext::oneapi::experimental::properties_tag{}), + std::forward(Rest)...); + } else { + parallel_for(Range, + ext::oneapi::experimental::empty_properties_t{}, + std::forward(Rest)...); + } } template (); #endif + std::cout << "4\n"; + // static_assert(!std::is_same_v); //property + // already missing! detail::reduction_parallel_for(*this, Range, Properties, std::forward(Rest)...); } @@ -2247,9 +2280,18 @@ class __SYCL_EXPORT handler { typename... RestT> std::enable_if_t::value> parallel_for(nd_range Range, RestT &&...Rest) { - parallel_for(Range, - ext::oneapi::experimental::empty_properties_t{}, - std::forward(Rest)...); + const auto &KernelObj = (Rest, ...); + if constexpr (ext::oneapi::experimental::detail:: + HasKernelPropertiesGetMethod< + decltype(KernelObj)>::value) { + parallel_for( + Range, KernelObj.get(ext::oneapi::experimental::properties_tag{}), + std::forward(Rest)...); + } else { + parallel_for(Range, + ext::oneapi::experimental::empty_properties_t{}, + std::forward(Rest)...); + } } /// }@ diff --git a/sycl/test/check_device_code/extensions/properties/properties_kernel_device_has.cpp b/sycl/test/check_device_code/extensions/properties/properties_kernel_device_has.cpp index 018115a0ef2b..5f2933440b42 100644 --- a/sycl/test/check_device_code/extensions/properties/properties_kernel_device_has.cpp +++ b/sycl/test/check_device_code/extensions/properties/properties_kernel_device_has.cpp @@ -1,8 +1,5 @@ -// TODO: Currently using the -Wno-deprecated-declarations flag due to issue -// https://github.com/intel/llvm/issues/16320. Remove the flag once the issue is -// resolved. -// RUN: %clangxx -fsycl-device-only -S -Wno-deprecated-declarations -Xclang -emit-llvm -Xclang -disable-llvm-passes %s -o - | FileCheck %s --check-prefix CHECK-IR -// RUN: %clangxx -fsycl -fsyntax-only -Wno-deprecated-declarations -Xclang -verify %s +// RUN: %clangxx -fsycl-device-only -S -Xclang -emit-llvm -Xclang -disable-llvm-passes %s -o - | FileCheck %s --check-prefix CHECK-IR +// RUN: %clangxx -fsycl -fsyntax-only -Xclang -verify %s // expected-no-diagnostics #include @@ -58,6 +55,13 @@ struct TestKernelHasDevice_nd_item1_2 { auto get(properties_tag) const { return properties{device_has_all}; } }; +struct TestKernelHasDevice_work_group { + void operator()(group<1> G) const { + G.parallel_for_work_item([&](h_item<1>) {}); + } + auto get(properties_tag) const { return properties{device_has_all}; } +}; + int main() { queue Q; event Ev; @@ -65,8 +69,6 @@ int main() { range<1> R1{1}; nd_range<1> NDR1{R1, R1}; - constexpr auto Props = properties{device_has_all}; - auto Redu1 = reduction(nullptr, plus()); auto Redu2 = reduction(nullptr, multiplies()); @@ -85,20 +87,32 @@ int main() { Q.parallel_for(R1, {Ev}, TestKernelHasDevice_id1{}); // CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}WGSizeKernel6{{.*}}{{.*}} #[[DHAttr2:[0-9]+]] - Q.parallel_for(R1, Props, Redu1, [](id<1>, auto &) {}); + parallel_for(Q, R1, TestKernelHasDevice_id1_1{}, Redu1); // CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}WGSizeKernel7{{.*}}{{.*}} #[[DHAttr2]] - Q.parallel_for(R1, Ev, Props, Redu1, - [](id<1>, auto &) {}); + Q.submit([&](sycl::handler &CGH) { + CGH.depends_on(Ev); + parallel_for(Q, R1, TestKernelHasDevice_id1_1{}, + Redu1); + }); // CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}WGSizeKernel8{{.*}}{{.*}} #[[DHAttr2]] - Q.parallel_for(R1, {Ev}, Props, Redu1, - [](id<1>, auto &) {}); + Q.submit([&](sycl::handler &CGH) { + CGH.depends_on({Ev}); + parallel_for(Q, R1, TestKernelHasDevice_id1_1{}, + Redu1); + }); // CHECK-IR: spir_kernel void @{{.*}}WGSizeKernel9(){{.*}} #[[DHAttr2]] - Q.parallel_for(NDR1, Props, [](nd_item<1>) {}); + nd_launch(Q, NDR1, TestKernelHasDevice_nd_item1{}); // CHECK-IR: spir_kernel void @{{.*}}WGSizeKernel10(){{.*}} #[[DHAttr2]] - Q.parallel_for(NDR1, Ev, Props, [](nd_item<1>) {}); + Q.submit([&](sycl::handler &CGH) { + CGH.depends_on(Ev); + nd_launch(CGH, NDR1, TestKernelHasDevice_nd_item1{}); + }); // CHECK-IR: spir_kernel void @{{.*}}WGSizeKernel11(){{.*}} #[[DHAttr2]] - Q.parallel_for(NDR1, {Ev}, Props, [](nd_item<1>) {}); + Q.submit([&](sycl::handler &CGH) { + CGH.depends_on({Ev}); + nd_launch(CGH, NDR1, TestKernelHasDevice_nd_item1{}); + }); // CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}WGSizeKernel12{{.*}}{{.*}} #[[DHAttr2]] nd_launch(Q, NDR1, TestKernelHasDevice_nd_item1_1{}, @@ -144,8 +158,8 @@ int main() { // CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}WGSizeKernel20{{.*}}{{.*}} #[[DHAttr2]] Q.submit([&](handler &CGH) { - CGH.parallel_for(R1, Props, Redu1, [](id<1>, auto &) { - }); // note: this one still doesn't work + CGH.parallel_for(R1, Redu1, + TestKernelHasDevice_id1_1{}); }); // CHECK-IR: spir_kernel void @{{.*}}WGSizeKernel21(){{.*}} #[[DHAttr2]] @@ -154,15 +168,22 @@ int main() { TestKernelHasDevice_nd_item1{}); }); - // DUPLICATE, REMOVED + // CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}WGSizeKernel22{{.*}}{{.*}} #[[DHAttr2]] + Q.submit([&](handler &CGH) { + CGH.parallel_for(NDR1, Redu1, + TestKernelHasDevice_nd_item1_1{}); + }); - // DUPLICATE, REMOVED + // CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}WGSizeKernel23{{.*}}{{.*}} #[[DHAttr2]] + Q.submit([&](handler &CGH) { + CGH.parallel_for(NDR1, Redu1, Redu2, + TestKernelHasDevice_nd_item1_2{}); + }); // CHECK-IR: spir_kernel void @{{.*}}WGSizeKernel24(){{.*}} #[[DHAttr2]] Q.submit([&](handler &CGH) { CGH.parallel_for_work_group( - R1, Props, - [](group<1> G) { G.parallel_for_work_item([&](h_item<1>) {}); }); + R1, TestKernelHasDevice_work_group{}); }); return 0; From 530e3c3a5940054c3db2c130274a68faea8e6b96 Mon Sep 17 00:00:00 2001 From: "Hu, Peisen" Date: Tue, 4 Nov 2025 05:21:10 -0800 Subject: [PATCH 4/5] [SYCL] Remove unrelated changes Signed-off-by: Hu, Peisen --- sycl/include/sycl/handler.hpp | 6 ------ 1 file changed, 6 deletions(-) diff --git a/sycl/include/sycl/handler.hpp b/sycl/include/sycl/handler.hpp index 7553ac4351cd..0d5b53051f77 100644 --- a/sycl/include/sycl/handler.hpp +++ b/sycl/include/sycl/handler.hpp @@ -2152,7 +2152,6 @@ class __SYCL_EXPORT handler { throwIfGraphAssociated(); #endif - detail::reduction_parallel_for(*this, Range, Properties, std::forward(Rest)...); } @@ -2173,7 +2172,6 @@ class __SYCL_EXPORT handler { throwIfGraphAssociated(); #endif - detail::reduction_parallel_for(*this, Range, Properties, std::forward(Rest)...); } @@ -2194,7 +2192,6 @@ class __SYCL_EXPORT handler { throwIfGraphAssociated(); #endif - detail::reduction_parallel_for(*this, Range, Properties, std::forward(Rest)...); } @@ -2269,9 +2266,6 @@ class __SYCL_EXPORT handler { throwIfGraphAssociated(); #endif - std::cout << "4\n"; - // static_assert(!std::is_same_v); //property - // already missing! detail::reduction_parallel_for(*this, Range, Properties, std::forward(Rest)...); } From 969d606f86eadc07375d9291bfc160ddf4f58c3f Mon Sep 17 00:00:00 2001 From: "Hu, Peisen" Date: Tue, 4 Nov 2025 06:36:02 -0800 Subject: [PATCH 5/5] [SYCL] Address nd_range variant without reductions Signed-off-by: Hu, Peisen --- sycl/include/sycl/handler.hpp | 13 ++++++++++++- 1 file changed, 12 insertions(+), 1 deletion(-) diff --git a/sycl/include/sycl/handler.hpp b/sycl/include/sycl/handler.hpp index 0d5b53051f77..98742a4cd76b 100644 --- a/sycl/include/sycl/handler.hpp +++ b/sycl/include/sycl/handler.hpp @@ -2272,7 +2272,8 @@ class __SYCL_EXPORT handler { template - std::enable_if_t::value> + std::enable_if_t::value && + (sizeof...(RestT) > 1)> // variant with reductions parallel_for(nd_range Range, RestT &&...Rest) { const auto &KernelObj = (Rest, ...); if constexpr (ext::oneapi::experimental::detail:: @@ -2288,6 +2289,16 @@ class __SYCL_EXPORT handler { } } + template + std::enable_if_t::value && + (sizeof...(RestT) == 1)> // variant without reductions + parallel_for(nd_range Range, RestT &&...Rest) { + parallel_for(Range, + ext::oneapi::experimental::empty_properties_t{}, + std::forward(Rest)...); + } + /// }@ template