view(uvm): don't fence during view's allocation (#6005) #6008

romintomasetti · 2023-03-23T12:54:01Z

This PR is a fix for

Allocation of Kokkos::View given a space instance and Cuda UVM #6005

Basically, when CudaUVM is enabled, there is a fence before and after the view's allocation (whether the space instance was given via arg_prop or not).

However, the fence was not performed on the space instance the user provided (if any).

According to the "poll" below, it was decided to

enforce Kokkos semantics: Don't ever fence for the view initialization when an execution space instance is given

Therefore, any additional fence that was added for CudaUVM in view's allocation has been removed. Note that fencing still occurs in the backend implementations.

dalg24-jenkins · 2023-03-23T12:54:03Z

Can one of the admins verify this patch?

containers/src/Kokkos_DynRankView.hpp

core/src/Kokkos_View.hpp

containers/src/Kokkos_DynRankView.hpp

masterleinad · 2023-03-29T19:20:25Z

The developer meeting asked for more clarification of what we are discussing:

Consider

Kokkos::Cuda exec_space;
Kokkos::View<int*, Kokkos::CudaUVMSpace> view(Kokkos::view_alloc(exec_space, "view"), 10);
view[0] = 2;

should this code work or do we require the user to fence?
What about

Kokkos::Experimental::HPX exec_space;
Kokkos::View<int*, Kokkos::HostSpace> view(Kokkos::view_alloc(exec_space, "view"), 10);
view[0] = 2;

? Currently, we are fencing the default execution space instance(!) before allocation and after (potential initialization) only for Kokkos::CudaUVMSpace. Note that fences for the allocation itself are in the backend implementations.
I would think that we can get rid of the first fence for sure. Can we also get rid of the second one in case the user provided an execution space instance? Do we care to check for Kokkos::WithoutInitializing for this fence additionally?

dalg24 · 2023-03-29T19:31:07Z

The developer meeting asked for more clarification of what we are discussing:

Consider
Kokkos::Cuda exec_space;
Kokkos::View<int*, Kokkos::CudaUVMSpace> view(Kokkos::view_alloc(exec_space, "view"), 10);
view[0] = 2;
should this code work or do we require the user to fence? What about
Kokkos::Experimental::HPX exec_space;
Kokkos::View<int*, Kokkos::HostSpace> view(Kokkos::view_alloc(exec_space, "view"), 10);
view[0] = 2;

In my opinion, yes both need fences.

romintomasetti · 2023-03-31T10:23:37Z

The developer meeting asked for more clarification of what we are discussing:

Consider
Kokkos::Cuda exec_space;
Kokkos::View<int*, Kokkos::CudaUVMSpace> view(Kokkos::view_alloc(exec_space, "view"), 10);
view[0] = 2;
should this code work or do we require the user to fence? What about
Kokkos::Experimental::HPX exec_space;
Kokkos::View<int*, Kokkos::HostSpace> view(Kokkos::view_alloc(exec_space, "view"), 10);
view[0] = 2;
? Currently, we are fencing the default execution space instance(!) before allocation and after (potential initialization) only for Kokkos::CudaUVMSpace. Note that fences for the allocation itself are in the backend implementations. I would think that we can get rid of the first fence for sure. Can we also get rid of the second one in case the user provided an execution space instance? Do we care to check for Kokkos::WithoutInitializing for this fence additionally?

Hi @dalg24 @masterleinad !

We tried to think about this a bit more. Here is a table that show what we have in mind. Note the introduction of the opt-in Kokkos::Async flag, that allows the user to say "I'm good with async".

In our opinion, the behavior for a shared memory space should be similar to any other one.

space instance	`Kokkos::view_alloc`	expected behavior
no space instance	❌	Use default space instance, and fence it at the end of `Kokkos::View` constructor (view is 'ready')
no space instance	`Kokkos::Async`	Use default space instance, don't fence it at the end of the constructor (view might not be 'ready' on host, but is ready on the FIFO queue of the default `space` instance)
space instance provided	❌	Use provided `space` instance and fence it at the end of the constructor (view is 'ready')
space instance provided	`Kokkos::Async`	Use provided `space` instance and don't fence it (view might not be 'ready' on host but is ready on the FIFO queue of the provided `space` instance)

core/src/Cuda/Kokkos_CudaSpace.cpp

Rombur · 2023-03-31T14:14:28Z

Note the introduction of the opt-in Kokkos::Async flag, that allows the user to say "I'm good with async".

There is no need for this flag. If you give an execution space , you agree that the operation may be asynchronous. If you don't provide an execution space. we should fence for you. If you provide an execution space, then it's your responsibility.

masterleinad · 2023-03-31T14:53:57Z

Let's just do a simple poll that might decide pretty much already how we want to proceed:

enforce Kokkos semantics: Don't ever fence for the view initialization when an execution space instance is given 🚀
fence in some cases even if an execution space instance is given, e.g. for UVM allocations 🎉

romintomasetti · 2023-04-04T08:19:19Z

Hi @masterleinad ! 👋

I guess your poll is unanimous: Kokkos semantics should always be enforced, which in this case implies this rule:

Don't ever fence for the view initialization when an execution space instance is given

If you agree, I'll make the changes.

masterleinad · 2023-04-04T12:43:07Z

If you agree, I'll make the changes.

Yes, just delete the UVM fences.

containers/unit_tests/TestWithoutInitializing.hpp

core/src/Cuda/Kokkos_CudaSpace.cpp

core/src/Cuda/Kokkos_Cuda_Instance.cpp

core/unit_test/TestWithoutInitializing.hpp

romintomasetti · 2023-04-14T09:35:30Z

@masterleinad I guess the reason for failure of GCC-8.4.0 is the same as before. Do I need to worry about the failure of SYCL-OneAPI?

masterleinad · 2023-04-14T12:41:24Z

@masterleinad I guess the reason for failure of GCC-8.4.0 is the same as before. Do I need to worry about the failure of SYCL-OneAPI?

This doesn't really touch SYCL and we have seen the failures elsewhere. I wouldn't worry about it (but we might decide to merge a fix first and rerun CI just to make sure).

romintomasetti · 2023-04-19T14:33:35Z

@masterleinad The last action showed only GCC-8.4.0 failing. Should we proceed with the PR?

masterleinad · 2023-04-19T14:40:28Z

@masterleinad The last action showed only GCC-8.4.0 failing. Should we proceed with the PR?

I already approved it. So you need to bug other people (possibly someone with the power to merge it). 🙂

romintomasetti · 2023-04-19T14:42:48Z

Indeed, sorry @masterleinad ! I guess I need to bug @dalg24 and @PhilMiller 📨

maartenarnst · 2023-04-24T18:39:39Z

Hi @dalg24 and @PhilMiller, would you have a moment to look at this PR, and, if you agree with the changes, approve? Many thanks in advance!

containers/src/Kokkos_DualView.hpp

containers/unit_tests/TestWithoutInitializing.hpp

core/src/impl/Kokkos_HostSpace_deepcopy.hpp

dalg24 · 2023-05-01T17:23:53Z

Make sure you update the title line and the description of this PR to reflect the latest version.

PhilMiller · 2023-05-01T17:53:59Z

containers/unit_tests/TestWithoutInitializing.hpp

+//! @name Some tests are skipped for @c CudaUVM memory space.
+///@{
+#ifdef KOKKOS_ENABLE_CUDA
+#define GTEST_SKIP_IF_CUDAUVM_MEMORY_SPACE                            \


Other discussion aside, standard practice when defining a macro like this is to wrap the contents in do { ... } while (false) so that the use sites have to treat it like a statement with a trailing ;

It is now

#ifdef KOKKOS_ENABLE_CUDA #define GTEST_SKIP_IF_CUDAUVM_MEMORY_SPACE \ if constexpr (std::is_same_v<typename TEST_EXECSPACE::memory_space, \ Kokkos::CudaUVMSpace>) \ GTEST_SKIP() << "skipping since CudaUVMSpace requires additional fences" #else #define GTEST_SKIP_IF_CUDAUVM_MEMORY_SPACE do {} while(false) #endif

such that the caller must add the trailing ;. Is that what you meant? :)

