Rebase sycl-nliber2 into develop #3231

nliber · 2020-07-28T02:54:33Z

Work in progress

Includes a sycl_cp_wrapper to make Codeplays compute++ compiler behave like a single pass compiler. Kokkos::initialize will create a cl::sycl::queue

There are a couple different paths

Includes a sycl_cp_wrapper to make Codeplays compute++ compiler behave like a single pass compiler. Kokkos::initialize will create a cl::sycl::queue

(based on bjoo's code)

accelerator

Includes a sycl_cp_wrapper to make Codeplays compute++ compiler behave like a single pass compiler. Kokkos::initialize will create a cl::sycl::queue

For host, now use device.is_host() to determine which SYCL device to use, as this appears to work under both clang++ and icpx

masterleinad · 2020-09-23T16:37:50Z

Retest this please.

masterleinad · 2020-09-23T21:19:43Z

Retest this please.

masterleinad · 2020-09-23T21:50:08Z

Retest this please

masterleinad · 2020-09-23T23:14:48Z

Retest this please.

masterleinad · 2020-09-25T19:52:26Z

Retest this please.

masterleinad · 2020-09-25T20:51:04Z

OK. This is finally passing.

masterleinad · 2020-09-29T15:21:30Z

Makefile.kokkos

@@ -161,6 +165,8 @@ KOKKOS_INTERNAL_COMPILER_CLANG       := $(call kokkos_has_string,$(KOKKOS_CXX_VE
 KOKKOS_INTERNAL_COMPILER_APPLE_CLANG := $(call kokkos_has_string,$(KOKKOS_CXX_VERSION),Apple LLVM)
 KOKKOS_INTERNAL_COMPILER_HCC         := $(call kokkos_has_string,$(KOKKOS_CXX_VERSION),HCC)
 KOKKOS_INTERNAL_COMPILER_GCC         := $(call kokkos_has_string,$(KOKKOS_CXX_VERSION),GCC)
+KOKKOS_INTERNAL_COMPILER_COMPUTEPP   := $(call kokkos_has_string,$(KOKKOS_CXX_VERSION),Codeplay)


We don't have anything corresponding in the CMake configuration. Is this required?

masterleinad · 2020-09-29T15:36:30Z

bin/sycl_cp_wrapper

+    compute++ ${SYCLFLAGS} ${FLAGS} ${cpp_file}
+    ${host_compiler} ${FLAGS} -include ${sycl_file} ${cpp_file}
+  fi
+fi


Do we still need/want this file?

masterleinad · 2020-09-29T15:36:48Z

Makefile.targets

+	$(CXX) $(KOKKOS_CPPFLAGS) $(KOKKOS_CXXFLAGS) $(CXXFLAGS) -c $(KOKKOS_PATH)/core/src/SYCL/Kokkos_SYCL_Space.cpp
+Kokkos_SYCL_Instance.o: $(KOKKOS_CPP_DEPENDS) $(KOKKOS_PATH)/core/src/SYCL/Kokkos_SYCL_Instance.cpp
+	$(CXX) $(KOKKOS_CPPFLAGS) $(KOKKOS_CXXFLAGS) $(CXXFLAGS) -c $(KOKKOS_PATH)/core/src/SYCL/Kokkos_SYCL_Instance.cpp
+endif


So far I haven't checked the Makefile system.

masterleinad · 2020-09-29T15:45:50Z

core/src/Kokkos_Atomic.hpp

@@ -90,6 +90,10 @@

 #define KOKKOS_ENABLE_HIP_ATOMICS

+#elif defined(KOKKOS_ACTIVE_EXECUTION_MEMORY_SPACE_SYCL_GPU)
+
+#define KOKKOS_ENABLE_SYCL_ATOMICS


We still need to discuss what we do for the atomics but I don't think this needs to be part of this pull request.
Currently, KOKKOS_ACTIVE_EXECUTION_MEMORY_SPACE_SYCL_GPU is not defined and the build fails if we use KOKKOS_ACTIVE_EXECUTION_MEMORY_SPACE_SYCL (which is defined) instead.

masterleinad · 2020-09-29T15:47:34Z

core/src/Kokkos_Core_fwd.hpp

+class SYCLHostUSMSpace;    ///< Memory space on SYCL CPU as device
+class SYCLDeviceUSMSpace;  ///< Memory space on SYCL GPU as device
+class SYCLSharedUSMSpace;  ///< Memory space shared USM
+class SYCL;                ///< Execution space for SYCL GPU


We should discuss what the purpose of these three memory spaces are and when to use them.

masterleinad · 2020-09-29T18:58:54Z

core/src/SYCL/Kokkos_SYCL_Instance.cpp

+/* only compile this file if SYCL is enabled for Kokkos */
+#ifdef KOKKOS_ENABLE_SYCL
+
+//#include <SYCL/Kokkos_SYCL_Internal.hpp>


We don't need this include file?

masterleinad · 2020-09-29T18:59:42Z

core/src/SYCL/Kokkos_SYCL_Instance.cpp

+namespace Experimental {
+namespace Impl {
+
+int SYCLInternal::was_finalized = 0;


why is this not a bool?

masterleinad · 2020-09-29T19:17:57Z

core/src/SYCL/Kokkos_SYCL_Instance.cpp

+  return 8 * 64 * 40;  // 20480 kaveri
+#else
+  return 32 * 8 * 40;  // 81920 fiji and hawaii
+#endif


This looks wrong...

masterleinad · 2020-09-29T19:33:34Z

core/src/SYCL/Kokkos_SYCL_Instance.cpp

+
+void SYCL::impl_initialize(SYCL::SYCLDevice d) {
+  Impl::SYCLInternal::singleton().initialize(d.get_device());
+#if defined(KOKKOS_ENABLE_PROFILING)


This should go away...

masterleinad · 2020-09-29T19:44:55Z

core/src/SYCL/Kokkos_SYCL_KernelLaunch.hpp

+  isTriviallyCopyable<Driver>();
+  isTriviallyCopyable<decltype(driver.m_functor)>();
+  driver.m_policy.space().impl_internal_space_instance()->m_queue->wait();
+#ifndef SYCL_USE_BIND_LAUNCH


We don't need this (anymore), do we?

An indirect kernel is one where we have a functor that is not trivially copyable and so is explicitly constructed by the host in USM shared memory before being passed "by pointer" (inside a reference_wrapper) to SYCL parallel_for. This is to address the limitation that SYCL data types can only be implicitly copied to the device if they are trivially copyable.

Added support for non-trivially copyable kernels under SYCL

implementation to prepare for changing the USM memory type from shared to device. Current version is functionally equivalent (but not interface equivalent) to the previous version.

before allocating a new one.

tends to go together with default constructible and might as well do it while the design of the class is in my head), as that is much cleaner than wrapping it in an optional.

Fixed USMObjectMem move constructor so that its implementation is actually noexcept (like its declaration says it is). Note: this assumes the implementation of sycl::queue is just a shared_ptr underneath, which it is in all known implementations.

Use USM device memory for IndirectKernel

crtrott and others added 30 commits July 20, 2020 14:19

SYCL Initial commit adding the ability to compile a sycl code

b986a13

Includes a sycl_cp_wrapper to make Codeplays compute++ compiler behave like a single pass compiler. Kokkos::initialize will create a cl::sycl::queue

SYCL: enable simple parallel for

04f0640

Partly fix multipl kernel calls

3195afd

There are a couple different paths

Use templated wrapping launch class

314d9d3

SYCL Initial commit adding the ability to compile a sycl code

47a7cb2

Includes a sycl_cp_wrapper to make Codeplays compute++ compiler behave like a single pass compiler. Kokkos::initialize will create a cl::sycl::queue

SYCL Initial commit adding the ability to compile a sycl code

582edcf

Includes a sycl_cp_wrapper to make Codeplays compute++ compiler behave like a single pass compiler. Kokkos::initialize will create a cl::sycl::queue

WIP SyclHostUSMSpace (passes simple test)

4444c0d

Revert change for printing out make variables

8e78cff

SYCLHostUSMSpace now should pass Test02a

e3f9cd3

Added USM_memcpy for the deep_copy of SYCLHostUSMSpace views

e39897d

deep_copy between HostSpace and SYCLHostUSMSpace

d58b6e8

Additions and fixes for SYCLDeviceUSMSpace and SYCLSharedUSMSpace

8ee767c

Added SYCLDeviceUSMSpace and SYCLSharedUSMSpace

6e994a3

Cleanup

c7ffa0f

Initialize SYCL with the first GPU or accelerator

0375510

(based on bjoo's code)

Refactored so that can init with a specific device or the first GPU or

2dfdbdc

accelerator

Kokkos::SYCL::SelectDevice now defaults to the first GPU device

92657e7

Dump more device info out

d49673f

SYCL Initial commit adding the ability to compile a sycl code

d038435

Includes a sycl_cp_wrapper to make Codeplays compute++ compiler behave like a single pass compiler. Kokkos::initialize will create a cl::sycl::queue

Restored SharedAllocationTracker constructor

78e6d56

SYCL Initial commit adding the ability to compile a sycl code

5df2b64

Includes a sycl_cp_wrapper to make Codeplays compute++ compiler behave like a single pass compiler. Kokkos::initialize will create a cl::sycl::queue

SYCL Initial commit adding the ability to compile a sycl code

a4f52ac

Includes a sycl_cp_wrapper to make Codeplays compute++ compiler behave like a single pass compiler. Kokkos::initialize will create a cl::sycl::queue

In SYCL parallel_for, added wait() afterwards

c159276

Fix for deep_copy of Kokkos::View

2b4acbc

Cleanup

3ab25a1

Reworked how SYCL initialization was done

b9ce947

For host, now use device.is_host() to determine which SYCL device to use, as this appears to work under both clang++ and icpx

Removed SYCL::SelectDevice

04b584c

Removed SYCLDevice

f182d9e

Renamed SelectDevice2 to SYCLDevice

9de66ad

Changed passing sycl::device by value to by const reference

4d6f80f

masterleinad added 3 commits September 23, 2020 16:55

Only use one active SYCL memory space for now

1e7a8f4

Adapt output of ProfilingTestLibraryLoad

58465b9

Disable tests that ask for a GPU in the SYCL CI

c20a9d1

masterleinad added 2 commits September 23, 2020 17:30

Lower FEATURE_LEVEL for SYCL to 4

88aa004

Fix Serial incremental test

8fc70af

bd4 mentioned this pull request Sep 25, 2020

[SYCL] const correctness with custom containers intel/llvm#2538

Closed

masterleinad added 2 commits September 25, 2020 15:08

Remove some debug comments

3522edc

Merge remote-tracking branch 'upstream/develop' into sycl_newest

7d13423

masterleinad reviewed Sep 29, 2020

View reviewed changes

masterleinad added 4 commits September 29, 2020 17:03

Call std::abort() in some more places

31345ea

Address KOKKOS_ENABLE_PROFILING

f3a88c1

Abort on deep_copy

6355690

Revert change to Kokkos_Core.cpp

ded3e56

This was referenced Oct 1, 2020

Initial SYCL backend #3444

Closed

Initial SYCL backend rebased for backend refactoring #3447

Merged

nliber added 8 commits October 5, 2020 23:59

Merge pull request #7 from nliber/indirect-launch

145a9be

Added support for non-trivially copyable kernels under SYCL

WIP: Replaced IndirectKernel from using a vector to a hand rolled

c4b0f6b

implementation to prepare for changing the USM memory type from shared to device. Current version is functionally equivalent (but not interface equivalent) to the previous version.

For IndirectKernelMemory, changed from USM shared to USM device

0b0274a

Optimization for USMObject::reserve() to free the buffer it is holding

5708d9d

before allocating a new one.

Made USMObject default constructible (and moveable, as that

f41cb3f

tends to go together with default constructible and might as well do it while the design of the class is in my head), as that is much cleaner than wrapping it in an optional.

Renamed USMObject to USMObjectMem

0427232

Fixed USMObjectMem move constructor so that its implementation is actually noexcept (like its declaration says it is). Note: this assumes the implementation of sycl::queue is just a shared_ptr underneath, which it is in all known implementations.

Merge pull request #8 from nliber/indirect-launch

c08e5b2

Use USM device memory for IndirectKernel

crtrott closed this Dec 2, 2020

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

Rebase sycl-nliber2 into develop #3231

Rebase sycl-nliber2 into develop #3231

nliber commented Jul 28, 2020

masterleinad commented Sep 23, 2020

masterleinad commented Sep 23, 2020

masterleinad commented Sep 23, 2020

masterleinad commented Sep 23, 2020

masterleinad commented Sep 25, 2020

masterleinad commented Sep 25, 2020

masterleinad Sep 29, 2020

masterleinad Sep 29, 2020

masterleinad Sep 29, 2020

masterleinad Sep 29, 2020

masterleinad Sep 29, 2020

masterleinad Sep 29, 2020

masterleinad Sep 29, 2020

masterleinad Sep 29, 2020

masterleinad Sep 29, 2020

masterleinad Sep 29, 2020

Rebase sycl-nliber2 into develop #3231

Rebase sycl-nliber2 into develop #3231

Conversation

nliber commented Jul 28, 2020

masterleinad commented Sep 23, 2020

masterleinad commented Sep 23, 2020

masterleinad commented Sep 23, 2020

masterleinad commented Sep 23, 2020

masterleinad commented Sep 25, 2020

masterleinad commented Sep 25, 2020

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment