Skip to content

Commit a79a154

Browse files
yeahdongcnqnixsynapse
yeahdongcn
authored and
qnixsynapse
committed
musa: add GGML_UNUSED_VARS (ggml-org#15446)
Signed-off-by: Xiaodong Ye <[email protected]>
1 parent b161e6a commit a79a154

File tree

17 files changed

+113
-148
lines changed

17 files changed

+113
-148
lines changed

‎ggml/include/ggml.h‎

Lines changed: 7 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -244,6 +244,13 @@
244244
#define GGML_MROPE_SECTIONS 4
245245

246246
#define GGML_UNUSED(x) (void)(x)
247+
#ifdef __CUDACC__
248+
template<typename... Args>
249+
__host__ __device__ constexpr inline void ggml_unused_vars_impl(Args&&...) noexcept {}
250+
#define GGML_UNUSED_VARS(...) ggml_unused_vars_impl(__VA_ARGS__)
251+
#else
252+
#define GGML_UNUSED_VARS(...) do { (void)sizeof((__VA_ARGS__, 0)); } while(0)
253+
#endif // __CUDACC__
247254

248255
#define GGML_PAD(x, n) (((x) + (n) - 1) & ~((n) - 1))
249256

‎ggml/src/ggml-cuda/conv-transpose-1d.cu‎

Lines changed: 1 addition & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -34,10 +34,7 @@ static __global__ void conv_transpose_1d_kernel(
3434
}
3535
}
3636
dst[global_index] = accumulator;
37-
GGML_UNUSED(p0); GGML_UNUSED(d0); GGML_UNUSED(src0_ne3);
38-
GGML_UNUSED(src1_ne3); GGML_UNUSED(dst_ne3);
39-
GGML_UNUSED(src1_ne1); GGML_UNUSED(dst_ne1);
40-
GGML_UNUSED(src1_ne2); GGML_UNUSED(dst_ne2);
37+
GGML_UNUSED_VARS(p0, d0, src0_ne3, src1_ne3, dst_ne3, src1_ne1, dst_ne1, src1_ne2, dst_ne2);
4138
}
4239

4340
static void conv_transpose_1d_f32_f32_cuda(

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

Lines changed: 1 addition & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -71,9 +71,7 @@ static __global__ void dequantize_block_q8_0_f16(const void * __restrict__ vx, h
7171
y2[iy/2 + threadIdx.x] = __hmul2(make_half2(qs.x, qs.y), __half2half2(d));
7272
}
7373
#else
74-
GGML_UNUSED(vx);
75-
GGML_UNUSED(y);
76-
GGML_UNUSED(k);
74+
GGML_UNUSED_VARS(vx, y, k);
7775
NO_DEVICE_CODE;
7876
#endif // __CUDA_ARCH__ >= GGML_CUDA_CC_PASCAL
7977
}

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

Lines changed: 1 addition & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -134,8 +134,7 @@ void ggml_cuda_cpy_dest_ptrs_copy(ggml_cuda_graph * cuda_graph, char ** host_des
134134
CUDA_CHECK(cudaMemcpyAsync(cuda_graph->dest_ptrs_d, host_dest_ptrs, host_dest_ptrs_size*sizeof(char *), cudaMemcpyHostToDevice, stream));
135135
cuda_graph->graph_cpynode_index = 0; // reset index
136136
#else
137-
GGML_UNUSED(cuda_graph); GGML_UNUSED(host_dest_ptrs);
138-
GGML_UNUSED(host_dest_ptrs_size); GGML_UNUSED(stream);
137+
GGML_UNUSED_VARS(cuda_graph, host_dest_ptrs, host_dest_ptrs_size, stream);
139138
#endif
140139
}
141140

‎ggml/src/ggml-cuda/fattn-mma-f16.cuh‎

