[llvm] [Offload] Use flat array for cuLaunchKernel (PR #95116)

Joseph Huber via llvm-commits llvm-commits at lists.llvm.org
Wed Jun 12 04:39:46 PDT 2024


================
@@ -1285,11 +1289,16 @@ Error CUDAKernelTy::launchImpl(GenericDeviceTy &GenericDevice,
   uint32_t MaxDynCGroupMem =
       std::max(KernelArgs.DynCGroupMem, GenericDevice.getDynamicMemorySize());
 
-  CUresult Res =
-      cuLaunchKernel(Func, NumBlocks, /*gridDimY=*/1,
-                     /*gridDimZ=*/1, NumThreads,
-                     /*blockDimY=*/1, /*blockDimZ=*/1, MaxDynCGroupMem, Stream,
-                     (void **)Args, nullptr);
+  void *Config[] = {/* CU_LAUNCH_PARAM_BUFFER_POINTER */ (void *)0x01,
+                    LaunchParams.Data,
+                    /* CU_LAUNCH_PARAM_BUFFER_SIZE */ (void *)0x02,
+                    reinterpret_cast<void *>(&LaunchParams.Size),
+                    /* CU_LAUNCH_PARAM_END */ (void *)0x00};
----------------
jhuber6 wrote:

```suggestion
  void *Config[] = {/* CU_LAUNCH_PARAM_BUFFER_POINTER */ (void *)0x01,
                    LaunchParams.Data,
                    /* CU_LAUNCH_PARAM_BUFFER_SIZE */ (void *)0x02,
                    reinterpret_cast<void *>(&LaunchParams.Size),
                    /* CU_LAUNCH_PARAM_END */ (void *)0x00};
```
What's going on here? Shouldn't these just be added to `dynamic_cuda/cuda.h`?

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


More information about the llvm-commits mailing list