Skip to content

Commit 053b267

Browse files
vinser52vinser52
committed
Move NDRDescT and ArgDesc to separate headers
1 parent 2af08ff commit 053b267

File tree

4 files changed

+152
-109
lines changed

4 files changed

+152
-109
lines changed

‎sycl/source/detail/cg.hpp‎

Lines changed: 2 additions & 108 deletions
Original file line numberDiff line numberDiff line change
@@ -21,6 +21,8 @@
2121
#include <sycl/kernel_bundle.hpp> // for kernel_bundle_impl
2222

2323
#include <detail/device_kernel_info.hpp>
24+
#include <detail/kernel_arg_desc.hpp>
25+
#include <detail/ndrange_desc.hpp>
2426

2527
#include <assert.h> // for assert
2628
#include <memory> // for shared_ptr, unique_ptr
@@ -49,114 +51,6 @@ class stream_impl;
4951
class queue_impl;
5052
class kernel_bundle_impl;
5153

52-
// The structure represents kernel argument.
53-
class ArgDesc {
54-
public:
55-
ArgDesc(sycl::detail::kernel_param_kind_t Type, void *Ptr, int Size,
56-
int Index)
57-
: MType(Type), MPtr(Ptr), MSize(Size), MIndex(Index) {}
58-
59-
sycl::detail::kernel_param_kind_t MType;
60-
void *MPtr;
61-
int MSize;
62-
int MIndex;
63-
};
64-
65-
// The structure represents NDRange - global, local sizes, global offset and
66-
// number of dimensions.
67-
68-
// TODO: A lot of tests rely on particular values to be set for dimensions that
69-
// are not used. To clarify, for example, if a 2D kernel is invoked, in
70-
// NDRDescT, the value of index 2 in GlobalSize must be set to either 1 or 0
71-
// depending on which constructor is used for no clear reason.
72-
// Instead, only sensible defaults should be used and tests should be updated
73-
// to reflect this.
74-
class NDRDescT {
75-
76-
public:
77-
NDRDescT() = default;
78-
NDRDescT(const NDRDescT &Desc) = default;
79-
NDRDescT(NDRDescT &&Desc) = default;
80-
81-
template <int Dims_>
82-
NDRDescT(sycl::range<Dims_> N, bool SetNumWorkGroups) : Dims{size_t(Dims_)} {
83-
if (SetNumWorkGroups) {
84-
for (size_t I = 0; I < Dims_; ++I) {
85-
NumWorkGroups[I] = N[I];
86-
}
87-
} else {
88-
for (size_t I = 0; I < Dims_; ++I) {
89-
GlobalSize[I] = N[I];
90-
}
91-
92-
for (int I = Dims_; I < 3; ++I) {
93-
GlobalSize[I] = 1;
94-
}
95-
}
96-
}
97-
98-
template <int Dims_>
99-
NDRDescT(sycl::range<Dims_> NumWorkItems, sycl::range<Dims_> LocalSizes,
100-
sycl::id<Dims_> Offset)
101-
: Dims{size_t(Dims_)} {
102-
for (size_t I = 0; I < Dims_; ++I) {
103-
GlobalSize[I] = NumWorkItems[I];
104-
LocalSize[I] = LocalSizes[I];
105-
GlobalOffset[I] = Offset[I];
106-
}
107-
108-
for (int I = Dims_; I < 3; ++I) {
109-
LocalSize[I] = LocalSizes[0] ? 1 : 0;
110-
}
111-
112-
for (int I = Dims_; I < 3; ++I) {
113-
GlobalSize[I] = 1;
114-
}
115-
}
116-
117-
template <int Dims_>
118-
NDRDescT(sycl::range<Dims_> NumWorkItems, sycl::id<Dims_> Offset)
119-
: Dims{size_t(Dims_)} {
120-
for (size_t I = 0; I < Dims_; ++I) {
121-
GlobalSize[I] = NumWorkItems[I];
122-
GlobalOffset[I] = Offset[I];
123-
}
124-
}
125-
126-
template <int Dims_>
127-
NDRDescT(sycl::nd_range<Dims_> ExecutionRange)
128-
: NDRDescT(ExecutionRange.get_global_range(),
129-
ExecutionRange.get_local_range(),
130-
ExecutionRange.get_offset()) {}
131-
132-
template <int Dims_>
133-
NDRDescT(sycl::range<Dims_> Range)
134-
: NDRDescT(Range, /*SetNumWorkGroups=*/false) {}
135-
136-
template <int Dims_> void setClusterDimensions(sycl::range<Dims_> N) {
137-
if (this->Dims != size_t(Dims_)) {
138-
throw std::runtime_error(
139-
"Dimensionality of cluster, global and local ranges must be same");
140-
}
141-
142-
for (int I = 0; I < Dims_; ++I)
143-
ClusterDimensions[I] = N[I];
144-
}
145-
146-
NDRDescT &operator=(const NDRDescT &Desc) = default;
147-
NDRDescT &operator=(NDRDescT &&Desc) = default;
148-
149-
std::array<size_t, 3> GlobalSize{0, 0, 0};
150-
std::array<size_t, 3> LocalSize{0, 0, 0};
151-
std::array<size_t, 3> GlobalOffset{0, 0, 0};
152-
/// Number of workgroups, used to record the number of workgroups from the
153-
/// simplest form of parallel_for_work_group. If set, all other fields must be
154-
/// zero
155-
std::array<size_t, 3> NumWorkGroups{0, 0, 0};
156-
std::array<size_t, 3> ClusterDimensions{1, 1, 1};
157-
size_t Dims = 0;
158-
};
159-
16054
/// Base class for all types of command groups.
16155
class CG {
16256
public:

‎sycl/source/detail/kernel_arg_desc.hpp‎

Lines changed: 32 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,32 @@
1+
//==-------------- CG.hpp - SYCL standard header file ----------------------==//
2+
//
3+
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4+
// See https://llvm.org/LICENSE.txt for license information.
5+
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6+
//
7+
//===----------------------------------------------------------------------===//
8+
9+
#pragma once
10+
11+
#include <sycl/detail/kernel_desc.hpp>
12+
13+
namespace sycl {
14+
inline namespace _V1 {
15+
namespace detail {
16+
17+
// The structure represents kernel argument.
18+
class ArgDesc {
19+
public:
20+
ArgDesc(sycl::detail::kernel_param_kind_t Type, void *Ptr, int Size,
21+
int Index)
22+
: MType(Type), MPtr(Ptr), MSize(Size), MIndex(Index) {}
23+
24+
sycl::detail::kernel_param_kind_t MType;
25+
void *MPtr;
26+
int MSize;
27+
int MIndex;
28+
};
29+
30+
} // namespace detail
31+
} // namespace _V1
32+
} // namespace sycl

‎sycl/source/detail/kernel_data.hpp‎

Lines changed: 2 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -8,9 +8,10 @@
88

99
#pragma once
1010

11-
#include <detail/cg.hpp>
1211
#include <detail/device_kernel_info.hpp>
1312
#include <detail/graph/dynamic_impl.hpp>
13+
#include <detail/kernel_arg_desc.hpp>
14+
#include <detail/ndrange_desc.hpp>
1415

1516
#include <sycl/detail/kernel_desc.hpp>
1617

‎sycl/source/detail/ndrange_desc.hpp‎

Lines changed: 116 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,116 @@
1+
//==-------------- CG.hpp - SYCL standard header file ----------------------==//
2+
//
3+
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4+
// See https://llvm.org/LICENSE.txt for license information.
5+
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6+
//
7+
//===----------------------------------------------------------------------===//
8+
9+
#pragma once
10+
11+
#include <sycl/nd_range.hpp>
12+
#include <sycl/range.hpp>
13+
14+
#include <array>
15+
16+
namespace sycl {
17+
inline namespace _V1 {
18+
namespace detail {
19+
// The structure represents NDRange - global, local sizes, global offset and
20+
// number of dimensions.
21+
22+
// TODO: A lot of tests rely on particular values to be set for dimensions that
23+
// are not used. To clarify, for example, if a 2D kernel is invoked, in
24+
// NDRDescT, the value of index 2 in GlobalSize must be set to either 1 or 0
25+
// depending on which constructor is used for no clear reason.
26+
// Instead, only sensible defaults should be used and tests should be updated
27+
// to reflect this.
28+
class NDRDescT {
29+
30+
public:
31+
NDRDescT() = default;
32+
NDRDescT(const NDRDescT &Desc) = default;
33+
NDRDescT(NDRDescT &&Desc) = default;
34+
35+
template <int Dims_>
36+
NDRDescT(sycl::range<Dims_> N, bool SetNumWorkGroups) : Dims{size_t(Dims_)} {
37+
if (SetNumWorkGroups) {
38+
for (size_t I = 0; I < Dims_; ++I) {
39+
NumWorkGroups[I] = N[I];
40+
}
41+
} else {
42+
for (size_t I = 0; I < Dims_; ++I) {
43+
GlobalSize[I] = N[I];
44+
}
45+
46+
for (int I = Dims_; I < 3; ++I) {
47+
GlobalSize[I] = 1;
48+
}
49+
}
50+
}
51+
52+
template <int Dims_>
53+
NDRDescT(sycl::range<Dims_> NumWorkItems, sycl::range<Dims_> LocalSizes,
54+
sycl::id<Dims_> Offset)
55+
: Dims{size_t(Dims_)} {
56+
for (size_t I = 0; I < Dims_; ++I) {
57+
GlobalSize[I] = NumWorkItems[I];
58+
LocalSize[I] = LocalSizes[I];
59+
GlobalOffset[I] = Offset[I];
60+
}
61+
62+
for (int I = Dims_; I < 3; ++I) {
63+
LocalSize[I] = LocalSizes[0] ? 1 : 0;
64+
}
65+
66+
for (int I = Dims_; I < 3; ++I) {
67+
GlobalSize[I] = 1;
68+
}
69+
}
70+
71+
template <int Dims_>
72+
NDRDescT(sycl::range<Dims_> NumWorkItems, sycl::id<Dims_> Offset)
73+
: Dims{size_t(Dims_)} {
74+
for (size_t I = 0; I < Dims_; ++I) {
75+
GlobalSize[I] = NumWorkItems[I];
76+
GlobalOffset[I] = Offset[I];
77+
}
78+
}
79+
80+
template <int Dims_>
81+
NDRDescT(sycl::nd_range<Dims_> ExecutionRange)
82+
: NDRDescT(ExecutionRange.get_global_range(),
83+
ExecutionRange.get_local_range(),
84+
ExecutionRange.get_offset()) {}
85+
86+
template <int Dims_>
87+
NDRDescT(sycl::range<Dims_> Range)
88+
: NDRDescT(Range, /*SetNumWorkGroups=*/false) {}
89+
90+
template <int Dims_> void setClusterDimensions(sycl::range<Dims_> N) {
91+
if (this->Dims != size_t(Dims_)) {
92+
throw std::runtime_error(
93+
"Dimensionality of cluster, global and local ranges must be same");
94+
}
95+
96+
for (int I = 0; I < Dims_; ++I)
97+
ClusterDimensions[I] = N[I];
98+
}
99+
100+
NDRDescT &operator=(const NDRDescT &Desc) = default;
101+
NDRDescT &operator=(NDRDescT &&Desc) = default;
102+
103+
std::array<size_t, 3> GlobalSize{0, 0, 0};
104+
std::array<size_t, 3> LocalSize{0, 0, 0};
105+
std::array<size_t, 3> GlobalOffset{0, 0, 0};
106+
/// Number of workgroups, used to record the number of workgroups from the
107+
/// simplest form of parallel_for_work_group. If set, all other fields must be
108+
/// zero
109+
std::array<size_t, 3> NumWorkGroups{0, 0, 0};
110+
std::array<size_t, 3> ClusterDimensions{1, 1, 1};
111+
size_t Dims = 0;
112+
};
113+
114+
} // namespace detail
115+
} // namespace _V1
116+
} // namespace sycl

0 commit comments

Comments
 (0)