Lines changed: 18 additions & 27 deletions
Original file line numberDiff line numberDiff line change
@@ -767,14 +767,11 @@ static __device__ __forceinline__ void flash_attn_ext_f16_iter(
767767
}
768768
}
769769
#else
770-
GGML_UNUSED(Q_f2); GGML_UNUSED(K_h2); GGML_UNUSED(V_h2);
771-
GGML_UNUSED(mask_h2); GGML_UNUSED(dstk); GGML_UNUSED(dstk_fixup);
772-
GGML_UNUSED(scale); GGML_UNUSED(slope); GGML_UNUSED(logit_softcap);
773-
GGML_UNUSED(ne01); GGML_UNUSED(ne02); GGML_UNUSED(stride_K); GGML_UNUSED(stride_V);
774-
GGML_UNUSED(stride_mask); GGML_UNUSED(tile_K);
775-
GGML_UNUSED(tile_V); GGML_UNUSED(tile_mask); GGML_UNUSED(Q_B);
776-
GGML_UNUSED(VKQ_C); GGML_UNUSED(KQ_max); GGML_UNUSED(KQ_rowsum);
777-
GGML_UNUSED(kb0); GGML_UNUSED(tile_Q);
770+
GGML_UNUSED_VARS(Q_f2, K_h2, V_h2, mask_h2, dstk, dstk_fixup,
771+
scale, slope, logit_softcap, ne01, ne02,
772+
stride_K, stride_V, stride_mask,
773+
tile_Q, tile_K, tile_V, tile_mask,
774+
Q_B, VKQ_C, KQ_max, KQ_rowsum, kb0);
778775
NO_DEVICE_CODE;
779776
#endif // TURING_MMA_AVAILABLE
780777
}
@@ -1236,14 +1233,10 @@ static __device__ __forceinline__ void flash_attn_ext_f16_process_tile(
12361233
}
12371234
}
12381235
#else
1239-
GGML_UNUSED(Q_f2); GGML_UNUSED(K_h2); GGML_UNUSED(V_h2);
1240-
GGML_UNUSED(mask_h2); GGML_UNUSED(sinks_f);
1241-
GGML_UNUSED(dstk); GGML_UNUSED(dstk_fixup);
1242-
GGML_UNUSED(scale); GGML_UNUSED(slope); GGML_UNUSED(logit_softcap);
1243-
GGML_UNUSED(ne01); GGML_UNUSED(ne02);
1244-
GGML_UNUSED(stride_Q1); GGML_UNUSED(stride_Q2);
1245-
GGML_UNUSED(stride_K); GGML_UNUSED(stride_V); GGML_UNUSED(stride_mask);
1246-
GGML_UNUSED(jt); GGML_UNUSED(kb0_start); GGML_UNUSED(kb0_stop);
1236+
GGML_UNUSED_VARS(Q_f2, K_h2, V_h2, mask_h2, sinks_f, dstk, dstk_fixup,
1237+
scale, slope, logit_softcap, ne01, ne02,
1238+
stride_Q1, stride_Q2, stride_K, stride_V, stride_mask,
1239+
jt, kb0_start, kb0_stop);
12471240
NO_DEVICE_CODE;
12481241
#endif // TURING_MMA_AVAILABLE
12491242
}
@@ -1397,17 +1390,15 @@ static __global__ void flash_attn_ext_f16(
13971390
(Q_f2, K_h2, V_h2, mask_h2, sinks_f, dstk, dst_meta, scale, slope, logit_softcap,
13981391
ne01, ne02, stride_Q1, stride_Q2, stride_K, stride_V, stride_mask, jt, kb0_start_kernel, kb0_stop_kernel);
13991392
#else
1400-
GGML_UNUSED(Q); GGML_UNUSED(K); GGML_UNUSED(V); GGML_UNUSED(mask);
1401-
GGML_UNUSED(sinks); GGML_UNUSED(KV_max); GGML_UNUSED(dst); GGML_UNUSED(dst_meta);
1402-
GGML_UNUSED(scale); GGML_UNUSED(max_bias); GGML_UNUSED(m0); GGML_UNUSED(m1);
1403-
GGML_UNUSED(n_head_log2); GGML_UNUSED(logit_softcap);
1404-
GGML_UNUSED(ne00); GGML_UNUSED(ne01); GGML_UNUSED(ne02); GGML_UNUSED(ne03);
1405-
GGML_UNUSED(nb01); GGML_UNUSED(nb02); GGML_UNUSED(nb03);
1406-
GGML_UNUSED(ne10); GGML_UNUSED(ne11); GGML_UNUSED(ne12); GGML_UNUSED(ne13);
1407-
GGML_UNUSED(nb11); GGML_UNUSED(nb12); GGML_UNUSED(nb13);
1408-
GGML_UNUSED(nb21); GGML_UNUSED(nb22); GGML_UNUSED(nb23);
1409-
GGML_UNUSED(ne31); GGML_UNUSED(ne32); GGML_UNUSED(ne33);
1410-
GGML_UNUSED(nb31); GGML_UNUSED(nb32); GGML_UNUSED(nb33);
1393+
GGML_UNUSED_VARS(Q, K, V, mask, sinks, KV_max, dst, dst_meta, scale,
1394+
max_bias, m0, m1, n_head_log2, logit_softcap,
1395+
ne00, ne01, ne02, ne03,
1396+
nb01, nb02, nb03,
1397+
ne10, ne11, ne12, ne13,
1398+
nb11, nb12, nb13,
1399+
nb21, nb22, nb23,
1400+
ne31, ne32, ne33,
1401+
nb31, nb32, nb33);
14111402
NO_DEVICE_CODE;
14121403
#endif // defined(FLASH_ATTN_AVAILABLE) && defined(TURING_MMA_AVAILABLE)
14131404
}