I don't agree with Phil's suggestion. This only apply to macro that take arguments.
I would prefer to revert but I am not blocking on this.

I also preferred how it was before. I'll revert.

@dalg24 I think this is the last thread. If you agree with all the changes, this PR is good to go 🚀

romintomasetti · 2023-05-03T19:24:43Z

core/src/Cuda/Kokkos_CudaSpace.cpp

+      Impl::cuda_device_synchronize(
+          "Kokkos::Cuda: backend fence after async malloc");


@masterleinad This is just a comment about the clang-format. Do you think it would make sense (not in this PR of course) to lower the penalty of excess characters? (See https://clang.llvm.org/docs/ClangFormatStyleOptions.html#penaltyexcesscharacter)

I mean, sometimes a line with 87 characters might be nicer than the same line split weirdly across 2 lines 😄 This is just a question, not suggesting to do anything but gathering your opinion on this 😄

We need to update the clang-format version we are using to something more recent in the near future.
This will trigger changes throughout the code base. When we do so would be a time where we would consider changing the clang-format configuration. If that's something you are interested in you are welcome to explore options but this needs to be done in an issue elsewhere.
A word of warning, this is a time sink and at times can get controversial.

If I'm not mistaken, the PenaltyExcessCharacter is available from clang-format 3.7, and you're requiring version 8.0 in scripts/apply-clang-format. So I guess the option could be readily used.

Instead of modifying the whole code base when the clang format changes (either version or options), I think you should use git-clang-format to check for style only in modified files. This would avoid committing purely stylistic changes. That would avoid messing up the git history.

We allow disabling clang-format when appropriate. You will find all sort of examples throughout the code base such as

kokkos/core/src/impl/Kokkos_QuadPrecisionMath.hpp

Lines 46 to 67 in 4b6d971

// clang-format off

// Numeric distinguished value traits

KOKKOS_IMPL_SPECIALIZE_NUMERIC_TRAIT(infinity, __float128, __float128, HUGE_VALQ)

KOKKOS_IMPL_SPECIALIZE_NUMERIC_TRAIT(finite_min, __float128, __float128, -FLT128_MAX)

KOKKOS_IMPL_SPECIALIZE_NUMERIC_TRAIT(finite_max, __float128, __float128, FLT128_MAX)

KOKKOS_IMPL_SPECIALIZE_NUMERIC_TRAIT(epsilon, __float128, __float128, FLT128_EPSILON)

KOKKOS_IMPL_SPECIALIZE_NUMERIC_TRAIT(round_error, __float128, __float128, static_cast<__float128>(0.5))

KOKKOS_IMPL_SPECIALIZE_NUMERIC_TRAIT(norm_min, __float128, __float128, FLT128_MIN)

KOKKOS_IMPL_SPECIALIZE_NUMERIC_TRAIT(denorm_min, __float128, __float128, FLT128_DENORM_MIN)

KOKKOS_IMPL_SPECIALIZE_NUMERIC_TRAIT(quiet_NaN, __float128, __float128, __builtin_nanq(""))

KOKKOS_IMPL_SPECIALIZE_NUMERIC_TRAIT(signaling_NaN, __float128, __float128, __builtin_nansq(""))

// Numeric characteristics traits

KOKKOS_IMPL_SPECIALIZE_NUMERIC_TRAIT(digits, __float128, int, FLT128_MANT_DIG)

KOKKOS_IMPL_SPECIALIZE_NUMERIC_TRAIT(digits10, __float128, int, FLT128_DIG)

KOKKOS_IMPL_SPECIALIZE_NUMERIC_TRAIT(max_digits10, __float128, int, 36)

KOKKOS_IMPL_SPECIALIZE_NUMERIC_TRAIT(radix, __float128, int, 2)

KOKKOS_IMPL_SPECIALIZE_NUMERIC_TRAIT(min_exponent, __float128, int, FLT128_MIN_EXP)

KOKKOS_IMPL_SPECIALIZE_NUMERIC_TRAIT(max_exponent, __float128, int, FLT128_MAX_EXP)

KOKKOS_IMPL_SPECIALIZE_NUMERIC_TRAIT(min_exponent10, __float128, int, FLT128_MIN_10_EXP)

KOKKOS_IMPL_SPECIALIZE_NUMERIC_TRAIT(max_exponent10, __float128, int, FLT128_MAX_10_EXP)

// clang-format on

or

kokkos/core/src/OpenACC/Kokkos_OpenACC_ParallelFor_Range.hpp

Lines 35 to 37 in 4b6d971

// clang-format off

#pragma acc parallel loop gang(static:chunk_size) vector copyin(functor) async(async_arg)

// clang-format on

In principle you are right but in practice I think it would mean that the entire file would be reformatted as soon as someone touches it and this "spreads" the noise across the history of the repo instead of having one clean "break".

dalg24 · 2023-05-03T19:37:10Z

containers/unit_tests/TestWithoutInitializing.hpp

+            (event.descriptor().find(fence_event_message_for_zeromemset(
+                 TEST_EXECSPACE{})) != std::string::npos))


