From 4e94d77f3a068863042aef635ab5a2f8f6303387 Mon Sep 17 00:00:00 2001 From: Robert Cohn Date: Sat, 4 Nov 2023 10:01:36 -0500 Subject: [PATCH 1/8] test local_mdspans --- test/gtest/mhp/mdstar.cpp | 19 +++++++++++++++++++ 1 file changed, 19 insertions(+) diff --git a/test/gtest/mhp/mdstar.cpp b/test/gtest/mhp/mdstar.cpp index b9dc948469..7dec0a32be 100644 --- a/test/gtest/mhp/mdstar.cpp +++ b/test/gtest/mhp/mdstar.cpp @@ -261,6 +261,25 @@ TEST_F(Mdarray, Enumerate) { static_assert(dr::distributed_range); } +TEST_F(Mdarray, Slabs) { + // leading dimension decomp of 3d array creates slabs + xhp::distributed_mdarray mdarray(extents3d); + for (auto slab : dr::mhp::local_mdspans(mdarray)) { + for (std::size_t i = 0; i < slab.extent(0); i++) { + for (std::size_t j = 0; j < slab.extent(1); j++) { + for (std::size_t k = 0; k < slab.extent(2); k++) { + slab(i, j, k) = 1; + } + } + } + } + + EXPECT_EQ(mdarray.mdspan()(0, 0, 0), 1); + EXPECT_EQ( + mdarray.mdspan()(extents3d[0] - 1, extents3d[1] - 1, extents3d[2] - 1), + 1); +} + using Submdspan = Mdspan; TEST_F(Submdspan, StaticAssert) { From 47c6dd25a07c8b2dda6f3fe93eec223fd560d701 Mon Sep 17 00:00:00 2001 From: Robert Cohn Date: Sat, 4 Nov 2023 11:15:24 -0500 Subject: [PATCH 2/8] support 1-3 dimensions for mdspan foreach with sycl --- .github/workflows/README.rst | 26 +++++++++++------- include/dr/mhp/algorithms/md_for_each.hpp | 32 ++++++++++++++++------- test/gtest/mhp/mdstar.cpp | 15 +++++++++++ 3 files changed, 54 insertions(+), 19 deletions(-) diff --git a/.github/workflows/README.rst b/.github/workflows/README.rst index 8d98d0981b..89e4b31d67 100644 --- a/.github/workflows/README.rst +++ b/.github/workflows/README.rst @@ -9,24 +9,30 @@ DevCloud Runners ================ -We have 2 self-hosted runners on devcloud. We run in tmux so it will -continue to run after disconnecting. ``ssh`` to DevCloud:: +We have 2 self-hosted runners on devcloud. We run in an emacs daemon +so it will continue to run after disconnecting. ``ssh`` to DevCloud +and start an emacs daemon:: + + emacs --daemon + +Inside an emacs shell:: - tmux cd github/runner-1 ./run.sh -For the second runner, split the window with ``control-b "`` and start -the runner:: +Rename the shell to ``devcloud-1`` and start another emacs shell:: cd github/runner-2 ./run.sh -To check on a runner ``ssh`` to DevCloud and do:: +If you disconnect emacs and the runners will continue to run. To check +on a runner ``ssh`` to DevCloud and do:: - tmux ls + emacsclient -nw -To see the sessions. Usually there is just one name ``0``. Attach to -the session:: +And visit ``devcloud-1`` and ``devcloud-2`` buffers. - tmux attach -t 0 +``tmux`` is an alternative. I do not use it because devcloud allows 4 +logins via ``ssh``. Using ``tmux`` with 2 runners consumes 3 even +while it is detached. You may create a situation where you cannot log +in because ``tmux`` is consuming all the logins. diff --git a/include/dr/mhp/algorithms/md_for_each.hpp b/include/dr/mhp/algorithms/md_for_each.hpp index 6928ab62f2..ccd98d52c2 100644 --- a/include/dr/mhp/algorithms/md_for_each.hpp +++ b/include/dr/mhp/algorithms/md_for_each.hpp @@ -142,8 +142,15 @@ void for_each(F op, is_mdspan_view auto &&...drs) { auto invoke_index = [=](auto index) { // Transform mdspans into references auto references = detail::tie_transform( - operand_mdspans, [index](auto mdspan) -> decltype(auto) { - return mdspan(index[0], index[1]); + operand_mdspans, [mdspan0, index](auto mdspan) -> decltype(auto) { + static_assert(1 <= mdspan0.rank() && mdspan0.rank() <= 3); + if constexpr (mdspan0.rank() == 1) { + return mdspan(index[0]); + } else if constexpr (mdspan0.rank() == 2) { + return mdspan(index[0], index[1]); + } else if constexpr (mdspan0.rank() == 3) { + return mdspan(index[0], index[1], index[2]); + } }); static_assert( std::invocable || @@ -160,13 +167,20 @@ void for_each(F op, is_mdspan_view auto &&...drs) { } }; - // TODO: Extend sycl_utils.hpp to handle ranges > 1D. It uses - // ndrange and handles > 32 bits. - - dr::__detail::parallel_for( - mhp::sycl_queue(), - sycl::range(mdspan0.extent(0), mdspan0.extent(1)), invoke_index) - .wait(); + if constexpr (mdspan0.rank() == 1) { + auto range = sycl::range(mdspan0.extent(0)); + dr::__detail::parallel_for(mhp::sycl_queue(), range, invoke_index) + .wait(); + } else if constexpr (mdspan0.rank() == 2) { + auto range = sycl::range(mdspan0.extent(0), mdspan0.extent(1)); + dr::__detail::parallel_for(mhp::sycl_queue(), range, invoke_index) + .wait(); + } else if constexpr (mdspan0.rank() == 3) { + auto range = sycl::range(mdspan0.extent(0), mdspan0.extent(1), + mdspan0.extent(2)); + dr::__detail::parallel_for(mhp::sycl_queue(), range, invoke_index) + .wait(); + } #else assert(false); #endif diff --git a/test/gtest/mhp/mdstar.cpp b/test/gtest/mhp/mdstar.cpp index 7dec0a32be..29e2d0889d 100644 --- a/test/gtest/mhp/mdstar.cpp +++ b/test/gtest/mhp/mdstar.cpp @@ -280,6 +280,21 @@ TEST_F(Mdarray, Slabs) { 1); } +TEST_F(Mdarray, MdForEach3d) { + // leading dimension decomp of 3d array creates slabs + xhp::distributed_mdarray mdarray(extents3d); + std::vector local(mdarray.size(), 0); + // rng::iota(local, 0); + + auto set = [d1 = extents3d[1], d2 = extents3d[2]](auto index, auto v) { + auto &[o] = v; + o = index[0] * d1 * d2 + index[1] * d2 + index[2]; + }; + dr::mhp::for_each(set, mdarray); + + EXPECT_EQ(mdarray.view(), local) << mdrange_message(mdarray); +} + using Submdspan = Mdspan; TEST_F(Submdspan, StaticAssert) { From 9311c875ff49213c816b50e9ec01992824fd10d8 Mon Sep 17 00:00:00 2001 From: Robert Cohn Date: Sat, 4 Nov 2023 12:13:37 -0500 Subject: [PATCH 3/8] update --- scripts/devcloud-test.sh | 7 ++++--- 1 file changed, 4 insertions(+), 3 deletions(-) diff --git a/scripts/devcloud-test.sh b/scripts/devcloud-test.sh index f898dfeabd..87898f39ef 100755 --- a/scripts/devcloud-test.sh +++ b/scripts/devcloud-test.sh @@ -26,9 +26,10 @@ echo "::group::SHP GPU Test" ONEAPI_DEVICE_SELECTOR=level_zero:0 time ctest --test-dir build -L SHP echo "::endgroup::" -echo "::group::SHP CPU Test" -ONEAPI_DEVICE_SELECTOR=opencl:cpu time ctest --test-dir build -L SHP -echo "::endgroup::" +# disabled: very slow or fails when cryptominer is on devcloud +#echo "::group::SHP CPU Test" +#ONEAPI_DEVICE_SELECTOR=opencl:cpu time ctest --test-dir build -L SHP +#echo "::endgroup::" echo "::group::MHP GPU Test" ONEAPI_DEVICE_SELECTOR=level_zero:* time ctest --test-dir build -L MHP From 22efc346b41cb8c9a7db36b4d19cf696f67bb293 Mon Sep 17 00:00:00 2001 From: Robert Cohn Date: Sat, 4 Nov 2023 12:17:10 -0500 Subject: [PATCH 4/8] update --- test/gtest/mhp/mdstar.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/test/gtest/mhp/mdstar.cpp b/test/gtest/mhp/mdstar.cpp index 29e2d0889d..d036efdca9 100644 --- a/test/gtest/mhp/mdstar.cpp +++ b/test/gtest/mhp/mdstar.cpp @@ -284,7 +284,7 @@ TEST_F(Mdarray, MdForEach3d) { // leading dimension decomp of 3d array creates slabs xhp::distributed_mdarray mdarray(extents3d); std::vector local(mdarray.size(), 0); - // rng::iota(local, 0); + rng::iota(local, 0); auto set = [d1 = extents3d[1], d2 = extents3d[2]](auto index, auto v) { auto &[o] = v; From 2365d6fc54f0d460a9bf691cacf03686bf73dad1 Mon Sep 17 00:00:00 2001 From: Robert Cohn Date: Sat, 4 Nov 2023 13:29:33 -0500 Subject: [PATCH 5/8] update --- test/gtest/mhp/mdstar.cpp | 5 +++-- 1 file changed, 3 insertions(+), 2 deletions(-) diff --git a/test/gtest/mhp/mdstar.cpp b/test/gtest/mhp/mdstar.cpp index d036efdca9..b7278a656c 100644 --- a/test/gtest/mhp/mdstar.cpp +++ b/test/gtest/mhp/mdstar.cpp @@ -283,7 +283,7 @@ TEST_F(Mdarray, Slabs) { TEST_F(Mdarray, MdForEach3d) { // leading dimension decomp of 3d array creates slabs xhp::distributed_mdarray mdarray(extents3d); - std::vector local(mdarray.size(), 0); + std::vector local(extents3d[0] * extents3d[1] * extents3d[2], 0); rng::iota(local, 0); auto set = [d1 = extents3d[1], d2 = extents3d[2]](auto index, auto v) { @@ -292,7 +292,8 @@ TEST_F(Mdarray, MdForEach3d) { }; dr::mhp::for_each(set, mdarray); - EXPECT_EQ(mdarray.view(), local) << mdrange_message(mdarray); + EXPECT_EQ(xhp::views::take(mdarray.view(), local.size()), local) + << mdrange_message(mdarray); } using Submdspan = Mdspan; From 16f203d2fb5241f516368be02122cba8c87ace24 Mon Sep 17 00:00:00 2001 From: Robert Cohn Date: Sun, 5 Nov 2023 02:26:46 -0800 Subject: [PATCH 6/8] update --- test/gtest/mhp/mdstar.cpp | 5 +++++ 1 file changed, 5 insertions(+) diff --git a/test/gtest/mhp/mdstar.cpp b/test/gtest/mhp/mdstar.cpp index b7278a656c..1d47b0e824 100644 --- a/test/gtest/mhp/mdstar.cpp +++ b/test/gtest/mhp/mdstar.cpp @@ -262,6 +262,11 @@ TEST_F(Mdarray, Enumerate) { } TEST_F(Mdarray, Slabs) { + // local_mdspan is not accessible for device memory + if (options.count("device-memory")) { + return; + } + // leading dimension decomp of 3d array creates slabs xhp::distributed_mdarray mdarray(extents3d); for (auto slab : dr::mhp::local_mdspans(mdarray)) { From c6909a4f85a5c46a2d71e5786a4df415159e5835 Mon Sep 17 00:00:00 2001 From: Robert Cohn Date: Sun, 5 Nov 2023 02:29:16 -0800 Subject: [PATCH 7/8] update --- .github/workflows/pr.yml | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/.github/workflows/pr.yml b/.github/workflows/pr.yml index c3474b9090..02825f8ff0 100644 --- a/.github/workflows/pr.yml +++ b/.github/workflows/pr.yml @@ -20,7 +20,7 @@ env: jobs: checks: - runs-on: ${{ github.repository_owner == 'oneapi-src' && 'intel-ubuntu-latest' || 'ubuntu-latest' }} + runs-on: ['ubuntu-latest','intel-ubuntu-latest'] steps: - uses: actions/checkout@v4 - name: Ubuntu dependencies From feeafe1e32e7421d3cce497b173ef44f844ab49d Mon Sep 17 00:00:00 2001 From: Robert Cohn Date: Sun, 5 Nov 2023 02:30:35 -0800 Subject: [PATCH 8/8] update --- .github/workflows/pr.yml | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/.github/workflows/pr.yml b/.github/workflows/pr.yml index 02825f8ff0..c3474b9090 100644 --- a/.github/workflows/pr.yml +++ b/.github/workflows/pr.yml @@ -20,7 +20,7 @@ env: jobs: checks: - runs-on: ['ubuntu-latest','intel-ubuntu-latest'] + runs-on: ${{ github.repository_owner == 'oneapi-src' && 'intel-ubuntu-latest' || 'ubuntu-latest' }} steps: - uses: actions/checkout@v4 - name: Ubuntu dependencies