Skip to content

Commit 7a53018

Browse files
NexesenexNexesenex
committed
Revert "musa: fix build warnings (ggml-org#15258)"
This reverts commit a094f38.
1 parent 68d3134 commit 7a53018

File tree

10 files changed

+54
-32
lines changed

10 files changed

+54
-32
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 = *(const int32_t *) ((const char *) src2 + i1*sizeof(int32_t) + i2*nb21);
14+
const int i11 = *(int32_t *) ((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 *)((const char *)src0 + i1*nb01 + i2*nb02);
21-
const float * src1_row = (const float *)((const char *)src1 + i11*nb11);
20+
const float * src0_row = (const float *)((char *)src0 + i1*nb01 + i2*nb02);
21+
const float * src1_row = (const float *)((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: 5 additions & 7 deletions
Original file line numberDiff line numberDiff line change
@@ -1237,12 +1237,10 @@ 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(sinks_f);
1241-
GGML_UNUSED(dstk); GGML_UNUSED(dstk_fixup);
1240+
GGML_UNUSED(mask_h2); GGML_UNUSED(dstk); GGML_UNUSED(dstk_fixup);
12421241
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);
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);
12461244
GGML_UNUSED(jt); GGML_UNUSED(kb0_start); GGML_UNUSED(kb0_stop);
12471245
NO_DEVICE_CODE;
12481246
#endif // TURING_MMA_AVAILABLE
@@ -1397,8 +1395,8 @@ static __global__ void flash_attn_ext_f16(
13971395
(Q_f2, K_h2, V_h2, mask_h2, sinks_f, dstk, dst_meta, scale, slope, logit_softcap,
13981396
ne01, ne02, stride_Q1, stride_Q2, stride_K, stride_V, stride_mask, jt, kb0_start_kernel, kb0_stop_kernel);
13991397
#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);
1398+
GGML_UNUSED(Q); GGML_UNUSED(K); GGML_UNUSED(V); GGML_UNUSED(mask); GGML_UNUSED(sinks);
1399+
GGML_UNUSED(dst); GGML_UNUSED(dst_meta);
14021400
GGML_UNUSED(scale); GGML_UNUSED(max_bias); GGML_UNUSED(m0); GGML_UNUSED(m1);
14031401
GGML_UNUSED(n_head_log2); GGML_UNUSED(logit_softcap);
14041402
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);
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);
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);
305305
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);
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);
313313
NO_DEVICE_CODE;
314314
#endif // defined(FLASH_ATTN_AVAILABLE) && defined(FP16_AVAILABLE)
315315
}

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

Lines changed: 12 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -38,6 +38,17 @@ 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);
4152
NO_DEVICE_CODE;
4253
return;
4354
}
@@ -302,7 +313,7 @@ static __global__ void flash_attn_tile_ext_f32(
302313
}
303314
#else
304315
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);
316+
GGML_UNUSED(dst); GGML_UNUSED(dst_meta);
306317
GGML_UNUSED(scale); GGML_UNUSED(max_bias); GGML_UNUSED(m0); GGML_UNUSED(m1);
307318
GGML_UNUSED(n_head_log2); GGML_UNUSED(logit_softcap);
308319
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);
353-
GGML_UNUSED(sinks); GGML_UNUSED(KV_max); GGML_UNUSED(dst); GGML_UNUSED(dst_meta);
352+
GGML_UNUSED(Q); GGML_UNUSED(K); GGML_UNUSED(V); GGML_UNUSED(mask); GGML_UNUSED(sinks);
353+
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: 13 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -37,6 +37,17 @@ 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);
4051
NO_DEVICE_CODE;
4152
return;
4253
}
@@ -335,8 +346,8 @@ static __global__ void flash_attn_vec_ext_f32(
335346
}
336347
#else
337348
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);
349+
GGML_UNUSED(dst); GGML_UNUSED(dst_meta); GGML_UNUSED(scale);
350+
GGML_UNUSED(max_bias); GGML_UNUSED(m0); GGML_UNUSED(m1);
340351
GGML_UNUSED(n_head_log2); GGML_UNUSED(logit_softcap);
341352
GGML_UNUSED(ne00); GGML_UNUSED(ne01); GGML_UNUSED(ne02); GGML_UNUSED(ne03);
342353
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
@@ -471,9 +471,9 @@ 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);
474+
GGML_UNUSED(Q); GGML_UNUSED(K); GGML_UNUSED(V); GGML_UNUSED(mask); GGML_UNUSED(sinks);
475+
GGML_UNUSED(dst); GGML_UNUSED(dst_meta); GGML_UNUSED(scale);
476+
GGML_UNUSED(max_bias); GGML_UNUSED(m0); GGML_UNUSED(m1);
477477
GGML_UNUSED(n_head_log2); GGML_UNUSED(logit_softcap);
478478
GGML_UNUSED(ne00); GGML_UNUSED(ne01); GGML_UNUSED(ne02); GGML_UNUSED(ne03);
479479
GGML_UNUSED(ne10); GGML_UNUSED(ne11); GGML_UNUSED(ne12); GGML_UNUSED(ne13);

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

Lines changed: 4 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -151,6 +151,7 @@ static void mul_mat_f_cuda(
151151
cudaStream_t stream) {
152152
typedef tile<16, 8, T> tile_A;
153153
typedef tile< 8, 8, T> tile_B;
154+
typedef tile<16, 8, float> tile_C;
154155

155156
GGML_ASSERT(!ids && "mul_mat_id not implemented");
156157

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

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+
354358
const float * src1_d = (const float *) src1->data;
355359
const int32_t * ids_d = ids ? (const int32_t *) ids->data : nullptr;
356360
float * dst_d = (float *) dst->data;

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

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

28622862
const int i0 = (threadIdx.y / ntx) * (ntx*tile_C::I);
28632863
#if defined(TURING_MMA_AVAILABLE) || defined(AMD_MFMA_AVAILABLE)
28642864
static_assert(nwarps*tile_C::I == mmq_y, "nwarps*tile_C::I != mmq_y");
2865-
#else
2866-
GGML_UNUSED(nwarps);
28672865
#endif // defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE)
28682866

28692867
#pragma unroll

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

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -39,7 +39,7 @@ static __global__ void reduce_rows_f32(const float * __restrict__ x, float * __r
3939
}
4040
__syncthreads();
4141
sum = 0.0f;
42-
if (lane_id < (static_cast<int>(blockDim.x) / WARP_SIZE)) {
42+
if (lane_id < (blockDim.x / WARP_SIZE)) {
4343
sum = s_sum[lane_id];
4444
}
4545
sum = warp_reduce_sum(sum);

0 commit comments

Comments
 (0)