Skip to content

Commit 8cd30ba

Browse files
lalaluneclaude
lalalune
and
claude
committed
build(cuda): Gemma-clean CUDA backend — take upstream ggml-cuda over the merged tree
The merged tree's CUDA backend now COMPILES clean (248/248 targets, 0 errors) by aligning ggml-cuda to upstream's refactored type system (block_q8_1_layout template, ggml_cuda_kernel_launch params, fattn type-exhaustive instances). Drops the legacy QJL/PolarQuant/TBQ3_TCQ CUDA KV kernels + their custom-type template instances (q1_0_g128, fattn-vec tbq3_0/tbq4_0): these are head_dim=128 and Gemma-irrelevant (Gemma uses stock q8_0 KV — M6). Preserving them on CUDA would need a separate reconciliation of our custom GGML_TYPE_* integration with upstream's type refactor; out of scope for the Gemma cutover. Verified: builds with CMAKE_CUDA_ARCHITECTURES=90-virtual. Runtime on the local RTX 5080 (sm_120) still needs CUDA 12.8+/13 — the CUDA 12.0 runtime cannot enumerate Blackwell (confirmed: ggml_cuda_init 'no CUDA-capable device'); not a code issue. Metal/Vulkan kernels unchanged. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
1 parent d0f701b commit 8cd30ba

27 files changed

Lines changed: 317 additions & 2226 deletions

‎ggml/src/ggml-cuda/CMakeLists.txt‎

Lines changed: 0 additions & 39 deletions
Original file line numberDiff line numberDiff line change
@@ -56,13 +56,6 @@ if (CUDAToolkit_FOUND)
5656
endif()
5757
endif()
5858

59-
if(NOT DEFINED CMAKE_CUDA_COMPILER)
60-
if(WIN32)
61-
set(CMAKE_CUDA_COMPILER "${CUDAToolkit_BIN_DIR}/nvcc.exe")
62-
else()
63-
set(CMAKE_CUDA_COMPILER "${CUDAToolkit_BIN_DIR}/nvcc")
64-
endif()
65-
endif()
6659
enable_language(CUDA)
6760

6861
# TODO: Remove once CCCL 3.2 has been released and bundled with CUDA Toolkit
@@ -129,14 +122,6 @@ if (CUDAToolkit_FOUND)
129122
template-instances/fattn-vec-instance-q4_0-q4_0.cu
130123
template-instances/fattn-vec-instance-q8_0-q8_0.cu
131124
template-instances/fattn-vec-instance-bf16-bf16.cu)
132-
file(GLOB SRCS "template-instances/fattn-vec*tbq3_0-tbq3_0.cu")
133-
list(APPEND GGML_SOURCES_CUDA ${SRCS})
134-
file(GLOB SRCS "template-instances/fattn-vec*tbq3_0-tbq4_0.cu")
135-
list(APPEND GGML_SOURCES_CUDA ${SRCS})
136-
file(GLOB SRCS "template-instances/fattn-vec*tbq4_0-tbq3_0.cu")
137-
list(APPEND GGML_SOURCES_CUDA ${SRCS})
138-
file(GLOB SRCS "template-instances/fattn-vec*tbq4_0-tbq4_0.cu")
139-
list(APPEND GGML_SOURCES_CUDA ${SRCS})
140125
endif()
141126

