Skip to content

Commit 6ee27a4

Browse files
committed
Rename candle's moe_gemm_wmma -> candle_moe_gemm_wmma to avoid symbol collision with mistralrs-core
Both candle-kernels and mistralrs-core define an extern "C" void moe_gemm_wmma kernel with the same signature but incompatible implementations. When both crates are linked into the same binary (e.g. spiced with cuda + models features), ld.lld rejects the duplicate symbol. Prefix the candle side with 'candle_' so both can coexist; the mistralrs-core copy keeps the original name and call sites.
1 parent c87b9bc commit 6ee27a4

3 files changed

Lines changed: 3 additions & 3 deletions

File tree

candle-kernels/src/ffi.rs

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -2,7 +2,7 @@ use core::ffi::c_void;
22
#[allow(dead_code)]
33
extern "C" {
44
// for unquntized models
5-
pub fn moe_gemm_wmma(
5+
pub fn candle_moe_gemm_wmma(
66
input: *const c_void, // device pointer [size_m, size_k]
77
weights: *const c_void, // device pointer [num_experts, size_n, size_k]
88
sorted_token_ids: *const i32, // device pointer [size_m]

candle-kernels/src/moe/moe_wmma.cu

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -231,7 +231,7 @@ __global__ void moe_gemm_grouped_kernel(
231231
size_m, size_n, size_k \
232232
);\
233233

234-
extern "C" void moe_gemm_wmma(
234+
extern "C" void candle_moe_gemm_wmma(
235235
const void* input, // [size_m, size_k]
236236
const void* weights, // [num_experts, size_n, size_k]
237237
const int32_t* sorted_token_ids, // [size_m] (Device)

candle-nn/src/moe.rs

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -94,7 +94,7 @@ pub fn moe_gemm(
9494
use core::ffi::c_void;
9595

9696
unsafe {
97-
ffi::moe_gemm_wmma(
97+
ffi::candle_moe_gemm_wmma(
9898
input.device_ptr(input.stream()).0 as *const c_void, // [size_m, size_k]
9999
weights.device_ptr(weights.stream()).0 as *const c_void, // [num_experts, size_n, size_k]
100100
sorted_token_ids.device_ptr(sorted_token_ids.stream()).0 as *const i32,

0 commit comments

Comments
 (0)