‎ggml/src/ggml-cuda/fattn-tile-f16.cu‎

Lines changed: 9 additions & 11 deletions
Original file line numberDiff line numberDiff line change
@@ -299,17 +299,15 @@ static __global__ void flash_attn_tile_ext_f16(
299299
}
300300
}
301301
#else
302-
GGML_UNUSED(Q); GGML_UNUSED(K); GGML_UNUSED(V); GGML_UNUSED(mask);
303-
GGML_UNUSED(sinks); GGML_UNUSED(KV_max); GGML_UNUSED(dst); GGML_UNUSED(dst_meta);
304-
GGML_UNUSED(scale); GGML_UNUSED(max_bias); GGML_UNUSED(m0); GGML_UNUSED(m1);
305-
GGML_UNUSED(n_head_log2); GGML_UNUSED(logit_softcap);
306-
GGML_UNUSED(ne00); GGML_UNUSED(ne01); GGML_UNUSED(ne02); GGML_UNUSED(ne03);
307-
GGML_UNUSED(nb01); GGML_UNUSED(nb02); GGML_UNUSED(nb03);
308-
GGML_UNUSED(ne10); GGML_UNUSED(ne11); GGML_UNUSED(ne12); GGML_UNUSED(ne13);
309-
GGML_UNUSED(nb11); GGML_UNUSED(nb12); GGML_UNUSED(nb13);
310-
GGML_UNUSED(nb21); GGML_UNUSED(nb22); GGML_UNUSED(nb23);
311-
GGML_UNUSED(ne31); GGML_UNUSED(ne32); GGML_UNUSED(ne33);
312-
GGML_UNUSED(nb31); GGML_UNUSED(nb32); GGML_UNUSED(nb33);
302+
GGML_UNUSED_VARS(Q, K, V, mask, sinks, KV_max, dst, dst_meta, scale,
303+
max_bias, m0, m1, n_head_log2, logit_softcap,
304+
ne00, ne01, ne02, ne03,
305+
nb01, nb02, nb03,
306+
ne10, ne11, ne12, ne13,
307+
nb11, nb12, nb13,
308+
nb21, nb22, nb23,
309+
ne31, ne32, ne33,
310+
nb31, nb32, nb33);
313311
NO_DEVICE_CODE;
314312
#endif // defined(FLASH_ATTN_AVAILABLE) && defined(FP16_AVAILABLE)
315313
}

