diff --git a/common/arg.cpp b/common/arg.cpp index f5ffd8018..0ddb19025 100644 --- a/common/arg.cpp +++ b/common/arg.cpp @@ -2933,7 +2933,7 @@ common_params_context common_params_parser_init(common_params & params, llama_ex ).set_examples({LLAMA_EXAMPLE_TOKENIZE})); add_opt(common_arg( {"--stdin"}, - string_format("read the prompt from stdin (mutually exclusive with -f/--file and -p/--prompt) (default: %s)", params.tokenize_stdin ? "true" : "false"), + string_format("read the prompt from stdin (takes precedence over -f/--file and -p/--prompt) (default: %s)", params.tokenize_stdin ? "true" : "false"), [](common_params & params) { params.tokenize_stdin = true; } diff --git a/common/speculative.cpp b/common/speculative.cpp index f512d059a..f78f5383f 100644 --- a/common/speculative.cpp +++ b/common/speculative.cpp @@ -260,7 +260,10 @@ struct common_speculative_impl_draft_simple : public common_speculative_impl { bool process(const llama_batch & batch) override { auto * ctx_dft = params.ctx_dft; - const int ret = llama_decode(ctx_dft, batch); + llama_batch batch_dft = batch; + batch_dft.logits = nullptr; + + const int ret = llama_decode(ctx_dft, batch_dft); if (ret != 0) { SPC_ERR("failed to decode draft batch, ret = %d\n", ret); diff --git a/ggml/src/ggml-cpu/ggml-cpu.c b/ggml/src/ggml-cpu/ggml-cpu.c index 18e60ba41..4da7cb9a3 100644 --- a/ggml/src/ggml-cpu/ggml-cpu.c +++ b/ggml/src/ggml-cpu/ggml-cpu.c @@ -3695,7 +3695,8 @@ struct ggml_cplan ggml_graph_plan( } break; case GGML_OP_OUT_PROD: { - if (ggml_is_quantized(node->src[0]->type)) { + if (ggml_is_quantized(node->src[0]->type) || + node->src[0]->type == GGML_TYPE_F16) { cur = ggml_type_size(GGML_TYPE_F32) * node->src[0]->ne[0] * n_tasks; } } break; diff --git a/ggml/src/ggml-cpu/ggml-cpu.cpp b/ggml/src/ggml-cpu/ggml-cpu.cpp index f1b3d394c..fcb7db963 100644 --- a/ggml/src/ggml-cpu/ggml-cpu.cpp +++ b/ggml/src/ggml-cpu/ggml-cpu.cpp @@ -462,11 +462,12 @@ static bool ggml_backend_cpu_device_supports_op(ggml_backend_dev_t dev, const st return max_bias == 0.0f; } case GGML_OP_IM2COL_BACK: - return src0->type == GGML_TYPE_F32 && src1->type == GGML_TYPE_F32; + return src0->type == GGML_TYPE_F32 && (src1->type == GGML_TYPE_F32 || src1->type == GGML_TYPE_F16); case GGML_OP_GET_ROWS_BACK: return src0->type == GGML_TYPE_F32 || src0->type == GGML_TYPE_F16; case GGML_OP_OUT_PROD: - return (src0->type == GGML_TYPE_F32 || (ggml_is_quantized(src0->type) && src0->ne[2] == src1->ne[2] && src0->ne[3] == src1->ne[3])) && + return (src0->type == GGML_TYPE_F32 || + ((src0->type == GGML_TYPE_F16 || ggml_is_quantized(src0->type)) && src0->ne[2] == src1->ne[2] && src0->ne[3] == src1->ne[3])) && src1->type == GGML_TYPE_F32 && op->type == GGML_TYPE_F32; default: return true; diff --git a/ggml/src/ggml-cpu/ops.cpp b/ggml/src/ggml-cpu/ops.cpp index c49719374..8022ab49d 100644 --- a/ggml/src/ggml-cpu/ops.cpp +++ b/ggml/src/ggml-cpu/ops.cpp @@ -2081,8 +2081,8 @@ void ggml_compute_forward_concat( const ggml_tensor * src1 = dst->src[1]; if (ggml_is_quantized(src0->type)) { - GGML_ASSERT(ggml_is_contiguous(src0)); - GGML_ASSERT(ggml_is_contiguous(src1)); + GGML_ASSERT(ggml_is_contiguous_rows(src0)); + GGML_ASSERT(ggml_is_contiguous_rows(src1)); GGML_ASSERT(src0->ne[0] % ggml_blck_size(src0->type) == 0); GGML_ASSERT(src1->ne[0] % ggml_blck_size(src1->type) == 0); } @@ -4449,6 +4449,70 @@ static void ggml_compute_forward_out_prod_q_f32( } } +static void ggml_compute_forward_out_prod_f16_f32( + const ggml_compute_params * params, + ggml_tensor * dst) { + + const ggml_tensor * src0 = dst->src[0]; + const ggml_tensor * src1 = dst->src[1]; + + GGML_TENSOR_BINARY_OP_LOCALS; + + const int ith = params->ith; + const int nth = params->nth; + + GGML_ASSERT(src0->type == GGML_TYPE_F16); + GGML_ASSERT(src1->type == GGML_TYPE_F32); + GGML_ASSERT(dst->type == GGML_TYPE_F32); + + GGML_ASSERT(ne02 == ne12); + GGML_ASSERT(ne03 == ne13); + GGML_ASSERT(ne2 == ne12); + GGML_ASSERT(ne3 == ne13); + + GGML_ASSERT(nb00 == sizeof(ggml_fp16_t)); + GGML_ASSERT(nb0 == sizeof(float)); + + GGML_ASSERT(ne0 == ne00); + GGML_ASSERT(ne1 == ne10); + GGML_ASSERT(ne2 == ne02); + GGML_ASSERT(ne3 == ne03); + + if (ith == 0) { + ggml_vec_set_f32(ne0*ne1*ne2*ne3, (float *)dst->data, 0); + } + ggml_barrier(params->threadpool); + + const int64_t nr = ne1*ne2*ne3; + const int64_t dr = (nr + nth - 1)/nth; + const int64_t ir0 = dr*ith; + const int64_t ir1 = MIN(ir0 + dr, nr); + + float * wdata = (float *) params->wdata + (ne0 + CACHE_LINE_SIZE_F32) * ith; + + for (int64_t ir = ir0; ir < ir1; ++ir) { + const int64_t i3 = ir/(ne2*ne1); + const int64_t i2 = (ir - i3*ne2*ne1)/ne1; + const int64_t i1 = (ir - i3*ne2*ne1 - i2*ne1); + + const int64_t i02 = i2; + const int64_t i03 = i3; + + const int64_t i12 = i2; + const int64_t i13 = i3; + + float * d = (float *) ((char *) dst->data + (i1*nb1 + i2*nb2 + i3*nb3)); + + for (int64_t i01 = 0; i01 < ne01; ++i01) { + const int64_t i11 = i01; + ggml_fp16_t * s0 = (ggml_fp16_t *) ((char *) src0->data + (i01*nb01 + i02*nb02 + i03*nb03)); + float * s1 = (float *) ((char *) src1->data + (i1*nb10 + i11*nb11 + i12*nb12 + i13*nb13)); + ggml_fp16_to_fp32_row(s0, wdata, ne0); + ggml_vec_mad_f32(ne0, d, wdata, *s1); + } + } +} + void ggml_compute_forward_out_prod( const ggml_compute_params * params, ggml_tensor * dst) { @@ -4486,9 +4550,8 @@ void ggml_compute_forward_out_prod( } break; case GGML_TYPE_F16: { - GGML_ABORT("fatal error"); // todo - // ggml_compute_forward_out_prod_f16_f32(params, dst); - } + ggml_compute_forward_out_prod_f16_f32(params, dst); + } break; case GGML_TYPE_F32: { ggml_compute_forward_out_prod_f32(params, dst); @@ -6469,7 +6532,7 @@ void ggml_compute_forward_im2col_back_f32( const ggml_tensor * src1 = dst->src[1]; // convolution kernel GGML_ASSERT(src0->type == GGML_TYPE_F32); - GGML_ASSERT(src1->type == GGML_TYPE_F32); + GGML_ASSERT(src1->type == GGML_TYPE_F32 || src1->type == GGML_TYPE_F16); GGML_ASSERT( dst->type == GGML_TYPE_F32); GGML_TENSOR_BINARY_OP_LOCALS; diff --git a/ggml/src/ggml-cuda/common.cuh b/ggml/src/ggml-cuda/common.cuh index 2a5c188fa..3e43dbec6 100644 --- a/ggml/src/ggml-cuda/common.cuh +++ b/ggml/src/ggml-cuda/common.cuh @@ -1122,7 +1122,8 @@ struct ggml_cuda_type_traits { ////////////////////// struct ggml_cuda_device_info { - int device_count; + int device_count; // number of (possibly virtual) devices exposed to the rest of ggml + int physical_device_count; // number of physical CUDA devices actually present struct cuda_device_info { int cc; // compute capability @@ -1135,6 +1136,9 @@ struct ggml_cuda_device_info { size_t total_vram; int warp_size; // Number of threads in a dispatch bool supports_cooperative_launch; // whether cooperative launch is supported + int physical_device; // backing physical CUDA device for this (virtual) device + int physical_share_count; // number of (virtual) devices sharing this device's physical GPU + int virtual_index; // index of this (virtual) device among those sharing its physical GPU }; cuda_device_info devices[GGML_CUDA_MAX_DEVICES] = {}; diff --git a/ggml/src/ggml-cuda/concat.cu b/ggml/src/ggml-cuda/concat.cu index 276ee64e8..6df89013c 100644 --- a/ggml/src/ggml-cuda/concat.cu +++ b/ggml/src/ggml-cuda/concat.cu @@ -141,27 +141,25 @@ static __global__ void __launch_bounds__(CUDA_CONCAT_BLOCK_SIZE) template static void concat_cuda(const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst, int dim, cudaStream_t stream) { - if (ggml_is_contiguous(src0) && ggml_is_contiguous(src1)) { + if (dim != 3 && ggml_is_contiguous_to_3(src0) && ggml_is_contiguous_to_3(src1)) { const T * src0_d = (const T *) src0->data; const T * src1_d = (const T *) src1->data; T * dst_d = (T *) dst->data; - if (dim != 3) { - for (int64_t i3 = 0; i3 < dst->ne[3]; i3++) { - concat_cont_cuda( - src0_d + i3*(src0->nb[3] / sizeof(T)), - src1_d + i3*(src1->nb[3] / sizeof(T)), - dst_d + i3*( dst->nb[3] / sizeof(T)), - ggml_row_size(src0->type, src0->ne[0])/sizeof(T), src0->ne[1], src0->ne[2], - ggml_row_size(dst->type, dst->ne[0])/sizeof(T), dst->ne[1], dst->ne[2], dim, stream); - } - } else { - const size_t size0 = ggml_nbytes(src0); - const size_t size1 = ggml_nbytes(src1); - - CUDA_CHECK(cudaMemcpyAsync((char *) dst->data, src0->data, size0, cudaMemcpyDeviceToDevice, stream)); - CUDA_CHECK(cudaMemcpyAsync((char *) dst->data + size0, src1->data, size1, cudaMemcpyDeviceToDevice, stream)); + for (int64_t i3 = 0; i3 < dst->ne[3]; i3++) { + concat_cont_cuda( + src0_d + i3*(src0->nb[3] / sizeof(T)), + src1_d + i3*(src1->nb[3] / sizeof(T)), + dst_d + i3*( dst->nb[3] / sizeof(T)), + ggml_row_size(src0->type, src0->ne[0])/sizeof(T), src0->ne[1], src0->ne[2], + ggml_row_size(dst->type, dst->ne[0])/sizeof(T), dst->ne[1], dst->ne[2], dim, stream); } + } else if (dim == 3 && ggml_is_contiguous(src0) && ggml_is_contiguous(src1)) { + const size_t size0 = ggml_nbytes(src0); + const size_t size1 = ggml_nbytes(src1); + + CUDA_CHECK(cudaMemcpyAsync((char *) dst->data, src0->data, size0, cudaMemcpyDeviceToDevice, stream)); + CUDA_CHECK(cudaMemcpyAsync((char *) dst->data + size0, src1->data, size1, cudaMemcpyDeviceToDevice, stream)); } else { GGML_ASSERT(!ggml_is_quantized(src0->type)); @@ -208,12 +206,17 @@ void ggml_cuda_op_concat(ggml_backend_cuda_context & ctx, ggml_tensor * dst) { GGML_ASSERT(dst->type == src0->type); if (ggml_is_quantized(src0->type)) { - GGML_ASSERT(ggml_is_contiguous(src0)); - GGML_ASSERT(ggml_is_contiguous(src1)); + if (dim == 3) { + GGML_ASSERT(ggml_is_contiguous(src0)); + GGML_ASSERT(ggml_is_contiguous(src1)); + } else { + GGML_ASSERT(ggml_is_contiguous_to_3(src0)); + GGML_ASSERT(ggml_is_contiguous_to_3(src1)); + } GGML_ASSERT(src0->ne[0] % ggml_blck_size(src0->type) == 0); GGML_ASSERT(src1->ne[0] % ggml_blck_size(src1->type) == 0); - // if tensors are contiguous and ne[0] is multiple of the block size we can concat both tensors as byte tensors + // if first 3 dimensions are contiguous and ne[0] is multiple of the block size we can concat both tensors as byte tensors concat_cuda(src0, src1, dst, dim, stream); } else { GGML_ASSERT(ggml_blck_size(src0->type) == 1); diff --git a/ggml/src/ggml-cuda/ggml-cuda.cu b/ggml/src/ggml-cuda/ggml-cuda.cu index 06c6086de..1931888e6 100644 --- a/ggml/src/ggml-cuda/ggml-cuda.cu +++ b/ggml/src/ggml-cuda/ggml-cuda.cu @@ -67,6 +67,7 @@ bool g_mul_mat_q = true; #include "ggml-cuda/tri.cuh" #include "ggml-cuda/cumsum.cuh" #include "ggml-cuda/fill.cuh" +#include "ggml-cuda/lightning-indexer.cuh" #include "ggml.h" #include @@ -106,17 +107,27 @@ void ggml_cuda_error(const char * stmt, const char * func, const char * file, in GGML_ABORT(GGML_CUDA_NAME " error"); } +// map a (possibly virtual) device id to the physical CUDA device that backs it +static int ggml_cuda_get_physical_device(int device) { + const ggml_cuda_device_info & info = ggml_cuda_info(); + GGML_ASSERT(device >= 0 && device < info.device_count); + return info.devices[device].physical_device; +} + // this is faster on Windows // probably because the Windows CUDA libraries forget to make this check before invoking the drivers void ggml_cuda_set_device(int device) { + // translate the (possibly virtual) device id to the physical CUDA device that backs it + const int physical_device = ggml_cuda_get_physical_device(device); + int current_device; CUDA_CHECK(cudaGetDevice(¤t_device)); - if (device == current_device) { + if (physical_device == current_device) { return; } - CUDA_CHECK(cudaSetDevice(device)); + CUDA_CHECK(cudaSetDevice(physical_device)); } int ggml_cuda_get_device() { @@ -207,56 +218,102 @@ static int ggml_cuda_parse_id(char devName[]) { static ggml_cuda_device_info ggml_cuda_init() { ggml_cuda_device_info info = {}; - cudaError_t err = cudaGetDeviceCount(&info.device_count); + cudaError_t err = cudaGetDeviceCount(&info.physical_device_count); if (err != cudaSuccess) { GGML_LOG_ERROR("%s: failed to initialize " GGML_CUDA_NAME ": %s\n", __func__, cudaGetErrorString(err)); return info; } - GGML_ASSERT(info.device_count <= GGML_CUDA_MAX_DEVICES); + GGML_ASSERT(info.physical_device_count <= GGML_CUDA_MAX_DEVICES); + + // by default expose exactly the physical devices; GGML_CUDA_DEVICES can request a different + // number of (virtual) devices to emulate multi-GPU systems on a machine with fewer GPUs + info.device_count = info.physical_device_count; + + const char * devices_env = getenv("GGML_CUDA_DEVICES"); + if (devices_env != nullptr && info.physical_device_count > 0) { + const int requested = atoi(devices_env); + if (requested > 0) { + info.device_count = requested; + } else { + GGML_LOG_WARN("%s: ignoring invalid GGML_CUDA_DEVICES=\"%s\"\n", __func__, devices_env); + } + } + + if (info.device_count > GGML_CUDA_MAX_DEVICES) { + GGML_LOG_WARN("%s: requested %d devices, clamping to GGML_CUDA_MAX_DEVICES=%d\n", + __func__, info.device_count, GGML_CUDA_MAX_DEVICES); + info.device_count = GGML_CUDA_MAX_DEVICES; + } + + // map each (virtual) device to a backing physical device (round-robin), assign each its index + // among the (virtual) devices sharing that physical GPU, and store the per-physical share count + int physical_share_count[GGML_CUDA_MAX_DEVICES] = {}; + GGML_ASSERT(info.device_count == 0 || info.physical_device_count > 0); + for (int id = 0; id < info.device_count; ++id) { + info.devices[id].physical_device = id % info.physical_device_count; + info.devices[id].virtual_index = physical_share_count[info.devices[id].physical_device]++; + } int64_t total_vram = 0; - for (int id = 0; id < info.device_count; ++id) { + for (int id = 0; id < info.physical_device_count; ++id) { cudaDeviceProp prop; CUDA_CHECK(cudaGetDeviceProperties(&prop, id)); total_vram += prop.totalGlobalMem; } GGML_LOG_INFO("%s: found %d " GGML_CUDA_NAME " devices (Total VRAM: %zu MiB):\n", - __func__, info.device_count, (size_t)(total_vram / (1024 * 1024))); + __func__, info.physical_device_count, (size_t)(total_vram / (1024 * 1024))); + if (info.device_count != info.physical_device_count) { + GGML_LOG_INFO("%s: emulating %d virtual device(s) on %d physical device(s) (GGML_CUDA_DEVICES)\n", + __func__, info.device_count, info.physical_device_count); + } total_vram = 0; std::vector> turing_devices_without_mma; for (int id = 0; id < info.device_count; ++id) { + const int physical_id = info.devices[id].physical_device; + int device_vmm = 0; #if defined(GGML_USE_VMM) CUdevice device; - CU_CHECK(cuDeviceGet(&device, id)); + CU_CHECK(cuDeviceGet(&device, physical_id)); CU_CHECK(cuDeviceGetAttribute(&device_vmm, CU_DEVICE_ATTRIBUTE_VIRTUAL_MEMORY_MANAGEMENT_SUPPORTED, device)); if (device_vmm) { CUmemAllocationProp alloc_prop = {}; alloc_prop.type = CU_MEM_ALLOCATION_TYPE_PINNED; alloc_prop.location.type = CU_MEM_LOCATION_TYPE_DEVICE; - alloc_prop.location.id = id; + alloc_prop.location.id = physical_id; CU_CHECK(cuMemGetAllocationGranularity(&info.devices[id].vmm_granularity, &alloc_prop, CU_MEM_ALLOC_GRANULARITY_RECOMMENDED)); } #endif // defined(GGML_USE_VMM) info.devices[id].vmm = !!device_vmm; cudaDeviceProp prop; - CUDA_CHECK(cudaGetDeviceProperties(&prop, id)); + CUDA_CHECK(cudaGetDeviceProperties(&prop, physical_id)); + + // a virtual device owns only a share of its physical GPU's memory; report that share so the + // logged per-device VRAM sums to the physical total above. + GGML_ASSERT(physical_share_count[physical_id] > 0); + info.devices[id].physical_share_count = physical_share_count[physical_id]; + const size_t device_vram = prop.totalGlobalMem / info.devices[id].physical_share_count; + const size_t device_vram_mib = device_vram / (1024 * 1024); info.default_tensor_split[id] = total_vram; - total_vram += prop.totalGlobalMem; + total_vram += device_vram; +#if defined(GGML_USE_HIP) + info.devices[id].integrated = prop.integrated; +#else info.devices[id].integrated = false; // Temporarily disabled due to issues with corrupted output (e.g. #15034) +#endif info.devices[id].nsm = prop.multiProcessorCount; info.devices[id].smpb = prop.sharedMemPerBlock; info.devices[id].warp_size = prop.warpSize; #ifndef GGML_USE_MUSA int supports_coop_launch = 0; - CUDA_CHECK(cudaDeviceGetAttribute(&supports_coop_launch, cudaDevAttrCooperativeLaunch, id)); + CUDA_CHECK(cudaDeviceGetAttribute(&supports_coop_launch, cudaDevAttrCooperativeLaunch, physical_id)); info.devices[id].supports_cooperative_launch = !!supports_coop_launch; #else info.devices[id].supports_cooperative_launch = false; @@ -279,7 +336,7 @@ static ggml_cuda_device_info ggml_cuda_init() { GGML_LOG_INFO(" Device %d: %s, %s (0x%x), VMM: %s, Wave Size: %d, VRAM: %zu MiB\n", id, prop.name, prop.gcnArchName, info.devices[id].cc & 0xffff, device_vmm ? "yes" : "no", prop.warpSize, - (size_t)(prop.totalGlobalMem / (1024 * 1024))); + device_vram_mib); #elif defined(GGML_USE_MUSA) // FIXME: Ensure compatibility with varying warp sizes across different MUSA archs. info.devices[id].warp_size = 32; @@ -288,13 +345,13 @@ static ggml_cuda_device_info ggml_cuda_init() { info.devices[id].cc += prop.minor * 0x10; GGML_LOG_INFO(" Device %d: %s, compute capability %d.%d, VMM: %s, VRAM: %zu MiB\n", id, prop.name, prop.major, prop.minor, device_vmm ? "yes" : "no", - (size_t)(prop.totalGlobalMem / (1024 * 1024))); + device_vram_mib); #else info.devices[id].smpbo = prop.sharedMemPerBlockOptin; info.devices[id].cc = 100*prop.major + 10*prop.minor; GGML_LOG_INFO(" Device %d: %s, compute capability %d.%d, VMM: %s, VRAM: %zu MiB\n", id, prop.name, prop.major, prop.minor, device_vmm ? "yes" : "no", - (size_t)(prop.totalGlobalMem / (1024 * 1024))); + device_vram_mib); std::string device_name(prop.name); if (device_name == "NVIDIA GeForce MX450") { turing_devices_without_mma.push_back({ id, device_name }); @@ -309,7 +366,7 @@ static ggml_cuda_device_info ggml_cuda_init() { // TODO: Check for future drivers the default scheduling strategy and // remove this call again when cudaDeviceScheduleSpin is default. if (prop.major == 12 && prop.minor == 1) { - CUDA_CHECK(cudaSetDevice(id)); + CUDA_CHECK(cudaSetDevice(physical_id)); CUDA_CHECK(cudaSetDeviceFlags(cudaDeviceScheduleSpin)); } @@ -334,9 +391,9 @@ static ggml_cuda_device_info ggml_cuda_init() { // CUBLAS_CHECK(cublasLoggerConfigure(1, 1, 0, nullptr)); if (getenv("GGML_CUDA_P2P") != nullptr) { - for (int id = 0; id < info.device_count; ++id) { - ggml_cuda_set_device(id); - for (int id_other = 0; id_other < info.device_count; ++id_other) { + for (int id = 0; id < info.physical_device_count; ++id) { + CUDA_CHECK(cudaSetDevice(id)); + for (int id_other = 0; id_other < info.physical_device_count; ++id_other) { if (id == id_other) { continue; } @@ -479,6 +536,7 @@ struct ggml_cuda_pool_vmm : public ggml_cuda_pool { static const size_t CUDA_POOL_VMM_MAX_SIZE = 1ull << 35; // 32 GB int device; + int physical_device; CUdeviceptr pool_addr = 0; size_t pool_used = 0; size_t pool_size = 0; @@ -489,6 +547,7 @@ struct ggml_cuda_pool_vmm : public ggml_cuda_pool { explicit ggml_cuda_pool_vmm(int device) : device(device), + physical_device(ggml_cuda_get_physical_device(device)), granularity(ggml_cuda_info().devices[device].vmm_granularity) { } @@ -524,7 +583,7 @@ struct ggml_cuda_pool_vmm : public ggml_cuda_pool { CUmemAllocationProp prop = {}; prop.type = CU_MEM_ALLOCATION_TYPE_PINNED; prop.location.type = CU_MEM_LOCATION_TYPE_DEVICE; - prop.location.id = device; + prop.location.id = physical_device; CUmemGenericAllocationHandle handle; CU_CHECK(cuMemCreate(&handle, reserve_size, &prop, 0)); @@ -553,20 +612,28 @@ struct ggml_cuda_pool_vmm : public ggml_cuda_pool { // NCCL implicitly enables peer access (cudaDeviceEnablePeerAccess), and // GGML_CUDA_P2P enables it explicitly. Unlike cudaMalloc buffers, VMM // allocations do not become peer-accessible from that alone, so access - // must be granted explicitly here. + // must be granted explicitly here. With virtual devices, grant access + // on the backing *physical* devices (deduplicated, since several + // virtual devices can map to the same physical GPU). std::vector access_descs; + bool physical_seen[GGML_CUDA_MAX_DEVICES] = {}; const int device_count = ggml_cuda_info().device_count; for (int id = 0; id < device_count; ++id) { - if (id != device) { + const int id_physical = ggml_cuda_get_physical_device(id); + if (id_physical != physical_device) { int can_access_peer = 0; - CUDA_CHECK(cudaDeviceCanAccessPeer(&can_access_peer, id, device)); + CUDA_CHECK(cudaDeviceCanAccessPeer(&can_access_peer, id_physical, physical_device)); if (!can_access_peer) { continue; } } + if (physical_seen[id_physical]) { + continue; + } + physical_seen[id_physical] = true; CUmemAccessDesc access = {}; access.location.type = CU_MEM_LOCATION_TYPE_DEVICE; - access.location.id = id; + access.location.id = id_physical; access.flags = CU_MEM_ACCESS_FLAGS_PROT_READWRITE; access_descs.push_back(access); } @@ -575,7 +642,7 @@ struct ggml_cuda_pool_vmm : public ggml_cuda_pool { // set access for non P2P CUmemAccessDesc access = {}; access.location.type = CU_MEM_LOCATION_TYPE_DEVICE; - access.location.id = device; + access.location.id = physical_device; access.flags = CU_MEM_ACCESS_FLAGS_PROT_READWRITE; CU_CHECK(cuMemSetAccess(start_ptr, reserve_size, &access, 1)); } @@ -751,13 +818,17 @@ static bool ggml_backend_cuda_buffer_cpy_tensor(ggml_backend_buffer_t buffer, co if (ggml_backend_buffer_is_cuda(src->buffer)) { ggml_backend_cuda_buffer_context * src_ctx = (ggml_backend_cuda_buffer_context *)src->buffer->context; ggml_backend_cuda_buffer_context * dst_ctx = (ggml_backend_cuda_buffer_context *)dst->buffer->context; - if (src_ctx->device == dst_ctx->device) { + // compare the backing physical devices: distinct virtual devices may share one physical GPU, + // in which case a same-device copy (not a peer copy) is required + const int src_physical = ggml_cuda_get_physical_device(src_ctx->device); + const int dst_physical = ggml_cuda_get_physical_device(dst_ctx->device); + if (src_physical == dst_physical) { CUDA_CHECK(cudaMemcpyAsync(dst->data, src->data, ggml_nbytes(src), cudaMemcpyDeviceToDevice, cudaStreamPerThread)); } else { #ifdef GGML_CUDA_NO_PEER_COPY return false; #else - CUDA_CHECK(cudaMemcpyPeerAsync(dst->data, dst_ctx->device, src->data, src_ctx->device, ggml_nbytes(src), cudaStreamPerThread)); + CUDA_CHECK(cudaMemcpyPeerAsync(dst->data, dst_physical, src->data, src_physical, ggml_nbytes(src), cudaStreamPerThread)); #endif } CUDA_CHECK(cudaStreamSynchronize(cudaStreamPerThread)); @@ -1099,6 +1170,15 @@ static void ggml_backend_cuda_comm_init_internal(ggml_backend_cuda_comm_context static void ggml_backend_cuda_comm_init_nccl(ggml_backend_cuda_comm_context * ret) { #ifdef GGML_USE_NCCL + // Disabling NCCL path when CUDA virtual devices are in use since NCCL requires one distinct physical GPU per rank. + const ggml_cuda_device_info & info = ggml_cuda_info(); + if (info.device_count > info.physical_device_count) { + GGML_LOG_WARN("NCCL disabled: virtual devices in use; " + "falling back to internal AllReduce\n"); + ggml_backend_cuda_comm_init_internal(ret); + return; + } + const size_t n = ret->dev_ids.size(); ret->comms.resize(n); ncclResult_t rc = ncclCommInitAll(ret->comms.data(), (int) n, ret->dev_ids.data()); @@ -2261,6 +2341,9 @@ static bool ggml_cuda_compute_forward(ggml_backend_cuda_context & ctx, struct gg case GGML_OP_FILL: ggml_cuda_op_fill(ctx, dst); break; + case GGML_OP_LIGHTNING_INDEXER: + ggml_cuda_lightning_indexer(ctx, dst); + break; default: return false; } @@ -2359,13 +2442,17 @@ static bool ggml_backend_cuda_cpy_tensor_async(ggml_backend_t backend_src, ggml_ if (backend_src != backend_dst) { // copy on src stream - if (cuda_ctx_src->device == cuda_ctx_dst->device) { + // compare the backing physical devices: distinct virtual devices may share one physical GPU, + // in which case a same-device copy (not a peer copy) is required + const int src_physical = ggml_cuda_get_physical_device(cuda_ctx_src->device); + const int dst_physical = ggml_cuda_get_physical_device(cuda_ctx_dst->device); + if (src_physical == dst_physical) { CUDA_CHECK(cudaMemcpyAsync(dst->data, src->data, ggml_nbytes(dst), cudaMemcpyDeviceToDevice, cuda_ctx_src->stream())); } else { #ifdef GGML_CUDA_NO_PEER_COPY return false; #else - CUDA_CHECK(cudaMemcpyPeerAsync(dst->data, cuda_ctx_dst->device, src->data, cuda_ctx_src->device, ggml_nbytes(dst), cuda_ctx_src->stream())); + CUDA_CHECK(cudaMemcpyPeerAsync(dst->data, dst_physical, src->data, src_physical, ggml_nbytes(dst), cuda_ctx_src->stream())); #endif // GGML_CUDA_NO_PEER_COPY } @@ -3983,7 +4070,7 @@ static bool ggml_cuda_graph_set_enabled(ggml_backend_cuda_context * cuda_ctx, co ggml_cuda_graph * graph = cuda_ctx->cuda_graph(graph_key); if (graph->graph == nullptr) { - if (ggml_cuda_info().devices[cuda_ctx->device].cc < GGML_CUDA_CC_AMPERE) { + if (ggml_cuda_info().devices[cuda_ctx->device].cc < GGML_CUDA_CC_VOLTA) { if (!graph->disable_due_to_gpu_arch) { if(!cugraph_warned) { @@ -4359,16 +4446,38 @@ int ggml_backend_cuda_get_device_count() { return ggml_cuda_info().device_count; } -void ggml_backend_cuda_get_device_description(int device, char * description, size_t description_size) { +static std::string ggml_cuda_device_description(int device) { cudaDeviceProp prop; - CUDA_CHECK(cudaGetDeviceProperties(&prop, device)); - snprintf(description, description_size, "%s", prop.name); + CUDA_CHECK(cudaGetDeviceProperties(&prop, ggml_cuda_get_physical_device(device))); + + const ggml_cuda_device_info & info = ggml_cuda_info(); + std::string description = prop.name; + if (info.device_count > info.physical_device_count) { + description += " (physical device " + std::to_string(info.devices[device].physical_device) + + ", virtual device " + std::to_string(info.devices[device].virtual_index) + ")"; + } + return description; +} + +void ggml_backend_cuda_get_device_description(int device, char * description, size_t description_size) { + snprintf(description, description_size, "%s", ggml_cuda_device_description(device).c_str()); +} + +static int ggml_cuda_physical_device_share_count(int device) { + const ggml_cuda_device_info & info = ggml_cuda_info(); + GGML_ASSERT(device >= 0 && device < info.device_count); + return info.devices[device].physical_share_count; } void ggml_backend_cuda_get_device_memory(int device, size_t * free, size_t * total) { ggml_cuda_set_device(device); CUDA_CHECK(cudaMemGetInfo(free, total)); + + // virtual devices sharing one physical GPU share its memory pool; split it between them + const int share_count = ggml_cuda_physical_device_share_count(device); + *free /= share_count; + *total /= share_count; } bool ggml_backend_cuda_register_host_buffer(void * buffer, size_t size) { @@ -4519,7 +4628,7 @@ static void ggml_backend_cuda_device_get_memory(ggml_backend_dev_t dev, size_t * #if defined(__linux__) // Check if this is a UMA (Unified Memory Architecture) system cudaDeviceProp prop; - CUDA_CHECK(cudaGetDeviceProperties(&prop, ctx->device)); + CUDA_CHECK(cudaGetDeviceProperties(&prop, ggml_cuda_get_physical_device(ctx->device))); // Check if UMA is explicitly enabled via environment variable bool uma_env = getenv("GGML_CUDA_ENABLE_UNIFIED_MEMORY") != nullptr; @@ -4538,13 +4647,17 @@ static void ggml_backend_cuda_device_get_memory(ggml_backend_dev_t dev, size_t * } #endif // defined(__linux__) + // virtual devices sharing one physical GPU share its memory pool; split it between them + const int share_count = ggml_cuda_physical_device_share_count(ctx->device); + *free /= share_count; + *total /= share_count; } static enum ggml_backend_dev_type ggml_backend_cuda_device_get_type(ggml_backend_dev_t dev) { ggml_backend_cuda_device_context * ctx = (ggml_backend_cuda_device_context *) dev->context; cudaDeviceProp prop; - CUDA_CHECK(cudaGetDeviceProperties(&prop, ctx->device)); + CUDA_CHECK(cudaGetDeviceProperties(&prop, ggml_cuda_get_physical_device(ctx->device))); return prop.integrated ? GGML_BACKEND_DEVICE_TYPE_IGPU @@ -4829,13 +4942,23 @@ static bool ggml_backend_cuda_device_supports_op(ggml_backend_dev_t dev, const g { ggml_type src0_type = op->src[0]->type; ggml_type src1_type = op->src[1]->type; + const int32_t dim = op->op_params[0]; return src0_type == src1_type && src0_type == op->type && ( ( ggml_is_quantized(src0_type) && - ggml_is_contiguous(op->src[0]) && - ggml_is_contiguous(op->src[1]) && + ( + ( + dim == 3 && + ggml_is_contiguous(op->src[0]) && + ggml_is_contiguous(op->src[1]) + ) || ( + dim != 3 && + ggml_is_contiguous_to_3(op->src[0]) && + ggml_is_contiguous_to_3(op->src[1]) + ) + ) && op->src[0]->ne[0] % ggml_blck_size(src0_type) == 0 && op->src[1]->ne[0] % ggml_blck_size(src0_type) == 0 ) || ( @@ -4990,6 +5113,8 @@ static bool ggml_backend_cuda_device_supports_op(ggml_backend_dev_t dev, const g case GGML_OP_DIAG: case GGML_OP_SOLVE_TRI: return true; + case GGML_OP_LIGHTNING_INDEXER: + return ggml_cuda_lightning_indexer_supported(dev_ctx->device, op); default: return false; @@ -5193,18 +5318,24 @@ ggml_backend_reg_t ggml_backend_cuda_reg() { ggml_backend_cuda_reg_context * ctx = new ggml_backend_cuda_reg_context; const int min_batch_size = getenv("GGML_OP_OFFLOAD_MIN_BATCH") ? atoi(getenv("GGML_OP_OFFLOAD_MIN_BATCH")) : 32; - for (int i = 0; i < ggml_cuda_info().device_count; i++) { + const ggml_cuda_device_info & info = ggml_cuda_info(); + const bool virtual_devices = info.device_count > info.physical_device_count; + + for (int i = 0; i < info.device_count; i++) { + const int physical_id = info.devices[i].physical_device; + ggml_backend_cuda_device_context * dev_ctx = new ggml_backend_cuda_device_context; dev_ctx->device = i; dev_ctx->name = GGML_CUDA_NAME + std::to_string(i); - - cudaDeviceProp prop; - CUDA_CHECK(cudaGetDeviceProperties(&prop, i)); - dev_ctx->description = prop.name; + dev_ctx->description = ggml_cuda_device_description(i); char pci_bus_id[32] = {}; - CUDA_CHECK(cudaDeviceGetPCIBusId(pci_bus_id, sizeof(pci_bus_id), i)); + CUDA_CHECK(cudaDeviceGetPCIBusId(pci_bus_id, sizeof(pci_bus_id), physical_id)); dev_ctx->pci_bus_id = pci_bus_id; + if (virtual_devices) { + // make the pci bus id unique for virtual devices + dev_ctx->pci_bus_id += "-v" + std::to_string(i); + } for (char & c : dev_ctx->pci_bus_id) { c = std::tolower(c); } diff --git a/ggml/src/ggml-cuda/lightning-indexer.cu b/ggml/src/ggml-cuda/lightning-indexer.cu new file mode 100644 index 000000000..5edc967e0 --- /dev/null +++ b/ggml/src/ggml-cuda/lightning-indexer.cu @@ -0,0 +1,588 @@ +#include "common.cuh" +#include "lightning-indexer.cuh" +#include "fattn-common.cuh" +#include "convert.cuh" + +#if !defined(GGML_USE_HIP) && !defined(GGML_USE_MUSA) +#if defined(TURING_MMA_AVAILABLE) + +typedef union { + int2 i2; + half2 h2[2]; +} half4; + +// TODO add support for AMD cards via rocWMMA +#include +namespace wmma = nvcuda::wmma; + +template +static __global__ void lightning_indexer_kernel_wmma( + const float * Q, const char * K, const float * W, const half * M, float * dst, + int64_t n_stream, int64_t n_batch, int64_t n_kv, + size_t nb1, size_t nb2, size_t nb3, + size_t nbq1, size_t nbq2, size_t nbq3, + size_t nbk1, size_t nbk2, size_t nbk3, + size_t nbw1, size_t nbw2, size_t nbw3, + size_t nbm1, size_t nbm2, size_t nbm3, + int64_t nem3 + ) { + + constexpr int THREADS_PER_BLOCK = WARPS_PER_BLOCK * WARP_SIZE; + constexpr int HEADS_PER_INNER_LOOP = 8; + constexpr int K_EMBD_PER_INNER_LOOP = 16; + constexpr int N_EMBD_PADDED = N_EMBD + 8; + + const int i_batch = blockIdx.y; + const int i_stream = blockIdx.z; + const int i_warp = threadIdx.y; + const int i_lane = threadIdx.x; + const int tid = i_warp * WARP_SIZE + i_lane; + + // each block processes K_VECS_PER_BLOCK K vectors + const int start_kv = blockIdx.x * K_VECS_PER_BLOCK; + + const char * q_base = (const char *) Q + i_batch*nbq2 + i_stream*nbq3; + const float * w_base = (const float *) ((const char *) W + i_batch*nbw1 + i_stream*nbw3); + + // phase 1 - load weights and first Q tile to shared memory + + __shared__ float w_shared[N_HEAD]; + __shared__ int2 q_shared_h[HEADS_PER_INNER_LOOP][N_EMBD_PADDED / 4]; + + if (tid < N_HEAD) { + w_shared[tid] = w_base[tid]; + } + + // total number of half4 elements in HEADS_PER_INNER_LOOP x N_EMBD Q tile + constexpr int N_Q_TILE = HEADS_PER_INNER_LOOP * (N_EMBD / 4); + // number of registers needed in each thread to store Q tile in thread block + constexpr int N_Q_NEXT = (N_Q_TILE + THREADS_PER_BLOCK - 1) / THREADS_PER_BLOCK; + +#pragma unroll + for (int i_q = tid; i_q < N_Q_TILE; i_q += THREADS_PER_BLOCK) { + const int i_head = i_q / (N_EMBD / 4); + const int i_embd = i_q % (N_EMBD / 4); + const float4 q = *(const float4 *) (q_base + i_head*nbq1 + i_embd*sizeof(float4)); + half4 q_packed; + q_packed.h2[0] = __float22half2_rn(make_float2(q.x, q.y)); + q_packed.h2[1] = __float22half2_rn(make_float2(q.z, q.w)); + q_shared_h[i_head][i_embd] = q_packed.i2; + } + + // phase 2 - load (and dequantize if needed) K to shared mem + + __shared__ half2 k_shared_h[K_VECS_PER_BLOCK][N_EMBD_PADDED / 4][2]; + + constexpr int n_k = K_VECS_PER_BLOCK * (N_EMBD / 4); + + if constexpr (TYPE_K == GGML_TYPE_F16) { +#pragma unroll + for (int i_k = tid; i_k < n_k; i_k += THREADS_PER_BLOCK) { + const int i_k_vec = i_k / (N_EMBD / 4); + const int i_embd = i_k % (N_EMBD / 4); + const int i_kv = start_kv + i_k_vec; + if (i_kv < n_kv) { + const int2 * k_base = (const int2 *) ((const char *) K + i_kv*nbk2 + i_stream*nbk3); + *(int2*) &k_shared_h[i_k_vec][i_embd] = k_base[i_embd]; + } else { + *(int2*) &k_shared_h[i_k_vec][i_embd] = make_int2(0, 0); + } + } + } else { + constexpr dequantize_V_t dequantize_k = get_dequantize_V(); +#pragma unroll + for (int i_k = tid; i_k < n_k; i_k += THREADS_PER_BLOCK) { + const int i_k_vec = i_k / (N_EMBD / 4); + const int i_embd = i_k % (N_EMBD / 4); + const int i_kv = start_kv + i_k_vec; + if (i_kv < n_kv) { + const void * k_base = (const void *) ((const char *) K + i_kv*nbk2 + i_stream*nbk3); + dequantize_k(k_base, &k_shared_h[i_k_vec][i_embd][0], i_embd * 4); + } else { + *(int2*) &k_shared_h[i_k_vec][i_embd] = make_int2(0, 0); + } + } + } + + __syncthreads(); + + // phase 3 - calculate lightning indexer scores + + __shared__ float qk_shared[WARPS_PER_BLOCK][HEADS_PER_INNER_LOOP][K_VECS_PER_BLOCK]; + + // load K fragment + wmma::fragment frag_k; + wmma::load_matrix_sync(frag_k, (half*) &k_shared_h[0][i_warp * K_EMBD_PER_INNER_LOOP / 4], N_EMBD_PADDED); + + float score_k = 0.0f; + + for (int i_head_0 = 0; i_head_0 < N_HEAD; i_head_0 += HEADS_PER_INNER_LOOP) { + const int i_head_next = i_head_0 + HEADS_PER_INNER_LOOP; + + // we don't use accumulator for anything, fill it with zeros + wmma::fragment frag_acc; + wmma::fill_fragment(frag_acc, 0.0f); + + // load Q fragment + wmma::fragment frag_q; + wmma::load_matrix_sync(frag_q, (half*) &q_shared_h[0][i_warp * K_EMBD_PER_INNER_LOOP / 4], N_EMBD_PADDED); + + // preload next Q tile to registers during matrix multiplication + float4 q_next[N_Q_NEXT]; + + if (i_head_next < N_HEAD) { +#pragma unroll + for (int i_q = tid, i_q_next = 0; i_q < N_Q_TILE; i_q += THREADS_PER_BLOCK) { + const int i_head = i_head_next + i_q / (N_EMBD / 4); + const int i_embd = i_q % (N_EMBD / 4); + q_next[i_q_next++] = *(const float4 *) (q_base + i_head*nbq1 + i_embd*sizeof(float4)); + } + } + + // perform matrix multiplication + wmma::mma_sync(frag_acc, frag_q, frag_k, frag_acc); + wmma::store_matrix_sync((float*) &qk_shared[i_warp][0][0], frag_acc, K_VECS_PER_BLOCK, wmma::mem_row_major); + + // make sure all threads finished using q_shared_h so we can store next tile + __syncthreads(); + + // write preloaded Q tile to shared memory + if (i_head_next < N_HEAD) { +#pragma unroll + for (int i_q = tid, i_q_next = 0; i_q < N_Q_TILE; i_q += THREADS_PER_BLOCK) { + const int i_head = i_q / (N_EMBD / 4); + const int i_embd = i_q % (N_EMBD / 4); + half4 q_packed; + q_packed.h2[0] = __float22half2_rn(make_float2(q_next[i_q_next].x, q_next[i_q_next].y)); + q_packed.h2[1] = __float22half2_rn(make_float2(q_next[i_q_next].z, q_next[i_q_next].w)); + q_shared_h[i_head][i_embd] = q_packed.i2; + ++i_q_next; + } + } + + // accumulate QK multiplication results from all block warps + // (there are 256 threads in block and 256 matmul outputs) + // TODO it will break if WARP_SIZE is not 32 + const int h = tid / K_VECS_PER_BLOCK; + const int k = tid % K_VECS_PER_BLOCK; + const float w_val = w_shared[i_head_0 + h]; + + float sum = 0.0f; +#pragma unroll + for (int w = 0; w < WARPS_PER_BLOCK; ++w) { + sum += qk_shared[w][h][k]; + } + + // ReLU, weight + sum = sum > 0.0f ? sum : 0.0f; + sum *= w_val; + + // wait until qk_shared[0] is no longer used + __syncthreads(); + + // reuse qk_shared[0] for storing partial results + qk_shared[0][h][k] = sum; + + // wait until all threads write their results + __syncthreads(); + + // accumulate result over heads + if (tid < K_VECS_PER_BLOCK) { +#pragma unroll + for (int i_head = 0; i_head < HEADS_PER_INNER_LOOP; ++i_head) { + score_k += qk_shared[0][i_head][tid]; + } + } + + // make sure all threads finished using qk_shared + __syncthreads(); + } + + // phase 4 - store output to VRAM + + if (tid < K_VECS_PER_BLOCK) { + const int i_kv = start_kv + tid; + if (i_kv < n_kv) { + const half * m_base = (const half *) ((const char *) M + i_batch*nbm1 + (i_stream%nem3)*nbm3); + float * dst_base = (float *) ((char *) dst + i_batch*nb1 + i_stream*nb3); + dst_base[i_kv] = score_k + __half2float(m_base[i_kv]); + } + } +} + +#else // defined(TURING_MMA_AVAILABLE) + +template +static __global__ void lightning_indexer_kernel_wmma( + const float * Q, const char * K, const float * W, const half * M, float * dst, + int64_t n_stream, int64_t n_batch, int64_t n_kv, + size_t nb1, size_t nb2, size_t nb3, + size_t nbq1, size_t nbq2, size_t nbq3, + size_t nbk1, size_t nbk2, size_t nbk3, + size_t nbw1, size_t nbw2, size_t nbw3, + size_t nbm1, size_t nbm2, size_t nbm3, + int64_t nem3 + ) { + GGML_UNUSED_VARS(Q, K, W, M, dst, + n_stream, n_batch, n_kv, + nb1, nb2, nb3, + nbq1, nbq2, nbq3, + nbk1, nbk2, nbk3, + nbw1, nbw2, nbw3, + nem3); + NO_DEVICE_CODE; +} + +#endif // defined(TURING_MMA_AVAILABLE) +#endif // !defined(GGML_USE_HIP) && !defined(GGML_USE_MUSA) + +// TODO there is one ugly assumption used in this kernel - that WARP_SIZE is equal to 32 +// thanks to that one warp operating on float4 processes whole indexer K/Q vectors +// 32 * 4 = 128 (N_EMBD) + +template +static __global__ void lightning_indexer_kernel_vec( + const float * Q, const char * K, const float * W, const half * M, float * dst, + int64_t n_stream, int64_t n_batch, int64_t n_kv, + size_t nb1, size_t nb2, size_t nb3, + size_t nbq1, size_t nbq2, size_t nbq3, + size_t nbk1, size_t nbk2, size_t nbk3, + size_t nbw1, size_t nbw2, size_t nbw3, + size_t nbm1, size_t nbm2, size_t nbm3, + int64_t nem3 + ) { + + constexpr int K_VECS_PER_WARP = K_VECS_PER_BLOCK / WARPS_PER_BLOCK; + constexpr int THREADS_PER_BLOCK = WARPS_PER_BLOCK * WARP_SIZE; + + const int i_batch = blockIdx.y; + const int i_stream = blockIdx.z; + const int i_warp = threadIdx.y; + const int i_lane = threadIdx.x; + const int tid = i_warp * WARP_SIZE + i_lane; + + // each warp processes K_VECS_PER_WARP K vectors + const int start_kv_block = blockIdx.x * K_VECS_PER_BLOCK; + const int start_kv = start_kv_block + i_warp * K_VECS_PER_WARP; + + const char * q_base = (const char *) Q + i_batch*nbq2 + i_stream*nbq3; + const float * w_base = (const float *) ((const char *) W + i_batch*nbw1 + i_stream*nbw3); + + // phase 1 - load (and dequantize if needed) K to registers + + float4 k_reg_f[K_VECS_PER_WARP]; + + if constexpr (TYPE_K == GGML_TYPE_F32) { + // direct copy of float4 +#pragma unroll + for (int k = 0; k < K_VECS_PER_WARP; ++k) { + int i_kv = start_kv + k; + if (i_kv < n_kv) { + const float4 * k_base = (const float4 *) ((const char *) K + i_kv*nbk2 + i_stream*nbk3); + k_reg_f[k] = k_base[i_lane]; + } else { + k_reg_f[k] = make_float4(0, 0, 0, 0); + } + } + } else { + // dequantize remaining types to float + constexpr dequantize_V_t dequantize_k = get_dequantize_V(); +#pragma unroll + for (int k = 0; k < K_VECS_PER_WARP; ++k) { + int i_kv = start_kv + k; + if (i_kv < n_kv) { + const void * k_base = (const void *) ((const char *) K + i_kv*nbk2 + i_stream*nbk3); + dequantize_k(k_base, &k_reg_f[k], i_lane * 4); + } else { + k_reg_f[k] = make_float4(0, 0, 0, 0); + } + } + } + + float score_k[K_VECS_PER_WARP] = { 0.0f }; + + // load weights and Q only for N_HEAD_INNER heads at once to reduce shared memory usage + constexpr int N_HEAD_INNER = N_HEAD / 4; + + for (int i_head_0 = 0; i_head_0 < N_HEAD; i_head_0 += N_HEAD_INNER) { + // phase 2 - load weights and Q to shared memory + + __shared__ float w_shared[N_HEAD_INNER]; + __shared__ float4 q_shared_f[N_HEAD_INNER][N_EMBD / 4]; + + if (tid < N_HEAD_INNER) { + w_shared[tid] = w_base[i_head_0 + tid]; + } + + constexpr int n_q = N_HEAD_INNER * (N_EMBD / 4); +#pragma unroll + for (int i_q = tid; i_q < n_q; i_q += THREADS_PER_BLOCK) { + const int i_head_inner = i_q / (N_EMBD / 4); + const int i_head = i_head_0 + i_head_inner; + const int i_embd = i_q % (N_EMBD / 4); + q_shared_f[i_head_inner][i_embd] = *(const float4 *) (q_base + i_head*nbq1 + i_embd*sizeof(float4)); + } + + __syncthreads(); + + // phase 3 - calculate lightning indexer scores + + for (int i_head_inner = 0; i_head_inner < N_HEAD_INNER; ++i_head_inner) { + const float w_val = w_shared[i_head_inner]; + float qk[K_VECS_PER_WARP] = { 0.0f }; + + // dot product of floats + const float4 q_vec = q_shared_f[i_head_inner][i_lane]; + +#pragma unroll + for (int k = 0; k < K_VECS_PER_WARP; ++k) { + ggml_cuda_mad(qk[k], q_vec.x, k_reg_f[k].x); + ggml_cuda_mad(qk[k], q_vec.y, k_reg_f[k].y); + ggml_cuda_mad(qk[k], q_vec.z, k_reg_f[k].z); + ggml_cuda_mad(qk[k], q_vec.w, k_reg_f[k].w); + } + +#pragma unroll + for (int k = 0; k < K_VECS_PER_WARP; ++k) { + float sum = warp_reduce_sum(qk[k]); + + // ReLU, weight + if (i_lane == 0) { + sum = (sum > 0.0f) ? sum : 0.0f; + score_k[k] += sum * w_val; + } + } + } + + __syncthreads(); + } + + // phase 4 - store outputs to shared memory + + __shared__ float dst_shared[K_VECS_PER_BLOCK]; + + if (i_lane == 0) { +#pragma unroll + for (int k = 0; k < K_VECS_PER_WARP; ++k) { + dst_shared[i_warp * K_VECS_PER_WARP + k] = score_k[k]; + } + } + + __syncthreads(); + + // phase 5 - write from shared memory to VRAM in coalesced manner + + if (tid < K_VECS_PER_BLOCK) { + int i_kv = start_kv_block + tid; + if (i_kv < n_kv) { + const half * m_base = (const half *) ((const char *) M + i_batch*nbm1 + (i_stream%nem3)*nbm3); + float * dst_base = (float *) ((char *) dst + i_batch*nb1 + i_stream*nb3); + dst_base[i_kv] = dst_shared[tid] + __half2float(m_base[i_kv]); + } + } +} + +#define LIGHTNING_INDEXER_CASE(lightning_indexer_kernel, n_embd, n_head, K, type_K) \ + if (K->type == (type_K)) { \ + lightning_indexer_kernel \ + <<>>( \ + q_d, k_d, w_d, m_d, dst_d, \ + n_stream, n_batch, n_kv, \ + nb1, nb2, nb3, \ + nbq1, nbq2, nbq3, \ + nbk1, nbk2, nbk3, \ + nbw1, nbw2, nbw3, \ + nbm1, nbm2, nbm3, \ + nem3 \ + ); \ + } else + +void ggml_cuda_lightning_indexer(ggml_backend_cuda_context & ctx, ggml_tensor * dst) { + const ggml_tensor * q = dst->src[0]; + const ggml_tensor * k = dst->src[1]; + const ggml_tensor * w = dst->src[2]; // weights + const ggml_tensor * m = dst->src[3]; // mask + + GGML_ASSERT(dst->type == GGML_TYPE_F32); + GGML_ASSERT( q->type == GGML_TYPE_F32); + GGML_ASSERT( w->type == GGML_TYPE_F32); + GGML_ASSERT( m->type == GGML_TYPE_F16); + + GGML_TENSOR_LOCALS(int64_t, neq, q, ne) + GGML_TENSOR_LOCALS(size_t, nbq, q, nb) + GGML_TENSOR_LOCALS(int64_t, nek, k, ne) + GGML_TENSOR_LOCALS(size_t, nbk, k, nb) + GGML_TENSOR_LOCALS(int64_t, new, w, ne) + GGML_TENSOR_LOCALS(size_t, nbw, w, nb) + GGML_TENSOR_LOCALS(int64_t, nem, m, ne) + GGML_TENSOR_LOCALS(size_t, nbm, m, nb) + GGML_TENSOR_LOCALS(int64_t, ne, dst, ne) + GGML_TENSOR_LOCALS(size_t, nb, dst, nb) + + // input tensor rows must be contiguous + GGML_ASSERT(nbq0 == ggml_type_size(q->type)); + GGML_ASSERT(nbk0 == ggml_type_size(k->type)); + GGML_ASSERT(nbw0 == ggml_type_size(w->type)); + GGML_ASSERT(nbm0 == ggml_type_size(m->type)); + + // dst cannot be transposed or permuted + GGML_ASSERT(nb0 == sizeof(float)); + GGML_ASSERT(nb0 <= nb1); + GGML_ASSERT(nb1 <= nb2); + GGML_ASSERT(nb2 <= nb3); + + const int n_embd = q->ne[0]; + const int n_head = q->ne[1]; + const int n_batch = q->ne[2]; + const int n_stream = q->ne[3]; + const int n_kv = k->ne[2]; + + const float * q_d = (const float *) q->data; + const char * k_d = (const char *) k->data; + const float * w_d = (const float *) w->data; + const half * m_d = (const half *) m->data; + float * dst_d = ( float *) dst->data; + + const int device = ggml_cuda_get_device(); + const int cc = ggml_cuda_info().devices[device].cc; + + if (n_embd == 128 && n_head == 64) { +#if !defined(GGML_USE_HIP) && !defined(GGML_USE_MUSA) + if (GGML_CUDA_CC_IS_NVIDIA(cc) && turing_mma_available(cc) && k->type != GGML_TYPE_F32 && k->type != GGML_TYPE_BF16) { + // use wmma kernel + constexpr int K_VECS_PER_BLOCK = 32; + constexpr int WARPS_PER_BLOCK = 8; + + dim3 block(32, WARPS_PER_BLOCK); + int num_kv_blocks = (n_kv + (K_VECS_PER_BLOCK) - 1) / (K_VECS_PER_BLOCK); + dim3 grid(num_kv_blocks, n_batch, n_stream); + + LIGHTNING_INDEXER_CASE(lightning_indexer_kernel_wmma, 128, 64, k, GGML_TYPE_F16) + LIGHTNING_INDEXER_CASE(lightning_indexer_kernel_wmma, 128, 64, k, GGML_TYPE_Q4_0) + LIGHTNING_INDEXER_CASE(lightning_indexer_kernel_wmma, 128, 64, k, GGML_TYPE_Q4_1) + LIGHTNING_INDEXER_CASE(lightning_indexer_kernel_wmma, 128, 64, k, GGML_TYPE_Q5_0) + LIGHTNING_INDEXER_CASE(lightning_indexer_kernel_wmma, 128, 64, k, GGML_TYPE_Q5_1) + LIGHTNING_INDEXER_CASE(lightning_indexer_kernel_wmma, 128, 64, k, GGML_TYPE_Q8_0) + GGML_ABORT("fatal error"); + } else { +#else // !defined(GGML_USE_HIP) && !defined(GGML_USE_MUSA) + { +#endif // !defined(GGML_USE_HIP) && !defined(GGML_USE_MUSA) + // use vector kernel + constexpr int K_VECS_PER_WARP = 8; + constexpr int WARPS_PER_BLOCK = 8; + constexpr int K_VECS_PER_BLOCK = K_VECS_PER_WARP * WARPS_PER_BLOCK; + + dim3 block(32, WARPS_PER_BLOCK); + int num_kv_blocks = (n_kv + (K_VECS_PER_BLOCK) - 1) / (K_VECS_PER_BLOCK); + dim3 grid(num_kv_blocks, n_batch, n_stream); + + LIGHTNING_INDEXER_CASE(lightning_indexer_kernel_vec, 128, 64, k, GGML_TYPE_F16) + LIGHTNING_INDEXER_CASE(lightning_indexer_kernel_vec, 128, 64, k, GGML_TYPE_Q4_0) + LIGHTNING_INDEXER_CASE(lightning_indexer_kernel_vec, 128, 64, k, GGML_TYPE_Q4_1) + LIGHTNING_INDEXER_CASE(lightning_indexer_kernel_vec, 128, 64, k, GGML_TYPE_Q5_0) + LIGHTNING_INDEXER_CASE(lightning_indexer_kernel_vec, 128, 64, k, GGML_TYPE_Q5_1) + LIGHTNING_INDEXER_CASE(lightning_indexer_kernel_vec, 128, 64, k, GGML_TYPE_Q8_0) + LIGHTNING_INDEXER_CASE(lightning_indexer_kernel_vec, 128, 64, k, GGML_TYPE_BF16) + LIGHTNING_INDEXER_CASE(lightning_indexer_kernel_vec, 128, 64, k, GGML_TYPE_F32) + GGML_ABORT("fatal error"); + } + } else if (n_embd == 128 && n_head == 32) { +#if !defined(GGML_USE_HIP) && !defined(GGML_USE_MUSA) + if (GGML_CUDA_CC_IS_NVIDIA(cc) && turing_mma_available(cc) && k->type != GGML_TYPE_F32 && k->type != GGML_TYPE_BF16) { + // use wmma kernel + constexpr int K_VECS_PER_BLOCK = 32; + constexpr int WARPS_PER_BLOCK = 8; + + dim3 block(32, WARPS_PER_BLOCK); + int num_kv_blocks = (n_kv + (K_VECS_PER_BLOCK) - 1) / (K_VECS_PER_BLOCK); + dim3 grid(num_kv_blocks, n_batch, n_stream); + + LIGHTNING_INDEXER_CASE(lightning_indexer_kernel_wmma, 128, 32, k, GGML_TYPE_F16) + LIGHTNING_INDEXER_CASE(lightning_indexer_kernel_wmma, 128, 32, k, GGML_TYPE_Q4_0) + LIGHTNING_INDEXER_CASE(lightning_indexer_kernel_wmma, 128, 32, k, GGML_TYPE_Q4_1) + LIGHTNING_INDEXER_CASE(lightning_indexer_kernel_wmma, 128, 32, k, GGML_TYPE_Q5_0) + LIGHTNING_INDEXER_CASE(lightning_indexer_kernel_wmma, 128, 32, k, GGML_TYPE_Q5_1) + LIGHTNING_INDEXER_CASE(lightning_indexer_kernel_wmma, 128, 32, k, GGML_TYPE_Q8_0) + GGML_ABORT("fatal error"); + } else { +#else // !defined(GGML_USE_HIP) && !defined(GGML_USE_MUSA) + { +#endif // !defined(GGML_USE_HIP) && !defined(GGML_USE_MUSA) + // use vector kernel + constexpr int K_VECS_PER_WARP = 8; + constexpr int WARPS_PER_BLOCK = 8; + constexpr int K_VECS_PER_BLOCK = K_VECS_PER_WARP * WARPS_PER_BLOCK; + + dim3 block(32, WARPS_PER_BLOCK); + int num_kv_blocks = (n_kv + (K_VECS_PER_BLOCK) - 1) / (K_VECS_PER_BLOCK); + dim3 grid(num_kv_blocks, n_batch, n_stream); + + LIGHTNING_INDEXER_CASE(lightning_indexer_kernel_vec, 128, 32, k, GGML_TYPE_F16) + LIGHTNING_INDEXER_CASE(lightning_indexer_kernel_vec, 128, 32, k, GGML_TYPE_Q4_0) + LIGHTNING_INDEXER_CASE(lightning_indexer_kernel_vec, 128, 32, k, GGML_TYPE_Q4_1) + LIGHTNING_INDEXER_CASE(lightning_indexer_kernel_vec, 128, 32, k, GGML_TYPE_Q5_0) + LIGHTNING_INDEXER_CASE(lightning_indexer_kernel_vec, 128, 32, k, GGML_TYPE_Q5_1) + LIGHTNING_INDEXER_CASE(lightning_indexer_kernel_vec, 128, 32, k, GGML_TYPE_Q8_0) + LIGHTNING_INDEXER_CASE(lightning_indexer_kernel_vec, 128, 32, k, GGML_TYPE_BF16) + LIGHTNING_INDEXER_CASE(lightning_indexer_kernel_vec, 128, 32, k, GGML_TYPE_F32) + GGML_ABORT("fatal error"); + } + } else { + GGML_ABORT("fatal error"); + } +} + +bool ggml_cuda_lightning_indexer_supported(int device, const ggml_tensor * dst) { + GGML_UNUSED(device); + + const ggml_tensor * q = dst->src[0]; + const ggml_tensor * k = dst->src[1]; + const ggml_tensor * w = dst->src[2]; // weights + const ggml_tensor * m = dst->src[3]; // mask + + GGML_TENSOR_LOCALS(int64_t, neq, q, ne) + GGML_TENSOR_LOCALS(size_t, nbq, q, nb) + GGML_TENSOR_LOCALS(int64_t, nek, k, ne) + GGML_TENSOR_LOCALS(size_t, nbk, k, nb) + GGML_TENSOR_LOCALS(int64_t, new, w, ne) + GGML_TENSOR_LOCALS(size_t, nbw, w, nb) + GGML_TENSOR_LOCALS(int64_t, nem, m, ne) + GGML_TENSOR_LOCALS(size_t, nbm, m, nb) + GGML_TENSOR_LOCALS(int64_t, ne, dst, ne) + GGML_TENSOR_LOCALS(size_t, nb, dst, nb) + + if (neq0 != 128) { + return false; + } + + if (neq1 != 64 && neq1 != 32) { + return false; + } + + // alignment checks + for (const ggml_tensor * t : {q, k}) { + if (ggml_is_quantized(t->type)) { + continue; + } + for (size_t i = 1; i < GGML_MAX_DIMS; ++i) { + if (t->nb[i] % 16 != 0) { + return false; + } + } + } + + switch(k->type) { + case GGML_TYPE_F32: + case GGML_TYPE_BF16: + case GGML_TYPE_F16: + case GGML_TYPE_Q8_0: + case GGML_TYPE_Q5_1: + case GGML_TYPE_Q5_0: + case GGML_TYPE_Q4_1: + case GGML_TYPE_Q4_0: + return true; + default: + return false; + } +} diff --git a/ggml/src/ggml-cuda/lightning-indexer.cuh b/ggml/src/ggml-cuda/lightning-indexer.cuh new file mode 100644 index 000000000..f2fc95181 --- /dev/null +++ b/ggml/src/ggml-cuda/lightning-indexer.cuh @@ -0,0 +1,4 @@ +#include "common.cuh" + +void ggml_cuda_lightning_indexer(ggml_backend_cuda_context & ctx, ggml_tensor * dst); +bool ggml_cuda_lightning_indexer_supported(int device, const ggml_tensor * dst); diff --git a/ggml/src/ggml-cuda/mmf.cu b/ggml/src/ggml-cuda/mmf.cu index aad4c34aa..646a5899c 100644 --- a/ggml/src/ggml-cuda/mmf.cu +++ b/ggml/src/ggml-cuda/mmf.cu @@ -85,7 +85,7 @@ void ggml_cuda_mul_mat_f(ggml_backend_cuda_context & ctx, const ggml_tensor * sr GGML_ASSERT(sis1 > 0); ggml_cuda_launch_mm_ids_helper(ids_d, ids_src_compact_dev.get(), ids_dst_compact_dev.get(), expert_bounds_dev.get(), - static_cast(n_experts), static_cast(n_tokens), static_cast(n_expert_used), static_cast(ne11), si1, sis1, ctx.stream()); + static_cast(n_experts), static_cast(n_tokens), static_cast(n_expert_used), static_cast(ne11), si1, sis1, /*write_inverse =*/ false, ctx.stream()); CUDA_CHECK(cudaGetLastError()); ids_info.ids_src_compact = ids_src_compact_dev.get(); diff --git a/ggml/src/ggml-cuda/mmid.cu b/ggml/src/ggml-cuda/mmid.cu index 3c61e4595..f80442fbe 100644 --- a/ggml/src/ggml-cuda/mmid.cu +++ b/ggml/src/ggml-cuda/mmid.cu @@ -27,7 +27,7 @@ template __launch_bounds__(ggml_cuda_get_physical_warp_size(), 1) static __global__ void mm_ids_helper( const int32_t * __restrict__ ids, int32_t * __restrict__ ids_src1, int32_t * __restrict__ ids_dst, int32_t * __restrict__ expert_bounds, - const int n_tokens, const int n_expert_used_var, const int nchannels_y, const int si1, const int sis1) { + const int n_tokens, const int n_expert_used_var, const int nchannels_y, const int si1, const int sis1, const bool write_inverse) { constexpr int warp_size = ggml_cuda_get_physical_warp_size(); const int n_expert_used = n_expert_used_template == 0 ? n_expert_used_var : n_expert_used_template; const int expert = blockIdx.x; @@ -98,8 +98,13 @@ static __global__ void mm_ids_helper( const mm_ids_helper_store store_it = store[itc]; const int it = store_it.it(); const int iex_used = store_it.iex_used(); - ids_src1[nex_prev + itc] = it*sis1 + iex_used % nchannels_y; - ids_dst [nex_prev + itc] = it*n_expert_used + iex_used; + ids_dst[nex_prev + itc] = it*n_expert_used + iex_used; + // ids_src1 holds the forward map, or the inverse map (token slot -> compact row) for quant dedup + if (write_inverse) { + ids_src1[it*n_expert_used + iex_used] = nex_prev + itc; + } else { + ids_src1[nex_prev + itc] = it*sis1 + iex_used % nchannels_y; + } } if (threadIdx.x != 0) { @@ -118,7 +123,7 @@ static __global__ void mm_ids_helper( template static void launch_mm_ids_helper( const int32_t * __restrict__ ids, int32_t * __restrict__ ids_src1, int32_t * __restrict__ ids_dst, int32_t * __restrict__ expert_bounds, - const int n_experts, const int n_tokens, const int n_expert_used_var, const int nchannels_y, const int si1, const int sis1, cudaStream_t stream) { + const int n_experts, const int n_tokens, const int n_expert_used_var, const int nchannels_y, const int si1, const int sis1, const bool write_inverse, cudaStream_t stream) { GGML_ASSERT(n_tokens < (1 << 22) && "too few bits in mm_ids_helper_store"); GGML_ASSERT(n_expert_used_var < (1 << 10) && "too few bits in mm_ids_helper_store"); @@ -132,33 +137,33 @@ static void launch_mm_ids_helper( const size_t nbytes_shared = n_tokens*sizeof(mm_ids_helper_store); GGML_ASSERT(nbytes_shared <= smpbo); mm_ids_helper<<>> - (ids, ids_src1, ids_dst, expert_bounds, n_tokens, n_expert_used_var, nchannels_y, si1, sis1); + (ids, ids_src1, ids_dst, expert_bounds, n_tokens, n_expert_used_var, nchannels_y, si1, sis1, write_inverse); } void ggml_cuda_launch_mm_ids_helper( const int32_t * __restrict__ ids, int32_t * __restrict__ ids_src1, int32_t * __restrict__ ids_dst, int32_t * __restrict__ expert_bounds, - const int n_experts, const int n_tokens, const int n_expert_used, const int nchannels_y, const int si1, const int sis1, cudaStream_t stream) { + const int n_experts, const int n_tokens, const int n_expert_used, const int nchannels_y, const int si1, const int sis1, const bool write_inverse, cudaStream_t stream) { switch (n_expert_used) { case 2: - launch_mm_ids_helper< 2>(ids, ids_src1, ids_dst, expert_bounds, n_experts, n_tokens, n_expert_used, nchannels_y, si1, sis1, stream); + launch_mm_ids_helper< 2>(ids, ids_src1, ids_dst, expert_bounds, n_experts, n_tokens, n_expert_used, nchannels_y, si1, sis1, write_inverse, stream); break; case 4: - launch_mm_ids_helper< 4>(ids, ids_src1, ids_dst, expert_bounds, n_experts, n_tokens, n_expert_used, nchannels_y, si1, sis1, stream); + launch_mm_ids_helper< 4>(ids, ids_src1, ids_dst, expert_bounds, n_experts, n_tokens, n_expert_used, nchannels_y, si1, sis1, write_inverse, stream); break; case 6: - launch_mm_ids_helper< 6>(ids, ids_src1, ids_dst, expert_bounds, n_experts, n_tokens, n_expert_used, nchannels_y, si1, sis1, stream); + launch_mm_ids_helper< 6>(ids, ids_src1, ids_dst, expert_bounds, n_experts, n_tokens, n_expert_used, nchannels_y, si1, sis1, write_inverse, stream); break; case 8: - launch_mm_ids_helper< 8>(ids, ids_src1, ids_dst, expert_bounds, n_experts, n_tokens, n_expert_used, nchannels_y, si1, sis1, stream); + launch_mm_ids_helper< 8>(ids, ids_src1, ids_dst, expert_bounds, n_experts, n_tokens, n_expert_used, nchannels_y, si1, sis1, write_inverse, stream); break; case 16: - launch_mm_ids_helper<16>(ids, ids_src1, ids_dst, expert_bounds, n_experts, n_tokens, n_expert_used, nchannels_y, si1, sis1, stream); + launch_mm_ids_helper<16>(ids, ids_src1, ids_dst, expert_bounds, n_experts, n_tokens, n_expert_used, nchannels_y, si1, sis1, write_inverse, stream); break; case 32: - launch_mm_ids_helper<32>(ids, ids_src1, ids_dst, expert_bounds, n_experts, n_tokens, n_expert_used, nchannels_y, si1, sis1, stream); + launch_mm_ids_helper<32>(ids, ids_src1, ids_dst, expert_bounds, n_experts, n_tokens, n_expert_used, nchannels_y, si1, sis1, write_inverse, stream); break; default: - launch_mm_ids_helper< 0>(ids, ids_src1, ids_dst, expert_bounds, n_experts, n_tokens, n_expert_used, nchannels_y, si1, sis1, stream); + launch_mm_ids_helper< 0>(ids, ids_src1, ids_dst, expert_bounds, n_experts, n_tokens, n_expert_used, nchannels_y, si1, sis1, write_inverse, stream); break; } } diff --git a/ggml/src/ggml-cuda/mmid.cuh b/ggml/src/ggml-cuda/mmid.cuh index ac090aea9..74c2db433 100644 --- a/ggml/src/ggml-cuda/mmid.cuh +++ b/ggml/src/ggml-cuda/mmid.cuh @@ -2,4 +2,4 @@ void ggml_cuda_launch_mm_ids_helper( const int32_t * ids, int32_t * ids_src1, int32_t * ids_dst, int32_t * expert_bounds, - int n_experts, int n_tokens, int n_expert_used, int nchannels_y, int si1, int sis1, cudaStream_t stream); + int n_experts, int n_tokens, int n_expert_used, int nchannels_y, int si1, int sis1, bool write_inverse, cudaStream_t stream); diff --git a/ggml/src/ggml-cuda/mmq-load-tiles.cuh b/ggml/src/ggml-cuda/mmq-load-tiles.cuh index 3978b1baa..7fb242096 100644 --- a/ggml/src/ggml-cuda/mmq-load-tiles.cuh +++ b/ggml/src/ggml-cuda/mmq-load-tiles.cuh @@ -39,29 +39,37 @@ template static __device__ __forceinline_ } const block_q1_0 * bxi = (const block_q1_0 *) x + kbx0 + i*stride + kbx; - const int qs_offset = 4*kqsx; - const int qs0 = bxi->qs[qs_offset + 0] | (bxi->qs[qs_offset + 1] << 8) | - (bxi->qs[qs_offset + 2] << 16) | (bxi->qs[qs_offset + 3] << 24); - - int unpacked_bytes[8]; -#pragma unroll - for (int j = 0; j < 8; ++j) { - const int shift = j * 4; - const int bits4 = (qs0 >> shift) & 0x0F; - const int b0 = (bits4 & 0x01) ? 1 : -1; - const int b1 = (bits4 & 0x02) ? 1 : -1; - const int b2 = (bits4 & 0x04) ? 1 : -1; - const int b3 = (bits4 & 0x08) ? 1 : -1; - unpacked_bytes[j] = (b0 & 0xFF) | ((b1 & 0xFF) << 8) | ((b2 & 0xFF) << 16) | ((b3 & 0xFF) << 24); - } + const int16_t * qxi = (const int16_t *) bxi->qs + kqsx * 2; const int dst_offset = kbx*(scale_entries_per_block*QI8_0) + kqsx*QI8_0; #pragma unroll - for (int j = 0; j < 8; ++j) { + for (int j = 0; j < 2; ++j) { + const int q = qxi[j]; + + // unpack crumbs into nibble indices + const int n0 = __byte_perm(0x11100100, 0x11100100, q >> 0); // [0, 1, 4, 5] [ 8, 9, 12, 13] + const int n1 = __byte_perm(0x11100100, 0x11100100, q >> 2); // [2, 3, 6, 7] [10, 11, 14, 15] + // unpack nibbles into byte values + const int s0 = __byte_perm(0x01FF, 0x01FF, n0 >> 0); + const int s1 = __byte_perm(0x01FF, 0x01FF, n1 >> 0); + const int s2 = __byte_perm(0x01FF, 0x01FF, n0 >> 16); + const int s3 = __byte_perm(0x01FF, 0x01FF, n1 >> 16); + // unshuffle values + const int v0 = __byte_perm(s0, s1, 0x5410); + const int v1 = __byte_perm(s0, s1, 0x7632); + const int v2 = __byte_perm(s2, s3, 0x5410); + const int v3 = __byte_perm(s2, s3, 0x7632); + #if defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) - x_qs[i*sram_stride + dst_offset + j] = unpacked_bytes[j]; + x_qs[i*sram_stride + dst_offset + j*4+0] = v0; + x_qs[i*sram_stride + dst_offset + j*4+1] = v1; + x_qs[i*sram_stride + dst_offset + j*4+2] = v2; + x_qs[i*sram_stride + dst_offset + j*4+3] = v3; #else - x_qs[i*(2*MMQ_TILE_NE_K + 1) + dst_offset + j] = unpacked_bytes[j]; + x_qs[i*(2*MMQ_TILE_NE_K + 1) + dst_offset + j*4+0] = v0; + x_qs[i*(2*MMQ_TILE_NE_K + 1) + dst_offset + j*4+1] = v1; + x_qs[i*(2*MMQ_TILE_NE_K + 1) + dst_offset + j*4+2] = v2; + x_qs[i*(2*MMQ_TILE_NE_K + 1) + dst_offset + j*4+3] = v3; #endif // defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) } } diff --git a/ggml/src/ggml-cuda/mmq.cu b/ggml/src/ggml-cuda/mmq.cu index 34b6cc47a..0a1062dd1 100644 --- a/ggml/src/ggml-cuda/mmq.cu +++ b/ggml/src/ggml-cuda/mmq.cu @@ -122,11 +122,12 @@ void ggml_cuda_mul_mat_q( const bool fallback = ne01 % 128 != 0; - // TODO: tighter pool buffer size vs q8 path const bool use_native_fp4 = blackwell_mma_available(cc) && (src0->type == GGML_TYPE_MXFP4 || src0->type == GGML_TYPE_NVFP4); + const size_t y_block_size = use_native_fp4 ? sizeof(block_fp4_mmq) : sizeof(block_q8_1_mmq); + const size_t y_values_per_block = use_native_fp4 ? QK_FP4_MMQ : QK8_1_MMQ; if (!ids) { - const size_t nbytes_src1_q8_1 = ne13*ne12 * ne11*ne10_padded * sizeof(block_q8_1)/QK8_1 + + const size_t nbytes_src1_q8_1 = ne13*ne12 * ne11*ne10_padded * y_block_size/y_values_per_block + ggml_cuda_mmq_get_J_max(src0->type, fallback, cc, ne11) * sizeof(block_q8_1_mmq); ggml_cuda_pool_alloc src1_q8_1(ctx.pool(), nbytes_src1_q8_1); @@ -148,7 +149,7 @@ void ggml_cuda_mul_mat_q( // Stride depends on quantization format const int64_t s12 = use_native_fp4 ? - ne11 * ne10_padded * sizeof(block_fp4_mmq) / (QK_K * sizeof(int)) : // block_fp4_mmq holds 256 values + ne11 * ne10_padded * sizeof(block_fp4_mmq) / (QK_FP4_MMQ * sizeof(int)) : ne11 * ne10_padded * sizeof(block_q8_1) / (QK8_1 * sizeof(int)); const int64_t s13 = ne12*s12; @@ -174,17 +175,21 @@ void ggml_cuda_mul_mat_q( ggml_cuda_pool_alloc ids_dst(ctx.pool(), ne_get_rows); ggml_cuda_pool_alloc expert_bounds(ctx.pool(), ne02 + 1); + // gate/up activations are broadcast across experts (ne11 == 1): quantize each token once and + // scatter to its slots. ids_src1 then holds the inverse map (token slot -> compact row). + const bool dedup_bcast = ne11 == 1 && n_expert_used > 1; + { GGML_ASSERT(ids->nb[0] == ggml_element_size(ids)); const int si1 = ids->nb[1] / ggml_element_size(ids); const int sis1 = nb12 / nb11; ggml_cuda_launch_mm_ids_helper((const int32_t *) ids->data, ids_src1.get(), ids_dst.get(), expert_bounds.get(), - ne02, ne12, n_expert_used, ne11, si1, sis1, stream); + ne02, ne12, n_expert_used, ne11, si1, sis1, /*write_inverse =*/ dedup_bcast, stream); CUDA_CHECK(cudaGetLastError()); } - const size_t nbytes_src1_q8_1 = ne12*n_expert_used*ne10_padded * sizeof(block_q8_1)/QK8_1 + + const size_t nbytes_src1_q8_1 = ne12*n_expert_used*ne10_padded * y_block_size/y_values_per_block + ggml_cuda_mmq_get_J_max(src0->type, fallback, cc, ne11) * sizeof(block_q8_1_mmq); ggml_cuda_pool_alloc src1_q8_1(ctx.pool(), nbytes_src1_q8_1); @@ -197,7 +202,16 @@ void ggml_cuda_mul_mat_q( const int64_t s12 = src1->nb[2] / ts_src1; const int64_t s13 = src1->nb[3] / ts_src1; - if (use_native_fp4) { + if (dedup_bcast) { + // quantize each token once, scatter its block to all n_expert_used slots + if (use_native_fp4) { + quantize_scatter_mmq_fp4_cuda(src1_d, ids_src1.get(), src1_q8_1.get(), src0->type, ne10, + /*stride_token=*/s12, ne10_padded, ne12, ne11_flat, n_expert_used, stream); + } else { + quantize_scatter_mmq_q8_1_cuda(src1_d, ids_src1.get(), src1_q8_1.get(), src0->type, ne10, + /*stride_token=*/s12, ne10_padded, ne12, ne11_flat, n_expert_used, stream); + } + } else if (use_native_fp4) { quantize_mmq_fp4_cuda(src1_d, ids_src1.get(), src1_q8_1.get(), src0->type, ne10, s11, s12, s13, ne10_padded, ne11_flat, ne12_flat, ne13_flat, stream); } else { @@ -207,8 +221,8 @@ void ggml_cuda_mul_mat_q( CUDA_CHECK(cudaGetLastError()); } - static_assert(QK_K == 8 * QK_MXFP4, "QK_K needs to be 8 * QK_MXFP4"); - const int64_t s12 = use_native_fp4 ? ne11 * ne10_padded * sizeof(block_fp4_mmq) / (QK_K * sizeof(int)) : + static_assert(QK_FP4_MMQ == 8 * QK_MXFP4, "QK_FP4_MMQ needs to be 8 * QK_MXFP4"); + const int64_t s12 = use_native_fp4 ? ne11 * ne10_padded * sizeof(block_fp4_mmq) / (QK_FP4_MMQ * sizeof(int)) : ne11 * ne10_padded * sizeof(block_q8_1) / (QK8_1 * sizeof(int)); const int64_t s13 = ne12*s12; diff --git a/ggml/src/ggml-cuda/mmq.cuh b/ggml/src/ggml-cuda/mmq.cuh index ce7e4216c..df35902e2 100644 --- a/ggml/src/ggml-cuda/mmq.cuh +++ b/ggml/src/ggml-cuda/mmq.cuh @@ -22,6 +22,9 @@ enum mmq_q8_1_ds_layout { MMQ_Q8_1_DS_LAYOUT_D2S6, }; +static constexpr int QK8_1_MMQ = 4*QK8_1; +static constexpr int QK_FP4_MMQ = 2*QK8_1_MMQ; + struct block_q8_1_mmq { // The y float data is converted to a data layout that can simply be copied to shared memory as a contiguous block. // The y float data is first grouped as blocks of 128 values. @@ -40,7 +43,7 @@ struct block_q8_1_mmq { half d2s6[8]; // 1 16 bit scale per 64 values + 1 16 bit partial sum per 16 values for the first 96 values, // stored as d0,d1,s1,s2,s3,s4,s5 }; - int8_t qs[4*QK8_1]; // 128 values quantized to 8 bit each + int8_t qs[QK8_1_MMQ]; }; // this struct is used for fp4 data types (currently only used for Blackwell) @@ -48,10 +51,10 @@ struct block_q8_1_mmq { // nvfp4 has block size 16, each int32 of d4 contains 4 ue4m3 scales struct block_fp4_mmq { uint32_t d4[4]; - int8_t qs[4 * 32]; // 256 FP4 values packed as 4-bit pairs (2 per byte) + int8_t qs[QK_FP4_MMQ / 2]; }; -static_assert(sizeof(block_q8_1_mmq) == 4*QK8_1 + 4*sizeof(half2), "Unexpected block_q8_1_mmq size"); +static_assert(sizeof(block_q8_1_mmq) == QK8_1_MMQ + 4*sizeof(half2), "Unexpected block_q8_1_mmq size"); static_assert(sizeof(block_q8_1_mmq) == 4*sizeof(block_q8_1), "Unexpected block_q8_1_mmq size"); static_assert(sizeof(block_fp4_mmq) == sizeof(block_q8_1_mmq), "Unexpected block_fp4_mmq size"); @@ -834,9 +837,9 @@ static __device__ __forceinline__ void mul_mat_q_process_tile( #if defined(BLACKWELL_MMA_AVAILABLE) // FP4 tile stores 8 blocks - constexpr int ne_block = (type == GGML_TYPE_MXFP4 || type == GGML_TYPE_NVFP4) ? QK_K : 4 * QK8_1; + constexpr int ne_block = (type == GGML_TYPE_MXFP4 || type == GGML_TYPE_NVFP4) ? QK_FP4_MMQ : QK8_1_MMQ; #else - constexpr int ne_block = 4 * QK8_1; + constexpr int ne_block = QK8_1_MMQ; #endif // defined(BLACKWELL_MMA_AVAILABLE) constexpr int ITER_K = ggml_cuda_mmq_get_K_vram(type, J, fallback); diff --git a/ggml/src/ggml-cuda/quantize.cu b/ggml/src/ggml-cuda/quantize.cu index 39a500a17..af7fe105b 100644 --- a/ggml/src/ggml-cuda/quantize.cu +++ b/ggml/src/ggml-cuda/quantize.cu @@ -75,10 +75,12 @@ __device__ __forceinline__ uint8_t compute_e8m0_scale(float amax) { } +// scatter: grid over tokens, quantize once, write to all the token's compact rows +template static __global__ void quantize_mmq_nvfp4( const float * __restrict__ x, const int32_t * __restrict__ ids, void * __restrict__ vy, const int64_t ne00, const int64_t s01, const int64_t s02, const int64_t s03, - const int64_t ne0, const int64_t ne1, const int64_t ne2) { + const int64_t ne0, const int64_t ne1, const int64_t ne2, const int n_expert_used) { #if defined(BLACKWELL_MMA_AVAILABLE) const int64_t i0_base = ((int64_t) blockDim.x * blockIdx.y + threadIdx.x) * QK_NVFP4_SUB; @@ -86,25 +88,25 @@ static __global__ void quantize_mmq_nvfp4( return; } - const int64_t i1 = blockIdx.x; - const int64_t i2 = blockIdx.z % ne2; - const int64_t i3 = blockIdx.z / ne2; - const int64_t i01 = ids ? ids[i1] : i1; - const int64_t k_block = i0_base / QK_K; - const int64_t blocks_per_col = (ne0 + QK_K - 1) / QK_K; + const int64_t k_block = i0_base / QK_FP4_MMQ; + const int64_t blocks_per_col = (ne0 + QK_FP4_MMQ - 1) / QK_FP4_MMQ; if (k_block >= blocks_per_col) { return; } + const int sub = (i0_base % QK_FP4_MMQ) / QK_NVFP4_SUB; - const int64_t ib = blockIdx.z * ((int64_t) blocks_per_col * ne1) + k_block * ne1 + blockIdx.x; - block_fp4_mmq * y = (block_fp4_mmq *) vy; - block_fp4_mmq * yb = y + ib; - - const int sub = (i0_base % QK_K) / QK_NVFP4_SUB; + int64_t base_idx; + if constexpr (scatter) { + base_idx = (int64_t) blockIdx.x * s02; // one physical row per token + } else { + const int64_t i2 = blockIdx.z % ne2; + const int64_t i3 = blockIdx.z / ne2; + const int64_t i01 = ids ? ids[blockIdx.x] : blockIdx.x; + base_idx = i3 * s03 + i2 * s02 + i01 * s01; + } float vals_raw[QK_NVFP4_SUB]; float amax_raw = 0.0f; - const int64_t base_idx = i3 * s03 + i2 * s02 + i01 * s01; #pragma unroll for (int k = 0; k < QK_NVFP4_SUB; k++) { const int64_t i00 = i0_base + k; @@ -160,11 +162,27 @@ static __global__ void quantize_mmq_nvfp4( q1 |= (uint32_t) ggml_cuda_float_to_fp4_e2m1(vals_raw[k + 12], inv_scale) << (8 * k + 4); } - uint32_t * yqs = reinterpret_cast(yb->qs); - yqs[2 * sub + 0] = q0; - yqs[2 * sub + 1] = q1; - reinterpret_cast(yb->d4)[sub] = fp8_code; + block_fp4_mmq * y = (block_fp4_mmq *) vy; + if constexpr (scatter) { +#pragma unroll + for (int slot = 0; slot < n_expert_used; ++slot) { + const int64_t i = ids[(int64_t) blockIdx.x * n_expert_used + slot]; + block_fp4_mmq * yb = y + (k_block * ne1 + i); + uint32_t * yqs = reinterpret_cast(yb->qs); + yqs[2 * sub + 0] = q0; + yqs[2 * sub + 1] = q1; + reinterpret_cast(yb->d4)[sub] = fp8_code; + } + } else { + block_fp4_mmq * yb = y + (blockIdx.z * ((int64_t) blocks_per_col * ne1) + k_block * ne1 + blockIdx.x); + uint32_t * yqs = reinterpret_cast(yb->qs); + yqs[2 * sub + 0] = q0; + yqs[2 * sub + 1] = q1; + reinterpret_cast(yb->d4)[sub] = fp8_code; + } + GGML_UNUSED(n_expert_used); #else + GGML_UNUSED(n_expert_used); NO_DEVICE_CODE; // This is for Blackwell NVFP4 activations only. #endif // defined(BLACKWELL_MMA_AVAILABLE) @@ -172,6 +190,8 @@ static __global__ void quantize_mmq_nvfp4( // quantize values in the format mxfp4 is stored which is interleaved nibbles // i.e. a block a0-a31 is represented as a0a16,a1a17 ...a15a31 +// scatter: grid over tokens, quantize once, write to all the token's compact rows +template static __global__ void quantize_mmq_mxfp4(const float * __restrict__ x, const int32_t * __restrict__ ids, void * __restrict__ vy, @@ -181,7 +201,8 @@ static __global__ void quantize_mmq_mxfp4(const float * __restrict__ x, const int64_t s03, const int64_t ne0, const int ne1, - const int ne2) { + const int ne2, + const int n_expert_used) { constexpr int vals_per_scale = 32; constexpr int vals_per_warp = 2 * vals_per_scale; // Each warp processes 2 blocks of 32 = 64 values @@ -196,30 +217,27 @@ static __global__ void quantize_mmq_mxfp4(const float * __restrict__ x, return; } - const int64_t i1 = blockIdx.x; - const int64_t i2 = blockIdx.z % ne2; - const int64_t i3 = blockIdx.z / ne2; - - ggml_cuda_pdl_sync(); - const int64_t i01 = ids ? ids[i1] : i1; - const int64_t i02 = i2; - const int64_t i03 = i3; - - block_fp4_mmq * y = (block_fp4_mmq *) vy; - - const int64_t block_fp4_mmq_size = 8 * QK_MXFP4; // 256 values - const int64_t ib0 = blockIdx.z * ((int64_t) ne1 * (ne0 / block_fp4_mmq_size)); - const int64_t ib = ib0 + (warp_start_offset / block_fp4_mmq_size) * ne1 + blockIdx.x; + const int64_t block_fp4_mmq_size = QK_FP4_MMQ; + const int64_t k_block = warp_start_offset / block_fp4_mmq_size; const int64_t quad_idx_in_block = (warp_start_offset % block_fp4_mmq_size) / vals_per_warp; const int group_id = lane_id_32 / 4; const int lane_in_group = lane_id_32 % 4; const int base = group_id * 2; - char2 * yqs2 = (char2 *) y[ib].qs; - const int64_t base_pos = i03 * s03 + i02 * s02 + i01 * s01; + ggml_cuda_pdl_sync(); + int64_t base_pos; + if constexpr (scatter) { + base_pos = (int64_t) blockIdx.x * s02; // one physical row per token + } else { + const int64_t i2 = blockIdx.z % ne2; + const int64_t i3 = blockIdx.z / ne2; + const int64_t i01 = ids ? ids[blockIdx.x] : blockIdx.x; + base_pos = i3 * s03 + i2 * s02 + i01 * s01; + } uint8_t scales[2]; + char2 packed[2]; #pragma unroll for (int b = 0; b < 2; ++b) { @@ -244,11 +262,8 @@ static __global__ void quantize_mmq_mxfp4(const float * __restrict__ x, const float val2 = __shfl_sync(0xFFFFFFFF, scaled_val, base + 1, WARP_SIZE); const float val3 = __shfl_sync(0xFFFFFFFF, scaled_val, base + 17, WARP_SIZE); - if (lane_in_group == 0) { - __nv_fp4x4_e2m1 fp4_packed(make_float4(val0, val1, val2, val3)); - - yqs2[quad_idx_in_block * 16 + b * 8 + group_id] = *(char2 *) &fp4_packed; - } + __nv_fp4x4_e2m1 fp4_packed(make_float4(val0, val1, val2, val3)); + packed[b] = *(char2 *) &fp4_packed; #else // Fallback: manual FP4 conversion using LUT const uint8_t q_val = ggml_cuda_float_to_fp4_e2m1(xi, inv_s); @@ -258,26 +273,49 @@ static __global__ void quantize_mmq_mxfp4(const float * __restrict__ x, const uint8_t q_hi_0 = __shfl_sync(0xFFFFFFFF, q_val, base + 16, WARP_SIZE); const uint8_t q_hi_1 = __shfl_sync(0xFFFFFFFF, q_val, base + 17, WARP_SIZE); - if (lane_in_group == 0) { - char2 q; - q.x = (q_hi_0 << 4) | q_lo_0; - q.y = (q_hi_1 << 4) | q_lo_1; - yqs2[quad_idx_in_block * 16 + b * 8 + group_id] = q; - } + char2 q; + q.x = (q_hi_0 << 4) | q_lo_0; + q.y = (q_hi_1 << 4) | q_lo_1; + packed[b] = q; #endif // CUDART_VERSION >= 12080 } - if (lane_id_32 == 0) { - // Store 2 scales packed into 1 uint32 - y[ib].d4[quad_idx_in_block] = (scales[1] << 8) | scales[0]; + block_fp4_mmq * y = (block_fp4_mmq *) vy; + if constexpr (scatter) { +#pragma unroll + for (int slot = 0; slot < n_expert_used; ++slot) { + const int64_t i = ids[(int64_t) blockIdx.x * n_expert_used + slot]; + block_fp4_mmq * yb = y + (k_block * ne1 + i); + char2 * yqs2 = (char2 *) yb->qs; + if (lane_in_group == 0) { + yqs2[quad_idx_in_block * 16 + 0 * 8 + group_id] = packed[0]; + yqs2[quad_idx_in_block * 16 + 1 * 8 + group_id] = packed[1]; + } + if (lane_id_32 == 0) { + yb->d4[quad_idx_in_block] = (scales[1] << 8) | scales[0]; + } + } + } else { + const int64_t ib0 = blockIdx.z * ((int64_t) ne1 * (ne0 / block_fp4_mmq_size)); + block_fp4_mmq * yb = y + (ib0 + k_block * ne1 + blockIdx.x); + char2 * yqs2 = (char2 *) yb->qs; + if (lane_in_group == 0) { + yqs2[quad_idx_in_block * 16 + 0 * 8 + group_id] = packed[0]; + yqs2[quad_idx_in_block * 16 + 1 * 8 + group_id] = packed[1]; + } + if (lane_id_32 == 0) { + yb->d4[quad_idx_in_block] = (scales[1] << 8) | scales[0]; + } } + GGML_UNUSED(n_expert_used); } -template +// scatter: grid over tokens, quantize once, write to all the token's compact rows +template static __global__ void quantize_mmq_q8_1( const float * __restrict__ x, const int32_t * __restrict__ ids, void * __restrict__ vy, const int64_t ne00, const int64_t s01, const int64_t s02, const int64_t s03, - const int64_t ne0, const int ne1, const int ne2) { + const int64_t ne0, const int ne1, const int ne2, const int n_expert_used) { constexpr int vals_per_scale = ds_layout == MMQ_Q8_1_DS_LAYOUT_D2S6 ? 64 : 32; constexpr int vals_per_sum = ds_layout == MMQ_Q8_1_DS_LAYOUT_D2S6 ? 16 : 32; @@ -288,26 +326,27 @@ static __global__ void quantize_mmq_q8_1( return; } - const int64_t i1 = blockIdx.x; - const int64_t i2 = blockIdx.z % ne2; - const int64_t i3 = blockIdx.z / ne2; - const int64_t i00 = i0; ggml_cuda_pdl_sync(); - const int64_t i01 = ids ? ids[i1] : i1; - const int64_t i02 = i2; - const int64_t i03 = i3; + + int64_t base_idx; + if constexpr (scatter) { + base_idx = (int64_t) blockIdx.x * s02; // one physical row per token + } else { + const int64_t i2 = blockIdx.z % ne2; + const int64_t i3 = blockIdx.z / ne2; + const int64_t i01 = ids ? ids[blockIdx.x] : blockIdx.x; + base_idx = i3*s03 + i2*s02 + i01*s01; + } const float4 * x4 = (const float4 *) x; - block_q8_1_mmq * y = (block_q8_1_mmq *) vy; - const int64_t ib0 = blockIdx.z*((int64_t)gridDim.x*gridDim.y*blockDim.x/QK8_1); // first block of channel - const int64_t ib = ib0 + (i0 / (4*QK8_1))*ne1 + blockIdx.x; // block index in channel - const int64_t iqs = i0 % (4*QK8_1); // quant index in block + const int64_t k_block = i0 / QK8_1_MMQ; // column block in the channel + const int64_t iqs = i0 % QK8_1_MMQ; // quant index in block // Load 4 floats per thread and calculate max. abs. value between them: - const float4 xi = i0 < ne00 ? x4[(i03*s03 + i02*s02 + i01*s01 + i00)/4] : make_float4(0.0f, 0.0f, 0.0f, 0.0f); + const float4 xi = i0 < ne00 ? x4[(base_idx + i00)/4] : make_float4(0.0f, 0.0f, 0.0f, 0.0f); float amax = fabsf(xi.x); amax = fmaxf(amax, fabsf(xi.y)); amax = fmaxf(amax, fabsf(xi.z)); @@ -336,40 +375,41 @@ static __global__ void quantize_mmq_q8_1( q.y = roundf(xi.y*d_inv); q.z = roundf(xi.z*d_inv); q.w = roundf(xi.w*d_inv); - - // Write back 4 int8 values as a single 32 bit value for better memory bandwidth: - char4 * yqs4 = (char4 *) y[ib].qs; - yqs4[iqs/4] = q; - - if (ds_layout == MMQ_Q8_1_DS_LAYOUT_D2S6) { - if (iqs % 16 != 0 || iqs >= 96) { - return; - } - - y[ib].d2s6[2 + iqs/16] = sum; - - if (iqs % 64 != 0) { - return; - } - - const float d = 1.0f / d_inv; - - y[ib].d2s6[iqs/64] = d; - - return; - } - - if (iqs % 32 != 0) { - return; - } - const float d = 1.0f / d_inv; - if (ds_layout == MMQ_Q8_1_DS_LAYOUT_DS4) { - y[ib].ds4[iqs/32] = make_half2(d, sum); - } else { - y[ib].d4[iqs/32] = d; + // write the block once (normal) or to each of the token's compact rows (scatter) + const int nwrite = scatter ? n_expert_used : 1; +#pragma unroll + for (int slot = 0; slot < nwrite; ++slot) { + int64_t ib; + if constexpr (scatter) { + const int64_t i = ids[(int64_t) blockIdx.x * n_expert_used + slot]; + ib = k_block*ne1 + i; + } else { + const int64_t ib0 = blockIdx.z*((int64_t)gridDim.x*gridDim.y*blockDim.x/QK8_1); // first block of channel + ib = ib0 + k_block*ne1 + blockIdx.x; + } + + // Write back 4 int8 values as a single 32 bit value for better memory bandwidth: + char4 * yqs4 = (char4 *) y[ib].qs; + yqs4[iqs/4] = q; + + if (ds_layout == MMQ_Q8_1_DS_LAYOUT_D2S6) { + if (iqs % 16 == 0 && iqs < 96) { + y[ib].d2s6[2 + iqs/16] = sum; + if (iqs % 64 == 0) { + y[ib].d2s6[iqs/64] = d; + } + } + } else if (iqs % 32 == 0) { + if (ds_layout == MMQ_Q8_1_DS_LAYOUT_DS4) { + y[ib].ds4[iqs/32] = make_half2(d, sum); + } else { + y[ib].d4[iqs/32] = d; + } + } } + GGML_UNUSED(n_expert_used); } void quantize_row_q8_1_cuda( @@ -394,7 +434,7 @@ void quantize_mmq_q8_1_cuda( const int64_t ne00, const int64_t s01, const int64_t s02, const int64_t s03, const int64_t ne0, const int64_t ne1, const int64_t ne2, const int64_t ne3, cudaStream_t stream) { GGML_ASSERT(ne00 % 4 == 0); - GGML_ASSERT(ne0 % (4*QK8_1) == 0); + GGML_ASSERT(ne0 % QK8_1_MMQ == 0); // ne1 tends to assume the highest values, therefore use it as the "x" dimension of the CUDA grid: const int64_t block_num_y = (ne0 + 4*CUDA_QUANTIZE_BLOCK_SIZE_MMQ - 1) / (4*CUDA_QUANTIZE_BLOCK_SIZE_MMQ); @@ -402,16 +442,16 @@ void quantize_mmq_q8_1_cuda( const dim3 block_size(CUDA_QUANTIZE_BLOCK_SIZE_MMQ, 1, 1); switch (mmq_get_q8_1_ds_layout(type_src0)) { case MMQ_Q8_1_DS_LAYOUT_D4: - quantize_mmq_q8_1 - <<>>(x, ids, vy, ne00, s01, s02, s03, ne0, ne1, ne2); + quantize_mmq_q8_1 + <<>>(x, ids, vy, ne00, s01, s02, s03, ne0, ne1, ne2, /*n_expert_used=*/0); break; case MMQ_Q8_1_DS_LAYOUT_DS4: - quantize_mmq_q8_1 - <<>>(x, ids, vy, ne00, s01, s02, s03, ne0, ne1, ne2); + quantize_mmq_q8_1 + <<>>(x, ids, vy, ne00, s01, s02, s03, ne0, ne1, ne2, /*n_expert_used=*/0); break; case MMQ_Q8_1_DS_LAYOUT_D2S6: - quantize_mmq_q8_1 - <<>>(x, ids, vy, ne00, s01, s02, s03, ne0, ne1, ne2); + quantize_mmq_q8_1 + <<>>(x, ids, vy, ne00, s01, s02, s03, ne0, ne1, ne2, /*n_expert_used=*/0); break; default: GGML_ABORT("fatal error"); @@ -419,6 +459,62 @@ void quantize_mmq_q8_1_cuda( } } +// scatter=true reuses the quant kernel: grid over tokens, ids = inverse map (token slot -> compact row) +void quantize_scatter_mmq_q8_1_cuda( + const float * x, const int32_t * ids_src1_inv, void * vy, const ggml_type type_src0, + const int64_t ne00, const int64_t stride_token, const int64_t ne0, + const int64_t n_tokens, const int64_t nrows_dst, const int n_expert_used, cudaStream_t stream) { + GGML_ASSERT(ne00 % 4 == 0); + GGML_ASSERT(ne0 % QK8_1_MMQ == 0); + + const int64_t block_num_y = (ne0 + 4*CUDA_QUANTIZE_BLOCK_SIZE_MMQ - 1) / (4*CUDA_QUANTIZE_BLOCK_SIZE_MMQ); + const dim3 num_blocks(n_tokens, block_num_y, 1); + const dim3 block_size(CUDA_QUANTIZE_BLOCK_SIZE_MMQ, 1, 1); + switch (mmq_get_q8_1_ds_layout(type_src0)) { + case MMQ_Q8_1_DS_LAYOUT_D4: + quantize_mmq_q8_1<<>>( + x, ids_src1_inv, vy, ne00, /*s01=*/0, /*s02=*/stride_token, /*s03=*/0, ne0, /*ne1=*/(int) nrows_dst, /*ne2=*/1, n_expert_used); + break; + case MMQ_Q8_1_DS_LAYOUT_DS4: + quantize_mmq_q8_1<<>>( + x, ids_src1_inv, vy, ne00, /*s01=*/0, /*s02=*/stride_token, /*s03=*/0, ne0, /*ne1=*/(int) nrows_dst, /*ne2=*/1, n_expert_used); + break; + case MMQ_Q8_1_DS_LAYOUT_D2S6: + quantize_mmq_q8_1<<>>( + x, ids_src1_inv, vy, ne00, /*s01=*/0, /*s02=*/stride_token, /*s03=*/0, ne0, /*ne1=*/(int) nrows_dst, /*ne2=*/1, n_expert_used); + break; + default: + GGML_ABORT("fatal error"); + break; + } +} + +// scatter=true reuses the quant kernels: grid over tokens, ids = inverse map (token slot -> compact row) +void quantize_scatter_mmq_fp4_cuda( + const float * x, const int32_t * ids_src1_inv, void * vy, const ggml_type type_src0, + const int64_t ne00, const int64_t stride_token, const int64_t ne0, + const int64_t n_tokens, const int64_t nrows_dst, const int n_expert_used, cudaStream_t stream) { + GGML_ASSERT(ne0 > 0); + if (type_src0 == GGML_TYPE_NVFP4) { + GGML_ASSERT(ne00 % QK_NVFP4 == 0); + constexpr int nvfp4_block_size = 128; + const int64_t block_num_y = (ne0 + QK_NVFP4_SUB * nvfp4_block_size - 1) / (QK_NVFP4_SUB * nvfp4_block_size); + const dim3 block_size(nvfp4_block_size, 1, 1); + const dim3 num_blocks(n_tokens, block_num_y, 1); + quantize_mmq_nvfp4<<>>( + x, ids_src1_inv, vy, ne00, /*s01=*/0, /*s02=*/stride_token, /*s03=*/0, ne0, /*ne1=*/nrows_dst, /*ne2=*/1, n_expert_used); + } else { + GGML_ASSERT(type_src0 == GGML_TYPE_MXFP4); + constexpr int nwarps = 8; + constexpr int vals_per_block = nwarps * 2 * QK_MXFP4; + const int64_t block_num_y = (ne0 + vals_per_block - 1) / vals_per_block; + const dim3 block_size(WARP_SIZE, nwarps, 1); + const dim3 num_blocks(n_tokens, block_num_y, 1); + quantize_mmq_mxfp4<<>>( + x, ids_src1_inv, vy, ne00, /*s01=*/0, /*s02=*/stride_token, /*s03=*/0, ne0, /*ne1=*/(int) nrows_dst, /*ne2=*/1, n_expert_used); + } +} + void quantize_mmq_fp4_cuda( const float * x, const int32_t * ids, void * vy, const ggml_type type_src0, const int64_t ne00, const int64_t s01, const int64_t s02, const int64_t s03, @@ -432,8 +528,8 @@ void quantize_mmq_fp4_cuda( const int64_t block_num_y = (ne0 + QK_NVFP4_SUB * nvfp4_block_size - 1) / (QK_NVFP4_SUB * nvfp4_block_size); const dim3 block_size(nvfp4_block_size, 1, 1); const dim3 num_blocks(ne1, block_num_y, ne2 * ne3); - quantize_mmq_nvfp4<<>>( - x, ids, vy, ne00, s01, s02, s03, ne0, ne1, ne2); + quantize_mmq_nvfp4<<>>( + x, ids, vy, ne00, s01, s02, s03, ne0, ne1, ne2, /*n_expert_used=*/0); } else { GGML_ASSERT(ne0 % (2 * QK_MXFP4) == 0); @@ -445,6 +541,6 @@ void quantize_mmq_fp4_cuda( const dim3 num_blocks(ne1, block_num_y, ne2 * ne3); const dim3 block_size(WARP_SIZE, nwarps, 1); - quantize_mmq_mxfp4<<>>(x, ids, vy, ne00, s01, s02, s03, ne0, ne1, ne2); + quantize_mmq_mxfp4<<>>(x, ids, vy, ne00, s01, s02, s03, ne0, ne1, ne2, /*n_expert_used=*/0); } } diff --git a/ggml/src/ggml-cuda/quantize.cuh b/ggml/src/ggml-cuda/quantize.cuh index 768a3ae6d..ed47fa027 100644 --- a/ggml/src/ggml-cuda/quantize.cuh +++ b/ggml/src/ggml-cuda/quantize.cuh @@ -39,3 +39,28 @@ void quantize_mmq_fp4_cuda(const float * x, int64_t ne2, int64_t ne3, cudaStream_t stream); + +// quantize each token once and scatter the block to its compact rows (via the inverse map) +void quantize_scatter_mmq_fp4_cuda(const float * x, + const int32_t * ids_src1_inv, + void * vy, + ggml_type type_src0, + int64_t ne00, + int64_t stride_token, + int64_t ne0, + int64_t n_tokens, + int64_t nrows_dst, + int n_expert_used, + cudaStream_t stream); + +void quantize_scatter_mmq_q8_1_cuda(const float * x, + const int32_t * ids_src1_inv, + void * vy, + ggml_type type_src0, + int64_t ne00, + int64_t stride_token, + int64_t ne0, + int64_t n_tokens, + int64_t nrows_dst, + int n_expert_used, + cudaStream_t stream); diff --git a/ggml/src/ggml-cuda/vecdotq.cuh b/ggml/src/ggml-cuda/vecdotq.cuh index d1741cc8d..b9932bce9 100644 --- a/ggml/src/ggml-cuda/vecdotq.cuh +++ b/ggml/src/ggml-cuda/vecdotq.cuh @@ -681,35 +681,40 @@ static __device__ __forceinline__ float vec_dot_q1_0_q8_1( // Q8_1: 32 elements per block with individual scales // iqs selects which of the 4 chunks of 32 elements to process (0-3) - const float d1 = bq1_0->d; + const float d1 = bq1_0->d; + const int16_t * qs = (const int16_t *) bq1_0->qs + iqs * 2; // Process only the chunk specified by iqs const block_q8_1 * bq8_1_chunk = bq8_1 + iqs; - // Load 32 bits (4 bytes) for this chunk from Q1_0 - const int offset = iqs * 4; - const int v = bq1_0->qs[offset + 0] | (bq1_0->qs[offset + 1] << 8) | - (bq1_0->qs[offset + 2] << 16) | (bq1_0->qs[offset + 3] << 24); - - // Unpack 32 bits into 32 signed values (-1 or +1) - int vi_bytes[8]; -#pragma unroll - for (int j = 0; j < 8; ++j) { - const int shift = j * 4; - const int bits4 = (v >> shift) & 0x0F; - const int b0 = (bits4 & 0x01) ? 1 : -1; - const int b1 = (bits4 & 0x02) ? 1 : -1; - const int b2 = (bits4 & 0x04) ? 1 : -1; - const int b3 = (bits4 & 0x08) ? 1 : -1; - vi_bytes[j] = (b0 & 0xFF) | ((b1 & 0xFF) << 8) | ((b2 & 0xFF) << 16) | ((b3 & 0xFF) << 24); - } - - // Compute dot product for this 32-element chunk int sumi = 0; #pragma unroll - for (int j = 0; j < 8; ++j) { - const int u = get_int_b4(bq8_1_chunk->qs, j); - sumi = ggml_cuda_dp4a(vi_bytes[j], u, sumi); + for (int j = 0; j < 2; ++j) { + const int q = qs[j]; + + const int u0 = get_int_b4(bq8_1_chunk->qs, j*4+0); + const int u1 = get_int_b4(bq8_1_chunk->qs, j*4+1); + const int u2 = get_int_b4(bq8_1_chunk->qs, j*4+2); + const int u3 = get_int_b4(bq8_1_chunk->qs, j*4+3); + + // unpack crumbs into nibble indices + const int n0 = __byte_perm(0x11100100, 0x11100100, q >> 0); // [0, 1, 4, 5] [ 8, 9, 12, 13] + const int n1 = __byte_perm(0x11100100, 0x11100100, q >> 2); // [2, 3, 6, 7] [10, 11, 14, 15] + // unpack nibbles into byte values + const int s0 = __byte_perm(0x01FF, 0x01FF, n0 >> 0); + const int s1 = __byte_perm(0x01FF, 0x01FF, n1 >> 0); + const int s2 = __byte_perm(0x01FF, 0x01FF, n0 >> 16); + const int s3 = __byte_perm(0x01FF, 0x01FF, n1 >> 16); + // unshuffle values + const int v0 = __byte_perm(s0, s1, 0x5410); + const int v1 = __byte_perm(s0, s1, 0x7632); + const int v2 = __byte_perm(s2, s3, 0x5410); + const int v3 = __byte_perm(s2, s3, 0x7632); + + sumi = ggml_cuda_dp4a(v0, u0, sumi); + sumi = ggml_cuda_dp4a(v1, u1, sumi); + sumi = ggml_cuda_dp4a(v2, u2, sumi); + sumi = ggml_cuda_dp4a(v3, u3, sumi); } // Apply Q1_0's single scale and this chunk's Q8_1 scale diff --git a/ggml/src/ggml-metal/ggml-metal-device.cpp b/ggml/src/ggml-metal/ggml-metal-device.cpp index 870f93df0..270c1411a 100644 --- a/ggml/src/ggml-metal/ggml-metal-device.cpp +++ b/ggml/src/ggml-metal/ggml-metal-device.cpp @@ -1834,6 +1834,23 @@ ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_col2im_1d(ggml_m return res; } +ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_snake(ggml_metal_library_t lib, enum ggml_type type) { + GGML_ASSERT(type == GGML_TYPE_F32 || type == GGML_TYPE_F16 || type == GGML_TYPE_BF16); + + char base[256]; + char name[256]; + + snprintf(base, 256, "kernel_snake_%s", ggml_type_name(type)); + snprintf(name, 256, "%s", base); + + ggml_metal_pipeline_with_params res = ggml_metal_library_get_pipeline(lib, name); + if (!res.pipeline) { + res = ggml_metal_library_compile_pipeline(lib, base, name, nullptr); + } + + return res; +} + ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_conv_transpose_2d(ggml_metal_library_t lib, const ggml_tensor * op) { assert(op->op == GGML_OP_CONV_TRANSPOSE_2D); diff --git a/ggml/src/ggml-metal/ggml-metal-device.h b/ggml/src/ggml-metal/ggml-metal-device.h index 9d4aca121..b36fa8110 100644 --- a/ggml/src/ggml-metal/ggml-metal-device.h +++ b/ggml/src/ggml-metal/ggml-metal-device.h @@ -151,6 +151,7 @@ struct ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_im2col struct ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_conv_transpose_1d (ggml_metal_library_t lib, const struct ggml_tensor * op); struct ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_conv_transpose_2d (ggml_metal_library_t lib, const struct ggml_tensor * op); struct ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_col2im_1d (ggml_metal_library_t lib, const struct ggml_tensor * op); +struct ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_snake (ggml_metal_library_t lib, enum ggml_type type); struct ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_conv_2d (ggml_metal_library_t lib, const struct ggml_tensor * op); struct ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_conv_2d_dw (ggml_metal_library_t lib, const struct ggml_tensor * op, bool tiled); struct ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_conv_3d (ggml_metal_library_t lib, const struct ggml_tensor * op); diff --git a/ggml/src/ggml-metal/ggml-metal-impl.h b/ggml/src/ggml-metal/ggml-metal-impl.h index a45f3ac67..330278d00 100644 --- a/ggml/src/ggml-metal/ggml-metal-impl.h +++ b/ggml/src/ggml-metal/ggml-metal-impl.h @@ -616,6 +616,11 @@ typedef struct { int32_t p0; } ggml_metal_kargs_col2im_1d; +typedef struct { + int32_t T; + int32_t C; +} ggml_metal_kargs_snake; + typedef struct { int32_t IC; int32_t IH; diff --git a/ggml/src/ggml-metal/ggml-metal-ops.cpp b/ggml/src/ggml-metal/ggml-metal-ops.cpp index 805bc4093..c716f118f 100644 --- a/ggml/src/ggml-metal/ggml-metal-ops.cpp +++ b/ggml/src/ggml-metal/ggml-metal-ops.cpp @@ -3077,7 +3077,58 @@ int ggml_metal_op_flash_attn_ext(ggml_metal_op_t ctx, int idx) { return 1; } +// Snake activation autofuse: mul -> sin -> sqr -> mul -> add +static bool ggml_metal_op_can_fuse_snake(ggml_metal_op_t ctx, int idx) { + static constexpr ggml_op snake_ops[5] = { GGML_OP_MUL, GGML_OP_SIN, GGML_OP_SQR, GGML_OP_MUL, GGML_OP_ADD }; + + if (ctx->node(idx)->op != GGML_OP_MUL || !ctx->can_fuse(idx, snake_ops, 5)) { + return false; + } + + const ggml_tensor * mul0 = ctx->node(idx + 0); + const ggml_tensor * sin_node = ctx->node(idx + 1); + const ggml_tensor * sqr = ctx->node(idx + 2); + const ggml_tensor * mul1 = ctx->node(idx + 3); + const ggml_tensor * add = ctx->node(idx + 4); + + // x carries the full activation shape, a is the broadcast operand + const ggml_tensor * x = ggml_are_same_shape(mul0, mul0->src[0]) ? mul0->src[0] : mul0->src[1]; + const ggml_tensor * a = (x == mul0->src[0]) ? mul0->src[1] : mul0->src[0]; + + // mul1 reads sqr and inv_b in either operand order + const ggml_tensor * inv_b = (mul1->src[0] == sqr) ? mul1->src[1] : mul1->src[0]; + + // closure check: the trailing add reads the same x as the leading mul + const ggml_tensor * x_in_add = (add->src[0] == mul1) ? add->src[1] : add->src[0]; + + // x is in the supported whitelist and every chain intermediate shares x's type. + // a and inv_b bind as device const float * in the kernel, so they stay F32. + const bool types_ok = + (x->type == GGML_TYPE_F32 || x->type == GGML_TYPE_F16 || x->type == GGML_TYPE_BF16) && + (a->type == GGML_TYPE_F32) && (inv_b->type == GGML_TYPE_F32) && + (mul0->type == x->type) && (sin_node->type == x->type) && + (sqr->type == x->type) && (mul1->type == x->type) && + (add->type == x->type); + // a / inv_b collapse to [1, C, 1, 1], x and add stay 2D + const bool shape_ok = ggml_are_same_shape(a, inv_b) && a->ne[0] == 1 && a->ne[1] == x->ne[1]; + const bool dim_ok = + (x->ne[2] == 1) && (x->ne[3] == 1) && + (add->ne[2] == 1) && (add->ne[3] == 1) && + (a->ne[2] == 1) && (a->ne[3] == 1) && + (inv_b->ne[2] == 1) && (inv_b->ne[3] == 1); + // kernel reads x[idx] and a[c] / inv_b[c] linearly, so every operand is contiguous + const bool contig_ok = + ggml_is_contiguous(x) && ggml_is_contiguous(add) && + ggml_is_contiguous(a) && ggml_is_contiguous(inv_b); + + return types_ok && shape_ok && dim_ok && contig_ok && x_in_add == x; +} + int ggml_metal_op_bin(ggml_metal_op_t ctx, int idx) { + if (ctx->use_fusion && ggml_metal_op_can_fuse_snake(ctx, idx)) { + return ggml_metal_op_snake_fused(ctx, idx); + } + ggml_tensor * op = ctx->node(idx); ggml_metal_library_t lib = ctx->lib; @@ -3984,6 +4035,55 @@ int ggml_metal_op_col2im_1d(ggml_metal_op_t ctx, int idx) { return 1; } +// Dispatch the fused snake kernel from the matched mul -> sin -> sqr -> mul -> add chain. +// idx points at the leading mul. The caller has validated the chain. +int ggml_metal_op_snake_fused(ggml_metal_op_t ctx, int idx) { + ggml_metal_library_t lib = ctx->lib; + ggml_metal_encoder_t enc = ctx->enc; + + const ggml_tensor * mul0 = ctx->node(idx + 0); + const ggml_tensor * sqr = ctx->node(idx + 2); + const ggml_tensor * mul1 = ctx->node(idx + 3); + ggml_tensor * add = ctx->node(idx + 4); + + const ggml_tensor * x = ggml_are_same_shape(mul0, mul0->src[0]) ? mul0->src[0] : mul0->src[1]; + const ggml_tensor * a = (x == mul0->src[0]) ? mul0->src[1] : mul0->src[0]; + const ggml_tensor * inv_b = (mul1->src[0] == sqr) ? mul1->src[1] : mul1->src[0]; + + const int T = (int) x->ne[0]; + const int C = (int) x->ne[1]; + const int total = T * C; + + // the encode loop pre-checked the leading mul only, check the rest of the chain + for (int i = 1; i < 5; ++i) { + if (!ggml_metal_op_concurrency_check(ctx, ctx->node(idx + i))) { + ggml_metal_op_concurrency_reset(ctx); + + break; + } + } + + auto pipeline = ggml_metal_library_get_pipeline_snake(lib, x->type); + + ggml_metal_kargs_snake args = { + /*.T =*/ T, + /*.C =*/ C, + }; + + ggml_metal_encoder_set_pipeline(enc, pipeline); + ggml_metal_encoder_set_bytes (enc, &args, sizeof(args), 0); + ggml_metal_encoder_set_buffer (enc, ggml_metal_get_buffer_id(x), 1); + ggml_metal_encoder_set_buffer (enc, ggml_metal_get_buffer_id(a), 2); + ggml_metal_encoder_set_buffer (enc, ggml_metal_get_buffer_id(inv_b), 3); + ggml_metal_encoder_set_buffer (enc, ggml_metal_get_buffer_id(add), 4); + + const int nth = 256; + const int ntg = (total + nth - 1) / nth; + ggml_metal_encoder_dispatch_threadgroups(enc, ntg, 1, 1, nth, 1, 1); + + return 5; +} + int ggml_metal_op_conv_transpose_2d(ggml_metal_op_t ctx, int idx) { ggml_tensor * op = ctx->node(idx); diff --git a/ggml/src/ggml-metal/ggml-metal-ops.h b/ggml/src/ggml-metal/ggml-metal-ops.h index 0bebd836a..89a6ad82f 100644 --- a/ggml/src/ggml-metal/ggml-metal-ops.h +++ b/ggml/src/ggml-metal/ggml-metal-ops.h @@ -80,6 +80,7 @@ int ggml_metal_op_conv_3d (ggml_metal_op_t ctx, int idx); int ggml_metal_op_conv_transpose_1d (ggml_metal_op_t ctx, int idx); int ggml_metal_op_conv_transpose_2d (ggml_metal_op_t ctx, int idx); int ggml_metal_op_col2im_1d (ggml_metal_op_t ctx, int idx); +int ggml_metal_op_snake_fused (ggml_metal_op_t ctx, int idx); int ggml_metal_op_upscale (ggml_metal_op_t ctx, int idx); int ggml_metal_op_pad (ggml_metal_op_t ctx, int idx); int ggml_metal_op_pad_reflect_1d (ggml_metal_op_t ctx, int idx); diff --git a/ggml/src/ggml-metal/ggml-metal.metal b/ggml/src/ggml-metal/ggml-metal.metal index 38919ebad..969fddfa5 100644 --- a/ggml/src/ggml-metal/ggml-metal.metal +++ b/ggml/src/ggml-metal/ggml-metal.metal @@ -5406,6 +5406,35 @@ template [[host_name("kernel_col2im_1d_bf16")]] kernel void kernel_col2im_1d +kernel void kernel_snake( + constant ggml_metal_kargs_snake & args, + device const T * x, + device const float * a, + device const float * inv_b, + device T * dst, + uint tgpig [[threadgroup_position_in_grid]], + uint tpitg [[thread_position_in_threadgroup]], + uint ntg [[threads_per_threadgroup]]) { + + const int idx = tgpig * ntg + tpitg; + if (idx >= args.T * args.C) { + return; + } + + const int c = idx / args.T; // x is [T, C], a / inv_b collapse to [1, C] + const float xi = float(x[idx]); + const float si = sin(a[c] * xi); + dst[idx] = T(xi + si * si * inv_b[c]); +} + +template [[host_name("kernel_snake_f32")]] kernel void kernel_snake(constant ggml_metal_kargs_snake &, device const float *, device const float *, device const float *, device float *, uint, uint, uint); +template [[host_name("kernel_snake_f16")]] kernel void kernel_snake(constant ggml_metal_kargs_snake &, device const half *, device const float *, device const float *, device half *, uint, uint, uint); +#if defined(GGML_METAL_HAS_BF16) +template [[host_name("kernel_snake_bf16")]] kernel void kernel_snake(constant ggml_metal_kargs_snake &, device const bfloat *, device const float *, device const float *, device bfloat *, uint, uint, uint); +#endif + + typedef void (conv_transpose_2d_t)( constant ggml_metal_kargs_conv_transpose_2d & args, device const float * src0, diff --git a/ggml/src/ggml-vulkan/ggml-vulkan.cpp b/ggml/src/ggml-vulkan/ggml-vulkan.cpp index e9c423af6..def0e35d6 100644 --- a/ggml/src/ggml-vulkan/ggml-vulkan.cpp +++ b/ggml/src/ggml-vulkan/ggml-vulkan.cpp @@ -967,6 +967,7 @@ struct vk_device_struct { vk_pipeline pipeline_col2im_1d_f32; vk_pipeline pipeline_col2im_1d_f16; vk_pipeline pipeline_col2im_1d_bf16; + vk_pipeline pipeline_out_prod_f32; vk_pipeline pipeline_snake_f32; vk_pipeline pipeline_snake_f16; vk_pipeline pipeline_snake_bf16; @@ -5485,6 +5486,8 @@ static void ggml_vk_load_shaders(vk_device& device, vk_pipeline requested) { ggml_vk_create_pipeline(device, device->pipeline_col2im_1d_f16, "col2im_1d_f16", col2im_1d_f16_len, col2im_1d_f16_data, "main", 2, sizeof(vk_op_col2im_1d_push_constants), {256, 1, 1}, {}, 1, true); ggml_vk_create_pipeline(device, device->pipeline_col2im_1d_bf16, "col2im_1d_bf16", col2im_1d_bf16_len, col2im_1d_bf16_data, "main", 2, sizeof(vk_op_col2im_1d_push_constants), {256, 1, 1}, {}, 1, true); + ggml_vk_create_pipeline(device, device->pipeline_out_prod_f32, "out_prod_f32", out_prod_f32_len, out_prod_f32_data, "main", 3, sizeof(vk_op_binary_push_constants), {256, 1, 1}, {}, 1); + ggml_vk_create_pipeline(device, device->pipeline_snake_f32, "snake_f32", snake_f32_len, snake_f32_data, "main", 4, sizeof(vk_op_snake_push_constants), {256, 1, 1}, {}, 1); ggml_vk_create_pipeline(device, device->pipeline_snake_f16, "snake_f16", snake_f16_len, snake_f16_data, "main", 4, sizeof(vk_op_snake_push_constants), {256, 1, 1}, {}, 1); ggml_vk_create_pipeline(device, device->pipeline_snake_bf16, "snake_bf16", snake_bf16_len, snake_bf16_data, "main", 4, sizeof(vk_op_snake_push_constants), {256, 1, 1}, {}, 1); @@ -10778,6 +10781,11 @@ static vk_pipeline ggml_vk_op_get_pipeline(ggml_backend_vk_context * ctx, const return ctx->device->pipeline_add_id_f32; } return nullptr; + case GGML_OP_OUT_PROD: + if (src0->type == GGML_TYPE_F32 && src1->type == GGML_TYPE_F32 && dst->type == GGML_TYPE_F32) { + return ctx->device->pipeline_out_prod_f32; + } + return nullptr; case GGML_OP_CONCAT: { if (src0->type != src1->type || src0->type != dst->type) { return nullptr; @@ -11734,6 +11742,7 @@ static void ggml_vk_op_f32(ggml_backend_vk_context * ctx, vk_context& subctx, co case GGML_OP_DIV: case GGML_OP_MUL: case GGML_OP_ADD1: + case GGML_OP_OUT_PROD: case GGML_OP_ARANGE: case GGML_OP_FILL: case GGML_OP_SCALE: @@ -12047,6 +12056,24 @@ static void ggml_vk_add(ggml_backend_vk_context * ctx, vk_context& subctx, const }); } +static void ggml_vk_out_prod(ggml_backend_vk_context * ctx, vk_context& subctx, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst) { + const uint32_t src0_type_size = ggml_type_size(src0->type); + const uint32_t src1_type_size = ggml_type_size(src1->type); + const uint32_t dst_type_size = ggml_type_size(dst->type); + + ggml_vk_op_f32(ctx, subctx, src0, src1, nullptr, nullptr, dst, GGML_OP_OUT_PROD, { + (uint32_t)ggml_nelements(dst), + (uint32_t)src0->ne[0], (uint32_t)src0->ne[1], (uint32_t)src0->ne[2],(uint32_t)src0->ne[3], + (uint32_t)src0->nb[0] / src0_type_size, (uint32_t)src0->nb[1] / src0_type_size, (uint32_t)src0->nb[2] / src0_type_size, (uint32_t)src0->nb[3] / src0_type_size, + (uint32_t)src1->ne[0], (uint32_t)src1->ne[1], (uint32_t)src1->ne[2],(uint32_t)src1->ne[3], + (uint32_t)src1->nb[0] / src1_type_size, (uint32_t)src1->nb[1] / src1_type_size, (uint32_t)src1->nb[2] / src1_type_size, (uint32_t)src1->nb[3] / src1_type_size, + (uint32_t) dst->ne[0], (uint32_t) dst->ne[1], (uint32_t) dst->ne[2],(uint32_t) dst->ne[3], + (uint32_t) dst->nb[0] / dst_type_size, (uint32_t) dst->nb[1] / dst_type_size, (uint32_t) dst->nb[2] / dst_type_size, (uint32_t) dst->nb[3] / dst_type_size, + 0, + 0.0f, 0.0f, 0, + }); +} + static void ggml_vk_sub(ggml_backend_vk_context * ctx, vk_context& subctx, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst) { const uint32_t src0_type_size = ggml_type_size(src0->type); const uint32_t src1_type_size = ggml_type_size(src1->type); @@ -14834,6 +14861,9 @@ static bool ggml_vk_build_graph(ggml_backend_vk_context * ctx, ggml_cgraph * cgr ggml_vk_add(ctx, compute_ctx, src0, src1, node); } break; + case GGML_OP_OUT_PROD: + ggml_vk_out_prod(ctx, compute_ctx, src0, src1, node); + break; case GGML_OP_SUB: ggml_vk_sub(ctx, compute_ctx, src0, src1, node); @@ -17688,6 +17718,10 @@ static bool ggml_backend_vk_device_supports_op(ggml_backend_dev_t dev, const ggm case GGML_OP_OPT_STEP_ADAMW: case GGML_OP_OPT_STEP_SGD: return ggml_is_contiguous(op->src[0]) && op->src[0]->type == GGML_TYPE_F32; + case GGML_OP_OUT_PROD: + return ggml_is_contiguous(op->src[0]) && op->src[0]->type == GGML_TYPE_F32 + && ggml_is_contiguous(op->src[1]) && op->src[1]->type == GGML_TYPE_F32 + && op->type == GGML_TYPE_F32; case GGML_OP_LOG: case GGML_OP_TRI: case GGML_OP_DIAG: diff --git a/ggml/src/ggml-vulkan/vulkan-shaders/out_prod.comp b/ggml/src/ggml-vulkan/vulkan-shaders/out_prod.comp new file mode 100644 index 000000000..197316996 --- /dev/null +++ b/ggml/src/ggml-vulkan/vulkan-shaders/out_prod.comp @@ -0,0 +1,59 @@ +#version 450 + +#extension GL_EXT_shader_16bit_storage : require + +layout (push_constant) uniform parameter +{ + uint ne; + uint ne00; uint ne01; uint ne02; uint ne03; uint nb00; uint nb01; uint nb02; uint nb03; + uint ne10; uint ne11; uint ne12; uint ne13; uint nb10; uint nb11; uint nb12; uint nb13; + uint ne20; uint ne21; uint ne22; uint ne23; uint nb20; uint nb21; uint nb22; uint nb23; + uint misalign_offsets; + float param1; float param2; int param3; +} p; + +layout (binding = 0) readonly buffer A {float data_a[];}; +layout (binding = 1) readonly buffer B {float data_b[];}; +layout (binding = 2) writeonly buffer D {float data_d[];}; + +uint get_idx() { + return gl_GlobalInvocationID.z * 262144 + gl_GlobalInvocationID.y * 512 + gl_GlobalInvocationID.x; +} + +uint get_aoffset() { return p.misalign_offsets >> 16; } +uint get_boffset() { return (p.misalign_offsets >> 8) & 0xFF; } +uint get_doffset() { return p.misalign_offsets & 0xFF; } + +layout(local_size_x = 256, local_size_y = 1, local_size_z = 1) in; + +void main() { + uint idx = get_idx(); + if (idx >= p.ne) { + return; + } + + uint tmp = idx; + uint i0 = tmp % p.ne20; tmp /= p.ne20; + uint i1 = tmp % p.ne21; tmp /= p.ne21; + uint i2 = tmp % p.ne22; tmp /= p.ne22; + uint i3 = tmp; + + uint a_i0 = i0 % p.ne00; + uint a_i2 = i2 / (p.ne22 / p.ne02); + uint a_i3 = i3 / (p.ne23 / p.ne03); + + uint b_i0 = i1 % p.ne10; + uint b_i2 = i2; + uint b_i3 = i3; + + float sum = 0.0f; + uint K = p.ne01; + for (uint k = 0; k < K; k++) { + uint aoff = get_aoffset() + a_i3*p.nb03 + a_i2*p.nb02 + k*p.nb01 + a_i0*p.nb00; + uint boff = get_boffset() + b_i3*p.nb13 + b_i2*p.nb12 + k*p.nb11 + b_i0*p.nb10; + sum += data_a[aoff] * data_b[boff]; + } + + uint doff = get_doffset() + i3*p.nb23 + i2*p.nb22 + i1*p.nb21 + i0*p.nb20; + data_d[doff] = 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 c816dcc90..a9893a48c 100644 --- a/ggml/src/ggml-vulkan/vulkan-shaders/vulkan-shaders-gen.cpp +++ b/ggml/src/ggml-vulkan/vulkan-shaders/vulkan-shaders-gen.cpp @@ -1062,6 +1062,8 @@ void process_shaders() { } } + string_to_spv("out_prod_f32", "out_prod.comp", {}); + string_to_spv("timestep_embedding_f32", "timestep_embedding.comp", merge_maps(base_dict, {{"A_TYPE", "float"}, {"D_TYPE", "float"}})); string_to_spv("conv_transpose_1d_f32", "conv_transpose_1d.comp", {{"A_TYPE", "float"}, {"B_TYPE", "float"}, {"D_TYPE", "float"}}); diff --git a/src/llama-quant.cpp b/src/llama-quant.cpp index d30f8d91a..60caccadd 100644 --- a/src/llama-quant.cpp +++ b/src/llama-quant.cpp @@ -675,7 +675,7 @@ static ggml_type llama_tensor_get_type(quantize_state_impl & qs, const llama_mod ggml_type new_type = default_type; // get more optimal quantization type based on the tensor shape, layer, etc. - if (!params->pure && ggml_is_quantized(default_type)) { + if (ggml_is_quantized(default_type)) { // if the user provided tensor types - use those bool manual = false; if (!qs.tensor_type_patterns.empty()) { @@ -694,7 +694,7 @@ static ggml_type llama_tensor_get_type(quantize_state_impl & qs, const llama_mod } // if not manual - use the standard logic for choosing the quantization type based on the selected mixture - if (!manual) { + if (!manual && !params->pure) { new_type = llama_tensor_get_type_impl(qs, new_type, tensor, params->ftype, tm.category); } diff --git a/tools/server/server-http.cpp b/tools/server/server-http.cpp index 24a38452a..783b01b82 100644 --- a/tools/server/server-http.cpp +++ b/tools/server/server-http.cpp @@ -283,9 +283,9 @@ bool server_http_context::init(const common_params & params) { } else if (params.cors_origins == "localhost") { // special case: only reflect the Origin header if it is a localhost origin std::string origin = req.get_header_value("Origin"); - if (origin_is_localhost(origin)) { + if (!origin.empty() && origin_is_localhost(origin)) { res.set_header("Access-Control-Allow-Origin", origin); - } else { + } else if (!origin.empty()) { SRV_WRN("(CORS) skip non-localhost origin: %s\n", origin.c_str()); } } else { diff --git a/tools/server/server-tools.cpp b/tools/server/server-tools.cpp index a8216d7db..9eb57abae 100644 --- a/tools/server/server-tools.cpp +++ b/tools/server/server-tools.cpp @@ -289,7 +289,7 @@ struct server_tool_read_file : server_tool { {"function", { {"name", name}, {"description", "Read the contents of a file. Optionally specify a 1-based line range. " - "If append_loc is true, each line is prefixed with its line number (e.g. \"1\u2192 ...\")."}, + "If append_loc is true, each line is prefixed with its line number (e.g. \"1\u2192...\")."}, {"parameters", { {"type", "object"}, {"properties", { @@ -339,7 +339,7 @@ struct server_tool_read_file : server_tool { std::string out_line; if (append_loc) { - out_line = std::to_string(lineno) + "\u2192 " + line + "\n"; + out_line = std::to_string(lineno) + "\u2192" + line + "\n"; } else { out_line = line + "\n"; } diff --git a/tools/server/tests/unit/test_speculative.py b/tools/server/tests/unit/test_speculative.py index 84cd77e6f..c6568479c 100644 --- a/tools/server/tests/unit/test_speculative.py +++ b/tools/server/tests/unit/test_speculative.py @@ -12,8 +12,9 @@ def create_server(): server = ServerPreset.stories15m_moe() # set default values server.model_draft = download_file(MODEL_DRAFT_FILE_URL) - server.draft_min = 4 - server.draft_max = 8 + server.spec_type = "draft-simple" + server.spec_draft_n_min = 4 + server.spec_draft_n_max = 8 server.fa = "off" @@ -25,6 +26,7 @@ def fixture_create_server(): def test_with_and_without_draft(): global server server.model_draft = None # disable draft model + server.spec_type = None server.start() res = server.make_request("POST", "/completion", data={ "prompt": "I believe the meaning of life is", @@ -46,6 +48,7 @@ def test_with_and_without_draft(): "n_predict": 16, }) assert res.status_code == 200 + assert res.body["timings"]["draft_n"] > 0 content_draft = res.body["content"] assert content_no_draft == content_draft @@ -63,8 +66,8 @@ def test_different_draft_min_draft_max(): last_content = None for draft_min, draft_max in test_values: server.stop() - server.draft_min = draft_min - server.draft_max = draft_max + server.spec_draft_n_min = draft_min + server.spec_draft_n_max = draft_max server.start() res = server.make_request("POST", "/completion", data={ "prompt": "I believe the meaning of life is", diff --git a/tools/server/tests/utils.py b/tools/server/tests/utils.py index f4f0e61e6..5d5c873ac 100644 --- a/tools/server/tests/utils.py +++ b/tools/server/tests/utils.py @@ -95,6 +95,7 @@ class ServerProcess: no_models_autoload: bool | None = None lora_files: List[str] | None = None enable_ctx_shift: int | None = False + spec_type: str | None = None spec_draft_n_min: int | None = None spec_draft_n_max: int | None = None no_ui: bool | None = None @@ -226,6 +227,8 @@ class ServerProcess: server_args.extend(["--lora", lora_file]) if self.enable_ctx_shift: server_args.append("--context-shift") + if self.spec_type: + server_args.extend(["--spec-type", self.spec_type]) if self.api_key: server_args.extend(["--api-key", self.api_key]) if self.spec_draft_n_max: diff --git a/tools/ui/eslint.config.js b/tools/ui/eslint.config.js index b90376b6c..54679a05b 100644 --- a/tools/ui/eslint.config.js +++ b/tools/ui/eslint.config.js @@ -29,6 +29,13 @@ export default ts.config( // This app uses hash-based routing (#/) where resolve() from $app/paths does not apply 'svelte/no-navigation-without-resolve': 'off', + // Snippet bodies often ignore one or more of the parent's params + // (e.g. `{#snippet children(_meta, ctx)}` when only ctx is read). + '@typescript-eslint/no-unused-vars': [ + 'error', + { argsIgnorePattern: '^_', varsIgnorePattern: '^_' } + ], + // Enforce empty line at end of file 'eol-last': 'error' } diff --git a/tools/ui/src/app.css b/tools/ui/src/app.css index 8c4056477..f9b544beb 100644 --- a/tools/ui/src/app.css +++ b/tools/ui/src/app.css @@ -193,6 +193,33 @@ -ms-overflow-style: none; scrollbar-width: none; } + + .shimmer-text { + background: linear-gradient( + 90deg, + var(--muted-foreground), + var(--foreground), + var(--muted-foreground) + ); + background-size: 200% 100%; + background-clip: text; + -webkit-background-clip: text; + -webkit-text-fill-color: transparent; + font-weight: 500; + animation: shimmer 1s linear infinite; + } + + @keyframes shimmer { + to { + background-position: -200% 0; + } + } + + @media (prefers-reduced-motion: reduce) { + .shimmer-text { + animation: none; + } + } } .mermaidTooltip { diff --git a/tools/ui/src/lib/components/app/actions/ActionIconCopyToClipboard.svelte b/tools/ui/src/lib/components/app/actions/ActionIconCopyToClipboard.svelte index 999f0cba9..9b7b370ad 100644 --- a/tools/ui/src/lib/components/app/actions/ActionIconCopyToClipboard.svelte +++ b/tools/ui/src/lib/components/app/actions/ActionIconCopyToClipboard.svelte @@ -1,4 +1,5 @@ + +{#if isRouter} + { + const status = modelsStore.getModelStatus(modelId); + + if (status !== ServerModelStatus.LOADED) { + pendingModel = modelId; + + try { + await modelsStore.loadModel(modelId); + } finally { + pendingModel = null; + } + } + + onRegenerate(modelName); + return true; + }} + /> +{:else} + +{/if} diff --git a/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageAssistant/ChatMessageAssistantProcessingInfo.svelte b/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageAssistant/ChatMessageAssistantProcessingInfo.svelte new file mode 100644 index 000000000..356512ecb --- /dev/null +++ b/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageAssistant/ChatMessageAssistantProcessingInfo.svelte @@ -0,0 +1,25 @@ + + +
+
+ + {modelLoadingText ?? + processingState.getPromptProgressText() ?? + processingState.getProcessingMessage() ?? + 'Processing...'} + +
+
diff --git a/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageAssistant/ChatMessageAssistantRawOutput.svelte b/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageAssistant/ChatMessageAssistantRawOutput.svelte new file mode 100644 index 000000000..30ce16be9 --- /dev/null +++ b/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageAssistant/ChatMessageAssistantRawOutput.svelte @@ -0,0 +1,33 @@ + + +
{rawOutputContent || ''}
+ + diff --git a/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageAssistant/ChatMessageAssistantStatistics.svelte b/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageAssistant/ChatMessageAssistantStatistics.svelte new file mode 100644 index 000000000..4cc4080c3 --- /dev/null +++ b/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageAssistant/ChatMessageAssistantStatistics.svelte @@ -0,0 +1,40 @@ + + +{#if showMessageStats && message.timings && message.timings.predicted_n && message.timings.predicted_ms} + {@const agentic = message.timings.agentic} + +{:else if isLoading && showMessageStats} + {@const liveStats = processingState.getLiveProcessingStats()} + {@const genStats = processingState.getLiveGenerationStats()} + + {#if genStats} + + {/if} +{/if} diff --git a/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/ChatMessageToolCallBlock.svelte b/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/ChatMessageToolCallBlock.svelte new file mode 100644 index 000000000..b1daedfc8 --- /dev/null +++ b/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/ChatMessageToolCallBlock.svelte @@ -0,0 +1,66 @@ + + +{#if isSearchCall} + +{:else if section.toolName === BuiltInTool.GET_DATETIME} + +{:else if section.toolName === BuiltInTool.READ_FILE} + +{:else if section.toolName === BuiltInTool.EDIT_FILE} + +{:else if section.toolName === BuiltInTool.WRITE_FILE} + +{:else if section.toolName === BuiltInTool.EXEC_SHELL_COMMAND} + +{:else if section.toolName === BuiltInTool.FILE_GLOB_SEARCH} + +{:else if section.toolName === BuiltInTool.GREP_SEARCH} + +{:else if section.toolName === BuiltInTool.RUN_JAVASCRIPT} + +{:else} + +{/if} diff --git a/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/ChatMessageToolCallBlockDefault.svelte b/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/ChatMessageToolCallBlockDefault.svelte new file mode 100644 index 000000000..acf2de12a --- /dev/null +++ b/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/ChatMessageToolCallBlockDefault.svelte @@ -0,0 +1,124 @@ + + + + {#snippet children(_meta, ctx)} + {#if ctx.isStreamingCall} +
+ Input + {#if ctx.isStreaming} + + {/if} +
+ {#if section.toolArgs} + + {:else if ctx.isStreaming} +
+ Receiving arguments... +
+ {:else} +
+ Response was truncated +
+ {/if} + {:else} + {@const showInput = Boolean(section.toolArgs)} + {#if showInput} +
+ Input +
+ + {/if} +
+ Output + {#if ctx.isPending} + + {/if} +
+ {#if ctx.isPending} +
+ Waiting for result... +
+ {:else if section.toolResult} + {#if outputKind === ToolResultKind.JSON} + + {:else if outputKind === ToolResultKind.MARKDOWN} + + {:else} +
+ {#each parsedLines as line, i (i)} +
+ {line.text} +
+ {#if line.image} + {line.image.name} + {/if} + {/each} +
+ {/if} + {:else} +
No output
+ {/if} + {/if} + {/snippet} + diff --git a/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/ChatMessageToolCallBlockEditFile.svelte b/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/ChatMessageToolCallBlockEditFile.svelte new file mode 100644 index 000000000..6f30060f5 --- /dev/null +++ b/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/ChatMessageToolCallBlockEditFile.svelte @@ -0,0 +1,166 @@ + + + + {#snippet titleSnippet()} + Edit file + {editFileMeta?.filePath} + {#if editFileMeta?.errorMessage} + (failed) + {/if} + {/snippet} + + {#snippet children(meta, _ctx)} + {#if meta?.errorMessage} +
+ + {meta.errorMessage} +
+ {:else if meta && meta.edits.length > 0} + {#each editDiffs as diffLines, ei (ei)} +
+
+ Edit {ei + 1} of {meta.edits.length} +
+
+
+ {#each diffLines as line, li (li)} +
+ {line.oldLine ?? ''} + {prefixFor(line.kind)} + {line.newLine ?? ''} + {line.text || ' '} +
+ {/each} +
+
+
+ {/each} +
+ {#if meta.resultMessage} + {meta.resultMessage}{meta.editsApplied != null ? RESULT_STAT_SEPARATOR : ''}{/if} + {#if meta.editsApplied != null} + {meta.editsApplied} + {meta.editsApplied === 1 ? 'edit' : 'edits'} applied + {/if} +
+ {:else} +
No edits
+ {/if} + {/snippet} + + + diff --git a/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/ChatMessageToolCallBlockExecShellCommand.svelte b/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/ChatMessageToolCallBlockExecShellCommand.svelte new file mode 100644 index 000000000..5de801d39 --- /dev/null +++ b/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/ChatMessageToolCallBlockExecShellCommand.svelte @@ -0,0 +1,293 @@ + + +{#snippet execShellTitle()} + {#if highlightedCommandHtml} + {@html highlightedCommandHtml} + {:else} + {execShellMeta?.command} + {/if} +{/snippet} + + + {#snippet titleSnippet()} + {@render execShellTitle()} + {/snippet} + + {#snippet children(_meta, ctx)} + {#if ctx.isPending} +
+ + Running... +
+ {:else if execShellError} +
+ + {execShellError} +
+ {:else if section.toolResult} +
+ {#each outputLines as line, i (i)} +
{line.text}
+ {#if line.image} + {line.image.name} + {/if} + {/each} + + {#if isExitCodeFinalLine && execShellExitStatus} +
+ {#if execShellExitStatus.timedOut} + + timed out + · + exit {execShellExitStatus.code} + {:else if execShellExitStatus.code === 0} + + exit 0 + {:else} + + exit {execShellExitStatus.code} + {/if} +
+ {/if} +
+ {/if} + {/snippet} + + + diff --git a/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/ChatMessageToolCallBlockFileGlobSearch.svelte b/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/ChatMessageToolCallBlockFileGlobSearch.svelte new file mode 100644 index 000000000..ad082039f --- /dev/null +++ b/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/ChatMessageToolCallBlockFileGlobSearch.svelte @@ -0,0 +1,61 @@ + + + + {#snippet titleSnippet()} + {#if fileGlobMeta} + {fileGlobMeta.include === '**' ? 'List files' : 'Search files'}  + {#if fileGlobMeta.include !== '**'} + {fileGlobMeta.include} + {/if} +  in  + {fileGlobMeta.path} + {/if} + {/snippet} + + {#snippet children(meta, ctx)} + {#if ctx.isPending} +
+ Searching... +
+ {:else if meta?.errorMessage} +
+ + {meta.errorMessage} +
+ {:else if meta && meta.matches.length > 0} +
+ {#each meta.matches as match, i (i)} +
{match}
+ {/each} +
+
+ Total matches: {meta.totalMatches ?? meta.matches.length} +
+ {:else} +
No matches
+
+ Total matches: {meta?.totalMatches ?? 0} +
+ {/if} + {/snippet} + diff --git a/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/ChatMessageToolCallBlockGetDatetime.svelte b/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/ChatMessageToolCallBlockGetDatetime.svelte new file mode 100644 index 000000000..e0c701dea --- /dev/null +++ b/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/ChatMessageToolCallBlockGetDatetime.svelte @@ -0,0 +1,57 @@ + + +
+ + {#if showSpinner} + Current time + + {:else if dateMeta.errorMessage} + Current time  + - {dateMeta.errorMessage} + {:else if dateMeta.dateString} + Current time is  + {dateMeta.dateString} + {:else} + Current time + {/if} +
diff --git a/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/ChatMessageToolCallBlockGrepSearch.svelte b/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/ChatMessageToolCallBlockGrepSearch.svelte new file mode 100644 index 000000000..afb06fef7 --- /dev/null +++ b/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/ChatMessageToolCallBlockGrepSearch.svelte @@ -0,0 +1,67 @@ + + + + {#snippet titleSnippet()} + {#if grepMeta} + Search for  + {grepMeta.pattern} +  in  + {grepMeta.path} + {/if} + {/snippet} + + {#snippet children(meta, ctx)} + {#if ctx.isPending} +
+ Searching... +
+ {:else if meta?.errorMessage} +
+ + {meta.errorMessage} +
+ {:else if meta && meta.matches.length > 0} +
+ {#each meta.matches as match, mi (mi)} +
+ {match.file} + {#if meta.showLineNumbers && match.line != null} + :{match.line} + {/if} + : + {match.content} +
+ {/each} +
+
+ Total matches: {meta.totalMatches ?? meta.matches.length} + {#if meta.showLineNumbers} +  (with line numbers) + {/if} +
+ {:else} +
No matches
+
+ Total matches: {meta?.totalMatches ?? 0} +
+ {/if} + {/snippet} + diff --git a/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/ChatMessageToolCallBlockReadFile.svelte b/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/ChatMessageToolCallBlockReadFile.svelte new file mode 100644 index 000000000..a99ff9cee --- /dev/null +++ b/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/ChatMessageToolCallBlockReadFile.svelte @@ -0,0 +1,44 @@ + + + + {#snippet titleSnippet()} + Read file + {readFileMeta?.fileName} + {#if readFileMeta?.lineRange} +  (lines {readFileMeta.lineRange.start}-{readFileMeta.lineRange.end}) + {/if} + {/snippet} + + {#snippet children(_meta, _ctx)} + {#if section.toolResult} + + {:else} +
+ Waiting for file content... +
+ {/if} + {/snippet} + diff --git a/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/ChatMessageToolCallBlockRunJavascript.svelte b/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/ChatMessageToolCallBlockRunJavascript.svelte new file mode 100644 index 000000000..707d83d73 --- /dev/null +++ b/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/ChatMessageToolCallBlockRunJavascript.svelte @@ -0,0 +1,69 @@ + + + + {#snippet children(meta, ctx)} + {#if ctx.isPending} +
Running...
+ {:else if meta?.errorMessage} +
+ + {meta.errorMessage} +
+
+ +
+ {:else if meta} + +
+ + Console + {#if meta.timeoutMs != null} + · timeout {meta.timeoutMs} ms + {/if} +
+ {#if section.toolResult} +
+ +
+ {:else} +
No output
+ {/if} + {/if} + {/snippet} + diff --git a/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/ChatMessageToolCallBlockSearchResults.svelte b/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/ChatMessageToolCallBlockSearchResults.svelte new file mode 100644 index 000000000..e4b4adf15 --- /dev/null +++ b/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/ChatMessageToolCallBlockSearchResults.svelte @@ -0,0 +1,167 @@ + + +{#snippet pill(result: SearchResult)} + {@const faviconUrl = faviconForUrl(result.url)} + {@const safeUrl = sanitizeExternalUrl(result.url)} + {@const showHoverCard = safeUrl !== null && hasDetails(result)} + {#if safeUrl} + + + {#if faviconUrl} + + {:else} + + {/if} + {result.title} + + {#if showHoverCard} + {@const publishDate = formatPublishDate(result.published)} + {@const host = hostFor(safeUrl)} + +
+ {result.title} + {#if publishDate || result.author} +
+ {#if publishDate} + {publishDate} + {/if} + {#if publishDate && result.author} + · + {/if} + {#if result.author} + {result.author} + {/if} +
+ {/if} + {#if result.highlights} +

+ {result.highlights} +

+ {/if} + {#if host} +
{host}
+ {/if} +
+ + {/if} + + {/if} +{/snippet} + + + {#if results.length > 0} +
+ {#each results as result (result.url)} + {@render pill(result)} + {/each} +
+ {:else if showSpinner} +
+ + Searching... +
+ {:else} +
No results
+ {/if} + diff --git a/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/ChatMessageToolCallBlockWriteFile.svelte b/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/ChatMessageToolCallBlockWriteFile.svelte new file mode 100644 index 000000000..eda067662 --- /dev/null +++ b/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/ChatMessageToolCallBlockWriteFile.svelte @@ -0,0 +1,55 @@ + + + + {#snippet titleSnippet()} + Write file + {writeFileMeta?.filePath} + {#if writeFileMeta?.errorMessage} + (failed) + {/if} + {/snippet} + + {#snippet children(meta, ctx)} + {#if meta?.errorMessage} +
+ + {meta.errorMessage} +
+ {:else if meta} + +
+ {#if meta.resultMessage} + {meta.resultMessage}{meta.bytesWritten != null ? RESULT_STAT_SEPARATOR : ''}{/if} + {#if meta.bytesWritten != null} + {meta.bytesWritten} + bytes + {/if} +
+ {/if} + {/snippet} + diff --git a/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/ToolCallBlock.svelte b/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/ToolCallBlock.svelte new file mode 100644 index 000000000..a17a74e16 --- /dev/null +++ b/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/ToolCallBlock.svelte @@ -0,0 +1,129 @@ + + + + {@render children(meta, { + isStreaming, + isPending, + isStreamingCall, + isCodeStreaming + })} + diff --git a/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/parsers/_shared.ts b/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/parsers/_shared.ts new file mode 100644 index 000000000..6114f17b5 --- /dev/null +++ b/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/parsers/_shared.ts @@ -0,0 +1,49 @@ +// Helpers shared by the per-tool meta parsers under +// `src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/parsers/`. +// Each tool needs the same first three steps (tool-name check, +// args-present check, JSON parse) - keeping them here lets each parser +// stay focused on its own format quirks. + +import { BuiltInTool } from '$lib/enums'; +import { parsePartialJsonArgs } from '$lib/utils/parse-partial-json-args'; +import type { AgenticSection } from '$lib/utils/agentic'; + +/** + * Strict (final-state) JSON parser for a tool-args blob. Mirrors the + * behaviour the per-tool components used before extraction: an + * invalid JSON blob, a JSON array, or a JSON primitive all map to + * `null` so callers don't have to guard against surprise shapes. + */ +function parseFinalToolArgs(blob: string): Record | null { + try { + const parsed: unknown = JSON.parse(blob); + if (parsed && typeof parsed === 'object' && !Array.isArray(parsed)) { + return parsed as Record; + } + return null; + } catch { + return null; + } +} + +/** + * Parse a section's toolArgs against an expected tool name. Returns + * `null` when: + * - the section's toolName doesn't match (component isn't for this + * tool); + * - the section has no args yet (call hasn't started streaming); + * - or the args blob can't be parsed. + * + * Pass `{ partial: true }` for tools that need to render incrementally + * as each token lands (read_file, edit_file, write_file). + */ +export function parseToolArgs( + expected: BuiltInTool, + section: AgenticSection, + options: { partial?: boolean } = {} +): Record | null { + if (section.toolName !== expected || !section.toolArgs) return null; + return options.partial + ? parsePartialJsonArgs(section.toolArgs) + : parseFinalToolArgs(section.toolArgs); +} diff --git a/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/parsers/edit-file.ts b/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/parsers/edit-file.ts new file mode 100644 index 000000000..4bff25bb5 --- /dev/null +++ b/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/parsers/edit-file.ts @@ -0,0 +1,71 @@ +// Meta parser for `edit_file` tool calls. Reads the file path and the +// array of edits from the streamed args (partial JSON for incremental +// rendering), plus the result blob for `result` / `edits_applied` / +// `error` fields. + +import { BuiltInTool } from '$lib/enums'; +import { FILE_PATH_SEPARATOR_REGEX } from '$lib/constants'; +import { tryParseToolResultObject, type AgenticSection } from '$lib/utils'; +import { parseToolArgs } from './_shared'; + +export type EditFileEdit = { + oldText: string; + newText: string; +}; + +export type EditFileMeta = { + fileName: string; + filePath: string; + edits: EditFileEdit[]; + resultMessage?: string; + editsApplied?: number; + errorMessage?: string; +}; + +export function parseEditFileMeta(section: AgenticSection): EditFileMeta | null { + const args = parseToolArgs(BuiltInTool.EDIT_FILE, section, { partial: true }); + if (!args) return null; + + const rawPath = args.path ?? args.file_path ?? args.filePath; + if (typeof rawPath !== 'string' || !rawPath) return null; + + const fileName = rawPath.split(FILE_PATH_SEPARATOR_REGEX).pop() || rawPath; + + // Filter the streamed edits array strictly: each entry must be an + // object with a non-empty `old_text`. Edits without an old_text + // would diff against empty and render as a full re-write. + const rawEdits = Array.isArray(args.edits) ? args.edits : []; + const edits: EditFileEdit[] = []; + for (const e of rawEdits) { + if (!e || typeof e !== 'object' || Array.isArray(e)) continue; + const obj = e as Record; + const oldText = typeof obj.old_text === 'string' ? obj.old_text : ''; + if (!oldText) continue; + const newText = typeof obj.new_text === 'string' ? obj.new_text : ''; + edits.push({ oldText, newText }); + } + + const resultObj = tryParseToolResultObject(section.toolResult); + let resultMessage: string | undefined; + let editsApplied: number | undefined; + let errorMessage: string | undefined; + if (typeof resultObj?.error === 'string') { + errorMessage = resultObj.error; + } else if (resultObj) { + if (typeof resultObj.result === 'string') { + resultMessage = resultObj.result; + } + if (Number.isFinite(Number(resultObj.edits_applied))) { + editsApplied = Number(resultObj.edits_applied); + } + } + + return { + fileName, + filePath: rawPath, + edits, + resultMessage, + editsApplied, + errorMessage + }; +} diff --git a/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/parsers/exec-shell-command.ts b/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/parsers/exec-shell-command.ts new file mode 100644 index 000000000..e8adbd18b --- /dev/null +++ b/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/parsers/exec-shell-command.ts @@ -0,0 +1,23 @@ +// Meta parser for `exec_shell_command` tool calls. Surfaces the +// command text from args `command` / `cmd` / `shell_command` aliases. +// The exit-status and error parsing live in their own utilities +// (`parse-exec-shell-status.ts` / `parse-exec-shell-error.ts`) - this +// file only deals with what's strictly about *calling* the tool, since +// the error / exit status elide from call-section to result-section. + +import { BuiltInTool } from '$lib/enums'; +import type { AgenticSection } from '$lib/utils'; +import { parseToolArgs } from './_shared'; + +export type ExecShellCommandMeta = { + command: string; +}; + +export function parseExecShellCommandMeta(section: AgenticSection): ExecShellCommandMeta | null { + const args = parseToolArgs(BuiltInTool.EXEC_SHELL_COMMAND, section); + if (!args) return null; + + const commandRaw = args.command ?? args.cmd ?? args.shell_command; + if (typeof commandRaw !== 'string' || !commandRaw) return null; + return { command: commandRaw }; +} diff --git a/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/parsers/file-glob-search.ts b/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/parsers/file-glob-search.ts new file mode 100644 index 000000000..1ad92b74c --- /dev/null +++ b/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/parsers/file-glob-search.ts @@ -0,0 +1,58 @@ +// Meta parser for `file_glob_search` tool calls. Reads the path, +// include pattern, and optional exclude from the args (strict parsing) +// and the matches from the result blob. Like grep_search, the result +// parser keeps the original raw-text fallback for MCP servers that +// emit unparseable output. + +import { BuiltInTool } from '$lib/enums'; +import { splitSearchSummaryList, type AgenticSection } from '$lib/utils'; +import { parseToolArgs } from './_shared'; + +export type FileGlobSearchMeta = { + path: string; + include: string; + exclude?: string; + matches: string[]; + totalMatches?: number; + errorMessage?: string; +}; + +export function parseFileGlobSearchMeta(section: AgenticSection): FileGlobSearchMeta | null { + const args = parseToolArgs(BuiltInTool.FILE_GLOB_SEARCH, section); + if (!args) return null; + + const path = typeof args.path === 'string' ? args.path : ''; + const include = typeof args.include === 'string' && args.include ? args.include : '**'; + const exclude = typeof args.exclude === 'string' && args.exclude ? args.exclude : undefined; + if (!path) return null; + + let matches: string[] = []; + let totalMatches: number | undefined; + let errorMessage: string | undefined; + + const toolResultString = section.toolResult; + if (toolResultString) { + try { + const parsed: unknown = JSON.parse(toolResultString); + if (parsed && typeof parsed === 'object' && !Array.isArray(parsed)) { + const obj = parsed as Record; + if (typeof obj.error === 'string') { + errorMessage = obj.error; + } else if (typeof obj.plain_text_response === 'string') { + const split = splitSearchSummaryList(obj.plain_text_response, (total) => { + totalMatches = total; + }); + matches = split.lines; + } + } + } catch { + // See grep-search.ts: same fallback used there. + const split = splitSearchSummaryList(toolResultString, (total) => { + totalMatches = total; + }); + matches = split.lines; + } + } + + return { path, include, exclude, matches, totalMatches, errorMessage }; +} diff --git a/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/parsers/grep-search.ts b/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/parsers/grep-search.ts new file mode 100644 index 000000000..0e606e193 --- /dev/null +++ b/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/parsers/grep-search.ts @@ -0,0 +1,108 @@ +// Meta parser for `grep_search` tool calls. Reads the path/pattern +// triplet from args (strict parsing - we wait for the args to +// complete) and the matches from the result blob. The result parser +// keeps the original "scan result as raw text on JSON.parse failure" +// fallback so MCP servers that return unparseable output still get +// surfaced. + +import { BuiltInTool } from '$lib/enums'; +import { splitSearchSummaryList, type AgenticSection } from '$lib/utils'; +import { parseToolArgs } from './_shared'; + +export type GrepSearchMatch = { + file: string; + line?: number; + content: string; +}; + +export type GrepSearchMeta = { + path: string; + pattern: string; + include: string; + exclude?: string; + showLineNumbers: boolean; + matches: GrepSearchMatch[]; + totalMatches?: number; + errorMessage?: string; +}; + +export function parseGrepSearchMeta(section: AgenticSection): GrepSearchMeta | null { + const args = parseToolArgs(BuiltInTool.GREP_SEARCH, section); + if (!args) return null; + + const path = typeof args.path === 'string' ? args.path : ''; + const pattern = typeof args.pattern === 'string' ? args.pattern : ''; + if (!path || !pattern) return null; + + const include = typeof args.include === 'string' && args.include ? args.include : '**'; + const exclude = typeof args.exclude === 'string' && args.exclude ? args.exclude : undefined; + const showLineNumbers = args.return_line_numbers === true; + + let matches: GrepSearchMatch[] = []; + let totalMatches: number | undefined; + let errorMessage: string | undefined; + + const toolResultString = section.toolResult; + if (toolResultString) { + try { + const parsed: unknown = JSON.parse(toolResultString); + if (parsed && typeof parsed === 'object' && !Array.isArray(parsed)) { + const obj = parsed as Record; + if (typeof obj.error === 'string') { + errorMessage = obj.error; + } else if (typeof obj.plain_text_response === 'string') { + const split = splitSearchSummaryList(obj.plain_text_response, (total) => { + totalMatches = total; + }); + matches = split.lines.map((line) => parseGrepLine(line, showLineNumbers)); + } + } + } catch { + // Result wasn't JSON: keep behaviour for MCP servers that + // emit raw text and treat each line as a `:` + // (or `::`) match. + const split = splitSearchSummaryList(toolResultString, (total) => { + totalMatches = total; + }); + matches = split.lines.map((line) => parseGrepLine(line, showLineNumbers)); + } + } + + return { + path, + pattern, + include, + exclude, + showLineNumbers, + matches, + totalMatches, + errorMessage + }; +} + +function parseGrepLine(line: string, showLineNumbers: boolean): GrepSearchMatch { + // Server output: + // : when return_line_numbers=false + // :: when return_line_numbers=true + const firstColon = line.indexOf(':'); + if (firstColon === -1) { + return { file: line, content: '' }; + } + const file = line.slice(0, firstColon); + const tail = line.slice(firstColon + 1); + + if (!showLineNumbers) { + return { file, content: tail }; + } + + const secondColon = tail.indexOf(':'); + if (secondColon === -1) { + return { file, content: tail }; + } + const lineNum = parseInt(tail.slice(0, secondColon), 10); + return { + file, + line: Number.isFinite(lineNum) ? lineNum : undefined, + content: tail.slice(secondColon + 1) + }; +} diff --git a/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/parsers/read-file.ts b/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/parsers/read-file.ts new file mode 100644 index 000000000..d37dcf501 --- /dev/null +++ b/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/parsers/read-file.ts @@ -0,0 +1,52 @@ +// Meta parser for `read_file` tool calls. Reads the file path and an +// optional line range (either `start_line`+`end_line` or +// `start_line`+`line_count`). Args are parsed partially so a header +// can render incrementally as the file path streams in. + +import { BuiltInTool } from '$lib/enums'; +import { + DEFAULT_LANGUAGE, + FILE_PATH_SEPARATOR_REGEX, + TEXT_LANGUAGE_PREFIX_REGEX +} from '$lib/constants'; +import { getFileTypeByExtension, type AgenticSection } from '$lib/utils'; +import { parseToolArgs } from './_shared'; + +export type ReadFileMeta = { + fileName: string; + lineRange: { start: number; end: number } | null; + language: string; +}; + +export function parseReadFileMeta(section: AgenticSection): ReadFileMeta | null { + const args = parseToolArgs(BuiltInTool.READ_FILE, section, { partial: true }); + if (!args) return null; + + const rawPath = args.path ?? args.file_path ?? args.filePath; + if (typeof rawPath !== 'string' || !rawPath) return null; + + const fileName = rawPath.split(FILE_PATH_SEPARATOR_REGEX).pop() || rawPath; + + // Models emit range arguments under several aliases. Accept all to + // stay forgiving across prompt variations. + const startRaw = args.start_line ?? args.line_start ?? args.startLine ?? args.from_line; + const endRaw = args.end_line ?? args.line_end ?? args.endLine ?? args.to_line; + const countRaw = args.line_count ?? args.count ?? args.num_lines; + + let lineRange: { start: number; end: number } | null = null; + const sNum = Number(startRaw); + const eNum = Number(endRaw); + if (startRaw != null && endRaw != null && Number.isFinite(sNum) && Number.isFinite(eNum)) { + lineRange = { start: sNum, end: eNum }; + } else if (startRaw != null && countRaw != null) { + const cNum = Number(countRaw); + if (Number.isFinite(sNum) && Number.isFinite(cNum)) { + lineRange = { start: sNum, end: sNum + cNum - 1 }; + } + } + + const fileType = getFileTypeByExtension(fileName); + const language = fileType ? fileType.replace(TEXT_LANGUAGE_PREFIX_REGEX, '') : DEFAULT_LANGUAGE; + + return { fileName, lineRange, language }; +} diff --git a/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/parsers/run-javascript.ts b/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/parsers/run-javascript.ts new file mode 100644 index 000000000..9bcba8f03 --- /dev/null +++ b/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/parsers/run-javascript.ts @@ -0,0 +1,56 @@ +// Meta parser for `run_javascript` tool calls. Reads the JS code and +// optional timeout from args (strict parsing) and surfaces any error +// from the result blob. SandboxService.formatReply emits a JSON object +// containing an `error` field on failure, but a partial/non-JSON +// failure renders as a flat line beginning with `Error:`. Both shapes +// are handled. + +import { BuiltInTool } from '$lib/enums'; +import type { AgenticSection } from '$lib/utils'; +import { parseToolArgs } from './_shared'; + +export type RunJavascriptMeta = { + code: string; + timeoutMs?: number; + errorMessage?: string; +}; + +export function parseRunJavascriptMeta(section: AgenticSection): RunJavascriptMeta | null { + const args = parseToolArgs(BuiltInTool.RUN_JAVASCRIPT, section); + if (!args) return null; + + const code = typeof args.code === 'string' ? args.code : ''; + if (!code) return null; + + const timeoutRaw = Number(args.timeout_ms); + const timeoutMs = Number.isFinite(timeoutRaw) && timeoutRaw > 0 ? timeoutRaw : undefined; + + let errorMessage: string | undefined; + const toolResultString = section.toolResult; + if (toolResultString) { + // Branches matter here: a JSON object can carry `error`, but a + // JSON array always represents successful output (sandbox returns + // the array of values). Only when the result isn't a JSON object + // do we scan raw lines for the `Error:` prefix. + let parsedObject: Record | null = null; + try { + const parsed: unknown = JSON.parse(toolResultString); + if (parsed && typeof parsed === 'object' && !Array.isArray(parsed)) { + parsedObject = parsed as Record; + } + } catch { + parsedObject = null; + } + if (typeof parsedObject?.error === 'string') { + errorMessage = parsedObject.error; + } else if (!parsedObject) { + const errorLine = toolResultString + .split('\n') + .map((line) => line.trim()) + .find((line) => line.startsWith('Error:')); + if (errorLine) errorMessage = errorLine.slice('Error:'.length).trim(); + } + } + + return { code, timeoutMs, errorMessage }; +} diff --git a/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/parsers/write-file.ts b/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/parsers/write-file.ts new file mode 100644 index 000000000..95edc3d95 --- /dev/null +++ b/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/parsers/write-file.ts @@ -0,0 +1,54 @@ +// Meta parser for `write_file` tool calls. Reads the path/content from +// the streamed args (partial JSON so we can render before the call +// finishes) and surfaces `bytes`, `result`, and `error` from the +// result blob. + +import { BuiltInTool } from '$lib/enums'; +import { + DEFAULT_LANGUAGE, + FILE_PATH_SEPARATOR_REGEX, + TEXT_LANGUAGE_PREFIX_REGEX +} from '$lib/constants'; +import { getFileTypeByExtension, tryParseToolResultObject, type AgenticSection } from '$lib/utils'; +import { parseToolArgs } from './_shared'; + +export type WriteFileMeta = { + fileName: string; + filePath: string; + language: string; + content: string; + bytesWritten?: number; + resultMessage?: string; + errorMessage?: string; +}; + +export function parseWriteFileMeta(section: AgenticSection): WriteFileMeta | null { + const args = parseToolArgs(BuiltInTool.WRITE_FILE, section, { partial: true }); + if (!args) return null; + + // Tool contracts drifted over time: some models emit `path`, + // others `file_path` / `filePath`. Accept all three. + const rawPath = args.path ?? args.file_path ?? args.filePath; + if (typeof rawPath !== 'string' || !rawPath) return null; + + const fileName = rawPath.split(FILE_PATH_SEPARATOR_REGEX).pop() || rawPath; + const content = typeof args.content === 'string' ? args.content : ''; + const language = + getFileTypeByExtension(rawPath)?.replace(TEXT_LANGUAGE_PREFIX_REGEX, '') ?? DEFAULT_LANGUAGE; + + const resultObj = tryParseToolResultObject(section.toolResult); + const bytesWritten = + resultObj && Number.isFinite(Number(resultObj.bytes)) ? Number(resultObj.bytes) : undefined; + const resultMessage = typeof resultObj?.result === 'string' ? resultObj.result : undefined; + const errorMessage = typeof resultObj?.error === 'string' ? resultObj.error : undefined; + + return { + fileName, + filePath: rawPath, + language, + content, + bytesWritten, + resultMessage, + errorMessage + }; +} diff --git a/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessageActions/ChatMessageActionCard/ChatMessageActionCard.svelte b/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessageActions/ChatMessageActionCard/ChatMessageActionCard.svelte index 254031979..17d8e21d7 100644 --- a/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessageActions/ChatMessageActionCard/ChatMessageActionCard.svelte +++ b/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessageActions/ChatMessageActionCard/ChatMessageActionCard.svelte @@ -1,4 +1,5 @@ -{#snippet renderSection(section: (typeof sectionsParsed)[number], index: number)} +{#snippet renderSection(section: AgenticSection, index: number)} {#if section.type === AgenticSectionType.TEXT}
- {:else if section.type === AgenticSectionType.TOOL_CALL_STREAMING} - {@const streamingIcon = isStreaming ? Loader2 : Loader2} - {@const streamingIconClass = isStreaming ? 'h-4 w-4 animate-spin' : 'h-4 w-4'} - - toggleExpanded(index, section)} - > -
-
- Arguments: - - {#if isStreaming} - - {/if} -
- {#if section.toolArgs} - - {:else if isStreaming} -
- Receiving arguments... -
- {:else} -
- Response was truncated -
- {/if} -
- - {:else if section.type === AgenticSectionType.TOOL_CALL || section.type === AgenticSectionType.TOOL_CALL_PENDING} - {@const isPending = section.type === AgenticSectionType.TOOL_CALL_PENDING} - {@const toolIcon = isPending ? Loader2 : Wrench} - {@const toolIconClass = isPending ? 'h-4 w-4 animate-spin' : 'h-4 w-4'} - - + {:else if section.type === AgenticSectionType.TOOL_CALL || section.type === AgenticSectionType.TOOL_CALL_PENDING || section.type === AgenticSectionType.TOOL_CALL_STREAMING} + toggleExpanded(index, section)} - > - {#if section.toolArgs && section.toolArgs !== '{}'} -
-
Arguments:
- - -
- {/if} - -
-
- Result: - - {#if isPending} - - {/if} -
- {#if isPending} -
- Waiting for result... -
- {:else if section.toolResult} -
- {#each section.parsedLines as line, i (i)} -
- {line.text} -
- {#if line.image} - {line.image.name} - {/if} - {/each} -
- {:else} -
No output
- {/if} -
- - {:else if section.type === AgenticSectionType.REASONING} - {@const reasoningSubtitle = section.wasInterrupted - ? hasReasoningError - ? 'Error' - : 'Cancelled' - : isStreaming - ? '' - : undefined} - - toggleExpanded(index, section)} - > -
- {#if renderThinkingAsMarkdown} - - {:else} -
- {section.content} -
- {/if} -
- - {:else if section.type === AgenticSectionType.REASONING_PENDING} - {@const reasoningTitle = isStreaming ? 'Reasoning...' : 'Reasoning'} - {@const reasoningSubtitle = isStreaming ? '' : hasReasoningError ? 'Error' : 'Cancelled'} - - toggleExpanded(index, section)} - > -
- {#if renderThinkingAsMarkdown} - - {:else} -
- {section.content} -
- {/if} -
- + /> {/if} {/snippet} -
+
{#if turnGroups.length > 1} {#each turnGroups as turn, turnIndex (turnIndex)} {@const turnStats = message?.timings?.agentic?.perTurn?.[turnIndex]} -
+
{#each turn.sections as section, sIdx (turn.flatIndices[sIdx])} {@render renderSection(section, turn.flatIndices[sIdx])} {/each} - {#if turnStats && showMessageStats} -
+ {#if turnStats && showAgenticTurnStats} +
{/each} {:else} - {#each sectionsParsed as section, index (index)} + {#each sections as section, index (index)} {@render renderSection(section, index)} {/each} {/if} @@ -404,7 +267,6 @@ flex-direction: column; width: 100%; max-width: 48rem; - gap: 1rem; } .agentic-content > :global(*), diff --git a/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessageReasoningBlock.svelte b/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessageReasoningBlock.svelte new file mode 100644 index 000000000..833cae5db --- /dev/null +++ b/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessageReasoningBlock.svelte @@ -0,0 +1,151 @@ + + + +
+ {#if renderThinkingAsMarkdown} + + {:else} +
+ {section.content} +
+ {/if} +
+ + + diff --git a/tools/ui/src/lib/components/app/chat/ChatScreen/ChatScreenActionScrollDown.svelte b/tools/ui/src/lib/components/app/chat/ChatScreen/ChatScreenActionScrollDown.svelte index c470ac729..dca24afd4 100644 --- a/tools/ui/src/lib/components/app/chat/ChatScreen/ChatScreenActionScrollDown.svelte +++ b/tools/ui/src/lib/components/app/chat/ChatScreen/ChatScreenActionScrollDown.svelte @@ -1,4 +1,5 @@ @@ -76,59 +45,54 @@ open = value; onToggle?.(); }} - class="{className} my-0!" + class={cn('group/collapsible', 'my-0!', className)} > - - -
-
- {#if IconComponent} - - {/if} + +
+ {#if iconUrl} + + {:else if IconComponent} + + {/if} - {title} - - {#if subtitle} - {subtitle} - {/if} -
- - {#if displayedPreview && !showThoughtInProgress} -
-
- {displayedPreview} -
- {#if displayedOverflow > 0} - {displayedOverflow}+ chars - {/if} -
+ + {#if titleSnippet} + {@render titleSnippet()} + {:else} + {title} {/if} -
+ -
- + {#if subtitle} + {subtitle} + {/if} +
- Toggle content -
- + - -
+ Toggle content + + + +
+
{@render children()}
- - +
+ diff --git a/tools/ui/src/lib/components/app/content/CollapsibleTerminalBlock.svelte b/tools/ui/src/lib/components/app/content/CollapsibleTerminalBlock.svelte new file mode 100644 index 000000000..4370aea42 --- /dev/null +++ b/tools/ui/src/lib/components/app/content/CollapsibleTerminalBlock.svelte @@ -0,0 +1,97 @@ + + + { + open = value; + onToggle?.(); + }} + class={cn('group/collapsible', 'overflow-hidden rounded-md', className)} + style="background: var(--code-background); border: 1px solid color-mix(in oklch, var(--border) 30%, transparent);" +> + +
+ {#if iconUrl} + + {:else if IconComponent} + + {/if} + + + {#if titleSnippet} + {@render titleSnippet()} + {:else} + {title} + {/if} + + + {#if subtitle} + {subtitle} + {/if} +
+ + + + Toggle content + + + +
+ {@render children()} +
+ + diff --git a/tools/ui/src/lib/components/app/content/MarkdownContent/markdown-content.css b/tools/ui/src/lib/components/app/content/MarkdownContent/markdown-content.css index b0e04ca62..41813f4fd 100644 --- a/tools/ui/src/lib/components/app/content/MarkdownContent/markdown-content.css +++ b/tools/ui/src/lib/components/app/content/MarkdownContent/markdown-content.css @@ -19,8 +19,16 @@ line-height: 1.75; } +.markdown-content :global(.markdown-block:first-child p:first-child) { + margin-block-start: 0; +} + +.markdown-content :global(.markdown-block:last-child p:last-child) { + margin-block-end: 0; +} + .markdown-content :global(:is(h1, h2, h3, h4, h5, h6):first-child) { - margin-top: 0; + margin-top: 0.5rem; } /* Headers with consistent spacing */ diff --git a/tools/ui/src/lib/components/app/content/MermaidPreviewControls.svelte b/tools/ui/src/lib/components/app/content/MermaidPreviewControls.svelte index bb3185f40..39540e7a8 100644 --- a/tools/ui/src/lib/components/app/content/MermaidPreviewControls.svelte +++ b/tools/ui/src/lib/components/app/content/MermaidPreviewControls.svelte @@ -1,4 +1,5 @@
- + 
{@html highlightedHtml}
diff --git a/tools/ui/src/lib/components/app/content/index.ts b/tools/ui/src/lib/components/app/content/index.ts index 5d2884bb2..5cfdd1b9c 100644 --- a/tools/ui/src/lib/components/app/content/index.ts +++ b/tools/ui/src/lib/components/app/content/index.ts @@ -68,7 +68,6 @@ export { default as SyntaxHighlightedCode } from './SyntaxHighlightedCode.svelte * ```svelte * @@ -78,6 +77,22 @@ export { default as SyntaxHighlightedCode } from './SyntaxHighlightedCode.svelte */ export { default as CollapsibleContentBlock } from './CollapsibleContentBlock.svelte'; +/** + * **CollapsibleTerminalBlock** - Expandable content card with a terminal-style frame + * + * Same shape as CollapsibleContentBlock, but with a `code-background` + * fill, subtle border, and tightened padding suited for shell command + * output and similar dense / monospace content. + * + * @example + * ```svelte + * + * 
{output}
+ * + * ``` + */ +export { default as CollapsibleTerminalBlock } from './CollapsibleTerminalBlock.svelte'; + /** * **MermaidPreview** - Interactive Mermaid diagram viewer * diff --git a/tools/ui/src/lib/components/app/dialogs/DialogMcpResourcesBrowser.svelte b/tools/ui/src/lib/components/app/dialogs/DialogMcpResourcesBrowser.svelte index eb162a557..f741b544b 100644 --- a/tools/ui/src/lib/components/app/dialogs/DialogMcpResourcesBrowser.svelte +++ b/tools/ui/src/lib/components/app/dialogs/DialogMcpResourcesBrowser.svelte @@ -1,4 +1,5 @@ - + - Add New Server + Add New MCP Server + {#if recommendationsToShow.length > 0} +
+
+

Recommended Servers

+ +
+ +
+ {#each recommendationsToShow as recommendation (recommendation.id)} + handleRecommendationClick(recommendation.id)} + selected={selectedRecommendationId === recommendation.id} + dimmed={hasSelection && selectedRecommendationId !== recommendation.id} + /> + {/each} +
+
+ {/if} +
(newServerUseProxy = v)} urlError={newServerUrl ? newServerUrlError : null} id="new-server" + bind:wantsAuthorization={newServerWantsAuthorization} + required={authRequired} />
diff --git a/tools/ui/src/lib/components/app/dialogs/DialogModelNotAvailable.svelte b/tools/ui/src/lib/components/app/dialogs/DialogModelNotAvailable.svelte index a6c20291f..89d23cd4b 100644 --- a/tools/ui/src/lib/components/app/dialogs/DialogModelNotAvailable.svelte +++ b/tools/ui/src/lib/components/app/dialogs/DialogModelNotAvailable.svelte @@ -1,4 +1,5 @@ + + +
+ {#if activeIconUrl} + + {/if} + +

{server.name}

+
+ +

{server.description}

+ diff --git a/tools/ui/src/lib/components/app/mcp/McpServerForm.svelte b/tools/ui/src/lib/components/app/mcp/McpServerForm.svelte index 7f05d5fef..a7472add3 100644 --- a/tools/ui/src/lib/components/app/mcp/McpServerForm.svelte +++ b/tools/ui/src/lib/components/app/mcp/McpServerForm.svelte @@ -5,9 +5,14 @@ import type { KeyValuePair } from '$lib/types'; import { parseHeadersToArray, serializeHeaders } from '$lib/utils'; import { UrlProtocol } from '$lib/enums'; - import { MCP_SERVER_URL_PLACEHOLDER } from '$lib/constants'; + import { + AUTHORIZATION_HEADER, + BEARER_PREFIX, + CLI_FLAGS, + MCP_SERVER_URL_PLACEHOLDER, + REDACTED_HEADERS + } from '$lib/constants'; import { mcpStore } from '$lib/stores/mcp.svelte'; - import { CLI_FLAGS } from '$lib/constants'; interface Props { url: string; @@ -18,6 +23,22 @@ onUseProxyChange?: (useProxy: boolean) => void; urlError?: string | null; id?: string; + /** + * "Wants Authorization" is the user's *intent* to add a Bearer token + * (separate from `hasAuthorization` which reflects what's already in + * the headers). Bindable so a parent - e.g. the recommendation cards + * on the "Add New Server" dialog - can flip the switch on when the + * picked server ships a `needsAuthorization: true` flag. + */ + wantsAuthorization?: boolean; + /** + * Marks the "Authorization" field as required. Locks the toggle so the + * user can't dismiss it, and visually marks the field with a red + * asterisk. The parent is expected to gate its submit affordance on + * the bearer token actually being filled. Used by the "Add New Server" + * dialog for recommendations whose `needsAuthorization` flag is true. + */ + required?: boolean; } let { @@ -28,7 +49,9 @@ onHeadersChange, onUseProxyChange, urlError = null, - id = 'server' + id = 'server', + wantsAuthorization = $bindable(false), + required = false }: Props = $props(); let isWebSocket = $derived( @@ -38,14 +61,11 @@ let headerPairs = $derived(parseHeadersToArray(headers)); - const AUTHORIZATION_HEADER = 'Authorization'; - const BEARER_PREFIX = 'Bearer '; - // Heuristic: this dedicated UI only owns Authorization headers that already // carry a Bearer scheme. Anything else (e.g. Basic, raw tokens) stays in the // KV section so the user can still edit those values verbatim. const matchesAuthorizationKey = (key: string): boolean => - key.trim().toLowerCase() === AUTHORIZATION_HEADER.toLowerCase(); + REDACTED_HEADERS.has(key.trim().toLowerCase()); const isBearerScheme = (value: string): boolean => value.trim().toLowerCase().startsWith(BEARER_PREFIX.toLowerCase()); @@ -55,8 +75,6 @@ let hasAuthorization = $derived(headerPairs.some(ownedByBearerUi)); - let wantsAuthorization = $state(false); - let showAuthorization = $derived(hasAuthorization || wantsAuthorization); let urlInput: HTMLInputElement | null = $state(null); @@ -119,7 +137,7 @@
-
-