[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