‎ggml/src/ggml-cuda/fattn-tile-f32.cu‎

Lines changed: 18 additions & 11 deletions
Original file line numberDiff line numberDiff line change
@@ -38,6 +38,15 @@ static __global__ void flash_attn_tile_ext_f32(
3838
return;
3939
#endif // FP16_MMA_AVAILABLE
4040
if (use_logit_softcap && !(D == 128 || D == 256)) {
41+
GGML_UNUSED_VARS(Q, K, V, mask, sinks, KV_max, dst, dst_meta, scale,
42+
max_bias, m0, m1, n_head_log2, logit_softcap,
43+
ne00, ne01, ne02, ne03,
44+
nb01, nb02, nb03,
45+
ne10, ne11, ne12, ne13,
46+
nb11, nb12, nb13,
47+
nb21, nb22, nb23,
48+
ne31, ne32, ne33,
49+
nb31, nb32, nb33);
4150
NO_DEVICE_CODE;
4251
return;
4352
}
@@ -301,17 +310,15 @@ static __global__ void flash_attn_tile_ext_f32(
301310
}
302311
}
303312
#else
304-
GGML_UNUSED(Q); GGML_UNUSED(K); GGML_UNUSED(V); GGML_UNUSED(mask);
305-
GGML_UNUSED(sinks); GGML_UNUSED(KV_max); GGML_UNUSED(dst); GGML_UNUSED(dst_meta);
306-
GGML_UNUSED(scale); GGML_UNUSED(max_bias); GGML_UNUSED(m0); GGML_UNUSED(m1);
307-
GGML_UNUSED(n_head_log2); GGML_UNUSED(logit_softcap);
308-
GGML_UNUSED(ne00); GGML_UNUSED(ne01); GGML_UNUSED(ne02); GGML_UNUSED(ne03);
309-
GGML_UNUSED(nb01); GGML_UNUSED(nb02); GGML_UNUSED(nb03);
310-
GGML_UNUSED(ne10); GGML_UNUSED(ne11); GGML_UNUSED(ne12); GGML_UNUSED(ne13);
311-
GGML_UNUSED(nb11); GGML_UNUSED(nb12); GGML_UNUSED(nb13);
312-
GGML_UNUSED(nb21); GGML_UNUSED(nb22); GGML_UNUSED(nb23);
313-
GGML_UNUSED(ne31); GGML_UNUSED(ne32); GGML_UNUSED(ne33);
314-
GGML_UNUSED(nb31); GGML_UNUSED(nb32); GGML_UNUSED(nb33);
313+
GGML_UNUSED_VARS(Q, K, V, mask, sinks, KV_max, dst, dst_meta, scale,
314+
max_bias, m0, m1, n_head_log2, logit_softcap,
315+
ne00, ne01, ne02, ne03,
316+
nb01, nb02, nb03,
317+
ne10, ne11, ne12, ne13,
318+
nb11, nb12, nb13,
319+
nb21, nb22, nb23,
320+
ne31, ne32, ne33,
321+
nb31, nb32, nb33);
315322
NO_DEVICE_CODE;
316323
#endif // FLASH_ATTN_AVAILABLE
317324
}

‎ggml/src/ggml-cuda/fattn-vec-f16.cuh‎

