Skip to content

Commit fe831be

Browse files
yeahdongcnyeahdongcn
committed
musa: fix build warnings
Signed-off-by: Xiaodong Ye <[email protected]>
1 parent 25ff6f7 commit fe831be

File tree

9 files changed

+31
-54
lines changed

9 files changed

+31
-54
lines changed

‎ggml/src/ggml-cuda/add-id.cu‎

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -11,14 +11,14 @@ static __global__ void add_id_kernel(
1111
const int64_t i1 = blockIdx.x;
1212
const int64_t i2 = blockIdx.y;
1313

14-
const int i11 = *(int32_t *) ((char *) src2 + i1*sizeof(int32_t) + i2*nb21);
14+
const int i11 = *(const int32_t *) ((const char *) src2 + i1*sizeof(int32_t) + i2*nb21);
1515

1616
const size_t nb1 = ne0 * sizeof(float);
1717
const size_t nb2 = ne1 * nb1;
1818

1919
float * dst_row = (float *)((char *)dst + i1*nb1 + i2*nb2);
20-
const float * src0_row = (const float *)((char *)src0 + i1*nb01 + i2*nb02);
21-
const float * src1_row = (const float *)((char *)src1 + i11*nb11);
20+
const float * src0_row = (const float *)((const char *)src0 + i1*nb01 + i2*nb02);
21+
const float * src1_row = (const float *)((const char *)src1 + i11*nb11);
2222

2323
for (int64_t i0 = threadIdx.x; i0 < ne0; i0 += blockDim.x) {
2424
dst_row[i0] = src0_row[i0] + src1_row[i0];

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

Lines changed: 7 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -1237,10 +1237,12 @@ static __device__ __forceinline__ void flash_attn_ext_f16_process_tile(
12371237
}
12381238
#else
12391239
GGML_UNUSED(Q_f2); GGML_UNUSED(K_h2); GGML_UNUSED(V_h2);
1240-
GGML_UNUSED(mask_h2); GGML_UNUSED(dstk); GGML_UNUSED(dstk_fixup);
1240+
GGML_UNUSED(mask_h2); GGML_UNUSED(sinks_f);
1241+
GGML_UNUSED(dstk); GGML_UNUSED(dstk_fixup);
12411242
GGML_UNUSED(scale); GGML_UNUSED(slope); GGML_UNUSED(logit_softcap);
1242-
GGML_UNUSED(ne01); GGML_UNUSED(ne02); GGML_UNUSED(stride_Q1);
1243-
GGML_UNUSED(stride_Q2); GGML_UNUSED(stride_K); GGML_UNUSED(stride_V); GGML_UNUSED(stride_mask);
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);
12441246
GGML_UNUSED(jt); GGML_UNUSED(kb0_start); GGML_UNUSED(kb0_stop);
12451247
NO_DEVICE_CODE;
12461248
#endif // TURING_MMA_AVAILABLE
@@ -1395,8 +1397,8 @@ static __global__ void flash_attn_ext_f16(
13951397
(Q_f2, K_h2, V_h2, mask_h2, sinks_f, dstk, dst_meta, scale, slope, logit_softcap,
13961398
ne01, ne02, stride_Q1, stride_Q2, stride_K, stride_V, stride_mask, jt, kb0_start_kernel, kb0_stop_kernel);
13971399
#else
1398-
GGML_UNUSED(Q); GGML_UNUSED(K); GGML_UNUSED(V); GGML_UNUSED(mask); GGML_UNUSED(sinks);
1399-
GGML_UNUSED(dst); GGML_UNUSED(dst_meta);
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);
14001402
GGML_UNUSED(scale); GGML_UNUSED(max_bias); GGML_UNUSED(m0); GGML_UNUSED(m1);
14011403
GGML_UNUSED(n_head_log2); GGML_UNUSED(logit_softcap);
14021404
GGML_UNUSED(ne00); GGML_UNUSED(ne01); GGML_UNUSED(ne02); GGML_UNUSED(ne03);

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

Lines changed: 10 additions & 10 deletions
Original file line numberDiff line numberDiff line change
@@ -299,17 +299,17 @@ 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); GGML_UNUSED(sinks);
303-
GGML_UNUSED(dst); GGML_UNUSED(dst_meta); GGML_UNUSED(scale);
304-
GGML_UNUSED(max_bias); GGML_UNUSED(m0); GGML_UNUSED(m1);
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);
305305
GGML_UNUSED(n_head_log2); GGML_UNUSED(logit_softcap);
306-
GGML_UNUSED(ne00); GGML_UNUSED(ne01); GGML_UNUSED(ne02);
307-
GGML_UNUSED(ne03); GGML_UNUSED(ne10); GGML_UNUSED(ne11);
308-
GGML_UNUSED(ne12); GGML_UNUSED(ne13); GGML_UNUSED(ne31); GGML_UNUSED(ne32); GGML_UNUSED(ne33);
309-
GGML_UNUSED(nb31); GGML_UNUSED(nb32); GGML_UNUSED(nb33); GGML_UNUSED(nb01); GGML_UNUSED(nb02);
310-
GGML_UNUSED(nb03); GGML_UNUSED(nb11); GGML_UNUSED(nb12);
311-
GGML_UNUSED(nb13); GGML_UNUSED(nb21); GGML_UNUSED(nb22);
312-
GGML_UNUSED(nb23);
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);
313313
NO_DEVICE_CODE;
314314
#endif // defined(FLASH_ATTN_AVAILABLE) && defined(FP16_AVAILABLE)
315315
}

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