This change is unnecessary. Please undo.

containers/unit_tests/TestWithoutInitializing.hpp

dalg24 · 2023-05-03T19:40:11Z

containers/unit_tests/TestWithoutInitializing.hpp

+/**
+ * @name Some tests are skipped for @c CudaUVM memory space.
+ *
+ * @todo To be revised according to the future of @c KOKKOS_ENABLE_CUDA_UVM .
+ */
+///@{


Not requiring changes but these are 5 lines of comment (not even counting the closing comment) when one single line would do and we do not really use Doxygen so chances are it would not even generate correct documentation after other edit the file.

I've used /// instead of /** ... */. Tell me if it's good for you ;)

dalg24 · 2023-05-03T19:41:02Z

(almost there)

romintomasetti mentioned this pull request Mar 23, 2023

Allocation of Kokkos::View given a space instance and Cuda UVM #6005

Closed

romintomasetti force-pushed the cuda-uvm-space-instance-fence branch from 1f3c50f to cfc6ce4 Compare March 23, 2023 13:05

masterleinad reviewed Mar 23, 2023

View reviewed changes

containers/src/Kokkos_DynRankView.hpp Outdated Show resolved Hide resolved

containers/src/Kokkos_DynRankView.hpp Outdated Show resolved Hide resolved

romintomasetti force-pushed the cuda-uvm-space-instance-fence branch from cfc6ce4 to 2d8069e Compare March 23, 2023 14:00

