Implement `KOKKOS_IF_ON_{HOST,DEVICE}` macros #4660

dalg24 · 2022-01-11T16:49:34Z

Proposing a solution to #3404

masterleinad

I think I would prefer KOKKOS_IF_ON_HOST/DEVICE slightly but otherwise, this looks pretty good to me already.

… in the dimension runtime check in the View constructor

…ter on the host side

…ationRecord implementation

…fetView constructors

…n and improve error message

…accessible

Co-Authored-By: Christian Trott <crtrott@sandia.gov>

…SPACE_HOST

…ad nested #ifdef

…ORY_SPACE_CUDA

…> and drop macro guards

crtrott

Looks good: just need to discuss naming in meeting tomorrow.

crtrott · 2022-01-19T19:47:51Z

Name: IF_ON_HOST/DEVICE or IF_HOST/DEVICE
(F means IF_ON_HOST/DEVICE)
F N A
7 5 3

janciesko · 2022-01-19T20:00:34Z

Now that I think of it - since KOKKOS_IF_{HOST,DEVICE} refers to code (sources and files) itself which reside wherever, leaving it as proposed by Damien makes it sound like KOKKOS_IF_{HOST_CODE,DEVICE_CODE} w/o stating the obvious “CODE”, which I like more now.

dalg24 · 2022-01-20T03:25:17Z

Retest this please

masterleinad

Can you please comment on deprecating ActiveExecutionMemorySpace?
I left some comments but other than that, the pull request looks good to me.

Rebasing some pull requests will be fun. 🙂

core/src/Kokkos_Core_fwd.hpp

masterleinad · 2022-01-20T19:22:24Z

core/src/Kokkos_Core_fwd.hpp