Lines changed: 1 addition & 12 deletions
Original file line numberDiff line numberDiff line change
@@ -38,17 +38,6 @@ 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(Q); GGML_UNUSED(K); GGML_UNUSED(V); GGML_UNUSED(mask); GGML_UNUSED(sinks);
42-
GGML_UNUSED(dst); GGML_UNUSED(dst_meta);
43-
GGML_UNUSED(scale); GGML_UNUSED(max_bias); GGML_UNUSED(m0); GGML_UNUSED(m1);
44-
GGML_UNUSED(n_head_log2); GGML_UNUSED(logit_softcap);
45-
GGML_UNUSED(ne00); GGML_UNUSED(ne01); GGML_UNUSED(ne02); GGML_UNUSED(ne03);
46-
GGML_UNUSED(nb01); GGML_UNUSED(nb02); GGML_UNUSED(nb03);
47-
GGML_UNUSED(ne10); GGML_UNUSED(ne11); GGML_UNUSED(ne12); GGML_UNUSED(ne13);
48-
GGML_UNUSED(nb11); GGML_UNUSED(nb12); GGML_UNUSED(nb13);
49-
GGML_UNUSED(nb21); GGML_UNUSED(nb22); GGML_UNUSED(nb23);
50-
GGML_UNUSED(ne31); GGML_UNUSED(ne32); GGML_UNUSED(ne33);
51-
GGML_UNUSED(nb31); GGML_UNUSED(nb32); GGML_UNUSED(nb33);
5241
NO_DEVICE_CODE;
5342
return;
5443
}
@@ -313,7 +302,7 @@ static __global__ void flash_attn_tile_ext_f32(
313302
}
314303
#else
315304
GGML_UNUSED(Q); GGML_UNUSED(K); GGML_UNUSED(V); GGML_UNUSED(mask);
316-
GGML_UNUSED(dst); GGML_UNUSED(dst_meta);
305+
GGML_UNUSED(sinks); GGML_UNUSED(KV_max); GGML_UNUSED(dst); GGML_UNUSED(dst_meta);
317306
GGML_UNUSED(scale); GGML_UNUSED(max_bias); GGML_UNUSED(m0); GGML_UNUSED(m1);
318307
GGML_UNUSED(n_head_log2); GGML_UNUSED(logit_softcap);
319308
GGML_UNUSED(ne00); GGML_UNUSED(ne01); GGML_UNUSED(ne02); GGML_UNUSED(ne03);

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

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -349,8 +349,8 @@ 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); GGML_UNUSED(sinks);
353-
GGML_UNUSED(dst); GGML_UNUSED(dst_meta);
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);
354354
GGML_UNUSED(scale); GGML_UNUSED(max_bias); GGML_UNUSED(m0); GGML_UNUSED(m1);
355355
GGML_UNUSED(n_head_log2); GGML_UNUSED(logit_softcap);
356356
GGML_UNUSED(ne00); GGML_UNUSED(ne01); GGML_UNUSED(ne02); GGML_UNUSED(ne03);

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