Lines changed: 9 additions & 11 deletions
Original file line numberDiff line numberDiff line change
@@ -349,17 +349,15 @@ static __global__ void flash_attn_vec_ext_f16(
349349
dst_meta[((sequence*ne01 + ic0 + tid)*ne02 + head)*gridDim.y + blockIdx.y] = make_float2(kqmax[tid], kqsum[tid]);
350350
}
351351
#else
352-
GGML_UNUSED(Q); GGML_UNUSED(K); GGML_UNUSED(V); GGML_UNUSED(mask);
353-
GGML_UNUSED(sinks); GGML_UNUSED(KV_max); GGML_UNUSED(dst); GGML_UNUSED(dst_meta);
354-
GGML_UNUSED(scale); GGML_UNUSED(max_bias); GGML_UNUSED(m0); GGML_UNUSED(m1);
355-
GGML_UNUSED(n_head_log2); GGML_UNUSED(logit_softcap);
356-
GGML_UNUSED(ne00); GGML_UNUSED(ne01); GGML_UNUSED(ne02); GGML_UNUSED(ne03);
357-
GGML_UNUSED(nb01); GGML_UNUSED(nb02); GGML_UNUSED(nb03);
358-
GGML_UNUSED(ne10); GGML_UNUSED(ne11); GGML_UNUSED(ne12); GGML_UNUSED(ne13);
359-
GGML_UNUSED(nb11); GGML_UNUSED(nb12); GGML_UNUSED(nb13);
360-
GGML_UNUSED(nb21); GGML_UNUSED(nb22); GGML_UNUSED(nb23);
361-
GGML_UNUSED(ne31); GGML_UNUSED(ne32); GGML_UNUSED(ne33);
362-
GGML_UNUSED(nb31); GGML_UNUSED(nb32); GGML_UNUSED(nb33);
352+
GGML_UNUSED_VARS(Q, K, V, mask, sinks, KV_max, dst, dst_meta, scale,
353+
max_bias, m0, m1, n_head_log2, logit_softcap,
354+
ne00, ne01, ne02, ne03,
355+
nb01, nb02, nb03,
356+
ne10, ne11, ne12, ne13,
357+
nb11, nb12, nb13,
358+
nb21, nb22, nb23,
359+
ne31, ne32, ne33,
360+
nb31, nb32, nb33);
363361
NO_DEVICE_CODE;
364362
#endif // defined(FLASH_ATTN_AVAILABLE) && defined(FP16_AVAILABLE)
365363
}

‎ggml/src/ggml-cuda/fattn-vec-f32.cuh‎

Lines changed: 18 additions & 11 deletions
Original file line numberDiff line numberDiff line change
@@ -37,6 +37,15 @@ static __global__ void flash_attn_vec_ext_f32(
3737

3838
// Skip unused kernel variants for faster compilation:
3939
if (use_logit_softcap && !(D == 128 || D == 256)) {
40+
GGML_UNUSED_VARS(Q, K, V, mask, sinks, KV_max, dst, dst_meta, scale,
41+
max_bias, m0, m1, n_head_log2, logit_softcap,
42+
ne00, ne01, ne02, ne03,
43+
nb01, nb02, nb03,
44+
ne10, ne11, ne12, ne13,
45+
nb11, nb12, nb13,
46+
nb21, nb22, nb23,
47+
ne31, ne32, ne33,
48+
nb31, nb32, nb33);
4049
NO_DEVICE_CODE;
4150
return;
4251
}
@@ -334,17 +343,15 @@ static __global__ void flash_attn_vec_ext_f32(
334343
dst_meta[((sequence*ne01 + ic0 + tid)*ne02 + head)*gridDim.y + blockIdx.y] = make_float2(kqmax[tid], kqsum[tid]);
335344
}
336345
#else
337-
GGML_UNUSED(Q); GGML_UNUSED(K); GGML_UNUSED(V); GGML_UNUSED(mask);
338-
GGML_UNUSED(sinks); GGML_UNUSED(KV_max); GGML_UNUSED(dst); GGML_UNUSED(dst_meta);
339-
GGML_UNUSED(scale); GGML_UNUSED(max_bias); GGML_UNUSED(m0); GGML_UNUSED(m1);
340-
GGML_UNUSED(n_head_log2); GGML_UNUSED(logit_softcap);
341-
GGML_UNUSED(ne00); GGML_UNUSED(ne01); GGML_UNUSED(ne02); GGML_UNUSED(ne03);
342-
GGML_UNUSED(ne10); GGML_UNUSED(ne11); GGML_UNUSED(ne12); GGML_UNUSED(ne13);
343-
GGML_UNUSED(ne31); GGML_UNUSED(ne32); GGML_UNUSED(ne33);
344-
GGML_UNUSED(nb31); GGML_UNUSED(nb32); GGML_UNUSED(nb33);
345-
GGML_UNUSED(nb01); GGML_UNUSED(nb02); GGML_UNUSED(nb03);
346-
GGML_UNUSED(nb11); GGML_UNUSED(nb12); GGML_UNUSED(nb13);
347-
GGML_UNUSED(nb21); GGML_UNUSED(nb22); GGML_UNUSED(nb23);
346+
GGML_UNUSED_VARS(Q, K, V, mask, sinks, KV_max, dst, dst_meta, scale,
347+
max_bias, m0, m1, n_head_log2, logit_softcap,
348+
ne00, ne01, ne02, ne03,
349+
nb01, nb02, nb03,
350+
ne10, ne11, ne12, ne13,
351+
nb11, nb12, nb13,
352+
nb21, nb22, nb23,
353+
ne31, ne32, ne33,
354+
nb31, nb32, nb33);
348355
NO_DEVICE_CODE;
349356
#endif // FLASH_ATTN_AVAILABLE
350357
}