+  KOKKOS_DEPRECATED KOKKOS_FUNCTION static void check() {}
 };

 template <typename DstMemorySpace, typename SrcMemorySpace>
 struct verify_space<DstMemorySpace, SrcMemorySpace, false> {
-  KOKKOS_FUNCTION static void check() {
+  KOKKOS_DEPRECATED KOKKOS_FUNCTION static void check() {
    Kokkos::abort(
        "Kokkos::View ERROR: attempt to access inaccessible memory space");
  };
 };
+#endif


Highlighting deprecations.

masterleinad · 2022-01-20T19:22:35Z

core/src/Kokkos_Core_fwd.hpp

+#ifdef KOKKOS_ENABLE_DEPRECATED_CODE_3
 #define KOKKOS_RESTRICT_EXECUTION_TO_DATA(DATA_SPACE, DATA_PTR)        \
  Kokkos::Impl::verify_space<Kokkos::Impl::ActiveExecutionMemorySpace, \
                             DATA_SPACE>::check();

 #define KOKKOS_RESTRICT_EXECUTION_TO_(DATA_SPACE)                      \
  Kokkos::Impl::verify_space<Kokkos::Impl::ActiveExecutionMemorySpace, \
                             DATA_SPACE>::check();
+#endif


Highlighting deprecations.

masterleinad · 2022-01-20T22:03:26Z

core/src/Kokkos_MemoryPool.hpp

@@ -538,7 +538,7 @@ class MemoryPool {
 #else
    const uint32_t block_id_hint =
        (uint32_t)(Kokkos::Impl::clock_tic()
-#if defined(KOKKOS_ACTIVE_EXECUTION_MEMORY_SPACE_CUDA)
+#ifdef __CUDA_ARCH__  // FIXME_CUDA


Highlighting

core/src/Kokkos_MemoryPool.hpp

masterleinad · 2022-01-20T22:06:11Z

core/src/Kokkos_NumericTraits.hpp

+// No __host__ __device__ annotation because long double treated as double in
+// device code.  May be revisited later if that is not true any more.


Suggested change

// No __host__ __device__ annotation because long double treated as double in

// device code. May be revisited later if that is not true any more.

// FIXME No __host__ __device__ annotation because long double treated as double in

// device code. May be revisited later if that is not true any more.

This is not a FIXME

I think comments such as "Maybe revisited later" are entirely ignored if they don't have a FIXME with them.

FIXME_MAYBE if one gpu backend supports this one day and we really need it

masterleinad · 2022-01-20T22:30:05Z

core/src/impl/Kokkos_HostBarrier.hpp

+    KOKKOS_IF_ON_HOST((impl_wait_until_equal_host(ptr, v, active_wait);))
+
+    KOKKOS_IF_ON_DEVICE(((void)active_wait; while (!test_equal(ptr, v)){}))
+  }
+
+  static void impl_wait_until_equal_host(int* ptr, const int v,
+                                         bool active_wait = true) noexcept {


Why are specifically here defining a different function but noting other places?

Because of all the #ifdefs

masterleinad · 2022-01-20T22:36:53Z

core/src/impl/Kokkos_TaskBase.hpp

@@ -314,7 +314,7 @@ class Task : public TaskBase, public FunctorType {
    // If team then only one thread calls destructor.

    const bool only_one_thread =
-#if defined(KOKKOS_ACTIVE_EXECUTION_MEMORY_SPACE_CUDA)
+#ifdef __CUDA_ARCH__  // FIXME_CUDA


Highlighting

masterleinad · 2022-01-20T22:37:01Z

core/src/impl/Kokkos_TaskNode.hpp

@@ -669,7 +669,7 @@ class alignas(16) RunnableTask
    // If team then only one thread calls destructor.

    const bool only_one_thread =
-#if defined(KOKKOS_ACTIVE_EXECUTION_MEMORY_SPACE_CUDA)
+#ifdef __CUDA_ARCH__  // FIXME_CUDA


Highlighting

Co-Authored-By: Daniel Arndt <arndtd@ornl.gov>

ajpowelsnl · 2022-02-21T15:40:04Z

Morning @dalg24 -- Is it ok to close this PR ?

dalg24 · 2022-02-21T15:41:34Z

What do you mean close a PR? There is nothing to close it is already merged. You may remove the "Blocks Promotion" label if you like.

ajpowelsnl · 2022-02-21T15:48:39Z

What do you mean close a PR? There is nothing to close it is already merged. You may remove the "Blocks Promotion" label if you like.

OK, sounds good. Sorry for the clumsiness , just trying to figure out how to handle the flow of information, so that nothing slips through the cracks.

Remove use of ActiveExecutionMemorySpace for compatibility with disable of kokkos deprecated code See kokkos/kokkos#4668 and kokkos/kokkos#4660

dalg24 mentioned this pull request Jan 11, 2022

Drop KOKKOS_IMPL_ENABLE_OVERLOAD_HOST_DEVICE leftovers #4661

Merged

dalg24 force-pushed the if_on_host branch from eaa2662 to 22b05cf Compare January 11, 2022 21:46

This was referenced Jan 12, 2022

Cleanup unnecessary #ifdef KOKKOS_ACTIVE_EXECUTION_MEMORY_SPACE_HOST guards #4666

Merged

Deprecate/remove KOKKOS_ACTIVE_EXECUTION_MEMORY_SPACE_* macros and the ActiveExecutionMemorySpace alias #4668

Closed

masterleinad reviewed Jan 12, 2022

View reviewed changes

crtrott mentioned this pull request Jan 13, 2022

Fix failing detection of view params bugfix #4618

Merged

dalg24 added 6 commits January 14, 2022 11:51

Fix typos in comments

b85b6a6

Implement KOKKOS_IF_{DEVICE,HOST} macros

6ba83d1

Use KOKKOS_IF_{HOST,DEVICE} to get rid of the OpenMPTarget workaround…

aefaa35

… in the dimension runtime check in the View constructor

Use KOKKOS_IF_{HOST,DEVICE} to increment HostSharedPtr reference coun…

cad1cc6

…ter on the host side

Replace KOKKOS_IMPL_IF_ON_HOST OpenMPTarget workaround in SharedAlloc…

354bef1

…ationRecord implementation

Use NV_IF_TARGET macro

7b2fe12

dalg24 force-pushed the if_on_host branch 2 times, most recently from e0672ee to 1b09cfb Compare January 14, 2022 17:47

dalg24 added 8 commits January 14, 2022 17:02

Get rid of #ifdef KOKKOS_ACTIVE_EXECUTION_MEMORY_SPACE_HOST in all Of…

761a6e1

…fetView constructors

Remove commented code in OffsetView header

f02b9b3

Implement runtime_check_memory_access_violation<MemorySpace>()

66c8376

Update Kokkos::View runtime checks for memory access violation

e6072da

Update {DynRank,Offset}View runtime checks for memory access violatio…

10983c0

…n and improve error message

Refactor DynamicView runtime checks for memory access violation

4e46441

Refactor MemoryPool::superblock_state() code path if memory space is …

64cb0ce

…accessible

Fix spelling inter[-]relationship in comment

a149786

dalg24 force-pushed the if_on_host branch from 1b09cfb to 0b481a0 Compare January 14, 2022 22:02

dalg24 and others added 7 commits January 14, 2022 21:25

Deprecate Impl::verify_space and Impl::ActiveExecutionMemorySpace

c54cb5b

Co-Authored-By: Christian Trott <crtrott@sandia.gov>

Get rid of most occurrences of #ifdef KOKKOS_ACTIVE_EXECUTION_MEMORY_…

0da1270

…SPACE_HOST

Update commented code in view mapping unit test

a9623f5

Fix code with #ifdef KOKKOS_ACTIVE_EXECUTION_MEMORY_SPACE_HOST that h…

0570f72

…ad nested #ifdef

Update all host specializations of SharedAllocationRecord::allocate()

c25dc02

Temporary workaround to get rid of #ifdef KOKKOS_ACTIVE_EXECUTION_MEM…

382909a

…ORY_SPACE_CUDA

Drop __host__ __device__ annotation in reduction_identity<long double…

4bde4c2

…> and drop macro guards

dalg24 force-pushed the if_on_host branch from 00ad64f to 4bde4c2 Compare January 15, 2022 02:34

crtrott marked this pull request as ready for review January 19, 2022 00:39

crtrott approved these changes Jan 19, 2022

View reviewed changes

crtrott added the Blocks Promotion Overview issue for release-blocking bugs label Jan 19, 2022

nmm0 added this to In progress in Kokkos Release 3.6 via automation Jan 19, 2022

Per review rename KOKKOS_IF_{HOST,DEVICE} -> KOKKOS_IF_ON_{HOST,DEVICE}

beff2f0

dalg24 mentioned this pull request Jan 19, 2022

Refactor Algo::Level{2,3}::Blocked::mb() kokkos/kokkos-kernels#1265

Merged

dalg24 changed the title ~~Implement KOKKOS_IF_{HOST,DEVICE} macros~~ Implement KOKKOS_IF_ON_{HOST,DEVICE} macros Jan 19, 2022

masterleinad reviewed Jan 20, 2022

View reviewed changes

Prefer immediately invoked function expression

d4a7d45

Co-Authored-By: Daniel Arndt <arndtd@ornl.gov>

masterleinad approved these changes Jan 21, 2022

View reviewed changes

dalg24 merged commit c74499d into kokkos:develop Jan 21, 2022

Kokkos Release 3.6 automation moved this from In progress to Done Jan 21, 2022

dalg24 deleted the if_on_host branch January 21, 2022 16:00

dalg24 mentioned this pull request Jan 21, 2022

Rename KOKKOS_IF_{HOST,DEVICE} -> KOKKOS_IF_ON_{HOST,DEVICE} kokkos/kokkos-kernels#1278

Merged

mzuzek mentioned this pull request Jan 25, 2022

Fix out-of-sync Kokkos items kokkos/kokkos-kernels#1288

Closed

ajpowelsnl removed the Blocks Promotion Overview issue for release-blocking bugs label Feb 21, 2022

dalg24 mentioned this pull request Apr 7, 2022

Fix bug in explicit specialization of RuntimeCheckMemoryAccessViolation #4946

Merged

dalg24 mentioned this pull request May 6, 2022

Refactor BitOps to not rely on __CUDA_ARCH__ macro #5001

Merged

mzuzek mentioned this pull request Jul 29, 2022

Move SerialGEMV to KokkosBlas kokkos/kokkos-kernels#1433

Merged

12 tasks

ndellingwood added a commit to trilinos/Trilinos that referenced this pull request Aug 30, 2022

ifpack2: updates for kokkos deprecated code

7152460

Remove use of ActiveExecutionMemorySpace for compatibility with disable of kokkos deprecated code See kokkos/kokkos#4668 and kokkos/kokkos#4660

ndellingwood added a commit to trilinos/Trilinos that referenced this pull request Aug 30, 2022

ifpack2: updates for kokkos deprecated code

5c2eac3

Remove use of ActiveExecutionMemorySpace for compatibility with disable of kokkos deprecated code See kokkos/kokkos#4668 and kokkos/kokkos#4660

ndellingwood mentioned this pull request Aug 30, 2022

Kokkos + KokkosKernels Promotion To 3.7.00 trilinos/Trilinos#10962

Merged

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

Implement `KOKKOS_IF_ON_{HOST,DEVICE}` macros #4660

Implement `KOKKOS_IF_ON_{HOST,DEVICE}` macros #4660

dalg24 commented Jan 11, 2022

masterleinad left a comment

crtrott left a comment

crtrott commented Jan 19, 2022

janciesko commented Jan 19, 2022 •

edited

dalg24 commented Jan 20, 2022

masterleinad left a comment

masterleinad Jan 20, 2022

masterleinad Jan 20, 2022

masterleinad Jan 20, 2022

masterleinad Jan 20, 2022

dalg24 Jan 20, 2022

masterleinad Jan 20, 2022

dalg24 Jan 20, 2022

masterleinad Jan 20, 2022

dalg24 Jan 20, 2022

masterleinad Jan 20, 2022

masterleinad Jan 20, 2022

ajpowelsnl commented Feb 21, 2022

dalg24 commented Feb 21, 2022

ajpowelsnl commented Feb 21, 2022

		// No __host__ __device__ annotation because long double treated as double in
		// device code. May be revisited later if that is not true any more.

Implement KOKKOS_IF_ON_{HOST,DEVICE} macros #4660

Implement KOKKOS_IF_ON_{HOST,DEVICE} macros #4660

Conversation

dalg24 commented Jan 11, 2022

masterleinad left a comment

Choose a reason for hiding this comment

crtrott left a comment

Choose a reason for hiding this comment

crtrott commented Jan 19, 2022

janciesko commented Jan 19, 2022 • edited

dalg24 commented Jan 20, 2022

masterleinad left a 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

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

ajpowelsnl commented Feb 21, 2022

dalg24 commented Feb 21, 2022

ajpowelsnl commented Feb 21, 2022

Implement `KOKKOS_IF_ON_{HOST,DEVICE}` macros #4660

Implement `KOKKOS_IF_ON_{HOST,DEVICE}` macros #4660

janciesko commented Jan 19, 2022 •

edited