Update SYCL compiler and remove printf guarding with SYCL #614

aprokop · 2022-01-06T00:08:08Z

Fix #613.

masterleinad · 2022-01-06T01:28:23Z

You need to build Kokkos without deprecated code with this compiler version. It's a compiler bug.

aprokop · 2022-01-06T04:07:35Z

You need to build Kokkos without deprecated code with this compiler version. It's a compiler bug.

What is the option to disable deprecated code? -D Kokkos_ENABLE_DEPRECATED_CODE_3=OFF?

masterleinad · 2022-01-06T14:35:26Z

What is the option to disable deprecated code? -D Kokkos_ENABLE_DEPRECATED_CODE_3=OFF?

Yes, the configuration we were using in the Kokkos CI is
https://github.com/kokkos/kokkos/blob/fea35745dfad72decfc4516c4d026d4a0899af4e/.jenkins#L69-L83

masterleinad · 2022-01-06T14:37:12Z

Note that we will still get deprecation warnings unless we update to the latest Kokkos release.

aprokop · 2022-01-06T19:29:27Z

Kokkos 3.4 fails with

/scratch/kokkos/core/src/impl/Kokkos_Timer.hpp:58:15: error: no member named 'Timer' in namespace 'Kokkos'
using Kokkos::Timer;
      ~~~~~~~~^

masterleinad · 2022-01-06T19:33:15Z

Kokkos 3.4 fails with

That should not happen with deprecated code disabled...

aprokop · 2022-01-06T19:34:55Z

That should not happen with deprecated code disabled...

The deprecated code protection in kokkos/core/src/impl/Kokkos_Timer.hpp only appeared in Kokkos 3.5. In 3.4, it is an unprotected using Kokkos::Timer;.

aprokop · 2022-01-06T19:59:41Z

With Kokkos 3.5, the container builds. However, getting "unnamed lambda" errors from ArborX SYCL build:

In file included from /var/jenkins/workspace/ArborX_PR-614/build/test/tstQueryTreeDegenerate_BF.cpp:1:
In file included from /var/jenkins/workspace/ArborX_PR-614/src/ArborX_BruteForce.hpp:15:
In file included from /var/jenkins/workspace/ArborX_PR-614/src/details/ArborX_AccessTraits.hpp:15:
In file included from /var/jenkins/workspace/ArborX_PR-614/src/details/ArborX_Box.hpp:15:
In file included from /var/jenkins/workspace/ArborX_PR-614/src/details/ArborX_DetailsKokkosExtArithmeticTraits.hpp:15:
In file included from /opt/kokkos/include/Kokkos_Macros.hpp:110:
In file included from /opt/kokkos/include/KokkosCore_Config_SetupBackend.hpp:47:
In file included from /opt/kokkos/include/setup/Kokkos_Setup_SYCL.hpp:48:
In file included from /opt/sycl/bin/../include/sycl/CL/sycl.hpp:15:
In file included from /opt/sycl/bin/../include/sycl/CL/sycl/backend.hpp:26:
In file included from /opt/sycl/bin/../include/sycl/CL/sycl/queue.hpp:19:
/opt/sycl/bin/../include/sycl/CL/sycl/handler.hpp:1029:52: error: unnamed lambda 'sycl::detail::__pf_kernel_wrapper<oneapi::dpl::__par_backend_hetero::__parallel_sort_copy_back_kernel<oneapi::dpl::execution::DefaultKernelName, oneapi::dpl::__ranges::zip_view<oneapi::dpl::__ranges::guard_view<unsigned int *>, oneapi::dpl::__ranges::guard_view<unsigned int *>>, oneapi::dpl::__par_backend_hetero::__full_merge_kernel, (lambda at /var/jenkins/workspace/ArborX_PR-614/src/details/ArborX_DetailsSortUtils.hpp:180:7)>>' used
      kernel_parallel_for<KernelName, ElementType>(KernelFunc);

This is coming from this code:

ArborX/src/details/ArborX_DetailsSortUtils.hpp

Lines 178 to 180 in a1b1a00

    
           oneapi::dpl::sort( 
        
               policy, zipped_begin, zipped_begin + n, 
        
               [](auto lhs, auto rhs) { return std::get<0>(lhs) < std::get<0>(rhs); });

There is an fsycl-unnamed-lambda option, but it's ON by default, so unnamed lambdas should be allowed. Will try to enable them explicitly.

aprokop · 2022-01-06T20:20:58Z

Explicitly enabling -fsycl-unnamed-lambda did not help. Likely, because the message is confusing.

Trying a patch where lambda is not anonymous.

dalg24 · 2022-01-06T21:05:34Z

src/details/ArborX_DetailsSortUtils.hpp

-  oneapi::dpl::sort(
-      policy, zipped_begin, zipped_begin + n,
-      [](auto lhs, auto rhs) { return std::get<0>(lhs) < std::get<0>(rhs); });
+  oneapi::dpl::sort(policy, zipped_begin, zipped_begin + n, SYCLCompare{});


Sounds more like a bug with the toolchain. We do have -fsycl-unnamed-lambda in the compile line

