diff --git a/.github/workflows/build.yml b/.github/workflows/build.yml index 99fc8b0777..7df1e25867 100644 --- a/.github/workflows/build.yml +++ b/.github/workflows/build.yml @@ -1727,6 +1727,22 @@ jobs: vulkaninfo --summary GG_BUILD_VULKAN=1 bash ./ci/run.sh ~/results/llama.cpp ~/mnt/llama.cpp + ggml-ci-x64-linux-intel-vulkan: + runs-on: [self-hosted, Linux, X64, Intel] + + steps: + - name: Clone + id: checkout + uses: actions/checkout@v6 + with: + persist-credentials: false + + - name: Test + id: ggml-ci + run: | + vulkaninfo --summary + GG_BUILD_VULKAN=1 bash ./ci/run.sh ~/results/llama.cpp ~/mnt/llama.cpp + ggml-ci-arm64-cpu-kleidiai: runs-on: ubuntu-22.04-arm diff --git a/common/arg.cpp b/common/arg.cpp index 41da8563d6..10aa1b5e4f 100644 --- a/common/arg.cpp +++ b/common/arg.cpp @@ -732,23 +732,28 @@ static void common_params_print_completion(common_params_context & ctx_arg) { "llama-completion", "llama-convert-llama2c-to-ggml", "llama-cvector-generator", + "llama-debug", + "llama-diffusion-cli", "llama-embedding", "llama-eval-callback", "llama-export-lora", + "llama-finetune", + "llama-fit-params", + "llama-gemma3-cli", "llama-gen-docs", "llama-gguf", "llama-gguf-hash", "llama-gguf-split", - "llama-gritlm", + "llama-idle", "llama-imatrix", - "llama-infill", - "llama-mtmd-cli", - "llama-llava-clip-quantize-cli", + "llama-llava-cli", "llama-lookahead", "llama-lookup", "llama-lookup-create", "llama-lookup-merge", "llama-lookup-stats", + "llama-minicpmv-cli", + "llama-mtmd-cli", "llama-parallel", "llama-passkey", "llama-perplexity", @@ -2666,7 +2671,8 @@ common_params_context common_params_parser_init(common_params & params, llama_ex [](common_params & params, const std::string & value) { params.out_file = value; } - ).set_examples({LLAMA_EXAMPLE_IMATRIX, LLAMA_EXAMPLE_CVECTOR_GENERATOR, LLAMA_EXAMPLE_EXPORT_LORA, LLAMA_EXAMPLE_TTS, LLAMA_EXAMPLE_FINETUNE, LLAMA_EXAMPLE_RESULTS})); + ).set_examples({LLAMA_EXAMPLE_IMATRIX, LLAMA_EXAMPLE_CVECTOR_GENERATOR, LLAMA_EXAMPLE_EXPORT_LORA, LLAMA_EXAMPLE_TTS, LLAMA_EXAMPLE_FINETUNE, + LLAMA_EXAMPLE_RESULTS, LLAMA_EXAMPLE_EXPORT_GRAPH_OPS})); add_opt(common_arg( {"-ofreq", "--output-frequency"}, "N", string_format("output the imatrix every N iterations (default: %d)", params.n_out_freq), diff --git a/common/common.h b/common/common.h index c5645bba46..ee7a2d805e 100644 --- a/common/common.h +++ b/common/common.h @@ -105,6 +105,7 @@ enum llama_example { LLAMA_EXAMPLE_FINETUNE, LLAMA_EXAMPLE_FIT_PARAMS, LLAMA_EXAMPLE_RESULTS, + LLAMA_EXAMPLE_EXPORT_GRAPH_OPS, LLAMA_EXAMPLE_COUNT, }; diff --git a/convert_hf_to_gguf.py b/convert_hf_to_gguf.py index 6e3d3ad1dc..37834c78b8 100755 --- a/convert_hf_to_gguf.py +++ b/convert_hf_to_gguf.py @@ -10092,9 +10092,9 @@ class NemotronHModel(GraniteHybridModel): # Skip Multi-Token Prediction (MTP) tensors. These are used for # for speculative decoding but we don't include them in this model # conversion. See https://github.com/ggml-org/llama.cpp/pull/18886 - if "mtp" in name: + if name.startswith("mtp."): logger.info(f"gguf: Skipping MTP (Speculative) layer: {name}") - return [] + return if name.endswith("mixer.gate.e_score_correction_bias"): new_name = name.replace("e_score_correction_bias", "e_score_correction.bias") diff --git a/docs/backend/VirtGPU/development.md b/docs/backend/VirtGPU/development.md index ca2e47772a..c5c85dfb49 100644 --- a/docs/backend/VirtGPU/development.md +++ b/docs/backend/VirtGPU/development.md @@ -55,7 +55,8 @@ LLAMA_MAC_BUILD=$PWD/build/ggml-virtgpu-backend cmake -S . -B $LLAMA_MAC_BUILD \ -DGGML_NATIVE=OFF \ -DLLAMA_CURL=ON \ - -DGGML_REMOTINGBACKEND=ONLY \ + -DGGML_VIRTGPU=ON \ + -DGGML_VIRTGPU_BACKEND=ONLY \ -DGGML_METAL=ON TARGETS="ggml-metal" @@ -71,6 +72,7 @@ cmake --build $LLAMA_MAC_BUILD --parallel 8 --target $EXTRA_TARGETS ```bash # Build virglrenderer with APIR support mkdir virglrenderer +cd virglrenderer git clone https://gitlab.freedesktop.org/kpouget/virglrenderer -b main-macos src cd src @@ -95,7 +97,7 @@ mkdir llama.cpp git clone https://github.com/ggml-org/llama.cpp.git src cd src -LLAMA_LINUX_BUILD=$PWD//build-virtgpu +LLAMA_LINUX_BUILD=$PWD/build-virtgpu cmake -S . -B $LLAMA_LINUX_BUILD \ -DGGML_VIRTGPU=ON diff --git a/ggml/src/ggml-backend.cpp b/ggml/src/ggml-backend.cpp index bc57df20ba..22c656996c 100644 --- a/ggml/src/ggml-backend.cpp +++ b/ggml/src/ggml-backend.cpp @@ -1455,10 +1455,6 @@ static enum ggml_status ggml_backend_sched_compute_splits(ggml_backend_sched_t s int split_backend_id = split->backend_id; ggml_backend_t split_backend = sched->backends[split_backend_id]; - if (sched->events[split_backend_id][sched->cur_copy] == NULL) { - ggml_backend_synchronize(split_backend); - } - // copy the input tensors to the split backend for (int input_id = 0; input_id < split->n_inputs; input_id++) { ggml_backend_t input_backend = ggml_backend_sched_get_tensor_backend(sched, split->inputs[input_id]); @@ -1469,12 +1465,16 @@ static enum ggml_status ggml_backend_sched_compute_splits(ggml_backend_sched_t s // inputs from the user must be copied immediately to prevent the user overwriting the data before the copy is done if (sched->events[split_backend_id][sched->cur_copy] != NULL) { ggml_backend_event_synchronize(sched->events[split_backend_id][sched->cur_copy]); + } else { + ggml_backend_synchronize(split_backend); } - ggml_backend_tensor_copy_async(input_backend, split_backend, input, input_cpy); + ggml_backend_tensor_copy(input, input_cpy); } else { // wait for the split backend to finish using the input before overwriting it if (sched->events[split_backend_id][sched->cur_copy] != NULL) { ggml_backend_event_wait(split_backend, sched->events[split_backend_id][sched->cur_copy]); + } else { + ggml_backend_synchronize(split_backend); } // when offloading MoE weights, we can reduce the amount of data copied by copying only the experts that are used @@ -1578,10 +1578,6 @@ static enum ggml_status ggml_backend_sched_compute_splits(ggml_backend_sched_t s } } - if (sched->events[split_backend_id][sched->cur_copy] == NULL) { - ggml_backend_synchronize(split_backend); - } - if (!sched->callback_eval) { enum ggml_status ec = ggml_backend_graph_compute_async(split_backend, &split->graph); if (ec != GGML_STATUS_SUCCESS) { diff --git a/ggml/src/ggml-cuda/ggml-cuda.cu b/ggml/src/ggml-cuda/ggml-cuda.cu index cda275b8c5..9d2aacf4b2 100644 --- a/ggml/src/ggml-cuda/ggml-cuda.cu +++ b/ggml/src/ggml-cuda/ggml-cuda.cu @@ -2823,14 +2823,11 @@ static bool ggml_backend_cuda_cpy_tensor_async(ggml_backend_t backend_src, ggml_ ggml_backend_buffer_t buf_src = src->view_src ? src->view_src->buffer : src->buffer; ggml_backend_buffer_t buf_dst = dst->view_src ? dst->view_src->buffer : dst->buffer; - //enables async copies from CPU to CUDA, instead of only CUDA-to-CUDA - bool copy_from_host = ggml_backend_buffer_is_host(buf_src) && ggml_backend_dev_type(backend_src->device) == GGML_BACKEND_DEVICE_TYPE_CPU; - - if (!(copy_from_host || ggml_backend_is_cuda(backend_src)) || !ggml_backend_is_cuda(backend_dst)) { + if (!ggml_backend_is_cuda(backend_src) || !ggml_backend_is_cuda(backend_dst)) { return false; } - if (!(copy_from_host || ggml_backend_buffer_is_cuda(buf_src)) || !ggml_backend_buffer_is_cuda(dst->buffer)) { + if (!ggml_backend_buffer_is_cuda(src->buffer) || !ggml_backend_buffer_is_cuda(dst->buffer)) { return false; } @@ -2841,17 +2838,14 @@ static bool ggml_backend_cuda_cpy_tensor_async(ggml_backend_t backend_src, ggml_ ggml_backend_cuda_buffer_context * buf_ctx_src = (ggml_backend_cuda_buffer_context *)buf_src->context; ggml_backend_cuda_buffer_context * buf_ctx_dst = (ggml_backend_cuda_buffer_context *)buf_dst->context; - if ((copy_from_host && cuda_ctx_dst->device != buf_ctx_dst->device) || - !copy_from_host && (cuda_ctx_src->device != buf_ctx_src->device || cuda_ctx_dst->device != buf_ctx_dst->device)) { + if (cuda_ctx_src->device != buf_ctx_src->device || cuda_ctx_dst->device != buf_ctx_dst->device) { #ifndef NDEBUG GGML_LOG_DEBUG("%s: backend and buffer devices do not match\n", __func__); #endif return false; } - if (copy_from_host) { - CUDA_CHECK(cudaMemcpyAsync(dst->data, src->data, ggml_nbytes(dst), cudaMemcpyHostToDevice, cuda_ctx_dst->stream())); - } else if (backend_src != backend_dst) { + if (backend_src != backend_dst) { // copy on src stream if (cuda_ctx_src->device == cuda_ctx_dst->device) { CUDA_CHECK(cudaMemcpyAsync(dst->data, src->data, ggml_nbytes(dst), cudaMemcpyDeviceToDevice, cuda_ctx_src->stream())); diff --git a/ggml/src/ggml-metal/ggml-metal-context.m b/ggml/src/ggml-metal/ggml-metal-context.m index 855fd1adae..32d97cd5d0 100644 --- a/ggml/src/ggml-metal/ggml-metal-context.m +++ b/ggml/src/ggml-metal/ggml-metal-context.m @@ -554,7 +554,7 @@ enum ggml_status ggml_metal_graph_compute(ggml_metal_t ctx, struct ggml_cgraph * // enter here only when capturing in order to wait for all computation to finish // otherwise, we leave the graph to compute asynchronously - if (!use_capture && ctx->capture_started) { + if (use_capture && ctx->capture_started) { // wait for completion and check status of each command buffer // needed to detect if the device ran out-of-memory for example (#1881) { @@ -606,6 +606,8 @@ enum ggml_status ggml_metal_graph_compute(ggml_metal_t ctx, struct ggml_cgraph * [ctx->capture_scope endScope]; [[MTLCaptureManager sharedCaptureManager] stopCapture]; + + ctx->capture_started = false; } } diff --git a/ggml/src/ggml-metal/ggml-metal-device.cpp b/ggml/src/ggml-metal/ggml-metal-device.cpp index 15ae2e517d..72ad876d5e 100644 --- a/ggml/src/ggml-metal/ggml-metal-device.cpp +++ b/ggml/src/ggml-metal/ggml-metal-device.cpp @@ -1470,10 +1470,11 @@ ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_bin(ggml_metal_l const bool is_c4 = (op->src[0]->ne[0] % 4 == 0) && (op->src[1]->ne[0] % 4 == 0); + const bool is_cb = op->src[0]->ne[0] != op->src[1]->ne[0]; const bool is_rb = ggml_is_contiguous(op->src[0]) && ggml_is_contiguous(op->src[1]) && (ggml_nrows(op->src[1]) == 1) && ggml_nelements(op) < 65536; snprintf(base, 256, "kernel_bin_fuse_%s_%s_%s%s", t0_str, t1_str, t_str, is_c4 ? "_4" : ""); - snprintf(name, 256, "%s_op=%d_nf=%d_rb=%d", base, op_num, n_fuse, is_rb); + snprintf(name, 256, "%s_op=%d_nf=%d_rb=%d_cb=%d", base, op_num, n_fuse, is_rb, is_cb); ggml_metal_pipeline_with_params res = ggml_metal_library_get_pipeline(lib, name); if (!res.pipeline) { @@ -1482,6 +1483,7 @@ ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_bin(ggml_metal_l ggml_metal_cv_set_int16(cv, op_num, FC_BIN + 0); ggml_metal_cv_set_int16(cv, n_fuse, FC_BIN + 1); ggml_metal_cv_set_bool (cv, is_rb, FC_BIN + 2); + ggml_metal_cv_set_bool (cv, is_cb, FC_BIN + 3); res = ggml_metal_library_compile_pipeline(lib, base, name, cv); diff --git a/ggml/src/ggml-metal/ggml-metal-ops.cpp b/ggml/src/ggml-metal/ggml-metal-ops.cpp index 306dbcf366..c0bcad392b 100644 --- a/ggml/src/ggml-metal/ggml-metal-ops.cpp +++ b/ggml/src/ggml-metal/ggml-metal-ops.cpp @@ -3180,9 +3180,7 @@ int ggml_metal_op_bin(ggml_metal_op_t ctx, int idx) { ggml_metal_encoder_set_buffer (enc, bid_dst, 3); if (pipeline.cnt) { - const int n = pipeline.c4 ? ggml_nelements(op)/4 : ggml_nelements(op); - - ggml_metal_encoder_dispatch_threadgroups(enc, n, 1, 1, 1, 1, 1); + ggml_metal_encoder_dispatch_threadgroups(enc, args.ne0, ggml_nrows(op), 1, 1, 1, 1); } else { const int nth_max = MIN(256, ggml_metal_pipeline_max_theads_per_threadgroup(pipeline)); diff --git a/ggml/src/ggml-metal/ggml-metal.metal b/ggml/src/ggml-metal/ggml-metal.metal index 0b77d5349b..24a3092af2 100644 --- a/ggml/src/ggml-metal/ggml-metal.metal +++ b/ggml/src/ggml-metal/ggml-metal.metal @@ -1111,6 +1111,7 @@ template [[host_name("kernel_unary_f16_f16_4")]] kernel kernel_unary_t kernel_un constant short FC_bin_op [[function_constant(FC_BIN + 0)]]; constant short FC_bin_f [[function_constant(FC_BIN + 1)]]; constant bool FC_bin_rb [[function_constant(FC_BIN + 2)]]; +constant bool FC_bin_cb [[function_constant(FC_BIN + 3)]]; template kernel void kernel_bin_fuse_impl( @@ -1124,11 +1125,12 @@ kernel void kernel_bin_fuse_impl( #define FC_OP FC_bin_op #define FC_F FC_bin_f #define FC_RB FC_bin_rb +#define FC_CB FC_bin_cb if (FC_RB) { // row broadcast - const uint i0 = tgpig.x; - const uint i1 = i0%args.ne10; + const uint i0 = tgpig.y*args.ne00 + tgpig.x; + const uint i1 = FC_CB ? tgpig.x%args.ne10 : tgpig.x; device const T0 * src0_row = (device const T0 *) (src0); device T * dst_row = (device T *) (dst); @@ -1200,7 +1202,7 @@ kernel void kernel_bin_fuse_impl( device const T1 * src1_ptr = (device const T1 *) (src1 + args.o1[0] + i13*args.nb13 + i12*args.nb12 + i11*args.nb11); for (int i0 = tpitg.x; i0 < args.ne0; i0 += ntg.x) { - const int i10 = i0%args.ne10; + const int i10 = FC_CB ? i0%args.ne10 : i0; if (FC_OP == 0) { dst_ptr[i0] = src0_ptr[i0] + src1_ptr[i10]; @@ -1225,7 +1227,7 @@ kernel void kernel_bin_fuse_impl( } for (int i0 = tpitg.x; i0 < args.ne0; i0 += ntg.x) { - const int i10 = i0%args.ne10; + const int i10 = FC_CB ? i0%args.ne10 : i0; T res = src0_ptr[i0]; @@ -1261,6 +1263,7 @@ kernel void kernel_bin_fuse_impl( #undef FC_OP #undef FC_F #undef FC_RB +#undef FC_CB } typedef decltype(kernel_bin_fuse_impl) kernel_bin_fuse_t; diff --git a/ggml/src/ggml-opencl/CMakeLists.txt b/ggml/src/ggml-opencl/CMakeLists.txt index 70802c9c00..1f8250934b 100644 --- a/ggml/src/ggml-opencl/CMakeLists.txt +++ b/ggml/src/ggml-opencl/CMakeLists.txt @@ -132,6 +132,7 @@ set(GGML_OPENCL_KERNELS ssm_conv sub sum_rows + cumsum transpose concat tsembd diff --git a/ggml/src/ggml-opencl/ggml-opencl.cpp b/ggml/src/ggml-opencl/ggml-opencl.cpp index 67e4b9277f..e1dca6b4b4 100644 --- a/ggml/src/ggml-opencl/ggml-opencl.cpp +++ b/ggml/src/ggml-opencl/ggml-opencl.cpp @@ -547,6 +547,7 @@ struct ggml_backend_opencl_context { cl_kernel kernel_im2col_f32, kernel_im2col_f16; cl_kernel kernel_argsort_f32_i32; cl_kernel kernel_sum_rows_f32, kernel_sum_rows_f32_4; + cl_kernel kernel_cumsum_blk, kernel_cumsum_add; cl_kernel kernel_repeat_f32; cl_kernel kernel_pad; cl_kernel kernel_tanh_f32, kernel_tanh_f32_4, kernel_tanh_f32_nc; @@ -1927,6 +1928,24 @@ static void load_cl_kernels(ggml_backend_opencl_context *backend_ctx, ggml_cl_ve GGML_LOG_CONT("."); } + // cumsum + { +#ifdef GGML_OPENCL_EMBED_KERNELS + const std::string kernel_src { + #include "cumsum.cl.h" + }; +#else + const std::string kernel_src = read_file("cumsum.cl"); +#endif + cl_program prog; + prog = build_program_from_source(backend_ctx->context, backend_ctx->device, kernel_src.c_str(), compile_opts); + + CL_CHECK((backend_ctx->kernel_cumsum_blk = clCreateKernel(prog, "kernel_cumsum_blk", &err), err)); + CL_CHECK((backend_ctx->kernel_cumsum_add = clCreateKernel(prog, "kernel_cumsum_add", &err), err)); + GGML_LOG_CONT("."); + CL_CHECK(clReleaseProgram(prog)); + } + // sigmoid { #ifdef GGML_OPENCL_EMBED_KERNELS @@ -3803,6 +3822,8 @@ static bool ggml_opencl_supports_op(ggml_backend_dev_t dev, const struct ggml_te return cols <= max_workgroup_size && op->src[0]->type == GGML_TYPE_F32; } case GGML_OP_SUM_ROWS: + case GGML_OP_CUMSUM: + return op->src[0]->type == GGML_TYPE_F32 && ggml_is_contiguous(op->src[0]); case GGML_OP_MEAN: return op->src[0]->type == GGML_TYPE_F32; case GGML_OP_FLASH_ATTN_EXT: @@ -5775,19 +5796,12 @@ static void ggml_cl_get_rows(ggml_backend_t backend, const ggml_tensor * src0, c GGML_ASSERT(dst); GGML_ASSERT(dst->extra); - const int ne00 = src0->ne[0]; - const cl_ulong nb01 = src0->nb[1]; - const cl_ulong nb02 = src0->nb[2]; - const cl_ulong nb03 = src0->nb[3]; - const int ne10 = src1->ne[0]; - const cl_ulong nb10 = src1->nb[0]; - const int ne11 = src1->ne[1]; - const int ne12 = src1->ne[2]; - const cl_ulong nb11 = src1->nb[1]; - const cl_ulong nb12 = src1->nb[2]; - const cl_ulong nb1 = dst->nb[1]; - const cl_ulong nb2 = dst->nb[2]; - const cl_ulong nb3 = dst->nb[3]; + GGML_TENSOR_LOCALS(int, ne0, src0, ne); + GGML_TENSOR_LOCALS(cl_ulong, nb0, src0, nb); + GGML_TENSOR_LOCALS(int, ne1, src1, ne); + GGML_TENSOR_LOCALS(cl_ulong, nb1, src1, nb); + GGML_TENSOR_LOCALS(int, ne, dst, ne); + GGML_TENSOR_LOCALS(cl_ulong, nb, dst, nb); ggml_backend_opencl_context *backend_ctx = (ggml_backend_opencl_context *)backend->context; @@ -5833,8 +5847,14 @@ static void ggml_cl_get_rows(ggml_backend_t backend, const ggml_tensor * src0, c CL_CHECK(clSetKernelArg(kernel, 15, sizeof(cl_ulong), &nb2)); CL_CHECK(clSetKernelArg(kernel, 16, sizeof(cl_ulong), &nb3)); - size_t global_work_size[] = {(size_t)ne10*64, (size_t)ne11, (size_t)ne12}; - size_t local_work_size[] = {64, 1, 1}; + int max_workgroup_size = backend_ctx->get_kernel_workgroup_size(kernel); + int nth = 1; + while (nth < ne00 && 2*nth <= max_workgroup_size) { + nth *= 2; + } + + size_t global_work_size[] = {(size_t)ne10*nth, (size_t)ne11, (size_t)ne12}; + size_t local_work_size[] = {(size_t)nth, 1, 1}; backend_ctx->enqueue_ndrange_kernel(kernel, 3, global_work_size, local_work_size, dst); } @@ -11949,6 +11969,118 @@ static void ggml_cl_sum_rows(ggml_backend_t backend, const ggml_tensor * src0, c backend_ctx->enqueue_ndrange_kernel(kernel, 3, global_work_size, local_work_size, dst); } +static void ggml_cl_cumsum(ggml_backend_t backend, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst) { + GGML_ASSERT(src0); + GGML_ASSERT(src0->extra); + GGML_ASSERT(dst); + GGML_ASSERT(dst->extra); + GGML_UNUSED(src1); + + GGML_ASSERT(src0->nb[0] == ggml_type_size(src0->type)); + GGML_ASSERT(ggml_is_contiguous(src0)); + + ggml_backend_opencl_context *backend_ctx = (ggml_backend_opencl_context *)backend->context; + + ggml_tensor_extra_cl * extra0 = (ggml_tensor_extra_cl *)src0->extra; + ggml_tensor_extra_cl * extrad = (ggml_tensor_extra_cl *)dst->extra; + + cl_ulong offset0 = extra0->offset + src0->view_offs; + cl_ulong offsetd = extrad->offset + dst->view_offs; + + GGML_TENSOR_LOCALS(int, ne0, src0, ne); + GGML_TENSOR_LOCALS(cl_ulong, nb0, src0, nb); + + cl_kernel kernel = backend_ctx->kernel_cumsum_blk; + + int max_workgroup_size = backend_ctx->get_kernel_workgroup_size(kernel); + int nth = 1; + while (nth < ne00 && 2*nth <= max_workgroup_size) { + nth *= 2; + } + + GGML_ASSERT(ne00 <= nth*nth); + + const int net0 = CEIL_DIV(ne00, nth); + const int net1 = ne01; + const int net2 = ne02; + const int net3 = ne03; + + const cl_ulong nbt0 = sizeof(float); + const cl_ulong nbt1 = net0*nbt0; + const cl_ulong nbt2 = net1*nbt1; + const cl_ulong nbt3 = net2*nbt2; + + static ggml_cl_buffer tmp_buffer; + tmp_buffer.allocate(backend_ctx->context, net0*ne01*ne02*ne03*sizeof(float)); + + 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), &tmp_buffer.buffer)); + CL_CHECK(clSetKernelArg(kernel, 3, sizeof(cl_mem), &extrad->data_device)); + CL_CHECK(clSetKernelArg(kernel, 4, sizeof(cl_ulong), &offsetd)); + CL_CHECK(clSetKernelArg(kernel, 5, sizeof(int), &ne00)); + CL_CHECK(clSetKernelArg(kernel, 6, sizeof(int), &ne01)); + CL_CHECK(clSetKernelArg(kernel, 7, sizeof(int), &ne02)); + CL_CHECK(clSetKernelArg(kernel, 8, sizeof(int), &ne03)); + CL_CHECK(clSetKernelArg(kernel, 9, sizeof(cl_ulong), &nb00)); + 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), &net0)); + CL_CHECK(clSetKernelArg(kernel, 14, sizeof(int), &net1)); + CL_CHECK(clSetKernelArg(kernel, 15, sizeof(int), &net2)); + + size_t global_work_size[] = { (size_t)(nth*net0*ne01), (size_t)ne02, (size_t)ne03}; + size_t local_work_size[] = { (size_t)nth, 1, 1}; + + backend_ctx->enqueue_ndrange_kernel(kernel, 3, global_work_size, local_work_size, dst); + + if(ne00 > nth) { + // if a single workgroup cannot handle an entire row, each workgroup + // computes a partial sum and stores to dst, tmp_buffer contains the sum + // of the each workgroup; cumsum this buffer and add to the partial sums in dst + cl_ulong offsett = 0; + kernel = backend_ctx->kernel_cumsum_blk; + CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), &tmp_buffer.buffer)); + CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_ulong), &offsett)); + CL_CHECK(clSetKernelArg(kernel, 2, sizeof(cl_mem), &tmp_buffer.buffer)); + CL_CHECK(clSetKernelArg(kernel, 3, sizeof(cl_mem), &tmp_buffer.buffer)); + CL_CHECK(clSetKernelArg(kernel, 4, sizeof(cl_ulong), &offsett)); + CL_CHECK(clSetKernelArg(kernel, 5, sizeof(int), &net0)); + CL_CHECK(clSetKernelArg(kernel, 6, sizeof(int), &ne01)); + CL_CHECK(clSetKernelArg(kernel, 7, sizeof(int), &ne02)); + CL_CHECK(clSetKernelArg(kernel, 8, sizeof(int), &ne03)); + CL_CHECK(clSetKernelArg(kernel, 9, sizeof(cl_ulong), &nbt0)); + CL_CHECK(clSetKernelArg(kernel, 10, sizeof(cl_ulong), &nbt1)); + CL_CHECK(clSetKernelArg(kernel, 11, sizeof(cl_ulong), &nbt2)); + CL_CHECK(clSetKernelArg(kernel, 12, sizeof(cl_ulong), &nbt3)); + CL_CHECK(clSetKernelArg(kernel, 13, sizeof(int), &net0)); + CL_CHECK(clSetKernelArg(kernel, 14, sizeof(int), &net1)); + CL_CHECK(clSetKernelArg(kernel, 15, sizeof(int), &net2)); + + size_t global_work_size_1[] = { (size_t)net1*nth, (size_t)net2, (size_t)net3}; + size_t local_work_size_1[] = { (size_t)nth, 1, 1}; + backend_ctx->enqueue_ndrange_kernel(kernel, 3, global_work_size_1, local_work_size_1, dst); + + kernel = backend_ctx->kernel_cumsum_add; + CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), &tmp_buffer.buffer)); + CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_mem), &extrad->data_device)); + CL_CHECK(clSetKernelArg(kernel, 2, sizeof(cl_ulong), &offsetd)); + CL_CHECK(clSetKernelArg(kernel, 3, sizeof(int), &ne00)); + CL_CHECK(clSetKernelArg(kernel, 4, sizeof(int), &ne01)); + CL_CHECK(clSetKernelArg(kernel, 5, sizeof(int), &ne02)); + CL_CHECK(clSetKernelArg(kernel, 6, sizeof(int), &ne03)); + CL_CHECK(clSetKernelArg(kernel, 7, sizeof(int), &nbt0)); + CL_CHECK(clSetKernelArg(kernel, 8, sizeof(int), &nbt1)); + CL_CHECK(clSetKernelArg(kernel, 9, sizeof(int), &nbt2)); + CL_CHECK(clSetKernelArg(kernel, 10, sizeof(int), &nbt3)); + + size_t global_work_size_2[] = { (size_t)(nth*net0*ne01), (size_t)ne02, (size_t)ne03}; + size_t local_work_size_2[] = { (size_t)nth, 1, 1}; + backend_ctx->enqueue_ndrange_kernel(kernel, 3, global_work_size_2, local_work_size_2, dst); + } +} + static void ggml_cl_glu(ggml_backend_t backend, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst) { GGML_ASSERT(src0); GGML_ASSERT(src0->extra); @@ -12391,6 +12523,12 @@ bool ggml_cl_compute_forward(ggml_backend_t backend, struct ggml_tensor * tensor } func = ggml_cl_sum_rows; break; + case GGML_OP_CUMSUM: + if (!any_on_device) { + return false; + } + func = ggml_cl_cumsum; + break; case GGML_OP_FLASH_ATTN_EXT: if (!any_on_device) { return false; diff --git a/ggml/src/ggml-opencl/kernels/cumsum.cl b/ggml/src/ggml-opencl/kernels/cumsum.cl new file mode 100644 index 0000000000..edfb74b705 --- /dev/null +++ b/ggml/src/ggml-opencl/kernels/cumsum.cl @@ -0,0 +1,139 @@ +#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 + +// max workgroup size is usually 1024, this covers various subgroups sizes +#define MAX_SUBGROUPS 128 + +#ifdef INTEL_GPU +REQD_SUBGROUP_SIZE_32 +#elif defined (ADRENO_GPU) +REQD_SUBGROUP_SIZE_64 +#endif +kernel void kernel_cumsum_blk( + global char * src0, + ulong offset0, + global char * tmp, + global char * dst, + ulong offsetd, + int ne00, + int ne01, + int ne02, + int ne03, + ulong nb00, + ulong nb01, + ulong nb02, + ulong nb03, + uint net0, + uint net1, + uint net2 +) { + src0 = src0 + offset0; + dst = dst + offsetd; + + const int i3 = get_group_id(2); + const int i2 = get_group_id(1); + const int i1 = get_group_id(0); + + const int nth = get_local_size(0); + const int tid = get_local_id(0); + + const uint sg_size = get_sub_group_size(); + const uint sg_id = get_sub_group_id(); + const uint sg_lid = get_sub_group_local_id(); + + const int ib = i1 / ne01; + const int i00 = ib * nth; + const int i01 = i1 % ne01; + const int i02 = i2; + const int i03 = i3; + + global const float * src0_row = (global const float *)(src0 + i03*nb03 + i02*nb02 + i01*nb01); + global float * tmp_row = (global float *)tmp + net0 * i01 + net0 * net1 * i02 + net0 * net1 * net2 * i03; + global float * dst_row = (global float *)dst + i03*ne02*ne01*ne00 + i02*ne01*ne00 + i01*ne00; + + __local float partial[MAX_SUBGROUPS]; + + float v = 0.0f; + if (i00 + tid < ne00) { + v = src0_row[i00 + tid]; + } + + float s = sub_group_scan_inclusive_add(v); + if (sg_lid == sg_size - 1) { + partial[sg_id] = s; + } + barrier(CLK_LOCAL_MEM_FENCE); + + // NB: subgroup size should be larger than number of subgroups + // assuming max workgroup size of 1024, subgroup size should be >= 32 + if (sg_id == 0) { + float x = 0.0f; + if (sg_lid < get_num_sub_groups()) { + x = partial[sg_lid]; + } + float ex = sub_group_scan_exclusive_add(x); + if (sg_lid < get_num_sub_groups()) { + partial[sg_lid] = ex; + } + } + barrier(CLK_LOCAL_MEM_FENCE); + + s += partial[sg_id]; + + if (i00 + tid < ne00) { + dst_row[i00 + tid] = s; + } + if (ne00 > nth && tid == nth - 1) { + tmp_row[ib] = s; + } +} + +kernel void kernel_cumsum_add( + global char * tmp, + global char * dst, + ulong offsetd, + int ne00, + int ne01, + int ne02, + int ne03, + uint nbt0, + uint nbt1, + uint nbt2, + uint nbt3 +) { + dst = dst + offsetd; + + const int i3 = get_group_id(2); + const int i2 = get_group_id(1); + const int i1 = get_group_id(0); + + const int nth = get_local_size(0); + const int tid = get_local_id(0); + + const int ib = i1 / ne01; + if (ib == 0) { + return; + } + const int i00 = ib * nth; + const int i01 = i1 % ne01; + const int i02 = i2; + const int i03 = i3; + + global float * tmp_row = (global float *)(tmp + nbt1 * i01 + nbt2 * i02 + nbt3 * i03); + global float * dst_row = (global float *)dst + i03*ne02*ne01*ne00 + i02*ne01*ne00 + i01*ne00; + + if (i00 + tid < ne00) { + dst_row[i00 + tid] += tmp_row[ib - 1]; + } +} diff --git a/ggml/src/ggml-vulkan/ggml-vulkan.cpp b/ggml/src/ggml-vulkan/ggml-vulkan.cpp index 61d112c50a..3c81805b84 100644 --- a/ggml/src/ggml-vulkan/ggml-vulkan.cpp +++ b/ggml/src/ggml-vulkan/ggml-vulkan.cpp @@ -27,6 +27,7 @@ DispatchLoaderDynamic & ggml_vk_default_dispatcher(); #include #include #include +#include #include #include #include @@ -188,6 +189,11 @@ struct ggml_backend_vk_buffer_type_context { struct vk_queue; +struct vk_command_buffer { + vk::CommandBuffer buf; + bool in_use = false; +}; + // Stores command pool/buffers. There's an instance of this // for each (context,queue) pair and for each (device,queue) pair. struct vk_command_pool { @@ -195,10 +201,16 @@ struct vk_command_pool { void destroy(vk::Device& device); vk::CommandPool pool; - uint32_t cmd_buffer_idx; - std::vector cmd_buffers; + // Using deque so the pointers to command buffers + // remain valid even if we add more + std::deque cmd_buffers; vk_queue *q; + + size_t buffers_in_use() const { + return std::count_if(cmd_buffers.begin(), cmd_buffers.end(), + [](const auto& cb) { return cb.in_use; }); + } }; // Prevent simultaneous submissions to the same queue. @@ -813,6 +825,8 @@ struct vk_device_struct { vk_pipeline pipeline_pool2d_f32; vk_pipeline pipeline_rwkv_wkv6_f32; vk_pipeline pipeline_rwkv_wkv7_f32; + // [size_idx][kda] where size_idx: 0=d32, 1=d64, 2=d128 + vk_pipeline pipeline_gated_delta_net[3][2]; vk_pipeline pipeline_ssm_scan_f32_d128; vk_pipeline pipeline_ssm_scan_f32_d256; vk_pipeline pipeline_ssm_conv_f32; @@ -878,10 +892,12 @@ struct vk_device_struct { }; void vk_command_pool::init(vk_device& device, vk_queue *q_) { - cmd_buffer_idx = 0; + cmd_buffers.clear(); q = q_; - vk::CommandPoolCreateInfo command_pool_create_info(vk::CommandPoolCreateFlags(VK_COMMAND_POOL_CREATE_TRANSIENT_BIT), q->queue_family_index); + vk::CommandPoolCreateInfo command_pool_create_info( + vk::CommandPoolCreateFlags(VK_COMMAND_POOL_CREATE_TRANSIENT_BIT | VK_COMMAND_POOL_CREATE_RESET_COMMAND_BUFFER_BIT), + q->queue_family_index); pool = device->device.createCommandPool(command_pool_create_info); } @@ -929,6 +945,7 @@ struct vk_subbuffer { struct vk_event { vk::Event event; vk::Fence fence; + vk_command_buffer* cmd_buffer = nullptr; }; struct vk_semaphore { @@ -937,7 +954,7 @@ struct vk_semaphore { }; struct vk_submission { - vk::CommandBuffer buffer; + vk_command_buffer* buffer = nullptr; std::vector wait_semaphores; std::vector signal_semaphores; }; @@ -1439,6 +1456,18 @@ struct vk_op_rwkv_wkv7_push_constants { uint32_t C; uint32_t H; }; +struct vk_op_gated_delta_net_push_constants { + uint32_t H; + uint32_t n_tokens; + uint32_t n_seqs; + uint32_t s_off; + uint32_t sq1, sq2, sq3; + uint32_t sv1, sv2, sv3; + uint32_t sb1, sb2, sb3; + uint32_t neq1, rq3; + float scale; +}; + struct vk_op_ssm_scan_push_constants { uint32_t nb02, nb03, nb12, nb13; uint32_t nb21, nb22, nb31; @@ -2283,25 +2312,15 @@ static void ggml_pipeline_allocate_descriptor_sets(ggml_backend_vk_context * ctx } } -static vk::CommandBuffer ggml_vk_create_cmd_buffer(vk_device& device, vk_command_pool& p) { +static vk_command_buffer* ggml_vk_create_cmd_buffer(vk_device& device, vk_command_pool& p) { VK_LOG_DEBUG("ggml_vk_create_cmd_buffer()"); - - if (p.cmd_buffers.size() > p.cmd_buffer_idx) { - // Reuse command buffer - return p.cmd_buffers[p.cmd_buffer_idx++]; - } - vk::CommandBufferAllocateInfo command_buffer_alloc_info( p.pool, vk::CommandBufferLevel::ePrimary, 1); const std::vector cmd_buffers = device->device.allocateCommandBuffers(command_buffer_alloc_info); - auto buf = cmd_buffers.front(); - - p.cmd_buffers.push_back(buf); - p.cmd_buffer_idx++; - - return buf; + p.cmd_buffers.push_back({ cmd_buffers.front(), true }); + return &p.cmd_buffers[p.cmd_buffers.size()-1]; } static void ggml_vk_submit(vk_context& ctx, vk::Fence fence) { @@ -2368,7 +2387,7 @@ static void ggml_vk_submit(vk_context& ctx, vk::Fence fence) { tl_wait_semaphores[idx].data(), stage_flags[idx].data(), 1, - &submission.buffer, + &submission.buffer->buf, (uint32_t) submission.signal_semaphores.size(), tl_signal_semaphores[idx].data(), }; @@ -2492,7 +2511,11 @@ static void ggml_vk_command_pool_cleanup(vk_device& device, vk_command_pool& p) // Requires command buffers to be done device->device.resetCommandPool(p.pool); - p.cmd_buffer_idx = 0; + // Don't clear the command buffers and mark them as not in use. + // This allows us to reuse them + for (auto& cmd_buffer : p.cmd_buffers) { + cmd_buffer.in_use = false; + } } static void ggml_vk_queue_command_pools_cleanup(vk_device& device) { @@ -2501,10 +2524,10 @@ static void ggml_vk_queue_command_pools_cleanup(vk_device& device) { // Arbitrary frequency to cleanup/reuse command buffers static constexpr uint32_t cleanup_frequency = 10; - if (device->compute_queue.cmd_pool.cmd_buffer_idx >= cleanup_frequency) { + if (device->compute_queue.cmd_pool.buffers_in_use() >= cleanup_frequency) { ggml_vk_command_pool_cleanup(device, device->compute_queue.cmd_pool); } - if (device->transfer_queue.cmd_pool.cmd_buffer_idx >= cleanup_frequency) { + if (device->transfer_queue.cmd_pool.buffers_in_use() >= cleanup_frequency) { ggml_vk_command_pool_cleanup(device, device->transfer_queue.cmd_pool); } } @@ -2752,7 +2775,7 @@ static void ggml_vk_sync_buffers(ggml_backend_vk_context* ctx, vk_context& subct ctx->prealloc_x_need_sync = ctx->prealloc_y_need_sync = ctx->prealloc_split_k_need_sync = false; } - subctx->s->buffer.pipelineBarrier( + subctx->s->buffer->buf.pipelineBarrier( subctx->p->q->stage_flags, subctx->p->q->stage_flags, {}, @@ -2768,7 +2791,7 @@ static void ggml_vk_sync_buffers(ggml_backend_vk_context* ctx, vk_context& subct static void ggml_vk_set_event(vk_context& ctx, vk::Event& event) { VK_LOG_DEBUG("ggml_vk_set_event()"); - ctx->s->buffer.setEvent( + ctx->s->buffer->buf.setEvent( event, ctx->p->q->stage_flags ); @@ -2780,7 +2803,7 @@ static void ggml_vk_wait_events(vk_context& ctx, std::vector&& events return; } - ctx->s->buffer.waitEvents( + ctx->s->buffer->buf.waitEvents( events, ctx->p->q->stage_flags, ctx->p->q->stage_flags, @@ -4559,6 +4582,23 @@ static void ggml_vk_load_shaders(vk_device& device) { ggml_vk_create_pipeline(device, device->pipeline_rwkv_wkv7_f32, "rwkv_wkv7_f32", rwkv_wkv7_f32_len, rwkv_wkv7_f32_data, "main", 8, sizeof(vk_op_rwkv_wkv7_push_constants), {1, 1, 1}, {device->subgroup_size}, 1); + { + const uint32_t gdn_sizes[] = {32, 64, 128}; + const char * gdn_names[][2] = { + {"gated_delta_net_f32_d32", "gated_delta_net_f32_d32_kda"}, + {"gated_delta_net_f32_d64", "gated_delta_net_f32_d64_kda"}, + {"gated_delta_net_f32_d128", "gated_delta_net_f32_d128_kda"}, + }; + for (uint32_t si = 0; si < 3; si++) { + for (uint32_t kda = 0; kda < 2; kda++) { + ggml_vk_create_pipeline(device, device->pipeline_gated_delta_net[si][kda], + gdn_names[si][kda], gated_delta_net_f32_len, gated_delta_net_f32_data, + "main", 7, sizeof(vk_op_gated_delta_net_push_constants), + {1, 1, 1}, {gdn_sizes[si], kda}, 1); + } + } + } + if (device->subgroup_arithmetic && device->subgroup_require_full_support) { ggml_vk_create_pipeline(device, device->pipeline_ssm_scan_f32_d128, "ssm_scan_128_f32", ssm_scan_subgroup_f32_len, ssm_scan_subgroup_f32_data, "main", 8, sizeof(vk_op_ssm_scan_push_constants), {1, 1, 1}, {128, device->subgroup_size}, 1, true, true); ggml_vk_create_pipeline(device, device->pipeline_ssm_scan_f32_d256, "ssm_scan_256_f32", ssm_scan_subgroup_f32_len, ssm_scan_subgroup_f32_data, "main", 8, sizeof(vk_op_ssm_scan_push_constants), {1, 1, 1}, {256, device->subgroup_size}, 1, true, true); @@ -4567,7 +4607,7 @@ static void ggml_vk_load_shaders(vk_device& device) { ggml_vk_create_pipeline(device, device->pipeline_ssm_scan_f32_d256, "ssm_scan_256_f32", ssm_scan_f32_len, ssm_scan_f32_data, "main", 8, sizeof(vk_op_ssm_scan_push_constants), {1, 1, 1}, {256, device->subgroup_size, 16}, 1, true, true); } - ggml_vk_create_pipeline(device, device->pipeline_ssm_conv_f32, "ssm_conv_f32", ssm_conv_f32_len, ssm_conv_f32_data, "main", 3, sizeof(vk_op_ssm_conv_push_constants), {32, 1, 1}, {32}, 1); + ggml_vk_create_pipeline(device, device->pipeline_ssm_conv_f32, "ssm_conv_f32", ssm_conv_f32_len, ssm_conv_f32_data, "main", 3, sizeof(vk_op_ssm_conv_push_constants), {32, 16, 1}, {32, 16}, 1); ggml_vk_create_pipeline(device, device->pipeline_opt_step_adamw_f32, "opt_step_adamw_f32", opt_step_adamw_f32_len, opt_step_adamw_f32_data, "main", 5, sizeof(vk_op_push_constants), {512, 1, 1}, {}, 1); @@ -6348,13 +6388,24 @@ static vk_subbuffer ggml_vk_tensor_subbuffer( return vk_subbuffer{buffer, offset, size}; } +// Get a command buffer from pool. Create a new one if no reusable buffer is available +static vk_command_buffer* ggml_vk_get_or_create_cmd_buffer(vk_device& device, vk_command_pool& pool) { + for (auto& cmd_buffer : pool.cmd_buffers) { + if (!cmd_buffer.in_use) { + cmd_buffer.in_use = true; + return &cmd_buffer; + } + } + return ggml_vk_create_cmd_buffer(device, pool); +} + static vk_submission ggml_vk_begin_submission(vk_device& device, vk_command_pool& p, bool one_time = true) { vk_submission s; - s.buffer = ggml_vk_create_cmd_buffer(device, p); + s.buffer = ggml_vk_get_or_create_cmd_buffer(device, p); if (one_time) { - s.buffer.begin({ vk::CommandBufferUsageFlagBits::eOneTimeSubmit }); + s.buffer->buf.begin({ vk::CommandBufferUsageFlagBits::eOneTimeSubmit }); } else { - s.buffer.begin({ vk::CommandBufferUsageFlags{} }); + s.buffer->buf.begin({ vk::CommandBufferUsageFlags{} }); } return s; @@ -6407,18 +6458,18 @@ static void ggml_vk_dispatch_pipeline(ggml_backend_vk_context* ctx, vk_context& vk::WriteDescriptorSet write_descriptor_set{ descriptor_set, 0, 0, pipeline->parameter_count, vk::DescriptorType::eStorageBuffer, nullptr, descriptor_buffer_infos.begin() }; ctx->device->device.updateDescriptorSets({ write_descriptor_set }, {}); - subctx->s->buffer.pushConstants(pipeline->layout, vk::ShaderStageFlagBits::eCompute, 0, push_constant_size(push_constants), push_constant_data(push_constants)); - subctx->s->buffer.bindPipeline(vk::PipelineBindPoint::eCompute, pipeline->pipeline); - subctx->s->buffer.bindDescriptorSets(vk::PipelineBindPoint::eCompute, + subctx->s->buffer->buf.pushConstants(pipeline->layout, vk::ShaderStageFlagBits::eCompute, 0, push_constant_size(push_constants), push_constant_data(push_constants)); + subctx->s->buffer->buf.bindPipeline(vk::PipelineBindPoint::eCompute, pipeline->pipeline); + subctx->s->buffer->buf.bindDescriptorSets(vk::PipelineBindPoint::eCompute, pipeline->layout, 0, { descriptor_set }, {}); - subctx->s->buffer.dispatch(wg0, wg1, wg2); + subctx->s->buffer->buf.dispatch(wg0, wg1, wg2); } static void ggml_vk_end_submission(vk_submission& s, std::vector wait_semaphores, std::vector signal_semaphores) { - s.buffer.end(); + s.buffer->buf.end(); s.wait_semaphores = std::move(wait_semaphores); s.signal_semaphores = std::move(signal_semaphores); @@ -6430,7 +6481,7 @@ static void ggml_vk_ctx_end(vk_context& ctx) { return; } - ctx->s->buffer.end(); + ctx->s->buffer->buf.end(); ctx->s = nullptr; } @@ -6584,7 +6635,7 @@ static void ggml_vk_buffer_write_nc_async(ggml_backend_vk_context * ctx, vk_cont } ggml_vk_sync_buffers(ctx, subctx); - subctx->s->buffer.copyBuffer(buf->buffer, dst->buffer, slices); + subctx->s->buffer->buf.copyBuffer(buf->buffer, dst->buffer, slices); return; } @@ -6599,7 +6650,7 @@ static void ggml_vk_buffer_write_nc_async(ggml_backend_vk_context * ctx, vk_cont VkBufferCopy buf_copy{ 0, offset, copy_size }; ggml_vk_sync_buffers(ctx, subctx); - vkCmdCopyBuffer(subctx->s->buffer, (VkBuffer)staging->buffer, (VkBuffer)dst->buffer, 1, &buf_copy); + vkCmdCopyBuffer(subctx->s->buffer->buf, (VkBuffer)staging->buffer, (VkBuffer)dst->buffer, 1, &buf_copy); for (uint64_t i3 = 0; i3 < ne3; i3++) { for (uint64_t i2 = 0; i2 < ne2; i2++) { @@ -6648,7 +6699,7 @@ static bool ggml_vk_buffer_write_2d_async(vk_context subctx, vk_buffer& dst, siz } ggml_vk_sync_buffers(nullptr, subctx); - subctx->s->buffer.copyBuffer(buf->buffer, dst->buffer, slices); + subctx->s->buffer->buf.copyBuffer(buf->buffer, dst->buffer, slices); return true; } VK_LOG_DEBUG("STAGING"); @@ -6670,7 +6721,7 @@ static bool ggml_vk_buffer_write_2d_async(vk_context subctx, vk_buffer& dst, siz copy_size}; ggml_vk_sync_buffers(nullptr, subctx); - vkCmdCopyBuffer(subctx->s->buffer, (VkBuffer)staging_buffer->buffer, (VkBuffer)dst->buffer, 1, &buf_copy); + vkCmdCopyBuffer(subctx->s->buffer->buf, (VkBuffer)staging_buffer->buffer, (VkBuffer)dst->buffer, 1, &buf_copy); if (width == spitch) { deferred_memcpy((uint8_t *)staging_buffer->ptr, src, width * height, &subctx->in_memcpys); @@ -6756,7 +6807,7 @@ static bool ggml_vk_buffer_read_2d_async(vk_context subctx, vk_buffer& src, size if (buf != nullptr) { // Memory is pinned, use as staging buffer ggml_vk_sync_buffers(nullptr, subctx); - subctx->s->buffer.copyBuffer(src->buffer, buf->buffer, slices); + subctx->s->buffer->buf.copyBuffer(src->buffer, buf->buffer, slices); return true; } @@ -6774,7 +6825,7 @@ static bool ggml_vk_buffer_read_2d_async(vk_context subctx, vk_buffer& src, size vk_buffer& staging_buffer = src->device->sync_staging; ggml_vk_sync_buffers(nullptr, subctx); - subctx->s->buffer.copyBuffer(src->buffer, staging_buffer->buffer, slices); + subctx->s->buffer->buf.copyBuffer(src->buffer, staging_buffer->buffer, slices); deferred_memcpy(dst, staging_buffer->ptr, copy_size, &subctx->out_memcpys); return true; @@ -6821,7 +6872,7 @@ static void ggml_vk_buffer_copy_async(vk_context& ctx, vk_buffer& dst, size_t ds VkBufferCopy bc{ src_offset, dst_offset, size }; - vkCmdCopyBuffer(ctx->s->buffer, (VkBuffer)src->buffer, (VkBuffer)dst->buffer, 1, &bc); + vkCmdCopyBuffer(ctx->s->buffer->buf, (VkBuffer)src->buffer, (VkBuffer)dst->buffer, 1, &bc); } static void ggml_vk_buffer_copy(vk_buffer& dst, size_t dst_offset, vk_buffer& src, size_t src_offset, size_t size) { @@ -6859,7 +6910,7 @@ static void ggml_vk_buffer_memset_async(vk_context& ctx, vk_buffer& dst, size_t } // Fall back to GPU fillBuffer for non-UMA or non-host-visible buffers - ctx->s->buffer.fillBuffer(dst->buffer, offset, size, c); + ctx->s->buffer->buf.fillBuffer(dst->buffer, offset, size, c); } static void ggml_vk_buffer_memset(vk_buffer& dst, size_t offset, uint32_t c, size_t size) { @@ -6874,7 +6925,7 @@ static void ggml_vk_buffer_memset(vk_buffer& dst, size_t offset, uint32_t c, siz std::lock_guard guard(dst->device->mutex); vk_context subctx = ggml_vk_create_temporary_context(dst->device->transfer_queue.cmd_pool); ggml_vk_ctx_begin(dst->device, subctx); - subctx->s->buffer.fillBuffer(dst->buffer, offset, size, c); + subctx->s->buffer->buf.fillBuffer(dst->buffer, offset, size, c); ggml_vk_ctx_end(subctx); ggml_vk_submit(subctx, dst->device->fence); @@ -8820,7 +8871,7 @@ static void ggml_vk_flash_attn(ggml_backend_vk_context * ctx, vk_context& subctx } // Only use mask opt when the mask is fairly large. This hasn't been tuned extensively. - bool use_mask_opt = mask && nem1 >= 32 && nem0 * nem1 > 32768; + bool use_mask_opt = mask && nem1 >= 32 && nem0 * nem1 > 32768 && nem0 >= tuning_params.block_cols * 16; vk_fa_pipeline_state fa_pipeline_state = get_fa_pipeline_state(ctx->device, tuning_params, HSK, HSV, aligned, f32acc, mask != nullptr, use_mask_opt, logit_softcap != 0); @@ -9478,6 +9529,20 @@ static vk_pipeline ggml_vk_op_get_pipeline(ggml_backend_vk_context * ctx, const return ctx->device->pipeline_rwkv_wkv7_f32; } return nullptr; + case GGML_OP_GATED_DELTA_NET: + if (src0->type == GGML_TYPE_F32 && dst->type == GGML_TYPE_F32) { + const uint32_t S_v = dst->src[2]->ne[0]; + const uint32_t kda = (dst->src[3]->ne[0] == (int64_t)S_v) ? 1 : 0; + uint32_t si; + switch (S_v) { + case 32: si = 0; break; + case 64: si = 1; break; + case 128: si = 2; break; + default: return nullptr; + } + return ctx->device->pipeline_gated_delta_net[si][kda]; + } + return nullptr; case GGML_OP_SSM_SCAN: if (src0->type == GGML_TYPE_F32 && dst->type == GGML_TYPE_F32) { const uint32_t d_state = src0->ne[0]; @@ -10308,6 +10373,59 @@ static void ggml_vk_rwkv_wkv7(ggml_backend_vk_context * ctx, vk_context& subctx, ); } +static void ggml_vk_gated_delta_net(ggml_backend_vk_context * ctx, vk_context& subctx, ggml_tensor * dst) { + const ggml_tensor * src_q = dst->src[0]; + const ggml_tensor * src_v = dst->src[2]; + const ggml_tensor * src_beta = dst->src[4]; + + GGML_ASSERT(dst->buffer != nullptr); + + const uint32_t S_v = (uint32_t)src_v->ne[0]; + const uint32_t H = (uint32_t)src_v->ne[1]; + const uint32_t n_tokens = (uint32_t)src_v->ne[2]; + const uint32_t n_seqs = (uint32_t)src_v->ne[3]; + + const uint32_t s_off = S_v * H * n_tokens * n_seqs; + + vk_pipeline pipeline = ggml_vk_op_get_pipeline(ctx, dst->src[0], dst->src[1], dst->src[2], dst, dst->op); + GGML_ASSERT(pipeline != nullptr); + + ggml_pipeline_request_descriptor_sets(ctx, pipeline, 1); + + vk_subbuffer dst_buf = ggml_vk_tensor_subbuffer(ctx, dst); + vk_subbuffer src_buf[6] = {}; + for (int i = 0; i < 6; i++) { + src_buf[i] = ggml_vk_tensor_subbuffer(ctx, dst->src[i]); + } + + const uint32_t sq1 = (uint32_t)(src_q->nb[1] / sizeof(float)); + const uint32_t sq2 = (uint32_t)(src_q->nb[2] / sizeof(float)); + const uint32_t sq3 = (uint32_t)(src_q->nb[3] / sizeof(float)); + const uint32_t sv1 = (uint32_t)(src_v->nb[1] / sizeof(float)); + const uint32_t sv2 = (uint32_t)(src_v->nb[2] / sizeof(float)); + const uint32_t sv3 = (uint32_t)(src_v->nb[3] / sizeof(float)); + const uint32_t sb1 = (uint32_t)(src_beta->nb[1] / sizeof(float)); + const uint32_t sb2 = (uint32_t)(src_beta->nb[2] / sizeof(float)); + const uint32_t sb3 = (uint32_t)(src_beta->nb[3] / sizeof(float)); + + const uint32_t neq1 = (uint32_t)src_q->ne[1]; + const uint32_t rq3 = (uint32_t)(src_v->ne[3] / src_q->ne[3]); + + const float scale = 1.0f / sqrtf((float)S_v); + const vk_op_gated_delta_net_push_constants pc = { + H, n_tokens, n_seqs, s_off, + sq1, sq2, sq3, + sv1, sv2, sv3, + sb1, sb2, sb3, + neq1, rq3, + scale + }; + + ggml_vk_dispatch_pipeline(ctx, subctx, pipeline, + {src_buf[0], src_buf[1], src_buf[2], src_buf[3], src_buf[4], src_buf[5], dst_buf}, + pc, { H, n_seqs, 1u }); +} + static void ggml_vk_ssm_scan(ggml_backend_vk_context * ctx, vk_context& subctx, ggml_tensor * dst) { const ggml_tensor * src0 = dst->src[0]; const ggml_tensor * src1 = dst->src[1]; @@ -12682,7 +12800,7 @@ static bool ggml_vk_build_graph(ggml_backend_vk_context * ctx, ggml_cgraph * cgr if (vk_perf_logger_enabled && vk_perf_logger_concurrent) { ctx->query_node_idx[ctx->query_idx] = node_idx; - compute_ctx->s->buffer.writeTimestamp(vk::PipelineStageFlagBits::eAllCommands, ctx->query_pool, ctx->query_idx++); + compute_ctx->s->buffer->buf.writeTimestamp(vk::PipelineStageFlagBits::eAllCommands, ctx->query_pool, ctx->query_idx++); } } // Add all fused nodes to the unsynchronized lists. @@ -13024,6 +13142,11 @@ static bool ggml_vk_build_graph(ggml_backend_vk_context * ctx, ggml_cgraph * cgr break; + case GGML_OP_GATED_DELTA_NET: + ggml_vk_gated_delta_net(ctx, compute_ctx, node); + + break; + case GGML_OP_SSM_SCAN: ggml_vk_ssm_scan(ctx, compute_ctx, node); @@ -13521,7 +13644,7 @@ static void ggml_backend_vk_set_tensor_async(ggml_backend_t backend, ggml_tensor buffer_cpy.dstOffset = dst_offset; buffer_cpy.size = size; - cpy_ctx->s->buffer.copyBuffer(ctx->sync_staging->buffer, buf->buffer, { buffer_cpy }); + cpy_ctx->s->buffer->buf.copyBuffer(ctx->sync_staging->buffer, buf->buffer, { buffer_cpy }); deferred_memcpy(ctx->sync_staging->ptr, data, size, &cpy_ctx->in_memcpys); ggml_vk_synchronize(ctx); } @@ -13555,7 +13678,7 @@ static void ggml_backend_vk_get_tensor_async(ggml_backend_t backend, const ggml_ buffer_cpy.dstOffset = 0; buffer_cpy.size = size; - compute_ctx->s->buffer.copyBuffer(buf->buffer, ctx->sync_staging->buffer, { buffer_cpy }); + compute_ctx->s->buffer->buf.copyBuffer(buf->buffer, ctx->sync_staging->buffer, { buffer_cpy }); deferred_memcpy(data, ctx->sync_staging->ptr, size, &compute_ctx->out_memcpys); ggml_vk_synchronize(ctx); } @@ -13633,8 +13756,12 @@ static void ggml_vk_synchronize(ggml_backend_vk_context * ctx) { } vk_context compute_ctx; + vk_command_buffer* cmd_buf = nullptr; if (do_transfer) { compute_ctx = ctx->compute_ctx.lock(); + if (compute_ctx->s) { + cmd_buf = compute_ctx->s->buffer; + } ggml_vk_ctx_end(compute_ctx); @@ -13668,6 +13795,9 @@ static void ggml_vk_synchronize(ggml_backend_vk_context * ctx) { } ggml_vk_wait_for_fence(ctx); ctx->submit_pending = false; + if (cmd_buf) { + cmd_buf->in_use = false; + } } if (do_transfer) { @@ -14157,7 +14287,7 @@ static ggml_status ggml_backend_vk_graph_compute(ggml_backend_t backend, ggml_cg GGML_ASSERT(ctx->compute_ctx.expired()); compute_ctx = ggml_vk_get_compute_ctx(ctx); ctx->query_idx = 0; - compute_ctx->s->buffer.writeTimestamp(vk::PipelineStageFlagBits::eAllCommands, ctx->query_pool, ctx->query_idx++); + compute_ctx->s->buffer->buf.writeTimestamp(vk::PipelineStageFlagBits::eAllCommands, ctx->query_pool, ctx->query_idx++); } ctx->prealloc_y_last_pipeline_used = nullptr; @@ -14393,7 +14523,7 @@ static ggml_status ggml_backend_vk_graph_compute(ggml_backend_t backend, ggml_cg // track a single node/fusion for the current query ctx->query_nodes[ctx->query_idx] = cgraph->nodes[i]; ctx->query_fusion_names[ctx->query_idx] = fusion_string; - compute_ctx->s->buffer.writeTimestamp(vk::PipelineStageFlagBits::eAllCommands, ctx->query_pool, ctx->query_idx++); + compute_ctx->s->buffer->buf.writeTimestamp(vk::PipelineStageFlagBits::eAllCommands, ctx->query_pool, ctx->query_idx++); } else { // track a fusion string and number of fused ops for the current node_idx ctx->query_fusion_names[i] = fusion_string; @@ -14726,6 +14856,7 @@ static void ggml_backend_vk_event_record(ggml_backend_t backend, ggml_backend_ev ggml_vk_submit_transfer_ctx(ctx); vk_context compute_ctx = ggml_vk_get_compute_ctx(ctx); + auto* cmd_buf = compute_ctx->s->buffer; // retrieve pointer before it gets reset // the backend interface doesn't have an explicit reset, so reset it here // before we record the command to set it @@ -14738,6 +14869,7 @@ static void ggml_backend_vk_event_record(ggml_backend_t backend, ggml_backend_ev ggml_vk_submit(compute_ctx, {vkev->fence}); ctx->submit_pending = true; + vkev->cmd_buffer = cmd_buf; ctx->compute_ctx.reset(); } @@ -15426,6 +15558,19 @@ static bool ggml_backend_vk_device_supports_op(ggml_backend_dev_t dev, const ggm case GGML_OP_RWKV_WKV6: case GGML_OP_RWKV_WKV7: return true; // all inputs are contiguous, see ggml.c + case GGML_OP_GATED_DELTA_NET: + { + const uint32_t S_v = op->src[2]->ne[0]; + if (S_v != 32 && S_v != 64 && S_v != 128) { + return false; + } + for (int i = 0; i < 6; i++) { + if (op->src[i] == nullptr || op->src[i]->type != GGML_TYPE_F32) { + return false; + } + } + return op->type == GGML_TYPE_F32; + } case GGML_OP_SSM_SCAN: { for (int i = 0; i < 6; i++) { @@ -15557,6 +15702,10 @@ static void ggml_backend_vk_device_event_synchronize(ggml_backend_dev_t dev, ggm vk_event *vkev = (vk_event *)event->context; VK_CHECK(device->device.waitForFences({ vkev->fence }, true, UINT64_MAX), "event_synchronize"); + // Finished using current command buffer so we flag for reuse + if (vkev->cmd_buffer) { + vkev->cmd_buffer->in_use = false; + } } static vk_buffer ggml_vk_buffer_from_host_ptr(vk_device & device, void * ptr, size_t size) { @@ -16028,7 +16177,7 @@ static void ggml_vk_check_results_0(ggml_backend_vk_context * ctx, ggml_cgraph * tensor_clone = ggml_arange(ggml_ctx, start, stop, step); } else if (tensor->op == GGML_OP_FILL) { const float value = ggml_get_op_params_f32(tensor, 0); - tensor_clone = ggml_fill(ggml_ctx, tensor_clone, value); + tensor_clone = ggml_fill(ggml_ctx, src_clone[0], value); } else if (tensor->op == GGML_OP_SQR) { tensor_clone = ggml_sqr(ggml_ctx, src_clone[0]); } else if (tensor->op == GGML_OP_SQRT) { @@ -16299,6 +16448,9 @@ static void ggml_vk_check_results_0(ggml_backend_vk_context * ctx, ggml_cgraph * } else if (tensor->op == GGML_OP_RWKV_WKV7) { tensor_clone = ggml_rwkv_wkv7(ggml_ctx, src_clone[0], src_clone[1], src_clone[2], src_clone[3], src_clone[4], src_clone[5], src_clone[6]); + } else if (tensor->op == GGML_OP_GATED_DELTA_NET) { + tensor_clone = ggml_gated_delta_net(ggml_ctx, src_clone[0], src_clone[1], + src_clone[2], src_clone[3], src_clone[4], src_clone[5]); } else if (tensor->op == GGML_OP_OPT_STEP_ADAMW) { src_clone[0]->flags = tensor->src[0]->flags; tensor_clone = ggml_opt_step_adamw(ggml_ctx, src_clone[0], src_clone[1], diff --git a/ggml/src/ggml-vulkan/vulkan-shaders/flash_attn_mask_opt.comp b/ggml/src/ggml-vulkan/vulkan-shaders/flash_attn_mask_opt.comp index 8c92c1adcd..0e41770806 100644 --- a/ggml/src/ggml-vulkan/vulkan-shaders/flash_attn_mask_opt.comp +++ b/ggml/src/ggml-vulkan/vulkan-shaders/flash_attn_mask_opt.comp @@ -33,6 +33,61 @@ layout (push_constant) uniform parameter { shared float minsh[NUM_SUBGROUPS]; shared float maxsh[NUM_SUBGROUPS]; +float FLT_MAX_OVER_2 = uintBitsToFloat(0x7EFFFFFF); + +void loadvec4(inout uint result, const uint i0, const uint i1, const uint i2, const uint i3, const bool need_bounds_check) { + const uint tid = gl_LocalInvocationIndex; + + [[unroll]] for (uint block_x = 0; block_x < 16; ++block_x) { + float min_v = FLT_MAX_OVER_2; + float max_v = -FLT_MAX_OVER_2; + [[unroll]] for (uint i = 0; i < Br * Bc / 4; i += BLOCK_SIZE) { + uint j0 = (i + tid) % (Bc / 4); + uint j1 = (i + tid) / (Bc / 4); + + j0 *= 4; + j0 += (i0 * 16 + block_x) * Bc; + j1 += i1 * Br; + + if (!need_bounds_check || j0 + 3 < nem0) { + vec4 f = vec4(data_av4[(j0 + j1 * nbm1 + i2 * nbm2 + i3 * nbm3) / 4]); + [[unroll]] for (int c = 0; c < 4; ++c) { + min_v = min(min_v, f[c]); + max_v = max(max_v, f[c]); + } + } else { + [[unroll]] for (int c = 0; c < 4; ++c) { + if (j0 + c < nem0) { + float f = float(data_a[j0 + j1 * nbm1 + i2 * nbm2 + i3 * nbm3]); + min_v = min(min_v, f); + max_v = max(max_v, f); + } + } + } + } + min_v = subgroupMin(min_v); + max_v = subgroupMax(max_v); + if (gl_SubgroupInvocationID == 0) { + minsh[gl_SubgroupID] = min_v; + maxsh[gl_SubgroupID] = max_v; + } + barrier(); + if (tid == 0) { + [[unroll]] for (uint i = 0; i < NUM_SUBGROUPS; ++i) { + min_v = min(min_v, minsh[i]); + max_v = max(max_v, maxsh[i]); + } + if (max_v <= -FLT_MAX_OVER_2) { + result |= 1 << (2*block_x); + } + if (min_v == 0.0f && max_v == 0.0f) { + result |= 2 << (2*block_x); + } + } + barrier(); + } +} + // For each Br x Bc block of the mask (input) buffer, read all values and check // if it's all -inf or all zero. Write out a two-bit code indicating which it is // (or zero for neither). Each workgroup processes 16 tiles and writes out a @@ -48,50 +103,15 @@ void main() { const uint i2 = gl_WorkGroupID.z % nem2; const uint i3 = gl_WorkGroupID.z / nem2; - float FLT_MAX_OVER_2 = uintBitsToFloat(0x7EFFFFFF); - uint result = 0; // Fast path for fully in-bounds blocks where we can do f16vec4 loads if ((nem0 % Bc) == 0 && (nem1 % Br) == 0 && ((Br * Bc) % (BLOCK_SIZE * 4)) == 0) { - [[unroll]] for (uint block_x = 0; block_x < 16; ++block_x) { - float min_v = FLT_MAX_OVER_2; - float max_v = -FLT_MAX_OVER_2; - [[unroll]] for (uint i = 0; i < Br * Bc / 4; i += BLOCK_SIZE) { - uint j0 = (i + tid) % (Bc / 4); - uint j1 = (i + tid) / (Bc / 4); - - j0 *= 4; - j0 += (i0 * 16 + block_x) * Bc; - j1 += i1 * Br; - - vec4 f = vec4(data_av4[(j0 + j1 * nbm1 + i2 * nbm2 + i3 * nbm3) / 4]); - [[unroll]] for (int c = 0; c < 4; ++c) { - min_v = min(min_v, f[c]); - max_v = max(max_v, f[c]); - } - } - min_v = subgroupMin(min_v); - max_v = subgroupMax(max_v); - if (gl_SubgroupInvocationID == 0) { - minsh[gl_SubgroupID] = min_v; - maxsh[gl_SubgroupID] = max_v; - } - barrier(); - if (tid == 0) { - [[unroll]] for (uint i = 0; i < NUM_SUBGROUPS; ++i) { - min_v = min(min_v, minsh[i]); - max_v = max(max_v, maxsh[i]); - } - if (max_v <= -FLT_MAX_OVER_2) { - result |= 1 << (2*block_x); - } - if (min_v == 0.0f && max_v == 0.0f) { - result |= 2 << (2*block_x); - } - } - barrier(); + if ((i0 + 1) * 16 * Bc <= nem0) { + loadvec4(result, i0, i1, i2, i3, false); + } else { + loadvec4(result, i0, i1, i2, i3, true); } } else { [[unroll]] for (uint block_x = 0; block_x < 16; ++block_x) { diff --git a/ggml/src/ggml-vulkan/vulkan-shaders/gated_delta_net.comp b/ggml/src/ggml-vulkan/vulkan-shaders/gated_delta_net.comp new file mode 100644 index 0000000000..1fdf889e82 --- /dev/null +++ b/ggml/src/ggml-vulkan/vulkan-shaders/gated_delta_net.comp @@ -0,0 +1,128 @@ +#version 450 + +#extension GL_EXT_control_flow_attributes : require + +layout(constant_id = 0) const uint S_V = 128; +layout(constant_id = 1) const uint KDA = 0; + +layout(local_size_x_id = 0, local_size_y = 1, local_size_z = 1) in; + +layout(push_constant) uniform Parameters { + uint H; + uint n_tokens; + uint n_seqs; + uint s_off; + uint sq1, sq2, sq3; + uint sv1, sv2, sv3; + uint sb1, sb2, sb3; + uint neq1, rq3; + float scale; +}; + +layout(binding = 0) readonly buffer QBuf { FLOAT_TYPE data_q[]; }; +layout(binding = 1) readonly buffer KBuf { FLOAT_TYPE data_k[]; }; +layout(binding = 2) readonly buffer VBuf { FLOAT_TYPE data_v[]; }; +layout(binding = 3) readonly buffer GBuf { FLOAT_TYPE data_g[]; }; +layout(binding = 4) readonly buffer BetaBuf { FLOAT_TYPE data_beta[]; }; +layout(binding = 5) readonly buffer StateBuf { FLOAT_TYPE data_state[]; }; +layout(binding = 6) buffer DstBuf { FLOAT_TYPE data_dst[]; }; + +shared FLOAT_TYPE s_k[S_V]; +shared FLOAT_TYPE s_q[S_V]; +shared FLOAT_TYPE s_g[S_V]; // KDA only: cached exp(g[i]) + +void main() { + const uint head_id = gl_WorkGroupID.x; + const uint seq_id = gl_WorkGroupID.y; + const uint col = gl_LocalInvocationID.x; + + const uint iq1 = head_id % neq1; + const uint iq3 = seq_id / rq3; + + const uint state_size = S_V * S_V; + const uint state_base = (seq_id * H + head_id) * state_size; + + FLOAT_TYPE state[S_V]; + [[unroll]] for (uint i = 0; i < S_V; i++) { + state[i] = FLOAT_TYPE(data_state[state_base + i * S_V + col]); + } + + uint attn_off = (seq_id * n_tokens * H + head_id) * S_V; + + for (uint t = 0; t < n_tokens; t++) { + const uint q_off = iq3 * sq3 + t * sq2 + iq1 * sq1; + const uint k_off = q_off; + const uint v_off = seq_id * sv3 + t * sv2 + head_id * sv1; + + s_q[col] = FLOAT_TYPE(data_q[q_off + col]); + s_k[col] = FLOAT_TYPE(data_k[k_off + col]); + + const uint gb_off = seq_id * sb3 + t * sb2 + head_id * sb1; + + if (KDA != 0) { + const uint g_base = gb_off * S_V; + s_g[col] = exp(FLOAT_TYPE(data_g[g_base + col])); + } + + barrier(); + + const FLOAT_TYPE v_val = FLOAT_TYPE(data_v[v_off + col]); + const FLOAT_TYPE beta_val = FLOAT_TYPE(data_beta[gb_off]); + + if (KDA == 0) { + const FLOAT_TYPE g_val = exp(FLOAT_TYPE(data_g[gb_off])); + + FLOAT_TYPE kv_col = 0.0; + [[unroll]] for (uint i = 0; i < S_V; i += 4) { + kv_col += dot( + vec4(state[i], state[i+1], state[i+2], state[i+3]), + vec4(s_k[i], s_k[i+1], s_k[i+2], s_k[i+3]) + ); + } + + FLOAT_TYPE delta_col = (v_val - g_val * kv_col) * beta_val; + + FLOAT_TYPE attn_col = 0.0; + [[unroll]] for (uint i = 0; i < S_V; i += 4) { + vec4 sv = vec4(state[i], state[i+1], state[i+2], state[i+3]); + vec4 kv = vec4(s_k[i], s_k[i+1], s_k[i+2], s_k[i+3]); + sv = g_val * sv + kv * delta_col; + state[i] = sv.x; state[i+1] = sv.y; state[i+2] = sv.z; state[i+3] = sv.w; + + attn_col += dot(sv, vec4(s_q[i], s_q[i+1], s_q[i+2], s_q[i+3])); + } + + data_dst[attn_off + col] = attn_col * scale; + } else { + FLOAT_TYPE kv_col = 0.0; + [[unroll]] for (uint i = 0; i < S_V; i += 4) { + vec4 gv = vec4(s_g[i], s_g[i+1], s_g[i+2], s_g[i+3]); + vec4 sv = vec4(state[i], state[i+1], state[i+2], state[i+3]); + vec4 kv = vec4(s_k[i], s_k[i+1], s_k[i+2], s_k[i+3]); + kv_col += dot(gv * sv, kv); + } + + FLOAT_TYPE delta_col = (v_val - kv_col) * beta_val; + + FLOAT_TYPE attn_col = 0.0; + [[unroll]] for (uint i = 0; i < S_V; i += 4) { + vec4 gv = vec4(s_g[i], s_g[i+1], s_g[i+2], s_g[i+3]); + vec4 sv = vec4(state[i], state[i+1], state[i+2], state[i+3]); + vec4 kv = vec4(s_k[i], s_k[i+1], s_k[i+2], s_k[i+3]); + sv = gv * sv + kv * delta_col; + state[i] = sv.x; state[i+1] = sv.y; state[i+2] = sv.z; state[i+3] = sv.w; + + attn_col += dot(sv, vec4(s_q[i], s_q[i+1], s_q[i+2], s_q[i+3])); + } + + data_dst[attn_off + col] = attn_col * scale; + } + + attn_off += S_V * H; + barrier(); + } + + [[unroll]] for (uint i = 0; i < S_V; i++) { + data_dst[s_off + state_base + i * S_V + col] = state[i]; + } +} diff --git a/ggml/src/ggml-vulkan/vulkan-shaders/l2_norm.comp b/ggml/src/ggml-vulkan/vulkan-shaders/l2_norm.comp index 7d0a1de0df..f9af46744d 100644 --- a/ggml/src/ggml-vulkan/vulkan-shaders/l2_norm.comp +++ b/ggml/src/ggml-vulkan/vulkan-shaders/l2_norm.comp @@ -36,7 +36,7 @@ void main() { barrier(); } - const FLOAT_TYPE scale = inversesqrt(max(sum[0], FLOAT_TYPE(p.param1))); + const FLOAT_TYPE scale = 1.0f / max(sqrt(sum[0]), FLOAT_TYPE(p.param1)); [[unroll]] for (uint i0 = tid; i0 < p.ne00; i0 += BLOCK_SIZE) { data_d[i3*p.nb13 + i2*p.nb12 + i1*p.nb11 + i0] = D_TYPE(scale * FLOAT_TYPE(data_a[i3*p.nb03 + i2*p.nb02 + i1*p.nb01 + i0])); diff --git a/ggml/src/ggml-vulkan/vulkan-shaders/ssm_conv.comp b/ggml/src/ggml-vulkan/vulkan-shaders/ssm_conv.comp index d62696bcfa..6802b1fc95 100644 --- a/ggml/src/ggml-vulkan/vulkan-shaders/ssm_conv.comp +++ b/ggml/src/ggml-vulkan/vulkan-shaders/ssm_conv.comp @@ -5,8 +5,9 @@ #include "types.glsl" layout(constant_id = 0) const uint BLOCK_SIZE = 32; +layout(constant_id = 1) const uint TOKENS_PER_WG = 16; -layout(local_size_x_id = 0, local_size_y = 1, local_size_z = 1) in; +layout(local_size_x_id = 0, local_size_y_id = 1, local_size_z = 1) in; layout(binding = 0) readonly buffer Src0 { float src0[]; }; layout(binding = 1) readonly buffer Src1 { float src1[]; }; @@ -20,25 +21,30 @@ layout(push_constant) uniform PushConstants { }; void main() { - const uint global_thread_id = gl_GlobalInvocationID.x; - const uint i2 = gl_WorkGroupID.y; + const uint i1 = gl_GlobalInvocationID.x; + const uint i2 = gl_WorkGroupID.y * TOKENS_PER_WG + gl_LocalInvocationID.y; const uint i3 = gl_WorkGroupID.z; - if (global_thread_id >= nr || i2 >= n_t || i3 >= n_s) { + if (i1 >= nr || i2 >= n_t || i3 >= n_s) { return; } - const uint i1 = global_thread_id; const uint src0_base = i3 * (nb02 / 4) + i2 + i1 * (nb01 / 4); const uint src1_base = i1 * (nb11 / 4); - const uint dst_idx = i3 * (dst_nb2 / 4) + i2 * (dst_nb1 / 4) + i1; float sum = 0.0; - [[unroll]] for (uint i0 = 0; i0 < nc; i0++) { - const uint src0_idx = src0_base + i0; - const uint src1_idx = src1_base + i0; - sum += src0[src0_idx] * src1[src1_idx]; + + if (nc == 4) { + sum = dot( + vec4(src0[src0_base], src0[src0_base + 1], src0[src0_base + 2], src0[src0_base + 3]), + vec4(src1[src1_base], src1[src1_base + 1], src1[src1_base + 2], src1[src1_base + 3]) + ); + } else { + [[unroll]] for (uint i0 = 0; i0 < nc; i0++) { + sum += src0[src0_base + i0] * src1[src1_base + i0]; + } } + const uint dst_idx = i3 * (dst_nb2 / 4) + i2 * (dst_nb1 / 4) + i1; dst[dst_idx] = sum; } diff --git a/ggml/src/ggml-vulkan/vulkan-shaders/vulkan-shaders-gen.cpp b/ggml/src/ggml-vulkan/vulkan-shaders/vulkan-shaders-gen.cpp index fb8941232b..4b00ba3deb 100644 --- a/ggml/src/ggml-vulkan/vulkan-shaders/vulkan-shaders-gen.cpp +++ b/ggml/src/ggml-vulkan/vulkan-shaders/vulkan-shaders-gen.cpp @@ -987,6 +987,8 @@ void process_shaders() { string_to_spv("rwkv_wkv7_f32", "wkv7.comp", merge_maps(base_dict, {{"A_TYPE", "float"}})); + string_to_spv("gated_delta_net_f32", "gated_delta_net.comp", merge_maps(base_dict, {{"FLOAT_TYPE", "float"}})); + string_to_spv("opt_step_adamw_f32", "opt_step_adamw.comp", merge_maps(base_dict, {{"A_TYPE", "float"}})); string_to_spv("opt_step_sgd_f32", "opt_step_sgd.comp", merge_maps(base_dict, {{"A_TYPE", "float"}})); diff --git a/scripts/sync_vendor.py b/scripts/sync_vendor.py index a8bb6c4ffc..75d4a5ff61 100755 --- a/scripts/sync_vendor.py +++ b/scripts/sync_vendor.py @@ -5,7 +5,7 @@ import os import sys import subprocess -HTTPLIB_VERSION = "refs/tags/v0.37.0" +HTTPLIB_VERSION = "refs/tags/v0.37.1" vendor = { "https://github.com/nlohmann/json/releases/latest/download/json.hpp": "vendor/nlohmann/json.hpp", diff --git a/src/llama-context.cpp b/src/llama-context.cpp index 0be9493910..14dccac5b5 100644 --- a/src/llama-context.cpp +++ b/src/llama-context.cpp @@ -7,6 +7,7 @@ #include "llama-memory.h" #include "llama-mmap.h" #include "llama-model.h" +#include "llama-ext.h" #include #include @@ -341,6 +342,14 @@ llama_context::llama_context( if (cparams.pipeline_parallel) { LLAMA_LOG_INFO("%s: pipeline parallelism enabled\n", __func__); + + if (!graph_reuse_disable) { + // TODO: figure out a way to make graph reuse work with pipeline parallelism + // ref: https://github.com/ggml-org/llama.cpp/pull/20463 + LLAMA_LOG_WARN("%s: graph reuse is currently not compatible with pipeline parallelism - disabling\n", __func__); + + graph_reuse_disable = true; + } } sched_reserve(); @@ -3129,6 +3138,19 @@ uint32_t llama_get_sampled_probs_count_ith(llama_context * ctx, int32_t i) { return static_cast(ctx->get_sampled_probs_count(i)); } +struct ggml_cgraph * llama_graph_reserve( + struct llama_context * ctx, + uint32_t n_tokens, + uint32_t n_seqs, + uint32_t n_outputs) { + auto * memory = ctx->get_memory(); + llama_memory_context_ptr mctx; + if (memory) { + mctx = memory->init_full(); + } + return ctx->graph_reserve(n_tokens, n_seqs, n_outputs, mctx.get()); +} + // llama adapter API int32_t llama_set_adapters_lora( diff --git a/src/llama-ext.h b/src/llama-ext.h new file mode 100644 index 0000000000..13ced783b4 --- /dev/null +++ b/src/llama-ext.h @@ -0,0 +1,12 @@ +#pragma once + +#include "llama-context.h" +#include "ggml.h" +#include "stdint.h" + +// Reserve a new compute graph. It is valid until the next call to llama_graph_reserve. +LLAMA_API struct ggml_cgraph * llama_graph_reserve( + struct llama_context * ctx, + uint32_t n_tokens, + uint32_t n_seqs, + uint32_t n_outputs); diff --git a/src/llama-grammar.cpp b/src/llama-grammar.cpp index 3b7a625234..aac0d41f2b 100644 --- a/src/llama-grammar.cpp +++ b/src/llama-grammar.cpp @@ -1160,13 +1160,13 @@ struct llama_grammar * llama_grammar_init_impl( // if there is a grammar, parse it // rules will be empty (default) if there are parse errors if (!parser.parse(grammar_str) || parser.rules.empty()) { - fprintf(stderr, "%s: failed to parse grammar\n", __func__); + LLAMA_LOG_ERROR("failed to parse grammar\n"); return nullptr; } - // Ensure that there is a "root" node. - if (parser.symbol_ids.find("root") == parser.symbol_ids.end()) { - fprintf(stderr, "%s: grammar does not contain a 'root' symbol\n", __func__); + // Ensure that the grammar contains the start symbol + if (parser.symbol_ids.find(grammar_root) == parser.symbol_ids.end()) { + LLAMA_LOG_ERROR("grammar does not contain a '%s' symbol\n", grammar_root); return nullptr; } @@ -1195,7 +1195,7 @@ struct llama_grammar * llama_grammar_init_impl( continue; } if (llama_grammar_detect_left_recursion(vec_rules, i, &rules_visited, &rules_in_progress, &rules_may_be_empty)) { - LLAMA_LOG_ERROR("unsupported grammar, left recursion detected for nonterminal at index %zu", i); + LLAMA_LOG_ERROR("unsupported grammar, left recursion detected for nonterminal at index %zu\n", i); return nullptr; } } diff --git a/tests/CMakeLists.txt b/tests/CMakeLists.txt index bb0f0ef0ed..9582164b58 100644 --- a/tests/CMakeLists.txt +++ b/tests/CMakeLists.txt @@ -260,6 +260,7 @@ endif() set(LLAMA_TEST_NAME test-mtmd-c-api) llama_build_and_test(test-mtmd-c-api.c) target_link_libraries(${LLAMA_TEST_NAME} PRIVATE mtmd) +unset(LLAMA_TEST_NAME) # GGUF model data fetcher library for tests that need real model metadata # Only compile when cpp-httplib has SSL support (CPPHTTPLIB_OPENSSL_SUPPORT) @@ -284,4 +285,5 @@ target_link_libraries(${TEST_TARGET} PRIVATE llama) llama_build_and_test(test-alloc.cpp) target_include_directories(test-alloc PRIVATE ${PROJECT_SOURCE_DIR}/ggml/src) - +llama_build(export-graph-ops.cpp) +target_include_directories(export-graph-ops PRIVATE ${PROJECT_SOURCE_DIR}/ggml/src) diff --git a/tests/export-graph-ops.cpp b/tests/export-graph-ops.cpp new file mode 100644 index 0000000000..754089d068 --- /dev/null +++ b/tests/export-graph-ops.cpp @@ -0,0 +1,169 @@ +#include "arg.h" +#include "common.h" +#include "log.h" +#include "llama.h" +#include "../src/llama-ext.h" +#include "ggml.h" + +#include +#include +#include +#include +#include + +struct input_tensor { + ggml_type type; + std::array ne; + std::array nb; + + input_tensor(ggml_type type, int64_t * ne, size_t * nb): type(type) { + memcpy(this->ne.data(), ne, 4 * sizeof(int64_t)); + memcpy(this->nb.data(), nb, 4 * sizeof(size_t)); + } + + bool operator<(const input_tensor &b) const { + return std::tie(type, ne, nb) < + std::tie(b.type, b.ne, b.nb); + } + + void serialize(std::ostream& out) const { + out << type << ' '; + for (size_t i = 0; i < 4; i++) { + out << ne[i] << ' '; + } + for (size_t i = 0; i < 4; i++) { + out << nb[i] << ' '; + } + } +}; + +struct test_object { + ggml_op op; + ggml_type type; + std::array ne; + std::vector op_params; + std::vector sources; + std::string name; + + void serialize(std::ostream& out) const { + out << op << ' ' << type << ' '; + for (size_t i = 0; i < 4; i++) { + out << ne[i] << ' '; + } + + out << op_params.size() << ' '; + for (size_t i = 0; i < op_params.size(); i++) { + out << op_params[i] << ' '; + } + + out << sources.size() << ' '; + for (size_t s = 0; s < sources.size(); s++) { + sources[s].serialize(out); + } + + if (!name.empty()) { + out << name; + } else { + out << '-'; + } + + out << '\n'; + } + + bool operator<(const test_object &b) const { + return std::tie(op, type, ne, op_params, sources) < + std::tie(b.op, b.type, b.ne, b.op_params, b.sources); + } +}; + +static void extract_graph_ops(ggml_cgraph * cgraph, const char * label, std::set & tests) { + int n_nodes = ggml_graph_n_nodes(cgraph); + int n_skipped = 0; + int n_before = (int) tests.size(); + for (int i = 0; i < n_nodes; i++) { + ggml_tensor * node = ggml_graph_node(cgraph, i); + + if (node->op == GGML_OP_NONE || node->op == GGML_OP_VIEW || node->op == GGML_OP_RESHAPE || node->op == GGML_OP_PERMUTE || node->op == GGML_OP_TRANSPOSE) { + n_skipped++; + continue; + } + + test_object test; + + test.op = node->op; + test.type = node->type; + memcpy(&test.ne, node->ne, 4 * sizeof(int64_t)); + + test.op_params.resize(GGML_MAX_OP_PARAMS / sizeof(int32_t)); + memcpy(test.op_params.data(), node->op_params, GGML_MAX_OP_PARAMS); + + for (size_t s = 0; s < GGML_MAX_SRC; s++) { + if (node->src[s] == nullptr) { + break; + } + + test.sources.emplace_back(node->src[s]->type, node->src[s]->ne, node->src[s]->nb); + } + + test.name = node->name; + tests.insert(test); + } + + int n_new = (int) tests.size() - n_before; + LOG_INF("%s: %d unique ops, %d total nodes, %d skipped (view ops)\n", + label, n_new, n_nodes, n_skipped); +} + +int main(int argc, char ** argv) { + common_params params; + params.out_file = "tests.txt"; + + if (!common_params_parse(argc, argv, params, LLAMA_EXAMPLE_EXPORT_GRAPH_OPS)) { + return 1; + } + + common_init(); + + // Load CPU-only + ggml_backend_dev_t cpu_device = ggml_backend_dev_by_type(GGML_BACKEND_DEVICE_TYPE_CPU); + params.devices = { cpu_device, nullptr }; + params.fit_params = false; + params.n_gpu_layers = 0; + + params.warmup = false; + + auto init_result = common_init_from_params(params); + + llama_context * ctx = init_result->context(); + + const uint32_t n_seqs = llama_n_seq_max(ctx); + const uint32_t n_tokens = std::min(llama_n_ctx(ctx), llama_n_ubatch(ctx)); + + std::set tests; + + auto * gf_pp = llama_graph_reserve(ctx, n_tokens, n_seqs, n_tokens); + if (!gf_pp) { + throw std::runtime_error("failed to reserve prompt processing graph"); + } + extract_graph_ops(gf_pp, "pp", tests); + + auto * gf_tg = llama_graph_reserve(ctx, n_seqs, n_seqs, n_seqs); + if (!gf_tg) { + throw std::runtime_error("failed to reserve token generation graph"); + } + extract_graph_ops(gf_tg, "tg", tests); + + LOG_INF("%d unique ops total\n", (int) tests.size()); + + std::ofstream f(params.out_file); + + if (!f.is_open()) { + throw std::runtime_error("Unable to open output file"); + } + + for (const auto& test : tests) { + test.serialize(f); + } + + return 0; +} diff --git a/tests/test-backend-ops.cpp b/tests/test-backend-ops.cpp index 58d67d97f8..abf914faa1 100644 --- a/tests/test-backend-ops.cpp +++ b/tests/test-backend-ops.cpp @@ -31,10 +31,12 @@ #include #include #include +#include #include #include #include #include +#include #include #include #include @@ -6648,6 +6650,236 @@ struct test_diag : public test_case { } }; +// Deserializable generic test case +struct input_tensor { + ggml_type type; + std::array ne; + std::array nb; // strides (0 = use default contiguous strides) +}; + +static bool is_non_contiguous(const input_tensor & src) { + if (src.nb[0] == 0) { + return false; + } + const size_t default_nb0 = ggml_type_size(src.type); + const size_t default_nb1 = default_nb0 * (src.ne[0] / ggml_blck_size(src.type)); + const size_t default_nb2 = default_nb1 * src.ne[1]; + const size_t default_nb3 = default_nb2 * src.ne[2]; + return src.nb[0] != default_nb0 || + src.nb[1] != default_nb1 || + src.nb[2] != default_nb2 || + src.nb[3] != default_nb3; +} + +static std::string var_to_str(const std::vector& sources) { + std::ostringstream oss; + bool first = true; + for (const auto& src : sources) { + if (!first) oss << ","; + oss << ggml_type_name(src.type) << "[" << src.ne[0] << "," << src.ne[1] << "," << src.ne[2] << "," << src.ne[3] << "]"; + if (is_non_contiguous(src)) { + oss << "nb[" << src.nb[0] << "," << src.nb[1] << "," << src.nb[2] << "," << src.nb[3] << "]"; + } + first = false; + } + return oss.str(); +} + +static std::string var_to_str(const std::array& params) { + std::ostringstream oss; + oss << "["; + bool first = true; + for (size_t i = 0; i < params.size(); ++i) { + if (params[i] != 0) { + if (!first) oss << ","; + oss << i << ":" << params[i]; + first = false; + } + } + oss << "]"; + return oss.str(); +} + + +struct test_generic_op : public test_case { + const ggml_op op; + const ggml_type type; + const std::array ne; + const std::array op_params; + + const std::vector sources; + const std::string name; + + std::string vars() override { + if (name.empty()) { + return VARS_TO_STR4(type, ne, op_params, sources); + } + + return VARS_TO_STR5(name, type, ne, op_params, sources); + } + + test_generic_op(ggml_op op, ggml_type type, std::array ne, + std::array op_params, + std::vector sources, std::string name = "") + : op(op), type(type), ne(ne), op_params(op_params), sources(sources), name(std::move(name)) {} + + ggml_tensor * build_graph(ggml_context * ctx) override { + const size_t source_count = std::min(sources.size(), (size_t)GGML_MAX_SRC); + + std::array source_tensors; + for (size_t i = 0; i < source_count; ++i) { + const input_tensor& src = sources[i]; + + if (is_non_contiguous(src)) { + size_t total_size; + const size_t blck_size = ggml_blck_size(src.type); + if (blck_size == 1) { + total_size = ggml_type_size(src.type); + for (int d = 0; d < 4; d++) { + total_size += (src.ne[d] - 1) * src.nb[d]; + } + } else { + total_size = src.ne[0] * src.nb[0] / blck_size; + for (int d = 1; d < 4; d++) { + total_size += (src.ne[d] - 1) * src.nb[d]; + } + } + + // Convert bytes to elements, padded to block size for quantized types + const size_t type_size = ggml_type_size(src.type); + size_t backing_elements = (total_size * blck_size + type_size - 1) / type_size; + backing_elements = ((backing_elements + blck_size - 1) / blck_size) * blck_size; + ggml_tensor * backing = ggml_new_tensor_1d(ctx, src.type, backing_elements); + source_tensors[i] = ggml_view_4d(ctx, backing, + src.ne[0], src.ne[1], src.ne[2], src.ne[3], + src.nb[1], src.nb[2], src.nb[3], 0); + // nb[0] does not get set by view_4d, so set it manually + source_tensors[i]->nb[0] = src.nb[0]; + } else { + source_tensors[i] = ggml_new_tensor_4d(ctx, src.type, src.ne[0], src.ne[1], src.ne[2], src.ne[3]); + } + } + + // Ops with an inplace flag create a view of src[0] as their output. + bool inplace = false; + if (op == GGML_OP_SET || op == GGML_OP_ACC) { + inplace = op_params[4] != 0; + } else if (op == GGML_OP_ADD_REL_POS) { + inplace = op_params[0] != 0; + } + + ggml_tensor * out; + if (inplace && source_count > 0) { + out = ggml_view_tensor(ctx, source_tensors[0]); + } else { + out = ggml_new_tensor_4d(ctx, type, ne[0], ne[1], ne[2], ne[3]); + } + out->op = op; + for (size_t i = 0; i < source_count; ++i) { + out->src[i] = source_tensors[i]; + } + + memcpy(out->op_params, op_params.data(), GGML_MAX_OP_PARAMS); + ggml_set_name(out, "out"); + + return out; + } + + double max_nmse_err() override { + switch (op) { + case GGML_OP_MUL_MAT: + case GGML_OP_MUL_MAT_ID: + case GGML_OP_OUT_PROD: + case GGML_OP_CONV_TRANSPOSE_2D: + case GGML_OP_IM2COL: + case GGML_OP_CONV_2D: + case GGML_OP_CONV_3D: + case GGML_OP_SET_ROWS: + case GGML_OP_CPY: + return 5e-4; + case GGML_OP_SOFT_MAX: + return 1e-6; + case GGML_OP_RWKV_WKV7: + return 5e-3; + case GGML_OP_FLASH_ATTN_EXT: + { + // Scale error with kv length to account for accumulating floating point error + const int64_t kv = sources[1].ne[1]; + return 5e-4 * std::max(1.0, kv / 20000.0); + } + default: + return 1e-7; + } + } + + void initialize_tensors(ggml_context * ctx) override { + ggml_tensor * out = ggml_get_tensor(ctx, "out"); + + std::random_device rd; + std::default_random_engine rng(rd()); + + for (size_t i = 0; i < sources.size() && i < GGML_MAX_SRC; i++) { + ggml_tensor * t = out->src[i]; + if (!t) { + break; + } + + // FLASH_ATTN_EXT: src[3] is the KQ mask + if (op == GGML_OP_FLASH_ATTN_EXT && i == 3) { + init_tensor_kq_mask(t); + continue; + } + + if (t->type == GGML_TYPE_I32 || t->type == GGML_TYPE_I64) { + if (op == GGML_OP_GET_ROWS || op == GGML_OP_GET_ROWS_BACK) { + const int64_t num_rows = sources[0].ne[1]; + const int64_t nels = ggml_nelements(t); + std::vector data(nels); + std::uniform_int_distribution dist(0, num_rows - 1); + for (int64_t i = 0; i < nels; i++) { + data[i] = dist(rng); + } + ggml_backend_tensor_set(t, data.data(), 0, nels * sizeof(int32_t)); + } else if (op == GGML_OP_SET_ROWS) { + init_set_rows_row_ids(t, ne[1]); + } else if (op == GGML_OP_ROPE) { + const int mode = op_params[2]; + const int64_t nels = (mode & GGML_ROPE_TYPE_MROPE) ? ne[2] * 4 : ne[2]; + std::vector data(nels); + std::uniform_int_distribution dist(0, ne[2] - 1); + for (int64_t i = 0; i < nels; i++) { + data[i] = dist(rng); + } + ggml_backend_tensor_set(t, data.data(), 0, nels * sizeof(int32_t)); + } else if (op == GGML_OP_MUL_MAT_ID || op == GGML_OP_ADD_ID) { + const int64_t n_expert = (op == GGML_OP_MUL_MAT_ID) ? sources[0].ne[2] : sources[1].ne[1]; + for (int64_t r = 0; r < ggml_nrows(t); r++) { + std::vector data(t->ne[0]); + for (int32_t i = 0; i < t->ne[0]; i++) { + data[i] = i % n_expert; + } + std::shuffle(data.begin(), data.end(), rng); + ggml_backend_tensor_set(t, data.data(), r * t->nb[1], t->ne[0] * sizeof(int32_t)); + } + } else if (op == GGML_OP_SSM_SCAN) { + for (int64_t r = 0; r < ggml_nrows(t); r++) { + std::vector data(t->ne[0]); + for (int32_t i = 0; i < t->ne[0]; i++) { + data[i] = i; + } + std::shuffle(data.begin(), data.end(), rng); + ggml_backend_tensor_set(t, data.data(), r * t->nb[1], t->ne[0] * sizeof(int32_t)); + } + } else { + init_tensor_uniform(t); + } + } else { + init_tensor_uniform(t); + } + } + } +}; + enum llm_norm_type { LLM_NORM, @@ -7656,7 +7888,7 @@ static std::vector> make_test_cases_eval() { test_cases.emplace_back(new test_softcap(GGML_TYPE_F32, {10, 10, 10, 10}, 50.0f)); test_cases.emplace_back(new test_silu_back()); - for (float eps : { 0.0f, 1e-6f, 1e-4f, 1e-1f }) { + for (float eps : { 0.0f, 1e-6f, 1e-4f, 1e-1f, 10.f }) { for (uint32_t n : { 64, 1025 }) { for (bool v : { false, true }) { test_cases.emplace_back(new test_norm(GGML_TYPE_F32, { n, 5, 4, 3 }, v, eps)); @@ -8731,11 +8963,92 @@ static std::vector> make_test_cases_perf() { test_cases.emplace_back(new test_acc(GGML_TYPE_F32, {256, 17, 2, 3}, {128, 16, 2, 3}, 2)); test_cases.emplace_back(new test_acc(GGML_TYPE_F32, {256, 17, 2, 3}, {64, 16, 2, 3}, 3)); + // GATED_DELTA_NET: realistic model configurations + // TG: n_seq_tokens=1 (autoregressive) + test_cases.emplace_back(new test_gated_delta_net(GGML_TYPE_F32, 32, 128, 1, 1)); // Qwen3.5-like: 32 heads, d=128 + test_cases.emplace_back(new test_gated_delta_net(GGML_TYPE_F32, 16, 64, 1, 1)); // smaller model + test_cases.emplace_back(new test_gated_delta_net(GGML_TYPE_F32, 32, 128, 1, 1, 1, false, true)); // KDA + // PP: n_seq_tokens=64,256 (prompt processing) + test_cases.emplace_back(new test_gated_delta_net(GGML_TYPE_F32, 32, 128, 64, 1)); // PP-64 + test_cases.emplace_back(new test_gated_delta_net(GGML_TYPE_F32, 32, 128, 256, 1)); // PP-256 + test_cases.emplace_back(new test_gated_delta_net(GGML_TYPE_F32, 32, 128, 512, 1)); // PP-512 + test_cases.emplace_back(new test_gated_delta_net(GGML_TYPE_F32, 32, 128, 1024, 1)); // PP-1024 + // Small model configs (fewer heads = less GPU occupancy for autoregressive) + test_cases.emplace_back(new test_gated_delta_net(GGML_TYPE_F32, 4, 128, 64, 1)); // 4h PP-64 + test_cases.emplace_back(new test_gated_delta_net(GGML_TYPE_F32, 4, 128, 256, 1)); // 4h PP-256 + test_cases.emplace_back(new test_gated_delta_net(GGML_TYPE_F32, 4, 128, 512, 1)); // 4h PP-512 + test_cases.emplace_back(new test_gated_delta_net(GGML_TYPE_F32, 4, 128, 1024, 1)); // 4h PP-1024 + test_cases.emplace_back(new test_gated_delta_net(GGML_TYPE_F32, 32, 128, 64, 1, 1, false, true)); // KDA PP-64 + + return test_cases; +} + +static std::vector> make_test_cases_from_file(const char * path) { + std::ifstream f(path); + + if (!f.is_open()) { + throw std::runtime_error("Unable to read test file"); + } + + std::vector> test_cases; + + std::string line; + + while (std::getline(f, line)) { + std::istringstream iss(line); + + ggml_op op; + ggml_type type; + std::array ne; + std::array op_params = {}; + std::string name; + uint64_t tmp; + + iss >> tmp; + op = (ggml_op)tmp; + iss >> tmp; + type = (ggml_type)tmp; + + for (size_t i = 0; i < 4; i++) { + iss >> ne[i]; + } + + iss >> tmp; + for (size_t i = 0; i < tmp && i < op_params.size(); i++) { + iss >> op_params[i]; + } + + iss >> tmp; + + size_t num_src = std::min((uint64_t)GGML_MAX_SRC, tmp); + std::vector sources(num_src); + for (size_t i = 0; i < num_src; i++) { + input_tensor& src = sources[i]; + iss >> tmp; + src.type = (ggml_type)tmp; + + for (size_t i = 0; i < 4; i++) { + iss >> src.ne[i]; + } + for (size_t i = 0; i < 4; i++) { + iss >> src.nb[i]; + } + } + + iss >> name; + + if (name.length() == 1 && name[0] == '-') { + name = ""; + } + + test_cases.emplace_back(new test_generic_op(op, type, ne, op_params, sources, std::move(name))); + } + return test_cases; } static bool test_backend(ggml_backend_t backend, test_mode mode, const char * op_names_filter, const char * params_filter, - printer * output_printer) { + printer * output_printer, const char * test_file_path) { auto filter_test_cases = [](std::vector> & test_cases, const char * params_filter) { if (params_filter == nullptr) { return; @@ -8753,9 +9066,26 @@ static bool test_backend(ggml_backend_t backend, test_mode mode, const char * op } }; + std::vector> test_cases; + + if (test_file_path == nullptr) { + switch (mode) { + case MODE_TEST: + case MODE_GRAD: + case MODE_SUPPORT: + test_cases = make_test_cases_eval(); + break; + case MODE_PERF: + test_cases = make_test_cases_perf(); + break; + } + } else { + test_cases = make_test_cases_from_file(test_file_path); + } + + filter_test_cases(test_cases, params_filter); + if (mode == MODE_TEST) { - auto test_cases = make_test_cases_eval(); - filter_test_cases(test_cases, params_filter); ggml_backend_t backend_cpu = ggml_backend_init_by_type(GGML_BACKEND_DEVICE_TYPE_CPU, NULL); if (backend_cpu == NULL) { test_operation_info info("", "", "CPU"); @@ -8795,8 +9125,6 @@ static bool test_backend(ggml_backend_t backend, test_mode mode, const char * op } if (mode == MODE_GRAD) { - auto test_cases = make_test_cases_eval(); - filter_test_cases(test_cases, params_filter); size_t n_ok = 0; for (auto & test : test_cases) { if (test->eval_grad(backend, op_names_filter, output_printer)) { @@ -8809,8 +9137,6 @@ static bool test_backend(ggml_backend_t backend, test_mode mode, const char * op } if (mode == MODE_PERF) { - auto test_cases = make_test_cases_perf(); - filter_test_cases(test_cases, params_filter); for (auto & test : test_cases) { test->eval_perf(backend, op_names_filter, output_printer); } @@ -8818,9 +9144,6 @@ static bool test_backend(ggml_backend_t backend, test_mode mode, const char * op } if (mode == MODE_SUPPORT) { - auto test_cases = make_test_cases_eval(); - filter_test_cases(test_cases, params_filter); - // Filter out fusion cases test_cases.erase( std::remove_if(test_cases.begin(), test_cases.end(), [](const std::unique_ptr & tc) { @@ -8939,7 +9262,8 @@ static void show_test_coverage() { } static void usage(char ** argv) { - printf("Usage: %s [mode] [-o ] [-b ] [-p ] [--output ] [--list-ops] [--show-coverage]\n", argv[0]); + printf("Usage: %s [mode] [-o ] [-b ] [-p ] [--output ] [--list-ops]", argv[0]); + printf(" [--show-coverage] [--test-file ]\n"); printf(" valid modes:\n"); printf(" - test (default, compare with CPU backend for correctness)\n"); printf(" - grad (compare gradients from backpropagation with method of finite differences)\n"); @@ -8950,6 +9274,7 @@ static void usage(char ** argv) { printf(" --output specifies output format (default: console, options: console, sql, csv)\n"); printf(" --list-ops lists all available GGML operations\n"); printf(" --show-coverage shows test coverage\n"); + printf(" --test-file reads test operators from a test file generated by llama-export-graph-ops\n"); } int main(int argc, char ** argv) { @@ -8958,6 +9283,7 @@ int main(int argc, char ** argv) { const char * op_names_filter = nullptr; const char * backend_filter = nullptr; const char * params_filter = nullptr; + const char * test_file_path = nullptr; for (int i = 1; i < argc; i++) { if (strcmp(argv[i], "test") == 0) { @@ -9005,6 +9331,13 @@ int main(int argc, char ** argv) { } else if (strcmp(argv[i], "--show-coverage") == 0) { show_test_coverage(); return 0; + } else if (strcmp(argv[i], "--test-file") == 0) { + if (i + 1 < argc) { + test_file_path = argv[++i]; + } else { + usage(argv); + return 1; + } } else { usage(argv); return 1; @@ -9057,7 +9390,7 @@ int main(int argc, char ** argv) { false, "", ggml_backend_dev_description(dev), total / 1024 / 1024, free / 1024 / 1024, true)); - bool ok = test_backend(backend, mode, op_names_filter, params_filter, output_printer.get()); + bool ok = test_backend(backend, mode, op_names_filter, params_filter, output_printer.get(), test_file_path); if (ok) { n_ok++; diff --git a/tests/test-grammar-integration.cpp b/tests/test-grammar-integration.cpp index 7aa7e58a5c..526470a224 100644 --- a/tests/test-grammar-integration.cpp +++ b/tests/test-grammar-integration.cpp @@ -15,8 +15,12 @@ using json = nlohmann::ordered_json; +static llama_grammar * build_grammar_with_root(const std::string & grammar_str, const char * grammar_root) { + return llama_grammar_init_impl(nullptr, grammar_str.c_str(), grammar_root, false, nullptr, 0, nullptr, 0); +} + static llama_grammar * build_grammar(const std::string & grammar_str) { - return llama_grammar_init_impl(nullptr, grammar_str.c_str(), "root", false, nullptr, 0, nullptr, 0); + return build_grammar_with_root(grammar_str, "root"); } static bool test_build_grammar_fails(const std::string & grammar_str) { @@ -860,6 +864,36 @@ static void test_failure_left_recursion() { fprintf(stderr, " ✅︎ Passed\n"); } +static void test_failure_missing_root_symbol() { + fprintf(stderr, "⚫ Testing missing root symbol:\n"); + + const std::string grammar_str = R"""( + root ::= "foobar" + )"""; + + llama_grammar * failure_result = build_grammar_with_root(grammar_str, "nonexistent"); + assert(failure_result == nullptr); + + fprintf(stderr, " ✅︎ Passed\n"); +} + +static void test_custom_root_symbol_check() { + fprintf(stderr, "⚫ Testing custom root symbol check:\n"); + + const std::string custom_root_grammar_str = R"""( + foobar ::= "foobar" + )"""; + + llama_grammar * failure_result = build_grammar_with_root(custom_root_grammar_str, "root"); + assert(failure_result == nullptr); + + llama_grammar * success_result = build_grammar_with_root(custom_root_grammar_str, "foobar"); + assert(success_result != nullptr); + llama_grammar_free_impl(success_result); + + fprintf(stderr, " ✅︎ Passed\n"); +} + static void test_json_schema() { // Note that this is similar to the regular grammar tests, // but we convert each json schema to a grammar before parsing. @@ -1433,6 +1467,8 @@ int main() { test_failure_missing_root(); test_failure_missing_reference(); test_failure_left_recursion(); + test_failure_missing_root_symbol(); + test_custom_root_symbol_check(); test_json_schema(); fprintf(stdout, "All tests passed.\n"); return 0; diff --git a/tools/server/public/index.html.gz b/tools/server/public/index.html.gz index 3d0991dde0..493058aa01 100644 Binary files a/tools/server/public/index.html.gz and b/tools/server/public/index.html.gz differ diff --git a/tools/server/tests/unit/test_template.py b/tools/server/tests/unit/test_template.py index e5185fcbfa..43a356020d 100644 --- a/tools/server/tests/unit/test_template.py +++ b/tools/server/tests/unit/test_template.py @@ -11,6 +11,7 @@ sys.path.insert(0, str(path)) import datetime from utils import * +from typing import Literal server: ServerProcess @@ -23,24 +24,24 @@ def create_server(): @pytest.mark.parametrize("tools", [None, [], [TEST_TOOL]]) -@pytest.mark.parametrize("template_name,reasoning_budget,expected_end", [ - ("deepseek-ai-DeepSeek-R1-Distill-Qwen-32B", None, "\n"), - ("deepseek-ai-DeepSeek-R1-Distill-Qwen-32B", -1, "\n"), - ("deepseek-ai-DeepSeek-R1-Distill-Qwen-32B", 0, "\n"), +@pytest.mark.parametrize("template_name,reasoning,expected_end", [ + ("deepseek-ai-DeepSeek-R1-Distill-Qwen-32B", "on", "\n"), + ("deepseek-ai-DeepSeek-R1-Distill-Qwen-32B","auto", "\n"), + ("deepseek-ai-DeepSeek-R1-Distill-Qwen-32B", "off", "\n"), - ("Qwen-Qwen3-0.6B", -1, "<|im_start|>assistant\n"), - ("Qwen-Qwen3-0.6B", 0, "<|im_start|>assistant\n\n\n\n\n"), + ("Qwen-Qwen3-0.6B","auto", "<|im_start|>assistant\n"), + ("Qwen-Qwen3-0.6B", "off", "<|im_start|>assistant\n\n\n\n\n"), - ("Qwen-QwQ-32B", -1, "<|im_start|>assistant\n\n"), - ("Qwen-QwQ-32B", 0, "<|im_start|>assistant\n\n"), + ("Qwen-QwQ-32B","auto", "<|im_start|>assistant\n\n"), + ("Qwen-QwQ-32B", "off", "<|im_start|>assistant\n\n"), - ("CohereForAI-c4ai-command-r7b-12-2024-tool_use", -1, "<|START_OF_TURN_TOKEN|><|CHATBOT_TOKEN|>"), - ("CohereForAI-c4ai-command-r7b-12-2024-tool_use", 0, "<|START_OF_TURN_TOKEN|><|CHATBOT_TOKEN|><|START_THINKING|><|END_THINKING|>"), + ("CohereForAI-c4ai-command-r7b-12-2024-tool_use","auto", "<|START_OF_TURN_TOKEN|><|CHATBOT_TOKEN|>"), + ("CohereForAI-c4ai-command-r7b-12-2024-tool_use", "off", "<|START_OF_TURN_TOKEN|><|CHATBOT_TOKEN|><|START_THINKING|><|END_THINKING|>"), ]) -def test_reasoning_budget(template_name: str, reasoning_budget: int | None, expected_end: str, tools: list[dict]): +def test_reasoning(template_name: str, reasoning: Literal['on', 'off', 'auto'] | None, expected_end: str, tools: list[dict]): global server server.jinja = True - server.reasoning_budget = reasoning_budget + server.reasoning = reasoning server.chat_template_file = f'../../../models/templates/{template_name}.jinja' server.start() diff --git a/tools/server/tests/utils.py b/tools/server/tests/utils.py index db357d876b..c6fe11261f 100644 --- a/tools/server/tests/utils.py +++ b/tools/server/tests/utils.py @@ -95,7 +95,7 @@ class ServerProcess: no_webui: bool | None = None jinja: bool | None = None reasoning_format: Literal['deepseek', 'none', 'nothink'] | None = None - reasoning_budget: int | None = None + reasoning: Literal['on', 'off', 'auto'] | None = None chat_template: str | None = None chat_template_file: str | None = None server_path: str | None = None @@ -225,8 +225,8 @@ class ServerProcess: server_args.append("--no-jinja") if self.reasoning_format is not None: server_args.extend(("--reasoning-format", self.reasoning_format)) - if self.reasoning_budget is not None: - server_args.extend(("--reasoning-budget", self.reasoning_budget)) + if self.reasoning is not None: + server_args.extend(("--reasoning", self.reasoning)) if self.chat_template: server_args.extend(["--chat-template", self.chat_template]) if self.chat_template_file: diff --git a/tools/server/webui/src/lib/components/app/chat/ChatForm/ChatFormActions/ChatFormActions.svelte b/tools/server/webui/src/lib/components/app/chat/ChatForm/ChatFormActions/ChatFormActions.svelte index 8501776933..2ad830e18f 100644 --- a/tools/server/webui/src/lib/components/app/chat/ChatForm/ChatFormActions/ChatFormActions.svelte +++ b/tools/server/webui/src/lib/components/app/chat/ChatForm/ChatFormActions/ChatFormActions.svelte @@ -62,15 +62,12 @@ chatStore.getConversationModel(activeMessages() as DatabaseMessage[]) ); - let previousConversationModel: string | null = null; - $effect(() => { - if (conversationModel && conversationModel !== previousConversationModel) { - previousConversationModel = conversationModel; - - if (!isRouter || modelsStore.isModelLoaded(conversationModel)) { - modelsStore.selectModelByName(conversationModel); - } + if (conversationModel) { + modelsStore.selectModelByName(conversationModel); + } else if (isRouter && modelsStore.loadedModelIds.length > 0) { + const first = modelOptions().find((m) => modelsStore.loadedModelIds.includes(m.model)); + if (first) modelsStore.selectModelById(first.id); } }); diff --git a/vendor/cpp-httplib/httplib.cpp b/vendor/cpp-httplib/httplib.cpp index c8f88d87df..71a5f00567 100644 --- a/vendor/cpp-httplib/httplib.cpp +++ b/vendor/cpp-httplib/httplib.cpp @@ -4424,7 +4424,8 @@ get_range_offset_and_length(Range r, size_t content_length) { assert(r.first <= r.second && r.second < static_cast(content_length)); (void)(content_length); - return std::make_pair(r.first, static_cast(r.second - r.first) + 1); + return std::make_pair(static_cast(r.first), + static_cast(r.second - r.first) + 1); } std::string make_content_range_header_field( @@ -8616,11 +8617,17 @@ ClientImpl::open_stream(const std::string &method, const std::string &path, handle.body_reader_.stream = handle.stream_; handle.body_reader_.payload_max_length = payload_max_length_; - auto content_length_str = handle.response->get_header_value("Content-Length"); - if (!content_length_str.empty()) { + if (handle.response->has_header("Content-Length")) { + bool is_invalid = false; + auto content_length = detail::get_header_value_u64( + handle.response->headers, "Content-Length", 0, 0, is_invalid); + if (is_invalid) { + handle.error = Error::Read; + handle.response.reset(); + return handle; + } handle.body_reader_.has_content_length = true; - handle.body_reader_.content_length = - static_cast(std::stoull(content_length_str)); + handle.body_reader_.content_length = content_length; } auto transfer_encoding = diff --git a/vendor/cpp-httplib/httplib.h b/vendor/cpp-httplib/httplib.h index ac1908f421..e01b3550ba 100644 --- a/vendor/cpp-httplib/httplib.h +++ b/vendor/cpp-httplib/httplib.h @@ -8,28 +8,8 @@ #ifndef CPPHTTPLIB_HTTPLIB_H #define CPPHTTPLIB_HTTPLIB_H -#define CPPHTTPLIB_VERSION "0.37.0" -#define CPPHTTPLIB_VERSION_NUM "0x002500" - -/* - * Platform compatibility check - */ - -#if defined(_WIN32) && !defined(_WIN64) -#if defined(_MSC_VER) -#pragma message( \ - "cpp-httplib doesn't support 32-bit Windows. Please use a 64-bit compiler.") -#else -#warning \ - "cpp-httplib doesn't support 32-bit Windows. Please use a 64-bit compiler." -#endif -#elif defined(__SIZEOF_POINTER__) && __SIZEOF_POINTER__ < 8 -#warning \ - "cpp-httplib doesn't support 32-bit platforms. Please use a 64-bit compiler." -#elif defined(__SIZEOF_SIZE_T__) && __SIZEOF_SIZE_T__ < 8 -#warning \ - "cpp-httplib doesn't support platforms where size_t is less than 64 bits." -#endif +#define CPPHTTPLIB_VERSION "0.37.1" +#define CPPHTTPLIB_VERSION_NUM "0x002501" #ifdef _WIN32 #if defined(_WIN32_WINNT) && _WIN32_WINNT < 0x0A00 @@ -2797,7 +2777,7 @@ inline size_t get_header_value_u64(const Headers &headers, std::advance(it, static_cast(id)); if (it != rng.second) { if (is_numeric(it->second)) { - return std::strtoull(it->second.data(), nullptr, 10); + return static_cast(std::strtoull(it->second.data(), nullptr, 10)); } else { is_invalid_value = true; }