dalg24 reviewed Mar 23, 2023

View reviewed changes

core/src/Kokkos_View.hpp Outdated Show resolved Hide resolved

dalg24 reviewed Mar 23, 2023

View reviewed changes

core/src/Kokkos_View.hpp Outdated Show resolved Hide resolved

romintomasetti force-pushed the cuda-uvm-space-instance-fence branch 3 times, most recently from d26e6cc to e6d2ee2 Compare March 23, 2023 14:57

masterleinad reviewed Mar 23, 2023

View reviewed changes

core/src/Kokkos_View.hpp Outdated Show resolved Hide resolved

containers/src/Kokkos_DynRankView.hpp Outdated Show resolved Hide resolved

romintomasetti force-pushed the cuda-uvm-space-instance-fence branch 3 times, most recently from ad29115 to 33bed43 Compare March 24, 2023 08:40

romintomasetti force-pushed the cuda-uvm-space-instance-fence branch from 33bed43 to 391ccd3 Compare March 31, 2023 11:34

romintomasetti commented Mar 31, 2023

View reviewed changes

core/src/Cuda/Kokkos_CudaSpace.cpp Show resolved Hide resolved

romintomasetti force-pushed the cuda-uvm-space-instance-fence branch from 391ccd3 to 47f4b6b Compare April 4, 2023 14:20

masterleinad reviewed Apr 4, 2023

View reviewed changes

containers/unit_tests/TestWithoutInitializing.hpp Outdated Show resolved Hide resolved

romintomasetti force-pushed the cuda-uvm-space-instance-fence branch from 47f4b6b to 9b5e667 Compare April 4, 2023 15:08

PhilMiller reviewed Apr 4, 2023

View reviewed changes