Possibly. I'm not sure sure what to think of intel/llvm#4769. Lets see if it fixes it. If not, I'll try to put the latest nightly and check.

aprokop · 2022-01-06T23:41:26Z

Cannot build Kokkos 3.5 with the latest SYCL nightly (20220106):

/scratch/kokkos/core/src/desul/atomics/SYCLConversions.hpp:48:20: error: no member named 'ONEAPI' in namespace 'sycl'
  static constexpr DESUL_SYCL_NAMESPACE::memory_scope value = DESUL_SYCL_NAMESPACE::memory_scope::device;
                   ^~~~~~~~~~~~~~~~~~~~
/scratch/kokkos/core/src/desul/atomics/Macros.hpp:38:36: note: expanded from macro 'DESUL_SYCL_NAMESPACE'
#define DESUL_SYCL_NAMESPACE sycl::ONEAPI
                             ~~~~~~^

aprokop · 2022-01-07T18:00:10Z

This seems too brittle to continue trying to figure out. We could probably work around with a custom comparator, but that's simply another workaround. I think we should not spend any more time on this, and retry it in 6 months or a year. Maybe SYCL compiler would be more stable by then.

masterleinad · 2022-01-07T18:02:25Z

To me, it looks like we are almost there. I just don't have any idea what all this linking with omp is about. In the end, it makes sense to me to update after the next Kokkos release.

masterleinad · 2022-01-10T17:09:49Z

Let's see if setting the host implementation for oneDPL to "serial" helps, see https://github.com/oneapi-src/oneDPL/blob/main/cmake/README.md#using-onedpl-package.

aprokop · 2022-01-10T17:31:25Z

Let's see if setting the host implementation for oneDPL to "serial" helps, see oneapi-src/oneDPL@main/cmake/README.md#using-onedpl-package.

But even if it helps, I assume we need a parallel backend (which backend does it actually use on the Intel GPU, "openmp"?).

masterleinad · 2022-01-10T18:41:47Z

But even if it helps, I assume we need a parallel backend (which backend does it actually use on the Intel GPU, "openmp"?).

We are explicitly requesting dpcpp_only at compile time. My understanding is that ONEDPL_PAR_BACKEND really only affects host execution which we don't use anyway, also see https://github.com/oneapi-src/oneDPL/blob/394ae95fd935ac3532b129d920135433f42bc9eb/include/oneapi/dpl/pstl/parallel_backend.h.

It seems that we are now down to some performance regressions.

aprokop · 2022-01-10T19:28:45Z

It seems that we are now down to some performance regressions.

What do you mean?

masterleinad · 2022-01-10T19:37:27Z

What do you mean?

This pull request now compiles and runs but ArborX_QueryTree_Test and ArborX_BruteForce_Example are timing out.

masterleinad · 2022-01-12T15:31:08Z

We were running into a deadlock with the fences in the SYCL ThreadVectorRange implementation for which there is no good solution, yet (with the current compiler version). The workaround is don't required on Intel GPUs interestingly.

masterleinad · 2022-01-14T18:03:23Z

For reference, we need intel/llvm#4904 or intel/llvm#4903 (including an actual implementation) to fix the current behavior in Kokkos.

aprokop

Other than two minor comments, I think this is fine.

aprokop · 2022-01-24T17:11:49Z

.jenkins/continuous.groovy

-                                    -D CMAKE_CXX_FLAGS="-Wpedantic -Wall -Wextra -Wno-unknown-cuda-version -fsycl -fsycl-targets=nvptx64-nvidia-cuda-sycldevice" \
+                                    -D CMAKE_CXX_FLAGS="-Wpedantic -Wall -Wextra -Wno-unknown-cuda-version" \


Do we really need to remove these?

No, we don't need to remove these but there is no reason to keep it IMHO. These flags are(or should be) exported by Kokkos.

.jenkins/continuous.groovy

aprokop · 2022-01-24T17:23:46Z

@masterleinad OK, can you please squash the commits and restore other testing? Or, let me know and I'll do this myself.

…ely on Kokkos compiler flags

masterleinad · 2022-01-24T17:54:18Z

@masterleinad OK, can you please squash the commits and restore other testing? Or, let me know and I'll do this myself.

Done.

aprokop · 2022-01-24T19:29:15Z

retest this please

docker/Dockerfile.sycl

aprokop · 2022-01-24T20:11:26Z

Huh, just noticed that SYCL build time dropped from 30 minutes to 11.

dalg24

I'd prefer if you add the comment about the Kokkos git commit hash used when building the Docker image

aprokop added the testing Anything to do with tests and CI label Jan 6, 2022

aprokop changed the title ~~Update SYCL compiler and remove printf workarounds~~ Update SYCL compiler and remove printf guarding with SYCL Jan 6, 2022

aprokop force-pushed the sycl_update branch from 7571466 to a76dd3e Compare January 6, 2022 20:18

aprokop force-pushed the sycl_update branch 2 times, most recently from e0200b2 to a07bf60 Compare January 6, 2022 21:04

dalg24 reviewed Jan 6, 2022

View reviewed changes

