Repository navigation
NVRTC-compiled generated kernels missing extern "C" — cuModuleGetFunction can't find C++-mangled symbols - #120
Closed
AlekSimpson wants to merge 7 commits into
Closed
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-116_FixNvrtcKernelsMissingExternC
The codegen pipeline's runtime-compiled CUDA kernels (population _tick/_deliver functions from gpu_source.cpp, plus the four fixed master-kernel functions in master_kernel.cpp) were emitted as plain __global__ void, which nvcc/NVRTC C++-mangle by default. backend.cpp's compile_kernel() -> cuModuleGetFunction looks the kernel up by its plain, unmangled name, so every one of these lookups failed with "named symbol not found". Wrap each CUDA-path __global__ kernel signature in extern "C" (MSL template untouched -- Metal's newLibrary lookup has no equivalent mangling concern). 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
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.
Wraps every CUDA-path runtime-compiled kernel signature in
extern "C"(1 generic site insrc/nml/gpu_source.cpp'srender_signature_and_prologue, covering the generated_tick/_deliverpopulation kernels, plus the 4 fixed master kernels insrc/nml/master_kernel.cpp), socuModuleGetFunctioncan resolve their plain unmangled names instead of nvcc/NVRTC's defaultC++ name-mangled ones. MSL emission and
__device__helpers untouched.Eliminates all 41
cuModuleGetFunction failed: named symbol not foundfailures with zeroregressions (322 passed, unchanged). Note: this ticket's bug and #119's (NVRTC missing include
path for
vector_types.h) were entangled —AssembledModelthrows on the first kernel-compilefailure, so #119's bug was previously masked behind this one. With this fix alone, those ~41 tests
now fail via #119 instead (compiling further before hitting the next real bug) — expected, and
#119 has been updated to reflect its true ~42-test blast radius. The raw pass count won't move above
322 until #119 also lands.
Branched off #114 (merged with #108/#109), so this diff shows all four prior fixes too until they
merge first.
Acceptance criteria
extern "C".cuModuleGetFunction failederror fully eliminated (0 occurrences, down from 41 tests).Reviewer summary: PASS. All 5 emission sites confirmed correct and complete (grepped the whole
tree for any missed
__global__ voidstring-emission site — none found;kernels.cu'sbuild-time-compiled kernels correctly need no
extern "C"since they're launched vialaunch_*wrappers, not looked up by name). Spot-checked 4 representative tests and confirmed their failure
mode shifted from "named symbol not found" to the expected #119
vector_types.herror, not somenew third failure kind. No math/semantics change, no Metal regression.
Closes #116