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

Kevin Sala Penades via llvm-commits llvm-commits at lists.llvm.org
Wed Mar 4 15:06:53 PST 2026


================
@@ -1495,9 +1502,31 @@ 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);
+  // Validate cooperative launch requirements if needed
+  if (KernelArgs.Flags.Cooperative) {
+    CUDADeviceTy &CUDADevice = static_cast<CUDADeviceTy &>(GenericDevice);
+
+    uint32_t SupportsCooperative = 0;
+    if (auto Err = CUDADevice.getDeviceAttr(
----------------
kevinsala wrote:

Can you store this information in the device object (e.g., in the generic device class)? This is the critical path for cooperative kernel launches.

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


More information about the llvm-commits mailing list