core/src/Cuda/Kokkos_CudaSpace.cpp Show resolved Hide resolved

PhilMiller reviewed Apr 4, 2023

View reviewed changes

core/src/Cuda/Kokkos_Cuda_Instance.cpp Outdated Show resolved Hide resolved

PhilMiller reviewed Apr 4, 2023

View reviewed changes

core/unit_test/TestWithoutInitializing.hpp Outdated Show resolved Hide resolved

romintomasetti mentioned this pull request Apr 5, 2023

Encapsulate CudaAPI calls in CudaInternal, call cudaSetDevice for thread safety #5989

Closed

4 tasks

masterleinad approved these changes Apr 13, 2023

View reviewed changes

romintomasetti force-pushed the cuda-uvm-space-instance-fence branch from 46e4030 to ef07355 Compare April 19, 2023 12:19

romintomasetti requested review from dalg24 and PhilMiller April 19, 2023 14:41

dalg24 reviewed Apr 26, 2023

View reviewed changes

containers/src/Kokkos_DualView.hpp Show resolved Hide resolved

containers/unit_tests/TestWithoutInitializing.hpp Show resolved Hide resolved

core/src/impl/Kokkos_HostSpace_deepcopy.hpp Outdated Show resolved Hide resolved

PhilMiller reviewed May 1, 2023

View reviewed changes

romintomasetti force-pushed the cuda-uvm-space-instance-fence branch from ef07355 to e395d5f Compare May 2, 2023 08:04