142127
ggml_add_backend_library(ggml-cuda
@@ -170,30 +155,6 @@ if (CUDAToolkit_FOUND)
170155
add_compile_definitions(GGML_CUDA_NO_PEER_COPY)
171156
endif()
172157

173-
# Eliza fork: opt-in CUDA paths for QJL / PolarQuant / TBQ3_TCQ. Default ON
174-
# so the standard `-DGGML_CUDA=ON` build picks them up; the .cu files are
175-
# always globbed in (above) but their bodies are gated by these macros.
176-
if (GGML_CUDA_QJL)
177-
add_compile_definitions(GGML_CUDA_QJL)
178-
endif()
179-
if (GGML_CUDA_POLARQUANT)
180-
add_compile_definitions(GGML_CUDA_POLARQUANT)
181-
endif()
182-
if (GGML_CUDA_TBQ3_TCQ)
183-
add_compile_definitions(GGML_CUDA_TBQ3_TCQ)
184-
endif()
185-
186-
# ELIZA-CUDA-FUSED-ATTN-QJL
187-
# Fused QJL-K + TBQ-V attention (packages/inference/cuda/fused-attn-qjl-tbq.cu,
188-
# staged in by patchCudaKernels). Body is #ifdef GGML_CUDA_FUSED_ATTN_QJL; this
189-
# flips that define on when -DGGML_CUDA_FUSED_ATTN_QJL=ON is passed. Same shape
190-
# as the GGML_CUDA_QJL / POLARQUANT / TBQ3_TCQ blocks above. Optional kernel
191-
# (packages/inference/AGENTS.md §3) — off by default.
192-
if (GGML_CUDA_FUSED_ATTN_QJL)
193-
add_compile_definitions(GGML_CUDA_FUSED_ATTN_QJL)
194-
message(STATUS "ggml-cuda: GGML_CUDA_FUSED_ATTN_QJL enabled (fused QJL-K + TBQ-V attention)")
195-
endif()
196-
197158
if (GGML_STATIC)
198159
if (WIN32)
199160
# As of 12.3.1 CUDA Toolkit for Windows does not offer a static cublas library

‎ggml/src/ggml-cuda/common.cuh‎

Lines changed: 10 additions & 43 deletions
Original file line numberDiff line numberDiff line change
@@ -40,20 +40,6 @@
4040
#include "vendors/cuda.h"
4141
#endif // defined(GGML_USE_HIP)
4242

43-
template<int q8_1_layout_block_size>
44-
struct block_q8_1_layout {
45-
static_assert(q8_1_layout_block_size % QK8_1 == 0, "q8_1 layout block size must contain whole q8_1 blocks");
46-
47-
static constexpr int q8_1_blocks = q8_1_layout_block_size / QK8_1;
48-
49-
// Scales for all q8_1 blocks in the layout group are stored before the quantized values.
50-
half2 ds[q8_1_blocks];
51-
int32_t qs[q8_1_layout_block_size / sizeof(int32_t)];
52-
};
53-
54-
static_assert(sizeof(block_q8_1_layout<QK8_1>) == sizeof(block_q8_1), "Unexpected block_q8_1 layout size");
55-
static_assert(sizeof(block_q8_1_layout<4 * QK8_1>) == 4 * sizeof(block_q8_1), "Unexpected q8_1 x4 layout size");
56-
5743
#define STRINGIZE_IMPL(...) #__VA_ARGS__
5844
#define STRINGIZE(...) STRINGIZE_IMPL(__VA_ARGS__)
5945

@@ -979,13 +965,6 @@ struct ggml_cuda_type_traits<GGML_TYPE_Q1_0> {
979965
static constexpr int qi = QI1_0;
980966
};
981967

982-
template<>
983-
struct ggml_cuda_type_traits<GGML_TYPE_Q1_0_g32> {
984-
static constexpr int qk = QK1_0_g32;
985-
static constexpr int qr = QR1_0_g32;
986-
static constexpr int qi = QI1_0_g32;
987-
};
988-
989968
template<>
990969
struct ggml_cuda_type_traits<GGML_TYPE_Q4_0> {
991970
static constexpr int qk = QK4_0;
@@ -1021,13 +1000,6 @@ struct ggml_cuda_type_traits<GGML_TYPE_Q8_0> {
10211000
static constexpr int qi = QI8_0;
10221001
};
10231002

1024-
template<>
1025-
struct ggml_cuda_type_traits<GGML_TYPE_Q1_0_g128> {
1026-
static constexpr int qk = QK1_0_g128;
1027-
static constexpr int qr = QR1_0_g128;
1028-
static constexpr int qi = QI1_0_g128;
1029-
};
1030-
10311003
template<>
10321004
struct ggml_cuda_type_traits<GGML_TYPE_MXFP4> {
10331005
static constexpr int qk = QK_MXFP4;
@@ -1171,7 +1143,7 @@ int ggml_cuda_get_device();
11711143
struct ggml_cuda_pool {
11721144
virtual ~ggml_cuda_pool() = default;
11731145

1174-
virtual void * alloc(size_t size, size_t * actual_size, bool overallocate = false) = 0;
1146+
virtual void * alloc(size_t size, size_t * actual_size) = 0;
11751147
virtual void free(void * ptr, size_t size) = 0;
11761148
};
11771149

@@ -1197,16 +1169,16 @@ struct ggml_cuda_pool_alloc {
11971169
}
11981170

11991171
// size is in number of elements
1200-
T * alloc(size_t size, bool overallocate = false) {
1172+
T * alloc(size_t size) {
12011173
GGML_ASSERT(pool != nullptr);
12021174
GGML_ASSERT(ptr == nullptr);
1203-
ptr = (T *) pool->alloc(size * sizeof(T), &this->actual_size, overallocate);
1175+
ptr = (T *) pool->alloc(size * sizeof(T), &this->actual_size);
12041176
return ptr;
12051177
}
12061178

1207-
T * alloc(ggml_cuda_pool & pool, size_t size, bool overallocate = false) {
1179+
T * alloc(ggml_cuda_pool & pool, size_t size) {
12081180
this->pool = &pool;
1209-
return alloc(size, overallocate);
1181+
return alloc(size);
12101182
}
12111183

12121184
T * get() {
@@ -1272,6 +1244,10 @@ struct ggml_cuda_concurrent_event {
12721244
int n_streams = 0;
12731245
std::unordered_map<const ggml_tensor *, int> stream_mapping;
12741246

1247+
// Original order of nodes in this concurrent region (before interleaving)
1248+
// Used to restore grouping for fusion within streams
1249+
std::vector<const ggml_tensor *> original_order;
1250+
12751251
const ggml_tensor * join_node;
12761252

12771253
ggml_cuda_concurrent_event() = default;
@@ -1294,6 +1270,7 @@ struct ggml_cuda_concurrent_event {
12941270
, fork_event(other.fork_event)
12951271
, n_streams(other.n_streams)
12961272
, stream_mapping(std::move(other.stream_mapping))
1273+
, original_order(std::move(other.original_order))
12971274
, join_node(other.join_node) {
12981275
other.fork_event = nullptr;
12991276
}
@@ -1316,11 +1293,6 @@ struct ggml_cuda_concurrent_event {
13161293
const int64_t t_start = (int64_t) t->data;
13171294
const int64_t t_end = t_start + ggml_nbytes(t);
13181295

1319-
// skip empty tensors
1320-
if (t_end == t_start) {
1321-
continue;
1322-
}
1323-
13241296
// skip tensors that overlap with join_node's buffer.
13251297
if ((t_start <= join_start && join_start < t_end) || (join_start <= t_start && t_start < join_end)) {
13261298
continue;
@@ -1342,11 +1314,6 @@ struct ggml_cuda_concurrent_event {
13421314
const int64_t t_start = (int64_t) t->data;
13431315
const int64_t t_end = t_start + ggml_nbytes(t);
13441316

1345-
// skip empty tensors
1346-
if (t_end == t_start) {
1347-
continue;
1348-
}
1349-
13501317
// skip tensors that overlap with join_node's buffer
13511318
if ((t_start <= join_start && join_start < t_end) || (join_start <= t_start && t_start < join_end)) {
13521319
continue;

‎ggml/src/ggml-cuda/convert.cu‎

Lines changed: 0 additions & 40 deletions
Original file line numberDiff line numberDiff line change
@@ -713,10 +713,6 @@ to_fp16_cuda_t ggml_get_to_fp16_cuda(ggml_type type) {
713713
switch (type) {
714714
case GGML_TYPE_Q1_0:
715715
return dequantize_block_cont_cuda<QK1_0, QR1_0, dequantize_q1_0>;
716-
case GGML_TYPE_Q1_0_g32:
717-
return dequantize_block_cont_cuda<QK1_0_g32, QR1_0_g32, dequantize_q1_0_g32>;
718-
case GGML_TYPE_Q1_0_g128:
719-
return dequantize_block_cont_cuda<QK1_0_g128, QR1_0_g128, dequantize_q1_0_g128>;
720716
case GGML_TYPE_Q4_0:
721717
return dequantize_row_q4_0_cuda;
722718
case GGML_TYPE_Q4_1:
@@ -730,10 +726,6 @@ to_fp16_cuda_t ggml_get_to_fp16_cuda(ggml_type type) {
730726
return dequantize_block_q8_0_f16_cuda;
731727
}
732728
return dequantize_block_cont_cuda<QK8_0, QR8_0, dequantize_q8_0>;
733-
case GGML_TYPE_TBQ3_0:
734-
return dequantize_block_cont_cuda<QK_TBQ, QR_TBQ3, dequantize_tbq3_0>;
735-
case GGML_TYPE_TBQ4_0:
736-
return dequantize_block_cont_cuda<QK_TBQ, QR_TBQ4, dequantize_tbq4_0>;
737729
case GGML_TYPE_Q2_K:
738730
return dequantize_row_q2_K_cuda;
739731
case GGML_TYPE_Q3_K:
@@ -779,10 +771,6 @@ to_fp32_cuda_t ggml_get_to_fp32_cuda(ggml_type type) {
779771
switch (type) {
780772
case GGML_TYPE_Q1_0:
781773
return dequantize_block_cont_cuda<QK1_0, QR1_0, dequantize_q1_0>;
782-
case GGML_TYPE_Q1_0_g32:
783-
return dequantize_block_cont_cuda<QK1_0_g32, QR1_0_g32, dequantize_q1_0_g32>;
784-
case GGML_TYPE_Q1_0_g128:
785-
return dequantize_block_cont_cuda<QK1_0_g128, QR1_0_g128, dequantize_q1_0_g128>;
786774
case GGML_TYPE_Q4_0:
787775
return dequantize_row_q4_0_cuda;
788776
case GGML_TYPE_Q4_1:
@@ -793,10 +781,6 @@ to_fp32_cuda_t ggml_get_to_fp32_cuda(ggml_type type) {
793781
return dequantize_block_cont_cuda<QK5_1, QR5_1, dequantize_q5_1>;
794782
case GGML_TYPE_Q8_0:
795783
return dequantize_block_cont_cuda<QK8_0, QR8_0, dequantize_q8_0>;
796-
case GGML_TYPE_TBQ3_0:
797-
return dequantize_block_cont_cuda<QK_TBQ, QR_TBQ3, dequantize_tbq3_0>;
798-
case GGML_TYPE_TBQ4_0:
799-
return dequantize_block_cont_cuda<QK_TBQ, QR_TBQ4, dequantize_tbq4_0>;
800784
case GGML_TYPE_Q2_K:
801785
return dequantize_row_q2_K_cuda;
802786
case GGML_TYPE_Q3_K:
@@ -844,10 +828,6 @@ to_fp16_nc_cuda_t ggml_get_to_fp16_nc_cuda(ggml_type type) {
844828
return convert_unary_cuda<float>;
845829
case GGML_TYPE_Q1_0:
846830
return dequantize_block_cuda<QK1_0, QR1_0, dequantize_q1_0>;
847-
case GGML_TYPE_Q1_0_g32:
848-
return dequantize_block_cuda<QK1_0_g32, QR1_0_g32, dequantize_q1_0_g32>;
849-
case GGML_TYPE_Q1_0_g128:
850-
return dequantize_block_cuda<QK1_0_g128, QR1_0_g128, dequantize_q1_0_g128>;
851831
case GGML_TYPE_Q4_0:
852832
return dequantize_block_cuda<QK4_0, QR4_0, dequantize_q4_0>;
853833
case GGML_TYPE_Q4_1:
@@ -858,10 +838,6 @@ to_fp16_nc_cuda_t ggml_get_to_fp16_nc_cuda(ggml_type type) {
858838
return dequantize_block_cuda<QK5_1, QR5_1, dequantize_q5_1>;
859839
case GGML_TYPE_Q8_0:
860840
return dequantize_block_cuda<QK8_0, QR8_0, dequantize_q8_0>;
861-
case GGML_TYPE_TBQ3_0:
862-
return dequantize_block_cuda<QK_TBQ, QR_TBQ3, dequantize_tbq3_0>;
863-
case GGML_TYPE_TBQ4_0:
864-
return dequantize_block_cuda<QK_TBQ, QR_TBQ4, dequantize_tbq4_0>;
865841
case GGML_TYPE_BF16:
866842
return convert_unary_cuda<nv_bfloat16>;
867843
default:
@@ -875,10 +851,6 @@ to_bf16_nc_cuda_t ggml_get_to_bf16_nc_cuda(ggml_type type) {
875851
return convert_unary_cuda<float, nv_bfloat16>;
876852
case GGML_TYPE_Q1_0:
877853
return dequantize_block_cuda<QK1_0, QR1_0, dequantize_q1_0>;
878-
case GGML_TYPE_Q1_0_g32:
879-
return dequantize_block_cuda<QK1_0_g32, QR1_0_g32, dequantize_q1_0_g32>;
880-
case GGML_TYPE_Q1_0_g128:
881-
return dequantize_block_cuda<QK1_0_g128, QR1_0_g128, dequantize_q1_0_g128>;
882854
case GGML_TYPE_Q4_0:
883855
return dequantize_block_cuda<QK4_0, QR4_0, dequantize_q4_0>;
884856
case GGML_TYPE_Q4_1:
@@ -889,10 +861,6 @@ to_bf16_nc_cuda_t ggml_get_to_bf16_nc_cuda(ggml_type type) {
889861
return dequantize_block_cuda<QK5_1, QR5_1, dequantize_q5_1>;
890862
case GGML_TYPE_Q8_0:
891863
return dequantize_block_cuda<QK8_0, QR8_0, dequantize_q8_0>;
892-
case GGML_TYPE_TBQ3_0:
893-
return dequantize_block_cuda<QK_TBQ, QR_TBQ3, dequantize_tbq3_0>;
894-
case GGML_TYPE_TBQ4_0:
895-
return dequantize_block_cuda<QK_TBQ, QR_TBQ4, dequantize_tbq4_0>;
896864
case GGML_TYPE_F16:
897865
return convert_unary_cuda<half, nv_bfloat16>;
898866
default:
@@ -906,10 +874,6 @@ to_fp32_nc_cuda_t ggml_get_to_fp32_nc_cuda(ggml_type type) {
906874
return convert_unary_cuda<half, float>;
907875
case GGML_TYPE_Q1_0:
908876
return dequantize_block_cuda<QK1_0, QR1_0, dequantize_q1_0>;
909-
case GGML_TYPE_Q1_0_g32:
910-
return dequantize_block_cuda<QK1_0_g32, QR1_0_g32, dequantize_q1_0_g32>;
911-
case GGML_TYPE_Q1_0_g128:
912-
return dequantize_block_cuda<QK1_0_g128, QR1_0_g128, dequantize_q1_0_g128>;
913877
case GGML_TYPE_Q4_0:
914878
return dequantize_block_cuda<QK4_0, QR4_0, dequantize_q4_0>;
915879
case GGML_TYPE_Q4_1:
@@ -920,10 +884,6 @@ to_fp32_nc_cuda_t ggml_get_to_fp32_nc_cuda(ggml_type type) {
920884
return dequantize_block_cuda<QK5_1, QR5_1, dequantize_q5_1>;
921885
case GGML_TYPE_Q8_0:
922886
return dequantize_block_cuda<QK8_0, QR8_0, dequantize_q8_0>;
923-
case GGML_TYPE_TBQ3_0:
924-
return dequantize_block_cuda<QK_TBQ, QR_TBQ3, dequantize_tbq3_0>;
925-
case GGML_TYPE_TBQ4_0:
926-
return dequantize_block_cuda<QK_TBQ, QR_TBQ4, dequantize_tbq4_0>;
927887
case GGML_TYPE_BF16:
928888
return convert_unary_cuda<nv_bfloat16, float>;
929889
default:

‎ggml/src/ggml-cuda/cpy-utils.cuh‎

Lines changed: 0 additions & 9 deletions
Original file line numberDiff line numberDiff line change
@@ -2,7 +2,6 @@
22

33
#include "ggml-common.h"
44
#include "convert.cuh"
5-
#include "turboquant.cuh"
65

76
static __device__ __forceinline__ int best_index_int8(int n, const int8_t * val, float x) {
87
if (x <= val[0]) return 0;
@@ -208,14 +207,6 @@ static __device__ void cpy_blck_f32_q8_0(const char * cxi, char * cdsti) {
208207
quantize_f32_q8_0_block((const float *)cxi, (block_q8_0 *)cdsti);
209208
}
210209

211-
static __device__ void cpy_blck_f32_tbq3_0(const char * cxi, char * cdsti) {
212-
quantize_f32_tbq3_0_block((const float *) cxi, (block_tbq3_0 *) cdsti);
213-
}
214-
215-
static __device__ void cpy_blck_f32_tbq4_0(const char * cxi, char * cdsti) {
216-
quantize_f32_tbq4_0_block((const float *) cxi, (block_tbq4_0 *) cdsti);
217-
}
218-
219210
static __device__ void cpy_blck_f32_iq4_nl(const char * cxi, char * cdsti) {
220211
quantize_f32_iq4_nl_block((const float *)cxi, (block_iq4_nl *)cdsti);
221212
}

0 commit comments

Comments
 (0)