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

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.

Author: AidanBeltonS

sycl/include/CL/sycl/detail/pi.def

@@ -136,6 +136,9 @@ _PI_API(piextKernelSetArgSampler)
 
 _PI_API(piextPluginGetOpaqueData)
 
+_PI_API(piPluginGetLastError)
+
 _PI_API(piTearDown)
 
+
 #undef _PI_API

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

@@ -117,6 +117,9 @@ typedef enum {
   PI_IMAGE_FORMAT_NOT_SUPPORTED = CL_IMAGE_FORMAT_NOT_SUPPORTED,
   PI_MEM_OBJECT_ALLOCATION_FAILURE = CL_MEM_OBJECT_ALLOCATION_FAILURE,
   PI_LINK_PROGRAM_FAILURE = CL_LINK_PROGRAM_FAILURE,
+  PI_PLUGIN_SPECIFIC_ERROR = -996, ///< PI_PLUGIN_SPECIFIC_ERROR indicates
+                                   ///< that an backend spcific error or
+                                   ///< warning has been emitted by the plugin.
   PI_COMMAND_EXECUTION_FAILURE =
       -997, ///< PI_COMMAND_EXECUTION_FAILURE indicates an error occurred
             ///< during command enqueue or execution.
@@ -1796,6 +1799,18 @@ __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 the
+/// plugin owns the error message string. The string is thread-local. As a
+/// result, different threads may return different errors. A message is
+/// overwritten by the following error or warning that is produced within the
+/// given thread. The memory is cleaned up at the end of the thread's lifetime.
+///
+/// \return PI_SUCCESS if plugin is indicating non-fatal warning. Any other
+/// error code indicates that plugin considers this to be a fatal error and the
+/// runtime must handle it or end the application.
+__SYCL_EXPORT pi_result piPluginGetLastError(char **message);
+
 struct _pi_plugin {
   // PI version supported by host passed to the plugin. The Plugin
   // checks and writes the appropriate Function Pointers in

sycl/plugins/cuda/pi_cuda.cpp

@@ -57,6 +57,24 @@ pi_result map_error(CUresult result) {
   }
 }
 
+// Global variables for PI_PLUGIN_SPECIFIC_ERROR
+constexpr size_t MaxMessageSize = 256;
+thread_local pi_result ErrorMessageCode = PI_SUCCESS;
+thread_local char ErrorMessage[MaxMessageSize];
+
+// Utility function for setting a message and warning
+static void setErrorMessage(const char *message, pi_result error_code) {
+  assert(strlen(message) <= MaxMessageSize);
+  strcpy(ErrorMessage, message);
+  ErrorMessageCode = error_code;
+}
+
+// Returns plugin specific error and warning messages
+pi_result cuda_piPluginGetLastError(char **message) {
+  *message = &ErrorMessage[0];
+  return ErrorMessageCode;
+}
+
 // Iterates over the event wait list, returns correct pi_result error codes.
 // Invokes the callback for the latest event of each queue in the wait list.
 // The callback must take a single pi_event argument and return a pi_result.
@@ -4729,13 +4747,20 @@ pi_result cuda_piextUSMEnqueuePrefetch(pi_queue queue, const void *ptr,
                                        const pi_event *events_waitlist,
                                        pi_event *event) {
 
-// CUDA has an issue with cuMemPrefetchAsync returning cudaErrorInvalidDevice
-// for Windows machines
-// TODO: Remove when fix is found
-#ifdef _MSC_VER
-  cl::sycl::detail::pi::die(
-      "cuda_piextUSMEnqueuePrefetch does not currently work on Windows");
-#endif
+  // Certain cuda devices and Windows do not have support for some Unified
+  // Memory features. cuMemPrefetchAsync requires concurrent memory access
+  // for managed memory. Therfore, ignore prefetch hint if concurrent managed
+  // memory access is not available.
+  int isConcurrentManagedAccessAvailable = 0;
+  cuDeviceGetAttribute(&isConcurrentManagedAccessAvailable,
+                       CU_DEVICE_ATTRIBUTE_CONCURRENT_MANAGED_ACCESS,
+                       queue->get_context()->get_device()->get());
+  if (!isConcurrentManagedAccessAvailable) {
+    setErrorMessage("Prefetch hint ignored as device does not support "
+                    "concurrent managed access",
+                    PI_SUCCESS);
+    return PI_PLUGIN_SPECIFIC_ERROR;
+  }
 
   // flags is currently unused so fail if set
   if (flags != 0)
@@ -5083,6 +5108,7 @@ pi_result piPluginInit(pi_plugin *PluginInit) {
 
   _PI_CL(piextKernelSetArgMemObj, cuda_piextKernelSetArgMemObj)
   _PI_CL(piextKernelSetArgSampler, cuda_piextKernelSetArgSampler)
+  _PI_CL(piPluginGetLastError, cuda_piPluginGetLastError)
   _PI_CL(piTearDown, cuda_piTearDown)
 
 #undef _PI_CL

sycl/plugins/esimd_emulator/pi_esimd_emulator.cpp

@@ -141,6 +141,25 @@ static std::mutex *PiESimdSurfaceMapLock = new std::mutex;
 // For PI_DEVICE_INFO_DRIVER_VERSION info
 static char ESimdEmuVersionString[32];
 
+// Global variables for PI_PLUGIN_SPECIFIC_ERROR
+constexpr size_t MaxMessageSize = 256;
+thread_local pi_result ErrorMessageCode = PI_SUCCESS;
+thread_local char ErrorMessage[MaxMessageSize];
+
+// Utility function for setting a message and warning
+[[maybe_unused]] static void setErrorMessage(const char *message,
+                                             pi_result error_code) {
+  assert(strlen(message) <= MaxMessageSize);
+  strcpy(ErrorMessage, message);
+  ErrorMessageCode = error_code;
+}
+
+// Returns plugin specific error and warning messages
+pi_result piPluginGetLastError(char **message) {
+  *message = &ErrorMessage[0];
+  return ErrorMessageCode;
+}
+
 using IDBuilder = sycl::detail::Builder;
 
 template <int NDims>

sycl/plugins/hip/pi_hip.cpp

@@ -112,6 +112,25 @@ pi_result map_error(hipError_t result) {
   }
 }
 
+// Global variables for PI_PLUGIN_SPECIFIC_ERROR
+constexpr size_t MaxMessageSize = 256;
+thread_local pi_result ErrorMessageCode = PI_SUCCESS;
+thread_local char ErrorMessage[MaxMessageSize];
+
+// Utility function for setting a message and warning
+[[maybe_unused]] static void setErrorMessage(const char *message,
+                                             pi_result error_code) {
+  assert(strlen(message) <= MaxMessageSize);
+  strcpy(ErrorMessage, message);
+  ErrorMessageCode = error_code;
+}
+
+// Returns plugin specific error and warning messages
+pi_result hip_piPluginGetLastError(char **message) {
+  *message = &ErrorMessage[0];
+  return ErrorMessageCode;
+}
+
 // Iterates over the event wait list, returns correct pi_result error codes.
 // Invokes the callback for the latest event of each queue in the wait list.
 // The callback must take a single pi_event argument and return a pi_result.
@@ -4989,6 +5008,7 @@ pi_result piPluginInit(pi_plugin *PluginInit) {
 
   _PI_CL(piextKernelSetArgMemObj, hip_piextKernelSetArgMemObj)
   _PI_CL(piextKernelSetArgSampler, hip_piextKernelSetArgSampler)
+  _PI_CL(piPluginGetLastError, hip_piPluginGetLastError)
   _PI_CL(piTearDown, hip_piTearDown)
 
 #undef _PI_CL

sycl/plugins/level_zero/pi_level_zero.cpp

@@ -589,6 +589,25 @@ inline void zeParseError(ze_result_t ZeError, const char *&ErrorString) {
   } // switch
 }
 
+// Global variables for PI_PLUGIN_SPECIFIC_ERROR
+constexpr size_t MaxMessageSize = 256;
+thread_local pi_result ErrorMessageCode = PI_SUCCESS;
+thread_local char ErrorMessage[MaxMessageSize];
+
+// Utility function for setting a message and warning
+[[maybe_unused]] static void setErrorMessage(const char *message,
+                                             pi_result error_code) {
+  assert(strlen(message) <= MaxMessageSize);
+  strcpy(ErrorMessage, message);
+  ErrorMessageCode = error_code;
+}
+
+// Returns plugin specific error and warning messages
+pi_result piPluginGetLastError(char **message) {
+  *message = &ErrorMessage[0];
+  return ErrorMessageCode;
+}
+
 ze_result_t ZeCall::doCall(ze_result_t ZeResult, const char *ZeName,
                            const char *ZeArgs, bool TraceError) {
   zePrint("ZE ---> %s%s\n", ZeName, ZeArgs);

sycl/plugins/opencl/pi_opencl.cpp

@@ -71,6 +71,25 @@ CONSTFIX char clGetDeviceFunctionPointerName[] =
 
 #undef CONSTFIX
 
+// Global variables for PI_PLUGIN_SPECIFIC_ERROR
+constexpr size_t MaxMessageSize = 256;
+thread_local pi_result ErrorMessageCode = PI_SUCCESS;
+thread_local char ErrorMessage[MaxMessageSize];
+
+// Utility function for setting a message and warning
+[[maybe_unused]] static void setErrorMessage(const char *message,
+                                             pi_result error_code) {
+  assert(strlen(message) <= MaxMessageSize);
+  strcpy(ErrorMessage, message);
+  ErrorMessageCode = error_code;
+}
+
+// Returns plugin specific error and warning messages
+pi_result piPluginGetLastError(char **message) {
+  *message = &ErrorMessage[0];
+  return ErrorMessageCode;
+}
+
 // USM helper function to get an extension function pointer
 template <const char *FuncName, typename T>
 static pi_result getExtFuncFromContext(pi_context context, T *fptr) {
@@ -1543,6 +1562,7 @@ pi_result piPluginInit(pi_plugin *PluginInit) {
 
   _PI_CL(piextKernelSetArgMemObj, piextKernelSetArgMemObj)
   _PI_CL(piextKernelSetArgSampler, piextKernelSetArgSampler)
+  _PI_CL(piPluginGetLastError, piPluginGetLastError)
   _PI_CL(piTearDown, piTearDown)
 
 #undef _PI_CL

sycl/source/detail/common.cpp

@@ -218,6 +218,8 @@ const char *stringifyErrorCode(cl_int error) {
         */
   case PI_FUNCTION_ADDRESS_IS_NOT_AVAILABLE:
     return "Function exists but address is not available";
+  case PI_PLUGIN_SPECIFIC_ERROR:
+    return "The plugin has emitted a backend specific error";
   case PI_COMMAND_EXECUTION_FAILURE:
     return "Command failed to enqueue/execute";
   default:

sycl/source/detail/plugin.hpp

@@ -12,6 +12,7 @@
 #include <CL/sycl/detail/pi.hpp>
 #include <CL/sycl/detail/type_traits.hpp>
 #include <CL/sycl/stl.hpp>
+#include <detail/config.hpp>
 #include <detail/plugin_printers.hpp>
 #include <memory>
 #include <mutex>
@@ -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 = nullptr;
+      pi_result = call_nocheck<PiApiKind::piPluginGetLastError>(&message);
+
+      // If the warning level is greater then 2 emit the message
+      if (detail::SYCLConfig<detail::SYCL_RT_WARNING_LEVEL>::get() >= 2)
+        std::clog << message << std::endl;
+
+      // If it is a warning do not throw code
+      if (pi_result == PI_SUCCESS)
+        return;
+    }
     __SYCL_CHECK_OCL_CODE_THROW(pi_result, Exception);
   }
 
   /// \throw SYCL 2020 exception(errc) if pi_result is not PI_SUCCESS
   template <sycl::errc errc> void checkPiResult(RT::PiResult pi_result) const {
+    if (pi_result == PI_PLUGIN_SPECIFIC_ERROR) {
+      char *message = nullptr;
+      pi_result = call_nocheck<PiApiKind::piPluginGetLastError>(&message);
+
+      // If the warning level is greater then 2 emit the message
+      if (detail::SYCLConfig<detail::SYCL_RT_WARNING_LEVEL>::get() >= 2)
+        std::clog << message << std::endl;
+
+      // If it is a warning do not throw code
+      if (pi_result == PI_SUCCESS)
+        return;
+    }
     __SYCL_CHECK_CODE_THROW_VIA_ERRC(pi_result, errc);
   }
 

sycl/test/abi/pi_level_zero_symbol_check.dump

@@ -78,6 +78,7 @@ piSamplerCreate
 piSamplerGetInfo
 piSamplerRelease
 piSamplerRetain
+piPluginGetLastError
 piTearDown
 piclProgramCreateWithSource
 piextContextCreateWithNativeHandle

sycl/test/abi/pi_opencl_symbol_check.dump

@@ -26,6 +26,7 @@ piProgramCreateWithBinary
 piProgramLink
 piQueueCreate
 piSamplerCreate
+piPluginGetLastError
 piTearDown
 piclProgramCreateWithSource
 piextContextCreateWithNativeHandle

sycl/tools/sycl-trace/pi_trace_collector.cpp

@@ -111,6 +111,8 @@ static std::string getResult(pi_result Res) {
     return "PI_COMMAND_EXECUTION_FAILURE";
   case PI_FUNCTION_ADDRESS_IS_NOT_AVAILABLE:
     return "PI_FUNCTION_ADDRESS_IS_NOT_AVAILABLE";
+  case PI_PLUGIN_SPECIFIC_ERROR:
+    return "PI_PLUGIN_SPECIFIC_ERROR";
   case PI_ERROR_UNKNOWN:
     return "PI_ERROR_UNKNOWN";
   }