From a2ef8f08715843899313dc78caeb2ee093006667 Mon Sep 17 00:00:00 2001 From: jgrusewski Date: Mon, 16 Mar 2026 09:59:32 +0100 Subject: [PATCH] fix(cuda): guard DQN-specific q_forward_dueling_warp_shmem behind #if defined() Non-DQN kernels (curiosity, PPO, epsilon-greedy, backtest-PPO) include common_device_functions.cuh but don't define SHARED_H1/SHARED_H2/VALUE_H/ADV_H. The unguarded q_forward_dueling_warp_shmem() references these constants, causing nvcc compilation failures (undefined identifiers + cascading "user-defined literal operator not found" errors). Wraps the function in #if defined(SHARED_H1) && defined(SHARED_H2) && defined(VALUE_H) && defined(ADV_H) ... #endif so it's only compiled when the DQN architecture constants are injected. Co-Authored-By: Claude Opus 4.6 --- crates/ml/src/cuda_pipeline/common_device_functions.cuh | 5 +++++ 1 file changed, 5 insertions(+) diff --git a/crates/ml/src/cuda_pipeline/common_device_functions.cuh b/crates/ml/src/cuda_pipeline/common_device_functions.cuh index 42e99435c..f9d172007 100644 --- a/crates/ml/src/cuda_pipeline/common_device_functions.cuh +++ b/crates/ml/src/cuda_pipeline/common_device_functions.cuh @@ -1632,6 +1632,9 @@ __device__ __forceinline__ int simulate_fill_check( } while (0) #endif +/* Guard: only compile when DQN architecture constants are defined */ +#if defined(SHARED_H1) && defined(SHARED_H2) && defined(VALUE_H) && defined(ADV_H) + /** * Warp-cooperative dueling Q-network forward pass (clean, no noise). * @@ -1711,3 +1714,5 @@ __device__ void q_forward_dueling_warp_shmem( q_values[i] = value + adv[i] - adv_mean; } } + +#endif /* SHARED_H1 && SHARED_H2 && VALUE_H && ADV_H */