Skip to content

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

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Jump to bottom
Closed
yeahdongcn opened this issue Apr 1, 2025 · 6 comments
Closed

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

yeahdongcn opened this issue Apr 1, 2025 · 6 comments
Labels
bug Something isn't working

Comments

@yeahdongcn
Copy link
Contributor

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
Copy link
Contributor Author

@A3shTnT

This issue appears to have been introduced in: #10558

@A3shTnT
Copy link
Contributor

A3shTnT commented Apr 1, 2025

I'll fix it later this week.

@CISC CISC added bug Something isn't working and removed bug-unconfirmed labels Apr 1, 2025
@A3shTnT
Copy link
Contributor

A3shTnT commented Apr 1, 2025

@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
Copy link
Contributor Author

@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
Copy link
Contributor

A3shTnT commented Apr 1, 2025

@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
Copy link
Contributor Author

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.

This was referenced Apr 2, 2025
Closed
Merged
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Assignees
No one assigned
Labels
bug Something isn't working
Projects
None yet
Milestone
No milestone
Development

No branches or pull requests
3 participants
@CISC @yeahdongcn @A3shTnT