[llvm] [offload] Add properties parameter to olLaunchKernel (PR #184343)

[Łukasz Plewa via llvm-commits]

================
@@ -1495,9 +1514,44 @@ Error CUDAKernelTy::launchImpl(GenericDeviceTy &GenericDevice,
     MaxDynCGroupMemLimit = MaxDynCGroupMem;
   }
 
-  CUresult Res = cuLaunchKernel(Func, NumBlocks[0], NumBlocks[1], NumBlocks[2],
-                                NumThreads[0], NumThreads[1], NumThreads[2],
-                                MaxDynCGroupMem, Stream, nullptr, Config);
+  CUresult Res;
+  if (KernelArgs.Flags.Cooperative) {
+    CUDADeviceTy &CUDADevice = static_cast<CUDADeviceTy &>(GenericDevice);
+
+    uint32_t SupportsCooperative = 0;
+    if (auto Err = CUDADevice.getDeviceAttr(
+            CU_DEVICE_ATTRIBUTE_COOPERATIVE_LAUNCH, SupportsCooperative))
+      return Err;
+
+    if (!SupportsCooperative) {
+      return Plugin::error(ErrorCode::UNSUPPORTED,
+                           "Device does not support cooperative launch");
+    }
+
+    CUlaunchAttribute CoopAttr;
+    CoopAttr.id = CU_LAUNCH_ATTRIBUTE_COOPERATIVE;
+    CoopAttr.value.cooperative = 1;
+
+    CUlaunchConfig LaunchConfig;
----------------
lplewa wrote:

done

https://github.com/llvm/llvm-project/pull/184343