‎ggml/src/ggml-cuda/fattn-wmma-f16.cu‎

Lines changed: 9 additions & 10 deletions
Original file line numberDiff line numberDiff line change
@@ -471,16 +471,15 @@ static __global__ void flash_attn_ext_f16(
471471
dst_meta[j_dst_unrolled] = dst_meta_val;
472472
}
473473
#else
474-
GGML_UNUSED(Q); GGML_UNUSED(K); GGML_UNUSED(V); GGML_UNUSED(mask);
475-
GGML_UNUSED(sinks); GGML_UNUSED(KV_max); GGML_UNUSED(dst); GGML_UNUSED(dst_meta);
476-
GGML_UNUSED(scale); GGML_UNUSED(max_bias); GGML_UNUSED(m0); GGML_UNUSED(m1);
477-
GGML_UNUSED(n_head_log2); GGML_UNUSED(logit_softcap);
478-
GGML_UNUSED(ne00); GGML_UNUSED(ne01); GGML_UNUSED(ne02); GGML_UNUSED(ne03);
479-
GGML_UNUSED(ne10); GGML_UNUSED(ne11); GGML_UNUSED(ne12); GGML_UNUSED(ne13);
480-
GGML_UNUSED(ne31); GGML_UNUSED(ne32); GGML_UNUSED(ne33); GGML_UNUSED(nb31);
481-
GGML_UNUSED(nb32); GGML_UNUSED(nb33); GGML_UNUSED(nb01); GGML_UNUSED(nb02);
482-
GGML_UNUSED(nb03); GGML_UNUSED(nb11); GGML_UNUSED(nb12); GGML_UNUSED(nb13);
483-
GGML_UNUSED(nb21); GGML_UNUSED(nb22); GGML_UNUSED(nb23);
474+
GGML_UNUSED_VARS(Q, K, V, mask, sinks, KV_max, dst, dst_meta, scale,
475+
max_bias, m0, m1, n_head_log2, logit_softcap,
476+
ne00, ne01, ne02, ne03,
477+
nb01, nb02, nb03,
478+
ne10, ne11, ne12, ne13,
479+
nb11, nb12, nb13,
480+
nb21, nb22, nb23,
481+
ne31, ne32, ne33,
482+
nb31, nb32, nb33);
484483
NO_DEVICE_CODE;
485484
#endif // defined(FLASH_ATTN_AVAILABLE) && (__CUDA_ARCH__ == GGML_CUDA_CC_VOLTA || (defined(GGML_HIP_ROCWMMA_FATTN) && defined(FP16_MMA_AVAILABLE)))
486485
}

0 commit comments

Comments
 (0)