[SYCL] fallback mmvq

[airMeng]

• I have read the contributing guidelines
• Self-reported review complexity:
  • Low
  • Medium
  • High

There is a bug in SYCL MMVQ implementation, stuck when evaluate accuracy. fallback mmvq to avoid the stuck

./bin/llama-perplexity  --hellaswag -m ~/Meta-Llama-3-8B-Instruct-Q4_K_S.gguf -s 0 -ngl 99 --hellaswag-tasks 40 -f ../hellaswag_val_full.txt

[qnixsynapse]

I am getting performance regression with this PR ( from 35 tokens/sec to 4 tokens/sec on Arc A750 from the latest master branch) Reverting this fixes the issue. Subscribing myself just in case for checking out later.

[airMeng]

I am getting performance regression with this PR ( from 35 tokens/sec to 4 tokens/sec on Arc A750 from the latest master branch) Reverting this fixes the issue. Subscribing myself just in case for checking out later.

@qnixsynapse what the command you are using?

[qnixsynapse]

This:

build/bin/llama-server -c 8192 -t 4 -m meta-llama-3.1-8b-instruct-iq4_xs-imat.gguf --no-mmap -ngl 99 -a "LLaMA" -b 128 --port 8080 --log-disable -s 0 -mg 0 -sm none --log-format text -dt 0.1 --metrics

I haven't fully tested it yet, just identified the offending patch and reverted.

[qnixsynapse]

This is still causing performance regression on quantized models (esp iq4_xs) here.

stuck when evaluate accuracy. fallback mmvq to avoid the stuck

Can you please elaborate on why this is necessary? If Nvidia is unaffected by it, we can make it Nvidia exclusive.

Here is what I thought of doing:

bool use_mul_mat_vec_q =  ggml_is_quantized(src0->type)
       && src1->type == GGML_TYPE_F32 && dst->type == GGML_TYPE_F32
       && src1->ne[1] <= MMVQ_MAX_BATCH_SIZE;
         
   
 if (ctx.stream()->get_backend() == sycl::backend::ext_oneapi_cuda) {
            use_mul_mat_vec_q = use_mul_mat_vec_q && (src1->ne[1] > MMVQ_MIN_BATCH_SIZE);
  }

[qnixsynapse]

Update: With this PR, the boolean variable, "use_mul_mat_vec_q" is always false:

But reverting this PR, the value changes to true sometimes:

cc : @NeoZhangJianyu Please look at this.

[airMeng]

@qnixsynapse sorry for the carelessness, could you open a PR to revert it?

[qnixsynapse]

@airMeng Sure, no problem.