[SYCL][CUDA] Improvements to CUDA device selection

sycl/include/CL/sycl/device_selector.hpp

@@ -19,6 +19,10 @@ namespace sycl {
 class device;
 
 class __SYCL_EXPORT device_selector {
+protected:
+  // SYCL 1.2.1 defines a negative score to reject a device from selection
+  static constexpr int REJECT_DEVICE_SCORE = -1;
+
 public:
   virtual ~device_selector() = default;
 

sycl/plugins/cuda/pi_cuda.cpp

@@ -684,7 +684,7 @@ pi_result cuda_piPlatformGetInfo(pi_platform platform,
   switch (param_name) {
   case PI_PLATFORM_INFO_NAME:
     return getInfo(param_value_size, param_value, param_value_size_ret,
-                   "NVIDIA CUDA");
+                   "NVIDIA CUDA BACKEND");
   case PI_PLATFORM_INFO_VENDOR:
     return getInfo(param_value_size, param_value, param_value_size_ret,
                    "NVIDIA Corporation");

sycl/source/detail/context_impl.cpp

@@ -41,7 +41,8 @@ context_impl::context_impl(const vector_class<cl::sycl::device> Devices,
     DeviceIds.push_back(getSyclObjImpl(D)->getHandleRef());
   }
 
-  if (MPlatform->is_cuda()) {
+  const auto Backend = getPlugin().getBackend();
+  if (Backend == backend::cuda) {
 #if USE_PI_CUDA
     const pi_context_properties props[] = {PI_CONTEXT_PROPERTIES_CUDA_PRIMARY,
                                            UseCUDAPrimaryContext, 0};

sycl/source/detail/platform_impl.cpp

@@ -20,6 +20,33 @@ __SYCL_INLINE_NAMESPACE(cl) {
 namespace sycl {
 namespace detail {
 
+static bool IsBannedPlatform(platform Platform) {
+  // The NVIDIA OpenCL platform is currently not compatible with DPC++
+  // since it is only 1.2 but gets selected by default in many systems
+  // There is also no support on the PTX backend for OpenCL consumption,
+  // and there have been some internal reports.
+  // To avoid problems on default users and deployment of DPC++ on platforms
+  // where CUDA is available, the OpenCL support is disabled.
+  //
+  auto IsNVIDIAOpenCL = [](platform Platform) {
+    if (Platform.is_host())
+      return false;
+
+    const bool HasCUDA = Platform.get_info<info::platform::name>().find(
+                             "NVIDIA CUDA") != std::string::npos;
+    const auto Backend =
+        detail::getSyclObjImpl(Platform)->getPlugin().getBackend();
+    const bool IsCUDAOCL = (HasCUDA && Backend == backend::opencl);
+    if (detail::pi::trace(detail::pi::TraceLevel::PI_TRACE_ALL) && IsCUDAOCL) {
+      std::cout << "SYCL_PI_TRACE[all]: "
+                << "NVIDIA CUDA OpenCL platform found but is not compatible."
+                << std::endl;
+    }
+    return IsCUDAOCL;
+  };
+  return IsNVIDIAOpenCL(Platform);
+}
+
 vector_class<platform> platform_impl::get_platforms() {
   vector_class<platform> Platforms;
   vector_class<plugin> Plugins = RT::initialize();
@@ -39,7 +66,8 @@ vector_class<platform> platform_impl::get_platforms() {
         platform Platform = detail::createSyclObjFromImpl<platform>(
             std::make_shared<platform_impl>(PiPlatform, Plugins[i]));
         // Skip platforms which do not contain requested device types
-        if (!Platform.get_devices(ForcedType).empty())
+        if (!Platform.get_devices(ForcedType).empty() &&
+            !IsBannedPlatform(Platform))
           Platforms.push_back(Platform);
       }
     }

sycl/source/detail/platform_impl.hpp

@@ -73,14 +73,6 @@ class platform_impl {
   /// \return true if this SYCL platform is a host platform.
   bool is_host() const { return MHostPlatform; };
 
-  bool is_cuda() const {
-    const string_class CUDA_PLATFORM_STRING = "NVIDIA CUDA";
-    const string_class PlatformName =
-        get_platform_info<string_class, info::platform::name>::get(MPlatform,
-                                                                   getPlugin());
-    return PlatformName == CUDA_PLATFORM_STRING;
-  }
-
   /// \return an instance of OpenCL cl_platform_id.
   cl_platform_id get() const {
     if (is_host())

sycl/source/detail/program_manager/program_manager.cpp

@@ -86,29 +86,10 @@ static RT::PiProgram createBinaryProgram(const ContextImplPtr Context,
 
   RT::PiProgram Program;
 
-  bool IsCUDA = false;
-
   // TODO: Implement `piProgramCreateWithBinary` to not require extra logic for
   //       the CUDA backend.
-#if USE_PI_CUDA
-  // All devices in a context are from the same platform.
-  RT::PiDevice Device = getFirstDevice(Context);
-  RT::PiPlatform Platform = nullptr;
-  Plugin.call<PiApiKind::piDeviceGetInfo>(Device, PI_DEVICE_INFO_PLATFORM, sizeof(Platform),
-                           &Platform, nullptr);
-  size_t PlatformNameSize = 0u;
-  Plugin.call<PiApiKind::piPlatformGetInfo>(Platform, PI_PLATFORM_INFO_NAME, 0u, nullptr,
-                             &PlatformNameSize);
-  std::vector<char> PlatformName(PlatformNameSize, '\0');
-  Plugin.call<PiApiKind::piPlatformGetInfo>(Platform, PI_PLATFORM_INFO_NAME,
-                             PlatformName.size(), PlatformName.data(), nullptr);
-  if (PlatformNameSize > 0u &&
-      std::strncmp(PlatformName.data(), "NVIDIA CUDA", PlatformNameSize) == 0) {
-    IsCUDA = true;
-  }
-#endif // USE_PI_CUDA
-
-  if (IsCUDA) {
+  const auto Backend = Context->getPlugin().getBackend();
+  if (Backend == backend::cuda) {
     // TODO: Reemplace CreateWithSource with CreateWithBinary in CUDA backend
     const char *SignedData = reinterpret_cast<const char *>(Data);
     Plugin.call<PiApiKind::piclProgramCreateWithSource>(Context->getHandleRef(), 1 /*one binary*/, &SignedData,
@@ -259,6 +240,13 @@ RetT *getOrBuild(KernelProgramCache &KPCache, KeyT &&CacheKey,
 
 static bool isDeviceBinaryTypeSupported(const context &C,
                                         RT::PiDeviceBinaryType Format) {
+  const backend ContextBackend =
+      detail::getSyclObjImpl(C)->getPlugin().getBackend();
+
+  // The CUDA backend cannot use SPIRV
+  if (ContextBackend == backend::cuda && Format == PI_DEVICE_BINARY_TYPE_SPIRV)
+    return false;
+
   // All formats except PI_DEVICE_BINARY_TYPE_SPIRV are supported.
   if (Format != PI_DEVICE_BINARY_TYPE_SPIRV)
     return true;
@@ -272,8 +260,7 @@ static bool isDeviceBinaryTypeSupported(const context &C,
   }
 
   // OpenCL 2.1 and greater require clCreateProgramWithIL
-  backend CBackend = (detail::getSyclObjImpl(C)->getPlugin()).getBackend();
-  if ((CBackend == backend::opencl) &&
+  if ((ContextBackend == backend::opencl) &&
       C.get_platform().get_info<info::platform::version>() >= "2.1")
     return true;
 
@@ -337,7 +324,7 @@ RT::PiProgram ProgramManager::createPIProgram(const RTDeviceBinaryImage &Img,
 
   if (!isDeviceBinaryTypeSupported(Context, Format))
     throw feature_not_supported(
-        "Online compilation is not supported in this context",
+        "SPIR-V online compilation is not supported in this context",
         PI_INVALID_OPERATION);
 
   // Load the image

sycl/source/device_selector.cpp

@@ -30,28 +30,36 @@ static bool isDeviceOfPreferredSyclBe(const device &Device) {
 
 device device_selector::select_device() const {
   vector_class<device> devices = device::get_devices();
-  int score = -1;
+  int score = REJECT_DEVICE_SCORE;
   const device *res = nullptr;
+
   for (const auto &dev : devices) {
     int dev_score = (*this)(dev);
+
     if (detail::pi::trace(detail::pi::TraceLevel::PI_TRACE_ALL)) {
       string_class PlatformVersion = dev.get_info<info::device::platform>()
                                          .get_info<info::platform::version>();
       string_class DeviceName = dev.get_info<info::device::name>();
       std::cout << "SYCL_PI_TRACE[all]: "
-                << "select_device(): -> score = " << score << std::endl
+                << "select_device(): -> score = " << score
+                << ((score == REJECT_DEVICE_SCORE) ? "(REJECTED)" : " ")
+                << std::endl
                 << "SYCL_PI_TRACE[all]: "
                 << "  platform: " << PlatformVersion << std::endl
                 << "SYCL_PI_TRACE[all]: "
                 << "  device: " << DeviceName << std::endl;
     }
 
+    // Device is discarded if is marked with REJECT_DEVICE_SCORE
+    if (dev_score == REJECT_DEVICE_SCORE)
+      continue;
+
     // SYCL spec says: "If more than one device receives the high score then
     // one of those tied devices will be returned, but which of the devices
     // from the tied set is to be returned is not defined". Here we give a
     // preference to the device of the preferred BE.
     //
-    if (score < dev_score ||
+    if ((score < dev_score) ||
         (score == dev_score && isDeviceOfPreferredSyclBe(dev))) {
       res = &dev;
       score = dev_score;
@@ -79,7 +87,7 @@ device device_selector::select_device() const {
 
 int default_selector::operator()(const device &dev) const {
 
-  int Score = -1;
+  int Score = REJECT_DEVICE_SCORE;
 
   // Give preference to device of SYCL BE.
   if (isDeviceOfPreferredSyclBe(dev))
@@ -102,7 +110,8 @@ int default_selector::operator()(const device &dev) const {
 }
 
 int gpu_selector::operator()(const device &dev) const {
-  int Score = -1;
+  int Score = REJECT_DEVICE_SCORE;
+
   if (dev.is_gpu()) {
     Score = 1000;
     // Give preference to device of SYCL BE.
@@ -113,7 +122,7 @@ int gpu_selector::operator()(const device &dev) const {
 }
 
 int cpu_selector::operator()(const device &dev) const {
-  int Score = -1;
+  int Score = REJECT_DEVICE_SCORE;
   if (dev.is_cpu()) {
     Score = 1000;
     // Give preference to device of SYCL BE.
@@ -124,7 +133,7 @@ int cpu_selector::operator()(const device &dev) const {
 }
 
 int accelerator_selector::operator()(const device &dev) const {
-  int Score = -1;
+  int Score = REJECT_DEVICE_SCORE;
   if (dev.is_accelerator()) {
     Score = 1000;
     // Give preference to device of SYCL BE.
@@ -135,7 +144,7 @@ int accelerator_selector::operator()(const device &dev) const {
 }
 
 int host_selector::operator()(const device &dev) const {
-  int Score = -1;
+  int Score = REJECT_DEVICE_SCORE;
   if (dev.is_host()) {
     Score = 1000;
     // Give preference to device of SYCL BE.

sycl/tools/get_device_count_by_type.cpp

@@ -21,6 +21,7 @@
 
 #include <algorithm>
 #include <cstdlib>
+#include <cstring>
 #include <iostream>
 #include <sstream>
 #include <string>
@@ -88,6 +89,17 @@ static bool queryOpenCL(cl_device_type deviceType, cl_uint &deviceCount,
   }
 
   for (cl_uint i = 0; i < platformCount; i++) {
+    const size_t MAX_PLATFORM_VENDOR = 100u;
+    char info[MAX_PLATFORM_VENDOR];
+    // get platform attribute value
+    clGetPlatformInfo(platforms[i], CL_PLATFORM_VENDOR, MAX_PLATFORM_VENDOR,
+                      info, NULL);
+    const auto IsNVIDIAOpenCL = strstr(info, "NVIDIA") != NULL;
+    if (IsNVIDIAOpenCL) {
+      // Ignore NVIDIA OpenCL platform for testing
+      continue;
+    }
+
     cl_uint deviceCountPart = 0;
     iRet =
         clGetDeviceIDs(platforms[i], deviceType, 0, nullptr, &deviceCountPart);