[SYCL] Support newer non linear quantization

[qnixsynapse]

cc: @abhilash1910 @airMeng

Also important!

[airMeng]

we haven't supported iq3_xxs, but you can try to implement yourself based on CUDA code. Feel free to ping us here if you encounter any issues

llama.cpp/ggml-sycl.cpp

Lines 14970 to 14972 in 15499eb

	if (a->type == GGML_TYPE_IQ3_XXS) {
	return false;
	}

[qnixsynapse]

I see. Thanks. Should I close this then?

[airMeng]

Also important!

Thank you for informing!

I see. Thanks. Should I close this then?

up to you

[qnixsynapse]

I am reopening this. I may/may not work on this since I am not a C person. I wish I could set the label to "help-wanted".

[qnixsynapse]

@ggerganov Thank you!

[airMeng]

#5862

Can you try this PR and report the results?

[qnixsynapse]

Got a ggml assert while trying to test an IQ3_XXS quantized gguf:

...
...
get_memory_info: [warning] ext_intel_free_memory is not supported (export/set ZES_ENABLE_SYSMAN=1 to support), use total memory as free memory
...
...
GGML_ASSERT: llama.cpp/ggml-sycl.cpp:14097: to_fp16_sycl != nullptr

The other non IQ quants looking good although I didn't understand the warning that I am getting about ext_intel_free_memory.

[airMeng]

what model are you using? can you share a link?

[qnixsynapse]

This time I used Eris 7B to test IQ3.

And for linear quants I used Mistral 7B.

[NeoZhangJianyu]

Got a ggml assert while trying to test an IQ3_XXS quantized gguf:

...
...
get_memory_info: [warning] ext_intel_free_memory is not supported (export/set ZES_ENABLE_SYSMAN=1 to support), use total memory as free memory
...
...
GGML_ASSERT: llama.cpp/ggml-sycl.cpp:14097: to_fp16_sycl != nullptr

The other non IQ quants looking good although I didn't understand the warning that I am getting about ext_intel_free_memory.

export/set ZES_ENABLE_SYSMAN=1 to support ext_intel_free_memory() to get free memory.
It's not important if your GPU is not shared.

[abhilash1910]

This time I used Eris 7B to test IQ3.

And for linear quants I used Mistral 7B.

@akarshanbiswas Could you try 3.75 bit /1.5 bit (#5886 )and let us know for Erin model.

[qnixsynapse]

@abhilash1910 Failed with main: ggml-sycl.cpp:2968: dpct::detail::device_memory<const unsigned int, dpct::global, 1>::device_memory(const sycl::range<Dimension> &, std::initializer_list<value_t> &&) [T = const unsigned int, Memory = dpct::global, Dimension = 1]: Assertion init_list.size() <= in_range.size()' failed.`

Edit: Looks like everything is broken.

Edit 2: Confirmed that PR is the culprit!

[NeoZhangJianyu]

mistral-7b-instruct-v0.2.Q4_K_M.gguf is passed by other test.
Please try with this model.

[qnixsynapse]

The regression was fixed after grid space increment.

[NeoZhangJianyu]

I have asked the developer to fix it.

[qnixsynapse]

Thank you.

[airMeng]

shall be fixed in #6521

[qnixsynapse]

Awesome!! Really thank you everyone for all the hard work. I will try testing a model today with this patch (hopefully)!.

[qnixsynapse]

Update: Just tested a iq4_XXS gemma 7B(8.5B) 1.1 model which was released recently and it is working great!!!

Closing this issue as solved for now!!

[abhilash1910]

Thanks for confirming @akarshanbiswas .