diff --git a/ggml/src/ggml-opencl/CMakeLists.txt b/ggml/src/ggml-opencl/CMakeLists.txt index 046e42acd..f905370a3 100644 --- a/ggml/src/ggml-opencl/CMakeLists.txt +++ b/ggml/src/ggml-opencl/CMakeLists.txt @@ -124,6 +124,12 @@ set(GGML_OPENCL_KERNELS gemm_noshuffle_iq1_s_f32 gemm_noshuffle_iq1_s_q8_1_dp4a gemv_noshuffle_iq1_s_f32 + gemm_noshuffle_iq3_s_f32 + gemm_noshuffle_iq3_s_q8_1_dp4a + gemv_noshuffle_iq3_s_f32 + gemm_noshuffle_iq3_xxs_f32 + gemm_noshuffle_iq3_xxs_q8_1_dp4a + gemv_noshuffle_iq3_xxs_f32 gemm_noshuffle_q3_k_q8_1_dp4a gemv_noshuffle_q3_k_f32 mul_mv_q2_k_f32 @@ -150,6 +156,14 @@ set(GGML_OPENCL_KERNELS mul_mv_iq1_s_f32 mul_mv_iq1_s_f32_flat mul_mm_iq1_s_f32_l4_lm + iq3_xxs_grid + mul_mv_iq3_xxs_f32 + mul_mv_iq3_xxs_f32_flat + mul_mm_iq3_xxs_f32_l4_lm + iq3_s_grid + mul_mv_iq3_s_f32 + mul_mv_iq3_s_f32_flat + mul_mm_iq3_s_f32_l4_lm mul_mv_q3_k_f32_flat mul_mm_q3_k_f32_l4_lm mul_mv_q2_k_f32_flat diff --git a/ggml/src/ggml-opencl/ggml-opencl.cpp b/ggml/src/ggml-opencl/ggml-opencl.cpp index 5b05c9b36..6cfd0c0f9 100644 --- a/ggml/src/ggml-opencl/ggml-opencl.cpp +++ b/ggml/src/ggml-opencl/ggml-opencl.cpp @@ -1051,6 +1051,28 @@ struct ggml_backend_opencl_context { ggml_cl_buffer iq1_s_grid_buffer; cl_mem iq1_s_grid_image = nullptr; cl_kernel kernel_mul_mv_iq1_s_f32; + cl_kernel kernel_convert_block_iq3_xxs, kernel_restore_block_iq3_xxs; + cl_kernel kernel_mul_mv_iq3_xxs_f32_flat; + cl_kernel kernel_mul_mm_iq3_xxs_f32_l4_lm; + cl_kernel kernel_gemv_noshuffle_iq3_xxs_f32 = nullptr; + cl_kernel kernel_gemm_noshuffle_iq3_xxs_f32 = nullptr; + cl_kernel kernel_gemm_noshuffle_iq3_xxs_q8_1_dp4a = nullptr; + cl_kernel kernel_gemm_noshuffle_iq3_xxs_q8_1_dp4a_narrow = nullptr; + int iq3_xxs_mv_nsg = 4; + ggml_cl_buffer iq3_xxs_grid_buffer; + cl_mem iq3_xxs_grid_image = nullptr; + cl_kernel kernel_mul_mv_iq3_xxs_f32; + cl_kernel kernel_convert_block_iq3_s, kernel_restore_block_iq3_s; + cl_kernel kernel_mul_mv_iq3_s_f32_flat; + cl_kernel kernel_mul_mm_iq3_s_f32_l4_lm; + cl_kernel kernel_gemv_noshuffle_iq3_s_f32 = nullptr; + cl_kernel kernel_gemm_noshuffle_iq3_s_f32 = nullptr; + cl_kernel kernel_gemm_noshuffle_iq3_s_q8_1_dp4a = nullptr; + cl_kernel kernel_gemm_noshuffle_iq3_s_q8_1_dp4a_narrow = nullptr; + int iq3_s_mv_nsg = 4; + ggml_cl_buffer iq3_s_grid_buffer; + cl_mem iq3_s_grid_image = nullptr; + cl_kernel kernel_mul_mv_iq3_s_f32; cl_kernel kernel_mul_mv_q4_K_f32; cl_kernel kernel_mul_mv_q4_K_f32_flat; cl_kernel kernel_mul_mv_q5_K_f32; @@ -1429,6 +1451,12 @@ struct ggml_backend_opencl_context { if (iq1_s_grid_image) { CL_CHECK(clReleaseMemObject(iq1_s_grid_image)); } + if (iq3_xxs_grid_image) { + CL_CHECK(clReleaseMemObject(iq3_xxs_grid_image)); + } + if (iq3_s_grid_image) { + CL_CHECK(clReleaseMemObject(iq3_s_grid_image)); + } if (iq2_xxs_grid_image) { CL_CHECK(clReleaseMemObject(iq2_xxs_grid_image)); } @@ -2130,6 +2158,10 @@ static void load_cl_kernels(ggml_backend_opencl_context *backend_ctx) { CL_CHECK((backend_ctx->kernel_restore_block_iq1_s = clCreateKernel(backend_ctx->program_cvt, "kernel_restore_block_iq1_s", &err), err)); CL_CHECK((backend_ctx->kernel_convert_block_iq1_m = clCreateKernel(backend_ctx->program_cvt, "kernel_convert_block_iq1_m", &err), err)); CL_CHECK((backend_ctx->kernel_restore_block_iq1_m = clCreateKernel(backend_ctx->program_cvt, "kernel_restore_block_iq1_m", &err), err)); + CL_CHECK((backend_ctx->kernel_convert_block_iq3_xxs = clCreateKernel(backend_ctx->program_cvt, "kernel_convert_block_iq3_xxs", &err), err)); + CL_CHECK((backend_ctx->kernel_restore_block_iq3_xxs = clCreateKernel(backend_ctx->program_cvt, "kernel_restore_block_iq3_xxs", &err), err)); + CL_CHECK((backend_ctx->kernel_convert_block_iq3_s = clCreateKernel(backend_ctx->program_cvt, "kernel_convert_block_iq3_s", &err), err)); + CL_CHECK((backend_ctx->kernel_restore_block_iq3_s = clCreateKernel(backend_ctx->program_cvt, "kernel_restore_block_iq3_s", &err), err)); CL_CHECK((backend_ctx->kernel_convert_block_q3_K = clCreateKernel(backend_ctx->program_cvt, "kernel_convert_block_q3_K", &err), err)); CL_CHECK((backend_ctx->kernel_restore_block_q3_K = clCreateKernel(backend_ctx->program_cvt, "kernel_restore_block_q3_K", &err), err)); CL_CHECK((backend_ctx->kernel_convert_block_q2_K = clCreateKernel(backend_ctx->program_cvt, "kernel_convert_block_q2_K", &err), err)); @@ -2894,6 +2926,112 @@ static void load_cl_kernels(ggml_backend_opencl_context *backend_ctx) { GGML_LOG_CONT("."); } +#ifdef GGML_OPENCL_EMBED_KERNELS + const std::string iq3_s_grid_src { + #include "iq3_s_grid.cl.h" + }; +#else + const std::string iq3_s_grid_src = read_file("iq3_s_grid.cl"); +#endif + + // mul_mv_iq3_s_f32_flat + { +#ifdef GGML_OPENCL_EMBED_KERNELS + const std::string kernel_src { + #include "mul_mv_iq3_s_f32_flat.cl.h" + }; +#else + const std::string kernel_src = read_file("mul_mv_iq3_s_f32_flat.cl"); +#endif + cl_program prog = build_program_from_source(backend_ctx, (iq3_s_grid_src + kernel_src).c_str(), compile_opts); + CL_CHECK((backend_ctx->kernel_mul_mv_iq3_s_f32_flat = clCreateKernel(prog, "kernel_mul_mv_iq3_s_f32_flat", &err), err)); + CL_CHECK(clReleaseProgram(prog)); + GGML_LOG_CONT("."); + } + + // mul_mm_iq3_s_f32_l4_lm + { +#ifdef GGML_OPENCL_EMBED_KERNELS + const std::string kernel_src { + #include "mul_mm_iq3_s_f32_l4_lm.cl.h" + }; +#else + const std::string kernel_src = read_file("mul_mm_iq3_s_f32_l4_lm.cl"); +#endif + cl_program prog = build_program_from_source(backend_ctx, (iq3_s_grid_src + kernel_src).c_str(), compile_opts); + CL_CHECK((backend_ctx->kernel_mul_mm_iq3_s_f32_l4_lm = clCreateKernel(prog, "kernel_mul_mm_iq3_s_f32_l4_lm", &err), err)); + CL_CHECK(clReleaseProgram(prog)); + GGML_LOG_CONT("."); + } + + // mul_mv_iq3_s_f32 + { +#ifdef GGML_OPENCL_EMBED_KERNELS + const std::string kernel_src { + #include "mul_mv_iq3_s_f32.cl.h" + }; +#else + const std::string kernel_src = read_file("mul_mv_iq3_s_f32.cl"); +#endif + cl_program prog = build_program_from_source(backend_ctx, (iq3_s_grid_src + kernel_src).c_str(), compile_opts); + CL_CHECK((backend_ctx->kernel_mul_mv_iq3_s_f32 = clCreateKernel(prog, "kernel_mul_mv_iq3_s_f32", &err), err)); + CL_CHECK(clReleaseProgram(prog)); + GGML_LOG_CONT("."); + } + +#ifdef GGML_OPENCL_EMBED_KERNELS + const std::string iq3_xxs_grid_src { + #include "iq3_xxs_grid.cl.h" + }; +#else + const std::string iq3_xxs_grid_src = read_file("iq3_xxs_grid.cl"); +#endif + + // mul_mv_iq3_xxs_f32_flat + { +#ifdef GGML_OPENCL_EMBED_KERNELS + const std::string kernel_src { + #include "mul_mv_iq3_xxs_f32_flat.cl.h" + }; +#else + const std::string kernel_src = read_file("mul_mv_iq3_xxs_f32_flat.cl"); +#endif + cl_program prog = build_program_from_source(backend_ctx, (iq3_xxs_grid_src + kernel_src).c_str(), compile_opts); + CL_CHECK((backend_ctx->kernel_mul_mv_iq3_xxs_f32_flat = clCreateKernel(prog, "kernel_mul_mv_iq3_xxs_f32_flat", &err), err)); + CL_CHECK(clReleaseProgram(prog)); + GGML_LOG_CONT("."); + } + + // mul_mm_iq3_xxs_f32_l4_lm + { +#ifdef GGML_OPENCL_EMBED_KERNELS + const std::string kernel_src { + #include "mul_mm_iq3_xxs_f32_l4_lm.cl.h" + }; +#else + const std::string kernel_src = read_file("mul_mm_iq3_xxs_f32_l4_lm.cl"); +#endif + cl_program prog = build_program_from_source(backend_ctx, (iq3_xxs_grid_src + kernel_src).c_str(), compile_opts); + CL_CHECK((backend_ctx->kernel_mul_mm_iq3_xxs_f32_l4_lm = clCreateKernel(prog, "kernel_mul_mm_iq3_xxs_f32_l4_lm", &err), err)); + CL_CHECK(clReleaseProgram(prog)); + GGML_LOG_CONT("."); + } + + // mul_mv_iq3_xxs_f32 + { +#ifdef GGML_OPENCL_EMBED_KERNELS + const std::string kernel_src { + #include "mul_mv_iq3_xxs_f32.cl.h" + }; +#else + const std::string kernel_src = read_file("mul_mv_iq3_xxs_f32.cl"); +#endif + cl_program prog = build_program_from_source(backend_ctx, (iq3_xxs_grid_src + kernel_src).c_str(), compile_opts); + CL_CHECK((backend_ctx->kernel_mul_mv_iq3_xxs_f32 = clCreateKernel(prog, "kernel_mul_mv_iq3_xxs_f32", &err), err)); + CL_CHECK(clReleaseProgram(prog)); + GGML_LOG_CONT("."); + } + // mul_mv_q3_k_f32 { #ifdef GGML_OPENCL_EMBED_KERNELS @@ -4959,6 +5097,180 @@ static void load_cl_kernels(ggml_backend_opencl_context *backend_ctx) { GGML_LOG_CONT("."); } + // gemm_noshuffle_iq3_s_q8_1_dp4a + { +#ifdef GGML_OPENCL_EMBED_KERNELS + const std::string kernel_src { + #include "gemm_noshuffle_iq3_s_q8_1_dp4a.cl.h" + }; +#else + const std::string kernel_src = read_file("gemm_noshuffle_iq3_s_q8_1_dp4a.cl"); +#endif + cl_program prog = build_program_from_source(backend_ctx, (iq3_s_grid_src + kernel_src).c_str(), compile_opts); + CL_CHECK((backend_ctx->kernel_gemm_noshuffle_iq3_s_q8_1_dp4a = clCreateKernel(prog, "kernel_gemm_noshuffle_iq3_s_q8_1_dp4a", &err), err)); + CL_CHECK(clReleaseProgram(prog)); + prog = build_program_from_source(backend_ctx, (iq3_s_grid_src + kernel_src).c_str(), compile_opts + " -DTILESIZE_N=8"); + CL_CHECK((backend_ctx->kernel_gemm_noshuffle_iq3_s_q8_1_dp4a_narrow = clCreateKernel(prog, "kernel_gemm_noshuffle_iq3_s_q8_1_dp4a", &err), err)); + CL_CHECK(clReleaseProgram(prog)); + GGML_LOG_CONT("."); + } + + // gemm_noshuffle_iq3_s_f32 + { +#ifdef GGML_OPENCL_EMBED_KERNELS + const std::string kernel_src { + #include "gemm_noshuffle_iq3_s_f32.cl.h" + }; +#else + const std::string kernel_src = read_file("gemm_noshuffle_iq3_s_f32.cl"); +#endif + cl_program prog = build_program_from_source(backend_ctx, (iq3_s_grid_src + kernel_src).c_str(), compile_opts); + CL_CHECK((backend_ctx->kernel_gemm_noshuffle_iq3_s_f32 = clCreateKernel(prog, "kernel_gemm_noshuffle_iq3_s_f32", &err), err)); + CL_CHECK(clReleaseProgram(prog)); + GGML_LOG_CONT("."); + } + + // gemv_noshuffle_iq3_s_f32 + { +#ifdef GGML_OPENCL_EMBED_KERNELS + const std::string kernel_src { + #include "gemv_noshuffle_iq3_s_f32.cl.h" + }; +#else + const std::string kernel_src = read_file("gemv_noshuffle_iq3_s_f32.cl"); +#endif + if (const char * env = getenv("GGML_OPENCL_IQ3_S_MV_NSG")) { + const int nsg = atoi(env); + if (nsg == 1 || nsg == 2 || nsg == 4 || nsg == 8) { + backend_ctx->iq3_s_mv_nsg = nsg; + } + } + for (;;) { + const std::string opts = compile_opts + " -DIQ3_S_MV_NSG=" + std::to_string(backend_ctx->iq3_s_mv_nsg); + cl_program prog = build_program_from_source(backend_ctx, (iq3_s_grid_src + kernel_src).c_str(), opts); + cl_kernel kernel = clCreateKernel(prog, "kernel_gemv_noshuffle_iq3_s_f32", &err); + CL_CHECK(err); + size_t max_workgroup; + CL_CHECK(clGetKernelWorkGroupInfo(kernel, backend_ctx->device, CL_KERNEL_WORK_GROUP_SIZE, sizeof(max_workgroup), &max_workgroup, nullptr)); + if (64*(size_t)backend_ctx->iq3_s_mv_nsg <= max_workgroup) { + backend_ctx->kernel_gemv_noshuffle_iq3_s_f32 = kernel; + cl_kernel export_grid = clCreateKernel(prog, "kernel_iq3_s_grid_export", &err); + CL_CHECK(err); + backend_ctx->iq3_s_grid_buffer.allocate(backend_ctx->context, 512*sizeof(cl_uint)); + CL_CHECK(clSetKernelArg(export_grid, 0, sizeof(cl_mem), &backend_ctx->iq3_s_grid_buffer.buffer)); + size_t global[] = { 512 }; + size_t local[] = { 64 }; + cl_event evt; + CL_CHECK(clEnqueueNDRangeKernel(backend_ctx->queue, export_grid, 1, nullptr, global, local, 0, nullptr, &evt)); + CL_CHECK(clWaitForEvents(1, &evt)); + CL_CHECK(clReleaseEvent(evt)); + CL_CHECK(clReleaseKernel(export_grid)); + CL_CHECK(clReleaseProgram(prog)); + + cl_image_format format = { CL_R, CL_UNSIGNED_INT32 }; + cl_image_desc desc = {}; + desc.image_type = CL_MEM_OBJECT_IMAGE1D_BUFFER; + desc.image_width = 512; + desc.buffer = backend_ctx->iq3_s_grid_buffer.buffer; + backend_ctx->iq3_s_grid_image = clCreateImage(backend_ctx->context, CL_MEM_READ_ONLY, &format, &desc, nullptr, &err); + CL_CHECK(err); + break; + } + CL_CHECK(clReleaseKernel(kernel)); + CL_CHECK(clReleaseProgram(prog)); + GGML_ASSERT(backend_ctx->iq3_s_mv_nsg > 1); + backend_ctx->iq3_s_mv_nsg /= 2; + } + GGML_LOG_CONT("."); + } + + // gemm_noshuffle_iq3_xxs_q8_1_dp4a + { +#ifdef GGML_OPENCL_EMBED_KERNELS + const std::string kernel_src { + #include "gemm_noshuffle_iq3_xxs_q8_1_dp4a.cl.h" + }; +#else + const std::string kernel_src = read_file("gemm_noshuffle_iq3_xxs_q8_1_dp4a.cl"); +#endif + cl_program prog = build_program_from_source(backend_ctx, (iq3_xxs_grid_src + kernel_src).c_str(), compile_opts); + CL_CHECK((backend_ctx->kernel_gemm_noshuffle_iq3_xxs_q8_1_dp4a = clCreateKernel(prog, "kernel_gemm_noshuffle_iq3_xxs_q8_1_dp4a", &err), err)); + CL_CHECK(clReleaseProgram(prog)); + prog = build_program_from_source(backend_ctx, (iq3_xxs_grid_src + kernel_src).c_str(), compile_opts + " -DTILESIZE_N=8"); + CL_CHECK((backend_ctx->kernel_gemm_noshuffle_iq3_xxs_q8_1_dp4a_narrow = clCreateKernel(prog, "kernel_gemm_noshuffle_iq3_xxs_q8_1_dp4a", &err), err)); + CL_CHECK(clReleaseProgram(prog)); + GGML_LOG_CONT("."); + } + + // gemm_noshuffle_iq3_xxs_f32 + { +#ifdef GGML_OPENCL_EMBED_KERNELS + const std::string kernel_src { + #include "gemm_noshuffle_iq3_xxs_f32.cl.h" + }; +#else + const std::string kernel_src = read_file("gemm_noshuffle_iq3_xxs_f32.cl"); +#endif + cl_program prog = build_program_from_source(backend_ctx, (iq3_xxs_grid_src + kernel_src).c_str(), compile_opts); + CL_CHECK((backend_ctx->kernel_gemm_noshuffle_iq3_xxs_f32 = clCreateKernel(prog, "kernel_gemm_noshuffle_iq3_xxs_f32", &err), err)); + CL_CHECK(clReleaseProgram(prog)); + GGML_LOG_CONT("."); + } + + // gemv_noshuffle_iq3_xxs_f32 + { +#ifdef GGML_OPENCL_EMBED_KERNELS + const std::string kernel_src { + #include "gemv_noshuffle_iq3_xxs_f32.cl.h" + }; +#else + const std::string kernel_src = read_file("gemv_noshuffle_iq3_xxs_f32.cl"); +#endif + if (const char * env = getenv("GGML_OPENCL_IQ3_XXS_MV_NSG")) { + const int nsg = atoi(env); + if (nsg == 1 || nsg == 2 || nsg == 4 || nsg == 8) { + backend_ctx->iq3_xxs_mv_nsg = nsg; + } + } + for (;;) { + const std::string opts = compile_opts + " -DIQ3_XXS_MV_NSG=" + std::to_string(backend_ctx->iq3_xxs_mv_nsg); + cl_program prog = build_program_from_source(backend_ctx, (iq3_xxs_grid_src + kernel_src).c_str(), opts); + cl_kernel kernel = clCreateKernel(prog, "kernel_gemv_noshuffle_iq3_xxs_f32", &err); + CL_CHECK(err); + size_t max_workgroup; + CL_CHECK(clGetKernelWorkGroupInfo(kernel, backend_ctx->device, CL_KERNEL_WORK_GROUP_SIZE, sizeof(max_workgroup), &max_workgroup, nullptr)); + if (64*(size_t)backend_ctx->iq3_xxs_mv_nsg <= max_workgroup) { + backend_ctx->kernel_gemv_noshuffle_iq3_xxs_f32 = kernel; + cl_kernel export_grid = clCreateKernel(prog, "kernel_iq3_xxs_grid_export", &err); + CL_CHECK(err); + backend_ctx->iq3_xxs_grid_buffer.allocate(backend_ctx->context, 256*sizeof(cl_uint)); + CL_CHECK(clSetKernelArg(export_grid, 0, sizeof(cl_mem), &backend_ctx->iq3_xxs_grid_buffer.buffer)); + size_t global[] = { 256 }; + size_t local[] = { 64 }; + cl_event evt; + CL_CHECK(clEnqueueNDRangeKernel(backend_ctx->queue, export_grid, 1, nullptr, global, local, 0, nullptr, &evt)); + CL_CHECK(clWaitForEvents(1, &evt)); + CL_CHECK(clReleaseEvent(evt)); + CL_CHECK(clReleaseKernel(export_grid)); + CL_CHECK(clReleaseProgram(prog)); + + cl_image_format format = { CL_R, CL_UNSIGNED_INT32 }; + cl_image_desc desc = {}; + desc.image_type = CL_MEM_OBJECT_IMAGE1D_BUFFER; + desc.image_width = 256; + desc.buffer = backend_ctx->iq3_xxs_grid_buffer.buffer; + backend_ctx->iq3_xxs_grid_image = clCreateImage(backend_ctx->context, CL_MEM_READ_ONLY, &format, &desc, nullptr, &err); + CL_CHECK(err); + break; + } + CL_CHECK(clReleaseKernel(kernel)); + CL_CHECK(clReleaseProgram(prog)); + GGML_ASSERT(backend_ctx->iq3_xxs_mv_nsg > 1); + backend_ctx->iq3_xxs_mv_nsg /= 2; + } + GGML_LOG_CONT("."); + } + // gemm_noshuffle_iq2_s_q8_1_dp4a { #ifdef GGML_OPENCL_EMBED_KERNELS @@ -8800,10 +9112,11 @@ static void transpose_2d_as_32b( ggml_backend_opencl_context * backend_ctx, cl_mem src, cl_mem dst, size_t size, cl_int stride, cl_int rows, - bool blocking = true + bool blocking = true, + bool auto_local = false ) { transpose_2d(backend_ctx, backend_ctx->kernel_transpose_32_buf, - src, dst, size, stride, rows, blocking); + src, dst, size, stride, rows, blocking, auto_local); } #endif // GGML_OPENCL_USE_ADRENO_KERNELS @@ -9354,6 +9667,49 @@ struct ggml_tensor_extra_cl_q3_K { } }; +struct ggml_tensor_extra_cl_iq3_xxs { + bool noshuffle = false; + cl_mem qs = nullptr; // grid indices + cl_mem sas = nullptr; // packed sign codes and scales + cl_mem d = nullptr; // super-block scales + + size_t size_qs = 0, size_sas = 0, size_d = 0; + + ~ggml_tensor_extra_cl_iq3_xxs() { reset(); } + + void reset() { + noshuffle = false; + if (qs != nullptr) { CL_CHECK(clReleaseMemObject(qs)); qs = nullptr; } + if (sas != nullptr) { CL_CHECK(clReleaseMemObject(sas)); sas = nullptr; } + if (d != nullptr) { CL_CHECK(clReleaseMemObject(d)); d = nullptr; } + size_qs = size_sas = size_d = 0; + } +}; + +struct ggml_tensor_extra_cl_iq3_s { + bool noshuffle = false; + cl_mem qs = nullptr; // grid index low bits + cl_mem sg = nullptr; // sign bits + cl_mem qh = nullptr; // grid index high bits + cl_mem sc = nullptr; // packed scales + cl_mem d = nullptr; // super-block scales + + size_t size_qs = 0, size_sg = 0, size_qh = 0, size_sc = 0, size_d = 0; + + ~ggml_tensor_extra_cl_iq3_s() { reset(); } + + void reset() { + noshuffle = false; + if (qs != nullptr) { CL_CHECK(clReleaseMemObject(qs)); qs = nullptr; } + if (sg != nullptr) { CL_CHECK(clReleaseMemObject(sg)); sg = nullptr; } + if (qh != nullptr) { CL_CHECK(clReleaseMemObject(qh)); qh = nullptr; } + if (sc != nullptr) { CL_CHECK(clReleaseMemObject(sc)); sc = nullptr; } + if (d != nullptr) { CL_CHECK(clReleaseMemObject(d)); d = nullptr; } + size_qs = size_sg = size_qh = size_sc = size_d = 0; + } +}; + + struct ggml_tensor_extra_cl_iq2_xxs { bool noshuffle = false; cl_mem qs = nullptr; // grid indices @@ -10588,6 +10944,42 @@ static bool ggml_cl_q3_k_use_noshuffle(const ggml_backend_opencl_context * backe #endif } +static bool ggml_cl_iq3_s_use_noshuffle(const ggml_backend_opencl_context * backend_ctx, const ggml_tensor * tensor) { +#if defined(GGML_OPENCL_SOA_Q) && defined(GGML_OPENCL_USE_ADRENO_KERNELS) + if (const char * env = getenv("GGML_OPENCL_IQ3_S_NOSHUFFLE")) { + if (env[0] == '0') { + return false; + } + } + return tensor->type == GGML_TYPE_IQ3_S && + backend_ctx->gpu_family == ADRENO && + ggml_is_contiguous(tensor) && !tensor->view_src && + use_adreno_kernels(backend_ctx, tensor); +#else + GGML_UNUSED(backend_ctx); + GGML_UNUSED(tensor); + return false; +#endif +} + +static bool ggml_cl_iq3_xxs_use_noshuffle(const ggml_backend_opencl_context * backend_ctx, const ggml_tensor * tensor) { +#if defined(GGML_OPENCL_SOA_Q) && defined(GGML_OPENCL_USE_ADRENO_KERNELS) + if (const char * env = getenv("GGML_OPENCL_IQ3_XXS_NOSHUFFLE")) { + if (env[0] == '0') { + return false; + } + } + return tensor->type == GGML_TYPE_IQ3_XXS && + backend_ctx->gpu_family == ADRENO && + ggml_is_contiguous(tensor) && !tensor->view_src && + use_adreno_kernels(backend_ctx, tensor); +#else + GGML_UNUSED(backend_ctx); + GGML_UNUSED(tensor); + return false; +#endif +} + static bool ggml_cl_iq1_m_use_noshuffle(const ggml_backend_opencl_context * backend_ctx, const ggml_tensor * tensor) { #if defined(GGML_OPENCL_SOA_Q) && defined(GGML_OPENCL_USE_ADRENO_KERNELS) if (const char * env = getenv("GGML_OPENCL_IQ1_M_NOSHUFFLE")) { @@ -11273,6 +11665,18 @@ static bool ggml_opencl_supports_op(ggml_backend_dev_t dev, const struct ggml_te return false; } return op->src[1]->type == GGML_TYPE_F32 && ggml_is_contiguous_rows(op->src[1]); + } else if (op->src[0]->type == GGML_TYPE_IQ3_S) { + const ggml_tensor * weight = op->src[0]->view_src; + if (weight && ggml_cl_iq3_s_use_noshuffle(backend_ctx, weight)) { + return false; + } + return op->src[1]->type == GGML_TYPE_F32 && ggml_is_contiguous_rows(op->src[1]); + } else if (op->src[0]->type == GGML_TYPE_IQ3_XXS) { + const ggml_tensor * weight = op->src[0]->view_src; + if (weight && ggml_cl_iq3_xxs_use_noshuffle(backend_ctx, weight)) { + return false; + } + return op->src[1]->type == GGML_TYPE_F32 && ggml_is_contiguous_rows(op->src[1]); } else if (op->src[0]->type == GGML_TYPE_Q3_K) { const ggml_tensor * weight = op->src[0]->view_src; if (weight && ggml_cl_q3_k_use_noshuffle(backend_ctx, weight)) { @@ -11686,6 +12090,12 @@ struct ggml_backend_opencl_buffer_context { for (ggml_tensor_extra_cl_iq1_s * e : temp_tensor_extras_iq1_s) { delete e; } + for (ggml_tensor_extra_cl_iq3_s * e : temp_tensor_extras_iq3_s) { + delete e; + } + for (ggml_tensor_extra_cl_iq3_xxs * e : temp_tensor_extras_iq3_xxs) { + delete e; + } for (ggml_tensor_extra_cl_q6_K * e : temp_tensor_extras_q6_K_in_use) { delete e; } @@ -11713,6 +12123,12 @@ struct ggml_backend_opencl_buffer_context { for (ggml_tensor_extra_cl_iq1_s * e : temp_tensor_extras_iq1_s_in_use) { delete e; } + for (ggml_tensor_extra_cl_iq3_s * e : temp_tensor_extras_iq3_s_in_use) { + delete e; + } + for (ggml_tensor_extra_cl_iq3_xxs * e : temp_tensor_extras_iq3_xxs_in_use) { + delete e; + } for (ggml_tensor_extra_cl_q5_K * e : temp_tensor_extras_q5_K) { delete e; } @@ -12006,6 +12422,36 @@ struct ggml_backend_opencl_buffer_context { return extra; } + ggml_tensor_extra_cl_iq3_s * ggml_opencl_alloc_temp_tensor_extra_iq3_s() { + ggml_tensor_extra_cl_iq3_s * extra; + if (temp_tensor_extras_iq3_s.empty()) { + extra = new ggml_tensor_extra_cl_iq3_s(); + } else { + extra = temp_tensor_extras_iq3_s.back(); + temp_tensor_extras_iq3_s.pop_back(); + } + + temp_tensor_extras_iq3_s_in_use.push_back(extra); + + extra->reset(); + return extra; + } + + ggml_tensor_extra_cl_iq3_xxs * ggml_opencl_alloc_temp_tensor_extra_iq3_xxs() { + ggml_tensor_extra_cl_iq3_xxs * extra; + if (temp_tensor_extras_iq3_xxs.empty()) { + extra = new ggml_tensor_extra_cl_iq3_xxs(); + } else { + extra = temp_tensor_extras_iq3_xxs.back(); + temp_tensor_extras_iq3_xxs.pop_back(); + } + + temp_tensor_extras_iq3_xxs_in_use.push_back(extra); + + extra->reset(); + return extra; + } + ggml_tensor_extra_cl_q6_K * ggml_opencl_alloc_temp_tensor_extra_q6_K() { ggml_tensor_extra_cl_q6_K * extra; if (temp_tensor_extras_q6_K.empty()) { @@ -12122,6 +12568,16 @@ struct ggml_backend_opencl_buffer_context { } temp_tensor_extras_iq1_s_in_use.clear(); + for (ggml_tensor_extra_cl_iq3_s * e : temp_tensor_extras_iq3_s_in_use) { + temp_tensor_extras_iq3_s.push_back(e); + } + temp_tensor_extras_iq3_s_in_use.clear(); + + for (ggml_tensor_extra_cl_iq3_xxs * e : temp_tensor_extras_iq3_xxs_in_use) { + temp_tensor_extras_iq3_xxs.push_back(e); + } + temp_tensor_extras_iq3_xxs_in_use.clear(); + q8_0_soa_tensors.clear(); q4_0_soa_tensors.clear(); } @@ -12158,6 +12614,7 @@ struct ggml_backend_opencl_buffer_context { std::vector temp_tensor_extras_iq2_s; std::vector temp_tensor_extras_iq2_xs; std::vector temp_tensor_extras_iq2_xxs; + std::vector temp_tensor_extras_iq3_xxs; std::vector temp_tensor_extras_q2_K_in_use; std::vector temp_tensor_extras_q3_K_in_use; std::vector temp_tensor_extras_iq2_s_in_use; @@ -12169,6 +12626,9 @@ struct ggml_backend_opencl_buffer_context { std::vector temp_tensor_extras_iq1_m; std::vector temp_tensor_extras_iq1_m_in_use; std::vector temp_tensor_extras_iq1_s_in_use; + std::vector temp_tensor_extras_iq3_s; + std::vector temp_tensor_extras_iq3_s_in_use; + std::vector temp_tensor_extras_iq3_xxs_in_use; std::vector temp_tensor_extras_q6_K; std::vector temp_tensor_extras_q6_K_in_use; @@ -12413,7 +12873,7 @@ static void ggml_backend_opencl_buffer_set_tensor(ggml_backend_buffer_t buffer, const int k = tensor->ne[0]; transpose_2d_as_8b(backend_ctx, extra->qs, extra->qs, size_qs, k/4, m, true, true); transpose_2d_as_8b(backend_ctx, extra->sc, extra->sc, size_sc, k/16, m, true, true); - transpose_2d(backend_ctx, backend_ctx->kernel_transpose_32_buf, extra->dm, extra->dm, size_dm, k/(int)blck, m, true, true); + transpose_2d_as_32b(backend_ctx, extra->dm, extra->dm, size_dm, k/(int)blck, m, true, true); } #endif @@ -12496,8 +12956,8 @@ static void ggml_backend_opencl_buffer_set_tensor(ggml_backend_buffer_t buffer, const int k = tensor->ne[0]; transpose_2d_as_8b(backend_ctx, extra->qs, extra->qs, size_qs, k/4, m, true, true); transpose_2d_as_8b(backend_ctx, extra->hm, extra->hm, size_hm, k/8, m, true, true); - transpose_2d(backend_ctx, backend_ctx->kernel_transpose_32_buf, extra->sc, extra->sc, size_sc, 3*k/(int)blck, m, true, true); - transpose_2d(backend_ctx, backend_ctx->kernel_transpose_16_buf, extra->d, extra->d, size_d, k/(int)blck, m, true, true); + transpose_2d_as_32b(backend_ctx, extra->sc, extra->sc, size_sc, 3*k/(int)blck, m, true, true); + transpose_2d_as_16b(backend_ctx, extra->d, extra->d, size_d, k/(int)blck, m, true, true); } #endif @@ -12510,6 +12970,168 @@ static void ggml_backend_opencl_buffer_set_tensor(ggml_backend_buffer_t buffer, return; } + if (tensor->type == GGML_TYPE_IQ3_S) { + GGML_ASSERT(offset == 0 && size == ggml_nbytes(tensor)); + ggml_tensor_extra_cl * extra_orig = (ggml_tensor_extra_cl *)tensor->extra; + GGML_ASSERT(extra_orig && "Tensors in OpenCL backend should have been allocated and initialized"); + + ggml_backend_opencl_buffer_context * ctx = (ggml_backend_opencl_buffer_context *) buffer->context; + ggml_tensor_extra_cl_iq3_s * extra = ctx->ggml_opencl_alloc_temp_tensor_extra_iq3_s(); + extra->noshuffle = ggml_cl_iq3_s_use_noshuffle(backend_ctx, tensor); + + const size_t blck = (size_t)ggml_blck_size(tensor->type); + const size_t n_blk = ggml_nelements(tensor) / blck; + + const size_t size_qs = n_blk * 64; + const size_t size_sg = n_blk * 32; + const size_t size_qh = n_blk * 8; + const size_t size_sc = n_blk * 4; + const size_t size_d = n_blk * 2; + GGML_ASSERT(size_qs + size_sg + size_qh + size_sc + size_d == ggml_nbytes(tensor) && "Incorrect tensor size"); + + cl_int err; + cl_mem data_device = clCreateBuffer(context, CL_MEM_READ_WRITE, + ggml_nbytes(tensor), NULL, &err); + CL_CHECK(err); + CL_CHECK(clEnqueueWriteBuffer( + queue, data_device, CL_TRUE, 0, + ggml_nbytes(tensor), data, 0, NULL, NULL)); + + cl_buffer_region region; + region.origin = align_to(extra_orig->offset + tensor->view_offs + offset, backend_ctx->alignment); + region.size = size_qs; + CL_CHECK((extra->qs = clCreateSubBuffer(extra_orig->data_device, CL_MEM_READ_WRITE, CL_BUFFER_CREATE_TYPE_REGION, ®ion, &err), err)); + + region.origin = align_to(region.origin + size_qs, backend_ctx->alignment); + region.size = size_sg; + CL_CHECK((extra->sg = clCreateSubBuffer(extra_orig->data_device, CL_MEM_READ_WRITE, CL_BUFFER_CREATE_TYPE_REGION, ®ion, &err), err)); + + region.origin = align_to(region.origin + size_sg, backend_ctx->alignment); + region.size = size_qh; + CL_CHECK((extra->qh = clCreateSubBuffer(extra_orig->data_device, CL_MEM_READ_WRITE, CL_BUFFER_CREATE_TYPE_REGION, ®ion, &err), err)); + + region.origin = align_to(region.origin + size_qh, backend_ctx->alignment); + region.size = size_sc; + CL_CHECK((extra->sc = clCreateSubBuffer(extra_orig->data_device, CL_MEM_READ_WRITE, CL_BUFFER_CREATE_TYPE_REGION, ®ion, &err), err)); + + region.origin = align_to(region.origin + size_sc, backend_ctx->alignment); + region.size = size_d; + CL_CHECK((extra->d = clCreateSubBuffer(extra_orig->data_device, CL_MEM_READ_WRITE, CL_BUFFER_CREATE_TYPE_REGION, ®ion, &err), err)); + + cl_kernel kernel = backend_ctx->kernel_convert_block_iq3_s; + cl_ulong nb_arg = (cl_ulong)n_blk; + CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), &data_device)); + CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_mem), &extra->qs)); + CL_CHECK(clSetKernelArg(kernel, 2, sizeof(cl_mem), &extra->sg)); + CL_CHECK(clSetKernelArg(kernel, 3, sizeof(cl_mem), &extra->qh)); + CL_CHECK(clSetKernelArg(kernel, 4, sizeof(cl_mem), &extra->sc)); + CL_CHECK(clSetKernelArg(kernel, 5, sizeof(cl_mem), &extra->d)); + CL_CHECK(clSetKernelArg(kernel, 6, sizeof(cl_ulong), &nb_arg)); + + size_t global_work_size[] = { (size_t)CEIL_DIV(n_blk, 64)*64, 1, 1 }; + size_t local_work_size[] = { 64, 1, 1 }; + + cl_event evt; + CL_CHECK(clEnqueueNDRangeKernel(queue, kernel, 3, NULL, global_work_size, local_work_size, 0, NULL, &evt)); + CL_CHECK(clWaitForEvents(1, &evt)); + CL_CHECK(clReleaseEvent(evt)); + CL_CHECK(clReleaseMemObject(data_device)); + +#ifdef GGML_OPENCL_USE_ADRENO_KERNELS + if (extra->noshuffle) { + const int m = tensor->ne[1]; + const int k = tensor->ne[0]; + transpose_2d_as_8b(backend_ctx, extra->qs, extra->qs, size_qs, k/4, m, true, true); + transpose_2d_as_8b(backend_ctx, extra->sg, extra->sg, size_sg, k/8, m, true, true); + transpose_2d_as_8b(backend_ctx, extra->qh, extra->qh, size_qh, k/32, m, true, true); + transpose_2d_as_8b(backend_ctx, extra->sc, extra->sc, size_sc, k/64, m, true, true); + transpose_2d_as_16b(backend_ctx, extra->d, extra->d, size_d, k/256, m, true, true); + } +#endif + + extra->size_qs = size_qs; + extra->size_sg = size_sg; + extra->size_qh = size_qh; + extra->size_sc = size_sc; + extra->size_d = size_d; + + tensor->extra = extra; + return; + } + + if (tensor->type == GGML_TYPE_IQ3_XXS) { + GGML_ASSERT(offset == 0 && size == ggml_nbytes(tensor)); + ggml_tensor_extra_cl * extra_orig = (ggml_tensor_extra_cl *)tensor->extra; + GGML_ASSERT(extra_orig && "Tensors in OpenCL backend should have been allocated and initialized"); + + ggml_backend_opencl_buffer_context * ctx = (ggml_backend_opencl_buffer_context *) buffer->context; + ggml_tensor_extra_cl_iq3_xxs * extra = ctx->ggml_opencl_alloc_temp_tensor_extra_iq3_xxs(); + extra->noshuffle = ggml_cl_iq3_xxs_use_noshuffle(backend_ctx, tensor); + + const size_t blck = (size_t)ggml_blck_size(tensor->type); + const size_t n_blk = ggml_nelements(tensor) / blck; + + const size_t size_qs = n_blk * 64; + const size_t size_sas = n_blk * 32; + const size_t size_d = n_blk * 2; + GGML_ASSERT(size_qs + size_sas + size_d == ggml_nbytes(tensor) && "Incorrect tensor size"); + + cl_int err; + cl_mem data_device = clCreateBuffer(context, CL_MEM_READ_WRITE, + ggml_nbytes(tensor), NULL, &err); + CL_CHECK(err); + CL_CHECK(clEnqueueWriteBuffer( + queue, data_device, CL_TRUE, 0, + ggml_nbytes(tensor), data, 0, NULL, NULL)); + + cl_buffer_region region; + region.origin = align_to(extra_orig->offset + tensor->view_offs + offset, backend_ctx->alignment); + region.size = size_qs; + CL_CHECK((extra->qs = clCreateSubBuffer(extra_orig->data_device, CL_MEM_READ_WRITE, CL_BUFFER_CREATE_TYPE_REGION, ®ion, &err), err)); + + region.origin = align_to(region.origin + size_qs, backend_ctx->alignment); + region.size = size_sas; + CL_CHECK((extra->sas = clCreateSubBuffer(extra_orig->data_device, CL_MEM_READ_WRITE, CL_BUFFER_CREATE_TYPE_REGION, ®ion, &err), err)); + + region.origin = align_to(region.origin + size_sas, backend_ctx->alignment); + region.size = size_d; + CL_CHECK((extra->d = clCreateSubBuffer(extra_orig->data_device, CL_MEM_READ_WRITE, CL_BUFFER_CREATE_TYPE_REGION, ®ion, &err), err)); + + cl_kernel kernel = backend_ctx->kernel_convert_block_iq3_xxs; + cl_ulong nb_arg = (cl_ulong)n_blk; + CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), &data_device)); + CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_mem), &extra->qs)); + CL_CHECK(clSetKernelArg(kernel, 2, sizeof(cl_mem), &extra->sas)); + CL_CHECK(clSetKernelArg(kernel, 3, sizeof(cl_mem), &extra->d)); + CL_CHECK(clSetKernelArg(kernel, 4, sizeof(cl_ulong), &nb_arg)); + + size_t global_work_size[] = { (size_t)CEIL_DIV(n_blk, 64)*64, 1, 1 }; + size_t local_work_size[] = { 64, 1, 1 }; + + cl_event evt; + CL_CHECK(clEnqueueNDRangeKernel(queue, kernel, 3, NULL, global_work_size, local_work_size, 0, NULL, &evt)); + CL_CHECK(clWaitForEvents(1, &evt)); + CL_CHECK(clReleaseEvent(evt)); + CL_CHECK(clReleaseMemObject(data_device)); + +#ifdef GGML_OPENCL_USE_ADRENO_KERNELS + if (extra->noshuffle) { + const int m = tensor->ne[1]; + const int k = tensor->ne[0]; + transpose_2d_as_8b(backend_ctx, extra->qs, extra->qs, size_qs, k/4, m, true, true); + transpose_2d_as_32b(backend_ctx, extra->sas, extra->sas, size_sas, k/32, m, true, true); + transpose_2d_as_16b(backend_ctx, extra->d, extra->d, size_d, k/256, m, true, true); + } +#endif + + extra->size_qs = size_qs; + extra->size_sas = size_sas; + extra->size_d = size_d; + + tensor->extra = extra; + return; + } + if (tensor->type == GGML_TYPE_IQ1_M) { GGML_ASSERT(offset == 0 && size == ggml_nbytes(tensor)); ggml_tensor_extra_cl * extra_orig = (ggml_tensor_extra_cl *)tensor->extra; @@ -12571,7 +13193,7 @@ static void ggml_backend_opencl_buffer_set_tensor(ggml_backend_buffer_t buffer, const int k = tensor->ne[0]; transpose_2d_as_8b(backend_ctx, extra->qs, extra->qs, size_qs, k/8, m, true, true); transpose_2d_as_8b(backend_ctx, extra->qh, extra->qh, size_qh, k/16, m, true, true); - transpose_2d(backend_ctx, backend_ctx->kernel_transpose_16_buf, extra->sc, extra->sc, size_sc, k/64, m, true, true); + transpose_2d_as_16b(backend_ctx, extra->sc, extra->sc, size_sc, k/64, m, true, true); } #endif @@ -12643,8 +13265,8 @@ static void ggml_backend_opencl_buffer_set_tensor(ggml_backend_buffer_t buffer, const int m = tensor->ne[1]; const int k = tensor->ne[0]; transpose_2d_as_8b(backend_ctx, extra->qs, extra->qs, size_qs, k/8, m, true, true); - transpose_2d(backend_ctx, backend_ctx->kernel_transpose_16_buf, extra->qh, extra->qh, size_qh, k/32, m, true, true); - transpose_2d(backend_ctx, backend_ctx->kernel_transpose_16_buf, extra->d, extra->d, size_d, k/(int)blck, m, true, true); + transpose_2d_as_16b(backend_ctx, extra->qh, extra->qh, size_qh, k/32, m, true, true); + transpose_2d_as_16b(backend_ctx, extra->d, extra->d, size_d, k/(int)blck, m, true, true); } #endif @@ -12731,7 +13353,7 @@ static void ggml_backend_opencl_buffer_set_tensor(ggml_backend_buffer_t buffer, transpose_2d_as_8b(backend_ctx, extra->sg, extra->sg, size_sg, k/8, m, true, true); transpose_2d_as_8b(backend_ctx, extra->qh, extra->qh, size_qh, k/32, m, true, true); transpose_2d_as_8b(backend_ctx, extra->sc, extra->sc, size_sc, k/32, m, true, true); - transpose_2d(backend_ctx, backend_ctx->kernel_transpose_16_buf, extra->d, extra->d, size_d, k/(int)blck, m, true, true); + transpose_2d_as_16b(backend_ctx, extra->d, extra->d, size_d, k/(int)blck, m, true, true); } #endif @@ -12804,9 +13426,9 @@ static void ggml_backend_opencl_buffer_set_tensor(ggml_backend_buffer_t buffer, if (extra->noshuffle) { const int m = tensor->ne[1]; const int k = tensor->ne[0]; - transpose_2d(backend_ctx, backend_ctx->kernel_transpose_16_buf, extra->qs, extra->qs, size_qs, k/8, m, true, true); + transpose_2d_as_16b(backend_ctx, extra->qs, extra->qs, size_qs, k/8, m, true, true); transpose_2d_as_8b(backend_ctx, extra->sc, extra->sc, size_sc, k/32, m, true, true); - transpose_2d(backend_ctx, backend_ctx->kernel_transpose_16_buf, extra->d, extra->d, size_d, k/(int)blck, m, true, true); + transpose_2d_as_16b(backend_ctx, extra->d, extra->d, size_d, k/(int)blck, m, true, true); } #endif @@ -12878,8 +13500,8 @@ static void ggml_backend_opencl_buffer_set_tensor(ggml_backend_buffer_t buffer, const int m = tensor->ne[1]; const int k = tensor->ne[0]; transpose_2d_as_8b(backend_ctx, extra->qs, extra->qs, size_qs, k/8, m, true, true); - transpose_2d(backend_ctx, backend_ctx->kernel_transpose_32_buf, extra->sas, extra->sas, size_sas, k/32, m, true, true); - transpose_2d(backend_ctx, backend_ctx->kernel_transpose_16_buf, extra->d, extra->d, size_d, k/(int)blck, m, true, true); + transpose_2d_as_32b(backend_ctx, extra->sas, extra->sas, size_sas, k/32, m, true, true); + transpose_2d_as_16b(backend_ctx, extra->d, extra->d, size_d, k/(int)blck, m, true, true); } #endif @@ -14671,7 +15293,7 @@ static void ggml_backend_opencl_buffer_get_tensor(ggml_backend_buffer_t buffer, const int k = tensor->ne[0]; transpose_2d_as_8b(backend_ctx, extra->qs, qs, extra->size_qs, m, k/4, true, true); transpose_2d_as_8b(backend_ctx, extra->sc, sc, extra->size_sc, m, k/16, true, true); - transpose_2d(backend_ctx, backend_ctx->kernel_transpose_32_buf, extra->dm, dm, extra->size_dm, m, k/(int)blck, true, true); + transpose_2d_as_32b(backend_ctx, extra->dm, dm, extra->size_dm, m, k/(int)blck, true, true); } #endif cl_kernel kernel = extra->noshuffle ? backend_ctx->kernel_restore_block_q2_k_noshuffle @@ -14719,8 +15341,8 @@ static void ggml_backend_opencl_buffer_get_tensor(ggml_backend_buffer_t buffer, const int k = tensor->ne[0]; transpose_2d_as_8b(backend_ctx, extra->qs, qs, extra->size_qs, m, k/4, true, true); transpose_2d_as_8b(backend_ctx, extra->hm, hm, extra->size_hm, m, k/8, true, true); - transpose_2d(backend_ctx, backend_ctx->kernel_transpose_32_buf, extra->sc, sc, extra->size_sc, m, 3*k/(int)blck, true, true); - transpose_2d(backend_ctx, backend_ctx->kernel_transpose_16_buf, extra->d, d, extra->size_d, m, k/(int)blck, true, true); + transpose_2d_as_32b(backend_ctx, extra->sc, sc, extra->size_sc, m, 3*k/(int)blck, true, true); + transpose_2d_as_16b(backend_ctx, extra->d, d, extra->size_d, m, k/(int)blck, true, true); } #endif cl_kernel kernel = extra->noshuffle ? backend_ctx->kernel_restore_block_q3_k_noshuffle @@ -14742,6 +15364,118 @@ static void ggml_backend_opencl_buffer_get_tensor(ggml_backend_buffer_t buffer, return; } + if (tensor->type == GGML_TYPE_IQ3_S) { + if (tensor->view_src) { + offset += tensor->view_offs; + tensor = tensor->view_src; + } + ggml_tensor_extra_cl_iq3_s * extra = (ggml_tensor_extra_cl_iq3_s *)tensor->extra; + + const size_t blck = (size_t)ggml_blck_size(tensor->type); + + ggml_cl_buffer buf_unpacked; + buf_unpacked.allocate(backend_ctx->context, ggml_nbytes(tensor)); + + const size_t n_blk = (size_t)ggml_nelements(tensor) / blck; + cl_ulong nb_arg = (cl_ulong)n_blk; + + cl_mem qs = extra->qs, sg = extra->sg, qh = extra->qh, sc = extra->sc, d = extra->d; +#ifdef GGML_OPENCL_USE_ADRENO_KERNELS + ggml_cl_buffer buf_qs, buf_sg, buf_qh, buf_sc, buf_d; + if (extra->noshuffle) { + buf_qs.allocate(context, extra->size_qs); + buf_sg.allocate(context, extra->size_sg); + buf_qh.allocate(context, extra->size_qh); + buf_sc.allocate(context, extra->size_sc); + buf_d.allocate(context, extra->size_d); + + qs = buf_qs.buffer; + sg = buf_sg.buffer; + qh = buf_qh.buffer; + sc = buf_sc.buffer; + d = buf_d.buffer; + + const int m = tensor->ne[1]; + const int k = tensor->ne[0]; + transpose_2d_as_8b(backend_ctx, extra->qs, qs, extra->size_qs, m, k/4, true, true); + transpose_2d_as_8b(backend_ctx, extra->sg, sg, extra->size_sg, m, k/8, true, true); + transpose_2d_as_8b(backend_ctx, extra->qh, qh, extra->size_qh, m, k/32, true, true); + transpose_2d_as_8b(backend_ctx, extra->sc, sc, extra->size_sc, m, k/64, true, true); + transpose_2d_as_16b(backend_ctx, extra->d, d, extra->size_d, m, k/256, true, true); + } +#endif + + cl_kernel kernel = backend_ctx->kernel_restore_block_iq3_s; + CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), &qs)); + CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_mem), &sg)); + CL_CHECK(clSetKernelArg(kernel, 2, sizeof(cl_mem), &qh)); + CL_CHECK(clSetKernelArg(kernel, 3, sizeof(cl_mem), &sc)); + CL_CHECK(clSetKernelArg(kernel, 4, sizeof(cl_mem), &d)); + CL_CHECK(clSetKernelArg(kernel, 5, sizeof(cl_mem), &buf_unpacked.buffer)); + CL_CHECK(clSetKernelArg(kernel, 6, sizeof(cl_ulong), &nb_arg)); + + size_t gws[] = { (n_blk + 63) / 64 * 64 }; + size_t lws[] = { 64 }; + cl_event evt; + CL_CHECK(clEnqueueNDRangeKernel(queue, kernel, 1, NULL, gws, lws, 0, NULL, &evt)); + CL_CHECK(clWaitForEvents(1, &evt)); + CL_CHECK(clReleaseEvent(evt)); + CL_CHECK(clEnqueueReadBuffer(queue, buf_unpacked.buffer, CL_TRUE, offset, size, data, 0, NULL, NULL)); + return; + } + + if (tensor->type == GGML_TYPE_IQ3_XXS) { + if (tensor->view_src) { + offset += tensor->view_offs; + tensor = tensor->view_src; + } + ggml_tensor_extra_cl_iq3_xxs * extra = (ggml_tensor_extra_cl_iq3_xxs *)tensor->extra; + + const size_t blck = (size_t)ggml_blck_size(tensor->type); + + ggml_cl_buffer buf_unpacked; + buf_unpacked.allocate(backend_ctx->context, ggml_nbytes(tensor)); + + const size_t n_blk = (size_t)ggml_nelements(tensor) / blck; + cl_ulong nb_arg = (cl_ulong)n_blk; + + cl_mem qs = extra->qs, sas = extra->sas, d = extra->d; +#ifdef GGML_OPENCL_USE_ADRENO_KERNELS + ggml_cl_buffer buf_qs, buf_sas, buf_d; + if (extra->noshuffle) { + buf_qs.allocate(context, extra->size_qs); + buf_sas.allocate(context, extra->size_sas); + buf_d.allocate(context, extra->size_d); + + qs = buf_qs.buffer; + sas = buf_sas.buffer; + d = buf_d.buffer; + + const int m = tensor->ne[1]; + const int k = tensor->ne[0]; + transpose_2d_as_8b(backend_ctx, extra->qs, qs, extra->size_qs, m, k/4, true, true); + transpose_2d_as_32b(backend_ctx, extra->sas, sas, extra->size_sas, m, k/32, true, true); + transpose_2d_as_16b(backend_ctx, extra->d, d, extra->size_d, m, k/256, true, true); + } +#endif + + cl_kernel kernel = backend_ctx->kernel_restore_block_iq3_xxs; + CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), &qs)); + CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_mem), &sas)); + CL_CHECK(clSetKernelArg(kernel, 2, sizeof(cl_mem), &d)); + CL_CHECK(clSetKernelArg(kernel, 3, sizeof(cl_mem), &buf_unpacked.buffer)); + CL_CHECK(clSetKernelArg(kernel, 4, sizeof(cl_ulong), &nb_arg)); + + size_t gws[] = { (n_blk + 63) / 64 * 64 }; + size_t lws[] = { 64 }; + cl_event evt; + CL_CHECK(clEnqueueNDRangeKernel(queue, kernel, 1, NULL, gws, lws, 0, NULL, &evt)); + CL_CHECK(clWaitForEvents(1, &evt)); + CL_CHECK(clReleaseEvent(evt)); + CL_CHECK(clEnqueueReadBuffer(queue, buf_unpacked.buffer, CL_TRUE, offset, size, data, 0, NULL, NULL)); + return; + } + if (tensor->type == GGML_TYPE_IQ1_M) { if (tensor->view_src) { offset += tensor->view_offs; @@ -14771,7 +15505,7 @@ static void ggml_backend_opencl_buffer_get_tensor(ggml_backend_buffer_t buffer, const int k = tensor->ne[0]; transpose_2d_as_8b(backend_ctx, extra->qs, qs, extra->size_qs, m, k/8, true, true); transpose_2d_as_8b(backend_ctx, extra->qh, qh, extra->size_qh, m, k/16, true, true); - transpose_2d(backend_ctx, backend_ctx->kernel_transpose_16_buf, extra->sc, sc, extra->size_sc, m, k/64, true, true); + transpose_2d_as_16b(backend_ctx, extra->sc, sc, extra->size_sc, m, k/64, true, true); } #endif cl_kernel kernel = backend_ctx->kernel_restore_block_iq1_m; @@ -14819,8 +15553,8 @@ static void ggml_backend_opencl_buffer_get_tensor(ggml_backend_buffer_t buffer, const int m = tensor->ne[1]; const int k = tensor->ne[0]; transpose_2d_as_8b(backend_ctx, extra->qs, qs, extra->size_qs, m, k/8, true, true); - transpose_2d(backend_ctx, backend_ctx->kernel_transpose_16_buf, extra->qh, qh, extra->size_qh, m, k/32, true, true); - transpose_2d(backend_ctx, backend_ctx->kernel_transpose_16_buf, extra->d, d, extra->size_d, m, k/(int)blck, true, true); + transpose_2d_as_16b(backend_ctx, extra->qh, qh, extra->size_qh, m, k/32, true, true); + transpose_2d_as_16b(backend_ctx, extra->d, d, extra->size_d, m, k/(int)blck, true, true); } #endif cl_kernel kernel = backend_ctx->kernel_restore_block_iq1_s; @@ -14875,7 +15609,7 @@ static void ggml_backend_opencl_buffer_get_tensor(ggml_backend_buffer_t buffer, transpose_2d_as_8b(backend_ctx, extra->sg, sg, extra->size_sg, m, k/8, true, true); transpose_2d_as_8b(backend_ctx, extra->qh, qh, extra->size_qh, m, k/32, true, true); transpose_2d_as_8b(backend_ctx, extra->sc, sc, extra->size_sc, m, k/32, true, true); - transpose_2d(backend_ctx, backend_ctx->kernel_transpose_16_buf, extra->d, d, extra->size_d, m, k/(int)blck, true, true); + transpose_2d_as_16b(backend_ctx, extra->d, d, extra->size_d, m, k/(int)blck, true, true); } #endif cl_kernel kernel = backend_ctx->kernel_restore_block_iq2_s; @@ -14924,9 +15658,9 @@ static void ggml_backend_opencl_buffer_get_tensor(ggml_backend_buffer_t buffer, d = buf_d.buffer; const int m = tensor->ne[1]; const int k = tensor->ne[0]; - transpose_2d(backend_ctx, backend_ctx->kernel_transpose_16_buf, extra->qs, qs, extra->size_qs, m, k/8, true, true); + transpose_2d_as_16b(backend_ctx, extra->qs, qs, extra->size_qs, m, k/8, true, true); transpose_2d_as_8b(backend_ctx, extra->sc, sc, extra->size_sc, m, k/32, true, true); - transpose_2d(backend_ctx, backend_ctx->kernel_transpose_16_buf, extra->d, d, extra->size_d, m, k/(int)blck, true, true); + transpose_2d_as_16b(backend_ctx, extra->d, d, extra->size_d, m, k/(int)blck, true, true); } #endif cl_kernel kernel = backend_ctx->kernel_restore_block_iq2_xs; @@ -14974,8 +15708,8 @@ static void ggml_backend_opencl_buffer_get_tensor(ggml_backend_buffer_t buffer, const int m = tensor->ne[1]; const int k = tensor->ne[0]; transpose_2d_as_8b(backend_ctx, extra->qs, qs, extra->size_qs, m, k/8, true, true); - transpose_2d(backend_ctx, backend_ctx->kernel_transpose_32_buf, extra->sas, sas, extra->size_sas, m, k/32, true, true); - transpose_2d(backend_ctx, backend_ctx->kernel_transpose_16_buf, extra->d, d, extra->size_d, m, k/(int)blck, true, true); + transpose_2d_as_32b(backend_ctx, extra->sas, sas, extra->size_sas, m, k/32, true, true); + transpose_2d_as_16b(backend_ctx, extra->d, d, extra->size_d, m, k/(int)blck, true, true); } #endif cl_kernel kernel = backend_ctx->kernel_restore_block_iq2_xxs; @@ -27671,6 +28405,124 @@ static bool ggml_cl_mul_mat_q3_K_dp4a(ggml_backend_t backend, const ggml_tensor return true; } +static bool ggml_cl_mul_mat_iq3_s_dp4a(ggml_backend_t backend, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst) { + ggml_backend_opencl_context * backend_ctx = (ggml_backend_opencl_context *)backend->context; + ggml_tensor_extra_cl_iq3_s * extra0 = (ggml_tensor_extra_cl_iq3_s *)src0->extra; + ggml_tensor_extra_cl * extra1 = (ggml_tensor_extra_cl *)src1->extra; + ggml_tensor_extra_cl * extrad = (ggml_tensor_extra_cl *)dst->extra; + + cl_ulong offset1 = extra1->offset + src1->view_offs; + cl_ulong offsetd = extrad->offset + dst->view_offs; + const int m = src0->ne[1]; + const int n = src1->ne[1]; + const int k = src0->ne[0]; + + const int tile_n = n <= 8 ? 8 : 32; + cl_kernel gemm = n <= 8 ? backend_ctx->kernel_gemm_noshuffle_iq3_s_q8_1_dp4a_narrow + : backend_ctx->kernel_gemm_noshuffle_iq3_s_q8_1_dp4a; + const char * env = getenv("GGML_OPENCL_IQ3_S_DENSE_DP4A"); + if (!gemm || !backend_ctx->has_integer_dot || (env && env[0] == '0') || + n <= 1 || !ggml_is_contiguous(src1) || offset1 % backend_ctx->alignment != 0) { + return false; + } + + cl_int err; + cl_buffer_region region = { (size_t)offset1, (size_t)k*n*sizeof(float) }; + cl_mem activation = clCreateSubBuffer(extra1->data_device, 0, CL_BUFFER_CREATE_TYPE_REGION, ®ion, &err); + CL_CHECK(err); + const size_t n_blocks = (size_t)n*(k/32); + backend_ctx->prealloc_moe_qa.allocate(backend_ctx->context, (size_t)n*k); + backend_ctx->prealloc_moe_da.allocate(backend_ctx->context, n_blocks*sizeof(cl_half)); + backend_ctx->prealloc_moe_sa.allocate(backend_ctx->context, n_blocks*sizeof(cl_half)); + + cl_int nb = n_blocks; + cl_kernel quant = backend_ctx->kernel_quant_a_q8_1; + CL_CHECK(clSetKernelArg(quant, 0, sizeof(cl_mem), &activation)); + CL_CHECK(clSetKernelArg(quant, 1, sizeof(cl_mem), &backend_ctx->prealloc_moe_qa.buffer)); + CL_CHECK(clSetKernelArg(quant, 2, sizeof(cl_mem), &backend_ctx->prealloc_moe_da.buffer)); + CL_CHECK(clSetKernelArg(quant, 3, sizeof(cl_mem), &backend_ctx->prealloc_moe_sa.buffer)); + CL_CHECK(clSetKernelArg(quant, 4, sizeof(cl_int), &nb)); + size_t q_global[] = { CEIL_DIV(n_blocks, 64)*64 }; + size_t q_local[] = { 64 }; + backend_ctx->enqueue_ndrange_kernel(quant, 1, q_global, q_local, dst); + + CL_CHECK(clSetKernelArg(gemm, 0, sizeof(cl_mem), &extra0->qs)); + CL_CHECK(clSetKernelArg(gemm, 1, sizeof(cl_mem), &extra0->sg)); + CL_CHECK(clSetKernelArg(gemm, 2, sizeof(cl_mem), &extra0->qh)); + CL_CHECK(clSetKernelArg(gemm, 3, sizeof(cl_mem), &extra0->sc)); + CL_CHECK(clSetKernelArg(gemm, 4, sizeof(cl_mem), &extra0->d)); + CL_CHECK(clSetKernelArg(gemm, 5, sizeof(cl_mem), &backend_ctx->prealloc_moe_qa.buffer)); + CL_CHECK(clSetKernelArg(gemm, 6, sizeof(cl_mem), &backend_ctx->prealloc_moe_da.buffer)); + CL_CHECK(clSetKernelArg(gemm, 7, sizeof(cl_mem), &extrad->data_device)); + CL_CHECK(clSetKernelArg(gemm, 8, sizeof(cl_ulong), &offsetd)); + CL_CHECK(clSetKernelArg(gemm, 9, sizeof(int), &m)); + CL_CHECK(clSetKernelArg(gemm, 10, sizeof(int), &n)); + CL_CHECK(clSetKernelArg(gemm, 11, sizeof(int), &k)); + size_t global[] = { 64, (size_t)CEIL_DIV(m, 64), (size_t)CEIL_DIV(n, tile_n) }; + size_t local[] = { 64, 1, 1 }; + backend_ctx->enqueue_ndrange_kernel(gemm, 3, global, local, dst); + CL_CHECK(clReleaseMemObject(activation)); + return true; +} + +static bool ggml_cl_mul_mat_iq3_xxs_dp4a(ggml_backend_t backend, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst) { + ggml_backend_opencl_context * backend_ctx = (ggml_backend_opencl_context *)backend->context; + ggml_tensor_extra_cl_iq3_xxs * extra0 = (ggml_tensor_extra_cl_iq3_xxs *)src0->extra; + ggml_tensor_extra_cl * extra1 = (ggml_tensor_extra_cl *)src1->extra; + ggml_tensor_extra_cl * extrad = (ggml_tensor_extra_cl *)dst->extra; + + cl_ulong offset1 = extra1->offset + src1->view_offs; + cl_ulong offsetd = extrad->offset + dst->view_offs; + const int m = src0->ne[1]; + const int n = src1->ne[1]; + const int k = src0->ne[0]; + + const int tile_n = n <= 8 ? 8 : 32; + cl_kernel gemm = n <= 8 ? backend_ctx->kernel_gemm_noshuffle_iq3_xxs_q8_1_dp4a_narrow + : backend_ctx->kernel_gemm_noshuffle_iq3_xxs_q8_1_dp4a; + const char * env = getenv("GGML_OPENCL_IQ3_XXS_DENSE_DP4A"); + if (!gemm || !backend_ctx->has_integer_dot || (env && env[0] == '0') || + n <= 1 || !ggml_is_contiguous(src1) || offset1 % backend_ctx->alignment != 0) { + return false; + } + + cl_int err; + cl_buffer_region region = { (size_t)offset1, (size_t)k*n*sizeof(float) }; + cl_mem activation = clCreateSubBuffer(extra1->data_device, 0, CL_BUFFER_CREATE_TYPE_REGION, ®ion, &err); + CL_CHECK(err); + const size_t n_blocks = (size_t)n*(k/32); + backend_ctx->prealloc_moe_qa.allocate(backend_ctx->context, (size_t)n*k); + backend_ctx->prealloc_moe_da.allocate(backend_ctx->context, n_blocks*sizeof(cl_half)); + backend_ctx->prealloc_moe_sa.allocate(backend_ctx->context, n_blocks*sizeof(cl_half)); + + cl_int nb = n_blocks; + cl_kernel quant = backend_ctx->kernel_quant_a_q8_1; + CL_CHECK(clSetKernelArg(quant, 0, sizeof(cl_mem), &activation)); + CL_CHECK(clSetKernelArg(quant, 1, sizeof(cl_mem), &backend_ctx->prealloc_moe_qa.buffer)); + CL_CHECK(clSetKernelArg(quant, 2, sizeof(cl_mem), &backend_ctx->prealloc_moe_da.buffer)); + CL_CHECK(clSetKernelArg(quant, 3, sizeof(cl_mem), &backend_ctx->prealloc_moe_sa.buffer)); + CL_CHECK(clSetKernelArg(quant, 4, sizeof(cl_int), &nb)); + size_t q_global[] = { CEIL_DIV(n_blocks, 64)*64 }; + size_t q_local[] = { 64 }; + backend_ctx->enqueue_ndrange_kernel(quant, 1, q_global, q_local, dst); + + CL_CHECK(clSetKernelArg(gemm, 0, sizeof(cl_mem), &extra0->qs)); + CL_CHECK(clSetKernelArg(gemm, 1, sizeof(cl_mem), &extra0->sas)); + CL_CHECK(clSetKernelArg(gemm, 2, sizeof(cl_mem), &extra0->d)); + CL_CHECK(clSetKernelArg(gemm, 3, sizeof(cl_mem), &backend_ctx->prealloc_moe_qa.buffer)); + CL_CHECK(clSetKernelArg(gemm, 4, sizeof(cl_mem), &backend_ctx->prealloc_moe_da.buffer)); + CL_CHECK(clSetKernelArg(gemm, 5, sizeof(cl_mem), &extrad->data_device)); + CL_CHECK(clSetKernelArg(gemm, 6, sizeof(cl_ulong), &offsetd)); + CL_CHECK(clSetKernelArg(gemm, 7, sizeof(int), &m)); + CL_CHECK(clSetKernelArg(gemm, 8, sizeof(int), &n)); + CL_CHECK(clSetKernelArg(gemm, 9, sizeof(int), &k)); + size_t global[] = { 64, (size_t)CEIL_DIV(m, 64), (size_t)CEIL_DIV(n, tile_n) }; + size_t local[] = { 64, 1, 1 }; + backend_ctx->enqueue_ndrange_kernel(gemm, 3, global, local, dst); + CL_CHECK(clReleaseMemObject(activation)); + return true; +} + static bool ggml_cl_mul_mat_iq1_m_dp4a(ggml_backend_t backend, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst) { ggml_backend_opencl_context * backend_ctx = (ggml_backend_opencl_context *)backend->context; ggml_tensor_extra_cl_iq1_m * extra0 = (ggml_tensor_extra_cl_iq1_m *)src0->extra; @@ -28166,6 +29018,210 @@ static void ggml_cl_mul_mat_q3_k_f32_adreno(ggml_backend_t backend, const ggml_t CL_CHECK(clReleaseMemObject(b_sub_buf)); } +static void ggml_cl_mul_mat_iq3_s_f32_adreno(ggml_backend_t backend, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst) { + ggml_backend_opencl_context * backend_ctx = (ggml_backend_opencl_context *)backend->context; + ggml_tensor_extra_cl_iq3_s * extra0 = (ggml_tensor_extra_cl_iq3_s *)src0->extra; + ggml_tensor_extra_cl * extra1 = (ggml_tensor_extra_cl *)src1->extra; + ggml_tensor_extra_cl * extrad = (ggml_tensor_extra_cl *)dst->extra; + + cl_ulong offset1 = extra1->offset + src1->view_offs; + cl_ulong offsetd = extrad->offset + dst->view_offs; + const int m = src0->ne[1]; + const int n = src1->ne[1]; + const int k = src0->ne[0]; + + if (n == 1 || !ggml_is_contiguous(src1) || offset1 % backend_ctx->alignment != 0) { + cl_kernel gemv = backend_ctx->kernel_gemv_noshuffle_iq3_s_f32; + const int stride_d = dst->ne[0]; + const int stride_b = src1->nb[1]/sizeof(float); + CL_CHECK(clSetKernelArg(gemv, 0, sizeof(cl_mem), &backend_ctx->iq3_s_grid_image)); + CL_CHECK(clSetKernelArg(gemv, 1, sizeof(cl_mem), &extra0->qs)); + CL_CHECK(clSetKernelArg(gemv, 2, sizeof(cl_mem), &extra0->sg)); + CL_CHECK(clSetKernelArg(gemv, 3, sizeof(cl_mem), &extra0->qh)); + CL_CHECK(clSetKernelArg(gemv, 4, sizeof(cl_mem), &extra0->sc)); + CL_CHECK(clSetKernelArg(gemv, 5, sizeof(cl_mem), &extra0->d)); + CL_CHECK(clSetKernelArg(gemv, 6, sizeof(cl_mem), &extra1->data_device)); + CL_CHECK(clSetKernelArg(gemv, 7, sizeof(cl_ulong), &offset1)); + CL_CHECK(clSetKernelArg(gemv, 8, sizeof(cl_mem), &extrad->data_device)); + CL_CHECK(clSetKernelArg(gemv, 9, sizeof(cl_ulong), &offsetd)); + CL_CHECK(clSetKernelArg(gemv, 10, sizeof(int), &k)); + CL_CHECK(clSetKernelArg(gemv, 11, sizeof(int), &m)); + CL_CHECK(clSetKernelArg(gemv, 12, sizeof(int), &stride_b)); + CL_CHECK(clSetKernelArg(gemv, 13, sizeof(int), &stride_d)); + const size_t nsg = backend_ctx->iq3_s_mv_nsg; + size_t global[] = { (size_t)CEIL_DIV(m, 128)*64, (size_t)n*nsg, 1 }; + size_t local[] = { 64, nsg, 1 }; + backend_ctx->enqueue_ndrange_kernel(gemv, 3, global, local, dst); + return; + } + + if (ggml_cl_mul_mat_iq3_s_dp4a(backend, src0, src1, dst)) { + return; + } + + cl_context context = backend_ctx->context; + cl_int err; + cl_buffer_region region = { (size_t)offset1, (size_t)k*n*sizeof(float) }; + cl_mem b_sub_buf = clCreateSubBuffer(extra1->data_device, 0, CL_BUFFER_CREATE_TYPE_REGION, ®ion, &err); + CL_CHECK(err); + + cl_image_format img_fmt = { CL_RGBA, CL_FLOAT }; + cl_image_desc img_desc = {}; + img_desc.image_type = CL_MEM_OBJECT_IMAGE1D_BUFFER; + img_desc.image_width = (size_t)k*n/4; + img_desc.buffer = b_sub_buf; + cl_mem b_img = clCreateImage(context, CL_MEM_READ_ONLY, &img_fmt, &img_desc, nullptr, &err); + CL_CHECK(err); + + const int padded_n = CEIL_DIV(n, 8)*8; + region.origin = 0; + region.size = (size_t)k*padded_n*sizeof(cl_half); + backend_ctx->prealloc_act_trans.allocate(context, region.size); + cl_mem b_sub_buf_trans = clCreateSubBuffer(backend_ctx->prealloc_act_trans.buffer, 0, CL_BUFFER_CREATE_TYPE_REGION, ®ion, &err); + CL_CHECK(err); + + img_fmt = { CL_RGBA, CL_HALF_FLOAT }; + img_desc.image_width = (size_t)k*padded_n/4; + img_desc.buffer = b_sub_buf_trans; + cl_mem b_img_trans = clCreateImage(context, 0, &img_fmt, &img_desc, nullptr, &err); + CL_CHECK(err); + + const int height_b = std::max(n/4, 1); + const int width_b = k/4; + const int padded_height_b = padded_n/4; + cl_kernel kernel = backend_ctx->kernel_transpose_32_16; + CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), &b_img)); + CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_mem), &b_img_trans)); + CL_CHECK(clSetKernelArg(kernel, 2, sizeof(int), &height_b)); + CL_CHECK(clSetKernelArg(kernel, 3, sizeof(int), &width_b)); + CL_CHECK(clSetKernelArg(kernel, 4, sizeof(int), &padded_height_b)); + size_t t_global[] = { (size_t)width_b, (size_t)padded_height_b }; + size_t t_local[] = { 1, 16 }; + backend_ctx->enqueue_ndrange_kernel(kernel, 2, t_global, t_local, dst); + + kernel = backend_ctx->kernel_gemm_noshuffle_iq3_s_f32; + CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), &extra0->qs)); + CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_mem), &extra0->sg)); + CL_CHECK(clSetKernelArg(kernel, 2, sizeof(cl_mem), &extra0->qh)); + CL_CHECK(clSetKernelArg(kernel, 3, sizeof(cl_mem), &extra0->sc)); + CL_CHECK(clSetKernelArg(kernel, 4, sizeof(cl_mem), &extra0->d)); + CL_CHECK(clSetKernelArg(kernel, 5, sizeof(cl_mem), &b_img_trans)); + CL_CHECK(clSetKernelArg(kernel, 6, sizeof(cl_mem), &extrad->data_device)); + CL_CHECK(clSetKernelArg(kernel, 7, sizeof(cl_ulong), &offsetd)); + CL_CHECK(clSetKernelArg(kernel, 8, sizeof(int), &m)); + CL_CHECK(clSetKernelArg(kernel, 9, sizeof(int), &padded_n)); + CL_CHECK(clSetKernelArg(kernel, 10, sizeof(int), &k)); + CL_CHECK(clSetKernelArg(kernel, 11, sizeof(int), &n)); + size_t global[] = { (size_t)CEIL_DIV(n, 8), (size_t)CEIL_DIV(m, 512)*128, 1 }; + size_t local[] = { 1, 128, 1 }; + backend_ctx->enqueue_ndrange_kernel(kernel, 3, global, local, dst); + + CL_CHECK(clReleaseMemObject(b_img_trans)); + CL_CHECK(clReleaseMemObject(b_sub_buf_trans)); + CL_CHECK(clReleaseMemObject(b_img)); + CL_CHECK(clReleaseMemObject(b_sub_buf)); +} + +static void ggml_cl_mul_mat_iq3_xxs_f32_adreno(ggml_backend_t backend, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst) { + ggml_backend_opencl_context * backend_ctx = (ggml_backend_opencl_context *)backend->context; + ggml_tensor_extra_cl_iq3_xxs * extra0 = (ggml_tensor_extra_cl_iq3_xxs *)src0->extra; + ggml_tensor_extra_cl * extra1 = (ggml_tensor_extra_cl *)src1->extra; + ggml_tensor_extra_cl * extrad = (ggml_tensor_extra_cl *)dst->extra; + + cl_ulong offset1 = extra1->offset + src1->view_offs; + cl_ulong offsetd = extrad->offset + dst->view_offs; + const int m = src0->ne[1]; + const int n = src1->ne[1]; + const int k = src0->ne[0]; + + if (n == 1 || !ggml_is_contiguous(src1) || offset1 % backend_ctx->alignment != 0) { + cl_kernel gemv = backend_ctx->kernel_gemv_noshuffle_iq3_xxs_f32; + const int stride_d = dst->ne[0]; + const int stride_b = src1->nb[1]/sizeof(float); + CL_CHECK(clSetKernelArg(gemv, 0, sizeof(cl_mem), &backend_ctx->iq3_xxs_grid_image)); + CL_CHECK(clSetKernelArg(gemv, 1, sizeof(cl_mem), &extra0->qs)); + CL_CHECK(clSetKernelArg(gemv, 2, sizeof(cl_mem), &extra0->sas)); + CL_CHECK(clSetKernelArg(gemv, 3, sizeof(cl_mem), &extra0->d)); + CL_CHECK(clSetKernelArg(gemv, 4, sizeof(cl_mem), &extra1->data_device)); + CL_CHECK(clSetKernelArg(gemv, 5, sizeof(cl_ulong), &offset1)); + CL_CHECK(clSetKernelArg(gemv, 6, sizeof(cl_mem), &extrad->data_device)); + CL_CHECK(clSetKernelArg(gemv, 7, sizeof(cl_ulong), &offsetd)); + CL_CHECK(clSetKernelArg(gemv, 8, sizeof(int), &k)); + CL_CHECK(clSetKernelArg(gemv, 9, sizeof(int), &m)); + CL_CHECK(clSetKernelArg(gemv, 10, sizeof(int), &stride_b)); + CL_CHECK(clSetKernelArg(gemv, 11, sizeof(int), &stride_d)); + const size_t nsg = backend_ctx->iq3_xxs_mv_nsg; + size_t global[] = { (size_t)CEIL_DIV(m, 128)*64, (size_t)n*nsg, 1 }; + size_t local[] = { 64, nsg, 1 }; + backend_ctx->enqueue_ndrange_kernel(gemv, 3, global, local, dst); + return; + } + + if (ggml_cl_mul_mat_iq3_xxs_dp4a(backend, src0, src1, dst)) { + return; + } + + cl_context context = backend_ctx->context; + cl_int err; + cl_buffer_region region = { (size_t)offset1, (size_t)k*n*sizeof(float) }; + cl_mem b_sub_buf = clCreateSubBuffer(extra1->data_device, 0, CL_BUFFER_CREATE_TYPE_REGION, ®ion, &err); + CL_CHECK(err); + + cl_image_format img_fmt = { CL_RGBA, CL_FLOAT }; + cl_image_desc img_desc = {}; + img_desc.image_type = CL_MEM_OBJECT_IMAGE1D_BUFFER; + img_desc.image_width = (size_t)k*n/4; + img_desc.buffer = b_sub_buf; + cl_mem b_img = clCreateImage(context, CL_MEM_READ_ONLY, &img_fmt, &img_desc, nullptr, &err); + CL_CHECK(err); + + const int padded_n = CEIL_DIV(n, 8)*8; + region.origin = 0; + region.size = (size_t)k*padded_n*sizeof(cl_half); + backend_ctx->prealloc_act_trans.allocate(context, region.size); + cl_mem b_sub_buf_trans = clCreateSubBuffer(backend_ctx->prealloc_act_trans.buffer, 0, CL_BUFFER_CREATE_TYPE_REGION, ®ion, &err); + CL_CHECK(err); + + img_fmt = { CL_RGBA, CL_HALF_FLOAT }; + img_desc.image_width = (size_t)k*padded_n/4; + img_desc.buffer = b_sub_buf_trans; + cl_mem b_img_trans = clCreateImage(context, 0, &img_fmt, &img_desc, nullptr, &err); + CL_CHECK(err); + + const int height_b = std::max(n/4, 1); + const int width_b = k/4; + const int padded_height_b = padded_n/4; + cl_kernel kernel = backend_ctx->kernel_transpose_32_16; + CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), &b_img)); + CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_mem), &b_img_trans)); + CL_CHECK(clSetKernelArg(kernel, 2, sizeof(int), &height_b)); + CL_CHECK(clSetKernelArg(kernel, 3, sizeof(int), &width_b)); + CL_CHECK(clSetKernelArg(kernel, 4, sizeof(int), &padded_height_b)); + size_t t_global[] = { (size_t)width_b, (size_t)padded_height_b }; + size_t t_local[] = { 1, 16 }; + backend_ctx->enqueue_ndrange_kernel(kernel, 2, t_global, t_local, dst); + + kernel = backend_ctx->kernel_gemm_noshuffle_iq3_xxs_f32; + CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), &extra0->qs)); + CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_mem), &extra0->sas)); + CL_CHECK(clSetKernelArg(kernel, 2, sizeof(cl_mem), &extra0->d)); + CL_CHECK(clSetKernelArg(kernel, 3, sizeof(cl_mem), &b_img_trans)); + CL_CHECK(clSetKernelArg(kernel, 4, sizeof(cl_mem), &extrad->data_device)); + CL_CHECK(clSetKernelArg(kernel, 5, sizeof(cl_ulong), &offsetd)); + CL_CHECK(clSetKernelArg(kernel, 6, sizeof(int), &m)); + CL_CHECK(clSetKernelArg(kernel, 7, sizeof(int), &padded_n)); + CL_CHECK(clSetKernelArg(kernel, 8, sizeof(int), &k)); + CL_CHECK(clSetKernelArg(kernel, 9, sizeof(int), &n)); + size_t global[] = { (size_t)CEIL_DIV(n, 8), (size_t)CEIL_DIV(m, 512)*128, 1 }; + size_t local[] = { 1, 128, 1 }; + backend_ctx->enqueue_ndrange_kernel(kernel, 3, global, local, dst); + + CL_CHECK(clReleaseMemObject(b_img_trans)); + CL_CHECK(clReleaseMemObject(b_sub_buf_trans)); + CL_CHECK(clReleaseMemObject(b_img)); + CL_CHECK(clReleaseMemObject(b_sub_buf)); +} + static void ggml_cl_mul_mat_iq1_m_f32_adreno(ggml_backend_t backend, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst) { ggml_backend_opencl_context * backend_ctx = (ggml_backend_opencl_context *)backend->context; ggml_tensor_extra_cl_iq1_m * extra0 = (ggml_tensor_extra_cl_iq1_m *)src0->extra; @@ -28792,6 +29848,8 @@ static void ggml_cl_mul_mat(ggml_backend_t backend, const ggml_tensor * src0, co ggml_tensor_extra_cl_iq2_xxs * extra0_iq2_xxs = (ggml_tensor_extra_cl_iq2_xxs *)soa0_src->extra; ggml_tensor_extra_cl_iq1_m * extra0_iq1_m = (ggml_tensor_extra_cl_iq1_m *)soa0_src->extra; ggml_tensor_extra_cl_iq1_s * extra0_iq1_s = (ggml_tensor_extra_cl_iq1_s *)soa0_src->extra; + ggml_tensor_extra_cl_iq3_s * extra0_iq3_s = (ggml_tensor_extra_cl_iq3_s *)soa0_src->extra; + ggml_tensor_extra_cl_iq3_xxs * extra0_iq3_xxs = (ggml_tensor_extra_cl_iq3_xxs *)soa0_src->extra; ggml_tensor_extra_cl_q4_K * extra0_q4_K = (ggml_tensor_extra_cl_q4_K *)soa0_src->extra; ggml_tensor_extra_cl_q5_K * extra0_q5_K = (ggml_tensor_extra_cl_q5_K *)soa0_src->extra; ggml_tensor_extra_cl_q6_K * extra0_q6_K = (ggml_tensor_extra_cl_q6_K *)soa0_src->extra; @@ -29205,6 +30263,18 @@ static void ggml_cl_mul_mat(ggml_backend_t backend, const ggml_tensor * src0, co return; } + // iq3_s x fp32 + if (src0t == GGML_TYPE_IQ3_S && src1t == GGML_TYPE_F32 && extra0_iq3_s->noshuffle) { + ggml_cl_mul_mat_iq3_s_f32_adreno(backend, src0, src1, dst); + return; + } + + // iq3_xxs x fp32 + if (src0t == GGML_TYPE_IQ3_XXS && src1t == GGML_TYPE_F32 && extra0_iq3_xxs->noshuffle) { + ggml_cl_mul_mat_iq3_xxs_f32_adreno(backend, src0, src1, dst); + return; + } + // iq1_m x fp32 if (src0t == GGML_TYPE_IQ1_M && src1t == GGML_TYPE_F32 && extra0_iq1_m->noshuffle) { ggml_cl_mul_mat_iq1_m_f32_adreno(backend, src0, src1, dst); @@ -29873,6 +30943,96 @@ static void ggml_cl_mul_mat(ggml_backend_t backend, const ggml_tensor * src0, co backend_ctx->enqueue_ndrange_kernel(kernel, 3, global_work_size, local_work_size, dst); return; } + case GGML_TYPE_IQ3_S: { + if (ne11 < 32) { + break; + } + if (!ggml_is_contiguous(src0) || !ggml_is_contiguous(src1)) { + break; + } + + kernel = backend_ctx->kernel_mul_mm_iq3_s_f32_l4_lm; + nth0 = (backend_ctx->gpu_family == INTEL) ? 64 : 128; + offset0 = src0->view_offs; + + int batch_stride_a = ne00*ne01; + int batch_stride_b = ne10*ne11; + int batch_stride_d = ne0*ne1; + + CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), &extra0_iq3_s->qs)); + CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_mem), &extra0_iq3_s->sg)); + CL_CHECK(clSetKernelArg(kernel, 2, sizeof(cl_mem), &extra0_iq3_s->qh)); + CL_CHECK(clSetKernelArg(kernel, 3, sizeof(cl_mem), &extra0_iq3_s->sc)); + CL_CHECK(clSetKernelArg(kernel, 4, sizeof(cl_mem), &extra0_iq3_s->d)); + CL_CHECK(clSetKernelArg(kernel, 5, sizeof(cl_ulong), &offset0)); + CL_CHECK(clSetKernelArg(kernel, 6, sizeof(cl_mem), &extra1->data_device)); + CL_CHECK(clSetKernelArg(kernel, 7, sizeof(cl_ulong), &offset1)); + CL_CHECK(clSetKernelArg(kernel, 8, sizeof(cl_mem), &extrad->data_device)); + CL_CHECK(clSetKernelArg(kernel, 9, sizeof(cl_ulong), &offsetd)); + CL_CHECK(clSetKernelArg(kernel, 10, sizeof(int), &ne00)); + CL_CHECK(clSetKernelArg(kernel, 11, sizeof(int), &ne01)); + CL_CHECK(clSetKernelArg(kernel, 12, sizeof(int), &ne02)); + CL_CHECK(clSetKernelArg(kernel, 13, sizeof(int), &ne11)); + CL_CHECK(clSetKernelArg(kernel, 14, sizeof(int), &ne12)); + CL_CHECK(clSetKernelArg(kernel, 15, sizeof(int), &ne00)); + CL_CHECK(clSetKernelArg(kernel, 16, sizeof(int), &ne10)); + CL_CHECK(clSetKernelArg(kernel, 17, sizeof(int), &ne0)); + CL_CHECK(clSetKernelArg(kernel, 18, sizeof(int), &batch_stride_a)); + CL_CHECK(clSetKernelArg(kernel, 19, sizeof(int), &batch_stride_b)); + CL_CHECK(clSetKernelArg(kernel, 20, sizeof(int), &batch_stride_d)); + CL_CHECK(clSetKernelArg(kernel, 21, sizeof(int), &r2)); + CL_CHECK(clSetKernelArg(kernel, 22, sizeof(int), &r3)); + + size_t global_work_size[] = {(size_t)(CEIL_DIV(ne01, 64)*nth0), (size_t)CEIL_DIV(ne11, 64), (size_t)ne12*ne13}; + size_t local_work_size[] = {(size_t)nth0, 1, 1}; + + backend_ctx->enqueue_ndrange_kernel(kernel, 3, global_work_size, local_work_size, dst); + return; + } + case GGML_TYPE_IQ3_XXS: { + if (ne11 < 32) { + break; + } + if (!ggml_is_contiguous(src0) || !ggml_is_contiguous(src1)) { + break; + } + + kernel = backend_ctx->kernel_mul_mm_iq3_xxs_f32_l4_lm; + nth0 = (backend_ctx->gpu_family == INTEL) ? 64 : 128; + offset0 = src0->view_offs; + + int batch_stride_a = ne00*ne01; + int batch_stride_b = ne10*ne11; + int batch_stride_d = ne0*ne1; + + CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), &extra0_iq3_xxs->qs)); + CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_mem), &extra0_iq3_xxs->sas)); + CL_CHECK(clSetKernelArg(kernel, 2, sizeof(cl_mem), &extra0_iq3_xxs->d)); + CL_CHECK(clSetKernelArg(kernel, 3, sizeof(cl_ulong), &offset0)); + CL_CHECK(clSetKernelArg(kernel, 4, sizeof(cl_mem), &extra1->data_device)); + CL_CHECK(clSetKernelArg(kernel, 5, sizeof(cl_ulong), &offset1)); + CL_CHECK(clSetKernelArg(kernel, 6, sizeof(cl_mem), &extrad->data_device)); + CL_CHECK(clSetKernelArg(kernel, 7, sizeof(cl_ulong), &offsetd)); + CL_CHECK(clSetKernelArg(kernel, 8, sizeof(int), &ne00)); + CL_CHECK(clSetKernelArg(kernel, 9, sizeof(int), &ne01)); + CL_CHECK(clSetKernelArg(kernel, 10, sizeof(int), &ne02)); + CL_CHECK(clSetKernelArg(kernel, 11, sizeof(int), &ne11)); + CL_CHECK(clSetKernelArg(kernel, 12, sizeof(int), &ne12)); + CL_CHECK(clSetKernelArg(kernel, 13, sizeof(int), &ne00)); + CL_CHECK(clSetKernelArg(kernel, 14, sizeof(int), &ne10)); + CL_CHECK(clSetKernelArg(kernel, 15, sizeof(int), &ne0)); + CL_CHECK(clSetKernelArg(kernel, 16, sizeof(int), &batch_stride_a)); + CL_CHECK(clSetKernelArg(kernel, 17, sizeof(int), &batch_stride_b)); + CL_CHECK(clSetKernelArg(kernel, 18, sizeof(int), &batch_stride_d)); + CL_CHECK(clSetKernelArg(kernel, 19, sizeof(int), &r2)); + CL_CHECK(clSetKernelArg(kernel, 20, sizeof(int), &r3)); + + size_t global_work_size[] = {(size_t)(CEIL_DIV(ne01, 64)*nth0), (size_t)CEIL_DIV(ne11, 64), (size_t)ne12*ne13}; + size_t local_work_size[] = {(size_t)nth0, 1, 1}; + + backend_ctx->enqueue_ndrange_kernel(kernel, 3, global_work_size, local_work_size, dst); + return; + } case GGML_TYPE_IQ1_M: { if (ne11 < 32) { break; @@ -31200,6 +32360,134 @@ static void ggml_cl_mul_mat(ggml_backend_t backend, const ggml_tensor * src0, co #endif // GGML_OPENCL_SOA_Q break; } + case GGML_TYPE_IQ3_XXS: { + if (backend_ctx->gpu_family == INTEL) { + nth0 = 16; + nth1 = 1; + ndst = 4; + } else if (backend_ctx->gpu_family == ADRENO) { + nth0 = 64; + nth1 = 1; + ndst = 4; + } else { + GGML_ASSERT(false && "TODO: Unknown GPU"); + } + +#ifdef GGML_OPENCL_SOA_Q + kernel = backend_ctx->kernel_mul_mv_iq3_xxs_f32_flat; + offset0 = src0->view_offs; + + CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), &extra0_iq3_xxs->qs)); + CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_mem), &extra0_iq3_xxs->sas)); + CL_CHECK(clSetKernelArg(kernel, 2, sizeof(cl_mem), &extra0_iq3_xxs->d)); + CL_CHECK(clSetKernelArg(kernel, 3, sizeof(cl_ulong), &offset0)); + CL_CHECK(clSetKernelArg(kernel, 4, sizeof(cl_mem), &extra1->data_device)); + CL_CHECK(clSetKernelArg(kernel, 5, sizeof(cl_ulong), &offset1)); + CL_CHECK(clSetKernelArg(kernel, 6, sizeof(cl_mem), &extrad->data_device)); + CL_CHECK(clSetKernelArg(kernel, 7, sizeof(cl_ulong), &offsetd)); + CL_CHECK(clSetKernelArg(kernel, 8, sizeof(int), &ne00)); + CL_CHECK(clSetKernelArg(kernel, 9, sizeof(int), &ne01)); + CL_CHECK(clSetKernelArg(kernel, 10, sizeof(cl_ulong), &nb01)); + CL_CHECK(clSetKernelArg(kernel, 11, sizeof(cl_ulong), &nb02)); + CL_CHECK(clSetKernelArg(kernel, 12, sizeof(cl_ulong), &nb03)); + CL_CHECK(clSetKernelArg(kernel, 13, sizeof(int), &ne12)); + CL_CHECK(clSetKernelArg(kernel, 14, sizeof(cl_ulong), &nb11)); + CL_CHECK(clSetKernelArg(kernel, 15, sizeof(cl_ulong), &nb12)); + CL_CHECK(clSetKernelArg(kernel, 16, sizeof(cl_ulong), &nb13)); + CL_CHECK(clSetKernelArg(kernel, 17, sizeof(int), &ne0)); + CL_CHECK(clSetKernelArg(kernel, 18, sizeof(int), &ne1)); + CL_CHECK(clSetKernelArg(kernel, 19, sizeof(int), &r2)); + CL_CHECK(clSetKernelArg(kernel, 20, sizeof(int), &r3)); +#else + kernel = backend_ctx->kernel_mul_mv_iq3_xxs_f32; + + CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), &extra0->data_device)); + CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_ulong), &offset0)); + CL_CHECK(clSetKernelArg(kernel, 2, sizeof(cl_mem), &extra1->data_device)); + CL_CHECK(clSetKernelArg(kernel, 3, sizeof(cl_ulong), &offset1)); + CL_CHECK(clSetKernelArg(kernel, 4, sizeof(cl_mem), &extrad->data_device)); + CL_CHECK(clSetKernelArg(kernel, 5, sizeof(cl_ulong), &offsetd)); + CL_CHECK(clSetKernelArg(kernel, 6, sizeof(int), &ne00)); + CL_CHECK(clSetKernelArg(kernel, 7, sizeof(int), &ne01)); + CL_CHECK(clSetKernelArg(kernel, 8, sizeof(cl_ulong), &nb01)); + CL_CHECK(clSetKernelArg(kernel, 9, sizeof(cl_ulong), &nb02)); + CL_CHECK(clSetKernelArg(kernel, 10, sizeof(cl_ulong), &nb03)); + CL_CHECK(clSetKernelArg(kernel, 11, sizeof(int), &ne12)); + CL_CHECK(clSetKernelArg(kernel, 12, sizeof(cl_ulong), &nb11)); + CL_CHECK(clSetKernelArg(kernel, 13, sizeof(cl_ulong), &nb12)); + CL_CHECK(clSetKernelArg(kernel, 14, sizeof(cl_ulong), &nb13)); + CL_CHECK(clSetKernelArg(kernel, 15, sizeof(int), &ne0)); + CL_CHECK(clSetKernelArg(kernel, 16, sizeof(int), &ne1)); + CL_CHECK(clSetKernelArg(kernel, 17, sizeof(int), &r2)); + CL_CHECK(clSetKernelArg(kernel, 18, sizeof(int), &r3)); +#endif // GGML_OPENCL_SOA_Q + break; + } + case GGML_TYPE_IQ3_S: { + if (backend_ctx->gpu_family == INTEL) { + nth0 = 16; + nth1 = 1; + ndst = 4; + } else if (backend_ctx->gpu_family == ADRENO) { + nth0 = 64; + nth1 = 1; + ndst = 4; + } else { + GGML_ASSERT(false && "TODO: Unknown GPU"); + } + +#ifdef GGML_OPENCL_SOA_Q + kernel = backend_ctx->kernel_mul_mv_iq3_s_f32_flat; + offset0 = src0->view_offs; + + CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), &extra0_iq3_s->qs)); + CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_mem), &extra0_iq3_s->sg)); + CL_CHECK(clSetKernelArg(kernel, 2, sizeof(cl_mem), &extra0_iq3_s->qh)); + CL_CHECK(clSetKernelArg(kernel, 3, sizeof(cl_mem), &extra0_iq3_s->sc)); + CL_CHECK(clSetKernelArg(kernel, 4, sizeof(cl_mem), &extra0_iq3_s->d)); + CL_CHECK(clSetKernelArg(kernel, 5, sizeof(cl_ulong), &offset0)); + CL_CHECK(clSetKernelArg(kernel, 6, sizeof(cl_mem), &extra1->data_device)); + CL_CHECK(clSetKernelArg(kernel, 7, sizeof(cl_ulong), &offset1)); + CL_CHECK(clSetKernelArg(kernel, 8, sizeof(cl_mem), &extrad->data_device)); + CL_CHECK(clSetKernelArg(kernel, 9, sizeof(cl_ulong), &offsetd)); + CL_CHECK(clSetKernelArg(kernel, 10, sizeof(int), &ne00)); + CL_CHECK(clSetKernelArg(kernel, 11, sizeof(int), &ne01)); + CL_CHECK(clSetKernelArg(kernel, 12, sizeof(cl_ulong), &nb01)); + CL_CHECK(clSetKernelArg(kernel, 13, sizeof(cl_ulong), &nb02)); + CL_CHECK(clSetKernelArg(kernel, 14, sizeof(cl_ulong), &nb03)); + CL_CHECK(clSetKernelArg(kernel, 15, sizeof(int), &ne12)); + CL_CHECK(clSetKernelArg(kernel, 16, sizeof(cl_ulong), &nb11)); + CL_CHECK(clSetKernelArg(kernel, 17, sizeof(cl_ulong), &nb12)); + CL_CHECK(clSetKernelArg(kernel, 18, sizeof(cl_ulong), &nb13)); + CL_CHECK(clSetKernelArg(kernel, 19, sizeof(int), &ne0)); + CL_CHECK(clSetKernelArg(kernel, 20, sizeof(int), &ne1)); + CL_CHECK(clSetKernelArg(kernel, 21, sizeof(int), &r2)); + CL_CHECK(clSetKernelArg(kernel, 22, sizeof(int), &r3)); +#else + kernel = backend_ctx->kernel_mul_mv_iq3_s_f32; + + CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), &extra0->data_device)); + CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_ulong), &offset0)); + CL_CHECK(clSetKernelArg(kernel, 2, sizeof(cl_mem), &extra1->data_device)); + CL_CHECK(clSetKernelArg(kernel, 3, sizeof(cl_ulong), &offset1)); + CL_CHECK(clSetKernelArg(kernel, 4, sizeof(cl_mem), &extrad->data_device)); + CL_CHECK(clSetKernelArg(kernel, 5, sizeof(cl_ulong), &offsetd)); + CL_CHECK(clSetKernelArg(kernel, 6, sizeof(int), &ne00)); + CL_CHECK(clSetKernelArg(kernel, 7, sizeof(int), &ne01)); + CL_CHECK(clSetKernelArg(kernel, 8, sizeof(cl_ulong), &nb01)); + CL_CHECK(clSetKernelArg(kernel, 9, sizeof(cl_ulong), &nb02)); + CL_CHECK(clSetKernelArg(kernel, 10, sizeof(cl_ulong), &nb03)); + CL_CHECK(clSetKernelArg(kernel, 11, sizeof(int), &ne12)); + CL_CHECK(clSetKernelArg(kernel, 12, sizeof(cl_ulong), &nb11)); + CL_CHECK(clSetKernelArg(kernel, 13, sizeof(cl_ulong), &nb12)); + CL_CHECK(clSetKernelArg(kernel, 14, sizeof(cl_ulong), &nb13)); + CL_CHECK(clSetKernelArg(kernel, 15, sizeof(int), &ne0)); + CL_CHECK(clSetKernelArg(kernel, 16, sizeof(int), &ne1)); + CL_CHECK(clSetKernelArg(kernel, 17, sizeof(int), &r2)); + CL_CHECK(clSetKernelArg(kernel, 18, sizeof(int), &r3)); +#endif // GGML_OPENCL_SOA_Q + break; + } case GGML_TYPE_IQ1_M: { if (backend_ctx->gpu_family == INTEL) { nth0 = 16; @@ -31825,7 +33113,7 @@ static void ggml_cl_mul_mat(ggml_backend_t backend, const ggml_tensor * src0, co src0t == GGML_TYPE_Q1_0 || src0t == GGML_TYPE_Q2_0 || src0t == GGML_TYPE_IQ4_NL || - src0t == GGML_TYPE_Q2_K || src0t == GGML_TYPE_Q3_K || src0t == GGML_TYPE_IQ2_S || src0t == GGML_TYPE_IQ2_XS || src0t == GGML_TYPE_IQ2_XXS || src0t == GGML_TYPE_IQ1_S || src0t == GGML_TYPE_IQ1_M) { + src0t == GGML_TYPE_Q2_K || src0t == GGML_TYPE_Q3_K || src0t == GGML_TYPE_IQ2_S || src0t == GGML_TYPE_IQ2_XS || src0t == GGML_TYPE_IQ2_XXS || src0t == GGML_TYPE_IQ1_S || src0t == GGML_TYPE_IQ1_M || src0t == GGML_TYPE_IQ3_S || src0t == GGML_TYPE_IQ3_XXS) { // Each SIMD group produces N_DST values in the result. Assuming each // workgroup has N_SIMDGROUP SIMD groups, then each workgroup will // produce N_DST*N_SIMDGROUP values in the result. Hence, the grid size diff --git a/ggml/src/ggml-opencl/kernels/cvt.cl b/ggml/src/ggml-opencl/kernels/cvt.cl index 582d2e59e..143f7c24d 100644 --- a/ggml/src/ggml-opencl/kernels/cvt.cl +++ b/ggml/src/ggml-opencl/kernels/cvt.cl @@ -3178,3 +3178,108 @@ kernel void kernel_restore_block_iq1_m( dst[i].scales[2*j + 1] = (uchar)(sc >> 8); } } + +struct block_iq3_s { + half d; + uchar qs[QK_K/4]; + uchar qh[QK_K/32]; + uchar signs[QK_K/8]; + uchar scales[QK_K/64]; +}; + +kernel void kernel_convert_block_iq3_s( + global struct block_iq3_s * src0, + global uchar * dst_qs, + global uchar * dst_sg, + global uchar * dst_qh, + global uchar * dst_sc, + global half * dst_d, + ulong n_blocks +) { + const ulong i = get_global_id(0); + if (i >= n_blocks) { return; } + dst_d[i] = src0[i].d; + for (int j = 0; j < 64; ++j) { + dst_qs[64*i + j] = src0[i].qs[j]; + } + for (int j = 0; j < 32; ++j) { + dst_sg[32*i + j] = src0[i].signs[j]; + } + for (int j = 0; j < 8; ++j) { + dst_qh[8*i + j] = src0[i].qh[j]; + } + for (int j = 0; j < 4; ++j) { + dst_sc[4*i + j] = src0[i].scales[j]; + } +} + +kernel void kernel_restore_block_iq3_s( + global uchar * src_qs, + global uchar * src_sg, + global uchar * src_qh, + global uchar * src_sc, + global half * src_d, + global struct block_iq3_s * dst, + ulong n_blocks +) { + const ulong i = get_global_id(0); + if (i >= n_blocks) { return; } + dst[i].d = src_d[i]; + for (int j = 0; j < 64; ++j) { + dst[i].qs[j] = src_qs[64*i + j]; + } + for (int j = 0; j < 32; ++j) { + dst[i].signs[j] = src_sg[32*i + j]; + } + for (int j = 0; j < 8; ++j) { + dst[i].qh[j] = src_qh[8*i + j]; + } + for (int j = 0; j < 4; ++j) { + dst[i].scales[j] = src_sc[4*i + j]; + } +} + +struct block_iq3_xxs { + half d; + uchar qs[3*QK_K/8]; +}; + +kernel void kernel_convert_block_iq3_xxs( + global struct block_iq3_xxs * src0, + global uchar * dst_qs, + global uint * dst_sas, + global half * dst_d, + ulong n_blocks +) { + const ulong i = get_global_id(0); + if (i >= n_blocks) { return; } + dst_d[i] = src0[i].d; + for (int j = 0; j < 64; ++j) { + dst_qs[64*i + j] = src0[i].qs[j]; + } + for (int j = 0; j < 8; ++j) { + global ushort * q = (global ushort *)(src0[i].qs + 64 + 4*j); + dst_sas[8*i + j] = (uint)q[0] | ((uint)q[1] << 16); + } +} + +kernel void kernel_restore_block_iq3_xxs( + global uchar * src_qs, + global uint * src_sas, + global half * src_d, + global struct block_iq3_xxs * dst, + ulong n_blocks +) { + const ulong i = get_global_id(0); + if (i >= n_blocks) { return; } + dst[i].d = src_d[i]; + for (int j = 0; j < 64; ++j) { + dst[i].qs[j] = src_qs[64*i + j]; + } + for (int j = 0; j < 8; ++j) { + global ushort * q = (global ushort *)(dst[i].qs + 64 + 4*j); + uint sas = src_sas[8*i + j]; + q[0] = (ushort)sas; + q[1] = (ushort)(sas >> 16); + } +} diff --git a/ggml/src/ggml-opencl/kernels/gemm_noshuffle_iq3_s_f32.cl b/ggml/src/ggml-opencl/kernels/gemm_noshuffle_iq3_s_f32.cl new file mode 100644 index 000000000..1a15ff86c --- /dev/null +++ b/ggml/src/ggml-opencl/kernels/gemm_noshuffle_iq3_s_f32.cl @@ -0,0 +1,92 @@ +#pragma OPENCL EXTENSION cl_khr_fp16 : enable + +#ifdef cl_qcom_reqd_sub_group_size +#pragma OPENCL EXTENSION cl_qcom_reqd_sub_group_size : enable +__attribute__((qcom_reqd_sub_group_size("full"))) +#endif +kernel void kernel_gemm_noshuffle_iq3_s_f32( + global const uchar * src0_qs, + global const uchar * src0_sg, + global const uchar * src0_qh, + global const uchar * src0_sc, + global const half * src0_d, + read_only image1d_buffer_t src1, + global float * dst, + ulong offsetd, + int m, + int n, + int k, + int n_no_padding +) { + dst = (global float *)((global char *)dst + offsetd); + const int gy = get_global_id(0); + const int row = get_global_id(1) * 4; + const int n_4 = n / 4; + if (row >= m) { + return; + } + + half8 c0 = 0, c1 = 0, c2 = 0, c3 = 0; + for (int i = 0; i < k; i += 16) { + const uchar4 sc = vload4(0, src0_sc + row + (i / 64) * m); + const half4 d = vload4(0, src0_d + row + (i / 256) * m); + const float4 scale_f = convert_float4(d) * (1.0f + 2.0f*convert_float4((sc >> (uchar4)(4*((i / 32) & 1))) & (uchar4)(15))); + const half4 scale = convert_half4(scale_f); + + for (int l = 0; l < 16; l += 4) { + const int ki = i + l; + const uint4 qs = convert_uint4(vload4(0, src0_qs + row + (ki / 4) * m)); + const uint4 qh = convert_uint4(vload4(0, src0_qh + row + (ki / 32) * m)); + const uchar4 sg = vload4(0, src0_sg + row + (ki / 8) * m); + const uint4 gi = qs | (((qh >> (uint4)(((ki / 4) & 7))) & (uint4)(1)) << (uint4)(8)); + const uint4 grid = (uint4)(iq3_s_grid[gi.s0], iq3_s_grid[gi.s1], + iq3_s_grid[gi.s2], iq3_s_grid[gi.s3]); + #pragma unroll + for (int j = 0; j < 4; ++j) { + half8 b; + b.s0123 = read_imageh(src1, gy * 2 + (ki + j) * n_4); + b.s4567 = read_imageh(src1, gy * 2 + 1 + (ki + j) * n_4); + const int4 magnitude = convert_int4((grid >> (uint4)(8*j)) & (uint4)(255)); + const int4 signs = convert_int4((sg >> (uchar4)((ki & 7) + j)) & (uchar4)(1)); + const half4 w = convert_half4(magnitude * ((int4)(1) - 2*signs)) * scale; + c0 += b * w.s0; + c1 += b * w.s1; + c2 += b * w.s2; + c3 += b * w.s3; + } + } + } + + int idx = gy * 8 * m + row; + if (idx + 3 < m * n_no_padding) { + vstore4((float4)(c0.s0, c1.s0, c2.s0, c3.s0), 0, dst + idx); + idx += m; + } + if (idx + 3 < m * n_no_padding) { + vstore4((float4)(c0.s1, c1.s1, c2.s1, c3.s1), 0, dst + idx); + idx += m; + } + if (idx + 3 < m * n_no_padding) { + vstore4((float4)(c0.s2, c1.s2, c2.s2, c3.s2), 0, dst + idx); + idx += m; + } + if (idx + 3 < m * n_no_padding) { + vstore4((float4)(c0.s3, c1.s3, c2.s3, c3.s3), 0, dst + idx); + idx += m; + } + if (idx + 3 < m * n_no_padding) { + vstore4((float4)(c0.s4, c1.s4, c2.s4, c3.s4), 0, dst + idx); + idx += m; + } + if (idx + 3 < m * n_no_padding) { + vstore4((float4)(c0.s5, c1.s5, c2.s5, c3.s5), 0, dst + idx); + idx += m; + } + if (idx + 3 < m * n_no_padding) { + vstore4((float4)(c0.s6, c1.s6, c2.s6, c3.s6), 0, dst + idx); + idx += m; + } + if (idx + 3 < m * n_no_padding) { + vstore4((float4)(c0.s7, c1.s7, c2.s7, c3.s7), 0, dst + idx); + } +} diff --git a/ggml/src/ggml-opencl/kernels/gemm_noshuffle_iq3_s_q8_1_dp4a.cl b/ggml/src/ggml-opencl/kernels/gemm_noshuffle_iq3_s_q8_1_dp4a.cl new file mode 100644 index 000000000..43cb38f72 --- /dev/null +++ b/ggml/src/ggml-opencl/kernels/gemm_noshuffle_iq3_s_q8_1_dp4a.cl @@ -0,0 +1,154 @@ +#pragma OPENCL EXTENSION cl_khr_fp16 : enable +#pragma OPENCL EXTENSION cl_khr_subgroups : enable +#ifdef cl_khr_integer_dot_product +#pragma OPENCL EXTENSION cl_khr_integer_dot_product : enable +#endif + +#define QK_K 256 + +#ifndef TILESIZE_N +#define TILESIZE_N 32 +#endif + +inline uint iq3_s_pack(uint gv, uint sg, uint base) { + int v0 = (int)((gv >> 0) & 0xFF); if (sg & (1u << (base + 0))) { v0 = -v0; } + int v1 = (int)((gv >> 8) & 0xFF); if (sg & (1u << (base + 1))) { v1 = -v1; } + int v2 = (int)((gv >> 16) & 0xFF); if (sg & (1u << (base + 2))) { v2 = -v2; } + int v3 = (int)((gv >> 24) & 0xFF); if (sg & (1u << (base + 3))) { v3 = -v3; } + return ((uint)v0 & 0xFFu) | (((uint)v1 & 0xFFu) << 8) + | (((uint)v2 & 0xFFu) << 16) | (((uint)v3 & 0xFFu) << 24); +} + +#define IQ3_S_GEMM_GRID(i) iq3_s_grid[(i)] + +inline uint iq3_s_load(__global const uchar * qs, uint qsb, uint qhb, uint m, + uint u, uint sgv, uint base) { + const uint gidx = (uint)qs[qsb + u * m] | (((qhb >> u) & 1u) << 8); + return iq3_s_pack(IQ3_S_GEMM_GRID(gidx), sgv, base); +} + +// Pass staged operands by value; local dp4a operands miscompile on X2. +inline int dot8_q8a_v(uint8 qw, uint4 a0, uint4 a1) { + int r = 0; + r = dot_acc_sat_4x8packed_ss_int(qw.s0, a0.x, r); + r = dot_acc_sat_4x8packed_ss_int(qw.s1, a0.y, r); + r = dot_acc_sat_4x8packed_ss_int(qw.s2, a0.z, r); + r = dot_acc_sat_4x8packed_ss_int(qw.s3, a0.w, r); + r = dot_acc_sat_4x8packed_ss_int(qw.s4, a1.x, r); + r = dot_acc_sat_4x8packed_ss_int(qw.s5, a1.y, r); + r = dot_acc_sat_4x8packed_ss_int(qw.s6, a1.z, r); + r = dot_acc_sat_4x8packed_ss_int(qw.s7, a1.w, r); + return r; +} + +__attribute__((qcom_wave_pair_mode(1))) +kernel void kernel_gemm_noshuffle_iq3_s_q8_1_dp4a( + __global const uchar * src0_qs, + __global const uchar * src0_sg, + __global const uchar * src0_qh, + __global const uchar * src0_sc, + __global const half * src0_d, + __global const uint * src1_qa, + __global const half * src1_da, + __global float * dst, + ulong offsetd, + int m, + int n_no_padding, + int k +) { + dst = (global float *)((global char *)dst + offsetd); + + const uint lid = get_local_id(0); + const uint block_id_m = get_global_id(1); + const uint block_id_n = get_global_id(2); + + const uint row = block_id_m * 64 + lid; + const uint col_base = block_id_n * TILESIZE_N; + const bool row_valid = row < (uint)m; + const uint rrow = row_valid ? row : 0; + + const uint k_u = (uint)k >> 2; + const uint k_b = (uint)k >> 5; + + __local uint4 sh_qa4[TILESIZE_N][2]; + __local half sh_d[TILESIZE_N]; + +#define NGROUPS (TILESIZE_N / 4) + float4 acc[NGROUPS]; + #pragma unroll + for (int g = 0; g < NGROUPS; ++g) acc[g] = (float4)(0.0f); + + for (uint step = 0; step < (uint)k; step += 32) { + const uint sub = step >> 5; + + const uint ib = sub >> 3; + const uint sb = sub & 7u; + + const uint scb = (uint)src0_sc[rrow + (ib*(QK_K/64) + (sb >> 1)) * (uint)m]; + const uint nib = (sb & 1u) ? (scb >> 4) : (scb & 0xFu); + const float d_w = (float)src0_d[rrow + ib * (uint)m] * (float)(1 + 2*nib); + + const uint qhb = (uint)src0_qh[rrow + (ib*(QK_K/32) + sb) * (uint)m]; + const uint qsb = rrow + (ib*(QK_K/4) + 8*sb) * (uint)m; + const uint sgb = rrow + (ib*(QK_K/8) + 4*sb) * (uint)m; + + const uint sg0 = (uint)src0_sg[sgb + 0 * (uint)m]; + const uint sg1 = (uint)src0_sg[sgb + 1 * (uint)m]; + const uint sg2 = (uint)src0_sg[sgb + 2 * (uint)m]; + const uint sg3 = (uint)src0_sg[sgb + 3 * (uint)m]; + + uint8 qw; + qw.s0 = iq3_s_load(src0_qs, qsb, qhb, (uint)m, 0u, sg0, 0u); + qw.s1 = iq3_s_load(src0_qs, qsb, qhb, (uint)m, 1u, sg0, 4u); + qw.s2 = iq3_s_load(src0_qs, qsb, qhb, (uint)m, 2u, sg1, 0u); + qw.s3 = iq3_s_load(src0_qs, qsb, qhb, (uint)m, 3u, sg1, 4u); + qw.s4 = iq3_s_load(src0_qs, qsb, qhb, (uint)m, 4u, sg2, 0u); + qw.s5 = iq3_s_load(src0_qs, qsb, qhb, (uint)m, 5u, sg2, 4u); + qw.s6 = iq3_s_load(src0_qs, qsb, qhb, (uint)m, 6u, sg3, 0u); + qw.s7 = iq3_s_load(src0_qs, qsb, qhb, (uint)m, 7u, sg3, 4u); + + for (uint idx = lid; idx < TILESIZE_N * 2; idx += 64) { + const uint t = idx >> 1; + const uint v = idx & 1; + const uint c = col_base + t; + sh_qa4[t][v] = (c < (uint)n_no_padding) + ? vload4(0, src1_qa + c * k_u + (step >> 2) + (v << 2)) + : (uint4)(0u); + } + if (lid < TILESIZE_N) { + const uint c = col_base + lid; + sh_d[lid] = (c < (uint)n_no_padding) ? src1_da[c * k_b + sub] : (half)0; + } + barrier(CLK_LOCAL_MEM_FENCE); + +#define LD4(arr, b) ((float4)((float)arr[(b)+0], (float)arr[(b)+1], (float)arr[(b)+2], (float)arr[(b)+3])) + #pragma unroll + for (int g = 0; g < NGROUPS; ++g) { + const int b = g * 4; + float4 rf; +#define TDOT4(T) dot8_q8a_v(qw, (uint4)(sh_qa4[T][0]), (uint4)(sh_qa4[T][1])) + rf.s0 = (float)TDOT4(b+0); rf.s1 = (float)TDOT4(b+1); + rf.s2 = (float)TDOT4(b+2); rf.s3 = (float)TDOT4(b+3); + acc[g] += d_w * LD4(sh_d, b) * rf; + } +#undef TDOT4 +#undef LD4 + barrier(CLK_LOCAL_MEM_FENCE); + } + + if (!row_valid) { + return; + } + + #pragma unroll + for (int g = 0; g < NGROUPS; ++g) { + const uint b = (uint)(g * 4); + const float4 a = acc[g]; + const uint c0 = col_base + b; + if (c0 + 0 < (uint)n_no_padding) dst[(c0 + 0) * (uint)m + row] = a.s0; + if (c0 + 1 < (uint)n_no_padding) dst[(c0 + 1) * (uint)m + row] = a.s1; + if (c0 + 2 < (uint)n_no_padding) dst[(c0 + 2) * (uint)m + row] = a.s2; + if (c0 + 3 < (uint)n_no_padding) dst[(c0 + 3) * (uint)m + row] = a.s3; + } +#undef NGROUPS +} diff --git a/ggml/src/ggml-opencl/kernels/gemm_noshuffle_iq3_xxs_f32.cl b/ggml/src/ggml-opencl/kernels/gemm_noshuffle_iq3_xxs_f32.cl new file mode 100644 index 000000000..af58e850c --- /dev/null +++ b/ggml/src/ggml-opencl/kernels/gemm_noshuffle_iq3_xxs_f32.cl @@ -0,0 +1,89 @@ +#pragma OPENCL EXTENSION cl_khr_fp16 : enable + +#ifdef cl_qcom_reqd_sub_group_size +#pragma OPENCL EXTENSION cl_qcom_reqd_sub_group_size : enable +__attribute__((qcom_reqd_sub_group_size("full"))) +#endif +kernel void kernel_gemm_noshuffle_iq3_xxs_f32( + global const uchar * src0_qs, + global const uint * src0_sas, + global const half * src0_d, + read_only image1d_buffer_t src1, + global float * dst, + ulong offsetd, + int m, + int n, + int k, + int n_no_padding +) { + dst = (global float *)((global char *)dst + offsetd); + const int gy = get_global_id(0); + const int row = get_global_id(1) * 4; + const int n_4 = n / 4; + if (row >= m) { + return; + } + + half8 c0 = 0, c1 = 0, c2 = 0, c3 = 0; + for (int i = 0; i < k; i += 16) { + const uint4 sas = vload4(0, src0_sas + row + (i / 32) * m); + const half4 d = vload4(0, src0_d + row + (i / 256) * m); + const float4 scale_f = convert_float4(d) * 0.5f * (0.5f + convert_float4(sas >> (uint4)(28))); + const half4 scale = convert_half4(scale_f); + + for (int l = 0; l < 16; l += 4) { + const int ki = i + l; + const uint4 gi = convert_uint4(vload4(0, src0_qs + row + (ki / 4) * m)); + const uint4 code = (sas >> (uint4)(7*((ki / 8) & 3))) & (uint4)(127); + const uint4 sg = code | ((popcount(code) & (uint4)(1)) << (uint4)(7)); + const uint4 grid = (uint4)(iq3_xxs_grid[gi.s0], iq3_xxs_grid[gi.s1], + iq3_xxs_grid[gi.s2], iq3_xxs_grid[gi.s3]); + #pragma unroll + for (int j = 0; j < 4; ++j) { + half8 b; + b.s0123 = read_imageh(src1, gy * 2 + (ki + j) * n_4); + b.s4567 = read_imageh(src1, gy * 2 + 1 + (ki + j) * n_4); + const int4 magnitude = convert_int4((grid >> (uint4)(8*j)) & (uint4)(255)); + const int4 signs = convert_int4((sg >> (uint4)((ki & 7) + j)) & (uint4)(1)); + const half4 w = convert_half4(magnitude * ((int4)(1) - 2*signs)) * scale; + c0 += b * w.s0; + c1 += b * w.s1; + c2 += b * w.s2; + c3 += b * w.s3; + } + } + } + + int idx = gy * 8 * m + row; + if (idx + 3 < m * n_no_padding) { + vstore4((float4)(c0.s0, c1.s0, c2.s0, c3.s0), 0, dst + idx); + idx += m; + } + if (idx + 3 < m * n_no_padding) { + vstore4((float4)(c0.s1, c1.s1, c2.s1, c3.s1), 0, dst + idx); + idx += m; + } + if (idx + 3 < m * n_no_padding) { + vstore4((float4)(c0.s2, c1.s2, c2.s2, c3.s2), 0, dst + idx); + idx += m; + } + if (idx + 3 < m * n_no_padding) { + vstore4((float4)(c0.s3, c1.s3, c2.s3, c3.s3), 0, dst + idx); + idx += m; + } + if (idx + 3 < m * n_no_padding) { + vstore4((float4)(c0.s4, c1.s4, c2.s4, c3.s4), 0, dst + idx); + idx += m; + } + if (idx + 3 < m * n_no_padding) { + vstore4((float4)(c0.s5, c1.s5, c2.s5, c3.s5), 0, dst + idx); + idx += m; + } + if (idx + 3 < m * n_no_padding) { + vstore4((float4)(c0.s6, c1.s6, c2.s6, c3.s6), 0, dst + idx); + idx += m; + } + if (idx + 3 < m * n_no_padding) { + vstore4((float4)(c0.s7, c1.s7, c2.s7, c3.s7), 0, dst + idx); + } +} diff --git a/ggml/src/ggml-opencl/kernels/gemm_noshuffle_iq3_xxs_q8_1_dp4a.cl b/ggml/src/ggml-opencl/kernels/gemm_noshuffle_iq3_xxs_q8_1_dp4a.cl new file mode 100644 index 000000000..2ade6e2cf --- /dev/null +++ b/ggml/src/ggml-opencl/kernels/gemm_noshuffle_iq3_xxs_q8_1_dp4a.cl @@ -0,0 +1,137 @@ +#pragma OPENCL EXTENSION cl_khr_fp16 : enable +#pragma OPENCL EXTENSION cl_khr_subgroups : enable +#ifdef cl_khr_integer_dot_product +#pragma OPENCL EXTENSION cl_khr_integer_dot_product : enable +#endif + +#define QK_K 256 + +#ifndef TILESIZE_N +#define TILESIZE_N 32 +#endif + +inline uint iq3_xxs_pack(uint gv, uint sg, uint base) { + int v0 = (int)((gv >> 0) & 0xFF); if (sg & (1u << (base + 0))) { v0 = -v0; } + int v1 = (int)((gv >> 8) & 0xFF); if (sg & (1u << (base + 1))) { v1 = -v1; } + int v2 = (int)((gv >> 16) & 0xFF); if (sg & (1u << (base + 2))) { v2 = -v2; } + int v3 = (int)((gv >> 24) & 0xFF); if (sg & (1u << (base + 3))) { v3 = -v3; } + return ((uint)v0 & 0xFFu) | (((uint)v1 & 0xFFu) << 8) + | (((uint)v2 & 0xFFu) << 16) | (((uint)v3 & 0xFFu) << 24); +} + +// Pass staged operands by value; local dp4a operands miscompile on X2. +inline int dot8_q8a_v(uint8 qw, uint4 a0, uint4 a1) { + int r = 0; + r = dot_acc_sat_4x8packed_ss_int(qw.s0, a0.x, r); + r = dot_acc_sat_4x8packed_ss_int(qw.s1, a0.y, r); + r = dot_acc_sat_4x8packed_ss_int(qw.s2, a0.z, r); + r = dot_acc_sat_4x8packed_ss_int(qw.s3, a0.w, r); + r = dot_acc_sat_4x8packed_ss_int(qw.s4, a1.x, r); + r = dot_acc_sat_4x8packed_ss_int(qw.s5, a1.y, r); + r = dot_acc_sat_4x8packed_ss_int(qw.s6, a1.z, r); + r = dot_acc_sat_4x8packed_ss_int(qw.s7, a1.w, r); + return r; +} + +__attribute__((qcom_wave_pair_mode(1))) +kernel void kernel_gemm_noshuffle_iq3_xxs_q8_1_dp4a( + __global const uchar * src0_qs, + __global const uint * src0_sas, + __global const half * src0_d, + __global const uint * src1_qa, + __global const half * src1_da, + __global float * dst, + ulong offsetd, + int m, + int n_no_padding, + int k +) { + dst = (global float *)((global char *)dst + offsetd); + + const uint lid = get_local_id(0); + const uint block_id_m = get_global_id(1); + const uint block_id_n = get_global_id(2); + + const uint row = block_id_m * 64 + lid; + const uint col_base = block_id_n * TILESIZE_N; + const bool row_valid = row < (uint)m; + const uint rrow = row_valid ? row : 0; + + const uint k_u = (uint)k >> 2; + const uint k_b = (uint)k >> 5; + + __local uint4 sh_qa4[TILESIZE_N][2]; + __local half sh_d[TILESIZE_N]; + +#define NGROUPS (TILESIZE_N / 4) + float4 acc[NGROUPS]; + #pragma unroll + for (int g = 0; g < NGROUPS; ++g) acc[g] = (float4)(0.0f); + + for (uint step = 0; step < (uint)k; step += 32) { + const uint sub = step >> 5; + + const uint ib = sub >> 3; + const uint aux = src0_sas[rrow + sub*(uint)m]; + const float d_w = (float)src0_d[rrow + ib*(uint)m] * (0.5f + (float)(aux >> 28)) * 0.5f; + const uint sg0 = iq3_xxs_signs( aux & 127u); + const uint sg1 = iq3_xxs_signs((aux >> 7) & 127u); + const uint sg2 = iq3_xxs_signs((aux >> 14) & 127u); + const uint sg3 = iq3_xxs_signs((aux >> 21) & 127u); + const uint qsb = rrow + sub*8u*(uint)m; + + uint8 qw; + qw.s0 = iq3_xxs_pack(iq3_xxs_grid[src0_qs[qsb + 0u*(uint)m]], sg0, 0u); + qw.s1 = iq3_xxs_pack(iq3_xxs_grid[src0_qs[qsb + 1u*(uint)m]], sg0, 4u); + qw.s2 = iq3_xxs_pack(iq3_xxs_grid[src0_qs[qsb + 2u*(uint)m]], sg1, 0u); + qw.s3 = iq3_xxs_pack(iq3_xxs_grid[src0_qs[qsb + 3u*(uint)m]], sg1, 4u); + qw.s4 = iq3_xxs_pack(iq3_xxs_grid[src0_qs[qsb + 4u*(uint)m]], sg2, 0u); + qw.s5 = iq3_xxs_pack(iq3_xxs_grid[src0_qs[qsb + 5u*(uint)m]], sg2, 4u); + qw.s6 = iq3_xxs_pack(iq3_xxs_grid[src0_qs[qsb + 6u*(uint)m]], sg3, 0u); + qw.s7 = iq3_xxs_pack(iq3_xxs_grid[src0_qs[qsb + 7u*(uint)m]], sg3, 4u); + + for (uint idx = lid; idx < TILESIZE_N * 2; idx += 64) { + const uint t = idx >> 1; + const uint v = idx & 1; + const uint c = col_base + t; + sh_qa4[t][v] = (c < (uint)n_no_padding) + ? vload4(0, src1_qa + c * k_u + (step >> 2) + (v << 2)) + : (uint4)(0u); + } + if (lid < TILESIZE_N) { + const uint c = col_base + lid; + sh_d[lid] = (c < (uint)n_no_padding) ? src1_da[c * k_b + sub] : (half)0; + } + barrier(CLK_LOCAL_MEM_FENCE); + +#define LD4(arr, b) ((float4)((float)arr[(b)+0], (float)arr[(b)+1], (float)arr[(b)+2], (float)arr[(b)+3])) + #pragma unroll + for (int g = 0; g < NGROUPS; ++g) { + const int b = g * 4; + float4 rf; +#define TDOT4(T) dot8_q8a_v(qw, (uint4)(sh_qa4[T][0]), (uint4)(sh_qa4[T][1])) + rf.s0 = (float)TDOT4(b+0); rf.s1 = (float)TDOT4(b+1); + rf.s2 = (float)TDOT4(b+2); rf.s3 = (float)TDOT4(b+3); + acc[g] += d_w * LD4(sh_d, b) * rf; + } +#undef TDOT4 +#undef LD4 + barrier(CLK_LOCAL_MEM_FENCE); + } + + if (!row_valid) { + return; + } + + #pragma unroll + for (int g = 0; g < NGROUPS; ++g) { + const uint b = (uint)(g * 4); + const float4 a = acc[g]; + const uint c0 = col_base + b; + if (c0 + 0 < (uint)n_no_padding) dst[(c0 + 0) * (uint)m + row] = a.s0; + if (c0 + 1 < (uint)n_no_padding) dst[(c0 + 1) * (uint)m + row] = a.s1; + if (c0 + 2 < (uint)n_no_padding) dst[(c0 + 2) * (uint)m + row] = a.s2; + if (c0 + 3 < (uint)n_no_padding) dst[(c0 + 3) * (uint)m + row] = a.s3; + } +#undef NGROUPS +} diff --git a/ggml/src/ggml-opencl/kernels/gemv_noshuffle_iq3_s_f32.cl b/ggml/src/ggml-opencl/kernels/gemv_noshuffle_iq3_s_f32.cl new file mode 100644 index 000000000..e94b84a86 --- /dev/null +++ b/ggml/src/ggml-opencl/kernels/gemv_noshuffle_iq3_s_f32.cl @@ -0,0 +1,128 @@ +#pragma OPENCL EXTENSION cl_khr_fp16 : enable + +#define QK_K 256 + +#ifndef IQ3_S_MV_NSG +#define IQ3_S_MV_NSG 4 +#endif + +inline float4 iq3_s_vals(uint gv, uint sg, uint base) { + + const uint s = sg >> base; + const uint4 sgn = (uint4)(s << 31, s << 30, s << 29, s << 28) & 0x80000000u; + return as_float4(as_uint4(convert_float4(as_uchar4(gv))) ^ sgn); +} + +kernel void kernel_gemv_noshuffle_iq3_s_f32( + __read_only image1d_buffer_t grid_img, + global const uchar * src0_qs, + global const uchar * src0_sg, + global const uchar * src0_qh, + global const uchar * src0_sc, + global const half * src0_d, + global const float * src1, + ulong offset1, + global float * dst, + ulong offsetd, + int ne00, + int ne01, + int ne10, + int ne0 +) { + src1 = (global const float *)((global const char *)src1 + offset1); + dst = (global float *)((global char *)dst + offsetd); + + const uint m = (uint)ne01; + const uint K = (uint)ne00; + const uint nsb = K / QK_K; + + const uint lid = get_local_id(0); + const uint sgi = get_local_id(1); + const uint col = get_group_id(1); + + global const float * y = src1 + (ulong)col * (uint)ne10; + +#define IQ3_S_GRID(i) (read_imageui(grid_img, (int)(i)).x) + + const uint mh = m >> 1; + const uint j = get_group_id(0) * 64u + lid; + const uint row = j << 1; + + float sumf = 0.f; + float sumf1 = 0.f; + + if (j < mh) { + global const ushort * qsu = (global const ushort *)src0_qs; + global const ushort * qhu = (global const ushort *)src0_qh; + global const ushort * sgu = (global const ushort *)src0_sg; + global const ushort * scu = (global const ushort *)src0_sc; + + for (uint ib = sgi; ib < nsb; ib += IQ3_S_MV_NSG) { + const half2 dh = vload2(j + ib * mh, src0_d); + const float d0 = (float)dh.s0; + const float d1 = (float)dh.s1; + + const uint scbase = j + ib * 4u * mh; + ushort sc4[4]; + sc4[0] = scu[scbase + 0u * mh]; + sc4[1] = scu[scbase + 1u * mh]; + sc4[2] = scu[scbase + 2u * mh]; + sc4[3] = scu[scbase + 3u * mh]; + + float acc0 = 0.f, acc1 = 0.f; + for (uint sb = 0; sb < 8u; ++sb) { + const uint scv = (uint)sc4[sb >> 1]; + const uint nib0 = (sb & 1u) ? ((scv >> 4) & 0xFu) : ( scv & 0xFu); + const uint nib1 = (sb & 1u) ? ((scv >> 12) & 0xFu) : ((scv >> 8) & 0xFu); + + const uint qhv = (uint)qhu[j + (ib * 8u + sb) * mh]; + const uint qh0 = qhv & 0xFFu; + const uint qh1 = qhv >> 8; + + const uint grp = ib * 64u + sb * 8u; + const uint qsb = j + grp * mh; + const uint sgb = j + (ib * 32u + sb * 4u) * mh; + + float a0 = 0.f, a1 = 0.f; + for (uint u = 0; u < 8u; ++u) { + const uint qsv = (uint)qsu[qsb + u * mh]; + const uint sgv = (uint)sgu[sgb + (u >> 1) * mh]; + const uint g0 = ( qsv & 0xFFu) | (((qh0 >> u) & 1u) << 8); + const uint g1 = ((qsv >> 8) & 0xFFu) | (((qh1 >> u) & 1u) << 8); + const uint base = (u & 1u) * 4u; + const float4 yv = vload4(grp + u, y); + const uint gw0 = IQ3_S_GRID(g0); + const uint gw1 = IQ3_S_GRID(g1); + a0 += dot(yv, iq3_s_vals(gw0, sgv & 0xFFu, base)); + a1 += dot(yv, iq3_s_vals(gw1, (sgv >> 8) & 0xFFu, base)); + } + acc0 += (float)(1u + 2u * nib0) * a0; + acc1 += (float)(1u + 2u * nib1) * a1; + } + sumf += d0 * acc0; + sumf1 += d1 * acc1; + } + } + +#if IQ3_S_MV_NSG > 1 + __local float2 part[IQ3_S_MV_NSG][64]; + part[sgi][lid] = (float2)(sumf, sumf1); + barrier(CLK_LOCAL_MEM_FENCE); + if (sgi != 0) { + return; + } + for (uint s = 1; s < IQ3_S_MV_NSG; ++s) { + const float2 p = part[s][lid]; + sumf += p.s0; + sumf1 += p.s1; + } +#endif + + if (j < mh) { + vstore2((float2)(sumf, sumf1), 0, dst + (ulong)col * (uint)ne0 + row); + } +} + +kernel void kernel_iq3_s_grid_export(global uint * out) { + out[get_global_id(0)] = iq3_s_grid[get_global_id(0)]; +} diff --git a/ggml/src/ggml-opencl/kernels/gemv_noshuffle_iq3_xxs_f32.cl b/ggml/src/ggml-opencl/kernels/gemv_noshuffle_iq3_xxs_f32.cl new file mode 100644 index 000000000..c6a5dbeac --- /dev/null +++ b/ggml/src/ggml-opencl/kernels/gemv_noshuffle_iq3_xxs_f32.cl @@ -0,0 +1,120 @@ +#pragma OPENCL EXTENSION cl_khr_fp16 : enable + +#define QK_K 256 + +#ifndef IQ3_XXS_MV_NSG +#define IQ3_XXS_MV_NSG 4 +#endif + +inline float4 iq3_xxs_vals(uint gv, uint sgv, uint base) { + const uint s = sgv >> base; + const uint4 sgn = (uint4)(s << 31, s << 30, s << 29, s << 28) & 0x80000000u; + return as_float4(as_uint4(convert_float4(as_uchar4(gv))) ^ sgn); +} + +kernel void kernel_gemv_noshuffle_iq3_xxs_f32( + __read_only image1d_buffer_t grid_img, + global const uchar * src0_qs, + global const uint * src0_sas, + global const half * src0_d, + global const float * src1, + ulong offset1, + global float * dst, + ulong offsetd, + int ne00, + int ne01, + int ne10, + int ne0 +) { + src1 = (global const float *)((global const char *)src1 + offset1); + dst = (global float *)((global char *)dst + offsetd); + + const uint m = (uint)ne01; + const uint K = (uint)ne00; + const uint nsb = K / QK_K; + + const uint lid = get_local_id(0); + const uint sgi = get_local_id(1); + const uint col = get_group_id(1); + + global const float * y = src1 + (ulong)col * (uint)ne10; + +#define IQ3_XXS_GRID(i) (read_imageui(grid_img, (int)(i)).x) + + const uint mh = m >> 1; + const uint j = get_group_id(0) * 64u + lid; + const uint row = j << 1; + + float sumf = 0.f; + float sumf1 = 0.f; + + if (j < mh) { + global const ushort * qsu = (global const ushort *)src0_qs; + + for (uint ib = sgi; ib < nsb; ib += IQ3_XXS_MV_NSG) { + const half2 dh = vload2(j + ib * mh, src0_d); + const float d0 = (float)dh.s0; + const float d1 = (float)dh.s1; + + float acc0 = 0.f, acc1 = 0.f; + for (uint sb = 0; sb < 8u; ++sb) { + const uint sub = ib * 8u + sb; + const uint2 aux = vload2(j + sub * mh, src0_sas); + + const uint grp = sub * 8u; + const uint qsb = j + grp * mh; + + float a0 = 0.f, a1 = 0.f; + for (uint u = 0; u < 4u; ++u) { + const uint qsv0 = (uint)qsu[qsb + (2u*u)*mh]; + const uint qsv1 = (uint)qsu[qsb + (2u*u + 1u)*mh]; + const uint sg0 = iq3_xxs_signs((aux.s0 >> (7u * u)) & 127u); + const uint sg1 = iq3_xxs_signs((aux.s1 >> (7u * u)) & 127u); + + const uint g00 = qsv0 & 0xFFu; + const uint g01 = qsv1 & 0xFFu; + const uint g10 = (qsv0 >> 8) & 0xFFu; + const uint g11 = (qsv1 >> 8) & 0xFFu; + + const float4 yl = vload4(sub * 8u + u * 2u + 0u, y); + const float4 yh = vload4(sub * 8u + u * 2u + 1u, y); + + a0 += dot(yl, iq3_xxs_vals(IQ3_XXS_GRID(g00), sg0, 0u)) + + dot(yh, iq3_xxs_vals(IQ3_XXS_GRID(g01), sg0, 4u)); + a1 += dot(yl, iq3_xxs_vals(IQ3_XXS_GRID(g10), sg1, 0u)) + + dot(yh, iq3_xxs_vals(IQ3_XXS_GRID(g11), sg1, 4u)); + } + acc0 += (0.5f + (float)(aux.s0 >> 28)) * a0; + acc1 += (0.5f + (float)(aux.s1 >> 28)) * a1; + } + sumf += d0 * 0.5f * acc0; + sumf1 += d1 * 0.5f * acc1; + } + } + +#if IQ3_XXS_MV_NSG > 1 + __local float2 part[IQ3_XXS_MV_NSG][64]; + part[sgi][lid] = (float2)(sumf, sumf1); + barrier(CLK_LOCAL_MEM_FENCE); + if (sgi != 0) { + return; + } + for (uint s = 1; s < IQ3_XXS_MV_NSG; ++s) { + const float2 p = part[s][lid]; + sumf += p.s0; + sumf1 += p.s1; + } +#endif + + if (j < mh) { + vstore2((float2)(sumf, sumf1), 0, dst + (ulong)col * (uint)ne0 + row); + } +#undef IQ3_XXS_GRID +} + +kernel void kernel_iq3_xxs_grid_export(global uint * out) { + const uint i = get_global_id(0); + if (i < 256u) { + out[i] = iq3_xxs_grid[i]; + } +} diff --git a/ggml/src/ggml-opencl/kernels/iq3_s_grid.cl b/ggml/src/ggml-opencl/kernels/iq3_s_grid.cl new file mode 100644 index 000000000..7ffd2c4fe --- /dev/null +++ b/ggml/src/ggml-opencl/kernels/iq3_s_grid.cl @@ -0,0 +1,66 @@ +constant uint iq3_s_grid[512] = { + 0x01010101, 0x01010103, 0x01010105, 0x0101010b, 0x0101010f, 0x01010301, 0x01010303, 0x01010305, + 0x01010309, 0x0101030d, 0x01010501, 0x01010503, 0x0101050b, 0x01010707, 0x01010901, 0x01010905, + 0x0101090b, 0x0101090f, 0x01010b03, 0x01010b07, 0x01010d01, 0x01010d05, 0x01010f03, 0x01010f09, + 0x01010f0f, 0x01030101, 0x01030103, 0x01030105, 0x01030109, 0x01030301, 0x01030303, 0x0103030b, + 0x01030501, 0x01030507, 0x0103050f, 0x01030703, 0x0103070b, 0x01030909, 0x01030d03, 0x01030d0b, + 0x01030f05, 0x01050101, 0x01050103, 0x0105010b, 0x0105010f, 0x01050301, 0x01050307, 0x0105030d, + 0x01050503, 0x0105050b, 0x01050701, 0x01050709, 0x01050905, 0x0105090b, 0x0105090f, 0x01050b03, + 0x01050b07, 0x01050f01, 0x01050f07, 0x01070107, 0x01070303, 0x0107030b, 0x01070501, 0x01070505, + 0x01070703, 0x01070707, 0x0107070d, 0x01070909, 0x01070b01, 0x01070b05, 0x01070d0f, 0x01070f03, + 0x01070f0b, 0x01090101, 0x01090307, 0x0109030f, 0x01090503, 0x01090509, 0x01090705, 0x01090901, + 0x01090907, 0x01090b03, 0x01090f01, 0x010b0105, 0x010b0109, 0x010b0501, 0x010b0505, 0x010b050d, + 0x010b0707, 0x010b0903, 0x010b090b, 0x010b090f, 0x010b0d0d, 0x010b0f07, 0x010d010d, 0x010d0303, + 0x010d0307, 0x010d0703, 0x010d0b05, 0x010d0f03, 0x010f0101, 0x010f0105, 0x010f0109, 0x010f0501, + 0x010f0505, 0x010f050d, 0x010f0707, 0x010f0b01, 0x010f0b09, 0x03010101, 0x03010103, 0x03010105, + 0x03010109, 0x03010301, 0x03010303, 0x03010307, 0x0301030b, 0x0301030f, 0x03010501, 0x03010505, + 0x03010703, 0x03010709, 0x0301070d, 0x03010b09, 0x03010b0d, 0x03010d03, 0x03010f05, 0x03030101, + 0x03030103, 0x03030107, 0x0303010d, 0x03030301, 0x03030309, 0x03030503, 0x03030701, 0x03030707, + 0x03030903, 0x03030b01, 0x03030b05, 0x03030f01, 0x03030f0d, 0x03050101, 0x03050305, 0x0305030b, + 0x0305030f, 0x03050501, 0x03050509, 0x03050705, 0x03050901, 0x03050907, 0x03050b0b, 0x03050d01, + 0x03050f05, 0x03070103, 0x03070109, 0x0307010f, 0x03070301, 0x03070307, 0x03070503, 0x0307050f, + 0x03070701, 0x03070709, 0x03070903, 0x03070d05, 0x03070f01, 0x03090107, 0x0309010b, 0x03090305, + 0x03090309, 0x03090703, 0x03090707, 0x03090905, 0x0309090d, 0x03090b01, 0x03090b09, 0x030b0103, + 0x030b0301, 0x030b0307, 0x030b0503, 0x030b0701, 0x030b0705, 0x030b0b03, 0x030d0501, 0x030d0509, + 0x030d050f, 0x030d0909, 0x030d090d, 0x030f0103, 0x030f0107, 0x030f0301, 0x030f0305, 0x030f0503, + 0x030f070b, 0x030f0903, 0x030f0d05, 0x030f0f01, 0x05010101, 0x05010103, 0x05010107, 0x0501010b, + 0x0501010f, 0x05010301, 0x05010305, 0x05010309, 0x0501030d, 0x05010503, 0x05010507, 0x0501050f, + 0x05010701, 0x05010705, 0x05010903, 0x05010907, 0x0501090b, 0x05010b01, 0x05010b05, 0x05010d0f, + 0x05010f01, 0x05010f07, 0x05010f0b, 0x05030101, 0x05030105, 0x05030301, 0x05030307, 0x0503030f, + 0x05030505, 0x0503050b, 0x05030703, 0x05030709, 0x05030905, 0x05030b03, 0x05050103, 0x05050109, + 0x0505010f, 0x05050503, 0x05050507, 0x05050701, 0x0505070f, 0x05050903, 0x05050b07, 0x05050b0f, + 0x05050f03, 0x05050f09, 0x05070101, 0x05070105, 0x0507010b, 0x05070303, 0x05070505, 0x05070509, + 0x05070703, 0x05070707, 0x05070905, 0x05070b01, 0x05070d0d, 0x05090103, 0x0509010f, 0x05090501, + 0x05090507, 0x05090705, 0x0509070b, 0x05090903, 0x05090f05, 0x05090f0b, 0x050b0109, 0x050b0303, + 0x050b0505, 0x050b070f, 0x050b0901, 0x050b0b07, 0x050b0f01, 0x050d0101, 0x050d0105, 0x050d010f, + 0x050d0503, 0x050d0b0b, 0x050d0d03, 0x050f010b, 0x050f0303, 0x050f050d, 0x050f0701, 0x050f0907, + 0x050f0b01, 0x07010105, 0x07010303, 0x07010307, 0x0701030b, 0x0701030f, 0x07010505, 0x07010703, + 0x07010707, 0x0701070b, 0x07010905, 0x07010909, 0x0701090f, 0x07010b03, 0x07010d07, 0x07010f03, + 0x07030103, 0x07030107, 0x0703010b, 0x07030309, 0x07030503, 0x07030507, 0x07030901, 0x07030d01, + 0x07030f05, 0x07030f0d, 0x07050101, 0x07050305, 0x07050501, 0x07050705, 0x07050709, 0x07050b01, + 0x07070103, 0x07070301, 0x07070309, 0x07070503, 0x07070507, 0x0707050f, 0x07070701, 0x07070903, + 0x07070907, 0x0707090f, 0x07070b0b, 0x07070f07, 0x07090107, 0x07090303, 0x0709030d, 0x07090505, + 0x07090703, 0x07090b05, 0x07090d01, 0x07090d09, 0x070b0103, 0x070b0301, 0x070b0305, 0x070b050b, + 0x070b0705, 0x070b0909, 0x070b0b0d, 0x070b0f07, 0x070d030d, 0x070d0903, 0x070f0103, 0x070f0107, + 0x070f0501, 0x070f0505, 0x070f070b, 0x09010101, 0x09010109, 0x09010305, 0x09010501, 0x09010509, + 0x0901050f, 0x09010705, 0x09010903, 0x09010b01, 0x09010f01, 0x09030105, 0x0903010f, 0x09030303, + 0x09030307, 0x09030505, 0x09030701, 0x0903070b, 0x09030907, 0x09030b03, 0x09030b0b, 0x09050103, + 0x09050107, 0x09050301, 0x0905030b, 0x09050503, 0x09050707, 0x09050901, 0x09050b0f, 0x09050d05, + 0x09050f01, 0x09070109, 0x09070303, 0x09070307, 0x09070501, 0x09070505, 0x09070703, 0x0907070b, + 0x09090101, 0x09090105, 0x09090509, 0x0909070f, 0x09090901, 0x09090f03, 0x090b010b, 0x090b010f, + 0x090b0503, 0x090b0d05, 0x090d0307, 0x090d0709, 0x090d0d01, 0x090f0301, 0x090f030b, 0x090f0701, + 0x090f0907, 0x090f0b03, 0x0b010105, 0x0b010301, 0x0b010309, 0x0b010505, 0x0b010901, 0x0b010909, + 0x0b01090f, 0x0b010b05, 0x0b010d0d, 0x0b010f09, 0x0b030103, 0x0b030107, 0x0b03010b, 0x0b030305, + 0x0b030503, 0x0b030705, 0x0b030f05, 0x0b050101, 0x0b050303, 0x0b050507, 0x0b050701, 0x0b05070d, + 0x0b050b07, 0x0b070105, 0x0b07010f, 0x0b070301, 0x0b07050f, 0x0b070909, 0x0b070b03, 0x0b070d0b, + 0x0b070f07, 0x0b090103, 0x0b090109, 0x0b090501, 0x0b090705, 0x0b09090d, 0x0b0b0305, 0x0b0b050d, + 0x0b0b0b03, 0x0b0b0b07, 0x0b0d0905, 0x0b0f0105, 0x0b0f0109, 0x0b0f0505, 0x0d010303, 0x0d010307, + 0x0d01030b, 0x0d010703, 0x0d010707, 0x0d010d01, 0x0d030101, 0x0d030501, 0x0d03050f, 0x0d030d09, + 0x0d050305, 0x0d050709, 0x0d050905, 0x0d050b0b, 0x0d050d05, 0x0d050f01, 0x0d070101, 0x0d070309, + 0x0d070503, 0x0d070901, 0x0d09050b, 0x0d090907, 0x0d090d05, 0x0d0b0101, 0x0d0b0107, 0x0d0b0709, + 0x0d0b0d01, 0x0d0d010b, 0x0d0d0901, 0x0d0f0303, 0x0d0f0307, 0x0f010101, 0x0f010109, 0x0f01010f, + 0x0f010501, 0x0f010505, 0x0f01070d, 0x0f010901, 0x0f010b09, 0x0f010d05, 0x0f030105, 0x0f030303, + 0x0f030509, 0x0f030907, 0x0f03090b, 0x0f050103, 0x0f050109, 0x0f050301, 0x0f05030d, 0x0f050503, + 0x0f050701, 0x0f050b03, 0x0f070105, 0x0f070705, 0x0f07070b, 0x0f070b07, 0x0f090103, 0x0f09010b, + 0x0f090307, 0x0f090501, 0x0f090b01, 0x0f0b0505, 0x0f0b0905, 0x0f0d0105, 0x0f0d0703, 0x0f0f0101 +}; diff --git a/ggml/src/ggml-opencl/kernels/iq3_xxs_grid.cl b/ggml/src/ggml-opencl/kernels/iq3_xxs_grid.cl new file mode 100644 index 000000000..4c79ddf1a --- /dev/null +++ b/ggml/src/ggml-opencl/kernels/iq3_xxs_grid.cl @@ -0,0 +1,38 @@ +constant uint iq3_xxs_grid[256] = { + 0x04040404, 0x04040414, 0x04040424, 0x04040c0c, 0x04040c1c, 0x04040c3e, 0x04041404, 0x04041414, + 0x04041c0c, 0x04042414, 0x04043e1c, 0x04043e2c, 0x040c040c, 0x040c041c, 0x040c0c04, 0x040c0c14, + 0x040c140c, 0x040c142c, 0x040c1c04, 0x040c1c14, 0x040c240c, 0x040c2c24, 0x040c3e04, 0x04140404, + 0x04140414, 0x04140424, 0x04140c0c, 0x04141404, 0x04141414, 0x04141c0c, 0x04141c1c, 0x04141c3e, + 0x04142c0c, 0x04142c3e, 0x04143e2c, 0x041c040c, 0x041c043e, 0x041c0c04, 0x041c0c14, 0x041c142c, + 0x041c3e04, 0x04240c1c, 0x04241c3e, 0x04242424, 0x04242c3e, 0x04243e1c, 0x04243e2c, 0x042c040c, + 0x042c043e, 0x042c1c14, 0x042c2c14, 0x04341c2c, 0x04343424, 0x043e0c04, 0x043e0c24, 0x043e0c34, + 0x043e241c, 0x043e340c, 0x0c04040c, 0x0c04041c, 0x0c040c04, 0x0c040c14, 0x0c04140c, 0x0c04141c, + 0x0c041c04, 0x0c041c14, 0x0c041c24, 0x0c04243e, 0x0c042c04, 0x0c0c0404, 0x0c0c0414, 0x0c0c0c0c, + 0x0c0c1404, 0x0c0c1414, 0x0c14040c, 0x0c14041c, 0x0c140c04, 0x0c140c14, 0x0c14140c, 0x0c141c04, + 0x0c143e14, 0x0c1c0404, 0x0c1c0414, 0x0c1c1404, 0x0c1c1c0c, 0x0c1c2434, 0x0c1c3434, 0x0c24040c, + 0x0c24042c, 0x0c242c04, 0x0c2c1404, 0x0c2c1424, 0x0c2c2434, 0x0c2c3e0c, 0x0c34042c, 0x0c3e1414, + 0x0c3e2404, 0x14040404, 0x14040414, 0x14040c0c, 0x14040c1c, 0x14041404, 0x14041414, 0x14041434, + 0x14041c0c, 0x14042414, 0x140c040c, 0x140c041c, 0x140c042c, 0x140c0c04, 0x140c0c14, 0x140c140c, + 0x140c1c04, 0x140c341c, 0x140c343e, 0x140c3e04, 0x14140404, 0x14140414, 0x14140c0c, 0x14140c3e, + 0x14141404, 0x14141414, 0x14141c3e, 0x14142404, 0x14142c2c, 0x141c040c, 0x141c0c04, 0x141c0c24, + 0x141c3e04, 0x141c3e24, 0x14241c2c, 0x14242c1c, 0x142c041c, 0x142c143e, 0x142c240c, 0x142c3e24, + 0x143e040c, 0x143e041c, 0x143e0c34, 0x143e242c, 0x1c04040c, 0x1c040c04, 0x1c040c14, 0x1c04140c, + 0x1c04141c, 0x1c042c04, 0x1c04342c, 0x1c043e14, 0x1c0c0404, 0x1c0c0414, 0x1c0c1404, 0x1c0c1c0c, + 0x1c0c2424, 0x1c0c2434, 0x1c14040c, 0x1c14041c, 0x1c140c04, 0x1c14142c, 0x1c142c14, 0x1c143e14, + 0x1c1c0c0c, 0x1c1c1c1c, 0x1c241c04, 0x1c24243e, 0x1c243e14, 0x1c2c0404, 0x1c2c0434, 0x1c2c1414, + 0x1c2c2c2c, 0x1c340c24, 0x1c341c34, 0x1c34341c, 0x1c3e1c1c, 0x1c3e3404, 0x24040424, 0x24040c3e, + 0x24041c2c, 0x24041c3e, 0x24042c1c, 0x24042c3e, 0x240c3e24, 0x24141404, 0x24141c3e, 0x24142404, + 0x24143404, 0x24143434, 0x241c043e, 0x241c242c, 0x24240424, 0x24242c0c, 0x24243424, 0x242c142c, + 0x242c241c, 0x242c3e04, 0x243e042c, 0x243e0c04, 0x243e0c14, 0x243e1c04, 0x2c040c14, 0x2c04240c, + 0x2c043e04, 0x2c0c0404, 0x2c0c0434, 0x2c0c1434, 0x2c0c2c2c, 0x2c140c24, 0x2c141c14, 0x2c143e14, + 0x2c1c0414, 0x2c1c2c1c, 0x2c240c04, 0x2c24141c, 0x2c24143e, 0x2c243e14, 0x2c2c0414, 0x2c2c1c0c, + 0x2c342c04, 0x2c3e1424, 0x2c3e2414, 0x34041424, 0x34042424, 0x34042434, 0x34043424, 0x340c140c, + 0x340c340c, 0x34140c3e, 0x34143424, 0x341c1c04, 0x341c1c34, 0x34242424, 0x342c042c, 0x342c2c14, + 0x34341c1c, 0x343e041c, 0x343e140c, 0x3e04041c, 0x3e04042c, 0x3e04043e, 0x3e040c04, 0x3e041c14, + 0x3e042c14, 0x3e0c1434, 0x3e0c2404, 0x3e140c14, 0x3e14242c, 0x3e142c14, 0x3e1c0404, 0x3e1c0c2c, + 0x3e1c1c1c, 0x3e1c3404, 0x3e24140c, 0x3e24240c, 0x3e2c0404, 0x3e2c0414, 0x3e2c1424, 0x3e341c04 +}; + +inline uint iq3_xxs_signs(uint code) { + return code | ((popcount(code) & 1u) << 7); +} diff --git a/ggml/src/ggml-opencl/kernels/mul_mm_iq3_s_f32_l4_lm.cl b/ggml/src/ggml-opencl/kernels/mul_mm_iq3_s_f32_l4_lm.cl new file mode 100644 index 000000000..e75c296ac --- /dev/null +++ b/ggml/src/ggml-opencl/kernels/mul_mm_iq3_s_f32_l4_lm.cl @@ -0,0 +1,193 @@ +#pragma OPENCL EXTENSION cl_khr_fp16 : enable + +#ifdef cl_intel_required_subgroup_size +#define INTEL_GPU 1 +#endif + +#define QK_K 256 + +struct block_iq3_s { + half d; + uchar qs[QK_K/4]; + uchar qh[QK_K/32]; + uchar signs[QK_K/8]; + uchar scales[QK_K/64]; +}; + +#define LOAD_VEC_A 4 +#define LOAD_VEC_B 4 + +#define BM 64 +#define BN 64 +#define BK 32 +#ifdef INTEL_GPU +#define TM 8 +#define TN 8 +#else +#define TM 4 +#define TN 8 +#endif + +kernel void kernel_mul_mm_iq3_s_f32_l4_lm( + global uchar * src0_qs, + global uchar * src0_sg, + global uchar * src0_qh, + global uchar * src0_sc, + global half * src0_d, + ulong offset0, + global float4 * src1, + ulong offset1, + global float * dst, + ulong offsetd, + + int ne00, + int ne01, + int ne02, + int ne11, + int ne12, + + int stride_a, + int stride_b, + int stride_d, + + int batch_stride_a, + int batch_stride_b, + int batch_stride_d, + + int r2, + int r3 +) { + src1 = (global float4*)((global char*)src1 + offset1); + dst = (global float *)((global char*)dst + offsetd); + + local float buf_a[BM * BK]; + local float buf_b[BN * BK]; + + const int batch_idx = get_global_id(2); + + const int i13 = batch_idx / ne12; + const int i12 = batch_idx % ne12; + + const int i03 = i13 / r3; + const int i02 = i12 / r2; + + const int batch_idx_a = i03 * ne02 + i02; + + const int ir = get_group_id(0); + const int ic = get_group_id(1); + + const int tid = get_local_id(0); + const int th_r = tid % (BM / TM); + const int th_c = tid / (BM / TM); + + const int loadr_a = get_local_id(0) % (BK / LOAD_VEC_A); + const int loadc_a = get_local_id(0) / (BK / LOAD_VEC_A); + const int loadr_b = get_local_id(0) % (BK / LOAD_VEC_B); + const int loadc_b = get_local_id(0) / (BK / LOAD_VEC_B); + + const int loadstride_a = get_local_size(0) * LOAD_VEC_A / BK; + const int loadstride_b = get_local_size(0) * LOAD_VEC_B / BK; + + // pos_a counts elements, not blocks - one iq3_s super block holds QK_K of them + int pos_a = batch_idx_a * batch_stride_a + ir * BM * stride_a; + int pos_b = (batch_idx * batch_stride_b + ic * BN * stride_b) / LOAD_VEC_B; + + float sums[TM * TN]; + float cache_a[TM]; + float cache_b[TN]; + + for (int i = 0; i < TM * TN; i++) { + sums[i] = 0.0f; + } + + for (int block = 0; block < ne00; block += BK) { + for (int l = 0; l < BM; l += loadstride_a) { + if (ir*BM + loadc_a + l < ne01) { + int idx = pos_a + (loadc_a + l) * stride_a + loadr_a * LOAD_VEC_A; + int ib = idx / QK_K; + ib += offset0/sizeof(struct block_iq3_s); + int e = idx % QK_K; + + int ib32 = e >> 5; + int rem = e & 31; + int lg = rem >> 3; + int j0 = rem & 7; + + uchar scb = src0_sc[4*ib + ib32/2]; + float db = (float)src0_d[ib] * (float)(1 + 2*((scb >> (4*(ib32 & 1))) & 15)); + int which = j0 >> 2; + uint gi = (uint)src0_qs[64*ib + 8*ib32 + 2*lg + which] + | (((uint)src0_qh[8*ib + ib32] << (8 - 2*lg - which)) & 256); + uchar sg = src0_sg[32*ib + 4*ib32 + lg]; + uint g = iq3_s_grid[gi]; + + float4 v1; + v1.s0 = db * (float)((g >> 0) & 0xFF) * ((sg & (1 << (j0+0))) ? -1.f : 1.f); + v1.s1 = db * (float)((g >> 8) & 0xFF) * ((sg & (1 << (j0+1))) ? -1.f : 1.f); + v1.s2 = db * (float)((g >> 16) & 0xFF) * ((sg & (1 << (j0+2))) ? -1.f : 1.f); + v1.s3 = db * (float)((g >> 24) & 0xFF) * ((sg & (1 << (j0+3))) ? -1.f : 1.f); + + buf_a[(loadr_a * LOAD_VEC_A + 0) * BM + loadc_a + l] = v1.s0; + buf_a[(loadr_a * LOAD_VEC_A + 1) * BM + loadc_a + l] = v1.s1; + buf_a[(loadr_a * LOAD_VEC_A + 2) * BM + loadc_a + l] = v1.s2; + buf_a[(loadr_a * LOAD_VEC_A + 3) * BM + loadc_a + l] = v1.s3; + } else { + buf_a[(loadr_a * LOAD_VEC_A + 0) * BM + loadc_a + l] = 0.0f; + buf_a[(loadr_a * LOAD_VEC_A + 1) * BM + loadc_a + l] = 0.0f; + buf_a[(loadr_a * LOAD_VEC_A + 2) * BM + loadc_a + l] = 0.0f; + buf_a[(loadr_a * LOAD_VEC_A + 3) * BM + loadc_a + l] = 0.0f; + } + } + + for (int l = 0; l < BN; l += loadstride_b) { + if (ic*BN + loadc_b + l < ne11) { + int idx = pos_b + (loadc_b + l) * stride_b / LOAD_VEC_B + loadr_b; + buf_b[(loadr_b * LOAD_VEC_B + 0) * BN + loadc_b + l] = src1[idx].s0; + buf_b[(loadr_b * LOAD_VEC_B + 1) * BN + loadc_b + l] = src1[idx].s1; + buf_b[(loadr_b * LOAD_VEC_B + 2) * BN + loadc_b + l] = src1[idx].s2; + buf_b[(loadr_b * LOAD_VEC_B + 3) * BN + loadc_b + l] = src1[idx].s3; + } else { + buf_b[(loadr_b * LOAD_VEC_B + 0) * BN + loadc_b + l] = 0.0f; + buf_b[(loadr_b * LOAD_VEC_B + 1) * BN + loadc_b + l] = 0.0f; + buf_b[(loadr_b * LOAD_VEC_B + 2) * BN + loadc_b + l] = 0.0f; + buf_b[(loadr_b * LOAD_VEC_B + 3) * BN + loadc_b + l] = 0.0f; + } + } + + barrier(CLK_LOCAL_MEM_FENCE); + + pos_a += BK; + pos_b += BK / LOAD_VEC_B; + + for (int i = 0; i < BK; i++) { + for (int j = 0; j < TM; j++) { + cache_a[j] = buf_a[(i) * BM + th_r * TM + j]; + } + + for (int j = 0; j < TN; j++) { + cache_b[j] = buf_b[(i) * BN + th_c * TN + j]; + } + + for (int cc = 0; cc < TN; cc++) { + for (int cr = 0; cr < TM; cr++) { + const int sums_idx = cc*TM + cr; + sums[sums_idx] = mad(cache_a[cr], cache_b[cc], sums[sums_idx]); + } + } + } + barrier(CLK_LOCAL_MEM_FENCE); + } + + const int dr = ir * BM + th_r * TM; + const int dc = ic * BN + th_c * TN; + + const ulong offsets = batch_idx * batch_stride_d; + + for (int cc = 0; cc < TN; cc++) { + for (int cr = 0; cr < TM; cr++) { + if (dr + cr < ne01 && dc + cc < ne11) { + dst[offsets + (dc + cc) * stride_d + dr + cr] = sums[cc * TM + cr]; + } + } + } +} diff --git a/ggml/src/ggml-opencl/kernels/mul_mm_iq3_xxs_f32_l4_lm.cl b/ggml/src/ggml-opencl/kernels/mul_mm_iq3_xxs_f32_l4_lm.cl new file mode 100644 index 000000000..3c50d6230 --- /dev/null +++ b/ggml/src/ggml-opencl/kernels/mul_mm_iq3_xxs_f32_l4_lm.cl @@ -0,0 +1,187 @@ +#pragma OPENCL EXTENSION cl_khr_fp16 : enable + +#ifdef cl_intel_required_subgroup_size +#define INTEL_GPU 1 +#endif + +#define QK_K 256 + +struct block_iq3_xxs { + half d; + uchar qs[3*QK_K/8]; +}; + +#define LOAD_VEC_A 4 +#define LOAD_VEC_B 4 + +#define BM 64 +#define BN 64 +#define BK 32 +#ifdef INTEL_GPU +#define TM 8 +#define TN 8 +#else +#define TM 4 +#define TN 8 +#endif + +kernel void kernel_mul_mm_iq3_xxs_f32_l4_lm( + global uchar * src0_qs, + global uint * src0_sas, + global half * src0_d, + ulong offset0, + global float4 * src1, + ulong offset1, + global float * dst, + ulong offsetd, + + int ne00, + int ne01, + int ne02, + int ne11, + int ne12, + + int stride_a, + int stride_b, + int stride_d, + + int batch_stride_a, + int batch_stride_b, + int batch_stride_d, + + int r2, + int r3 +) { + src1 = (global float4*)((global char*)src1 + offset1); + dst = (global float *)((global char*)dst + offsetd); + + local float buf_a[BM * BK]; + local float buf_b[BN * BK]; + + const int batch_idx = get_global_id(2); + + const int i13 = batch_idx / ne12; + const int i12 = batch_idx % ne12; + + const int i03 = i13 / r3; + const int i02 = i12 / r2; + + const int batch_idx_a = i03 * ne02 + i02; + + const int ir = get_group_id(0); + const int ic = get_group_id(1); + + const int tid = get_local_id(0); + const int th_r = tid % (BM / TM); + const int th_c = tid / (BM / TM); + + const int loadr_a = get_local_id(0) % (BK / LOAD_VEC_A); + const int loadc_a = get_local_id(0) / (BK / LOAD_VEC_A); + const int loadr_b = get_local_id(0) % (BK / LOAD_VEC_B); + const int loadc_b = get_local_id(0) / (BK / LOAD_VEC_B); + + const int loadstride_a = get_local_size(0) * LOAD_VEC_A / BK; + const int loadstride_b = get_local_size(0) * LOAD_VEC_B / BK; + + // pos_a counts elements, not blocks - one iq3_xxs super block holds QK_K of them + int pos_a = batch_idx_a * batch_stride_a + ir * BM * stride_a; + int pos_b = (batch_idx * batch_stride_b + ic * BN * stride_b) / LOAD_VEC_B; + + float sums[TM * TN]; + float cache_a[TM]; + float cache_b[TN]; + + for (int i = 0; i < TM * TN; i++) { + sums[i] = 0.0f; + } + + for (int block = 0; block < ne00; block += BK) { + for (int l = 0; l < BM; l += loadstride_a) { + if (ir*BM + loadc_a + l < ne01) { + int idx = pos_a + (loadc_a + l) * stride_a + loadr_a * LOAD_VEC_A; + int ib = idx / QK_K; + ib += offset0/sizeof(struct block_iq3_xxs); + int e = idx % QK_K; + + int ib32 = e >> 5; + int rem = e & 31; + int lg = rem >> 3; + int j0 = rem & 7; + + uint aux = src0_sas[8*ib + ib32]; + float db = (float)src0_d[ib] * (0.5f + (float)(aux >> 28)) * 0.5f; + + uint gi = src0_qs[64*ib + 8*ib32 + 2*lg + (j0 >> 2)]; + uint sg = iq3_xxs_signs((aux >> (7*lg)) & 127u); + uint g = iq3_xxs_grid[gi]; + + float4 v1; + v1.s0 = db * (float)((g >> 0) & 0xFF) * ((sg & (1 << (j0+0))) ? -1.f : 1.f); + v1.s1 = db * (float)((g >> 8) & 0xFF) * ((sg & (1 << (j0+1))) ? -1.f : 1.f); + v1.s2 = db * (float)((g >> 16) & 0xFF) * ((sg & (1 << (j0+2))) ? -1.f : 1.f); + v1.s3 = db * (float)((g >> 24) & 0xFF) * ((sg & (1 << (j0+3))) ? -1.f : 1.f); + + buf_a[(loadr_a * LOAD_VEC_A + 0) * BM + loadc_a + l] = v1.s0; + buf_a[(loadr_a * LOAD_VEC_A + 1) * BM + loadc_a + l] = v1.s1; + buf_a[(loadr_a * LOAD_VEC_A + 2) * BM + loadc_a + l] = v1.s2; + buf_a[(loadr_a * LOAD_VEC_A + 3) * BM + loadc_a + l] = v1.s3; + } else { + buf_a[(loadr_a * LOAD_VEC_A + 0) * BM + loadc_a + l] = 0.0f; + buf_a[(loadr_a * LOAD_VEC_A + 1) * BM + loadc_a + l] = 0.0f; + buf_a[(loadr_a * LOAD_VEC_A + 2) * BM + loadc_a + l] = 0.0f; + buf_a[(loadr_a * LOAD_VEC_A + 3) * BM + loadc_a + l] = 0.0f; + } + } + + for (int l = 0; l < BN; l += loadstride_b) { + if (ic*BN + loadc_b + l < ne11) { + int idx = pos_b + (loadc_b + l) * stride_b / LOAD_VEC_B + loadr_b; + buf_b[(loadr_b * LOAD_VEC_B + 0) * BN + loadc_b + l] = src1[idx].s0; + buf_b[(loadr_b * LOAD_VEC_B + 1) * BN + loadc_b + l] = src1[idx].s1; + buf_b[(loadr_b * LOAD_VEC_B + 2) * BN + loadc_b + l] = src1[idx].s2; + buf_b[(loadr_b * LOAD_VEC_B + 3) * BN + loadc_b + l] = src1[idx].s3; + } else { + buf_b[(loadr_b * LOAD_VEC_B + 0) * BN + loadc_b + l] = 0.0f; + buf_b[(loadr_b * LOAD_VEC_B + 1) * BN + loadc_b + l] = 0.0f; + buf_b[(loadr_b * LOAD_VEC_B + 2) * BN + loadc_b + l] = 0.0f; + buf_b[(loadr_b * LOAD_VEC_B + 3) * BN + loadc_b + l] = 0.0f; + } + } + + barrier(CLK_LOCAL_MEM_FENCE); + + pos_a += BK; + pos_b += BK / LOAD_VEC_B; + + for (int i = 0; i < BK; i++) { + for (int j = 0; j < TM; j++) { + cache_a[j] = buf_a[(i) * BM + th_r * TM + j]; + } + + for (int j = 0; j < TN; j++) { + cache_b[j] = buf_b[(i) * BN + th_c * TN + j]; + } + + for (int cc = 0; cc < TN; cc++) { + for (int cr = 0; cr < TM; cr++) { + const int sums_idx = cc*TM + cr; + sums[sums_idx] = mad(cache_a[cr], cache_b[cc], sums[sums_idx]); + } + } + } + barrier(CLK_LOCAL_MEM_FENCE); + } + + const int dr = ir * BM + th_r * TM; + const int dc = ic * BN + th_c * TN; + + const ulong offsets = batch_idx * batch_stride_d; + + for (int cc = 0; cc < TN; cc++) { + for (int cr = 0; cr < TM; cr++) { + if (dr + cr < ne01 && dc + cc < ne11) { + dst[offsets + (dc + cc) * stride_d + dr + cr] = sums[cc * TM + cr]; + } + } + } +} diff --git a/ggml/src/ggml-opencl/kernels/mul_mv_iq3_s_f32.cl b/ggml/src/ggml-opencl/kernels/mul_mv_iq3_s_f32.cl new file mode 100644 index 000000000..0d9b67b4d --- /dev/null +++ b/ggml/src/ggml-opencl/kernels/mul_mv_iq3_s_f32.cl @@ -0,0 +1,146 @@ +#pragma OPENCL EXTENSION cl_khr_fp16 : enable + +#ifdef cl_intel_required_subgroup_size +#pragma OPENCL EXTENSION cl_intel_required_subgroup_size : enable +#define INTEL_GPU 1 +#define REQD_SUBGROUP_SIZE_16 __attribute__((intel_reqd_sub_group_size(16))) +#define REQD_SUBGROUP_SIZE_32 __attribute__((intel_reqd_sub_group_size(32))) +#elif defined(cl_qcom_reqd_sub_group_size) +#pragma OPENCL EXTENSION cl_qcom_reqd_sub_group_size : enable +#define ADRENO_GPU 1 +#define REQD_SUBGROUP_SIZE_64 __attribute__((qcom_reqd_sub_group_size("half"))) +#define REQD_SUBGROUP_SIZE_128 __attribute__((qcom_reqd_sub_group_size("full"))) +#endif + +#define QK_K 256 +#define IQ3_S_N_SCALE (QK_K/64) + +typedef struct { + half d; + uchar qs[QK_K/4]; + uchar qh[QK_K/32]; + uchar signs[QK_K/8]; + uchar scales[IQ3_S_N_SCALE]; +} block_iq3_s; + +#undef N_DST +#undef N_SIMDGROUP +#undef N_SIMDWIDTH + +#ifdef INTEL_GPU +#define N_DST 4 +#define N_SIMDGROUP 1 +#define N_SIMDWIDTH 16 +#elif defined (ADRENO_GPU) +#define N_DST 4 +#define N_SIMDGROUP 1 +#define N_SIMDWIDTH 64 +#endif + +#undef BLOCK_STRIDE +// 8 threads cover one super block, one 32 element sub block each +#define BLOCK_STRIDE (N_SIMDWIDTH/8) + +#ifdef INTEL_GPU +REQD_SUBGROUP_SIZE_16 +#elif defined (ADRENO_GPU) +REQD_SUBGROUP_SIZE_64 +#endif +kernel void kernel_mul_mv_iq3_s_f32( + global char * src0, + ulong offset0, + global char * src1, + ulong offset1, + global char * dst, + ulong offsetd, + int ne00, + int ne01, + ulong nb01, + ulong nb02, + ulong nb03, + int ne12, + ulong nb11, + ulong nb12, + ulong nb13, + int ne0, + int ne1, + int r2, + int r3 +) { + src0 = src0 + offset0; + src1 = src1 + offset1; + dst = dst + offsetd; + + int ix = get_sub_group_local_id()/8; // super block index + int it = get_sub_group_local_id()%8; // sub block inside the super block + + int nb = ne00/QK_K; + + int r0 = get_group_id(0); + int r1 = get_group_id(1); + int im = get_group_id(2); + int first_row = (r0 * N_SIMDGROUP + get_sub_group_id()) * N_DST; + + int i12 = im%ne12; + int i13 = im/ne12; + + ulong offset_src0 = first_row*nb01 + (i12/r2)*nb02 + (i13/r3)*nb03; + ulong offset_src1 = r1*nb11 + (i12 )*nb12 + (i13 )*nb13; + + global block_iq3_s * x = (global block_iq3_s *) (src0 + offset_src0); + global float * y = (global float *) (src1 + offset_src1); + + float yl[32]; + float sumf[N_DST] = {0.f}; + float all_sum; + + global float * y4 = y + ix * QK_K + 32 * it; + + for (int ib = ix; ib < nb; ib += BLOCK_STRIDE) { + for (int i = 0; i < 32; ++i) { + yl[i] = y4[i]; + } + + global char * xrow = (global char *)(x + ib); + + // only ne01 output rows exist; reading past them can pull in an inf scale + for (int row = 0; row < N_DST && first_row + row < ne01; row++) { + global block_iq3_s * xb = (global block_iq3_s *)(xrow + row*nb01); + + uchar nib = (it & 1) ? (xb->scales[it >> 1] >> 4) : (xb->scales[it >> 1] & 0xf); + float db = (float)xb->d * (float)(1 + 2*nib); + + uchar qhb = xb->qh[it]; + global uchar * qsb = xb->qs + 8*it; + global uchar * sgb = xb->signs + 4*it; + + float acc = 0.f; + for (int l = 0; l < 4; ++l) { + uint i1 = (uint)qsb[2*l+0] | (((uint)qhb << (8-2*l)) & 256); + uint i2 = (uint)qsb[2*l+1] | (((uint)qhb << (7-2*l)) & 256); + uint g1 = iq3_s_grid[i1]; + uint g2 = iq3_s_grid[i2]; + uchar sg = sgb[l]; + for (int j = 0; j < 4; ++j) { + acc += yl[8*l+j+0] * (float)((g1 >> (8*j)) & 0xFF) * ((sg & (1 << j)) ? -1.f : 1.f); + acc += yl[8*l+j+4] * (float)((g2 >> (8*j)) & 0xFF) * ((sg & (1 << (j+4))) ? -1.f : 1.f); + } + } + + sumf[row] += db * acc; + } + + y4 += BLOCK_STRIDE * QK_K; + } + + global float * dst_f32 = (global float *) dst + im*ne0*ne1 + r1*ne0; + + for (int row = 0; row < N_DST; ++row) { + all_sum = sub_group_reduce_add(sumf[row]); + if (first_row + row < ne01) { + if (get_sub_group_local_id() == 0) { + dst_f32[first_row + row] = all_sum; + } + } + } +} diff --git a/ggml/src/ggml-opencl/kernels/mul_mv_iq3_s_f32_flat.cl b/ggml/src/ggml-opencl/kernels/mul_mv_iq3_s_f32_flat.cl new file mode 100644 index 000000000..0208a53c3 --- /dev/null +++ b/ggml/src/ggml-opencl/kernels/mul_mv_iq3_s_f32_flat.cl @@ -0,0 +1,147 @@ +#pragma OPENCL EXTENSION cl_khr_fp16 : enable + +#ifdef cl_intel_required_subgroup_size +#pragma OPENCL EXTENSION cl_intel_required_subgroup_size : enable +#define INTEL_GPU 1 +#define REQD_SUBGROUP_SIZE_16 __attribute__((intel_reqd_sub_group_size(16))) +#define REQD_SUBGROUP_SIZE_32 __attribute__((intel_reqd_sub_group_size(32))) +#elif defined(cl_qcom_reqd_sub_group_size) +#pragma OPENCL EXTENSION cl_qcom_reqd_sub_group_size : enable +#define ADRENO_GPU 1 +#define REQD_SUBGROUP_SIZE_64 __attribute__((qcom_reqd_sub_group_size("half"))) +#define REQD_SUBGROUP_SIZE_128 __attribute__((qcom_reqd_sub_group_size("full"))) +#endif + +#define QK_K 256 +#define IQ3_S_N_SCALE (QK_K/64) + +typedef struct { + half d; + uchar qs[QK_K/4]; + uchar qh[QK_K/32]; + uchar signs[QK_K/8]; + uchar scales[IQ3_S_N_SCALE]; +} block_iq3_s; + +#undef N_DST +#undef N_SIMDGROUP +#undef N_SIMDWIDTH + +#ifdef INTEL_GPU +#define N_DST 4 +#define N_SIMDGROUP 1 +#define N_SIMDWIDTH 16 +#elif defined (ADRENO_GPU) +#define N_DST 4 +#define N_SIMDGROUP 1 +#define N_SIMDWIDTH 64 +#endif + +#undef BLOCK_STRIDE +// 8 threads cover one super block, one 32 element sub block each +#define BLOCK_STRIDE (N_SIMDWIDTH/8) + +#ifdef INTEL_GPU +REQD_SUBGROUP_SIZE_16 +#elif defined (ADRENO_GPU) +REQD_SUBGROUP_SIZE_64 +#endif +kernel void kernel_mul_mv_iq3_s_f32_flat( + global uchar * src0_qs, + global uchar * src0_sg, + global uchar * src0_qh, + global uchar * src0_sc, + global half * src0_d, + ulong offset0, + global char * src1, + ulong offset1, + global char * dst, + ulong offsetd, + int ne00, + int ne01, + ulong nb01, + ulong nb02, + ulong nb03, + int ne12, + ulong nb11, + ulong nb12, + ulong nb13, + int ne0, + int ne1, + int r2, + int r3 +) { + src1 = src1 + offset1; + dst = dst + offsetd; + + int ix = get_sub_group_local_id()/8; // super block index + int it = get_sub_group_local_id()%8; // sub block inside the super block + + int nb = ne00/QK_K; + + int r0 = get_group_id(0); + int r1 = get_group_id(1); + int im = get_group_id(2); + int first_row = (r0 * N_SIMDGROUP + get_sub_group_id()) * N_DST; + + int i12 = im%ne12; + int i13 = im/ne12; + + ulong offset_src0 = (offset0 + first_row*nb01 + (i12/r2)*nb02 + (i13/r3)*nb03)/sizeof(block_iq3_s); + ulong row_stride = nb01/sizeof(block_iq3_s); + ulong offset_src1 = r1*nb11 + (i12 )*nb12 + (i13 )*nb13; + + global float * y = (global float *) (src1 + offset_src1); + + float yl[32]; + float sumf[N_DST] = {0.f}; + float all_sum; + + global float * y4 = y + ix * QK_K + 32 * it; + + for (int ib = ix; ib < nb; ib += BLOCK_STRIDE) { + for (int i = 0; i < 32; ++i) { + yl[i] = y4[i]; + } + + // only ne01 output rows exist; reading past them can pull in an inf scale + for (int row = 0; row < N_DST && first_row + row < ne01; row++) { + const ulong b = offset_src0 + ib + row*row_stride; + + uchar nib = (it & 1) ? (src0_sc[4*b + (it >> 1)] >> 4) : (src0_sc[4*b + (it >> 1)] & 0xf); + float db = (float)src0_d[b] * (float)(1 + 2*nib); + + uchar qhb = src0_qh[8*b + it]; + global uchar * qsb = src0_qs + 64*b + 8*it; + global uchar * sgb = src0_sg + 32*b + 4*it; + + float acc = 0.f; + for (int l = 0; l < 4; ++l) { + uint i1 = (uint)qsb[2*l+0] | (((uint)qhb << (8-2*l)) & 256); + uint i2 = (uint)qsb[2*l+1] | (((uint)qhb << (7-2*l)) & 256); + uint g1 = iq3_s_grid[i1]; + uint g2 = iq3_s_grid[i2]; + uchar sg = sgb[l]; + for (int j = 0; j < 4; ++j) { + acc += yl[8*l+j+0] * (float)((g1 >> (8*j)) & 0xFF) * ((sg & (1 << j)) ? -1.f : 1.f); + acc += yl[8*l+j+4] * (float)((g2 >> (8*j)) & 0xFF) * ((sg & (1 << (j+4))) ? -1.f : 1.f); + } + } + + sumf[row] += db * acc; + } + + y4 += BLOCK_STRIDE * QK_K; + } + + global float * dst_f32 = (global float *) dst + im*ne0*ne1 + r1*ne0; + + for (int row = 0; row < N_DST; ++row) { + all_sum = sub_group_reduce_add(sumf[row]); + if (first_row + row < ne01) { + if (get_sub_group_local_id() == 0) { + dst_f32[first_row + row] = all_sum; + } + } + } +} diff --git a/ggml/src/ggml-opencl/kernels/mul_mv_iq3_xxs_f32.cl b/ggml/src/ggml-opencl/kernels/mul_mv_iq3_xxs_f32.cl new file mode 100644 index 000000000..8355621d1 --- /dev/null +++ b/ggml/src/ggml-opencl/kernels/mul_mv_iq3_xxs_f32.cl @@ -0,0 +1,145 @@ +#pragma OPENCL EXTENSION cl_khr_fp16 : enable + +#ifdef cl_intel_required_subgroup_size +#pragma OPENCL EXTENSION cl_intel_required_subgroup_size : enable +#define INTEL_GPU 1 +#define REQD_SUBGROUP_SIZE_16 __attribute__((intel_reqd_sub_group_size(16))) +#define REQD_SUBGROUP_SIZE_32 __attribute__((intel_reqd_sub_group_size(32))) +#elif defined(cl_qcom_reqd_sub_group_size) +#pragma OPENCL EXTENSION cl_qcom_reqd_sub_group_size : enable +#define ADRENO_GPU 1 +#define REQD_SUBGROUP_SIZE_64 __attribute__((qcom_reqd_sub_group_size("half"))) +#define REQD_SUBGROUP_SIZE_128 __attribute__((qcom_reqd_sub_group_size("full"))) +#endif + +#define QK_K 256 + +// d, then 64 grid index bytes, then 8 uint32 of packed scale+signs +typedef struct { + half d; + uchar qs[3*QK_K/8]; +} block_iq3_xxs; + +#undef N_DST +#undef N_SIMDGROUP +#undef N_SIMDWIDTH + +#ifdef INTEL_GPU +#define N_DST 4 +#define N_SIMDGROUP 1 +#define N_SIMDWIDTH 16 +#elif defined (ADRENO_GPU) +#define N_DST 4 +#define N_SIMDGROUP 1 +#define N_SIMDWIDTH 64 +#endif + +#undef BLOCK_STRIDE +// 8 threads cover one super block, one 32 element sub block each +#define BLOCK_STRIDE (N_SIMDWIDTH/8) + +#ifdef INTEL_GPU +REQD_SUBGROUP_SIZE_16 +#elif defined (ADRENO_GPU) +REQD_SUBGROUP_SIZE_64 +#endif +kernel void kernel_mul_mv_iq3_xxs_f32( + global char * src0, + ulong offset0, + global char * src1, + ulong offset1, + global char * dst, + ulong offsetd, + int ne00, + int ne01, + ulong nb01, + ulong nb02, + ulong nb03, + int ne12, + ulong nb11, + ulong nb12, + ulong nb13, + int ne0, + int ne1, + int r2, + int r3 +) { + src0 = src0 + offset0; + src1 = src1 + offset1; + dst = dst + offsetd; + + int ix = get_sub_group_local_id()/8; // super block index + int it = get_sub_group_local_id()%8; // sub block inside the super block + + int nb = ne00/QK_K; + + int r0 = get_group_id(0); + int r1 = get_group_id(1); + int im = get_group_id(2); + int first_row = (r0 * N_SIMDGROUP + get_sub_group_id()) * N_DST; + + int i12 = im%ne12; + int i13 = im/ne12; + + ulong offset_src0 = first_row*nb01 + (i12/r2)*nb02 + (i13/r3)*nb03; + ulong offset_src1 = r1*nb11 + (i12 )*nb12 + (i13 )*nb13; + + global block_iq3_xxs * x = (global block_iq3_xxs *) (src0 + offset_src0); + global float * y = (global float *) (src1 + offset_src1); + + float yl[32]; + float sumf[N_DST] = {0.f}; + float all_sum; + + global float * y4 = y + ix * QK_K + 32 * it; + + for (int ib = ix; ib < nb; ib += BLOCK_STRIDE) { + for (int i = 0; i < 32; ++i) { + yl[i] = y4[i]; + } + + global char * xrow = (global char *)(x + ib); + + // only ne01 output rows exist; reading past them can pull in an inf scale + for (int row = 0; row < N_DST && first_row + row < ne01; row++) { + global block_iq3_xxs * xb = (global block_iq3_xxs *)(xrow + row*nb01); + + global uchar * qs = xb->qs + 8*it; + + // Blocks are 98 bytes, so packed signs are only 2-byte aligned. + global ushort * sw = (global ushort *)(xb->qs + 64 + 4*it); + uint aux32 = (uint)sw[0] | ((uint)sw[1] << 16); + + float db = (float)xb->d * (0.5f + (float)(aux32 >> 28)) * 0.5f; + + float acc = 0.f; + for (int l = 0; l < 4; ++l) { + uchar signs = iq3_xxs_signs((aux32 >> (7*l)) & 127u); + uint g1 = iq3_xxs_grid[qs[2*l+0]]; + uint g2 = iq3_xxs_grid[qs[2*l+1]]; + + for (int j = 0; j < 4; ++j) { + float s1 = (signs & (1 << j)) ? -1.f : 1.f; + float s2 = (signs & (1 << (j+4))) ? -1.f : 1.f; + acc += yl[8*l+j+0] * (float)((g1 >> (8*j)) & 0xFF) * s1; + acc += yl[8*l+j+4] * (float)((g2 >> (8*j)) & 0xFF) * s2; + } + } + + sumf[row] += db * acc; + } + + y4 += BLOCK_STRIDE * QK_K; + } + + global float * dst_f32 = (global float *) dst + im*ne0*ne1 + r1*ne0; + + for (int row = 0; row < N_DST; ++row) { + all_sum = sub_group_reduce_add(sumf[row]); + if (first_row + row < ne01) { + if (get_sub_group_local_id() == 0) { + dst_f32[first_row + row] = all_sum; + } + } + } +} diff --git a/ggml/src/ggml-opencl/kernels/mul_mv_iq3_xxs_f32_flat.cl b/ggml/src/ggml-opencl/kernels/mul_mv_iq3_xxs_f32_flat.cl new file mode 100644 index 000000000..974daf859 --- /dev/null +++ b/ggml/src/ggml-opencl/kernels/mul_mv_iq3_xxs_f32_flat.cl @@ -0,0 +1,142 @@ +#pragma OPENCL EXTENSION cl_khr_fp16 : enable + +#ifdef cl_intel_required_subgroup_size +#pragma OPENCL EXTENSION cl_intel_required_subgroup_size : enable +#define INTEL_GPU 1 +#define REQD_SUBGROUP_SIZE_16 __attribute__((intel_reqd_sub_group_size(16))) +#define REQD_SUBGROUP_SIZE_32 __attribute__((intel_reqd_sub_group_size(32))) +#elif defined(cl_qcom_reqd_sub_group_size) +#pragma OPENCL EXTENSION cl_qcom_reqd_sub_group_size : enable +#define ADRENO_GPU 1 +#define REQD_SUBGROUP_SIZE_64 __attribute__((qcom_reqd_sub_group_size("half"))) +#define REQD_SUBGROUP_SIZE_128 __attribute__((qcom_reqd_sub_group_size("full"))) +#endif + +#define QK_K 256 + +// d, then 64 grid index bytes, then 8 uint32 of packed scale+signs +typedef struct { + half d; + uchar qs[3*QK_K/8]; +} block_iq3_xxs; + +#undef N_DST +#undef N_SIMDGROUP +#undef N_SIMDWIDTH + +#ifdef INTEL_GPU +#define N_DST 4 +#define N_SIMDGROUP 1 +#define N_SIMDWIDTH 16 +#elif defined (ADRENO_GPU) +#define N_DST 4 +#define N_SIMDGROUP 1 +#define N_SIMDWIDTH 64 +#endif + +#undef BLOCK_STRIDE +// 8 threads cover one super block, one 32 element sub block each +#define BLOCK_STRIDE (N_SIMDWIDTH/8) + +#ifdef INTEL_GPU +REQD_SUBGROUP_SIZE_16 +#elif defined (ADRENO_GPU) +REQD_SUBGROUP_SIZE_64 +#endif +kernel void kernel_mul_mv_iq3_xxs_f32_flat( + global uchar * src0_qs, + global uint * src0_sas, + global half * src0_d, + ulong offset0, + global char * src1, + ulong offset1, + global char * dst, + ulong offsetd, + int ne00, + int ne01, + ulong nb01, + ulong nb02, + ulong nb03, + int ne12, + ulong nb11, + ulong nb12, + ulong nb13, + int ne0, + int ne1, + int r2, + int r3 +) { + src1 = src1 + offset1; + dst = dst + offsetd; + + int ix = get_sub_group_local_id()/8; // super block index + int it = get_sub_group_local_id()%8; // sub block inside the super block + + int nb = ne00/QK_K; + + int r0 = get_group_id(0); + int r1 = get_group_id(1); + int im = get_group_id(2); + int first_row = (r0 * N_SIMDGROUP + get_sub_group_id()) * N_DST; + + int i12 = im%ne12; + int i13 = im/ne12; + + ulong offset_src0 = (offset0 + first_row*nb01 + (i12/r2)*nb02 + (i13/r3)*nb03)/sizeof(block_iq3_xxs); + ulong row_stride = nb01/sizeof(block_iq3_xxs); + ulong offset_src1 = r1*nb11 + (i12 )*nb12 + (i13 )*nb13; + + global float * y = (global float *) (src1 + offset_src1); + + float yl[32]; + float sumf[N_DST] = {0.f}; + float all_sum; + + global float * y4 = y + ix * QK_K + 32 * it; + + for (int ib = ix; ib < nb; ib += BLOCK_STRIDE) { + for (int i = 0; i < 32; ++i) { + yl[i] = y4[i]; + } + + // only ne01 output rows exist; reading past them can pull in an inf scale + for (int row = 0; row < N_DST && first_row + row < ne01; row++) { + const ulong b = offset_src0 + ib + row*row_stride; + + global uchar * qs = src0_qs + 64*b + 8*it; + + uint aux32 = src0_sas[8*b + it]; + + float db = (float)src0_d[b] * (0.5f + (float)(aux32 >> 28)) * 0.5f; + + float acc = 0.f; + for (int l = 0; l < 4; ++l) { + uchar signs = iq3_xxs_signs((aux32 >> (7*l)) & 127u); + uint g1 = iq3_xxs_grid[qs[2*l+0]]; + uint g2 = iq3_xxs_grid[qs[2*l+1]]; + + for (int j = 0; j < 4; ++j) { + float s1 = (signs & (1 << j)) ? -1.f : 1.f; + float s2 = (signs & (1 << (j+4))) ? -1.f : 1.f; + acc += yl[8*l+j+0] * (float)((g1 >> (8*j)) & 0xFF) * s1; + acc += yl[8*l+j+4] * (float)((g2 >> (8*j)) & 0xFF) * s2; + } + } + + sumf[row] += db * acc; + } + + y4 += BLOCK_STRIDE * QK_K; + } + + global float * dst_f32 = (global float *) dst + im*ne0*ne1 + r1*ne0; + + for (int row = 0; row < N_DST; ++row) { + all_sum = sub_group_reduce_add(sumf[row]); + if (first_row + row < ne01) { + if (get_sub_group_local_id() == 0) { + dst_f32[first_row + row] = all_sum; + } + } + } +}