[SYCL][CUDA] Fix context creation property parsing

Align the property parsing of `piContextCreate` with the way OpenCL
properties are provided, i.e., a property list contains keys/IDs and
values, never just a key/ID.

Furthermore enable property lists that just contain a list terminating
`0`.

Introduce a PI type for context properties to untie from OpenCL.

Adapt the SYCL runtime and tests accordingly.

Signed-off-by: Bjoern Knafla <[email protected]>

Author: bjoernknafla

sycl/include/CL/sycl/backend/cuda.hpp

@@ -18,7 +18,9 @@ namespace cuda {
 
 // Mem Object info: Retrieve the raw CUDA pointer from a cl_mem
 #define PI_CUDA_RAW_POINTER (0xFF01)
-// Context creation: Use the primary context instead of a custom one
+// Context creation: Use a primary CUDA context instead of a custom one by
+//                   providing a property value of PI_TRUE for the following
+//                   property ID.
 #define PI_CONTEXT_PROPERTIES_CUDA_PRIMARY (0xFF02)
 
 // PI Command Queue using Default stream

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

@@ -232,6 +232,8 @@ typedef enum {
   PI_LOCAL_MEM_TYPE_GLOBAL = CL_GLOBAL
 } _pi_local_mem_type;
 
+typedef intptr_t pi_context_properties;
+
 // TODO: populate
 typedef enum {
   PI_CONTEXT_INFO_DEVICES = CL_CONTEXT_DEVICES,
@@ -627,12 +629,12 @@ pi_result piextGetDeviceFunctionPointer(pi_device device, pi_program program,
 //
 // Context
 //
-pi_result piContextCreate(
-    const cl_context_properties *properties, // TODO: untie from OpenCL
-    pi_uint32 num_devices, const pi_device *devices,
-    void (*pfn_notify)(const char *errinfo, const void *private_info, size_t cb,
-                       void *user_data),
-    void *user_data, pi_context *ret_context);
+pi_result piContextCreate(const pi_context_properties *properties,
+                          pi_uint32 num_devices, const pi_device *devices,
+                          void (*pfn_notify)(const char *errinfo,
+                                             const void *private_info,
+                                             size_t cb, void *user_data),
+                          void *user_data, pi_context *ret_context);
 
 pi_result piContextGetInfo(pi_context context, pi_context_info param_name,
                            size_t param_value_size, void *param_value,

sycl/plugins/cuda/pi_cuda.cpp

@@ -1124,12 +1124,30 @@ pi_result cuda_piDeviceGetInfo(pi_device device, pi_device_info param_name,
 }
 
 /* Context APIs */
-pi_result cuda_piContextCreate(const cl_context_properties *properties,
-                                pi_uint32 num_devices, const pi_device *devices,
-                                void (*pfn_notify)(const char *errinfo,
-                                                   const void *private_info,
-                                                   size_t cb, void *user_data),
-                                void *user_data, pi_context *retcontext) {
+
+/// Create a PI CUDA context.
+///
+/// By default creates a scoped context and keeps the last active CUDA context
+/// on top of the CUDA context stack.
+/// With the PI_CONTEXT_PROPERTIES_CUDA_PRIMARY key/id and a value of PI_TRUE
+/// creates a primary CUDA context and activates it on the CUDA context stack.
+///
+/// @param[in] properties 0 terminated array of key/id-value combinations. Can
+/// be nullptr. Only accepts property key/id PI_CONTEXT_PROPERTIES_CUDA_PRIMARY
+/// with a pi_bool value.
+/// @param[in] num_devices Number of devices to create the context for.
+/// @param[in] devices Devices to create the context for.
+/// @param[in] pfn_notify Callback, currently unused.
+/// @param[in] user_data User data for callback.
+/// @param[out] retcontext Set to created context on success.
+///
+/// @return PI_SUCCESS on success, otherwise an error return code.
+pi_result cuda_piContextCreate(const pi_context_properties *properties,
+                               pi_uint32 num_devices, const pi_device *devices,
+                               void (*pfn_notify)(const char *errinfo,
+                                                  const void *private_info,
+                                                  size_t cb, void *user_data),
+                               void *user_data, pi_context *retcontext) {
 
   assert(devices != nullptr);
   // TODO: How to implement context callback?
@@ -1141,31 +1159,51 @@ pi_result cuda_piContextCreate(const cl_context_properties *properties,
   assert(retcontext != nullptr);
   pi_result errcode_ret = PI_SUCCESS;
 
+  // Parse properties.
+  bool property_cuda_primary = false;
+  while (properties && (0 != *properties)) {
+    // Consume property ID.
+    pi_context_properties id = *properties;
+    ++properties;
+    // Consume property value.
+    pi_context_properties value = *properties;
+    ++properties;
+    switch (id) {
+    case PI_CONTEXT_PROPERTIES_CUDA_PRIMARY:
+      assert(value == PI_FALSE || value == PI_TRUE);
+      property_cuda_primary = static_cast<bool>(value);
+      break;
+    default:
+      // Unknown property.
+      assert(!"Unknown piContextCreate property in property list");
+      return PI_INVALID_VALUE;
+    }
+  }
+
   std::unique_ptr<_pi_context> piContextPtr{nullptr};
   try {
-    if (properties && *properties != PI_CONTEXT_PROPERTIES_CUDA_PRIMARY) {
-      throw pi_result(CL_INVALID_VALUE);
-    } else if (!properties) {
+    if (property_cuda_primary) {
+      // Use the CUDA primary context and assume that we want to use it
+      // immediately as we want to forge context switches.
+      CUcontext Ctxt;
+      errcode_ret = PI_CHECK_ERROR(
+          cuDevicePrimaryCtxRetain(&Ctxt, devices[0]->cuDevice_));
+      piContextPtr = std::unique_ptr<_pi_context>(
+          new _pi_context{_pi_context::kind::primary, Ctxt, *devices});
+      errcode_ret = PI_CHECK_ERROR(cuCtxPushCurrent(Ctxt));
+    } else {
+      // Create a scoped context.
       CUcontext newContext, current;
       PI_CHECK_ERROR(cuCtxGetCurrent(&current));
-      errcode_ret = PI_CHECK_ERROR(cuCtxCreate(&newContext, CU_CTX_MAP_HOST,
-                                            (*devices)->cuDevice_));
+      errcode_ret = PI_CHECK_ERROR(
+          cuCtxCreate(&newContext, CU_CTX_MAP_HOST, devices[0]->cuDevice_));
       piContextPtr = std::unique_ptr<_pi_context>(new _pi_context{
           _pi_context::kind::user_defined, newContext, *devices});
+      // For scoped contexts keep the last active CUDA one on top of the stack
+      // as `cuCtxCreate` replaces it implicitly otherwise.
       if (current != nullptr) {
-        // If there was an existing context on the thread we recover it
         PI_CHECK_ERROR(cuCtxSetCurrent(current));
       }
-    } else if (properties 
-                  && *properties == PI_CONTEXT_PROPERTIES_CUDA_PRIMARY) {
-      CUcontext Ctxt;
-      errcode_ret = PI_CHECK_ERROR(cuDevicePrimaryCtxRetain(
-                                     &Ctxt, (*devices)->cuDevice_));
-      piContextPtr = std::unique_ptr<_pi_context>(
-          new _pi_context{_pi_context::kind::primary, Ctxt, *devices});
-      errcode_ret = PI_CHECK_ERROR(cuCtxPushCurrent(Ctxt));
-    } else {
-      throw pi_result(CL_INVALID_VALUE);
     }
 
     *retcontext = piContextPtr.release();

sycl/plugins/opencl/pi_opencl.cpp

@@ -451,12 +451,12 @@ pi_result OCL(piextGetDeviceFunctionPointer)(pi_device device,
                                   function_pointer_ret));
 }
 
-pi_result OCL(piContextCreate)(
-    const cl_context_properties *properties, // TODO: untie from OpenCL
-    pi_uint32 num_devices, const pi_device *devices,
-    void (*pfn_notify)(const char *errinfo, const void *private_info, size_t cb,
-                       void *user_data1),
-    void *user_data, pi_context *retcontext) {
+pi_result OCL(piContextCreate)(const pi_context_properties *properties,
+                               pi_uint32 num_devices, const pi_device *devices,
+                               void (*pfn_notify)(const char *errinfo,
+                                                  const void *private_info,
+                                                  size_t cb, void *user_data1),
+                               void *user_data, pi_context *retcontext) {
   pi_result ret = PI_INVALID_OPERATION;
   *retcontext = cast<pi_context>(
       clCreateContext(properties, cast<cl_uint>(num_devices),

sycl/source/detail/context_impl.cpp

@@ -44,9 +44,8 @@ context_impl::context_impl(const vector_class<cl::sycl::device> Devices,
 
   if (MPlatform->is_cuda()) {
 #if USE_PI_CUDA
-    const cl_context_properties props[] = {
-        PI_CONTEXT_PROPERTIES_CUDA_PRIMARY,
-        0};
+    const pi_context_properties props[] = {PI_CONTEXT_PROPERTIES_CUDA_PRIMARY,
+                                           UseCUDAPrimaryContext, 0};
 
     getPlugin().call<PiApiKind::piContextCreate>(props, DeviceIds.size(), 
 	  	  DeviceIds.data(), nullptr, nullptr, &MContext);

sycl/unittests/pi/cuda/test_base_objects.cpp

@@ -73,7 +73,7 @@ TEST_F(DISABLED_CudaBaseObjectsTest, piContextCreate) {
   ASSERT_EQ(cuErr, CUDA_SUCCESS);
 }
 
-TEST_F(DISABLED_CudaBaseObjectsTest, piContextCreatePrimary) {
+TEST_F(DISABLED_CudaBaseObjectsTest, piContextCreatePrimaryTrue) {
   pi_uint32 numPlatforms = 0;
   pi_platform platform;
   pi_device device;
@@ -91,11 +91,12 @@ TEST_F(DISABLED_CudaBaseObjectsTest, piContextCreatePrimary) {
   ASSERT_EQ((Plugins[0].call_nocheck<detail::PiApiKind::piDevicesGet>(
                 platform, PI_DEVICE_TYPE_GPU, 1, &device, nullptr)),
             PI_SUCCESS);
-  cl_context_properties properties = PI_CONTEXT_PROPERTIES_CUDA_PRIMARY;
+  pi_context_properties properties[] = {PI_CONTEXT_PROPERTIES_CUDA_PRIMARY,
+                                        PI_TRUE, 0};
 
   pi_context ctxt;
   ASSERT_EQ((Plugins[0].call_nocheck<detail::PiApiKind::piContextCreate>(
-                &properties, 1, &device, nullptr, nullptr, &ctxt)),
+                properties, 1, &device, nullptr, nullptr, &ctxt)),
             PI_SUCCESS);
   EXPECT_NE(ctxt, nullptr);
   EXPECT_EQ(ctxt->get_device(), device);
@@ -118,6 +119,52 @@ TEST_F(DISABLED_CudaBaseObjectsTest, piContextCreatePrimary) {
       PI_SUCCESS);
 }
 
+TEST_F(DISABLED_CudaBaseObjectsTest, piContextCreatePrimaryFalse) {
+  pi_uint32 numPlatforms = 0;
+  pi_platform platform;
+  pi_device device;
+
+  ASSERT_EQ((Plugins[0].call_nocheck<detail::PiApiKind::piPlatformsGet>(
+                0, nullptr, &numPlatforms)),
+            PI_SUCCESS)
+      << "piPlatformsGet failed.\n";
+
+  ASSERT_EQ((Plugins[0].call_nocheck<detail::PiApiKind::piPlatformsGet>(
+                numPlatforms, &platform, nullptr)),
+            PI_SUCCESS)
+      << "piPlatformsGet failed.\n";
+
+  ASSERT_EQ((Plugins[0].call_nocheck<detail::PiApiKind::piDevicesGet>(
+                platform, PI_DEVICE_TYPE_GPU, 1, &device, nullptr)),
+            PI_SUCCESS);
+  pi_context_properties properties[] = {PI_CONTEXT_PROPERTIES_CUDA_PRIMARY,
+                                        PI_FALSE, 0};
+
+  pi_context ctxt;
+  ASSERT_EQ((Plugins[0].call_nocheck<detail::PiApiKind::piContextCreate>(
+                properties, 1, &device, nullptr, nullptr, &ctxt)),
+            PI_SUCCESS);
+  EXPECT_NE(ctxt, nullptr);
+  EXPECT_EQ(ctxt->get_device(), device);
+  EXPECT_FALSE(ctxt->is_primary());
+
+  // Retrieve the cuCtxt to check information is correct
+  CUcontext cudaContext = ctxt->get();
+  unsigned int version = 0;
+  CUresult cuErr = cuCtxGetApiVersion(cudaContext, &version);
+  ASSERT_EQ(cuErr, CUDA_SUCCESS);
+  EXPECT_EQ(version, LATEST_KNOWN_CUDA_DRIVER_API_VERSION);
+
+  // Current context in the stack?
+  CUcontext current;
+  cuErr = cuCtxGetCurrent(&current);
+  ASSERT_EQ(cuErr, CUDA_SUCCESS);
+  ASSERT_EQ(current, cudaContext);
+  ASSERT_EQ(
+      (Plugins[0].call_nocheck<detail::PiApiKind::piContextRelease>(ctxt)),
+      PI_SUCCESS);
+}
+
 TEST_F(DISABLED_CudaBaseObjectsTest, piContextCreateChildThread) {
   pi_uint32 numPlatforms = 0;
   pi_platform platform;