[SYCL][CUDA] Ignore cuda prefetch hint if not supported #5043

AidanBeltonS · 2021-11-29T16:50:29Z

Specific devices and OS's, like Windows, do not support concurrent managed memory. cudaPrefetchAsync requires concurrent managed access for unified memory. This PR removes the windows error message and replaces it with a check for concurrent managed access. As the SYCL prefetch operation is a hint, this can return a success.

Let me know if there is a preferred error code to throw. Also, if it is best that a user warning is printed to indicate that the hint is being ignored as the device does not support the operation.

steffenlarsen · 2021-11-29T17:24:12Z

sycl/plugins/cuda/pi_cuda.cpp

+                       CU_DEVICE_ATTRIBUTE_CONCURRENT_MANAGED_ACCESS,
+                       queue->get_context()->get_device()->get());
+  if (!isConcurrentManagedAccessAvailable)
+    return PI_SUCCESS;


I think a warning would make sense, but I wonder what the best way to do it would be. I am not sure printing is the best option - at least not from the PI implementation. Maybe a better solution is to add a new "error" code, such as PI_UNSUPPORTED_OPERATION, then the runtime can handle it accordingly.

On the other hand, prefetch is a "hint" so ignoring the operation entirely when we don't support it is arguably still valid.

+1 for ignoring a hint if it cannot be used/not yet implemented.

For letting user know what happens we could use SYCL_RT_WARNING_LEVEL which is being introduced here #4918

Would SYCL_RT_WARNING_LEVEL be available in PI plugins? Otherwise it would still have to inform the runtime somehow.

Interesting question. I tend to say that SYCL_RT_WARNING_LEVEL should be available for the whole SYCL stack. Propagating detailed warning(which not only says something is not done, but also says why) to the core SYCL RT could be problematic.
Tagging @v-klochkov

I have taken a look at using SYCL_RT_WARNING_LEVEL within the cuda pi. However, it is unfortunately not available with pi_cuda. I tried a few things to get access to SYCLConfig, and the warning level. I believe making it available within the plugin would not be simple, as you would have to avoid cyclical dependencies between libpi_cuda.so and libsycl.so.
The easiest way would be to get the environment variable within the cuda pi. However, this has the downside of code duplication. Is there a preference on how to go forward with this? Many thanks

@smaslov-intel, @v-klochkov, can you comment please?

I think it is OK to re-read the SYCL_RT_WARNING_LEVEL from plugins. We already do similar for SYCL_PI_TRACE inside Level-Zero plugin:

llvm/sycl/plugins/level_zero/pi_level_zero.cpp

Line 1822 in 8213321

static const char *PiTrace = std::getenv("SYCL_PI_TRACE");

On the other hand maybe it is a good time to add formal support for backend-specific warnings/errors. Something like returning fixed PI_PLUGIN_SPECIFIC_ERROR, which would indicate that some warning/error is generated and can be obtained by a subsequent call to [new] piPluginGetLastError.

@smaslov-intel I think the PI_PLUGIN_SPECIFIC_ERROR approach would seem the most reasonable way to emit a warning/error message. I have implemented it in 478a546 for the CUDA backend. Let me know if this approach seems okay and I will extend it to the other backends.

AerialMantis · 2022-02-01T14:53:48Z

@steffenlarsen @romanovvlad @v-klochkov just pinging this one as there's not been much activity for a while, I think we just need to decide how to handle the warning, it was suggested that we use SYCL_RT_WARNING_LEVEL, should we wait for #4918 to merge?

romanovvlad · 2022-02-01T19:21:48Z

I would suggest waiting unless there is urgency.
@v-klochkov Could you please comment on when the patch with SYCL_WARNING_LEVEL is planned to be merged?

v-klochkov · 2022-02-04T08:21:49Z

I would suggest waiting unless there is urgency. @v-klochkov Could you please comment on when the patch with SYCL_WARNING_LEVEL is planned to be merged?

GPU RT remains the same after beginning of November'2021. It is still: 21.46.21636. That RT is too old for #4918 . It is good idea to split that patch into 2 parts: a) adding the new env variable and b) adding the switch to VC/GPU backend. Adding (a) would unblock this patch here.

I am on vacation/sabbatical & traveling right now and temporarily cannot split the patch now even though that is very simple work .

AerialMantis · 2022-02-10T08:51:20Z

Thanks for the update, this isn't urgent so I think it can wait, though I did see that #4918 has been merged now, would this be sufficient to use SYCL_WARNING_LEVEL in this patch?

bader · 2022-02-10T10:11:15Z

