Skip to content

Commit 21b65bd

Browse files
committed
Implement kernel_dim and components range mappers
1 parent 687af1b commit 21b65bd

File tree

3 files changed

+91
-4
lines changed

3 files changed

+91
-4
lines changed

include/range_mapper.h

Lines changed: 45 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1,6 +1,7 @@
11
#pragma once
22

33
#include <stdexcept>
4+
#include <tuple>
45
#include <type_traits>
56

67
#include <CL/sycl.hpp>
@@ -159,7 +160,7 @@ namespace access {
159160
fixed(const subrange<BufferDims>& sr) : m_sr(sr) {}
160161

161162
template <int KernelDims>
162-
subrange<BufferDims> operator()(const chunk<KernelDims>&) const {
163+
subrange<BufferDims> operator()(const chunk<KernelDims>& /* chnk */) const {
163164
return m_sr;
164165
}
165166

@@ -216,6 +217,49 @@ namespace access {
216217
neighborhood(size_t, size_t)->neighborhood<2>;
217218
neighborhood(size_t, size_t, size_t)->neighborhood<3>;
218219

220+
struct kernel_dim {
221+
explicit kernel_dim(const int d) : m_dim(d) {}
222+
223+
template <int KernelDims>
224+
subrange<1> operator()(const chunk<KernelDims>& chnk) const {
225+
return {chnk.offset[m_dim], chnk.range[m_dim]};
226+
};
227+
228+
private:
229+
int m_dim;
230+
};
231+
232+
template <typename... ComponentMappers>
233+
struct components {
234+
constexpr static int buffer_dims = sizeof...(ComponentMappers);
235+
236+
explicit components(const ComponentMappers&... components) : m_mappers{components...} {}
237+
238+
template <int KernelDims>
239+
subrange<buffer_dims> operator()(const chunk<KernelDims>& chnk, const range<buffer_dims>& buffer_size) const {
240+
return apply(chnk, buffer_size, std::make_integer_sequence<int, buffer_dims>());
241+
}
242+
243+
private:
244+
std::tuple<ComponentMappers...> m_mappers;
245+
246+
template <int Dim, int KernelDims>
247+
void apply_component(const chunk<KernelDims>& chnk, const range<buffer_dims>& buffer_size, subrange<buffer_dims>& out_sr) const {
248+
static_assert(Dim < buffer_dims);
249+
const auto component_sr = detail::invoke_range_mapper_for_kernel(std::get<Dim>(m_mappers), chnk, range<1>(buffer_size[Dim]));
250+
out_sr.offset[Dim] = component_sr.offset[0];
251+
out_sr.range[Dim] = component_sr.range[0];
252+
}
253+
254+
template <int KernelDims, int... BufferDims>
255+
subrange<buffer_dims> apply(
256+
const chunk<KernelDims>& chnk, const range<buffer_dims>& buffer_size, std::integer_sequence<int, BufferDims...> /* seq */) const {
257+
subrange<buffer_dims> sr;
258+
(apply_component<BufferDims>(chnk, buffer_size, sr), ...);
259+
return sr;
260+
}
261+
};
262+
219263
} // namespace access
220264

221265
namespace experimental::access {

test/runtime_tests.cc

Lines changed: 45 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -184,6 +184,51 @@ namespace detail {
184184
}
185185
}
186186

187+
TEST_CASE("kernel_dim built-in range mapper behaves as expected", "[range-mapper]") {
188+
using celerity::access::kernel_dim;
189+
{
190+
range_mapper rm{kernel_dim(0), cl::sycl::access::mode::read, range<1>{128}};
191+
auto sr = rm.map_1(chunk<1>{{1}, {4}, {7}});
192+
CHECK(sr.offset == id<1>{1});
193+
CHECK(sr.range == range<1>{4});
194+
}
195+
{
196+
range_mapper rm{kernel_dim(1), cl::sycl::access::mode::read, range<1>{128}};
197+
auto sr = rm.map_1(chunk<3>{{1, 2, 3}, {4, 5, 6}, {7, 8, 9}});
198+
CHECK(sr.offset == id<1>{2});
199+
CHECK(sr.range == range<1>{5});
200+
}
201+
}
202+
203+
TEST_CASE("components built-in range mapper behaves as expected", "[range-mapper]") {
204+
using celerity::access::components;
205+
using celerity::access::kernel_dim;
206+
{
207+
range_mapper rm{components(all(), all(), all()), cl::sycl::access::mode::read, range<3>{128, 128, 128}};
208+
auto sr = rm.map_3(chunk<3>{{1, 2, 3}, {40, 50, 60}, {70, 80, 90}});
209+
CHECK(sr.offset == id<3>{0, 0, 0});
210+
CHECK(sr.range == range<3>{128, 128, 128});
211+
}
212+
{
213+
range_mapper rm{components(fixed(subrange<1>(19, 31)), fixed(subrange<1>(15, 44))), cl::sycl::access::mode::read, range<2>{128, 128}};
214+
auto sr = rm.map_2(chunk<3>{{1, 2, 3}, {40, 50, 60}, {70, 80, 90}});
215+
CHECK(sr.offset == id<2>{19, 15});
216+
CHECK(sr.range == range<2>{31, 44});
217+
}
218+
{
219+
range_mapper rm{components(kernel_dim(2), kernel_dim(0), kernel_dim(1)), cl::sycl::access::mode::read, range<3>{128, 128, 128}};
220+
auto sr = rm.map_3(chunk<3>{{1, 2, 3}, {40, 50, 60}, {70, 80, 90}});
221+
CHECK(sr.offset == id<3>{3, 1, 2});
222+
CHECK(sr.range == range<3>{60, 40, 50});
223+
}
224+
{
225+
range_mapper rm{components(kernel_dim(0), kernel_dim(0)), cl::sycl::access::mode::read, range<2>{128, 128}};
226+
auto sr = rm.map_2(chunk<1>{{1}, {40}, {70}});
227+
CHECK(sr.offset == id<2>{1, 1});
228+
CHECK(sr.range == range<2>{40, 40});
229+
}
230+
}
231+
187232
TEST_CASE("even_split built-in range mapper behaves as expected", "[range-mapper]") {
188233
{
189234
range_mapper rm{even_split<3>(), cl::sycl::access::mode::read, range<3>{128, 345, 678}};

website/pages/en/index.js

Lines changed: 1 addition & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -120,9 +120,7 @@ int main() {
120120
distr_queue q;
121121
q.submit([&](handler &cgh) {
122122
// (2) specify data access patterns to enable distributed execution
123-
accessor m(matrix, cgh, [size](chunk<1> chnk) {
124-
return subrange<2>({chnk.offset[0], 0}, {chnk.range[0], size});
125-
}, read_only);
123+
accessor m(matrix, cgh, access::components(access::kernel_dim(0), access::all()), read_only);
126124
accessor v(vector, cgh, access::one_to_one(), read_only);
127125
accessor r(result, cgh, access::one_to_one(), write_only, no_init);
128126

0 commit comments

Comments
 (0)