Skip to content

Commit 50b264d

Browse files
lbushi25lbushi25
committed
Add specialization of is_device_copyable in integration header
1 parent 2f80b12 commit 50b264d

File tree

3 files changed

+63
-14
lines changed

3 files changed

+63
-14
lines changed

‎clang/lib/Sema/SemaSYCL.cpp‎

Lines changed: 9 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -7293,6 +7293,15 @@ void SYCLIntegrationHeader::emit(raw_ostream &O) {
72937293
// contains information about the offset, size and parameter
72947294
// kind of every field inside the struct at any nesting level
72957295
// This facilitates setting the arguments in the runtime.
7296+
// We also define is_device_copyable trait to be true for this type to
7297+
// allow it being passed in device kernels.
7298+
O << "template <>\n";
7299+
O << "struct "
7300+
"sycl::detail::is_device_copyable<";
7301+
Policy.SuppressTagKeyword = true;
7302+
type.print(O, Policy);
7303+
O << ">: std::true_type {};\n";
7304+
72967305
O << "template <>\n";
72977306
O << "struct "
72987307
"sycl::ext::oneapi::experimental::detail::is_struct_with_special_"

‎clang/test/CodeGenSYCL/free_function_int_header.cpp‎

Lines changed: 50 additions & 11 deletions
Original file line numberDiff line numberDiff line change
@@ -299,6 +299,12 @@ struct SecondLevelAccessor {
299299
AccessorAndInt accAndInt;
300300
};
301301

302+
template <typename T>
303+
struct TemplatedAccessorStruct {
304+
sycl::accessor<T, 1, sycl::access::mode::read_write> acc;
305+
sycl::local_accessor<T, 1> lacc;
306+
};
307+
302308
[[__sycl_detail__::add_ir_attributes_function("sycl-single-task-kernel", 0)]]
303309
void ff_25(AccessorAndLocalAccessor arg1) {
304310
}
@@ -311,6 +317,10 @@ void ff_26(AccessorAndLocalAccessor arg1, SecondLevelAccessor arg2) {
311317
void ff_27(IntAndAccessor arg1, AccessorAndInt) {
312318
}
313319

320+
[[__sycl_detail__::add_ir_attributes_function("sycl-single-task-kernel", 0)]]
321+
void ff_28(TemplatedAccessorStruct<int> arg1) {
322+
}
323+
314324
// CHECK: const char* const kernel_names[] = {
315325
// CHECK-NEXT: {{.*}}__sycl_kernel_ff_2Piii
316326
// CHECK-NEXT: {{.*}}__sycl_kernel_ff_2Piiii
@@ -349,6 +359,7 @@ void ff_27(IntAndAccessor arg1, AccessorAndInt) {
349359
// CHECK-NEXT: {{.*}}__sycl_kernel_ff_2524AccessorAndLocalAccessor",
350360
// CHECK-NEXT: {{.*}}__sycl_kernel_ff_2624AccessorAndLocalAccessor19SecondLevelAccessor",
351361
// CHECK-NEXT: {{.*}}__sycl_kernel_ff_2714IntAndAccessor14AccessorAndInt",
362+
// CHECK-NEXT: {{.*}}__sycl_kernel_ff_2823TemplatedAccessorStructIiE",
352363

353364
// CHECK-NEXT: {{.*}}__sycl_kernel_ff_23i"
354365

@@ -494,6 +505,11 @@ void ff_27(IntAndAccessor arg1, AccessorAndInt) {
494505
// CHECK-NEXT: { kernel_param_kind_t::kind_accessor, 4062, 0 },
495506
// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 4, 12 },
496507

508+
// CHECK: //--- _Z19__sycl_kernel_ff_2823TemplatedAccessorStructIiE
509+
// CHECK-NEXT: { kernel_param_kind_t::kind_struct_with_special_type, 36, 0 },
510+
// CHECK-NEXT: { kernel_param_kind_t::kind_accessor, 4062, 0 },
511+
// CHECK-NEXT: { kernel_param_kind_t::kind_accessor, 4064, 12 },
512+
497513
// CHECK: //--- _Z19__sycl_kernel_ff_23i
498514
// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 4, 0 },
499515

@@ -1655,11 +1671,6 @@ void ff_27(IntAndAccessor arg1, AccessorAndInt) {
16551671
// CHECK-NEXT: sycl::detail::kernel_param_kind_t::kind_invalid };
16561672
// CHECK-NEXT: };
16571673

1658-
// CHECK: template <>
1659-
// CHECK-NEXT: struct sycl::ext::oneapi::experimental::detail::is_struct_with_special_type<AccessorAndInt> {
1660-
// CHECK-NEXT: inline static constexpr bool value = true;
1661-
// CHECK-NEXT: };
1662-
16631674
// CHECK: static constexpr auto __sycl_shim35() {
16641675
// CHECK-NEXT: return (void (*)(struct IntAndAccessor, struct AccessorAndInt))ff_27;
16651676
// CHECK-NEXT: }
@@ -1672,18 +1683,46 @@ void ff_27(IntAndAccessor arg1, AccessorAndInt) {
16721683
// CHECK-NEXT: static constexpr bool value = true;
16731684
// CHECK-NEXT: };
16741685

1686+
1687+
// CHECK: Definition of _Z19__sycl_kernel_ff_2823TemplatedAccessorStructIiE as a free function kernel
1688+
// CHECK: Forward declarations of kernel and its argument types:
1689+
// CHECK: template <typename T> struct TemplatedAccessorStruct;
1690+
// CHECK: void ff_28(TemplatedAccessorStruct<int> arg1);
1691+
// CHECK-NEXT: template <>
1692+
// CHECK-NEXT: struct sycl::ext::oneapi::experimental::detail::is_struct_with_special_type<TemplatedAccessorStruct<int>> {
1693+
// CHECK-NEXT: inline static constexpr bool value = true;
1694+
// CHECK-NEXT: static constexpr int offsets[] = { 0, 12, -1};
1695+
// CHECK-NEXT: static constexpr int sizes[] = { 4062, 4064, -1};
1696+
// CHECK-NEXT: static constexpr sycl::detail::kernel_param_kind_t kinds[] = {
1697+
// CHECK-NEXT: sycl::detail::kernel_param_kind_t::kind_accessor,
1698+
// CHECK-NEXT: sycl::detail::kernel_param_kind_t::kind_accessor,
1699+
// CHECK-NEXT sycl::detail::kernel_param_kind_t::kind_invalid };
1700+
// CHECK-NEXT: };
1701+
1702+
// CHECK: static constexpr auto __sycl_shim36() {
1703+
// CHECK-NEXT: return (void (*)(struct TemplatedAccessorStruct<int>))ff_28;
1704+
// CHECK-NEXT: }
1705+
1706+
// CHECK: struct ext::oneapi::experimental::is_kernel<__sycl_shim36()> {
1707+
// CHECK-NEXT: static constexpr bool value = true;
1708+
// CHECK-NEXT: };
1709+
// CHECK-NEXT: template <>
1710+
// CHECK-NEXT: struct ext::oneapi::experimental::is_single_task_kernel<__sycl_shim36()> {
1711+
// CHECK-NEXT: static constexpr bool value = true;
1712+
// CHECK-NEXT: };
1713+
16751714
// CHECK: Definition of _Z19__sycl_kernel_ff_23i as a free function kernel
16761715
// CHECK: Forward declarations of kernel and its argument types:
16771716
// CHECK: void ff_23(int arg);
1678-
// CHECK-NEXT: static constexpr auto __sycl_shim36() {
1717+
// CHECK-NEXT: static constexpr auto __sycl_shim37() {
16791718
// CHECK-NEXT: return (void (*)(int))ff_23;
16801719
// CHECK-NEXT: }
16811720

16821721
// CHECK: namespace sycl {
16831722
// CHECK-NEXT: inline namespace _V1 {
16841723
// CHECK-NEXT: namespace detail {
1685-
// CHECK-NEXT: //Free Function Kernel info specialization for shim36
1686-
// CHECK-NEXT: template <> struct FreeFunctionInfoData<__sycl_shim36()> {
1724+
// CHECK-NEXT: //Free Function Kernel info specialization for shim37
1725+
// CHECK-NEXT: template <> struct FreeFunctionInfoData<__sycl_shim37()> {
16871726
// CHECK-NEXT: __SYCL_DLL_LOCAL
16881727
// CHECK-NEXT: static constexpr unsigned getNumParams() { return 1; }
16891728
// CHECK-NEXT: __SYCL_DLL_LOCAL
@@ -1695,11 +1734,11 @@ void ff_27(IntAndAccessor arg1, AccessorAndInt) {
16951734

16961735
// CHECK: namespace sycl {
16971736
// CHECK-NEXT: template <>
1698-
// CHECK-NEXT: struct ext::oneapi::experimental::is_kernel<__sycl_shim36()> {
1737+
// CHECK-NEXT: struct ext::oneapi::experimental::is_kernel<__sycl_shim37()> {
16991738
// CHECK-NEXT: static constexpr bool value = true;
17001739
// CHECK-NEXT: };
17011740
// CHECK-NEXT: template <>
1702-
// CHECK-NEXT: struct ext::oneapi::experimental::is_single_task_kernel<__sycl_shim36()> {
1741+
// CHECK-NEXT: struct ext::oneapi::experimental::is_single_task_kernel<__sycl_shim37()> {
17031742
// CHECK-NEXT: static constexpr bool value = true;
17041743
// CHECK-NEXT: };
17051744

@@ -1713,7 +1752,7 @@ void ff_27(IntAndAccessor arg1, AccessorAndInt) {
17131752
// CHECK-NEXT: namespace detail {
17141753
// CHECK-NEXT: struct GlobalMapUpdater {
17151754
// CHECK-NEXT: GlobalMapUpdater() {
1716-
// CHECK-NEXT: sycl::detail::free_function_info_map::add(sycl::detail::kernel_names, sycl::detail::kernel_args_sizes, 36);
1755+
// CHECK-NEXT: sycl::detail::free_function_info_map::add(sycl::detail::kernel_names, sycl::detail::kernel_args_sizes, 37);
17171756
// CHECK-NEXT: }
17181757
// CHECK-NEXT: };
17191758
// CHECK-NEXT: static GlobalMapUpdater updater;

‎sycl/include/sycl/handler.hpp‎

Lines changed: 4 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -1638,10 +1638,11 @@ class __SYCL_EXPORT handler {
16381638
std::is_pointer_v<remove_cv_ref_t<T>>) // USM
16391639
|| is_same_type<cl_mem, T>::value // Interop
16401640
|| is_same_type<stream, T>::value // Stream
1641-
|| ext::oneapi::experimental::detail::is_struct_with_special_type<
1642-
remove_cv_ref_t<T>>::value; // Structs that contain special types
1641+
||
1642+
sycl::is_device_copyable_v<remove_cv_ref_t<T>>; // Structs that contain
1643+
// special types
16431644
};
1644-
1645+
16451646
/// Sets argument for OpenCL interoperability kernels.
16461647
///
16471648
/// Registers Arg passed as argument # ArgIndex.

0 commit comments

Comments
 (0)