#4918 has been merged now

#4918 is still open.

AerialMantis · 2022-02-10T11:08:13Z

Sorry pasted the wrong PR, I meant this one - #5319

v-klochkov · 2022-02-23T04:38:59Z

Sorry pasted the wrong PR, I meant this one - #5319

Confirmed. #5319 added/merged the new variable SYCL_RT_WARNING_LEVEL. It can be used in this PR.

sycl/plugins/cuda/pi_cuda.cpp

smaslov-intel · 2022-03-11T21:14:43Z

sycl/include/CL/sycl/detail/pi.h

+/// points to. \param is_warning is a bool indicating if the message is a
+/// non-failing error.
+__SYCL_EXPORT pi_result piPluginGetLastError(char *message, size_t message_size,
+                                             bool *is_warning);


Maybe we can not add "is_warning" but return PI_SUCCESS to indicate that upper RT should treat it as a warning , but return other PI error codes for something that is an error that SYCL RT should try to recover from or throw an exception back to the user.

Also we need a way for SYCL RT to know the size of message array. Or we can just return a pointer directly(without the copy) which can be used by the SYCL RT until the next PI call is done(assuming we take thread_local approach from another comment).

I have updated it to not use is_warning but return a user specified code, so PI_SUCCESS is treated as a warning while others will trigger an error.
Additionally it now uses thread_local so it returns a ptr to the message as suggested.

romanovvlad · 2022-03-23T11:47:17Z

sycl/source/detail/plugin.hpp