romintomasetti changed the title ~~view(uvm): fence if need in allocation (#6005)~~ view(uvm): don't fence during view's allocation (#6005) May 2, 2023

romintomasetti force-pushed the cuda-uvm-space-instance-fence branch 5 times, most recently from cb17816 to 0a47643 Compare May 2, 2023 16:07

romintomasetti commented May 3, 2023

View reviewed changes

dalg24 reviewed May 3, 2023

View reviewed changes

view(uvm): fence if need in allocation (kokkos#6005)

5c2d948

romintomasetti force-pushed the cuda-uvm-space-instance-fence branch from 0a47643 to 5c2d948 Compare May 4, 2023 06:56

dalg24 approved these changes May 4, 2023

View reviewed changes

dalg24 merged commit d6944df into kokkos:develop May 4, 2023
28 checks passed

dalg24 mentioned this pull request May 4, 2023

CHANGELOG: 4.1.0 #5902

Closed

romintomasetti mentioned this pull request May 16, 2023

Stokhos: Tests will fail after https://github.com/kokkos/kokkos/pull/6008 for shared spaces (Cuda UVM at least) trilinos/Trilinos#11890

Closed

ndellingwood mentioned this pull request May 23, 2023

Trilinos nightly test failures in Stokhos tests with Cuda builds, UVM enabled #6158

Closed

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

view(uvm): don't fence during view's allocation (#6005) #6008

view(uvm): don't fence during view's allocation (#6005) #6008

romintomasetti commented Mar 23, 2023 •

edited

dalg24-jenkins commented Mar 23, 2023

masterleinad commented Mar 29, 2023

dalg24 commented Mar 29, 2023

romintomasetti commented Mar 31, 2023

Rombur commented Mar 31, 2023

masterleinad commented Mar 31, 2023 •

edited

romintomasetti commented Apr 4, 2023

masterleinad commented Apr 4, 2023

romintomasetti commented Apr 14, 2023

masterleinad commented Apr 14, 2023

romintomasetti commented Apr 19, 2023

masterleinad commented Apr 19, 2023

romintomasetti commented Apr 19, 2023

maartenarnst commented Apr 24, 2023

dalg24 commented May 1, 2023

PhilMiller May 1, 2023

romintomasetti May 2, 2023 •

edited

dalg24 May 2, 2023

romintomasetti May 2, 2023

romintomasetti May 3, 2023

romintomasetti May 3, 2023

dalg24 May 3, 2023

romintomasetti May 4, 2023

dalg24 May 4, 2023

dalg24 May 3, 2023

dalg24 May 3, 2023

romintomasetti May 4, 2023

dalg24 commented May 3, 2023

		Impl::cuda_device_synchronize(
		"Kokkos::Cuda: backend fence after async malloc");

	// clang-format off
	// Numeric distinguished value traits
	KOKKOS_IMPL_SPECIALIZE_NUMERIC_TRAIT(infinity, __float128, __float128, HUGE_VALQ)
	KOKKOS_IMPL_SPECIALIZE_NUMERIC_TRAIT(finite_min, __float128, __float128, -FLT128_MAX)
	KOKKOS_IMPL_SPECIALIZE_NUMERIC_TRAIT(finite_max, __float128, __float128, FLT128_MAX)
	KOKKOS_IMPL_SPECIALIZE_NUMERIC_TRAIT(epsilon, __float128, __float128, FLT128_EPSILON)
	KOKKOS_IMPL_SPECIALIZE_NUMERIC_TRAIT(round_error, __float128, __float128, static_cast<__float128>(0.5))
	KOKKOS_IMPL_SPECIALIZE_NUMERIC_TRAIT(norm_min, __float128, __float128, FLT128_MIN)
	KOKKOS_IMPL_SPECIALIZE_NUMERIC_TRAIT(denorm_min, __float128, __float128, FLT128_DENORM_MIN)
	KOKKOS_IMPL_SPECIALIZE_NUMERIC_TRAIT(quiet_NaN, __float128, __float128, __builtin_nanq(""))
	KOKKOS_IMPL_SPECIALIZE_NUMERIC_TRAIT(signaling_NaN, __float128, __float128, __builtin_nansq(""))

	// Numeric characteristics traits
	KOKKOS_IMPL_SPECIALIZE_NUMERIC_TRAIT(digits, __float128, int, FLT128_MANT_DIG)
	KOKKOS_IMPL_SPECIALIZE_NUMERIC_TRAIT(digits10, __float128, int, FLT128_DIG)
	KOKKOS_IMPL_SPECIALIZE_NUMERIC_TRAIT(max_digits10, __float128, int, 36)
	KOKKOS_IMPL_SPECIALIZE_NUMERIC_TRAIT(radix, __float128, int, 2)
	KOKKOS_IMPL_SPECIALIZE_NUMERIC_TRAIT(min_exponent, __float128, int, FLT128_MIN_EXP)
	KOKKOS_IMPL_SPECIALIZE_NUMERIC_TRAIT(max_exponent, __float128, int, FLT128_MAX_EXP)
	KOKKOS_IMPL_SPECIALIZE_NUMERIC_TRAIT(min_exponent10, __float128, int, FLT128_MIN_10_EXP)
	KOKKOS_IMPL_SPECIALIZE_NUMERIC_TRAIT(max_exponent10, __float128, int, FLT128_MAX_10_EXP)
	// clang-format on

	// clang-format off
	#pragma acc parallel loop gang(static:chunk_size) vector copyin(functor) async(async_arg)
	// clang-format on

		(event.descriptor().find(fence_event_message_for_zeromemset(
		TEST_EXECSPACE{})) != std::string::npos))

view(uvm): don't fence during view's allocation (#6005) #6008

view(uvm): don't fence during view's allocation (#6005) #6008

Conversation

romintomasetti commented Mar 23, 2023 • edited

dalg24-jenkins commented Mar 23, 2023

masterleinad commented Mar 29, 2023

dalg24 commented Mar 29, 2023

romintomasetti commented Mar 31, 2023

Rombur commented Mar 31, 2023

masterleinad commented Mar 31, 2023 • edited

romintomasetti commented Apr 4, 2023

masterleinad commented Apr 4, 2023

romintomasetti commented Apr 14, 2023

masterleinad commented Apr 14, 2023

romintomasetti commented Apr 19, 2023

masterleinad commented Apr 19, 2023

romintomasetti commented Apr 19, 2023

maartenarnst commented Apr 24, 2023

dalg24 commented May 1, 2023

Choose a reason for hiding this comment

romintomasetti May 2, 2023 • edited

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

dalg24 commented May 3, 2023

romintomasetti commented Mar 23, 2023 •

edited

masterleinad commented Mar 31, 2023 •

edited

romintomasetti May 2, 2023 •

edited