Compile bug: compilation warnings (clang) Introduced in #10558

[yeahdongcn]

Git commit

250d795

Operating systems

Linux

GGML backends

Musa

Problem description & steps to reproduce

The MUSA backend (compiled using Clang) generates warnings in /ggml/src/ggml-cuda/ssm-conv.cu.

First Bad Commit

250d795

Compile command

cmake -B build -DGGML_MUSA=ON -DMUSA_ARCHITECTURES="21"
cmake --build build --config Release

Relevant log output

/home/xiaodongye/ws/ggml/llama.cpp/ggml/src/ggml-cuda/ssm-conv.cu:12:55: warning: cast from 'const float *' to 'char *' drops const qualifier [-Wcast-qual]
    const float * x_block = (const float *) ((char *) src0 + bidx * src0_nb2 + bidy * split_d_inner * src0_nb1);
                                                      ^
/home/xiaodongye/ws/ggml/llama.cpp/ggml/src/ggml-cuda/ssm-conv.cu:13:55: warning: cast from 'const float *' to 'char *' drops const qualifier [-Wcast-qual]
    const float * w_block = (const float *) ((char *) src1 + bidy * split_d_inner * src1_nb1);
                                                      ^
/home/xiaodongye/ws/ggml/llama.cpp/ggml/src/ggml-cuda/ssm-conv.cu:5:47: warning: unused parameter 'src0_nb0' [-Wunused-parameter]
                                    const int src0_nb0, const int src0_nb1, const int src0_nb2, const int src1_nb1,
                                              ^