aprokop force-pushed the sycl_update branch from a07bf60 to 9a1d749 Compare January 6, 2022 23:10

aprokop marked this pull request as draft January 7, 2022 00:26

masterleinad force-pushed the sycl_update branch from d8211be to a0de29c Compare January 10, 2022 17:08

masterleinad mentioned this pull request Jan 12, 2022

Use TeamVectorRange for filling tiles in BruteForce implementation #616

Merged

masterleinad mentioned this pull request Jan 14, 2022

Add a FIXME for SYCL and ThreadVectorRange parallel_for kokkos/kokkos#4676

Merged

aprokop commented Jan 24, 2022

View reviewed changes

masterleinad added 3 commits January 24, 2022 12:50

Update Jenkins SYCL CI to use serial host oneDPL implementation and r…

20e96b3

…ely on Kokkos compiler flags

Update docker container for SYCL

00dbd21

Remove __SYCL_DEVICE_ONLY__ guards for ptintf

d54e117

masterleinad force-pushed the sycl_update branch from c7c3b30 to d54e117 Compare January 24, 2022 17:54

aprokop marked this pull request as ready for review January 24, 2022 19:06

dalg24 reviewed Jan 24, 2022

View reviewed changes

docker/Dockerfile.sycl Show resolved Hide resolved

dalg24 approved these changes Jan 24, 2022

View reviewed changes

Comment on Kokkos commit in SYCL Dockerfile

11db57d

dalg24 merged commit 5513b63 into arborx:master Jan 24, 2022

aprokop deleted the sycl_update branch January 24, 2022 23:03

masterleinad mentioned this pull request Feb 7, 2022

Revert "Remove __SYCL_DEVICE_ONLY__ guards for ptintf" #625

Merged

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

Update SYCL compiler and remove printf guarding with SYCL #614

Update SYCL compiler and remove printf guarding with SYCL #614

aprokop commented Jan 6, 2022

masterleinad commented Jan 6, 2022

aprokop commented Jan 6, 2022

masterleinad commented Jan 6, 2022

masterleinad commented Jan 6, 2022

aprokop commented Jan 6, 2022

masterleinad commented Jan 6, 2022

aprokop commented Jan 6, 2022

aprokop commented Jan 6, 2022 •

edited

Loading

aprokop commented Jan 6, 2022

dalg24 Jan 6, 2022

aprokop Jan 6, 2022

aprokop commented Jan 6, 2022 •

edited

Loading

aprokop commented Jan 7, 2022

masterleinad commented Jan 7, 2022

masterleinad commented Jan 10, 2022

aprokop commented Jan 10, 2022

masterleinad commented Jan 10, 2022

aprokop commented Jan 10, 2022

masterleinad commented Jan 10, 2022

masterleinad commented Jan 12, 2022

masterleinad commented Jan 14, 2022

aprokop left a comment

aprokop Jan 24, 2022

masterleinad Jan 24, 2022

aprokop commented Jan 24, 2022

masterleinad commented Jan 24, 2022

aprokop commented Jan 24, 2022

aprokop commented Jan 24, 2022

dalg24 left a comment

		-D CMAKE_CXX_FLAGS="-Wpedantic -Wall -Wextra -Wno-unknown-cuda-version -fsycl -fsycl-targets=nvptx64-nvidia-cuda-sycldevice" \
		-D CMAKE_CXX_FLAGS="-Wpedantic -Wall -Wextra -Wno-unknown-cuda-version" \

Update SYCL compiler and remove printf guarding with SYCL #614

Update SYCL compiler and remove printf guarding with SYCL #614

Conversation

aprokop commented Jan 6, 2022

masterleinad commented Jan 6, 2022

aprokop commented Jan 6, 2022

masterleinad commented Jan 6, 2022

masterleinad commented Jan 6, 2022

aprokop commented Jan 6, 2022

masterleinad commented Jan 6, 2022

aprokop commented Jan 6, 2022

aprokop commented Jan 6, 2022 • edited Loading

aprokop commented Jan 6, 2022

dalg24 Jan 6, 2022

Choose a reason for hiding this comment

aprokop Jan 6, 2022

Choose a reason for hiding this comment

aprokop commented Jan 6, 2022 • edited Loading

aprokop commented Jan 7, 2022

masterleinad commented Jan 7, 2022

masterleinad commented Jan 10, 2022

aprokop commented Jan 10, 2022

masterleinad commented Jan 10, 2022

aprokop commented Jan 10, 2022

masterleinad commented Jan 10, 2022

masterleinad commented Jan 12, 2022

masterleinad commented Jan 14, 2022

aprokop left a comment

Choose a reason for hiding this comment

aprokop Jan 24, 2022

Choose a reason for hiding this comment

masterleinad Jan 24, 2022

Choose a reason for hiding this comment

aprokop commented Jan 24, 2022

masterleinad commented Jan 24, 2022

aprokop commented Jan 24, 2022

aprokop commented Jan 24, 2022

dalg24 left a comment

Choose a reason for hiding this comment

aprokop commented Jan 6, 2022 •

edited

Loading

aprokop commented Jan 6, 2022 •

edited

Loading