fix linker errors about multiply defined symbols in STF

NVIDIA · Oct 28, 2024 · 553adde · 553adde
1 parent 8f1691d
commit 553adde
Showing 1 changed file with 3 additions and 2 deletions.
diff --git a/cudax/include/cuda/experimental/__stf/places/exec/host/callback_queues.cuh b/cudax/include/cuda/experimental/__stf/places/exec/host/callback_queues.cuh
@@ -320,6 +320,7 @@ inline void cudagraph_callback_dispatcher(void* userData)
 }
 
 // There is likely a more efficient way in the current implementation of callbacks !
+template <int = 0> // template to make this `inline` without using `inline`, which nvcc dislikes
 __global__ void callback_completion_kernel(int* completion_flag)
 {
   // Loop until *completion_flag == 1
@@ -341,7 +342,7 @@ inline class cb* get_current_cb()
   return cudaCallbackStateCtx::instance().get_current_cb();
 }
 
-cudaError_t cudaCallbackSetStatus(int step, void* private_ptr)
+inline cudaError_t cudaCallbackSetStatus(int step, void* private_ptr)
 {
   class cb* current_cb = get_current_cb();
   assert(current_cb);
@@ -350,7 +351,7 @@ cudaError_t cudaCallbackSetStatus(int step, void* private_ptr)
   return cudaSuccess;
 }
 
-cudaError_t cudaCallbackGetStatus(int* step, void** private_ptr)
+inline cudaError_t cudaCallbackGetStatus(int* step, void** private_ptr)
 {
   class cb* current_cb = get_current_cb();
   assert(current_cb);