@@ -113,11 +114,35 @@ class plugin {
  /// \throw Exception if pi_result is not a PI_SUCCESS.
  template <typename Exception = cl::sycl::runtime_error>
  void checkPiResult(RT::PiResult pi_result) const {
+    if (pi_result == PI_PLUGIN_SPECIFIC_ERROR) {
+      char *message;


Suggested change

char *message;

char *message = nullptr;

I have made this change

romanovvlad · 2022-03-23T11:48:09Z

sycl/source/detail/plugin.hpp

@@ -100,7 +101,7 @@ class plugin {
  plugin &operator=(plugin &&other) noexcept = default;
  plugin(plugin &&other) noexcept = default;

-  ~plugin() = default;
+  ~plugin() {}


Could you please tell why this change is needed?

This change is no longer needed. I have undone it

romanovvlad · 2022-03-23T11:50:11Z

sycl/plugins/cuda/pi_cuda.cpp

@@ -57,6 +58,24 @@ pi_result map_error(CUresult result) {
  }
 }

+// Global variables for PI_PLUGIN_SPECIFIC_ERROR
+static const size_t MaxMessageSize = 256;


Suggested change

static const size_t MaxMessageSize = 256;

constexpr size_t MaxMessageSize = 256;

?

This change has been applied

romanovvlad · 2022-03-23T11:50:51Z

sycl/plugins/cuda/pi_cuda.cpp

@@ -14,6 +14,7 @@
 #include <CL/sycl/detail/cuda_definitions.hpp>
 #include <CL/sycl/detail/defines.hpp>
 #include <CL/sycl/detail/pi.hpp>
+#include <CL/sycl/detail/spinlock.hpp>


Is this header still needed?

No it is not needed any more, I have removed it.

romanovvlad · 2022-03-23T11:51:38Z

sycl/include/CL/sycl/detail/pi.h

+/// API to get Plugin specific warning and error messages.
+/// \param message is a pointer to an array of characters which will be filled
+/// with the error message. \param message_size is the size of the array message
+/// points to. \param is_warning is a bool indicating if the message is a
+/// non-failing error.


It seems the description should be updated.

The description has now been updated.

romanovvlad

LGTM

AidanBeltonS · 2022-04-27T16:36:33Z

@smaslov-intel would you be able to re-review the warning level changes?

smaslov-intel · 2022-04-28T23:26:39Z

sycl/include/CL/sycl/detail/pi.h

@@ -1769,6 +1772,13 @@ __SYCL_EXPORT pi_result piextPluginGetOpaqueData(void *opaque_data_param,
 /// \param PluginParameter placeholder for future use, currenly not used.
 __SYCL_EXPORT pi_result piTearDown(void *PluginParameter);

+/// API to get Plugin specific warning and error messages.
+/// \param message is a returned address to the first element in the message.


Please clarify in the API who owns the storage of the message and what it's lifetime is.

I have clarified the storage and lifetime of the message string

smaslov-intel

LGTM, just a small ask to clarify error message storage ownership.

smaslov-intel

Thanks!

AidanBeltonS · 2022-05-03T16:16:50Z

@pvchupin are there any other reviewers or is this okay to be merged?

1. cuda prefetch issue seems to be fixed by: intel#5043 2. Performance issues with assert seem to be fixed by: intel#4505 intel#4516

* Release notes for commit range f34ba2c..4043dda * Update known issues: 1. cuda prefetch issue seems to be fixed by: #5043 2. Performance issues with assert seem to be fixed by: #4505 #4516

AidanBeltonS requested a review from a team as a code owner November 29, 2021 16:50

AidanBeltonS requested a review from sergey-semenov November 29, 2021 16:50

steffenlarsen reviewed Nov 29, 2021

View reviewed changes

AerialMantis mentioned this pull request Dec 6, 2021

[CUDA] SYCL/USM/dep_events.cpp from llvm-test-suite fails with PI CUDA ERROR in plugins/cuda/pi_cuda.cpp:4674 on Windows #4809

Closed

bader linked an issue Dec 7, 2021 that may be closed by this pull request

[CUDA] SYCL/USM/dep_events.cpp from llvm-test-suite fails with PI CUDA ERROR in plugins/cuda/pi_cuda.cpp:4674 on Windows #4809

Closed

bader added the cuda CUDA back-end label Dec 7, 2021

AidanBeltonS force-pushed the windows-prefetch-error branch from 47c7b58 to 478a546 Compare March 11, 2022 10:19

AidanBeltonS requested review from a team as code owners March 11, 2022 10:19

romanovvlad requested changes Mar 11, 2022

View reviewed changes

sycl/plugins/cuda/pi_cuda.cpp Show resolved Hide resolved

smaslov-intel reviewed Mar 11, 2022

View reviewed changes

romanovvlad requested changes Mar 23, 2022

View reviewed changes

romanovvlad previously approved these changes Mar 24, 2022

View reviewed changes

AidanBeltonS dismissed romanovvlad’s stale review via 98be8c5 March 30, 2022 08:56

smaslov-intel reviewed Apr 28, 2022

View reviewed changes

smaslov-intel previously approved these changes Apr 28, 2022

View reviewed changes

AidanBeltonS dismissed smaslov-intel’s stale review via 244fb58 April 29, 2022 09:04

AidanBeltonS requested a review from a team as a code owner April 29, 2022 09:04

aidan.belton added 4 commits April 29, 2022 11:44

ignore cuda prefetch hint if not supported

102d30a

Add plugin specific errors

17b9ab7

Add include spinlock

1802792

clang format

3e32db2

aidan.belton added 5 commits April 29, 2022 11:44

Add thread local to message

85982d2

Update based on feedback

f91b247

Extend piPluginGetLastError to all backends

f2b6e08

update piPluginGetLastError comment

a053163

rebase and add PI_PLUGIN_SPECIFIC_ERROR to sycl-trace

98d20ee

AidanBeltonS force-pushed the windows-prefetch-error branch from 244fb58 to 98d20ee Compare April 29, 2022 12:27

update symbol check

57e1513

smaslov-intel approved these changes Apr 29, 2022

View reviewed changes

pvchupin merged commit 082929a into intel:sycl May 3, 2022

againull added a commit to againull/llvm that referenced this pull request Jun 17, 2022

Update known issues

37864b2

1. cuda prefetch issue seems to be fixed by: intel#5043 2. Performance issues with assert seem to be fixed by: intel#4505 intel#4516

	static const size_t MaxMessageSize = 256;
	constexpr size_t MaxMessageSize = 256;

[SYCL][CUDA] Ignore cuda prefetch hint if not supported #5043

[SYCL][CUDA] Ignore cuda prefetch hint if not supported #5043

Uh oh!

Conversation

AidanBeltonS commented Nov 29, 2021

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

AerialMantis commented Feb 1, 2022

Uh oh!

romanovvlad commented Feb 1, 2022

Uh oh!

v-klochkov commented Feb 4, 2022

Uh oh!

AerialMantis commented Feb 10, 2022

Uh oh!

bader commented Feb 10, 2022

Uh oh!

AerialMantis commented Feb 10, 2022

Uh oh!

v-klochkov commented Feb 23, 2022

Uh oh!

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

romanovvlad left a comment

Choose a reason for hiding this comment

Uh oh!

AidanBeltonS commented Apr 27, 2022

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

smaslov-intel left a comment

Choose a reason for hiding this comment

Uh oh!

smaslov-intel left a comment

Choose a reason for hiding this comment

Uh oh!

AidanBeltonS commented May 3, 2022

Uh oh!