Lines changed: 2 additions & 13 deletions
Original file line numberDiff line numberDiff line change
@@ -37,17 +37,6 @@ 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(Q); GGML_UNUSED(K); GGML_UNUSED(V); GGML_UNUSED(mask);
41-
GGML_UNUSED(dst); GGML_UNUSED(dst_meta); GGML_UNUSED(scale);
42-
GGML_UNUSED(max_bias); GGML_UNUSED(m0); GGML_UNUSED(m1);
43-
GGML_UNUSED(n_head_log2); GGML_UNUSED(logit_softcap);
44-
GGML_UNUSED(ne00); GGML_UNUSED(ne01); GGML_UNUSED(ne02);
45-
GGML_UNUSED(ne03); GGML_UNUSED(ne10); GGML_UNUSED(ne11);
46-
GGML_UNUSED(ne12); GGML_UNUSED(ne13); GGML_UNUSED(ne31); GGML_UNUSED(ne32); GGML_UNUSED(ne33);
47-
GGML_UNUSED(nb31); GGML_UNUSED(nb32); GGML_UNUSED(nb33); GGML_UNUSED(nb01); GGML_UNUSED(nb02);
48-
GGML_UNUSED(nb03); GGML_UNUSED(nb11); GGML_UNUSED(nb12);
49-
GGML_UNUSED(nb13); GGML_UNUSED(nb21); GGML_UNUSED(nb22);
50-
GGML_UNUSED(nb23);
5140
NO_DEVICE_CODE;
5241
return;
5342
}
@@ -346,8 +335,8 @@ static __global__ void flash_attn_vec_ext_f32(
346335
}
347336
#else
348337
GGML_UNUSED(Q); GGML_UNUSED(K); GGML_UNUSED(V); GGML_UNUSED(mask);
349-
GGML_UNUSED(dst); GGML_UNUSED(dst_meta); GGML_UNUSED(scale);
350-
GGML_UNUSED(max_bias); GGML_UNUSED(m0); GGML_UNUSED(m1);
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);
351340
GGML_UNUSED(n_head_log2); GGML_UNUSED(logit_softcap);
352341
GGML_UNUSED(ne00); GGML_UNUSED(ne01); GGML_UNUSED(ne02); GGML_UNUSED(ne03);
353342
GGML_UNUSED(ne10); GGML_UNUSED(ne11); GGML_UNUSED(ne12); GGML_UNUSED(ne13);

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

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -472,9 +472,9 @@ static __global__ void flash_attn_ext_f16(
472472
dst_meta[j_dst_unrolled] = dst_meta_val;
473473
}
474474
#else
475-
GGML_UNUSED(Q); GGML_UNUSED(K); GGML_UNUSED(V); GGML_UNUSED(mask); GGML_UNUSED(sinks);
476-
GGML_UNUSED(dst); GGML_UNUSED(dst_meta); GGML_UNUSED(scale);
477-
GGML_UNUSED(max_bias); GGML_UNUSED(m0); GGML_UNUSED(m1);
475+
GGML_UNUSED(Q); GGML_UNUSED(K); GGML_UNUSED(V); GGML_UNUSED(mask);
476+
GGML_UNUSED(sinks); GGML_UNUSED(KV_max); GGML_UNUSED(dst); GGML_UNUSED(dst_meta);
477+
GGML_UNUSED(scale); GGML_UNUSED(max_bias); GGML_UNUSED(m0); GGML_UNUSED(m1);
478478
GGML_UNUSED(n_head_log2); GGML_UNUSED(logit_softcap);
479479
GGML_UNUSED(ne00); GGML_UNUSED(ne01); GGML_UNUSED(ne02); GGML_UNUSED(ne03);
480480
GGML_UNUSED(ne10); GGML_UNUSED(ne11); GGML_UNUSED(ne12); GGML_UNUSED(ne13);

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

Lines changed: 0 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -17,7 +17,6 @@ static __global__ void mul_mat_f(
1717
#if !defined(GGML_USE_HIP) && !defined(GGML_USE_MUSA)
1818
typedef tile<16, 8, T> tile_A;
1919
typedef tile< 8, 8, T> tile_B;
20-
typedef tile<16, 8, float> tile_C;
2120

2221
constexpr int warp_size = ggml_cuda_get_physical_warp_size();
2322
constexpr int tile_k_padded = warp_size + 4;
@@ -151,7 +150,6 @@ static void mul_mat_f_cuda(
151150
cudaStream_t stream) {
152151
typedef tile<16, 8, T> tile_A;
153152
typedef tile< 8, 8, T> tile_B;
154-
typedef tile<16, 8, float> tile_C;
155153

156154
GGML_ASSERT(!ids && "mul_mat_id not implemented");
157155

@@ -352,9 +350,6 @@ void ggml_cuda_mul_mat_f(ggml_backend_cuda_context & ctx, const ggml_tensor * sr
352350
GGML_ASSERT(!ids || ids->nb[0] == ggml_type_size(ids->type));
353351
GGML_ASSERT( nb0 == ts_dst);
354352

355-
const int cc = ggml_cuda_info().devices[ggml_cuda_get_device()].cc;
356-
const enum ggml_prec prec = fast_fp16_available(cc) ? ggml_prec(dst->op_params[0]) : GGML_PREC_F32;
357-
358353
const float * src1_d = (const float *) src1->data;
359354
const int32_t * ids_d = ids ? (const int32_t *) ids->data : nullptr;
360355
float * dst_d = (float *) dst->data;

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

Lines changed: 3 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -2855,12 +2855,14 @@ static __device__ __forceinline__ void mmq_write_back_mma(
28552855
#else
28562856
typedef tile<16, 8, int> tile_C;
28572857
constexpr int rows_per_warp = 2 * granularity;
2858-
#endif
2858+
#endif // defined(AMD_MFMA_AVAILABLE)
28592859
constexpr int ntx = rows_per_warp/tile_C::I; // Number of x minitiles per warp.
28602860

28612861
const int i0 = (threadIdx.y / ntx) * (ntx*tile_C::I);
28622862
#if defined(TURING_MMA_AVAILABLE) || defined(AMD_MFMA_AVAILABLE)
28632863
static_assert(nwarps*tile_C::I == mmq_y, "nwarps*tile_C::I != mmq_y");
2864+
#else
2865+
GGML_UNUSED(nwarps);
28642866
#endif // defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE)
28652867

28662868
#pragma unroll

0 commit comments

Comments
 (0)