NVRTC compile error: could not open source file vector_types.h in TsodyksMarkram synapse test - #122
Open
AlekSimpson wants to merge 7 commits into
Open
AlekSimpson wants to merge 7 commits into
AlekSimpson wants to merge 7 commits into
Conversation
…gcc CUDA build types.h defines Vector<T> and UnorderedMap<K,V> aliases but relied on transitive includes from <string>/<any> that only hold on macOS clang/libc++, not Linux gcc/libstdc++. Add the two missing includes directly. Co-Authored-By: Claude Sonnet 5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_015JWoo8NdTRVcnAzxdiXJbJ
… build ambiguity CUDA's <vector_types.h> (pulled in via cuda_runtime.h) also declares a global ::float4, which collided with spikecorec::float4 wherever this file's bare `float4` was looked up outside the spikecorec namespace (the anonymous- namespace refit helpers, plus every other unqualified use for consistency). Pure name-resolution fix, no behavior change.
… header (ticket #108) backend.h only included <cuda.h> (the CUDA Driver API header) but called Runtime API functions/constants (cudaMemPrefetchAsync, cudaCpuDeviceId, cudaMemAdvise, cudaMemAdviseSetReadMostly) declared in <cuda_runtime.h>. Add <cuda_runtime.h> alongside <cuda.h>, which is still required for the Driver API types CUfunction/CUmodule used by KernelHandle.
…109) backend.cpp's CUDA branches tried to launch raw __global__ kernels directly with nvcc-only triple-chevron syntax from a file compiled by host g++, and referenced an undeclared `stream` identifier at every site. Implements all 10 launch_* functions declared in kernels.cuh (plus launch_step_no_active_ optimization, needed for gpu_step's active_set_optimization_enabled==false path, which kernels.cuh had no wrapper for) inside kernels.cu, each doing the config computation + <<<...>>> launch of its corresponding kernel, and rewires backend.cpp's 10 gpu_* CUDA branches to call them with stream=nullptr, keeping the existing post-launch synchronize_gpu_work() (required due to concurrentManagedAccess=0 on this Jetson). Removes the duplicate `s64 total_pairs` redeclarations in gpu_neighbor_weights/gpu_k2tree_get_neighbors_batch and the local spikecorec::LaunchConfig variables that collided with spikecorec::cuda::LaunchConfig. Also adds the missing `coefficients` parameter to kernels.cuh's launch_neighbor_weights declaration (already present on gpu_neighbor_weights/the kernel itself, just missing from the header) and const_casts gpu_step's const float4* U/V to the non-const pointers the mutating step kernels require. Co-Authored-By: Claude Sonnet 5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_015JWoo8NdTRVcnAzxdiXJbJ
….cpp to fix CUDA build ambiguity Same root cause and fix pattern as #107 (weight_matrix.cpp): bare float4 collides with CUDA's built-in ::float4 from <vector_types.h> wherever both types are visible via this file's `using namespace spikecorec;`. Co-Authored-By: Claude Sonnet 5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_015JWoo8NdTRVcnAzxdiXJbJ
…pers' into SC-119_FixNvrtcIncludePathForVectorTypes
…clude CUDA headers nvrtcCompileProgram was invoked with zero compile options, so NVRTC (unlike nvcc) had no include search path at all -- any generated CUDA source that #includes a CUDA header (e.g. gpu_source.cpp's `#include <vector_types.h>`, emitted whenever a codegen path needs float4 to walk the shared-basis U/V edge storage) failed with "could not open source file". Now passes -I<cuda>/include, baked in at compile time from the same CUDA_PATH the Makefile already uses for the host compiler, fixing this generally rather than special-casing one kernel. Co-Authored-By: Claude Sonnet 5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_015JWoo8NdTRVcnAzxdiXJbJ
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
1 participant
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
compile_kernelinsrc/core/backend.cppinvokednvrtcCompileProgramwith zero options; NVRTC(unlike nvcc) has no default include search path, so any generated CUDA source that
#includes aCUDA header — e.g.
<vector_types.h>, emitted bygpu_source.cpp:1411for any codegen path thatwalks shared-basis U/V edge storage via
float4(needs_edge_walk_anywhere) — failed with "couldnot open source file". This is a systemic gap, not specific to the one test it was first noticed
on: no generated kernel ever got a CUDA include path, other paths just didn't need one.
Fix passes
-I<cuda>/includeon everycompile_kernelcall, baked in at compile time via a newSPIKECOREC_CUDA_INCLUDE_DIRdefine derived from the Makefile's existingCUDA_PATH(with amatching in-source fallback default), mirroring the existing
SPIKECOREC_NML_STD_LIB_DIRpattern.This ticket turned out to be entangled with #116 (already merged as PR #120) —
AssembledModelthrows on the first kernel-compile failure, so before #116 landed, ~41 of #116's tests were also
silently blocked behind this bug once #116's
extern "C"issue was fixed. Verified: standalone,this fix eliminates all
vector_types.h/NVRTC-include errors with pass count holding at 322/59(the entangled test still fails via #116's error, expected since #116 isn't in this branch's
ancestry). Combined with #116, the suite goes to 363 passed / 18 failed, and
TsodyksMarkram.per_edge_state_evolves_correctly_across_repeated_presynaptic_spikespasses. Theremaining 18 failures are #117 (xcrun/MSL, macOS-only) and #118 (WeightMatrix numeric), untouched.
Branched off #114 (merged with #108/#109), so this diff shows those fixes too until they merge first.
Acceptance criteria
vector_types.h, and confirmed the gap is systemic.compile_kernelcalls get the include path), not narrowly.TsodyksMarkram...passes (confirmed once combined with NVRTC-compiled generated kernels missing extern "C" — cuModuleGetFunction can't find C++-mangled symbols #116).Reviewer summary: PASS. Independently reproduced both the standalone (322/59, zero
vector_types errors) and combined-with-#116 (363/18, TsodyksMarkram passing) results in a throwaway
scratch worktree, since removed. Code review: sound
CUDA_PATH-derived include path, sanefallback default, correct argument count into
nvrtcCompileProgram, exactly one call site in thecodebase (no missed spots), no scope creep.
Closes #119