Implement kernel_dim and components range mappers

Author: fknorr

include/range_mapper.h

@@ -1,6 +1,7 @@
 #pragma once
 
 #include <stdexcept>
+#include <tuple>
 #include <type_traits>
 
 #include <CL/sycl.hpp>
@@ -172,7 +173,7 @@ namespace access {
 		fixed(const subrange<BufferDims>& sr) : m_sr(sr) {}
 
 		template <int KernelDims>
-		subrange<BufferDims> operator()(const chunk<KernelDims>&) const {
+		subrange<BufferDims> operator()(const chunk<KernelDims>& /* chnk */) const {
 			return m_sr;
 		}
 
@@ -248,6 +249,49 @@ namespace access {
 	neighborhood(size_t, size_t)->neighborhood<2>;
 	neighborhood(size_t, size_t, size_t)->neighborhood<3>;
 
+	struct kernel_dim {
+		explicit kernel_dim(const int d) : m_dim(d) {}
+
+		template <int KernelDims>
+		subrange<1> operator()(const chunk<KernelDims>& chnk) const {
+			return {chnk.offset[m_dim], chnk.range[m_dim]};
+		};
+
+	  private:
+		int m_dim;
+	};
+
+	template <typename... ComponentMappers>
+	struct components {
+		constexpr static int buffer_dims = sizeof...(ComponentMappers);
+
+		explicit components(const ComponentMappers&... components) : m_mappers{components...} {}
+
+		template <int KernelDims>
+		subrange<buffer_dims> operator()(const chunk<KernelDims>& chnk, const range<buffer_dims>& buffer_size) const {
+			return apply(chnk, buffer_size, std::make_integer_sequence<int, buffer_dims>());
+		}
+
+	  private:
+		std::tuple<ComponentMappers...> m_mappers;
+
+		template <int Dim, int KernelDims>
+		void apply_component(const chunk<KernelDims>& chnk, const range<buffer_dims>& buffer_size, subrange<buffer_dims>& out_sr) const {
+			static_assert(Dim < buffer_dims);
+			const auto component_sr = detail::invoke_range_mapper_for_kernel(std::get<Dim>(m_mappers), chnk, range<1>(buffer_size[Dim]));
+			out_sr.offset[Dim] = component_sr.offset[0];
+			out_sr.range[Dim] = component_sr.range[0];
+		}
+
+		template <int KernelDims, int... BufferDims>
+		subrange<buffer_dims> apply(
+		    const chunk<KernelDims>& chnk, const range<buffer_dims>& buffer_size, std::integer_sequence<int, BufferDims...> /* seq */) const {
+			subrange<buffer_dims> sr;
+			(apply_component<BufferDims>(chnk, buffer_size, sr), ...);
+			return sr;
+		}
+	};
+
 } // namespace access
 
 namespace experimental::access {

test/runtime_tests.cc

@@ -185,6 +185,51 @@ namespace detail {
 		}
 	}
 
+	TEST_CASE("dim built-in range mapper behaves as expected", "[range-mapper]") {
+		using celerity::access::kernel_dim;
+		{
+			range_mapper rm{kernel_dim(0), cl::sycl::access::mode::read, range<1>{128}};
+			auto sr = rm.map_1(chunk<1>{{1}, {4}, {7}});
+			CHECK(sr.offset == id<1>{1});
+			CHECK(sr.range == range<1>{4});
+		}
+		{
+			range_mapper rm{kernel_dim(1), cl::sycl::access::mode::read, range<1>{128}};
+			auto sr = rm.map_1(chunk<3>{{1, 2, 3}, {4, 5, 6}, {7, 8, 9}});
+			CHECK(sr.offset == id<1>{2});
+			CHECK(sr.range == range<1>{5});
+		}
+	}
+
+	TEST_CASE("shuffle built-in range mapper behaves as expected", "[range-mapper]") {
+		using celerity::access::components;
+		using celerity::access::kernel_dim;
+		{
+			range_mapper rm{components(all(), all(), all()), cl::sycl::access::mode::read, range<3>{128, 128, 128}};
+			auto sr = rm.map_3(chunk<3>{{1, 2, 3}, {40, 50, 60}, {70, 80, 90}});
+			CHECK(sr.offset == id<3>{0, 0, 0});
+			CHECK(sr.range == range<3>{128, 128, 128});
+		}
+		{
+			range_mapper rm{components(fixed(subrange<1>(19, 31)), fixed(subrange<1>(15, 44))), cl::sycl::access::mode::read, range<2>{128, 128}};
+			auto sr = rm.map_2(chunk<3>{{1, 2, 3}, {40, 50, 60}, {70, 80, 90}});
+			CHECK(sr.offset == id<2>{19, 15});
+			CHECK(sr.range == range<2>{31, 44});
+		}
+		{
+			range_mapper rm{components(kernel_dim(2), kernel_dim(0), kernel_dim(1)), cl::sycl::access::mode::read, range<3>{128, 128, 128}};
+			auto sr = rm.map_3(chunk<3>{{1, 2, 3}, {40, 50, 60}, {70, 80, 90}});
+			CHECK(sr.offset == id<3>{3, 1, 2});
+			CHECK(sr.range == range<3>{60, 40, 50});
+		}
+		{
+			range_mapper rm{components(kernel_dim(0), kernel_dim(0)), cl::sycl::access::mode::read, range<2>{128, 128}};
+			auto sr = rm.map_2(chunk<1>{{1}, {40}, {70}});
+			CHECK(sr.offset == id<2>{1, 1});
+			CHECK(sr.range == range<2>{40, 40});
+		}
+	}
+
 	TEST_CASE("even_split built-in range mapper behaves as expected", "[range-mapper]") {
 		{
 			range_mapper rm{even_split<3>(), cl::sycl::access::mode::read, range<3>{128, 345, 678}};

website/pages/en/index.js

@@ -120,9 +120,7 @@ int main() {
     distr_queue q;
     q.submit([&](handler &cgh) {
         // (2) specify data access patterns to enable distributed execution
-        accessor m(matrix, cgh, [size](chunk<1> chnk) {
-            return subrange<2>({chnk.offset[0], 0}, {chnk.range[0], size});
-        }, read_only);
+        accessor m(matrix, cgh, access::components(access::kernel_dim(0), access::all()), read_only);
         accessor v(vector, cgh, access::one_to_one(), read_only);
         accessor r(result, cgh, access::one_to_one(), write_only, no_init);