/home/xiaodongye/ws/ggml/llama.cpp/ggml/src/ggml-cuda/ssm-conv.cu:7:47: warning: unused parameter 'nc' [-Wunused-parameter]
                                    const int nc, const int ncs, const int nr, const int n_t, const int n_s) {
                                              ^
/home/xiaodongye/ws/ggml/llama.cpp/ggml/src/ggml-cuda/ssm-conv.cu:7:61: warning: unused parameter 'ncs' [-Wunused-parameter]
                                    const int nc, const int ncs, const int nr, const int n_t, const int n_s) {
                                                            ^
/home/xiaodongye/ws/ggml/llama.cpp/ggml/src/ggml-cuda/ssm-conv.cu:7:76: warning: unused parameter 'nr' [-Wunused-parameter]
                                    const int nc, const int ncs, const int nr, const int n_t, const int n_s) {
                                                                           ^
/home/xiaodongye/ws/ggml/llama.cpp/ggml/src/ggml-cuda/ssm-conv.cu:7:105: warning: unused parameter 'n_s' [-Wunused-parameter]
                                    const int nc, const int ncs, const int nr, const int n_t, const int n_s) {
                                                                                                        ^
/home/xiaodongye/ws/ggml/llama.cpp/ggml/src/ggml-cuda/ssm-conv.cu:58:55: warning: cast from 'const float *' to 'char *' drops const qualifier [-Wcast-qual]
    const float * x_block = (const float *) ((char *) src0 + bidx * src0_nb2 + bidy * split_d_inner * src0_nb1 +
                                                      ^
/home/xiaodongye/ws/ggml/llama.cpp/ggml/src/ggml-cuda/ssm-conv.cu:60:55: warning: cast from 'const float *' to 'char *' drops const qualifier [-Wcast-qual]
    const float * w_block = (const float *) ((char *) src1 + bidy * split_d_inner * src1_nb1);
                                                      ^
/home/xiaodongye/ws/ggml/llama.cpp/ggml/src/ggml-cuda/ssm-conv.cu:51:96: warning: unused parameter 'nc' [-Wunused-parameter]
                                               const int dst_nb1, const int dst_nb2, const int nc, const int ncs,
                                                                                               ^
/home/xiaodongye/ws/ggml/llama.cpp/ggml/src/ggml-cuda/ssm-conv.cu:51:110: warning: unused parameter 'ncs' [-Wunused-parameter]
                                               const int dst_nb1, const int dst_nb2, const int nc, const int ncs,
                                                                                                             ^
/home/xiaodongye/ws/ggml/llama.cpp/ggml/src/ggml-cuda/ssm-conv.cu:52:58: warning: unused parameter 'nr' [-Wunused-parameter]
                                               const int nr, const int n_t, const int n_s) {
                                                         ^
/home/xiaodongye/ws/ggml/llama.cpp/ggml/src/ggml-cuda/ssm-conv.cu:52:87: warning: unused parameter 'n_s' [-Wunused-parameter]
                                               const int nr, const int n_t, const int n_s) {
                                                                                      ^
/home/xiaodongye/ws/ggml/llama.cpp/ggml/src/ggml-cuda/ssm-conv.cu:12:55: warning: cast from 'const float *' to 'char *' drops const qualifier [-Wcast-qual]
    const float * x_block = (const float *) ((char *) src0 + bidx * src0_nb2 + bidy * split_d_inner * src0_nb1);
                                                      ^
/home/xiaodongye/ws/ggml/llama.cpp/ggml/src/ggml-cuda/ssm-conv.cu:108:13: note: in instantiation of function template specialization 'ssm_conv_f32<128UL, 4UL>' requested here
            ssm_conv_f32<threads, 4><<<blocks, threads, 0, stream>>>(src0, src1, src0_nb0, src0_nb1, src0_nb2, src1_nb1,
            ^
/home/xiaodongye/ws/ggml/llama.cpp/ggml/src/ggml-cuda/ssm-conv.cu:13:55: warning: cast from 'const float *' to 'char *' drops const qualifier [-Wcast-qual]
    const float * w_block = (const float *) ((char *) src1 + bidy * split_d_inner * src1_nb1);
                                                      ^
/home/xiaodongye/ws/ggml/llama.cpp/ggml/src/ggml-cuda/ssm-conv.cu:24:23: warning: comparison of integers of different signs: 'int' and 'unsigned long' [-Wsign-compare]
    for (int j = 0; j < d_conv; j++) {
                    ~ ^ ~~~~~~
/home/xiaodongye/ws/ggml/llama.cpp/ggml/src/ggml-cuda/ssm-conv.cu:32:31: warning: comparison of integers of different signs: 'int' and 'unsigned long' [-Wsign-compare]
            for (int j = 0; j < d_conv; j++) {
                            ~ ^ ~~~~~~
/home/xiaodongye/ws/ggml/llama.cpp/ggml/src/ggml-cuda/ssm-conv.cu:40:27: warning: comparison of integers of different signs: 'int' and 'unsigned long' [-Wsign-compare]
        for (int j = 0; j < d_conv; j++) {
                        ~ ^ ~~~~~~
/home/xiaodongye/ws/ggml/llama.cpp/ggml/src/ggml-cuda/ssm-conv.cu:58:55: warning: cast from 'const float *' to 'char *' drops const qualifier [-Wcast-qual]
    const float * x_block = (const float *) ((char *) src0 + bidx * src0_nb2 + bidy * split_d_inner * src0_nb1 +
                                                      ^
/home/xiaodongye/ws/ggml/llama.cpp/ggml/src/ggml-cuda/ssm-conv.cu:118:13: note: in instantiation of function template specialization 'ssm_conv_long_token_f32<128UL, 4UL, 32UL>' requested here
            ssm_conv_long_token_f32<threads, 4, split_n_t>
            ^
/home/xiaodongye/ws/ggml/llama.cpp/ggml/src/ggml-cuda/ssm-conv.cu:60:55: warning: cast from 'const float *' to 'char *' drops const qualifier [-Wcast-qual]
    const float * w_block = (const float *) ((char *) src1 + bidy * split_d_inner * src1_nb1);
                                                      ^
/home/xiaodongye/ws/ggml/llama.cpp/ggml/src/ggml-cuda/ssm-conv.cu:72:23: warning: comparison of integers of different signs: 'int' and 'unsigned long' [-Wsign-compare]
    for (int j = 0; j < d_conv; j++) {
                    ~ ^ ~~~~~~
/home/xiaodongye/ws/ggml/llama.cpp/ggml/src/ggml-cuda/ssm-conv.cu:77:23: warning: comparison of integers of different signs: 'int' and 'unsigned long' [-Wsign-compare]
    for (int i = 0; i < split_n_t; i++) {
                    ~ ^ ~~~~~~~~~
/home/xiaodongye/ws/ggml/llama.cpp/ggml/src/ggml-cuda/ssm-conv.cu:78:34: warning: comparison of integers of different signs: 'unsigned long' and 'const int' [-Wsign-compare]
        if (bidz * split_n_t + i < n_t) {
            ~~~~~~~~~~~~~~~~~~~~ ^ ~~~
/home/xiaodongye/ws/ggml/llama.cpp/ggml/src/ggml-cuda/ssm-conv.cu:82:35: warning: comparison of integers of different signs: 'int' and 'unsigned long' [-Wsign-compare]
                for (int j = 0; j < d_conv; j++) {
                                ~ ^ ~~~~~~
/home/xiaodongye/ws/ggml/llama.cpp/ggml/src/ggml-cuda/ssm-conv.cu:90:31: warning: comparison of integers of different signs: 'int' and 'unsigned long' [-Wsign-compare]
            for (int j = 0; j < d_conv; j++) {
                            ~ ^ ~~~~~~
25 warnings generated when compiling for mp_21.
/home/xiaodongye/ws/ggml/llama.cpp/ggml/src/ggml-cuda/ssm-conv.cu:12:55: warning: cast from 'const float *' to 'char *' drops const qualifier [-Wcast-qual]
    const float * x_block = (const float *) ((char *) src0 + bidx * src0_nb2 + bidy * split_d_inner * src0_nb1);
                                                      ^
/home/xiaodongye/ws/ggml/llama.cpp/ggml/src/ggml-cuda/ssm-conv.cu:13:55: warning: cast from 'const float *' to 'char *' drops const qualifier [-Wcast-qual]
    const float * w_block = (const float *) ((char *) src1 + bidy * split_d_inner * src1_nb1);
                                                      ^
/home/xiaodongye/ws/ggml/llama.cpp/ggml/src/ggml-cuda/ssm-conv.cu:5:47: warning: unused parameter 'src0_nb0' [-Wunused-parameter]
                                    const int src0_nb0, const int src0_nb1, const int src0_nb2, const int src1_nb1,
                                              ^
/home/xiaodongye/ws/ggml/llama.cpp/ggml/src/ggml-cuda/ssm-conv.cu:7:47: warning: unused parameter 'nc' [-Wunused-parameter]
                                    const int nc, const int ncs, const int nr, const int n_t, const int n_s) {
                                              ^
/home/xiaodongye/ws/ggml/llama.cpp/ggml/src/ggml-cuda/ssm-conv.cu:7:61: warning: unused parameter 'ncs' [-Wunused-parameter]
                                    const int nc, const int ncs, const int nr, const int n_t, const int n_s) {
                                                            ^
/home/xiaodongye/ws/ggml/llama.cpp/ggml/src/ggml-cuda/ssm-conv.cu:7:76: warning: unused parameter 'nr' [-Wunused-parameter]
                                    const int nc, const int ncs, const int nr, const int n_t, const int n_s) {
                                                                           ^
/home/xiaodongye/ws/ggml/llama.cpp/ggml/src/ggml-cuda/ssm-conv.cu:7:105: warning: unused parameter 'n_s' [-Wunused-parameter]
                                    const int nc, const int ncs, const int nr, const int n_t, const int n_s) {
                                                                                                        ^
/home/xiaodongye/ws/ggml/llama.cpp/ggml/src/ggml-cuda/ssm-conv.cu:58:55: warning: cast from 'const float *' to 'char *' drops const qualifier [-Wcast-qual]
    const float * x_block = (const float *) ((char *) src0 + bidx * src0_nb2 + bidy * split_d_inner * src0_nb1 +
                                                      ^
/home/xiaodongye/ws/ggml/llama.cpp/ggml/src/ggml-cuda/ssm-conv.cu:60:55: warning: cast from 'const float *' to 'char *' drops const qualifier [-Wcast-qual]
    const float * w_block = (const float *) ((char *) src1 + bidy * split_d_inner * src1_nb1);
                                                      ^
/home/xiaodongye/ws/ggml/llama.cpp/ggml/src/ggml-cuda/ssm-conv.cu:51:96: warning: unused parameter 'nc' [-Wunused-parameter]
                                               const int dst_nb1, const int dst_nb2, const int nc, const int ncs,
                                                                                               ^
/home/xiaodongye/ws/ggml/llama.cpp/ggml/src/ggml-cuda/ssm-conv.cu:51:110: warning: unused parameter 'ncs' [-Wunused-parameter]
                                               const int dst_nb1, const int dst_nb2, const int nc, const int ncs,
                                                                                                             ^
/home/xiaodongye/ws/ggml/llama.cpp/ggml/src/ggml-cuda/ssm-conv.cu:52:58: warning: unused parameter 'nr' [-Wunused-parameter]
                                               const int nr, const int n_t, const int n_s) {
                                                         ^
/home/xiaodongye/ws/ggml/llama.cpp/ggml/src/ggml-cuda/ssm-conv.cu:52:87: warning: unused parameter 'n_s' [-Wunused-parameter]
                                               const int nr, const int n_t, const int n_s) {
                                                                                      ^
/home/xiaodongye/ws/ggml/llama.cpp/ggml/src/ggml-cuda/ssm-conv.cu:12:55: warning: cast from 'const float *' to 'char *' drops const qualifier [-Wcast-qual]
    const float * x_block = (const float *) ((char *) src0 + bidx * src0_nb2 + bidy * split_d_inner * src0_nb1);
                                                      ^
/home/xiaodongye/ws/ggml/llama.cpp/ggml/src/ggml-cuda/ssm-conv.cu:108:13: note: in instantiation of function template specialization 'ssm_conv_f32<128UL, 4UL>' requested here
            ssm_conv_f32<threads, 4><<<blocks, threads, 0, stream>>>(src0, src1, src0_nb0, src0_nb1, src0_nb2, src1_nb1,
            ^
/home/xiaodongye/ws/ggml/llama.cpp/ggml/src/ggml-cuda/ssm-conv.cu:13:55: warning: cast from 'const float *' to 'char *' drops const qualifier [-Wcast-qual]
    const float * w_block = (const float *) ((char *) src1 + bidy * split_d_inner * src1_nb1);
                                                      ^
/home/xiaodongye/ws/ggml/llama.cpp/ggml/src/ggml-cuda/ssm-conv.cu:24:23: warning: comparison of integers of different signs: 'int' and 'unsigned long' [-Wsign-compare]
    for (int j = 0; j < d_conv; j++) {
                    ~ ^ ~~~~~~
/home/xiaodongye/ws/ggml/llama.cpp/ggml/src/ggml-cuda/ssm-conv.cu:32:31: warning: comparison of integers of different signs: 'int' and 'unsigned long' [-Wsign-compare]
            for (int j = 0; j < d_conv; j++) {
                            ~ ^ ~~~~~~
/home/xiaodongye/ws/ggml/llama.cpp/ggml/src/ggml-cuda/ssm-conv.cu:40:27: warning: comparison of integers of different signs: 'int' and 'unsigned long' [-Wsign-compare]
        for (int j = 0; j < d_conv; j++) {
                        ~ ^ ~~~~~~
/home/xiaodongye/ws/ggml/llama.cpp/ggml/src/ggml-cuda/ssm-conv.cu:58:55: warning: cast from 'const float *' to 'char *' drops const qualifier [-Wcast-qual]
    const float * x_block = (const float *) ((char *) src0 + bidx * src0_nb2 + bidy * split_d_inner * src0_nb1 +
                                                      ^
/home/xiaodongye/ws/ggml/llama.cpp/ggml/src/ggml-cuda/ssm-conv.cu:118:13: note: in instantiation of function template specialization 'ssm_conv_long_token_f32<128UL, 4UL, 32UL>' requested here
            ssm_conv_long_token_f32<threads, 4, split_n_t>
            ^
/home/xiaodongye/ws/ggml/llama.cpp/ggml/src/ggml-cuda/ssm-conv.cu:60:55: warning: cast from 'const float *' to 'char *' drops const qualifier [-Wcast-qual]
    const float * w_block = (const float *) ((char *) src1 + bidy * split_d_inner * src1_nb1);
                                                      ^
/home/xiaodongye/ws/ggml/llama.cpp/ggml/src/ggml-cuda/ssm-conv.cu:72:23: warning: comparison of integers of different signs: 'int' and 'unsigned long' [-Wsign-compare]
    for (int j = 0; j < d_conv; j++) {
                    ~ ^ ~~~~~~
/home/xiaodongye/ws/ggml/llama.cpp/ggml/src/ggml-cuda/ssm-conv.cu:77:23: warning: comparison of integers of different signs: 'int' and 'unsigned long' [-Wsign-compare]
    for (int i = 0; i < split_n_t; i++) {
                    ~ ^ ~~~~~~~~~
/home/xiaodongye/ws/ggml/llama.cpp/ggml/src/ggml-cuda/ssm-conv.cu:78:34: warning: comparison of integers of different signs: 'unsigned long' and 'const int' [-Wsign-compare]
        if (bidz * split_n_t + i < n_t) {
            ~~~~~~~~~~~~~~~~~~~~ ^ ~~~
/home/xiaodongye/ws/ggml/llama.cpp/ggml/src/ggml-cuda/ssm-conv.cu:82:35: warning: comparison of integers of different signs: 'int' and 'unsigned long' [-Wsign-compare]
                for (int j = 0; j < d_conv; j++) {
                                ~ ^ ~~~~~~
/home/xiaodongye/ws/ggml/llama.cpp/ggml/src/ggml-cuda/ssm-conv.cu:90:31: warning: comparison of integers of different signs: 'int' and 'unsigned long' [-Wsign-compare]
            for (int j = 0; j < d_conv; j++) {
                            ~ ^ ~~~~~~
25 warnings generated when compiling for host.
[ 18%] Building CXX object ggml/src/ggml-musa/CMakeFiles/ggml-musa.dir/__/ggml-cuda/ssm-scan.cu.o
/home/xiaodongye/ws/ggml/llama.cpp/ggml/src/ggml-cuda/ssm-scan.cu:28:56: warning: cast from 'const float *' to 'char *' drops const qualifier [-Wcast-qual]
    const float * s0_block = (const float *) ((char *) src0 + bidx * src0_nb2 + bidy * splitD * src0_nb1);
                                                       ^
/home/xiaodongye/ws/ggml/llama.cpp/ggml/src/ggml-cuda/ssm-scan.cu:29:56: warning: cast from 'const float *' to 'char *' drops const qualifier [-Wcast-qual]
    const float * x_block  = (const float *) ((char *) src1 + (bidx * src1_nb2) + bidy * splitD * sizeof(float));
                                                       ^
/home/xiaodongye/ws/ggml/llama.cpp/ggml/src/ggml-cuda/ssm-scan.cu:30:56: warning: cast from 'const float *' to 'char *' drops const qualifier [-Wcast-qual]
    const float * dt_block = (const float *) ((char *) src2 + (bidx * src2_nb2) + bidy * splitD * sizeof(float));
                                                       ^
/home/xiaodongye/ws/ggml/llama.cpp/ggml/src/ggml-cuda/ssm-scan.cu:31:56: warning: cast from 'const float *' to 'char *' drops const qualifier [-Wcast-qual]
    const float * A_block  = (const float *) ((char *) src3 + bidy * splitD * src3_nb1);
                                                       ^
/home/xiaodongye/ws/ggml/llama.cpp/ggml/src/ggml-cuda/ssm-scan.cu:32:56: warning: cast from 'const float *' to 'char *' drops const qualifier [-Wcast-qual]
    const float * B_block  = (const float *) ((char *) src4 + (bidx * src4_nb2));
                                                       ^
/home/xiaodongye/ws/ggml/llama.cpp/ggml/src/ggml-cuda/ssm-scan.cu:33:56: warning: cast from 'const float *' to 'char *' drops const qualifier [-Wcast-qual]
    const float * C_block  = (const float *) ((char *) src5 + (bidx * src5_nb2));
                                                       ^
/home/xiaodongye/ws/ggml/llama.cpp/ggml/src/ggml-cuda/ssm-scan.cu:12:68: warning: unused parameter 'src1_nb0' [-Wunused-parameter]
                 const int src0_nb1, const int src0_nb2, const int src1_nb0, const int src1_nb1, const int src1_nb2,
                                                                   ^
/home/xiaodongye/ws/ggml/llama.cpp/ggml/src/ggml-cuda/ssm-scan.cu:13:48: warning: unused parameter 'src2_nb0' [-Wunused-parameter]
                 const int src1_nb3, const int src2_nb0, const int src2_nb1, const int src2_nb2, const int src3_nb1,
                                               ^
/home/xiaodongye/ws/ggml/llama.cpp/ggml/src/ggml-cuda/ssm-scan.cu:15:54: warning: unused parameter 'D' [-Wunused-parameter]
                 float * __restrict__ dst, const int D, const int L, const int B) {
                                                     ^
/home/xiaodongye/ws/ggml/llama.cpp/ggml/src/ggml-cuda/ssm-scan.cu:15:80: warning: unused parameter 'B' [-Wunused-parameter]
                 float * __restrict__ dst, const int D, const int L, const int B) {
                                                                               ^
/home/xiaodongye/ws/ggml/llama.cpp/ggml/src/ggml-cuda/ssm-scan.cu:28:56: warning: cast from 'const float *' to 'char *' drops const qualifier [-Wcast-qual]
    const float * s0_block = (const float *) ((char *) src0 + bidx * src0_nb2 + bidy * splitD * src0_nb1);
                                                       ^
/home/xiaodongye/ws/ggml/llama.cpp/ggml/src/ggml-cuda/ssm-scan.cu:100:9: note: in instantiation of function template specialization 'ssm_scan_f32<128UL, 16UL>' requested here
        ssm_scan_f32<128, 16><<<blocks, threads, smem_size, stream>>>(
        ^
/home/xiaodongye/ws/ggml/llama.cpp/ggml/src/ggml-cuda/ssm-scan.cu:29:56: warning: cast from 'const float *' to 'char *' drops const qualifier [-Wcast-qual]
    const float * x_block  = (const float *) ((char *) src1 + (bidx * src1_nb2) + bidy * splitD * sizeof(float));
                                                       ^
/home/xiaodongye/ws/ggml/llama.cpp/ggml/src/ggml-cuda/ssm-scan.cu:30:56: warning: cast from 'const float *' to 'char *' drops const qualifier [-Wcast-qual]
    const float * dt_block = (const float *) ((char *) src2 + (bidx * src2_nb2) + bidy * splitD * sizeof(float));
                                                       ^
/home/xiaodongye/ws/ggml/llama.cpp/ggml/src/ggml-cuda/ssm-scan.cu:31:56: warning: cast from 'const float *' to 'char *' drops const qualifier [-Wcast-qual]
    const float * A_block  = (const float *) ((char *) src3 + bidy * splitD * src3_nb1);
                                                       ^
/home/xiaodongye/ws/ggml/llama.cpp/ggml/src/ggml-cuda/ssm-scan.cu:32:56: warning: cast from 'const float *' to 'char *' drops const qualifier [-Wcast-qual]
    const float * B_block  = (const float *) ((char *) src4 + (bidx * src4_nb2));
                                                       ^
/home/xiaodongye/ws/ggml/llama.cpp/ggml/src/ggml-cuda/ssm-scan.cu:33:56: warning: cast from 'const float *' to 'char *' drops const qualifier [-Wcast-qual]
    const float * C_block  = (const float *) ((char *) src5 + (bidx * src5_nb2));
                                                       ^
/home/xiaodongye/ws/ggml/llama.cpp/ggml/src/ggml-cuda/ssm-scan.cu:49:27: warning: comparison of integers of different signs: 'int' and 'unsigned long' [-Wsign-compare]
        for (int i = 0; i < splitD / 4; i += 2) {
                        ~ ^ ~~~~~~~~~~
/home/xiaodongye/ws/ggml/llama.cpp/ggml/src/ggml-cuda/ssm-scan.cu:57:27: warning: comparison of integers of different signs: 'int' and 'unsigned long' [-Wsign-compare]
        for (int i = 0; i < splitD / 4; i += 2) {
                        ~ ^ ~~~~~~~~~~
/home/xiaodongye/ws/ggml/llama.cpp/ggml/src/ggml-cuda/ssm-scan.cu:73:27: warning: comparison of integers of different signs: 'int' and 'unsigned long' [-Wsign-compare]
        for (int j = 0; j < N; j++) {
                        ~ ^ ~
19 warnings generated when compiling for mp_21.
/home/xiaodongye/ws/ggml/llama.cpp/ggml/src/ggml-cuda/ssm-scan.cu:28:56: warning: cast from 'const float *' to 'char *' drops const qualifier [-Wcast-qual]
    const float * s0_block = (const float *) ((char *) src0 + bidx * src0_nb2 + bidy * splitD * src0_nb1);
                                                       ^
/home/xiaodongye/ws/ggml/llama.cpp/ggml/src/ggml-cuda/ssm-scan.cu:29:56: warning: cast from 'const float *' to 'char *' drops const qualifier [-Wcast-qual]
    const float * x_block  = (const float *) ((char *) src1 + (bidx * src1_nb2) + bidy * splitD * sizeof(float));
                                                       ^
/home/xiaodongye/ws/ggml/llama.cpp/ggml/src/ggml-cuda/ssm-scan.cu:30:56: warning: cast from 'const float *' to 'char *' drops const qualifier [-Wcast-qual]
    const float * dt_block = (const float *) ((char *) src2 + (bidx * src2_nb2) + bidy * splitD * sizeof(float));
                                                       ^
/home/xiaodongye/ws/ggml/llama.cpp/ggml/src/ggml-cuda/ssm-scan.cu:31:56: warning: cast from 'const float *' to 'char *' drops const qualifier [-Wcast-qual]
    const float * A_block  = (const float *) ((char *) src3 + bidy * splitD * src3_nb1);
                                                       ^
/home/xiaodongye/ws/ggml/llama.cpp/ggml/src/ggml-cuda/ssm-scan.cu:32:56: warning: cast from 'const float *' to 'char *' drops const qualifier [-Wcast-qual]
    const float * B_block  = (const float *) ((char *) src4 + (bidx * src4_nb2));
                                                       ^
/home/xiaodongye/ws/ggml/llama.cpp/ggml/src/ggml-cuda/ssm-scan.cu:33:56: warning: cast from 'const float *' to 'char *' drops const qualifier [-Wcast-qual]
    const float * C_block  = (const float *) ((char *) src5 + (bidx * src5_nb2));
                                                       ^
/home/xiaodongye/ws/ggml/llama.cpp/ggml/src/ggml-cuda/ssm-scan.cu:12:68: warning: unused parameter 'src1_nb0' [-Wunused-parameter]
                 const int src0_nb1, const int src0_nb2, const int src1_nb0, const int src1_nb1, const int src1_nb2,
                                                                   ^
/home/xiaodongye/ws/ggml/llama.cpp/ggml/src/ggml-cuda/ssm-scan.cu:13:48: warning: unused parameter 'src2_nb0' [-Wunused-parameter]
                 const int src1_nb3, const int src2_nb0, const int src2_nb1, const int src2_nb2, const int src3_nb1,
                                               ^
/home/xiaodongye/ws/ggml/llama.cpp/ggml/src/ggml-cuda/ssm-scan.cu:15:54: warning: unused parameter 'D' [-Wunused-parameter]
                 float * __restrict__ dst, const int D, const int L, const int B) {
                                                     ^
/home/xiaodongye/ws/ggml/llama.cpp/ggml/src/ggml-cuda/ssm-scan.cu:15:80: warning: unused parameter 'B' [-Wunused-parameter]
                 float * __restrict__ dst, const int D, const int L, const int B) {
                                                                               ^
/home/xiaodongye/ws/ggml/llama.cpp/ggml/src/ggml-cuda/ssm-scan.cu:28:56: warning: cast from 'const float *' to 'char *' drops const qualifier [-Wcast-qual]
    const float * s0_block = (const float *) ((char *) src0 + bidx * src0_nb2 + bidy * splitD * src0_nb1);
                                                       ^
/home/xiaodongye/ws/ggml/llama.cpp/ggml/src/ggml-cuda/ssm-scan.cu:100:9: note: in instantiation of function template specialization 'ssm_scan_f32<128UL, 16UL>' requested here
        ssm_scan_f32<128, 16><<<blocks, threads, smem_size, stream>>>(
        ^
/home/xiaodongye/ws/ggml/llama.cpp/ggml/src/ggml-cuda/ssm-scan.cu:29:56: warning: cast from 'const float *' to 'char *' drops const qualifier [-Wcast-qual]
    const float * x_block  = (const float *) ((char *) src1 + (bidx * src1_nb2) + bidy * splitD * sizeof(float));
                                                       ^
/home/xiaodongye/ws/ggml/llama.cpp/ggml/src/ggml-cuda/ssm-scan.cu:30:56: warning: cast from 'const float *' to 'char *' drops const qualifier [-Wcast-qual]
    const float * dt_block = (const float *) ((char *) src2 + (bidx * src2_nb2) + bidy * splitD * sizeof(float));
                                                       ^
/home/xiaodongye/ws/ggml/llama.cpp/ggml/src/ggml-cuda/ssm-scan.cu:31:56: warning: cast from 'const float *' to 'char *' drops const qualifier [-Wcast-qual]
    const float * A_block  = (const float *) ((char *) src3 + bidy * splitD * src3_nb1);
                                                       ^
/home/xiaodongye/ws/ggml/llama.cpp/ggml/src/ggml-cuda/ssm-scan.cu:32:56: warning: cast from 'const float *' to 'char *' drops const qualifier [-Wcast-qual]
    const float * B_block  = (const float *) ((char *) src4 + (bidx * src4_nb2));
                                                       ^
/home/xiaodongye/ws/ggml/llama.cpp/ggml/src/ggml-cuda/ssm-scan.cu:33:56: warning: cast from 'const float *' to 'char *' drops const qualifier [-Wcast-qual]
    const float * C_block  = (const float *) ((char *) src5 + (bidx * src5_nb2));
                                                       ^
/home/xiaodongye/ws/ggml/llama.cpp/ggml/src/ggml-cuda/ssm-scan.cu:49:27: warning: comparison of integers of different signs: 'int' and 'unsigned long' [-Wsign-compare]
        for (int i = 0; i < splitD / 4; i += 2) {
                        ~ ^ ~~~~~~~~~~
/home/xiaodongye/ws/ggml/llama.cpp/ggml/src/ggml-cuda/ssm-scan.cu:57:27: warning: comparison of integers of different signs: 'int' and 'unsigned long' [-Wsign-compare]
        for (int i = 0; i < splitD / 4; i += 2) {
                        ~ ^ ~~~~~~~~~~
/home/xiaodongye/ws/ggml/llama.cpp/ggml/src/ggml-cuda/ssm-scan.cu:73:27: warning: comparison of integers of different signs: 'int' and 'unsigned long' [-Wsign-compare]
        for (int j = 0; j < N; j++) {
                        ~ ^ ~
19 warnings generated when compiling for host.

[yeahdongcn]

@A3shTnT

This issue appears to have been introduced in: #10558

[A3shTnT]

I'll fix it later this week.

[A3shTnT]

@A3shTnT

This issue appears to have been introduced in: #10558

Is this a special case for Musa? It seems that nvcc does not report these warnings by default.

[yeahdongcn]

@A3shTnT
This issue appears to have been introduced in: #10558

Is this a special case for Musa? It seems that nvcc does not report these warnings by default.

This is likely compiler-related, as both MUSA and HIP backends use Clang, which reports these warnings, whereas NVCC does not by default.

[A3shTnT]

@A3shTnT
This issue appears to have been introduced in: #10558

Is this a special case for Musa? It seems that nvcc does not report these warnings by default.

This is likely compiler-related, as both MUSA and HIP backends use Clang, which reports these warnings, whereas NVCC does not by default.

Since I have never used Musa or HIP before, could you please compile my branch and check if there are any warnings before I submit a fix PR? If there are no issues, I will submit a new fix PR.

[yeahdongcn]

You can easily verify this in a Docker container using the following commands:

# Assuming `$PWD` is the root of your `llama.cpp` repository
docker run -it -v $PWD:/ws -w /ws mthreads/musa:rc3.1.1-devel-ubuntu22.04 bash    
apt update -y && apt install -y cmake ccache git
git config --global --add safe.directory /ws
cmake -B build -DGGML_MUSA=ON -DMUSA_ARCHITECTURES="21"
cmake --build build --config Release

BTW, llama.cpp CI now includes a MUSA build as well.