From ba06fe254132771ecc097e90bd0e6457dbb74842 Mon Sep 17 00:00:00 2001 From: Nikhil Jain Date: Tue, 10 Mar 2026 20:21:21 -0700 Subject: [PATCH 1/9] Batch resolution of timestamp queries instead of eagerly resolving to reduce gpu kernel profiling overhead --- ggml/src/ggml-webgpu/ggml-webgpu.cpp | 169 ++++++++++++++++++++------- 1 file changed, 124 insertions(+), 45 deletions(-) diff --git a/ggml/src/ggml-webgpu/ggml-webgpu.cpp b/ggml/src/ggml-webgpu/ggml-webgpu.cpp index ccc34cb153f..7bfb875d3a1 100644 --- a/ggml/src/ggml-webgpu/ggml-webgpu.cpp +++ b/ggml/src/ggml-webgpu/ggml-webgpu.cpp @@ -75,7 +75,8 @@ static inline void compute_2d_workgroups(uint32_t total_wg, uint32_t max_per_dim #endif // GGML_WEBGPU_CPU_PROFILE #ifdef GGML_WEBGPU_GPU_PROFILE -# define WEBGPU_NUM_TIMESTAMP_QUERY_BUFS 32 +# define WEBGPU_NUM_INFLIGHT_PROFILE_SUBS 32 // we can have maximum 64 in-flight futures at once and max 32 QuerySets +# define WEBGPU_NUM_TIMESTAMP_QUERY_BUFS WEBGPU_NUM_INFLIGHT_PROFILE_SUBS # define WEBGPU_TIMESTAMP_QUERY_BUF_SIZE_BYTES 16 // e.g. enough for two timestamps #endif @@ -205,18 +206,19 @@ struct webgpu_buf_pool { }; #ifdef GGML_WEBGPU_GPU_PROFILE -struct webgpu_gpu_profile_bufs { +struct webgpu_profile_bufs { + int id; wgpu::Buffer host_buf; wgpu::Buffer dev_buf; wgpu::QuerySet query_set; }; -// Holds a pool of parameter buffers for WebGPU operations -struct webgpu_gpu_profile_buf_pool { - std::vector free; - std::mutex mutex; +struct webgpu_profile_buf_pool { + std::vector free; + std::vector in_use; + std::mutex mutex; std::condition_variable cv; void init(wgpu::Device device, @@ -224,47 +226,108 @@ struct webgpu_gpu_profile_buf_pool { size_t buf_size, wgpu::BufferUsage dev_buf_usage, wgpu::BufferUsage host_buf_usage) { + std::lock_guard lock(mutex); + for (int i = 0; i < num_bufs; i++) { wgpu::Buffer host_buf; wgpu::Buffer dev_buf; - ggml_webgpu_create_buffer(device, host_buf, buf_size, host_buf_usage, "ggml_webgpu_host_profile_buf"); - ggml_webgpu_create_buffer(device, dev_buf, buf_size, dev_buf_usage, "ggml_webgpu_dev_profile_buf"); - // Create a query set for 2 timestamps + + ggml_webgpu_create_buffer( + device, host_buf, buf_size, host_buf_usage, + "ggml_webgpu_host_profile_buf"); + ggml_webgpu_create_buffer( + device, dev_buf, buf_size, dev_buf_usage, + "ggml_webgpu_dev_profile_buf"); + wgpu::QuerySetDescriptor ts_query_set_desc = {}; + ts_query_set_desc.type = wgpu::QueryType::Timestamp; + ts_query_set_desc.count = 2; - ts_query_set_desc.type = wgpu::QueryType::Timestamp; - ts_query_set_desc.count = 2; wgpu::QuerySet ts_query_set = device.CreateQuerySet(&ts_query_set_desc); - free.push_back({ host_buf, dev_buf, ts_query_set }); + free.push_back({ i, host_buf, dev_buf, ts_query_set }); } } - webgpu_gpu_profile_bufs alloc_bufs() { + webgpu_profile_bufs alloc_bufs() { std::unique_lock lock(mutex); cv.wait(lock, [this] { return !free.empty(); }); - webgpu_gpu_profile_bufs bufs = free.back(); + + in_use.push_back(free.back()); free.pop_back(); - return bufs; + + return in_use.back(); } - void free_bufs(std::vector bufs) { + void free_bufs(const std::vector & bufs) { std::lock_guard lock(mutex); - free.insert(free.end(), bufs.begin(), bufs.end()); + + for (const auto & buf : bufs) { + auto it = std::find_if( + in_use.begin(), in_use.end(), + [&](const webgpu_profile_bufs & x) { + return x.id == buf.id; + }); + + if (it != in_use.end()) { + free.push_back(*it); + in_use.erase(it); + } + } + cv.notify_all(); } + size_t num_free_bufs() { + std::lock_guard lock(mutex); + return free.size(); + } + + size_t num_in_use_bufs() { + std::lock_guard lock(mutex); + return in_use.size(); + } + void cleanup() { std::lock_guard lock(mutex); + for (auto & bufs : free) { - bufs.host_buf.Destroy(); - bufs.dev_buf.Destroy(); - bufs.query_set.Destroy(); + if (bufs.host_buf) { + if (bufs.host_buf.GetMapState() == wgpu::BufferMapState::Mapped) { + bufs.host_buf.Unmap(); + } + bufs.host_buf.Destroy(); + } + if (bufs.dev_buf) { + bufs.dev_buf.Destroy(); + } + if (bufs.query_set) { + bufs.query_set.Destroy(); + } + } + + for (auto & bufs : in_use) { + if (bufs.host_buf) { + if (bufs.host_buf.GetMapState() == wgpu::BufferMapState::Mapped) { + bufs.host_buf.Unmap(); + } + bufs.host_buf.Destroy(); + } + if (bufs.dev_buf) { + bufs.dev_buf.Destroy(); + } + if (bufs.query_set) { + bufs.query_set.Destroy(); + } } + free.clear(); + in_use.clear(); } - ~webgpu_gpu_profile_buf_pool() { this->cleanup(); } + ~webgpu_profile_buf_pool() { + cleanup(); + } }; #endif @@ -273,7 +336,7 @@ struct webgpu_command { wgpu::CommandBuffer commands; std::vector params_bufs; #ifdef GGML_WEBGPU_GPU_PROFILE - webgpu_gpu_profile_bufs timestamp_query_bufs; + webgpu_profile_bufs timestamp_query_bufs; std::string pipeline_name; #endif }; @@ -301,7 +364,7 @@ struct webgpu_global_context_struct { webgpu_capabilities capabilities; // Shared buffer to move data from device to host wgpu::Buffer get_tensor_staging_buf; - // Global mutex for pipeline and staging buffer, will be refactored to exclude pipeline caches. + // Global mutex for pipeline and staging buffer std::recursive_mutex mutex; webgpu_buf_pool memset_buf_pool; @@ -318,7 +381,8 @@ struct webgpu_global_context_struct { // Profiling: per-shader GPU time in ms std::unordered_map shader_gpu_time_ms; // Profiling: pool of timestamp query buffers (one per operation) - webgpu_gpu_profile_buf_pool timestamp_query_buf_pool; + webgpu_profile_buf_pool timestamp_query_buf_pool; + std::vector profiling_futures; #endif #ifdef GGML_WEBGPU_DEBUG @@ -491,10 +555,15 @@ static void ggml_backend_webgpu_wait_profile_futures(webgpu_global_context & uint64_t timeout_ms = block ? UINT64_MAX : 0; if (block) { while (!futures.empty()) { - auto waitStatus = ctx->instance.WaitAny(futures.size(), futures.data(), timeout_ms); - if (ggml_backend_webgpu_handle_wait_status(waitStatus)) { - ggml_backend_webgpu_erase_completed_futures(futures); + // We can only wait on 64 futures at a time + size_t count = std::min(WEBGPU_NUM_INFLIGHT_PROFILE_SUBS, futures.size()); + + auto waitStatus = ctx->instance.WaitAny(count, futures.data(), UINT64_MAX); + if (!ggml_backend_webgpu_handle_wait_status(waitStatus)) { + break; } + + ggml_backend_webgpu_erase_completed_futures(futures); } } else { auto waitStatus = ctx->instance.WaitAny(futures.size(), futures.data(), timeout_ms); @@ -516,9 +585,6 @@ static void ggml_backend_webgpu_wait(webgpu_global_context & ctx, while (subs.size() >= WEBGPU_MAX_INFLIGHT_SUBS_PER_THREAD) { auto waitStatus = ctx->instance.WaitAny(1, &subs[0].submit_done, UINT64_MAX); if (ggml_backend_webgpu_handle_wait_status(waitStatus)) { -#ifdef GGML_WEBGPU_GPU_PROFILE - ggml_backend_webgpu_wait_profile_futures(ctx, subs[0].profile_futures, true); -#endif subs.erase(subs.begin()); } } @@ -533,9 +599,6 @@ static void ggml_backend_webgpu_wait(webgpu_global_context & ctx, auto waitStatus = ctx->instance.WaitAny(1, &sub.submit_done, UINT64_MAX); ggml_backend_webgpu_handle_wait_status(waitStatus); } -#ifdef GGML_WEBGPU_GPU_PROFILE - ggml_backend_webgpu_wait_profile_futures(ctx, sub.profile_futures, true); -#endif } subs.clear(); } else { @@ -543,12 +606,7 @@ static void ggml_backend_webgpu_wait(webgpu_global_context & ctx, for (auto sub = subs.begin(); sub != subs.end();) { auto waitStatus = ctx->instance.WaitAny(1, &sub->submit_done, 0); ggml_backend_webgpu_handle_wait_status(waitStatus, true); -#ifdef GGML_WEBGPU_GPU_PROFILE - ggml_backend_webgpu_wait_profile_futures(ctx, sub->profile_futures, false); - if (sub->submit_done.completed && sub->profile_futures.empty()) { -#else if (sub->submit_done.completed) { -#endif sub = subs.erase(sub); } else { ++sub; @@ -595,7 +653,7 @@ static webgpu_submission ggml_backend_webgpu_submit(webgpu_global_context & std::vector params_bufs; webgpu_submission submission; #ifdef GGML_WEBGPU_GPU_PROFILE - std::vector> pipeline_name_and_ts_bufs; + std::vector> pipeline_name_and_ts_bufs; #endif for (const auto & command : commands) { @@ -634,7 +692,7 @@ static webgpu_submission ggml_backend_webgpu_submit(webgpu_global_context & // We can't unmap in here due to WebGPU reentrancy limitations. ctx->timestamp_query_buf_pool.free_bufs({ ts_bufs }); }); - submission.profile_futures.push_back({ f }); + ctx->profiling_futures.push_back({ f, false }); } #endif return submission; @@ -678,7 +736,19 @@ static webgpu_command ggml_backend_webgpu_build_multi( } #ifdef GGML_WEBGPU_GPU_PROFILE - webgpu_gpu_profile_bufs ts_bufs = ctx->timestamp_query_buf_pool.alloc_bufs(); + // Only copy timestamp query results if we have no more timestamp buffers available + if (ctx->timestamp_query_buf_pool.num_free_bufs() == 0) { + for (auto& ts_bufs : ctx->timestamp_query_buf_pool.in_use) { + encoder.ResolveQuerySet(ts_bufs.query_set, 0, 2, ts_bufs.dev_buf, 0); + encoder.CopyBufferToBuffer(ts_bufs.dev_buf, 0, ts_bufs.host_buf, 0, ts_bufs.host_buf.GetSize()); + } + + // buffers are freed within the future, and futures are deleted here + ggml_backend_webgpu_wait_profile_futures(ctx, + ctx->profiling_futures, true); + } + + webgpu_profile_bufs ts_bufs = ctx->timestamp_query_buf_pool.alloc_bufs(); if (ts_bufs.host_buf.GetMapState() == wgpu::BufferMapState::Mapped) { ts_bufs.host_buf.Unmap(); } @@ -698,10 +768,10 @@ static webgpu_command ggml_backend_webgpu_build_multi( } pass.End(); -#ifdef GGML_WEBGPU_GPU_PROFILE - encoder.ResolveQuerySet(ts_bufs.query_set, 0, 2, ts_bufs.dev_buf, 0); - encoder.CopyBufferToBuffer(ts_bufs.dev_buf, 0, ts_bufs.host_buf, 0, ts_bufs.host_buf.GetSize()); -#endif +// #ifdef GGML_WEBGPU_GPU_PROFILE +// encoder.ResolveQuerySet(ts_bufs.query_set, 0, 2, ts_bufs.dev_buf, 0); +// encoder.CopyBufferToBuffer(ts_bufs.dev_buf, 0, ts_bufs.host_buf, 0, ts_bufs.host_buf.GetSize()); +// #endif wgpu::CommandBuffer commands = encoder.Finish(); webgpu_command result = {}; @@ -785,6 +855,15 @@ static void ggml_backend_webgpu_free(ggml_backend_t backend) { #endif #ifdef GGML_WEBGPU_GPU_PROFILE + // resolve any stragglers + wgpu::CommandEncoder encoder = ctx->webgpu_ctx->global_ctx->device.CreateCommandEncoder(); + for (auto & ts_bufs : ctx->webgpu_ctx->global_ctx->timestamp_query_buf_pool.in_use) { + encoder.ResolveQuerySet(ts_bufs.query_set, 0, 2, ts_bufs.dev_buf, 0); + encoder.CopyBufferToBuffer(ts_bufs.dev_buf, 0, ts_bufs.host_buf, 0, ts_bufs.host_buf.GetSize()); + ggml_backend_webgpu_wait_profile_futures(ctx->webgpu_ctx->global_ctx, + ctx->webgpu_ctx->global_ctx->profiling_futures, true); + } + std::cout << "\n[ggml_webgpu gpu profiling summary]\n"; double total_gpu = 0.0; for (const auto & kv : ctx->webgpu_ctx->global_ctx->shader_gpu_time_ms) { From 753c98267cb4141ecbe20b17224545022703bd17 Mon Sep 17 00:00:00 2001 From: Nikhil Jain Date: Tue, 10 Mar 2026 20:24:43 -0700 Subject: [PATCH 2/9] Run clang-format --- ggml/src/ggml-webgpu/ggml-webgpu.cpp | 59 +++++++++++----------------- 1 file changed, 23 insertions(+), 36 deletions(-) diff --git a/ggml/src/ggml-webgpu/ggml-webgpu.cpp b/ggml/src/ggml-webgpu/ggml-webgpu.cpp index 7bfb875d3a1..bacc63a014c 100644 --- a/ggml/src/ggml-webgpu/ggml-webgpu.cpp +++ b/ggml/src/ggml-webgpu/ggml-webgpu.cpp @@ -75,8 +75,9 @@ static inline void compute_2d_workgroups(uint32_t total_wg, uint32_t max_per_dim #endif // GGML_WEBGPU_CPU_PROFILE #ifdef GGML_WEBGPU_GPU_PROFILE -# define WEBGPU_NUM_INFLIGHT_PROFILE_SUBS 32 // we can have maximum 64 in-flight futures at once and max 32 QuerySets -# define WEBGPU_NUM_TIMESTAMP_QUERY_BUFS WEBGPU_NUM_INFLIGHT_PROFILE_SUBS +# define WEBGPU_NUM_INFLIGHT_PROFILE_SUBS \ + 32 // we can have maximum 64 in-flight futures at once and max 32 QuerySets +# define WEBGPU_NUM_TIMESTAMP_QUERY_BUFS WEBGPU_NUM_INFLIGHT_PROFILE_SUBS # define WEBGPU_TIMESTAMP_QUERY_BUF_SIZE_BYTES 16 // e.g. enough for two timestamps #endif @@ -206,19 +207,20 @@ struct webgpu_buf_pool { }; #ifdef GGML_WEBGPU_GPU_PROFILE +// Holds a host and device buffer for profiling, +// along with a query_set to execute the timestampQuery struct webgpu_profile_bufs { - int id; + int id; wgpu::Buffer host_buf; wgpu::Buffer dev_buf; wgpu::QuerySet query_set; }; - struct webgpu_profile_buf_pool { std::vector free; std::vector in_use; - std::mutex mutex; + std::mutex mutex; std::condition_variable cv; void init(wgpu::Device device, @@ -232,16 +234,12 @@ struct webgpu_profile_buf_pool { wgpu::Buffer host_buf; wgpu::Buffer dev_buf; - ggml_webgpu_create_buffer( - device, host_buf, buf_size, host_buf_usage, - "ggml_webgpu_host_profile_buf"); - ggml_webgpu_create_buffer( - device, dev_buf, buf_size, dev_buf_usage, - "ggml_webgpu_dev_profile_buf"); + ggml_webgpu_create_buffer(device, host_buf, buf_size, host_buf_usage, "ggml_webgpu_host_profile_buf"); + ggml_webgpu_create_buffer(device, dev_buf, buf_size, dev_buf_usage, "ggml_webgpu_dev_profile_buf"); wgpu::QuerySetDescriptor ts_query_set_desc = {}; - ts_query_set_desc.type = wgpu::QueryType::Timestamp; - ts_query_set_desc.count = 2; + ts_query_set_desc.type = wgpu::QueryType::Timestamp; + ts_query_set_desc.count = 2; wgpu::QuerySet ts_query_set = device.CreateQuerySet(&ts_query_set_desc); @@ -263,11 +261,8 @@ struct webgpu_profile_buf_pool { std::lock_guard lock(mutex); for (const auto & buf : bufs) { - auto it = std::find_if( - in_use.begin(), in_use.end(), - [&](const webgpu_profile_bufs & x) { - return x.id == buf.id; - }); + auto it = std::find_if(in_use.begin(), in_use.end(), + [&](const webgpu_profile_bufs & x) { return x.id == buf.id; }); if (it != in_use.end()) { free.push_back(*it); @@ -325,9 +320,7 @@ struct webgpu_profile_buf_pool { in_use.clear(); } - ~webgpu_profile_buf_pool() { - cleanup(); - } + ~webgpu_profile_buf_pool() { cleanup(); } }; #endif @@ -337,7 +330,7 @@ struct webgpu_command { std::vector params_bufs; #ifdef GGML_WEBGPU_GPU_PROFILE webgpu_profile_bufs timestamp_query_bufs; - std::string pipeline_name; + std::string pipeline_name; #endif }; @@ -381,7 +374,7 @@ struct webgpu_global_context_struct { // Profiling: per-shader GPU time in ms std::unordered_map shader_gpu_time_ms; // Profiling: pool of timestamp query buffers (one per operation) - webgpu_profile_buf_pool timestamp_query_buf_pool; + webgpu_profile_buf_pool timestamp_query_buf_pool; std::vector profiling_futures; #endif @@ -738,15 +731,14 @@ static webgpu_command ggml_backend_webgpu_build_multi( #ifdef GGML_WEBGPU_GPU_PROFILE // Only copy timestamp query results if we have no more timestamp buffers available if (ctx->timestamp_query_buf_pool.num_free_bufs() == 0) { - for (auto& ts_bufs : ctx->timestamp_query_buf_pool.in_use) { + for (auto & ts_bufs : ctx->timestamp_query_buf_pool.in_use) { encoder.ResolveQuerySet(ts_bufs.query_set, 0, 2, ts_bufs.dev_buf, 0); encoder.CopyBufferToBuffer(ts_bufs.dev_buf, 0, ts_bufs.host_buf, 0, ts_bufs.host_buf.GetSize()); } - + // buffers are freed within the future, and futures are deleted here - ggml_backend_webgpu_wait_profile_futures(ctx, - ctx->profiling_futures, true); - } + ggml_backend_webgpu_wait_profile_futures(ctx, ctx->profiling_futures, true); + } webgpu_profile_bufs ts_bufs = ctx->timestamp_query_buf_pool.alloc_bufs(); if (ts_bufs.host_buf.GetMapState() == wgpu::BufferMapState::Mapped) { @@ -768,11 +760,6 @@ static webgpu_command ggml_backend_webgpu_build_multi( } pass.End(); -// #ifdef GGML_WEBGPU_GPU_PROFILE -// encoder.ResolveQuerySet(ts_bufs.query_set, 0, 2, ts_bufs.dev_buf, 0); -// encoder.CopyBufferToBuffer(ts_bufs.dev_buf, 0, ts_bufs.host_buf, 0, ts_bufs.host_buf.GetSize()); -// #endif - wgpu::CommandBuffer commands = encoder.Finish(); webgpu_command result = {}; result.commands = commands; @@ -860,10 +847,10 @@ static void ggml_backend_webgpu_free(ggml_backend_t backend) { for (auto & ts_bufs : ctx->webgpu_ctx->global_ctx->timestamp_query_buf_pool.in_use) { encoder.ResolveQuerySet(ts_bufs.query_set, 0, 2, ts_bufs.dev_buf, 0); encoder.CopyBufferToBuffer(ts_bufs.dev_buf, 0, ts_bufs.host_buf, 0, ts_bufs.host_buf.GetSize()); - ggml_backend_webgpu_wait_profile_futures(ctx->webgpu_ctx->global_ctx, - ctx->webgpu_ctx->global_ctx->profiling_futures, true); + ggml_backend_webgpu_wait_profile_futures(ctx->webgpu_ctx->global_ctx, + ctx->webgpu_ctx->global_ctx->profiling_futures, true); } - + std::cout << "\n[ggml_webgpu gpu profiling summary]\n"; double total_gpu = 0.0; for (const auto & kv : ctx->webgpu_ctx->global_ctx->shader_gpu_time_ms) { From 2457f41dcf750c550d0121cb9875a46aa69ec14c Mon Sep 17 00:00:00 2001 From: Nikhil Jain Date: Fri, 13 Mar 2026 23:18:33 -0700 Subject: [PATCH 3/9] Save state --- ggml/src/ggml-webgpu/ggml-webgpu.cpp | 76 ++++++++++++++++++++-------- 1 file changed, 55 insertions(+), 21 deletions(-) diff --git a/ggml/src/ggml-webgpu/ggml-webgpu.cpp b/ggml/src/ggml-webgpu/ggml-webgpu.cpp index bacc63a014c..4ef1cee281f 100644 --- a/ggml/src/ggml-webgpu/ggml-webgpu.cpp +++ b/ggml/src/ggml-webgpu/ggml-webgpu.cpp @@ -370,13 +370,13 @@ struct webgpu_global_context_struct { std::unordered_map cpu_detail_ms; #endif -#ifdef GGML_WEBGPU_GPU_PROFILE - // Profiling: per-shader GPU time in ms - std::unordered_map shader_gpu_time_ms; - // Profiling: pool of timestamp query buffers (one per operation) - webgpu_profile_buf_pool timestamp_query_buf_pool; - std::vector profiling_futures; -#endif +// #ifdef GGML_WEBGPU_GPU_PROFILE +// // Profiling: per-shader GPU time in ms +// std::unordered_map shader_gpu_time_ms; +// // Profiling: pool of timestamp query buffers (one per operation) +// webgpu_profile_buf_pool timestamp_query_buf_pool; +// std::vector profiling_futures; +// #endif #ifdef GGML_WEBGPU_DEBUG wgpu::Buffer debug_host_buf; @@ -430,6 +430,13 @@ struct webgpu_context_struct { std::map>> soft_max_pipelines; // mask_type, has_sink, inplace size_t memset_bytes_per_thread; +#ifdef GGML_WEBGPU_GPU_PROFILE + // Profiling: per-shader GPU time in ms + std::unordered_map shader_gpu_time_ms; + // Profiling: pool of timestamp query buffers (one per operation) + webgpu_profile_buf_pool timestamp_query_buf_pool; + std::vector profiling_futures; +#endif }; typedef std::shared_ptr webgpu_context; @@ -843,14 +850,42 @@ static void ggml_backend_webgpu_free(ggml_backend_t backend) { #ifdef GGML_WEBGPU_GPU_PROFILE // resolve any stragglers - wgpu::CommandEncoder encoder = ctx->webgpu_ctx->global_ctx->device.CreateCommandEncoder(); - for (auto & ts_bufs : ctx->webgpu_ctx->global_ctx->timestamp_query_buf_pool.in_use) { - encoder.ResolveQuerySet(ts_bufs.query_set, 0, 2, ts_bufs.dev_buf, 0); - encoder.CopyBufferToBuffer(ts_bufs.dev_buf, 0, ts_bufs.host_buf, 0, ts_bufs.host_buf.GetSize()); - ggml_backend_webgpu_wait_profile_futures(ctx->webgpu_ctx->global_ctx, - ctx->webgpu_ctx->global_ctx->profiling_futures, true); + auto & global_ctx = ctx->webgpu_ctx->global_ctx; + + // Copy out the current in-use buffers so callbacks/free_bufs() can't + // invalidate iteration while we're working. + std::vector stragglers; + { + std::lock_guard lock(global_ctx->timestamp_query_buf_pool.mutex); + stragglers = global_ctx->timestamp_query_buf_pool.in_use; + } + std::cout << "Stragglers size: " << stragglers.size() << "\n"; + if (!stragglers.empty()) { + wgpu::CommandEncoder encoder = global_ctx->device.CreateCommandEncoder(); + + for (auto & ts_bufs : stragglers) { + encoder.ResolveQuerySet(ts_bufs.query_set, 0, 2, ts_bufs.dev_buf, 0); + encoder.CopyBufferToBuffer( + ts_bufs.dev_buf, 0, + ts_bufs.host_buf, 0, + ts_bufs.host_buf.GetSize()); + } + + wgpu::CommandBuffer commands = encoder.Finish(); + global_ctx->queue.Submit(1, &commands); + + // If your backend needs event pumping for callbacks/progress, do it once here. + global_ctx->instance.ProcessEvents(); + + // Now wait for the profiling futures. These futures should already have been + // created by MapAsync when the profile buffers were originally queued. + ggml_backend_webgpu_wait_profile_futures( + global_ctx, + global_ctx->profiling_futures, + true); } + std::cout << "\n[ggml_webgpu gpu profiling summary]\n"; double total_gpu = 0.0; for (const auto & kv : ctx->webgpu_ctx->global_ctx->shader_gpu_time_ms) { @@ -2908,14 +2943,6 @@ static bool create_webgpu_device(ggml_backend_webgpu_reg_context * ctx) { wgpu::BufferUsage::CopySrc | wgpu::BufferUsage::MapWrite); ctx->webgpu_global_ctx->queue = ctx->webgpu_global_ctx->device.GetQueue(); -#ifdef GGML_WEBGPU_GPU_PROFILE - // Initialize buffer pool for timestamp queries, used for profiling - ctx->webgpu_global_ctx->timestamp_query_buf_pool.init( - ctx->webgpu_global_ctx->device, WEBGPU_NUM_TIMESTAMP_QUERY_BUFS, WEBGPU_TIMESTAMP_QUERY_BUF_SIZE_BYTES, - wgpu::BufferUsage::QueryResolve | wgpu::BufferUsage::CopySrc, - wgpu::BufferUsage::MapRead | wgpu::BufferUsage::CopyDst); -#endif - GGML_LOG_INFO( "ggml_webgpu: adapter_info: vendor_id: %u | vendor: %s | architecture: %s | device_id: %u | name: %s | " "device_desc: %s\n", @@ -2932,6 +2959,13 @@ static webgpu_context initialize_webgpu_context(ggml_backend_dev_t dev) { webgpu_ctx->param_buf_pool.init(webgpu_ctx->global_ctx->device, WEBGPU_NUM_PARAM_BUFS, WEBGPU_PARAMS_BUF_SIZE_BYTES, wgpu::BufferUsage::CopyDst | wgpu::BufferUsage::Uniform, wgpu::BufferUsage::CopySrc | wgpu::BufferUsage::MapWrite, true); +#ifdef GGML_WEBGPU_GPU_PROFILE + // Initialize buffer pool for timestamp queries, used for profiling + webgpu_ctx->timestamp_query_buf_pool.init( + webgpu_ctx->global_ctx->device, WEBGPU_NUM_TIMESTAMP_QUERY_BUFS, WEBGPU_TIMESTAMP_QUERY_BUF_SIZE_BYTES, + wgpu::BufferUsage::QueryResolve | wgpu::BufferUsage::CopySrc, + wgpu::BufferUsage::MapRead | wgpu::BufferUsage::CopyDst); +#endif ggml_webgpu_create_buffer(webgpu_ctx->global_ctx->device, webgpu_ctx->set_rows_dev_error_buf, WEBGPU_SET_ROWS_ERROR_BUF_SIZE_BYTES, wgpu::BufferUsage::Storage | wgpu::BufferUsage::CopySrc, "set_rows_dev_error_buf"); From bb8eafc2dcb3663d91edd4c24eeff7963701681c Mon Sep 17 00:00:00 2001 From: Nikhil Jain Date: Wed, 18 Mar 2026 11:12:07 -0700 Subject: [PATCH 4/9] Move profiling buffers and info to per-thread state --- ggml/src/ggml-webgpu/ggml-webgpu.cpp | 222 ++++++++++++++++++--------- 1 file changed, 148 insertions(+), 74 deletions(-) diff --git a/ggml/src/ggml-webgpu/ggml-webgpu.cpp b/ggml/src/ggml-webgpu/ggml-webgpu.cpp index 20efae6e8bd..9cac8b141ac 100644 --- a/ggml/src/ggml-webgpu/ggml-webgpu.cpp +++ b/ggml/src/ggml-webgpu/ggml-webgpu.cpp @@ -370,14 +370,6 @@ struct webgpu_global_context_struct { std::unordered_map cpu_detail_ms; #endif -// #ifdef GGML_WEBGPU_GPU_PROFILE -// // Profiling: per-shader GPU time in ms -// std::unordered_map shader_gpu_time_ms; -// // Profiling: pool of timestamp query buffers (one per operation) -// webgpu_profile_buf_pool timestamp_query_buf_pool; -// std::vector profiling_futures; -// #endif - #ifdef GGML_WEBGPU_DEBUG wgpu::Buffer debug_host_buf; wgpu::Buffer debug_dev_buf; @@ -586,9 +578,6 @@ static void ggml_backend_webgpu_wait(webgpu_global_context & ctx, while (blocking_wait) { auto waitStatus = ctx->instance.WaitAny(1, &subs[0].submit_done, 0); if (ggml_backend_webgpu_handle_wait_status(waitStatus, true)) { -#ifdef GGML_WEBGPU_GPU_PROFILE - ggml_backend_webgpu_wait_profile_futures(ctx, subs[0].profile_futures, true); -#endif subs.erase(subs.begin()); } blocking_wait = (block && !subs.empty()) || subs.size() >= WEBGPU_MAX_INFLIGHT_SUBS_PER_THREAD; @@ -602,12 +591,7 @@ static void ggml_backend_webgpu_wait(webgpu_global_context & ctx, for (auto sub = subs.begin(); sub != subs.end();) { auto waitStatus = ctx->instance.WaitAny(1, &sub->submit_done, 0); bool success = ggml_backend_webgpu_handle_wait_status(waitStatus, true); -#ifdef GGML_WEBGPU_GPU_PROFILE - ggml_backend_webgpu_wait_profile_futures(ctx, sub->profile_futures, false); - if (success && sub->profile_futures.empty()) { -#else if (success) { -#endif sub = subs.erase(sub); } else { ++sub; @@ -646,7 +630,35 @@ static void ggml_backend_webgpu_debug(webgpu_global_context & ctx) { } #endif -static webgpu_submission ggml_backend_webgpu_submit(webgpu_global_context & ctx, +// buffer operations don't need profiling and don't have the thread level storage required for GPU profiling +static webgpu_submission ggml_backend_webgpu_submit_no_profile(webgpu_global_context & ctx, + std::vector & commands, + webgpu_buf_pool & param_buf_pool) { + std::vector command_buffers; + std::vector params_bufs; + webgpu_submission submission; + + for (const auto & command : commands) { + command_buffers.push_back(command.commands); + params_bufs.insert(params_bufs.end(), command.params_bufs.begin(), command.params_bufs.end()); + } + ctx->queue.Submit(command_buffers.size(), command_buffers.data()); + + wgpu::Future p_f = ctx->queue.OnSubmittedWorkDone( + wgpu::CallbackMode::AllowSpontaneous, + [¶m_buf_pool, params_bufs](wgpu::QueueWorkDoneStatus status, wgpu::StringView message) { + if (status != wgpu::QueueWorkDoneStatus::Success) { + GGML_LOG_ERROR("ggml_webgpu: Failed to submit commands: %s\n", std::string(message).c_str()); + } + // Free the staged buffers + param_buf_pool.free_bufs(params_bufs); + }); + submission.submit_done = { p_f }; + + return submission; +} + +static webgpu_submission ggml_backend_webgpu_submit(webgpu_context & ctx, std::vector & commands, webgpu_buf_pool & param_buf_pool) { std::vector command_buffers; @@ -660,9 +672,9 @@ static webgpu_submission ggml_backend_webgpu_submit(webgpu_global_context & command_buffers.push_back(command.commands); params_bufs.insert(params_bufs.end(), command.params_bufs.begin(), command.params_bufs.end()); } - ctx->queue.Submit(command_buffers.size(), command_buffers.data()); + ctx->global_ctx->queue.Submit(command_buffers.size(), command_buffers.data()); - wgpu::Future p_f = ctx->queue.OnSubmittedWorkDone( + wgpu::Future p_f = ctx->global_ctx->queue.OnSubmittedWorkDone( wgpu::CallbackMode::AllowSpontaneous, [¶m_buf_pool, params_bufs](wgpu::QueueWorkDoneStatus status, wgpu::StringView message) { if (status != wgpu::QueueWorkDoneStatus::Success) { @@ -699,7 +711,7 @@ static webgpu_submission ggml_backend_webgpu_submit(webgpu_global_context & } static webgpu_command ggml_backend_webgpu_build_multi( - webgpu_global_context & ctx, + webgpu_context & ctx, webgpu_buf_pool & param_buf_pool, const std::vector & pipelines, const std::vector> & params_list, @@ -725,14 +737,15 @@ static webgpu_command ggml_backend_webgpu_build_multi( bind_group_desc.entryCount = entries.size(); bind_group_desc.entries = entries.data(); bind_group_desc.label = pipelines[i].name.c_str(); - bind_groups.push_back(ctx->device.CreateBindGroup(&bind_group_desc)); + bind_groups.push_back(ctx->global_ctx->device.CreateBindGroup(&bind_group_desc)); params_bufs_list.push_back(params_bufs); } - wgpu::CommandEncoder encoder = ctx->device.CreateCommandEncoder(); + wgpu::CommandEncoder encoder = ctx->global_ctx->device.CreateCommandEncoder(); for (size_t i = 0; i < params_bufs_list.size(); i++) { - ctx->queue.WriteBuffer(params_bufs_list[i], 0, params_list[i].data(), params_list[i].size() * sizeof(uint32_t)); + ctx->global_ctx->queue.WriteBuffer(params_bufs_list[i], 0, params_list[i].data(), + params_list[i].size() * sizeof(uint32_t)); } #ifdef GGML_WEBGPU_GPU_PROFILE @@ -744,7 +757,7 @@ static webgpu_command ggml_backend_webgpu_build_multi( } // buffers are freed within the future, and futures are deleted here - ggml_backend_webgpu_wait_profile_futures(ctx, ctx->profiling_futures, true); + ggml_backend_webgpu_wait_profile_futures(ctx->global_ctx, ctx->profiling_futures, true); } webgpu_profile_bufs ts_bufs = ctx->timestamp_query_buf_pool.alloc_bufs(); @@ -780,7 +793,60 @@ static webgpu_command ggml_backend_webgpu_build_multi( return result; } -static webgpu_command ggml_backend_webgpu_build(webgpu_global_context & ctx, +static webgpu_command ggml_backend_webgpu_build_multi_no_profile( + webgpu_global_context & ctx, + webgpu_buf_pool & param_buf_pool, + const std::vector & pipelines, + const std::vector> & params_list, + const std::vector> & bind_group_entries_list, + const std::vector> & workgroups_list) { + GGML_ASSERT(pipelines.size() == params_list.size()); + GGML_ASSERT(pipelines.size() == bind_group_entries_list.size()); + GGML_ASSERT(pipelines.size() == workgroups_list.size()); + + std::vector params_bufs_list; + std::vector bind_groups; + + for (size_t i = 0; i < pipelines.size(); i++) { + wgpu::Buffer params_bufs = param_buf_pool.alloc_bufs(); + + std::vector entries = bind_group_entries_list[i]; + uint32_t params_binding_num = entries.size(); + entries.push_back( + { .binding = params_binding_num, .buffer = params_bufs, .offset = 0, .size = params_bufs.GetSize() }); + + wgpu::BindGroupDescriptor bind_group_desc; + bind_group_desc.layout = pipelines[i].pipeline.GetBindGroupLayout(0); + bind_group_desc.entryCount = entries.size(); + bind_group_desc.entries = entries.data(); + bind_group_desc.label = pipelines[i].name.c_str(); + bind_groups.push_back(ctx->device.CreateBindGroup(&bind_group_desc)); + + params_bufs_list.push_back(params_bufs); + } + + wgpu::CommandEncoder encoder = ctx->device.CreateCommandEncoder(); + for (size_t i = 0; i < params_bufs_list.size(); i++) { + ctx->queue.WriteBuffer(params_bufs_list[i], 0, params_list[i].data(), params_list[i].size() * sizeof(uint32_t)); + } + + wgpu::ComputePassEncoder pass = encoder.BeginComputePass(); + for (size_t i = 0; i < pipelines.size(); i++) { + pass.SetPipeline(pipelines[i].pipeline); + pass.SetBindGroup(0, bind_groups[i]); + pass.DispatchWorkgroups(workgroups_list[i].first, workgroups_list[i].second, 1); + } + pass.End(); + + wgpu::CommandBuffer commands = encoder.Finish(); + webgpu_command result = {}; + result.commands = commands; + result.params_bufs = params_bufs_list; + result.num_kernels = pipelines.size(); + return result; +} + +static webgpu_command ggml_backend_webgpu_build(webgpu_context & ctx, webgpu_buf_pool & param_buf_pool, webgpu_pipeline & pipeline, std::vector params, @@ -795,6 +861,22 @@ static webgpu_command ggml_backend_webgpu_build(webgpu_global_context & { { wg_x, wg_y } }); } +// buffer operations don't need profiling and don't have the thread level storage required for GPU profiling +static webgpu_command ggml_backend_webgpu_build_no_profile(webgpu_global_context & ctx, + webgpu_buf_pool & param_buf_pool, + webgpu_pipeline & pipeline, + std::vector params, + std::vector bind_group_entries, + uint32_t wg_x, + uint32_t wg_y = 1) { + return ggml_backend_webgpu_build_multi_no_profile(ctx, param_buf_pool, + { + pipeline + }, + { std::move(params) }, { std::move(bind_group_entries) }, + { { wg_x, wg_y } }); +} + static void ggml_backend_webgpu_buffer_memset(webgpu_global_context & ctx, wgpu::Buffer & buf, uint32_t value, @@ -807,10 +889,10 @@ static void ggml_backend_webgpu_buffer_memset(webgpu_global_context & ctx, size_t bytes_per_wg = WEBGPU_MAX_WG_SIZE * ctx->capabilities.memset_bytes_per_thread; uint32_t wg_x = CEIL_DIV(size + 3, bytes_per_wg); - webgpu_command command = - ggml_backend_webgpu_build(ctx, ctx->memset_buf_pool, ctx->memset_pipelines[0], params, entries, wg_x); + webgpu_command command = ggml_backend_webgpu_build_no_profile(ctx, ctx->memset_buf_pool, ctx->memset_pipelines[0], + params, entries, wg_x); std::vector commands = { command }; - std::vector sub = { ggml_backend_webgpu_submit(ctx, commands, ctx->memset_buf_pool) }; + std::vector sub = { ggml_backend_webgpu_submit_no_profile(ctx, commands, ctx->memset_buf_pool) }; ggml_backend_webgpu_wait(ctx, sub); } @@ -856,19 +938,16 @@ static void ggml_backend_webgpu_free(ggml_backend_t backend) { // invalidate iteration while we're working. std::vector stragglers; { - std::lock_guard lock(global_ctx->timestamp_query_buf_pool.mutex); - stragglers = global_ctx->timestamp_query_buf_pool.in_use; + std::lock_guard lock(ctx->webgpu_ctx->timestamp_query_buf_pool.mutex); + stragglers = ctx->webgpu_ctx->timestamp_query_buf_pool.in_use; } - std::cout << "Stragglers size: " << stragglers.size() << "\n"; + std::cout << "Num stragglers: " << stragglers.size() << "\n"; if (!stragglers.empty()) { wgpu::CommandEncoder encoder = global_ctx->device.CreateCommandEncoder(); for (auto & ts_bufs : stragglers) { encoder.ResolveQuerySet(ts_bufs.query_set, 0, 2, ts_bufs.dev_buf, 0); - encoder.CopyBufferToBuffer( - ts_bufs.dev_buf, 0, - ts_bufs.host_buf, 0, - ts_bufs.host_buf.GetSize()); + encoder.CopyBufferToBuffer(ts_bufs.dev_buf, 0, ts_bufs.host_buf, 0, ts_bufs.host_buf.GetSize()); } wgpu::CommandBuffer commands = encoder.Finish(); @@ -879,21 +958,17 @@ static void ggml_backend_webgpu_free(ggml_backend_t backend) { // Now wait for the profiling futures. These futures should already have been // created by MapAsync when the profile buffers were originally queued. - ggml_backend_webgpu_wait_profile_futures( - global_ctx, - global_ctx->profiling_futures, - true); + ggml_backend_webgpu_wait_profile_futures(global_ctx, ctx->webgpu_ctx->profiling_futures, true); } - std::cout << "\n[ggml_webgpu gpu profiling summary]\n"; double total_gpu = 0.0; - for (const auto & kv : ctx->webgpu_ctx->global_ctx->shader_gpu_time_ms) { + for (const auto & kv : ctx->webgpu_ctx->shader_gpu_time_ms) { total_gpu += kv.second; } std::cout << "ggml_webgpu: total gpu time (all shaders): " << total_gpu << " ms\n"; std::cout << "\nggml_webgpu: gpu breakdown:\n"; - for (const auto & kv : ctx->webgpu_ctx->global_ctx->shader_gpu_time_ms) { + for (const auto & kv : ctx->webgpu_ctx->shader_gpu_time_ms) { double pct = (total_gpu > 0.0) ? (kv.second / total_gpu * 100.0) : 0.0; std::cout << "ggml_webgpu: " << kv.first << ": " << kv.second << " ms (" << std::fixed << std::setprecision(2) << pct << "%)\n"; @@ -989,8 +1064,8 @@ static webgpu_command ggml_webgpu_cpy(webgpu_context & ctx, ggml_tensor * src, g }; uint32_t wg_x = CEIL_DIV(ne, WEBGPU_MAX_WG_SIZE); - return ggml_backend_webgpu_build(ctx->global_ctx, ctx->param_buf_pool, ctx->cpy_pipelines[src->type][dst->type], - params, entries, wg_x); + return ggml_backend_webgpu_build(ctx, ctx->param_buf_pool, ctx->cpy_pipelines[src->type][dst->type], params, + entries, wg_x); } static webgpu_command ggml_webgpu_pad(webgpu_context & ctx, ggml_tensor * src, ggml_tensor * dst) { @@ -1045,7 +1120,7 @@ static webgpu_command ggml_webgpu_pad(webgpu_context & ctx, ggml_tensor * src, g }; uint32_t wg_x = CEIL_DIV(ne, decisions->wg_size); - return ggml_backend_webgpu_build(ctx->global_ctx, ctx->param_buf_pool, pipeline, params, entries, wg_x); + return ggml_backend_webgpu_build(ctx, ctx->param_buf_pool, pipeline, params, entries, wg_x); } static std::optional ggml_webgpu_set_rows(webgpu_context & ctx, @@ -1114,7 +1189,7 @@ static std::optional ggml_webgpu_set_rows(webgpu_context & ctx, threads = src->ne[0] * src->ne[1] * src->ne[2] * src->ne[3]; } uint32_t wg_x = CEIL_DIV(threads, decisions->wg_size); - return ggml_backend_webgpu_build(ctx->global_ctx, ctx->param_buf_pool, pipeline, params, entries, wg_x, 1); + return ggml_backend_webgpu_build(ctx, ctx->param_buf_pool, pipeline, params, entries, wg_x, 1); } // Workgroup size is a common constant @@ -1175,7 +1250,7 @@ static webgpu_command ggml_webgpu_get_rows(webgpu_context & ctx, uint32_t wg_x = CEIL_DIV(dst->ne[1] * dst->ne[2] * dst->ne[3], decisions->wg_size); - return ggml_backend_webgpu_build(ctx->global_ctx, ctx->param_buf_pool, pipeline, params, entries, wg_x); + return ggml_backend_webgpu_build(ctx, ctx->param_buf_pool, pipeline, params, entries, wg_x); } static webgpu_command ggml_webgpu_mul_mat(webgpu_context & ctx, @@ -1320,7 +1395,7 @@ static webgpu_command ggml_webgpu_mul_mat(webgpu_context & ctx, compute_2d_workgroups(total_wg, max_wg_per_dim, wg_x, wg_y); } - return ggml_backend_webgpu_build(ctx->global_ctx, ctx->param_buf_pool, pipeline, params, entries, wg_x, wg_y); + return ggml_backend_webgpu_build(ctx, ctx->param_buf_pool, pipeline, params, entries, wg_x, wg_y); } #ifndef __EMSCRIPTEN__ @@ -1428,7 +1503,7 @@ static webgpu_command ggml_webgpu_flash_attn(webgpu_context & ctx, uint32_t wg_per_head = CEIL_DIV(Q->ne[1], decisions->q_tile); uint32_t wg_x = wg_per_head * Q->ne[2] * Q->ne[3]; // wg per head * number of heads * number of batches - return ggml_backend_webgpu_build(ctx->global_ctx, ctx->param_buf_pool, pipeline, params, entries, wg_x); + return ggml_backend_webgpu_build(ctx, ctx->param_buf_pool, pipeline, params, entries, wg_x); } #endif @@ -1507,7 +1582,7 @@ static webgpu_command ggml_webgpu_unary_op(webgpu_context & ctx, ggml_tensor * s } uint32_t wg_x = CEIL_DIV(ne, decisions->wg_size); - return ggml_backend_webgpu_build(ctx->global_ctx, ctx->param_buf_pool, pipeline, params, entries, wg_x); + return ggml_backend_webgpu_build(ctx, ctx->param_buf_pool, pipeline, params, entries, wg_x); } static webgpu_command ggml_webgpu_binary_op(webgpu_context & ctx, @@ -1609,7 +1684,7 @@ static webgpu_command ggml_webgpu_binary_op(webgpu_context & ctx, } uint32_t wg_x = CEIL_DIV(ne, decisions->wg_size); - return ggml_backend_webgpu_build(ctx->global_ctx, ctx->param_buf_pool, pipeline, params, entries, wg_x); + return ggml_backend_webgpu_build(ctx, ctx->param_buf_pool, pipeline, params, entries, wg_x); } static webgpu_command ggml_webgpu_concat(webgpu_context & ctx, @@ -1665,7 +1740,7 @@ static webgpu_command ggml_webgpu_concat(webgpu_context & ctx, webgpu_pipeline pipeline = ctx->shader_lib->get_concat_pipeline(shader_lib_ctx); auto * decisions = static_cast(pipeline.context.get()); uint32_t wg_x = CEIL_DIV(ne, decisions->wg_size); - return ggml_backend_webgpu_build(ctx->global_ctx, ctx->param_buf_pool, pipeline, params, entries, wg_x); + return ggml_backend_webgpu_build(ctx, ctx->param_buf_pool, pipeline, params, entries, wg_x); } static webgpu_command ggml_webgpu_repeat(webgpu_context & ctx, ggml_tensor * src0, ggml_tensor * dst) { @@ -1707,7 +1782,7 @@ static webgpu_command ggml_webgpu_repeat(webgpu_context & ctx, ggml_tensor * src webgpu_pipeline pipeline = ctx->shader_lib->get_repeat_pipeline(shader_lib_ctx); auto * decisions = static_cast(pipeline.context.get()); uint32_t wg_x = CEIL_DIV(ne, decisions->wg_size); - return ggml_backend_webgpu_build(ctx->global_ctx, ctx->param_buf_pool, pipeline, params, entries, wg_x); + return ggml_backend_webgpu_build(ctx, ctx->param_buf_pool, pipeline, params, entries, wg_x); } static webgpu_command ggml_webgpu_rms_norm(webgpu_context & ctx, ggml_tensor * src, ggml_tensor * dst) { @@ -1742,8 +1817,8 @@ static webgpu_command ggml_webgpu_rms_norm(webgpu_context & ctx, ggml_tensor * s .size = ggml_webgpu_tensor_binding_size(ctx, dst) }); } - return ggml_backend_webgpu_build(ctx->global_ctx, ctx->param_buf_pool, ctx->rms_norm_pipelines[inplace], params, - entries, ggml_nrows(src)); + return ggml_backend_webgpu_build(ctx, ctx->param_buf_pool, ctx->rms_norm_pipelines[inplace], params, entries, + ggml_nrows(src)); } static webgpu_command ggml_webgpu_rope(webgpu_context & ctx, @@ -1835,7 +1910,7 @@ static webgpu_command ggml_webgpu_rope(webgpu_context & ctx, webgpu_pipeline pipeline = ctx->rope_pipelines[dst->type][has_freq_factor][inplace]; uint32_t wg_x = CEIL_DIV(ggml_nelements(dst), WEBGPU_MAX_WG_SIZE); - return ggml_backend_webgpu_build(ctx->global_ctx, ctx->param_buf_pool, pipeline, params, entries, wg_x); + return ggml_backend_webgpu_build(ctx, ctx->param_buf_pool, pipeline, params, entries, wg_x); } static webgpu_command ggml_webgpu_glu(webgpu_context & ctx, ggml_tensor * src0, ggml_tensor * src1, ggml_tensor * dst) { @@ -1887,7 +1962,7 @@ static webgpu_command ggml_webgpu_glu(webgpu_context & ctx, ggml_tensor * src0, webgpu_pipeline pipeline = ctx->glu_pipelines[ggml_get_glu_op(dst)][dst->type][split]; uint32_t wg_x = CEIL_DIV(ggml_nelements(dst), WEBGPU_MAX_WG_SIZE); - return ggml_backend_webgpu_build(ctx->global_ctx, ctx->param_buf_pool, pipeline, params, entries, wg_x); + return ggml_backend_webgpu_build(ctx, ctx->param_buf_pool, pipeline, params, entries, wg_x); } static webgpu_command ggml_webgpu_scale(webgpu_context & ctx, ggml_tensor * src, ggml_tensor * dst) { @@ -1938,7 +2013,7 @@ static webgpu_command ggml_webgpu_scale(webgpu_context & ctx, ggml_tensor * src, } uint32_t wg_x = CEIL_DIV(ggml_nelements(dst), decisions->wg_size); - return ggml_backend_webgpu_build(ctx->global_ctx, ctx->param_buf_pool, pipeline, params, entries, wg_x); + return ggml_backend_webgpu_build(ctx, ctx->param_buf_pool, pipeline, params, entries, wg_x); } static webgpu_command ggml_webgpu_soft_max(webgpu_context & ctx, @@ -2010,9 +2085,8 @@ static webgpu_command ggml_webgpu_soft_max(webgpu_context & ctx, .size = ggml_webgpu_tensor_binding_size(ctx, dst) }); } - return ggml_backend_webgpu_build(ctx->global_ctx, ctx->param_buf_pool, - ctx->soft_max_pipelines[mask_type][has_sink][inplace], params, entries, - ggml_nrows(dst)); + return ggml_backend_webgpu_build(ctx, ctx->param_buf_pool, ctx->soft_max_pipelines[mask_type][has_sink][inplace], + params, entries, ggml_nrows(dst)); } static webgpu_command ggml_webgpu_argmax(webgpu_context & ctx, ggml_tensor * src, ggml_tensor * dst) { @@ -2037,7 +2111,7 @@ static webgpu_command ggml_webgpu_argmax(webgpu_context & ctx, ggml_tensor * src webgpu_pipeline pipeline = ctx->shader_lib->get_argmax_pipeline(shader_lib_ctx); uint32_t wg_x = ggml_nelements(dst); - return ggml_backend_webgpu_build(ctx->global_ctx, ctx->param_buf_pool, pipeline, params, entries, wg_x); + return ggml_backend_webgpu_build(ctx, ctx->param_buf_pool, pipeline, params, entries, wg_x); } static webgpu_command ggml_webgpu_argsort(webgpu_context & ctx, ggml_tensor * src, ggml_tensor * dst) { @@ -2131,8 +2205,8 @@ static webgpu_command ggml_webgpu_argsort(webgpu_context & ctx, ggml_tensor * sr workgroups_list.push_back({ wg_x_init, wg_y_init }); if (merge_passes == 0) { - return ggml_backend_webgpu_build_multi(ctx->global_ctx, ctx->param_buf_pool, pipelines, params_list, - entries_list, workgroups_list); + return ggml_backend_webgpu_build_multi(ctx, ctx->param_buf_pool, pipelines, params_list, entries_list, + workgroups_list); } bool in_is_tmp = start_in_tmp; @@ -2193,7 +2267,7 @@ static webgpu_command ggml_webgpu_argsort(webgpu_context & ctx, ggml_tensor * sr in_is_tmp = !in_is_tmp; } - return ggml_backend_webgpu_build_multi(ctx->global_ctx, ctx->param_buf_pool, pipelines, params_list, entries_list, + return ggml_backend_webgpu_build_multi(ctx, ctx->param_buf_pool, pipelines, params_list, entries_list, workgroups_list); } @@ -2222,7 +2296,7 @@ static webgpu_command ggml_webgpu_cumsum(webgpu_context & ctx, ggml_tensor * src webgpu_pipeline pipeline = ctx->shader_lib->get_cumsum_pipeline(shader_lib_ctx); uint32_t wg_x = ggml_nrows(dst); - return ggml_backend_webgpu_build(ctx->global_ctx, ctx->param_buf_pool, pipeline, params, entries, wg_x); + return ggml_backend_webgpu_build(ctx, ctx->param_buf_pool, pipeline, params, entries, wg_x); } static webgpu_command ggml_webgpu_sum_rows(webgpu_context & ctx, ggml_tensor * src, ggml_tensor * dst) { @@ -2254,7 +2328,7 @@ static webgpu_command ggml_webgpu_sum_rows(webgpu_context & ctx, ggml_tensor * s webgpu_pipeline pipeline = ctx->shader_lib->get_sum_rows_pipeline(shader_lib_ctx); uint32_t wg_x = total_sum ? 1 : ggml_nrows(dst); - return ggml_backend_webgpu_build(ctx->global_ctx, ctx->param_buf_pool, pipeline, params, entries, wg_x); + return ggml_backend_webgpu_build(ctx, ctx->param_buf_pool, pipeline, params, entries, wg_x); } // Returns the encoded command, or std::nullopt if the operation is a no-op @@ -2364,7 +2438,7 @@ static ggml_status ggml_backend_webgpu_graph_compute(ggml_backend_t backend, str if (num_batched_kernels >= WEBGPU_COMMAND_SUBMIT_BATCH_SIZE) { num_batched_kernels = 0; - subs.push_back(ggml_backend_webgpu_submit(ctx->global_ctx, commands, ctx->param_buf_pool)); + subs.push_back(ggml_backend_webgpu_submit(ctx, commands, ctx->param_buf_pool)); // Process events and check for completed submissions ctx->global_ctx->instance.ProcessEvents(); ggml_backend_webgpu_wait(ctx->global_ctx, subs, false); @@ -2372,7 +2446,7 @@ static ggml_status ggml_backend_webgpu_graph_compute(ggml_backend_t backend, str } } if (!commands.empty()) { - subs.push_back(ggml_backend_webgpu_submit(ctx->global_ctx, commands, ctx->param_buf_pool)); + subs.push_back(ggml_backend_webgpu_submit(ctx, commands, ctx->param_buf_pool)); commands.clear(); } @@ -3005,10 +3079,10 @@ static webgpu_context initialize_webgpu_context(ggml_backend_dev_t dev) { wgpu::BufferUsage::CopySrc | wgpu::BufferUsage::MapWrite, true); #ifdef GGML_WEBGPU_GPU_PROFILE // Initialize buffer pool for timestamp queries, used for profiling - webgpu_ctx->timestamp_query_buf_pool.init( - webgpu_ctx->global_ctx->device, WEBGPU_NUM_TIMESTAMP_QUERY_BUFS, WEBGPU_TIMESTAMP_QUERY_BUF_SIZE_BYTES, - wgpu::BufferUsage::QueryResolve | wgpu::BufferUsage::CopySrc, - wgpu::BufferUsage::MapRead | wgpu::BufferUsage::CopyDst); + webgpu_ctx->timestamp_query_buf_pool.init(webgpu_ctx->global_ctx->device, WEBGPU_NUM_TIMESTAMP_QUERY_BUFS, + WEBGPU_TIMESTAMP_QUERY_BUF_SIZE_BYTES, + wgpu::BufferUsage::QueryResolve | wgpu::BufferUsage::CopySrc, + wgpu::BufferUsage::MapRead | wgpu::BufferUsage::CopyDst); #endif ggml_webgpu_create_buffer(webgpu_ctx->global_ctx->device, webgpu_ctx->set_rows_dev_error_buf, WEBGPU_SET_ROWS_ERROR_BUF_SIZE_BYTES, @@ -3069,7 +3143,7 @@ static ggml_backend_buffer_type_t ggml_backend_webgpu_device_get_buffer_type(ggm /* .is_host = */ NULL, // defaults to false }, /* .device = */ - dev, + dev, /* .context = */ NULL }; From 8344f6068c6dff5d80b2d698a5aadc9c1942e4ce Mon Sep 17 00:00:00 2001 From: Nikhil Jain Date: Wed, 18 Mar 2026 11:36:55 -0700 Subject: [PATCH 5/9] Clean up buffer path for profiling --- ggml/src/ggml-webgpu/ggml-webgpu.cpp | 277 ++++++++++----------------- 1 file changed, 99 insertions(+), 178 deletions(-) diff --git a/ggml/src/ggml-webgpu/ggml-webgpu.cpp b/ggml/src/ggml-webgpu/ggml-webgpu.cpp index 9cac8b141ac..bf0cbebeb51 100644 --- a/ggml/src/ggml-webgpu/ggml-webgpu.cpp +++ b/ggml/src/ggml-webgpu/ggml-webgpu.cpp @@ -630,35 +630,8 @@ static void ggml_backend_webgpu_debug(webgpu_global_context & ctx) { } #endif -// buffer operations don't need profiling and don't have the thread level storage required for GPU profiling -static webgpu_submission ggml_backend_webgpu_submit_no_profile(webgpu_global_context & ctx, - std::vector & commands, - webgpu_buf_pool & param_buf_pool) { - std::vector command_buffers; - std::vector params_bufs; - webgpu_submission submission; - - for (const auto & command : commands) { - command_buffers.push_back(command.commands); - params_bufs.insert(params_bufs.end(), command.params_bufs.begin(), command.params_bufs.end()); - } - ctx->queue.Submit(command_buffers.size(), command_buffers.data()); - - wgpu::Future p_f = ctx->queue.OnSubmittedWorkDone( - wgpu::CallbackMode::AllowSpontaneous, - [¶m_buf_pool, params_bufs](wgpu::QueueWorkDoneStatus status, wgpu::StringView message) { - if (status != wgpu::QueueWorkDoneStatus::Success) { - GGML_LOG_ERROR("ggml_webgpu: Failed to submit commands: %s\n", std::string(message).c_str()); - } - // Free the staged buffers - param_buf_pool.free_bufs(params_bufs); - }); - submission.submit_done = { p_f }; - - return submission; -} - -static webgpu_submission ggml_backend_webgpu_submit(webgpu_context & ctx, +static webgpu_submission ggml_backend_webgpu_submit(webgpu_global_context & global_ctx, + std::optional ctx_opt, std::vector & commands, webgpu_buf_pool & param_buf_pool) { std::vector command_buffers; @@ -672,9 +645,9 @@ static webgpu_submission ggml_backend_webgpu_submit(webgpu_context & command_buffers.push_back(command.commands); params_bufs.insert(params_bufs.end(), command.params_bufs.begin(), command.params_bufs.end()); } - ctx->global_ctx->queue.Submit(command_buffers.size(), command_buffers.data()); + global_ctx->queue.Submit(command_buffers.size(), command_buffers.data()); - wgpu::Future p_f = ctx->global_ctx->queue.OnSubmittedWorkDone( + wgpu::Future p_f = global_ctx->queue.OnSubmittedWorkDone( wgpu::CallbackMode::AllowSpontaneous, [¶m_buf_pool, params_bufs](wgpu::QueueWorkDoneStatus status, wgpu::StringView message) { if (status != wgpu::QueueWorkDoneStatus::Success) { @@ -686,32 +659,37 @@ static webgpu_submission ggml_backend_webgpu_submit(webgpu_context & submission.submit_done = { p_f }; #ifdef GGML_WEBGPU_GPU_PROFILE - for (const auto & command : commands) { - auto label = command.pipeline_name; - auto ts_bufs = command.timestamp_query_bufs; - - wgpu::Future f = ts_bufs.host_buf.MapAsync( - wgpu::MapMode::Read, 0, ts_bufs.host_buf.GetSize(), wgpu::CallbackMode::AllowSpontaneous, - [ctx, ts_bufs, label](wgpu::MapAsyncStatus status, wgpu::StringView message) { - if (status != wgpu::MapAsyncStatus::Success) { - GGML_LOG_ERROR("ggml_webgpu: Failed to map timestamp buffer: %s\n", std::string(message).c_str()); - } else { - const uint64_t * ts_data = (const uint64_t *) ts_bufs.host_buf.GetConstMappedRange(); - // WebGPU timestamps are in ns; convert to ms - double elapsed_ms = double(ts_data[1] - ts_data[0]) * 1e-6; - ctx->shader_gpu_time_ms[label] += elapsed_ms; - } - // We can't unmap in here due to WebGPU reentrancy limitations. - ctx->timestamp_query_buf_pool.free_bufs({ ts_bufs }); - }); - ctx->profiling_futures.push_back({ f, false }); + if (ctx_opt.has_value()) { + auto ctx = ctx_opt.value(); + for (const auto & command : commands) { + auto label = command.pipeline_name; + auto ts_bufs = command.timestamp_query_bufs; + + wgpu::Future f = ts_bufs.host_buf.MapAsync( + wgpu::MapMode::Read, 0, ts_bufs.host_buf.GetSize(), wgpu::CallbackMode::AllowSpontaneous, + [ctx, ts_bufs, label](wgpu::MapAsyncStatus status, wgpu::StringView message) { + if (status != wgpu::MapAsyncStatus::Success) { + GGML_LOG_ERROR("ggml_webgpu: Failed to map timestamp buffer: %s\n", + std::string(message).c_str()); + } else { + const uint64_t * ts_data = (const uint64_t *) ts_bufs.host_buf.GetConstMappedRange(); + // WebGPU timestamps are in ns; convert to ms + double elapsed_ms = double(ts_data[1] - ts_data[0]) * 1e-6; + ctx->shader_gpu_time_ms[label] += elapsed_ms; + } + // We can't unmap in here due to WebGPU reentrancy limitations. + ctx->timestamp_query_buf_pool.free_bufs({ ts_bufs }); + }); + ctx->profiling_futures.push_back({ f, false }); + } } #endif return submission; } static webgpu_command ggml_backend_webgpu_build_multi( - webgpu_context & ctx, + webgpu_global_context & global_ctx, + std::optional ctx_opt, webgpu_buf_pool & param_buf_pool, const std::vector & pipelines, const std::vector> & params_list, @@ -737,41 +715,49 @@ static webgpu_command ggml_backend_webgpu_build_multi( bind_group_desc.entryCount = entries.size(); bind_group_desc.entries = entries.data(); bind_group_desc.label = pipelines[i].name.c_str(); - bind_groups.push_back(ctx->global_ctx->device.CreateBindGroup(&bind_group_desc)); + bind_groups.push_back(global_ctx->device.CreateBindGroup(&bind_group_desc)); params_bufs_list.push_back(params_bufs); } - wgpu::CommandEncoder encoder = ctx->global_ctx->device.CreateCommandEncoder(); + wgpu::CommandEncoder encoder = global_ctx->device.CreateCommandEncoder(); for (size_t i = 0; i < params_bufs_list.size(); i++) { - ctx->global_ctx->queue.WriteBuffer(params_bufs_list[i], 0, params_list[i].data(), - params_list[i].size() * sizeof(uint32_t)); + global_ctx->queue.WriteBuffer(params_bufs_list[i], 0, params_list[i].data(), + params_list[i].size() * sizeof(uint32_t)); } + wgpu::ComputePassEncoder pass; #ifdef GGML_WEBGPU_GPU_PROFILE - // Only copy timestamp query results if we have no more timestamp buffers available - if (ctx->timestamp_query_buf_pool.num_free_bufs() == 0) { - for (auto & ts_bufs : ctx->timestamp_query_buf_pool.in_use) { - encoder.ResolveQuerySet(ts_bufs.query_set, 0, 2, ts_bufs.dev_buf, 0); - encoder.CopyBufferToBuffer(ts_bufs.dev_buf, 0, ts_bufs.host_buf, 0, ts_bufs.host_buf.GetSize()); + webgpu_profile_bufs ts_bufs; + if (ctx_opt.has_value()) { + auto ctx = ctx_opt.value(); + // Only copy timestamp query results if we have no more timestamp buffers available + if (ctx->timestamp_query_buf_pool.num_free_bufs() == 0) { + for (auto & ts_bufs : ctx->timestamp_query_buf_pool.in_use) { + encoder.ResolveQuerySet(ts_bufs.query_set, 0, 2, ts_bufs.dev_buf, 0); + encoder.CopyBufferToBuffer(ts_bufs.dev_buf, 0, ts_bufs.host_buf, 0, ts_bufs.host_buf.GetSize()); + } + + // buffers are freed within the future, and futures are deleted here + ggml_backend_webgpu_wait_profile_futures(global_ctx, ctx->profiling_futures, true); } - // buffers are freed within the future, and futures are deleted here - ggml_backend_webgpu_wait_profile_futures(ctx->global_ctx, ctx->profiling_futures, true); - } + ts_bufs = ctx->timestamp_query_buf_pool.alloc_bufs(); + if (ts_bufs.host_buf.GetMapState() == wgpu::BufferMapState::Mapped) { + ts_bufs.host_buf.Unmap(); + } - webgpu_profile_bufs ts_bufs = ctx->timestamp_query_buf_pool.alloc_bufs(); - if (ts_bufs.host_buf.GetMapState() == wgpu::BufferMapState::Mapped) { - ts_bufs.host_buf.Unmap(); + wgpu::PassTimestampWrites ts_writes = { .querySet = ts_bufs.query_set, + .beginningOfPassWriteIndex = 0, + .endOfPassWriteIndex = 1 }; + wgpu::ComputePassDescriptor pass_desc = { .timestampWrites = &ts_writes }; + pass = encoder.BeginComputePass(&pass_desc); + } else { + // buffer operations do not have the thread level storage for profiling, so they fall into this case + pass = encoder.BeginComputePass(); } - - wgpu::PassTimestampWrites ts_writes = { .querySet = ts_bufs.query_set, - .beginningOfPassWriteIndex = 0, - .endOfPassWriteIndex = 1 }; - wgpu::ComputePassDescriptor pass_desc = { .timestampWrites = &ts_writes }; - wgpu::ComputePassEncoder pass = encoder.BeginComputePass(&pass_desc); #else - wgpu::ComputePassEncoder pass = encoder.BeginComputePass(); + pass = encoder.BeginComputePass(); #endif for (size_t i = 0; i < pipelines.size(); i++) { pass.SetPipeline(pipelines[i].pipeline); @@ -786,74 +772,24 @@ static webgpu_command ggml_backend_webgpu_build_multi( result.params_bufs = params_bufs_list; result.num_kernels = pipelines.size(); #ifdef GGML_WEBGPU_GPU_PROFILE - result.timestamp_query_bufs = ts_bufs; - // TODO: handle multiple pipeline names - result.pipeline_name = pipelines.front().name; -#endif - return result; -} - -static webgpu_command ggml_backend_webgpu_build_multi_no_profile( - webgpu_global_context & ctx, - webgpu_buf_pool & param_buf_pool, - const std::vector & pipelines, - const std::vector> & params_list, - const std::vector> & bind_group_entries_list, - const std::vector> & workgroups_list) { - GGML_ASSERT(pipelines.size() == params_list.size()); - GGML_ASSERT(pipelines.size() == bind_group_entries_list.size()); - GGML_ASSERT(pipelines.size() == workgroups_list.size()); - - std::vector params_bufs_list; - std::vector bind_groups; - - for (size_t i = 0; i < pipelines.size(); i++) { - wgpu::Buffer params_bufs = param_buf_pool.alloc_bufs(); - - std::vector entries = bind_group_entries_list[i]; - uint32_t params_binding_num = entries.size(); - entries.push_back( - { .binding = params_binding_num, .buffer = params_bufs, .offset = 0, .size = params_bufs.GetSize() }); - - wgpu::BindGroupDescriptor bind_group_desc; - bind_group_desc.layout = pipelines[i].pipeline.GetBindGroupLayout(0); - bind_group_desc.entryCount = entries.size(); - bind_group_desc.entries = entries.data(); - bind_group_desc.label = pipelines[i].name.c_str(); - bind_groups.push_back(ctx->device.CreateBindGroup(&bind_group_desc)); - - params_bufs_list.push_back(params_bufs); + if (ctx_opt.has_value()) { + result.timestamp_query_bufs = ts_bufs; + // TODO: handle multiple pipeline names + result.pipeline_name = pipelines.front().name; } - - wgpu::CommandEncoder encoder = ctx->device.CreateCommandEncoder(); - for (size_t i = 0; i < params_bufs_list.size(); i++) { - ctx->queue.WriteBuffer(params_bufs_list[i], 0, params_list[i].data(), params_list[i].size() * sizeof(uint32_t)); - } - - wgpu::ComputePassEncoder pass = encoder.BeginComputePass(); - for (size_t i = 0; i < pipelines.size(); i++) { - pass.SetPipeline(pipelines[i].pipeline); - pass.SetBindGroup(0, bind_groups[i]); - pass.DispatchWorkgroups(workgroups_list[i].first, workgroups_list[i].second, 1); - } - pass.End(); - - wgpu::CommandBuffer commands = encoder.Finish(); - webgpu_command result = {}; - result.commands = commands; - result.params_bufs = params_bufs_list; - result.num_kernels = pipelines.size(); +#endif return result; } -static webgpu_command ggml_backend_webgpu_build(webgpu_context & ctx, +static webgpu_command ggml_backend_webgpu_build(webgpu_global_context & global_ctx, + std::optional ctx_opt, webgpu_buf_pool & param_buf_pool, webgpu_pipeline & pipeline, std::vector params, std::vector bind_group_entries, uint32_t wg_x, uint32_t wg_y = 1) { - return ggml_backend_webgpu_build_multi(ctx, param_buf_pool, + return ggml_backend_webgpu_build_multi(global_ctx, ctx_opt, param_buf_pool, { pipeline }, @@ -861,22 +797,6 @@ static webgpu_command ggml_backend_webgpu_build(webgpu_context & { { wg_x, wg_y } }); } -// buffer operations don't need profiling and don't have the thread level storage required for GPU profiling -static webgpu_command ggml_backend_webgpu_build_no_profile(webgpu_global_context & ctx, - webgpu_buf_pool & param_buf_pool, - webgpu_pipeline & pipeline, - std::vector params, - std::vector bind_group_entries, - uint32_t wg_x, - uint32_t wg_y = 1) { - return ggml_backend_webgpu_build_multi_no_profile(ctx, param_buf_pool, - { - pipeline - }, - { std::move(params) }, { std::move(bind_group_entries) }, - { { wg_x, wg_y } }); -} - static void ggml_backend_webgpu_buffer_memset(webgpu_global_context & ctx, wgpu::Buffer & buf, uint32_t value, @@ -889,10 +809,11 @@ static void ggml_backend_webgpu_buffer_memset(webgpu_global_context & ctx, size_t bytes_per_wg = WEBGPU_MAX_WG_SIZE * ctx->capabilities.memset_bytes_per_thread; uint32_t wg_x = CEIL_DIV(size + 3, bytes_per_wg); - webgpu_command command = ggml_backend_webgpu_build_no_profile(ctx, ctx->memset_buf_pool, ctx->memset_pipelines[0], - params, entries, wg_x); + webgpu_command command = ggml_backend_webgpu_build(ctx, std::nullopt, ctx->memset_buf_pool, + ctx->memset_pipelines[0], params, entries, wg_x); std::vector commands = { command }; - std::vector sub = { ggml_backend_webgpu_submit_no_profile(ctx, commands, ctx->memset_buf_pool) }; + std::vector sub = { ggml_backend_webgpu_submit(ctx, std::nullopt, commands, + ctx->memset_buf_pool) }; ggml_backend_webgpu_wait(ctx, sub); } @@ -941,7 +862,6 @@ static void ggml_backend_webgpu_free(ggml_backend_t backend) { std::lock_guard lock(ctx->webgpu_ctx->timestamp_query_buf_pool.mutex); stragglers = ctx->webgpu_ctx->timestamp_query_buf_pool.in_use; } - std::cout << "Num stragglers: " << stragglers.size() << "\n"; if (!stragglers.empty()) { wgpu::CommandEncoder encoder = global_ctx->device.CreateCommandEncoder(); @@ -1064,8 +984,8 @@ static webgpu_command ggml_webgpu_cpy(webgpu_context & ctx, ggml_tensor * src, g }; uint32_t wg_x = CEIL_DIV(ne, WEBGPU_MAX_WG_SIZE); - return ggml_backend_webgpu_build(ctx, ctx->param_buf_pool, ctx->cpy_pipelines[src->type][dst->type], params, - entries, wg_x); + return ggml_backend_webgpu_build(ctx->global_ctx, ctx, ctx->param_buf_pool, + ctx->cpy_pipelines[src->type][dst->type], params, entries, wg_x); } static webgpu_command ggml_webgpu_pad(webgpu_context & ctx, ggml_tensor * src, ggml_tensor * dst) { @@ -1120,7 +1040,7 @@ static webgpu_command ggml_webgpu_pad(webgpu_context & ctx, ggml_tensor * src, g }; uint32_t wg_x = CEIL_DIV(ne, decisions->wg_size); - return ggml_backend_webgpu_build(ctx, ctx->param_buf_pool, pipeline, params, entries, wg_x); + return ggml_backend_webgpu_build(ctx->global_ctx, ctx, ctx->param_buf_pool, pipeline, params, entries, wg_x); } static std::optional ggml_webgpu_set_rows(webgpu_context & ctx, @@ -1189,7 +1109,7 @@ static std::optional ggml_webgpu_set_rows(webgpu_context & ctx, threads = src->ne[0] * src->ne[1] * src->ne[2] * src->ne[3]; } uint32_t wg_x = CEIL_DIV(threads, decisions->wg_size); - return ggml_backend_webgpu_build(ctx, ctx->param_buf_pool, pipeline, params, entries, wg_x, 1); + return ggml_backend_webgpu_build(ctx->global_ctx, ctx, ctx->param_buf_pool, pipeline, params, entries, wg_x, 1); } // Workgroup size is a common constant @@ -1250,7 +1170,7 @@ static webgpu_command ggml_webgpu_get_rows(webgpu_context & ctx, uint32_t wg_x = CEIL_DIV(dst->ne[1] * dst->ne[2] * dst->ne[3], decisions->wg_size); - return ggml_backend_webgpu_build(ctx, ctx->param_buf_pool, pipeline, params, entries, wg_x); + return ggml_backend_webgpu_build(ctx->global_ctx, ctx, ctx->param_buf_pool, pipeline, params, entries, wg_x); } static webgpu_command ggml_webgpu_mul_mat(webgpu_context & ctx, @@ -1395,7 +1315,7 @@ static webgpu_command ggml_webgpu_mul_mat(webgpu_context & ctx, compute_2d_workgroups(total_wg, max_wg_per_dim, wg_x, wg_y); } - return ggml_backend_webgpu_build(ctx, ctx->param_buf_pool, pipeline, params, entries, wg_x, wg_y); + return ggml_backend_webgpu_build(ctx->global_ctx, ctx, ctx->param_buf_pool, pipeline, params, entries, wg_x, wg_y); } #ifndef __EMSCRIPTEN__ @@ -1503,7 +1423,7 @@ static webgpu_command ggml_webgpu_flash_attn(webgpu_context & ctx, uint32_t wg_per_head = CEIL_DIV(Q->ne[1], decisions->q_tile); uint32_t wg_x = wg_per_head * Q->ne[2] * Q->ne[3]; // wg per head * number of heads * number of batches - return ggml_backend_webgpu_build(ctx, ctx->param_buf_pool, pipeline, params, entries, wg_x); + return ggml_backend_webgpu_build(ctx->global_ctx, ctx, ctx->param_buf_pool, pipeline, params, entries, wg_x); } #endif @@ -1582,7 +1502,7 @@ static webgpu_command ggml_webgpu_unary_op(webgpu_context & ctx, ggml_tensor * s } uint32_t wg_x = CEIL_DIV(ne, decisions->wg_size); - return ggml_backend_webgpu_build(ctx, ctx->param_buf_pool, pipeline, params, entries, wg_x); + return ggml_backend_webgpu_build(ctx->global_ctx, ctx, ctx->param_buf_pool, pipeline, params, entries, wg_x); } static webgpu_command ggml_webgpu_binary_op(webgpu_context & ctx, @@ -1684,7 +1604,7 @@ static webgpu_command ggml_webgpu_binary_op(webgpu_context & ctx, } uint32_t wg_x = CEIL_DIV(ne, decisions->wg_size); - return ggml_backend_webgpu_build(ctx, ctx->param_buf_pool, pipeline, params, entries, wg_x); + return ggml_backend_webgpu_build(ctx->global_ctx, ctx, ctx->param_buf_pool, pipeline, params, entries, wg_x); } static webgpu_command ggml_webgpu_concat(webgpu_context & ctx, @@ -1740,7 +1660,7 @@ static webgpu_command ggml_webgpu_concat(webgpu_context & ctx, webgpu_pipeline pipeline = ctx->shader_lib->get_concat_pipeline(shader_lib_ctx); auto * decisions = static_cast(pipeline.context.get()); uint32_t wg_x = CEIL_DIV(ne, decisions->wg_size); - return ggml_backend_webgpu_build(ctx, ctx->param_buf_pool, pipeline, params, entries, wg_x); + return ggml_backend_webgpu_build(ctx->global_ctx, ctx, ctx->param_buf_pool, pipeline, params, entries, wg_x); } static webgpu_command ggml_webgpu_repeat(webgpu_context & ctx, ggml_tensor * src0, ggml_tensor * dst) { @@ -1782,7 +1702,7 @@ static webgpu_command ggml_webgpu_repeat(webgpu_context & ctx, ggml_tensor * src webgpu_pipeline pipeline = ctx->shader_lib->get_repeat_pipeline(shader_lib_ctx); auto * decisions = static_cast(pipeline.context.get()); uint32_t wg_x = CEIL_DIV(ne, decisions->wg_size); - return ggml_backend_webgpu_build(ctx, ctx->param_buf_pool, pipeline, params, entries, wg_x); + return ggml_backend_webgpu_build(ctx->global_ctx, ctx, ctx->param_buf_pool, pipeline, params, entries, wg_x); } static webgpu_command ggml_webgpu_rms_norm(webgpu_context & ctx, ggml_tensor * src, ggml_tensor * dst) { @@ -1817,8 +1737,8 @@ static webgpu_command ggml_webgpu_rms_norm(webgpu_context & ctx, ggml_tensor * s .size = ggml_webgpu_tensor_binding_size(ctx, dst) }); } - return ggml_backend_webgpu_build(ctx, ctx->param_buf_pool, ctx->rms_norm_pipelines[inplace], params, entries, - ggml_nrows(src)); + return ggml_backend_webgpu_build(ctx->global_ctx, ctx, ctx->param_buf_pool, ctx->rms_norm_pipelines[inplace], + params, entries, ggml_nrows(src)); } static webgpu_command ggml_webgpu_rope(webgpu_context & ctx, @@ -1910,7 +1830,7 @@ static webgpu_command ggml_webgpu_rope(webgpu_context & ctx, webgpu_pipeline pipeline = ctx->rope_pipelines[dst->type][has_freq_factor][inplace]; uint32_t wg_x = CEIL_DIV(ggml_nelements(dst), WEBGPU_MAX_WG_SIZE); - return ggml_backend_webgpu_build(ctx, ctx->param_buf_pool, pipeline, params, entries, wg_x); + return ggml_backend_webgpu_build(ctx->global_ctx, ctx, ctx->param_buf_pool, pipeline, params, entries, wg_x); } static webgpu_command ggml_webgpu_glu(webgpu_context & ctx, ggml_tensor * src0, ggml_tensor * src1, ggml_tensor * dst) { @@ -1962,7 +1882,7 @@ static webgpu_command ggml_webgpu_glu(webgpu_context & ctx, ggml_tensor * src0, webgpu_pipeline pipeline = ctx->glu_pipelines[ggml_get_glu_op(dst)][dst->type][split]; uint32_t wg_x = CEIL_DIV(ggml_nelements(dst), WEBGPU_MAX_WG_SIZE); - return ggml_backend_webgpu_build(ctx, ctx->param_buf_pool, pipeline, params, entries, wg_x); + return ggml_backend_webgpu_build(ctx->global_ctx, ctx, ctx->param_buf_pool, pipeline, params, entries, wg_x); } static webgpu_command ggml_webgpu_scale(webgpu_context & ctx, ggml_tensor * src, ggml_tensor * dst) { @@ -2013,7 +1933,7 @@ static webgpu_command ggml_webgpu_scale(webgpu_context & ctx, ggml_tensor * src, } uint32_t wg_x = CEIL_DIV(ggml_nelements(dst), decisions->wg_size); - return ggml_backend_webgpu_build(ctx, ctx->param_buf_pool, pipeline, params, entries, wg_x); + return ggml_backend_webgpu_build(ctx->global_ctx, ctx, ctx->param_buf_pool, pipeline, params, entries, wg_x); } static webgpu_command ggml_webgpu_soft_max(webgpu_context & ctx, @@ -2085,8 +2005,9 @@ static webgpu_command ggml_webgpu_soft_max(webgpu_context & ctx, .size = ggml_webgpu_tensor_binding_size(ctx, dst) }); } - return ggml_backend_webgpu_build(ctx, ctx->param_buf_pool, ctx->soft_max_pipelines[mask_type][has_sink][inplace], - params, entries, ggml_nrows(dst)); + return ggml_backend_webgpu_build(ctx->global_ctx, ctx, ctx->param_buf_pool, + ctx->soft_max_pipelines[mask_type][has_sink][inplace], params, entries, + ggml_nrows(dst)); } static webgpu_command ggml_webgpu_argmax(webgpu_context & ctx, ggml_tensor * src, ggml_tensor * dst) { @@ -2111,7 +2032,7 @@ static webgpu_command ggml_webgpu_argmax(webgpu_context & ctx, ggml_tensor * src webgpu_pipeline pipeline = ctx->shader_lib->get_argmax_pipeline(shader_lib_ctx); uint32_t wg_x = ggml_nelements(dst); - return ggml_backend_webgpu_build(ctx, ctx->param_buf_pool, pipeline, params, entries, wg_x); + return ggml_backend_webgpu_build(ctx->global_ctx, ctx, ctx->param_buf_pool, pipeline, params, entries, wg_x); } static webgpu_command ggml_webgpu_argsort(webgpu_context & ctx, ggml_tensor * src, ggml_tensor * dst) { @@ -2205,8 +2126,8 @@ static webgpu_command ggml_webgpu_argsort(webgpu_context & ctx, ggml_tensor * sr workgroups_list.push_back({ wg_x_init, wg_y_init }); if (merge_passes == 0) { - return ggml_backend_webgpu_build_multi(ctx, ctx->param_buf_pool, pipelines, params_list, entries_list, - workgroups_list); + return ggml_backend_webgpu_build_multi(ctx->global_ctx, ctx, ctx->param_buf_pool, pipelines, params_list, + entries_list, workgroups_list); } bool in_is_tmp = start_in_tmp; @@ -2267,8 +2188,8 @@ static webgpu_command ggml_webgpu_argsort(webgpu_context & ctx, ggml_tensor * sr in_is_tmp = !in_is_tmp; } - return ggml_backend_webgpu_build_multi(ctx, ctx->param_buf_pool, pipelines, params_list, entries_list, - workgroups_list); + return ggml_backend_webgpu_build_multi(ctx->global_ctx, ctx, ctx->param_buf_pool, pipelines, params_list, + entries_list, workgroups_list); } static webgpu_command ggml_webgpu_cumsum(webgpu_context & ctx, ggml_tensor * src, ggml_tensor * dst) { @@ -2296,7 +2217,7 @@ static webgpu_command ggml_webgpu_cumsum(webgpu_context & ctx, ggml_tensor * src webgpu_pipeline pipeline = ctx->shader_lib->get_cumsum_pipeline(shader_lib_ctx); uint32_t wg_x = ggml_nrows(dst); - return ggml_backend_webgpu_build(ctx, ctx->param_buf_pool, pipeline, params, entries, wg_x); + return ggml_backend_webgpu_build(ctx->global_ctx, ctx, ctx->param_buf_pool, pipeline, params, entries, wg_x); } static webgpu_command ggml_webgpu_sum_rows(webgpu_context & ctx, ggml_tensor * src, ggml_tensor * dst) { @@ -2328,7 +2249,7 @@ static webgpu_command ggml_webgpu_sum_rows(webgpu_context & ctx, ggml_tensor * s webgpu_pipeline pipeline = ctx->shader_lib->get_sum_rows_pipeline(shader_lib_ctx); uint32_t wg_x = total_sum ? 1 : ggml_nrows(dst); - return ggml_backend_webgpu_build(ctx, ctx->param_buf_pool, pipeline, params, entries, wg_x); + return ggml_backend_webgpu_build(ctx->global_ctx, ctx, ctx->param_buf_pool, pipeline, params, entries, wg_x); } // Returns the encoded command, or std::nullopt if the operation is a no-op @@ -2438,7 +2359,7 @@ static ggml_status ggml_backend_webgpu_graph_compute(ggml_backend_t backend, str if (num_batched_kernels >= WEBGPU_COMMAND_SUBMIT_BATCH_SIZE) { num_batched_kernels = 0; - subs.push_back(ggml_backend_webgpu_submit(ctx, commands, ctx->param_buf_pool)); + subs.push_back(ggml_backend_webgpu_submit(ctx->global_ctx, ctx, commands, ctx->param_buf_pool)); // Process events and check for completed submissions ctx->global_ctx->instance.ProcessEvents(); ggml_backend_webgpu_wait(ctx->global_ctx, subs, false); @@ -2446,7 +2367,7 @@ static ggml_status ggml_backend_webgpu_graph_compute(ggml_backend_t backend, str } } if (!commands.empty()) { - subs.push_back(ggml_backend_webgpu_submit(ctx, commands, ctx->param_buf_pool)); + subs.push_back(ggml_backend_webgpu_submit(ctx->global_ctx, ctx, commands, ctx->param_buf_pool)); commands.clear(); } From 89dfb1169ba0f8bea778fc039f785c62e1f33460 Mon Sep 17 00:00:00 2001 From: Nikhil Jain Date: Wed, 18 Mar 2026 12:01:54 -0700 Subject: [PATCH 6/9] Add GGML_UNUSED --- ggml/src/ggml-webgpu/ggml-webgpu.cpp | 5 +++++ 1 file changed, 5 insertions(+) diff --git a/ggml/src/ggml-webgpu/ggml-webgpu.cpp b/ggml/src/ggml-webgpu/ggml-webgpu.cpp index bf0cbebeb51..ef44145cda9 100644 --- a/ggml/src/ggml-webgpu/ggml-webgpu.cpp +++ b/ggml/src/ggml-webgpu/ggml-webgpu.cpp @@ -683,6 +683,8 @@ static webgpu_submission ggml_backend_webgpu_submit(webgpu_global_context & ctx->profiling_futures.push_back({ f, false }); } } +#else + GGML_UNUSED(ctx_opt); #endif return submission; } @@ -757,6 +759,7 @@ static webgpu_command ggml_backend_webgpu_build_multi( pass = encoder.BeginComputePass(); } #else + GGML_UNUSED(ctx_opt); pass = encoder.BeginComputePass(); #endif for (size_t i = 0; i < pipelines.size(); i++) { @@ -777,6 +780,8 @@ static webgpu_command ggml_backend_webgpu_build_multi( // TODO: handle multiple pipeline names result.pipeline_name = pipelines.front().name; } +#else + GGML_UNUSED(ctx_opt); #endif return result; } From 667848a7efdcf56bf4cc25a312f591dc14d5e797 Mon Sep 17 00:00:00 2001 From: Nikhil Jain Date: Thu, 19 Mar 2026 12:59:19 -0700 Subject: [PATCH 7/9] Set up debugging in CI --- .github/workflows/build.yml | 1025 +---------------------------------- 1 file changed, 4 insertions(+), 1021 deletions(-) diff --git a/.github/workflows/build.yml b/.github/workflows/build.yml index 6d500d30986..68f53ff3518 100644 --- a/.github/workflows/build.yml +++ b/.github/workflows/build.yml @@ -96,42 +96,6 @@ jobs: cd build ctest -L main -E "test-llama-archs" --verbose --timeout 900 - macOS-latest-x64: - runs-on: macos-15-intel - - steps: - - name: Clone - id: checkout - uses: actions/checkout@v6 - - - name: ccache - uses: ggml-org/ccache-action@v1.2.21 - with: - key: macOS-latest-x64 - evict-old-files: 1d - save: ${{ github.event_name == 'push' && github.ref == 'refs/heads/master' }} - - - name: Build - id: cmake_build - run: | - sysctl -a - # Metal is disabled due to intermittent failures with Github runners not having a GPU: - # https://github.com/ggml-org/llama.cpp/actions/runs/8635935781/job/23674807267#step:5:2313 - cmake -B build \ - -DCMAKE_BUILD_RPATH="@loader_path" \ - -DLLAMA_FATAL_WARNINGS=ON \ - -DLLAMA_BUILD_BORINGSSL=ON \ - -DGGML_METAL=OFF \ - -DGGML_RPC=ON \ - -DCMAKE_OSX_DEPLOYMENT_TARGET=13.3 - cmake --build build --config Release -j $(sysctl -n hw.logicalcpu) - - - name: Test - id: cmake_test - run: | - cd build - ctest -L main --verbose --timeout 900 - macOS-latest-arm64-webgpu: runs-on: macos-latest @@ -174,148 +138,6 @@ jobs: cd build ctest -L main --verbose --timeout 900 - ubuntu-cpu: - strategy: - matrix: - include: - - build: 'x64' - os: ubuntu-22.04 - - build: 'arm64' - os: ubuntu-22.04-arm - - build: 's390x' - os: ubuntu-24.04-s390x - - build: 'ppc64le' - os: ubuntu-24.04-ppc64le - - runs-on: ${{ matrix.os }} - - steps: - - name: Clone - id: checkout - uses: actions/checkout@v6 - - - name: ccache - if: ${{ matrix.build != 's390x' && matrix.build != 'ppc64le' }} - uses: ggml-org/ccache-action@v1.2.21 - with: - key: ubuntu-cpu-${{ matrix.build }} - evict-old-files: 1d - save: ${{ github.event_name == 'push' && github.ref == 'refs/heads/master' }} - - - name: Build Dependencies - id: build_depends - run: | - sudo apt-get update - sudo apt-get install -y --no-install-recommends \ - python3 python3-pip python3-dev \ - libjpeg-dev build-essential libssl-dev \ - git-lfs - - - name: Python Dependencies - id: python_depends - run: | - python3 -m pip install --upgrade pip - pip3 install ./gguf-py - - - name: Swap Endianness - id: endianness - if: ${{ matrix.build == 's390x' }} - run: | - for f in models/*.gguf; do - echo YES | python3 gguf-py/gguf/scripts/gguf_convert_endian.py $f big - done - - - name: Build - id: cmake_build - run: | - cmake -B build \ - -DLLAMA_FATAL_WARNINGS=ON \ - -DGGML_RPC=ON - cmake --build build --config Release -j $(nproc) - - - name: Test - id: cmake_test - run: | - cd build - ctest -L main --verbose --timeout 900 - - - name: Test llama2c conversion - id: llama2c_test - if: ${{ matrix.build != 's390x' }} - run: | - cd build - echo "Fetch tokenizer" - wget https://huggingface.co/karpathy/tinyllamas/resolve/main/stories260K/tok512.bin - echo "Fetch llama2c model" - wget https://huggingface.co/karpathy/tinyllamas/resolve/main/stories260K/stories260K.bin - ./bin/llama-convert-llama2c-to-ggml --copy-vocab-from-model ./tok512.bin --llama2c-model stories260K.bin --llama2c-output-model stories260K.gguf - ./bin/llama-completion -m stories260K.gguf -p "One day, Lily met a Shoggoth" -n 500 -c 256 - - - name: Test llama2c (s390x) - id: llama2c_test_s390x - if: ${{ matrix.build == 's390x' }} - run: | - cd build - echo "Fetch llama2c big-endian model" - wget https://huggingface.co/ggml-org/models/resolve/main/tinyllamas/stories260K-be.gguf - ./bin/llama-completion -m stories260K-be.gguf -p "One day, Lily met a Shoggoth" -n 500 -c 256 - - ubuntu-latest-rpc: - runs-on: ubuntu-latest - - continue-on-error: true - - steps: - - name: Clone - id: checkout - uses: actions/checkout@v6 - - - name: Dependencies - id: depends - run: | - sudo apt-get update - sudo apt-get install build-essential libssl-dev - - - name: Build - id: cmake_build - run: | - cmake -B build \ - -DGGML_RPC=ON - cmake --build build --config Release -j $(nproc) - - - name: Test - id: cmake_test - run: | - cd build - ctest -L main --verbose - - ubuntu-24-vulkan: - runs-on: ${{ 'ubuntu-24.04-arm' || 'ubuntu-24.04' }} - - steps: - - name: Clone - id: checkout - uses: actions/checkout@v6 - - - name: Dependencies - id: depends - run: | - sudo apt-get install -y glslc libvulkan-dev libssl-dev - - - name: Configure - id: cmake_configure - run: | - cmake -B build \ - -DCMAKE_BUILD_TYPE=RelWithDebInfo \ - -DGGML_BACKEND_DL=ON \ - -DGGML_CPU_ALL_VARIANTS=ON \ - -DGGML_VULKAN=ON - - - name: Build - id: cmake_build - run: | - cmake --build build -j $(nproc) - ubuntu-24-webgpu: runs-on: ubuntu-24.04 @@ -387,6 +209,10 @@ jobs: # This is using llvmpipe and runs slower than other backends ctest -L main --verbose --timeout 900 + - name: Setup tmate session + if: ${{ failure() }} + uses: mxschmitt/action-tmate@v3 + ubuntu-24-webgpu-wasm: runs-on: ${{ 'ubuntu-24.04-arm' || 'ubuntu-24.04' }} @@ -420,846 +246,3 @@ jobs: -DEMDAWNWEBGPU_DIR=emdawnwebgpu_pkg cmake --build build-wasm --target test-backend-ops -j $(nproc) - - ubuntu-22-hip: - runs-on: ubuntu-22.04 - container: rocm/dev-ubuntu-22.04:6.1.2 - - steps: - - name: Clone - id: checkout - uses: actions/checkout@v6 - - - name: Dependencies - id: depends - run: | - sudo apt-get update - sudo apt-get install -y build-essential git cmake rocblas-dev hipblas-dev libssl-dev rocwmma-dev - - - name: ccache - uses: ggml-org/ccache-action@v1.2.21 - with: - key: ubuntu-22-hip - evict-old-files: 1d - save: ${{ github.event_name == 'push' && github.ref == 'refs/heads/master' }} - - - name: Build with native CMake HIP support - id: cmake_build - run: | - cmake -B build -S . \ - -DCMAKE_HIP_COMPILER="$(hipconfig -l)/clang" \ - -DGGML_HIP_ROCWMMA_FATTN=ON \ - -DGGML_HIP=ON - cmake --build build --config Release -j $(nproc) - - ubuntu-22-musa: - runs-on: ubuntu-22.04 - container: mthreads/musa:rc4.3.0-devel-ubuntu22.04-amd64 - - steps: - - name: Clone - id: checkout - uses: actions/checkout@v6 - - - name: Dependencies - id: depends - run: | - apt-get update - apt-get install -y build-essential git cmake libssl-dev - - - name: ccache - uses: ggml-org/ccache-action@v1.2.21 - with: - key: ubuntu-22-musa - evict-old-files: 1d - save: ${{ github.event_name == 'push' && github.ref == 'refs/heads/master' }} - - - name: Build with native CMake MUSA support - id: cmake_build - run: | - cmake -B build -S . \ - -DGGML_MUSA=ON - cmake --build build --config Release -j $(nproc) - - ubuntu-22-sycl: - runs-on: ubuntu-22.04 - - continue-on-error: true - - steps: - - uses: actions/checkout@v6 - - - name: add oneAPI to apt - shell: bash - run: | - cd /tmp - wget https://apt.repos.intel.com/intel-gpg-keys/GPG-PUB-KEY-INTEL-SW-PRODUCTS.PUB - sudo apt-key add GPG-PUB-KEY-INTEL-SW-PRODUCTS.PUB - rm GPG-PUB-KEY-INTEL-SW-PRODUCTS.PUB - sudo add-apt-repository "deb https://apt.repos.intel.com/oneapi all main" - - - name: install oneAPI dpcpp compiler - shell: bash - run: | - sudo apt update - sudo apt install intel-oneapi-compiler-dpcpp-cpp libssl-dev - - - name: install oneAPI MKL library - shell: bash - run: | - sudo apt install intel-oneapi-mkl-devel - - - name: Clone - id: checkout - uses: actions/checkout@v6 - - - name: ccache - uses: ggml-org/ccache-action@v1.2.21 - with: - key: ubuntu-22-sycl - evict-old-files: 1d - save: ${{ github.event_name == 'push' && github.ref == 'refs/heads/master' }} - - - name: Build - id: cmake_build - run: | - source /opt/intel/oneapi/setvars.sh - cmake -B build \ - -DGGML_SYCL=ON \ - -DCMAKE_C_COMPILER=icx \ - -DCMAKE_CXX_COMPILER=icpx - cmake --build build --config Release -j $(nproc) - - ubuntu-22-sycl-fp16: - runs-on: ubuntu-22.04 - - continue-on-error: true - - steps: - - uses: actions/checkout@v6 - - - name: add oneAPI to apt - shell: bash - run: | - cd /tmp - wget https://apt.repos.intel.com/intel-gpg-keys/GPG-PUB-KEY-INTEL-SW-PRODUCTS.PUB - sudo apt-key add GPG-PUB-KEY-INTEL-SW-PRODUCTS.PUB - rm GPG-PUB-KEY-INTEL-SW-PRODUCTS.PUB - sudo add-apt-repository "deb https://apt.repos.intel.com/oneapi all main" - - - name: install oneAPI dpcpp compiler - shell: bash - run: | - sudo apt update - sudo apt install intel-oneapi-compiler-dpcpp-cpp libssl-dev - - - name: install oneAPI MKL library - shell: bash - run: | - sudo apt install intel-oneapi-mkl-devel - - - name: Clone - id: checkout - uses: actions/checkout@v6 - - - name: ccache - uses: ggml-org/ccache-action@v1.2.21 - with: - key: ubuntu-22-sycl-fp16 - evict-old-files: 1d - save: ${{ github.event_name == 'push' && github.ref == 'refs/heads/master' }} - - - name: Build - id: cmake_build - run: | - source /opt/intel/oneapi/setvars.sh - cmake -B build \ - -DGGML_SYCL=ON \ - -DCMAKE_C_COMPILER=icx \ - -DCMAKE_CXX_COMPILER=icpx \ - -DGGML_SYCL_F16=ON - cmake --build build --config Release -j $(nproc) - - ubuntu-24-openvino: - name: ubuntu-24-openvino-${{ matrix.openvino_device }} - strategy: - matrix: - include: - - variant: cpu - runner: '"ubuntu-24.04"' - openvino_device: "CPU" - - variant: gpu - runner: '["self-hosted","Linux","X64","Intel"]' - openvino_device: "GPU" - - runs-on: ${{ fromJSON(matrix.runner) }} - - env: - # Sync versions in build.yml, build-self-hosted.yml, release.yml, build-cache.yml, .devops/openvino.Dockerfile - OPENVINO_VERSION_MAJOR: "2026.0" - OPENVINO_VERSION_FULL: "2026.0.0.20965.c6d6a13a886" - - steps: - - name: Clone - id: checkout - uses: actions/checkout@v6 - - - name: ccache - if: runner.environment == 'github-hosted' - uses: ggml-org/ccache-action@v1.2.21 - with: - key: ubuntu-24-openvino-${{ matrix.variant }}-no-preset-v1 - evict-old-files: 1d - save: ${{ github.event_name == 'push' && github.ref == 'refs/heads/master' }} - - - name: Dependencies - id: depends - run: | - sudo apt-get update - sudo apt-get install -y build-essential libssl-dev libtbb12 cmake ninja-build python3-pip - sudo apt-get install -y ocl-icd-opencl-dev opencl-headers opencl-clhpp-headers intel-opencl-icd - - - name: Use OpenVINO Toolkit Cache - if: runner.environment == 'github-hosted' - uses: actions/cache@v5 - id: cache-openvino - with: - path: ./openvino_toolkit - key: openvino-toolkit-v${{ env.OPENVINO_VERSION_FULL }}-${{ runner.os }} - - - name: Setup OpenVINO Toolkit - if: steps.cache-openvino.outputs.cache-hit != 'true' - uses: ./.github/actions/linux-setup-openvino - with: - path: ./openvino_toolkit - version_major: ${{ env.OPENVINO_VERSION_MAJOR }} - version_full: ${{ env.OPENVINO_VERSION_FULL }} - - - name: Install OpenVINO dependencies - run: | - cd ./openvino_toolkit - chmod +x ./install_dependencies/install_openvino_dependencies.sh - echo "Y" | sudo -E ./install_dependencies/install_openvino_dependencies.sh - - - name: Build - id: cmake_build - run: | - source ./openvino_toolkit/setupvars.sh - cmake -B build/ReleaseOV -G Ninja \ - -DCMAKE_BUILD_TYPE=Release \ - -DGGML_OPENVINO=ON - cmake --build build/ReleaseOV --config Release -j $(nproc) - - - name: Test - id: cmake_test - # TODO: fix and re-enable the `test-llama-archs` test below - run: | - cd ${{ github.workspace }} - if [ "${{ matrix.openvino_device }}" = "GPU" ]; then - export GGML_OPENVINO_DEVICE=GPU - fi - ctest --test-dir build/ReleaseOV -L main -E "test-llama-archs" --verbose --timeout 2000 - - windows-latest: - runs-on: windows-2025 - - env: - OPENBLAS_VERSION: 0.3.23 - SDE_VERSION: 9.33.0-2024-01-07 - VULKAN_VERSION: 1.4.313.2 - - strategy: - matrix: - include: - - build: 'cpu-x64 (static)' - arch: 'x64' - defines: '-G "Ninja Multi-Config" -D CMAKE_TOOLCHAIN_FILE=cmake/x64-windows-llvm.cmake -DGGML_NATIVE=OFF -DLLAMA_BUILD_SERVER=ON -DGGML_RPC=ON -DBUILD_SHARED_LIBS=OFF' - - build: 'openblas-x64' - arch: 'x64' - defines: '-G "Ninja Multi-Config" -D CMAKE_TOOLCHAIN_FILE=cmake/x64-windows-llvm.cmake -DGGML_NATIVE=OFF -DLLAMA_BUILD_SERVER=ON -DGGML_RPC=ON -DGGML_BACKEND_DL=ON -DGGML_CPU_ALL_VARIANTS=ON -DGGML_OPENMP=OFF -DGGML_BLAS=ON -DGGML_BLAS_VENDOR=OpenBLAS -DBLAS_INCLUDE_DIRS="$env:RUNNER_TEMP/openblas/include" -DBLAS_LIBRARIES="$env:RUNNER_TEMP/openblas/lib/openblas.lib"' - - build: 'vulkan-x64' - arch: 'x64' - defines: '-DCMAKE_BUILD_TYPE=Release -DGGML_NATIVE=OFF -DLLAMA_BUILD_SERVER=ON -DGGML_RPC=ON -DGGML_BACKEND_DL=ON -DGGML_CPU_ALL_VARIANTS=ON -DGGML_VULKAN=ON' - - build: 'llvm-arm64' - arch: 'arm64' - defines: '-G "Ninja Multi-Config" -D CMAKE_TOOLCHAIN_FILE=cmake/arm64-windows-llvm.cmake -DGGML_NATIVE=OFF -DLLAMA_BUILD_SERVER=ON' - - build: 'llvm-arm64-opencl-adreno' - arch: 'arm64' - defines: '-G "Ninja Multi-Config" -D CMAKE_TOOLCHAIN_FILE=cmake/arm64-windows-llvm.cmake -DCMAKE_PREFIX_PATH="$env:RUNNER_TEMP/opencl-arm64-release" -DGGML_OPENCL=ON -DGGML_OPENCL_USE_ADRENO_KERNELS=ON' - - steps: - - name: Clone - id: checkout - uses: actions/checkout@v6 - - - name: ccache - uses: ggml-org/ccache-action@v1.2.21 - with: - key: windows-latest-${{ matrix.build }} - variant: ccache - evict-old-files: 1d - save: ${{ github.event_name == 'push' && github.ref == 'refs/heads/master' }} - - - name: Download OpenBLAS - id: get_openblas - if: ${{ matrix.build == 'openblas-x64' }} - run: | - curl.exe -o $env:RUNNER_TEMP/openblas.zip -L "https://github.com/xianyi/OpenBLAS/releases/download/v${env:OPENBLAS_VERSION}/OpenBLAS-${env:OPENBLAS_VERSION}-x64.zip" - curl.exe -o $env:RUNNER_TEMP/OpenBLAS.LICENSE.txt -L "https://github.com/xianyi/OpenBLAS/raw/v${env:OPENBLAS_VERSION}/LICENSE" - mkdir $env:RUNNER_TEMP/openblas - tar.exe -xvf $env:RUNNER_TEMP/openblas.zip -C $env:RUNNER_TEMP/openblas - $vcdir = $(vswhere -latest -products * -requires Microsoft.VisualStudio.Component.VC.Tools.x86.x64 -property installationPath) - $msvc = $(join-path $vcdir $('VC\Tools\MSVC\'+$(gc -raw $(join-path $vcdir 'VC\Auxiliary\Build\Microsoft.VCToolsVersion.default.txt')).Trim())) - $lib = $(join-path $msvc 'bin\Hostx64\x64\lib.exe') - & $lib /machine:x64 "/def:${env:RUNNER_TEMP}/openblas/lib/libopenblas.def" "/out:${env:RUNNER_TEMP}/openblas/lib/openblas.lib" /name:openblas.dll - - - name: Install Vulkan SDK - id: get_vulkan - if: ${{ matrix.build == 'vulkan-x64' }} - run: | - curl.exe -o $env:RUNNER_TEMP/VulkanSDK-Installer.exe -L "https://sdk.lunarg.com/sdk/download/${env:VULKAN_VERSION}/windows/vulkansdk-windows-X64-${env:VULKAN_VERSION}.exe" - & "$env:RUNNER_TEMP\VulkanSDK-Installer.exe" --accept-licenses --default-answer --confirm-command install - Add-Content $env:GITHUB_ENV "VULKAN_SDK=C:\VulkanSDK\${env:VULKAN_VERSION}" - Add-Content $env:GITHUB_PATH "C:\VulkanSDK\${env:VULKAN_VERSION}\bin" - - - name: Install Ninja - id: install_ninja - run: | - choco install ninja - - - name: Install OpenCL Headers and Libs - id: install_opencl - if: ${{ matrix.build == 'llvm-arm64-opencl-adreno' }} - run: | - git clone https://github.com/KhronosGroup/OpenCL-Headers - cd OpenCL-Headers - cmake -B build ` - -DBUILD_TESTING=OFF ` - -DOPENCL_HEADERS_BUILD_TESTING=OFF ` - -DOPENCL_HEADERS_BUILD_CXX_TESTS=OFF ` - -DCMAKE_INSTALL_PREFIX="$env:RUNNER_TEMP/opencl-arm64-release" - cmake --build build --target install - git clone https://github.com/KhronosGroup/OpenCL-ICD-Loader - cd OpenCL-ICD-Loader - cmake -B build-arm64-release ` - -A arm64 ` - -DCMAKE_PREFIX_PATH="$env:RUNNER_TEMP/opencl-arm64-release" ` - -DCMAKE_INSTALL_PREFIX="$env:RUNNER_TEMP/opencl-arm64-release" - cmake --build build-arm64-release --target install --config release - - - name: Build - id: cmake_build - run: | - cmake -S . -B build ${{ matrix.defines }} ` - -DLLAMA_BUILD_BORINGSSL=ON - cmake --build build --config Release -j ${env:NUMBER_OF_PROCESSORS} - - - name: Add libopenblas.dll - id: add_libopenblas_dll - if: ${{ matrix.build == 'openblas-x64' }} - run: | - cp $env:RUNNER_TEMP/openblas/bin/libopenblas.dll ./build/bin/Release/openblas.dll - cp $env:RUNNER_TEMP/OpenBLAS.LICENSE.txt ./build/bin/Release/OpenBLAS-${env:OPENBLAS_VERSION}.txt - - - name: Test - id: cmake_test - if: ${{ matrix.arch == 'x64' }} - run: | - cd build - ctest -L main -C Release --verbose --timeout 900 - - # TODO: disabled for now, consider adding tests for all CPU variants instead - # - name: Test (Intel SDE) - # id: cmake_test_sde - # if: ${{ matrix.build == 'avx512-x64' && env.HAS_AVX512F == '0' }} # use Intel SDE for AVX-512 emulation - # run: | - # curl.exe -o $env:RUNNER_TEMP/sde.tar.xz -L "https://downloadmirror.intel.com/813591/sde-external-${env:SDE_VERSION}-win.tar.xz" - # # for some weird reason windows tar doesn't like sde tar.xz - # 7z x "-o${env:RUNNER_TEMP}" $env:RUNNER_TEMP/sde.tar.xz - # 7z x "-o${env:RUNNER_TEMP}" $env:RUNNER_TEMP/sde.tar - # $sde = $(join-path $env:RUNNER_TEMP sde-external-${env:SDE_VERSION}-win/sde.exe) - # cd build - # $env:LLAMA_SKIP_TESTS_SLOW_ON_EMULATOR = 1 - # & $sde -future -- ctest -L main -C Release --verbose --timeout 900 - - ubuntu-latest-cuda: - runs-on: ubuntu-latest - container: nvidia/cuda:12.6.2-devel-ubuntu24.04 - - steps: - - name: Clone - id: checkout - uses: actions/checkout@v6 - - - name: Install dependencies - env: - DEBIAN_FRONTEND: noninteractive - run: | - apt update - apt install -y cmake build-essential ninja-build libgomp1 git libssl-dev - - - name: ccache - uses: ggml-org/ccache-action@v1.2.21 - with: - key: ubuntu-latest-cuda - evict-old-files: 1d - save: ${{ github.event_name == 'push' && github.ref == 'refs/heads/master' }} - - - name: Build with CMake - # TODO: Remove GGML_CUDA_CUB_3DOT2 flag once CCCL 3.2 is bundled within CTK and that CTK version is used in this project - run: | - cmake -S . -B build -G Ninja \ - -DLLAMA_FATAL_WARNINGS=ON \ - -DCMAKE_BUILD_TYPE=Release \ - -DCMAKE_CUDA_ARCHITECTURES=89-real \ - -DCMAKE_EXE_LINKER_FLAGS=-Wl,--allow-shlib-undefined \ - -DGGML_NATIVE=OFF \ - -DGGML_CUDA=ON \ - -DGGML_CUDA_CUB_3DOT2=ON - cmake --build build - - windows-2022-cuda: - runs-on: windows-2022 - - strategy: - matrix: - cuda: ['12.4'] - - steps: - - name: Clone - id: checkout - uses: actions/checkout@v6 - - - name: Install ccache - uses: ggml-org/ccache-action@v1.2.21 - with: - key: windows-cuda-${{ matrix.cuda }} - variant: ccache - evict-old-files: 1d - save: ${{ github.event_name == 'push' && github.ref == 'refs/heads/master' }} - - - name: Install Cuda Toolkit - uses: ./.github/actions/windows-setup-cuda - with: - cuda_version: ${{ matrix.cuda }} - - - name: Install Ninja - id: install_ninja - run: | - choco install ninja - - - name: Build - id: cmake_build - shell: cmd - # TODO: Remove GGML_CUDA_CUB_3DOT2 flag once CCCL 3.2 is bundled within CTK and that CTK version is used in this project - run: | - call "C:\Program Files\Microsoft Visual Studio\2022\Enterprise\VC\Auxiliary\Build\vcvarsall.bat" x64 - cmake -S . -B build -G "Ninja Multi-Config" ^ - -DLLAMA_BUILD_SERVER=ON ^ - -DLLAMA_BUILD_BORINGSSL=ON ^ - -DGGML_NATIVE=OFF ^ - -DGGML_BACKEND_DL=ON ^ - -DGGML_CPU_ALL_VARIANTS=ON ^ - -DGGML_CUDA=ON ^ - -DGGML_RPC=ON ^ - -DGGML_CUDA_CUB_3DOT2=ON - set /A NINJA_JOBS=%NUMBER_OF_PROCESSORS%-1 - cmake --build build --config Release -j %NINJA_JOBS% -t ggml - cmake --build build --config Release - - windows-latest-sycl: - runs-on: windows-2022 - - defaults: - run: - shell: bash - - env: - WINDOWS_BASEKIT_URL: https://registrationcenter-download.intel.com/akdlm/IRC_NAS/24751ead-ddc5-4479-b9e6-f9fe2ff8b9f2/intel-deep-learning-essentials-2025.2.1.25_offline.exe - WINDOWS_DPCPP_MKL: intel.oneapi.win.cpp-dpcpp-common:intel.oneapi.win.mkl.devel:intel.oneapi.win.dnnl:intel.oneapi.win.tbb.devel - ONEAPI_ROOT: "C:/Program Files (x86)/Intel/oneAPI" - steps: - - name: Clone - id: checkout - uses: actions/checkout@v6 - - - name: ccache - uses: ggml-org/ccache-action@v1.2.21 - with: - key: windows-latest-sycl - variant: ccache - evict-old-files: 1d - save: ${{ github.event_name == 'push' && github.ref == 'refs/heads/master' }} - - - name: Install - run: | - scripts/install-oneapi.bat $WINDOWS_BASEKIT_URL $WINDOWS_DPCPP_MKL - - # TODO: add ssl support ; we will also need to modify win-build-sycl.bat to accept user-specified args - - - name: Build - id: cmake_build - run: examples/sycl/win-build-sycl.bat - - windows-latest-hip: - runs-on: windows-2022 - - env: - # Make sure this is in sync with build-cache.yml - HIPSDK_INSTALLER_VERSION: "26.Q1" - - steps: - - name: Clone - id: checkout - uses: actions/checkout@v6 - - - name: Grab rocWMMA package - id: grab_rocwmma - run: | - curl -o rocwmma.deb "https://repo.radeon.com/rocm/apt/7.2/pool/main/r/rocwmma-dev/rocwmma-dev_2.2.0.70200-43~24.04_amd64.deb" - 7z x rocwmma.deb - 7z x data.tar - - - name: Use ROCm Installation Cache - uses: actions/cache@v5 - id: cache-rocm - with: - path: C:\Program Files\AMD\ROCm - key: rocm-${{ env.HIPSDK_INSTALLER_VERSION }}-${{ runner.os }} - - - name: Setup ROCm - if: steps.cache-rocm.outputs.cache-hit != 'true' - uses: ./.github/actions/windows-setup-rocm - with: - version: ${{ env.HIPSDK_INSTALLER_VERSION }} - - - name: Verify ROCm - id: verify - run: | - # Find and test ROCm installation - $clangPath = Get-ChildItem 'C:\Program Files\AMD\ROCm\*\bin\clang.exe' | Select-Object -First 1 - if (-not $clangPath) { - Write-Error "ROCm installation not found" - exit 1 - } - & $clangPath.FullName --version - - - name: Install ccache - uses: ggml-org/ccache-action@v1.2.21 - with: - key: ${{ github.job }} - evict-old-files: 1d - save: ${{ github.event_name == 'push' && github.ref == 'refs/heads/master' }} - - - name: Build - id: cmake_build - run: | - $env:HIP_PATH=$(Resolve-Path 'C:\Program Files\AMD\ROCm\*\bin\clang.exe' | split-path | split-path) - $env:CMAKE_PREFIX_PATH="${env:HIP_PATH}" - cmake -G "Unix Makefiles" -B build -S . ` - -DCMAKE_C_COMPILER="${env:HIP_PATH}\bin\clang.exe" ` - -DCMAKE_CXX_COMPILER="${env:HIP_PATH}\bin\clang++.exe" ` - -DCMAKE_CXX_FLAGS="-I$($PWD.Path.Replace('\', '/'))/opt/rocm-7.2.0/include/" ` - -DCMAKE_BUILD_TYPE=Release ` - -DLLAMA_BUILD_BORINGSSL=ON ` - -DROCM_DIR="${env:HIP_PATH}" ` - -DGGML_HIP=ON ` - -DGGML_HIP_ROCWMMA_FATTN=ON ` - -DGGML_RPC=ON - cmake --build build -j ${env:NUMBER_OF_PROCESSORS} - - ubuntu-cpu-riscv64-native: - runs-on: RISCV64 - - steps: - - name: Install dependencies - run: | - sudo apt-get update - - # Install necessary packages - sudo apt-get install -y libatomic1 libtsan2 gcc-14 g++-14 rustup cmake build-essential libssl-dev wget ccache git-lfs - - # Set gcc-14 and g++-14 as the default compilers - sudo update-alternatives --install /usr/bin/gcc gcc /usr/bin/gcc-14 100 - sudo update-alternatives --install /usr/bin/g++ g++ /usr/bin/g++-14 100 - sudo ln -sf /usr/bin/gcc-14 /usr/bin/gcc - sudo ln -sf /usr/bin/g++-14 /usr/bin/g++ - - # Install Rust stable version - rustup install stable - rustup default stable - - git lfs install - - - name: Clone - id: checkout - uses: actions/checkout@v6 - - - name: Check environment - run: | - uname -a - gcc --version - g++ --version - ldd --version - cmake --version - rustc --version - - - name: Setup ccache - run: | - # Set unique cache directory for this job - export CCACHE_DIR="$HOME/.ccache/cpu-cmake-rv64-native" - mkdir -p "$CCACHE_DIR" - - # Configure ccache for optimal performance - ccache --set-config=max_size=5G - ccache --set-config=compression=true - ccache --set-config=compression_level=6 - ccache --set-config=cache_dir="$CCACHE_DIR" - - # Enable more aggressive caching - ccache --set-config=sloppiness=file_macro,time_macros,include_file_mtime,include_file_ctime - ccache --set-config=hash_dir=false - - # Export for subsequent steps - echo "CCACHE_DIR=$CCACHE_DIR" >> $GITHUB_ENV - echo "PATH=/usr/lib/ccache:$PATH" >> $GITHUB_ENV - - - name: Build - id: cmake_build - run: | - cmake -B build \ - -DCMAKE_BUILD_TYPE=Release \ - -DGGML_OPENMP=OFF \ - -DLLAMA_BUILD_EXAMPLES=ON \ - -DLLAMA_BUILD_TOOLS=ON \ - -DLLAMA_BUILD_TESTS=ON \ - -DCMAKE_C_COMPILER_LAUNCHER=ccache \ - -DCMAKE_CXX_COMPILER_LAUNCHER=ccache \ - -DGGML_RPC=ON \ - -DCMAKE_C_COMPILER=riscv64-linux-gnu-gcc-14 \ - -DCMAKE_CXX_COMPILER=riscv64-linux-gnu-g++-14 - - cmake --build build --config Release -j $(nproc) - - - name: Test - id: cmake_test - run: | - cd build - ctest -L main --verbose --timeout 900 - - - name: Test llama2c conversion - id: llama2c_test - run: | - cd build - echo "Fetch tokenizer" - wget https://huggingface.co/karpathy/tinyllamas/resolve/main/stories260K/tok512.bin - echo "Fetch llama2c model" - wget https://huggingface.co/karpathy/tinyllamas/resolve/main/stories260K/stories260K.bin - ./bin/llama-convert-llama2c-to-ggml --copy-vocab-from-model ./tok512.bin --llama2c-model stories260K.bin --llama2c-output-model stories260K.gguf - ./bin/llama-completion -m stories260K.gguf -p "One day, Lily met a Shoggoth" -n 500 -c 256 - -# TODO: simplify the following workflows using a matrix -# TODO: run lighter CI on PRs and the full CI only on master (if needed) - ggml-ci-x64-cpu-low-perf: - runs-on: ubuntu-22.04 - - steps: - - name: Clone - id: checkout - uses: actions/checkout@v6 - - - name: ccache - uses: ggml-org/ccache-action@v1.2.21 - with: - key: ggml-ci-x64-cpu-low-perf - evict-old-files: 1d - save: ${{ github.event_name == 'push' && github.ref == 'refs/heads/master' }} - - - name: Dependencies - id: depends - run: | - sudo apt-get update - sudo apt-get install build-essential - - - name: Test - id: ggml-ci - run: | - LLAMA_ARG_THREADS=$(nproc) GG_BUILD_LOW_PERF=1 bash ./ci/run.sh ./tmp/results ./tmp/mnt - - ggml-ci-arm64-cpu-low-perf: - runs-on: ubuntu-22.04-arm - - steps: - - name: Clone - id: checkout - uses: actions/checkout@v6 - - - name: ccache - uses: ggml-org/ccache-action@v1.2.21 - with: - key: ggml-ci-arm64-cpu-low-perf - evict-old-files: 1d - save: ${{ github.event_name == 'push' && github.ref == 'refs/heads/master' }} - - - name: Dependencies - id: depends - run: | - sudo apt-get update - sudo apt-get install build-essential - - - name: Test - id: ggml-ci - run: | - LLAMA_ARG_THREADS=$(nproc) GG_BUILD_LOW_PERF=1 bash ./ci/run.sh ./tmp/results ./tmp/mnt - - ggml-ci-x64-cpu-high-perf: - runs-on: ubuntu-22.04 - - steps: - - name: Clone - id: checkout - uses: actions/checkout@v6 - - - name: ccache - uses: ggml-org/ccache-action@v1.2.21 - with: - key: ggml-ci-x64-cpu-high-perf - evict-old-files: 1d - save: ${{ github.event_name == 'push' && github.ref == 'refs/heads/master' }} - - - name: Dependencies - id: depends - run: | - sudo apt-get update - sudo apt-get install build-essential - - - name: Test - id: ggml-ci - run: | - LLAMA_ARG_THREADS=$(nproc) GG_BUILD_HIGH_PERF=1 bash ./ci/run.sh ./tmp/results ./tmp/mnt - - ggml-ci-arm64-cpu-high-perf: - runs-on: ubuntu-22.04-arm - - steps: - - name: Clone - id: checkout - uses: actions/checkout@v6 - - - name: ccache - uses: ggml-org/ccache-action@v1.2.21 - with: - key: ggml-ci-arm64-cpu-high-perf - evict-old-files: 1d - save: ${{ github.event_name == 'push' && github.ref == 'refs/heads/master' }} - - - name: Dependencies - id: depends - run: | - sudo apt-get update - sudo apt-get install build-essential - - - name: Test - id: ggml-ci - run: | - LLAMA_ARG_THREADS=$(nproc) GG_BUILD_HIGH_PERF=1 GG_BUILD_NO_SVE=1 GG_BUILD_NO_BF16=1 GG_BUILD_EXTRA_TESTS_0=1 bash ./ci/run.sh ./tmp/results ./tmp/mnt - - ggml-ci-arm64-cpu-high-perf-sve: - runs-on: ubuntu-22.04-arm - - steps: - - name: Clone - id: checkout - uses: actions/checkout@v6 - - - name: ccache - uses: ggml-org/ccache-action@v1.2.21 - with: - key: ggml-ci-arm64-cpu-high-perf-sve - evict-old-files: 1d - save: ${{ github.event_name == 'push' && github.ref == 'refs/heads/master' }} - - - name: Dependencies - id: depends - run: | - sudo apt-get update - sudo apt-get install build-essential - - - name: Test - id: ggml-ci - run: | - LLAMA_ARG_THREADS=$(nproc) GG_BUILD_NO_BF16=1 GG_BUILD_EXTRA_TESTS_0=1 bash ./ci/run.sh ./tmp/results ./tmp/mnt - - ggml-ci-arm64-cpu-kleidiai: - runs-on: ubuntu-22.04-arm - - steps: - - name: Clone - id: checkout - uses: actions/checkout@v6 - - - name: ccache - uses: ggml-org/ccache-action@v1.2.21 - with: - key: ggml-ci-arm64-cpu-kleidiai - evict-old-files: 1d - save: ${{ github.event_name == 'push' && github.ref == 'refs/heads/master' }} - - - name: Dependencies - id: depends - run: | - sudo apt-get update - sudo apt-get install -y build-essential - - - name: Test - id: ggml-ci - run: | - GG_BUILD_KLEIDIAI=1 GG_BUILD_EXTRA_TESTS_0=1 bash ./ci/run.sh ./tmp/results ./tmp/mnt - - ggml-ci-arm64-cpu-kleidiai-graviton4: - runs-on: ah-ubuntu_22_04-c8g_8x - - steps: - - name: Clone - id: checkout - uses: actions/checkout@v6 - - - name: Dependencies - id: depends - run: | - set -euxo pipefail - sudo apt-get update - sudo DEBIAN_FRONTEND=noninteractive NEEDRESTART_MODE=a \ - apt-get install -y \ - build-essential \ - python3-venv \ - gpg \ - wget \ - time \ - git-lfs - - git lfs install - - # install the latest cmake - sudo install -d /usr/share/keyrings - wget -O - https://apt.kitware.com/keys/kitware-archive-latest.asc \ - | gpg --dearmor \ - | sudo tee /usr/share/keyrings/kitware-archive-keyring.gpg >/dev/null - echo 'deb [signed-by=/usr/share/keyrings/kitware-archive-keyring.gpg] https://apt.kitware.com/ubuntu/ jammy main' \ - | sudo tee /etc/apt/sources.list.d/kitware.list - sudo apt-get update - sudo apt-get install -y cmake - - - name: ccache - uses: ggml-org/ccache-action@v1.2.21 - with: - key: ggml-ci-arm64-cpu-kleidiai-graviton4 - evict-old-files: 1d - save: ${{ github.event_name == 'push' && github.ref == 'refs/heads/master' }} - - - name: Test - id: ggml-ci - run: | - GG_BUILD_KLEIDIAI=1 \ - GG_BUILD_EXTRA_TESTS_0=1 \ - bash ./ci/run.sh ./tmp/results ./tmp/mnt From 4aaaea113724bb338508d69d0fe050d786156b22 Mon Sep 17 00:00:00 2001 From: Nikhil Jain Date: Thu, 19 Mar 2026 13:32:43 -0700 Subject: [PATCH 8/9] Remove CI debugging changes --- .github/workflows/build.yml | 1025 ++++++++++++++++++++++++++++++++++- 1 file changed, 1021 insertions(+), 4 deletions(-) diff --git a/.github/workflows/build.yml b/.github/workflows/build.yml index 68f53ff3518..6d500d30986 100644 --- a/.github/workflows/build.yml +++ b/.github/workflows/build.yml @@ -96,6 +96,42 @@ jobs: cd build ctest -L main -E "test-llama-archs" --verbose --timeout 900 + macOS-latest-x64: + runs-on: macos-15-intel + + steps: + - name: Clone + id: checkout + uses: actions/checkout@v6 + + - name: ccache + uses: ggml-org/ccache-action@v1.2.21 + with: + key: macOS-latest-x64 + evict-old-files: 1d + save: ${{ github.event_name == 'push' && github.ref == 'refs/heads/master' }} + + - name: Build + id: cmake_build + run: | + sysctl -a + # Metal is disabled due to intermittent failures with Github runners not having a GPU: + # https://github.com/ggml-org/llama.cpp/actions/runs/8635935781/job/23674807267#step:5:2313 + cmake -B build \ + -DCMAKE_BUILD_RPATH="@loader_path" \ + -DLLAMA_FATAL_WARNINGS=ON \ + -DLLAMA_BUILD_BORINGSSL=ON \ + -DGGML_METAL=OFF \ + -DGGML_RPC=ON \ + -DCMAKE_OSX_DEPLOYMENT_TARGET=13.3 + cmake --build build --config Release -j $(sysctl -n hw.logicalcpu) + + - name: Test + id: cmake_test + run: | + cd build + ctest -L main --verbose --timeout 900 + macOS-latest-arm64-webgpu: runs-on: macos-latest @@ -138,6 +174,148 @@ jobs: cd build ctest -L main --verbose --timeout 900 + ubuntu-cpu: + strategy: + matrix: + include: + - build: 'x64' + os: ubuntu-22.04 + - build: 'arm64' + os: ubuntu-22.04-arm + - build: 's390x' + os: ubuntu-24.04-s390x + - build: 'ppc64le' + os: ubuntu-24.04-ppc64le + + runs-on: ${{ matrix.os }} + + steps: + - name: Clone + id: checkout + uses: actions/checkout@v6 + + - name: ccache + if: ${{ matrix.build != 's390x' && matrix.build != 'ppc64le' }} + uses: ggml-org/ccache-action@v1.2.21 + with: + key: ubuntu-cpu-${{ matrix.build }} + evict-old-files: 1d + save: ${{ github.event_name == 'push' && github.ref == 'refs/heads/master' }} + + - name: Build Dependencies + id: build_depends + run: | + sudo apt-get update + sudo apt-get install -y --no-install-recommends \ + python3 python3-pip python3-dev \ + libjpeg-dev build-essential libssl-dev \ + git-lfs + + - name: Python Dependencies + id: python_depends + run: | + python3 -m pip install --upgrade pip + pip3 install ./gguf-py + + - name: Swap Endianness + id: endianness + if: ${{ matrix.build == 's390x' }} + run: | + for f in models/*.gguf; do + echo YES | python3 gguf-py/gguf/scripts/gguf_convert_endian.py $f big + done + + - name: Build + id: cmake_build + run: | + cmake -B build \ + -DLLAMA_FATAL_WARNINGS=ON \ + -DGGML_RPC=ON + cmake --build build --config Release -j $(nproc) + + - name: Test + id: cmake_test + run: | + cd build + ctest -L main --verbose --timeout 900 + + - name: Test llama2c conversion + id: llama2c_test + if: ${{ matrix.build != 's390x' }} + run: | + cd build + echo "Fetch tokenizer" + wget https://huggingface.co/karpathy/tinyllamas/resolve/main/stories260K/tok512.bin + echo "Fetch llama2c model" + wget https://huggingface.co/karpathy/tinyllamas/resolve/main/stories260K/stories260K.bin + ./bin/llama-convert-llama2c-to-ggml --copy-vocab-from-model ./tok512.bin --llama2c-model stories260K.bin --llama2c-output-model stories260K.gguf + ./bin/llama-completion -m stories260K.gguf -p "One day, Lily met a Shoggoth" -n 500 -c 256 + + - name: Test llama2c (s390x) + id: llama2c_test_s390x + if: ${{ matrix.build == 's390x' }} + run: | + cd build + echo "Fetch llama2c big-endian model" + wget https://huggingface.co/ggml-org/models/resolve/main/tinyllamas/stories260K-be.gguf + ./bin/llama-completion -m stories260K-be.gguf -p "One day, Lily met a Shoggoth" -n 500 -c 256 + + ubuntu-latest-rpc: + runs-on: ubuntu-latest + + continue-on-error: true + + steps: + - name: Clone + id: checkout + uses: actions/checkout@v6 + + - name: Dependencies + id: depends + run: | + sudo apt-get update + sudo apt-get install build-essential libssl-dev + + - name: Build + id: cmake_build + run: | + cmake -B build \ + -DGGML_RPC=ON + cmake --build build --config Release -j $(nproc) + + - name: Test + id: cmake_test + run: | + cd build + ctest -L main --verbose + + ubuntu-24-vulkan: + runs-on: ${{ 'ubuntu-24.04-arm' || 'ubuntu-24.04' }} + + steps: + - name: Clone + id: checkout + uses: actions/checkout@v6 + + - name: Dependencies + id: depends + run: | + sudo apt-get install -y glslc libvulkan-dev libssl-dev + + - name: Configure + id: cmake_configure + run: | + cmake -B build \ + -DCMAKE_BUILD_TYPE=RelWithDebInfo \ + -DGGML_BACKEND_DL=ON \ + -DGGML_CPU_ALL_VARIANTS=ON \ + -DGGML_VULKAN=ON + + - name: Build + id: cmake_build + run: | + cmake --build build -j $(nproc) + ubuntu-24-webgpu: runs-on: ubuntu-24.04 @@ -209,10 +387,6 @@ jobs: # This is using llvmpipe and runs slower than other backends ctest -L main --verbose --timeout 900 - - name: Setup tmate session - if: ${{ failure() }} - uses: mxschmitt/action-tmate@v3 - ubuntu-24-webgpu-wasm: runs-on: ${{ 'ubuntu-24.04-arm' || 'ubuntu-24.04' }} @@ -246,3 +420,846 @@ jobs: -DEMDAWNWEBGPU_DIR=emdawnwebgpu_pkg cmake --build build-wasm --target test-backend-ops -j $(nproc) + + ubuntu-22-hip: + runs-on: ubuntu-22.04 + container: rocm/dev-ubuntu-22.04:6.1.2 + + steps: + - name: Clone + id: checkout + uses: actions/checkout@v6 + + - name: Dependencies + id: depends + run: | + sudo apt-get update + sudo apt-get install -y build-essential git cmake rocblas-dev hipblas-dev libssl-dev rocwmma-dev + + - name: ccache + uses: ggml-org/ccache-action@v1.2.21 + with: + key: ubuntu-22-hip + evict-old-files: 1d + save: ${{ github.event_name == 'push' && github.ref == 'refs/heads/master' }} + + - name: Build with native CMake HIP support + id: cmake_build + run: | + cmake -B build -S . \ + -DCMAKE_HIP_COMPILER="$(hipconfig -l)/clang" \ + -DGGML_HIP_ROCWMMA_FATTN=ON \ + -DGGML_HIP=ON + cmake --build build --config Release -j $(nproc) + + ubuntu-22-musa: + runs-on: ubuntu-22.04 + container: mthreads/musa:rc4.3.0-devel-ubuntu22.04-amd64 + + steps: + - name: Clone + id: checkout + uses: actions/checkout@v6 + + - name: Dependencies + id: depends + run: | + apt-get update + apt-get install -y build-essential git cmake libssl-dev + + - name: ccache + uses: ggml-org/ccache-action@v1.2.21 + with: + key: ubuntu-22-musa + evict-old-files: 1d + save: ${{ github.event_name == 'push' && github.ref == 'refs/heads/master' }} + + - name: Build with native CMake MUSA support + id: cmake_build + run: | + cmake -B build -S . \ + -DGGML_MUSA=ON + cmake --build build --config Release -j $(nproc) + + ubuntu-22-sycl: + runs-on: ubuntu-22.04 + + continue-on-error: true + + steps: + - uses: actions/checkout@v6 + + - name: add oneAPI to apt + shell: bash + run: | + cd /tmp + wget https://apt.repos.intel.com/intel-gpg-keys/GPG-PUB-KEY-INTEL-SW-PRODUCTS.PUB + sudo apt-key add GPG-PUB-KEY-INTEL-SW-PRODUCTS.PUB + rm GPG-PUB-KEY-INTEL-SW-PRODUCTS.PUB + sudo add-apt-repository "deb https://apt.repos.intel.com/oneapi all main" + + - name: install oneAPI dpcpp compiler + shell: bash + run: | + sudo apt update + sudo apt install intel-oneapi-compiler-dpcpp-cpp libssl-dev + + - name: install oneAPI MKL library + shell: bash + run: | + sudo apt install intel-oneapi-mkl-devel + + - name: Clone + id: checkout + uses: actions/checkout@v6 + + - name: ccache + uses: ggml-org/ccache-action@v1.2.21 + with: + key: ubuntu-22-sycl + evict-old-files: 1d + save: ${{ github.event_name == 'push' && github.ref == 'refs/heads/master' }} + + - name: Build + id: cmake_build + run: | + source /opt/intel/oneapi/setvars.sh + cmake -B build \ + -DGGML_SYCL=ON \ + -DCMAKE_C_COMPILER=icx \ + -DCMAKE_CXX_COMPILER=icpx + cmake --build build --config Release -j $(nproc) + + ubuntu-22-sycl-fp16: + runs-on: ubuntu-22.04 + + continue-on-error: true + + steps: + - uses: actions/checkout@v6 + + - name: add oneAPI to apt + shell: bash + run: | + cd /tmp + wget https://apt.repos.intel.com/intel-gpg-keys/GPG-PUB-KEY-INTEL-SW-PRODUCTS.PUB + sudo apt-key add GPG-PUB-KEY-INTEL-SW-PRODUCTS.PUB + rm GPG-PUB-KEY-INTEL-SW-PRODUCTS.PUB + sudo add-apt-repository "deb https://apt.repos.intel.com/oneapi all main" + + - name: install oneAPI dpcpp compiler + shell: bash + run: | + sudo apt update + sudo apt install intel-oneapi-compiler-dpcpp-cpp libssl-dev + + - name: install oneAPI MKL library + shell: bash + run: | + sudo apt install intel-oneapi-mkl-devel + + - name: Clone + id: checkout + uses: actions/checkout@v6 + + - name: ccache + uses: ggml-org/ccache-action@v1.2.21 + with: + key: ubuntu-22-sycl-fp16 + evict-old-files: 1d + save: ${{ github.event_name == 'push' && github.ref == 'refs/heads/master' }} + + - name: Build + id: cmake_build + run: | + source /opt/intel/oneapi/setvars.sh + cmake -B build \ + -DGGML_SYCL=ON \ + -DCMAKE_C_COMPILER=icx \ + -DCMAKE_CXX_COMPILER=icpx \ + -DGGML_SYCL_F16=ON + cmake --build build --config Release -j $(nproc) + + ubuntu-24-openvino: + name: ubuntu-24-openvino-${{ matrix.openvino_device }} + strategy: + matrix: + include: + - variant: cpu + runner: '"ubuntu-24.04"' + openvino_device: "CPU" + - variant: gpu + runner: '["self-hosted","Linux","X64","Intel"]' + openvino_device: "GPU" + + runs-on: ${{ fromJSON(matrix.runner) }} + + env: + # Sync versions in build.yml, build-self-hosted.yml, release.yml, build-cache.yml, .devops/openvino.Dockerfile + OPENVINO_VERSION_MAJOR: "2026.0" + OPENVINO_VERSION_FULL: "2026.0.0.20965.c6d6a13a886" + + steps: + - name: Clone + id: checkout + uses: actions/checkout@v6 + + - name: ccache + if: runner.environment == 'github-hosted' + uses: ggml-org/ccache-action@v1.2.21 + with: + key: ubuntu-24-openvino-${{ matrix.variant }}-no-preset-v1 + evict-old-files: 1d + save: ${{ github.event_name == 'push' && github.ref == 'refs/heads/master' }} + + - name: Dependencies + id: depends + run: | + sudo apt-get update + sudo apt-get install -y build-essential libssl-dev libtbb12 cmake ninja-build python3-pip + sudo apt-get install -y ocl-icd-opencl-dev opencl-headers opencl-clhpp-headers intel-opencl-icd + + - name: Use OpenVINO Toolkit Cache + if: runner.environment == 'github-hosted' + uses: actions/cache@v5 + id: cache-openvino + with: + path: ./openvino_toolkit + key: openvino-toolkit-v${{ env.OPENVINO_VERSION_FULL }}-${{ runner.os }} + + - name: Setup OpenVINO Toolkit + if: steps.cache-openvino.outputs.cache-hit != 'true' + uses: ./.github/actions/linux-setup-openvino + with: + path: ./openvino_toolkit + version_major: ${{ env.OPENVINO_VERSION_MAJOR }} + version_full: ${{ env.OPENVINO_VERSION_FULL }} + + - name: Install OpenVINO dependencies + run: | + cd ./openvino_toolkit + chmod +x ./install_dependencies/install_openvino_dependencies.sh + echo "Y" | sudo -E ./install_dependencies/install_openvino_dependencies.sh + + - name: Build + id: cmake_build + run: | + source ./openvino_toolkit/setupvars.sh + cmake -B build/ReleaseOV -G Ninja \ + -DCMAKE_BUILD_TYPE=Release \ + -DGGML_OPENVINO=ON + cmake --build build/ReleaseOV --config Release -j $(nproc) + + - name: Test + id: cmake_test + # TODO: fix and re-enable the `test-llama-archs` test below + run: | + cd ${{ github.workspace }} + if [ "${{ matrix.openvino_device }}" = "GPU" ]; then + export GGML_OPENVINO_DEVICE=GPU + fi + ctest --test-dir build/ReleaseOV -L main -E "test-llama-archs" --verbose --timeout 2000 + + windows-latest: + runs-on: windows-2025 + + env: + OPENBLAS_VERSION: 0.3.23 + SDE_VERSION: 9.33.0-2024-01-07 + VULKAN_VERSION: 1.4.313.2 + + strategy: + matrix: + include: + - build: 'cpu-x64 (static)' + arch: 'x64' + defines: '-G "Ninja Multi-Config" -D CMAKE_TOOLCHAIN_FILE=cmake/x64-windows-llvm.cmake -DGGML_NATIVE=OFF -DLLAMA_BUILD_SERVER=ON -DGGML_RPC=ON -DBUILD_SHARED_LIBS=OFF' + - build: 'openblas-x64' + arch: 'x64' + defines: '-G "Ninja Multi-Config" -D CMAKE_TOOLCHAIN_FILE=cmake/x64-windows-llvm.cmake -DGGML_NATIVE=OFF -DLLAMA_BUILD_SERVER=ON -DGGML_RPC=ON -DGGML_BACKEND_DL=ON -DGGML_CPU_ALL_VARIANTS=ON -DGGML_OPENMP=OFF -DGGML_BLAS=ON -DGGML_BLAS_VENDOR=OpenBLAS -DBLAS_INCLUDE_DIRS="$env:RUNNER_TEMP/openblas/include" -DBLAS_LIBRARIES="$env:RUNNER_TEMP/openblas/lib/openblas.lib"' + - build: 'vulkan-x64' + arch: 'x64' + defines: '-DCMAKE_BUILD_TYPE=Release -DGGML_NATIVE=OFF -DLLAMA_BUILD_SERVER=ON -DGGML_RPC=ON -DGGML_BACKEND_DL=ON -DGGML_CPU_ALL_VARIANTS=ON -DGGML_VULKAN=ON' + - build: 'llvm-arm64' + arch: 'arm64' + defines: '-G "Ninja Multi-Config" -D CMAKE_TOOLCHAIN_FILE=cmake/arm64-windows-llvm.cmake -DGGML_NATIVE=OFF -DLLAMA_BUILD_SERVER=ON' + - build: 'llvm-arm64-opencl-adreno' + arch: 'arm64' + defines: '-G "Ninja Multi-Config" -D CMAKE_TOOLCHAIN_FILE=cmake/arm64-windows-llvm.cmake -DCMAKE_PREFIX_PATH="$env:RUNNER_TEMP/opencl-arm64-release" -DGGML_OPENCL=ON -DGGML_OPENCL_USE_ADRENO_KERNELS=ON' + + steps: + - name: Clone + id: checkout + uses: actions/checkout@v6 + + - name: ccache + uses: ggml-org/ccache-action@v1.2.21 + with: + key: windows-latest-${{ matrix.build }} + variant: ccache + evict-old-files: 1d + save: ${{ github.event_name == 'push' && github.ref == 'refs/heads/master' }} + + - name: Download OpenBLAS + id: get_openblas + if: ${{ matrix.build == 'openblas-x64' }} + run: | + curl.exe -o $env:RUNNER_TEMP/openblas.zip -L "https://github.com/xianyi/OpenBLAS/releases/download/v${env:OPENBLAS_VERSION}/OpenBLAS-${env:OPENBLAS_VERSION}-x64.zip" + curl.exe -o $env:RUNNER_TEMP/OpenBLAS.LICENSE.txt -L "https://github.com/xianyi/OpenBLAS/raw/v${env:OPENBLAS_VERSION}/LICENSE" + mkdir $env:RUNNER_TEMP/openblas + tar.exe -xvf $env:RUNNER_TEMP/openblas.zip -C $env:RUNNER_TEMP/openblas + $vcdir = $(vswhere -latest -products * -requires Microsoft.VisualStudio.Component.VC.Tools.x86.x64 -property installationPath) + $msvc = $(join-path $vcdir $('VC\Tools\MSVC\'+$(gc -raw $(join-path $vcdir 'VC\Auxiliary\Build\Microsoft.VCToolsVersion.default.txt')).Trim())) + $lib = $(join-path $msvc 'bin\Hostx64\x64\lib.exe') + & $lib /machine:x64 "/def:${env:RUNNER_TEMP}/openblas/lib/libopenblas.def" "/out:${env:RUNNER_TEMP}/openblas/lib/openblas.lib" /name:openblas.dll + + - name: Install Vulkan SDK + id: get_vulkan + if: ${{ matrix.build == 'vulkan-x64' }} + run: | + curl.exe -o $env:RUNNER_TEMP/VulkanSDK-Installer.exe -L "https://sdk.lunarg.com/sdk/download/${env:VULKAN_VERSION}/windows/vulkansdk-windows-X64-${env:VULKAN_VERSION}.exe" + & "$env:RUNNER_TEMP\VulkanSDK-Installer.exe" --accept-licenses --default-answer --confirm-command install + Add-Content $env:GITHUB_ENV "VULKAN_SDK=C:\VulkanSDK\${env:VULKAN_VERSION}" + Add-Content $env:GITHUB_PATH "C:\VulkanSDK\${env:VULKAN_VERSION}\bin" + + - name: Install Ninja + id: install_ninja + run: | + choco install ninja + + - name: Install OpenCL Headers and Libs + id: install_opencl + if: ${{ matrix.build == 'llvm-arm64-opencl-adreno' }} + run: | + git clone https://github.com/KhronosGroup/OpenCL-Headers + cd OpenCL-Headers + cmake -B build ` + -DBUILD_TESTING=OFF ` + -DOPENCL_HEADERS_BUILD_TESTING=OFF ` + -DOPENCL_HEADERS_BUILD_CXX_TESTS=OFF ` + -DCMAKE_INSTALL_PREFIX="$env:RUNNER_TEMP/opencl-arm64-release" + cmake --build build --target install + git clone https://github.com/KhronosGroup/OpenCL-ICD-Loader + cd OpenCL-ICD-Loader + cmake -B build-arm64-release ` + -A arm64 ` + -DCMAKE_PREFIX_PATH="$env:RUNNER_TEMP/opencl-arm64-release" ` + -DCMAKE_INSTALL_PREFIX="$env:RUNNER_TEMP/opencl-arm64-release" + cmake --build build-arm64-release --target install --config release + + - name: Build + id: cmake_build + run: | + cmake -S . -B build ${{ matrix.defines }} ` + -DLLAMA_BUILD_BORINGSSL=ON + cmake --build build --config Release -j ${env:NUMBER_OF_PROCESSORS} + + - name: Add libopenblas.dll + id: add_libopenblas_dll + if: ${{ matrix.build == 'openblas-x64' }} + run: | + cp $env:RUNNER_TEMP/openblas/bin/libopenblas.dll ./build/bin/Release/openblas.dll + cp $env:RUNNER_TEMP/OpenBLAS.LICENSE.txt ./build/bin/Release/OpenBLAS-${env:OPENBLAS_VERSION}.txt + + - name: Test + id: cmake_test + if: ${{ matrix.arch == 'x64' }} + run: | + cd build + ctest -L main -C Release --verbose --timeout 900 + + # TODO: disabled for now, consider adding tests for all CPU variants instead + # - name: Test (Intel SDE) + # id: cmake_test_sde + # if: ${{ matrix.build == 'avx512-x64' && env.HAS_AVX512F == '0' }} # use Intel SDE for AVX-512 emulation + # run: | + # curl.exe -o $env:RUNNER_TEMP/sde.tar.xz -L "https://downloadmirror.intel.com/813591/sde-external-${env:SDE_VERSION}-win.tar.xz" + # # for some weird reason windows tar doesn't like sde tar.xz + # 7z x "-o${env:RUNNER_TEMP}" $env:RUNNER_TEMP/sde.tar.xz + # 7z x "-o${env:RUNNER_TEMP}" $env:RUNNER_TEMP/sde.tar + # $sde = $(join-path $env:RUNNER_TEMP sde-external-${env:SDE_VERSION}-win/sde.exe) + # cd build + # $env:LLAMA_SKIP_TESTS_SLOW_ON_EMULATOR = 1 + # & $sde -future -- ctest -L main -C Release --verbose --timeout 900 + + ubuntu-latest-cuda: + runs-on: ubuntu-latest + container: nvidia/cuda:12.6.2-devel-ubuntu24.04 + + steps: + - name: Clone + id: checkout + uses: actions/checkout@v6 + + - name: Install dependencies + env: + DEBIAN_FRONTEND: noninteractive + run: | + apt update + apt install -y cmake build-essential ninja-build libgomp1 git libssl-dev + + - name: ccache + uses: ggml-org/ccache-action@v1.2.21 + with: + key: ubuntu-latest-cuda + evict-old-files: 1d + save: ${{ github.event_name == 'push' && github.ref == 'refs/heads/master' }} + + - name: Build with CMake + # TODO: Remove GGML_CUDA_CUB_3DOT2 flag once CCCL 3.2 is bundled within CTK and that CTK version is used in this project + run: | + cmake -S . -B build -G Ninja \ + -DLLAMA_FATAL_WARNINGS=ON \ + -DCMAKE_BUILD_TYPE=Release \ + -DCMAKE_CUDA_ARCHITECTURES=89-real \ + -DCMAKE_EXE_LINKER_FLAGS=-Wl,--allow-shlib-undefined \ + -DGGML_NATIVE=OFF \ + -DGGML_CUDA=ON \ + -DGGML_CUDA_CUB_3DOT2=ON + cmake --build build + + windows-2022-cuda: + runs-on: windows-2022 + + strategy: + matrix: + cuda: ['12.4'] + + steps: + - name: Clone + id: checkout + uses: actions/checkout@v6 + + - name: Install ccache + uses: ggml-org/ccache-action@v1.2.21 + with: + key: windows-cuda-${{ matrix.cuda }} + variant: ccache + evict-old-files: 1d + save: ${{ github.event_name == 'push' && github.ref == 'refs/heads/master' }} + + - name: Install Cuda Toolkit + uses: ./.github/actions/windows-setup-cuda + with: + cuda_version: ${{ matrix.cuda }} + + - name: Install Ninja + id: install_ninja + run: | + choco install ninja + + - name: Build + id: cmake_build + shell: cmd + # TODO: Remove GGML_CUDA_CUB_3DOT2 flag once CCCL 3.2 is bundled within CTK and that CTK version is used in this project + run: | + call "C:\Program Files\Microsoft Visual Studio\2022\Enterprise\VC\Auxiliary\Build\vcvarsall.bat" x64 + cmake -S . -B build -G "Ninja Multi-Config" ^ + -DLLAMA_BUILD_SERVER=ON ^ + -DLLAMA_BUILD_BORINGSSL=ON ^ + -DGGML_NATIVE=OFF ^ + -DGGML_BACKEND_DL=ON ^ + -DGGML_CPU_ALL_VARIANTS=ON ^ + -DGGML_CUDA=ON ^ + -DGGML_RPC=ON ^ + -DGGML_CUDA_CUB_3DOT2=ON + set /A NINJA_JOBS=%NUMBER_OF_PROCESSORS%-1 + cmake --build build --config Release -j %NINJA_JOBS% -t ggml + cmake --build build --config Release + + windows-latest-sycl: + runs-on: windows-2022 + + defaults: + run: + shell: bash + + env: + WINDOWS_BASEKIT_URL: https://registrationcenter-download.intel.com/akdlm/IRC_NAS/24751ead-ddc5-4479-b9e6-f9fe2ff8b9f2/intel-deep-learning-essentials-2025.2.1.25_offline.exe + WINDOWS_DPCPP_MKL: intel.oneapi.win.cpp-dpcpp-common:intel.oneapi.win.mkl.devel:intel.oneapi.win.dnnl:intel.oneapi.win.tbb.devel + ONEAPI_ROOT: "C:/Program Files (x86)/Intel/oneAPI" + steps: + - name: Clone + id: checkout + uses: actions/checkout@v6 + + - name: ccache + uses: ggml-org/ccache-action@v1.2.21 + with: + key: windows-latest-sycl + variant: ccache + evict-old-files: 1d + save: ${{ github.event_name == 'push' && github.ref == 'refs/heads/master' }} + + - name: Install + run: | + scripts/install-oneapi.bat $WINDOWS_BASEKIT_URL $WINDOWS_DPCPP_MKL + + # TODO: add ssl support ; we will also need to modify win-build-sycl.bat to accept user-specified args + + - name: Build + id: cmake_build + run: examples/sycl/win-build-sycl.bat + + windows-latest-hip: + runs-on: windows-2022 + + env: + # Make sure this is in sync with build-cache.yml + HIPSDK_INSTALLER_VERSION: "26.Q1" + + steps: + - name: Clone + id: checkout + uses: actions/checkout@v6 + + - name: Grab rocWMMA package + id: grab_rocwmma + run: | + curl -o rocwmma.deb "https://repo.radeon.com/rocm/apt/7.2/pool/main/r/rocwmma-dev/rocwmma-dev_2.2.0.70200-43~24.04_amd64.deb" + 7z x rocwmma.deb + 7z x data.tar + + - name: Use ROCm Installation Cache + uses: actions/cache@v5 + id: cache-rocm + with: + path: C:\Program Files\AMD\ROCm + key: rocm-${{ env.HIPSDK_INSTALLER_VERSION }}-${{ runner.os }} + + - name: Setup ROCm + if: steps.cache-rocm.outputs.cache-hit != 'true' + uses: ./.github/actions/windows-setup-rocm + with: + version: ${{ env.HIPSDK_INSTALLER_VERSION }} + + - name: Verify ROCm + id: verify + run: | + # Find and test ROCm installation + $clangPath = Get-ChildItem 'C:\Program Files\AMD\ROCm\*\bin\clang.exe' | Select-Object -First 1 + if (-not $clangPath) { + Write-Error "ROCm installation not found" + exit 1 + } + & $clangPath.FullName --version + + - name: Install ccache + uses: ggml-org/ccache-action@v1.2.21 + with: + key: ${{ github.job }} + evict-old-files: 1d + save: ${{ github.event_name == 'push' && github.ref == 'refs/heads/master' }} + + - name: Build + id: cmake_build + run: | + $env:HIP_PATH=$(Resolve-Path 'C:\Program Files\AMD\ROCm\*\bin\clang.exe' | split-path | split-path) + $env:CMAKE_PREFIX_PATH="${env:HIP_PATH}" + cmake -G "Unix Makefiles" -B build -S . ` + -DCMAKE_C_COMPILER="${env:HIP_PATH}\bin\clang.exe" ` + -DCMAKE_CXX_COMPILER="${env:HIP_PATH}\bin\clang++.exe" ` + -DCMAKE_CXX_FLAGS="-I$($PWD.Path.Replace('\', '/'))/opt/rocm-7.2.0/include/" ` + -DCMAKE_BUILD_TYPE=Release ` + -DLLAMA_BUILD_BORINGSSL=ON ` + -DROCM_DIR="${env:HIP_PATH}" ` + -DGGML_HIP=ON ` + -DGGML_HIP_ROCWMMA_FATTN=ON ` + -DGGML_RPC=ON + cmake --build build -j ${env:NUMBER_OF_PROCESSORS} + + ubuntu-cpu-riscv64-native: + runs-on: RISCV64 + + steps: + - name: Install dependencies + run: | + sudo apt-get update + + # Install necessary packages + sudo apt-get install -y libatomic1 libtsan2 gcc-14 g++-14 rustup cmake build-essential libssl-dev wget ccache git-lfs + + # Set gcc-14 and g++-14 as the default compilers + sudo update-alternatives --install /usr/bin/gcc gcc /usr/bin/gcc-14 100 + sudo update-alternatives --install /usr/bin/g++ g++ /usr/bin/g++-14 100 + sudo ln -sf /usr/bin/gcc-14 /usr/bin/gcc + sudo ln -sf /usr/bin/g++-14 /usr/bin/g++ + + # Install Rust stable version + rustup install stable + rustup default stable + + git lfs install + + - name: Clone + id: checkout + uses: actions/checkout@v6 + + - name: Check environment + run: | + uname -a + gcc --version + g++ --version + ldd --version + cmake --version + rustc --version + + - name: Setup ccache + run: | + # Set unique cache directory for this job + export CCACHE_DIR="$HOME/.ccache/cpu-cmake-rv64-native" + mkdir -p "$CCACHE_DIR" + + # Configure ccache for optimal performance + ccache --set-config=max_size=5G + ccache --set-config=compression=true + ccache --set-config=compression_level=6 + ccache --set-config=cache_dir="$CCACHE_DIR" + + # Enable more aggressive caching + ccache --set-config=sloppiness=file_macro,time_macros,include_file_mtime,include_file_ctime + ccache --set-config=hash_dir=false + + # Export for subsequent steps + echo "CCACHE_DIR=$CCACHE_DIR" >> $GITHUB_ENV + echo "PATH=/usr/lib/ccache:$PATH" >> $GITHUB_ENV + + - name: Build + id: cmake_build + run: | + cmake -B build \ + -DCMAKE_BUILD_TYPE=Release \ + -DGGML_OPENMP=OFF \ + -DLLAMA_BUILD_EXAMPLES=ON \ + -DLLAMA_BUILD_TOOLS=ON \ + -DLLAMA_BUILD_TESTS=ON \ + -DCMAKE_C_COMPILER_LAUNCHER=ccache \ + -DCMAKE_CXX_COMPILER_LAUNCHER=ccache \ + -DGGML_RPC=ON \ + -DCMAKE_C_COMPILER=riscv64-linux-gnu-gcc-14 \ + -DCMAKE_CXX_COMPILER=riscv64-linux-gnu-g++-14 + + cmake --build build --config Release -j $(nproc) + + - name: Test + id: cmake_test + run: | + cd build + ctest -L main --verbose --timeout 900 + + - name: Test llama2c conversion + id: llama2c_test + run: | + cd build + echo "Fetch tokenizer" + wget https://huggingface.co/karpathy/tinyllamas/resolve/main/stories260K/tok512.bin + echo "Fetch llama2c model" + wget https://huggingface.co/karpathy/tinyllamas/resolve/main/stories260K/stories260K.bin + ./bin/llama-convert-llama2c-to-ggml --copy-vocab-from-model ./tok512.bin --llama2c-model stories260K.bin --llama2c-output-model stories260K.gguf + ./bin/llama-completion -m stories260K.gguf -p "One day, Lily met a Shoggoth" -n 500 -c 256 + +# TODO: simplify the following workflows using a matrix +# TODO: run lighter CI on PRs and the full CI only on master (if needed) + ggml-ci-x64-cpu-low-perf: + runs-on: ubuntu-22.04 + + steps: + - name: Clone + id: checkout + uses: actions/checkout@v6 + + - name: ccache + uses: ggml-org/ccache-action@v1.2.21 + with: + key: ggml-ci-x64-cpu-low-perf + evict-old-files: 1d + save: ${{ github.event_name == 'push' && github.ref == 'refs/heads/master' }} + + - name: Dependencies + id: depends + run: | + sudo apt-get update + sudo apt-get install build-essential + + - name: Test + id: ggml-ci + run: | + LLAMA_ARG_THREADS=$(nproc) GG_BUILD_LOW_PERF=1 bash ./ci/run.sh ./tmp/results ./tmp/mnt + + ggml-ci-arm64-cpu-low-perf: + runs-on: ubuntu-22.04-arm + + steps: + - name: Clone + id: checkout + uses: actions/checkout@v6 + + - name: ccache + uses: ggml-org/ccache-action@v1.2.21 + with: + key: ggml-ci-arm64-cpu-low-perf + evict-old-files: 1d + save: ${{ github.event_name == 'push' && github.ref == 'refs/heads/master' }} + + - name: Dependencies + id: depends + run: | + sudo apt-get update + sudo apt-get install build-essential + + - name: Test + id: ggml-ci + run: | + LLAMA_ARG_THREADS=$(nproc) GG_BUILD_LOW_PERF=1 bash ./ci/run.sh ./tmp/results ./tmp/mnt + + ggml-ci-x64-cpu-high-perf: + runs-on: ubuntu-22.04 + + steps: + - name: Clone + id: checkout + uses: actions/checkout@v6 + + - name: ccache + uses: ggml-org/ccache-action@v1.2.21 + with: + key: ggml-ci-x64-cpu-high-perf + evict-old-files: 1d + save: ${{ github.event_name == 'push' && github.ref == 'refs/heads/master' }} + + - name: Dependencies + id: depends + run: | + sudo apt-get update + sudo apt-get install build-essential + + - name: Test + id: ggml-ci + run: | + LLAMA_ARG_THREADS=$(nproc) GG_BUILD_HIGH_PERF=1 bash ./ci/run.sh ./tmp/results ./tmp/mnt + + ggml-ci-arm64-cpu-high-perf: + runs-on: ubuntu-22.04-arm + + steps: + - name: Clone + id: checkout + uses: actions/checkout@v6 + + - name: ccache + uses: ggml-org/ccache-action@v1.2.21 + with: + key: ggml-ci-arm64-cpu-high-perf + evict-old-files: 1d + save: ${{ github.event_name == 'push' && github.ref == 'refs/heads/master' }} + + - name: Dependencies + id: depends + run: | + sudo apt-get update + sudo apt-get install build-essential + + - name: Test + id: ggml-ci + run: | + LLAMA_ARG_THREADS=$(nproc) GG_BUILD_HIGH_PERF=1 GG_BUILD_NO_SVE=1 GG_BUILD_NO_BF16=1 GG_BUILD_EXTRA_TESTS_0=1 bash ./ci/run.sh ./tmp/results ./tmp/mnt + + ggml-ci-arm64-cpu-high-perf-sve: + runs-on: ubuntu-22.04-arm + + steps: + - name: Clone + id: checkout + uses: actions/checkout@v6 + + - name: ccache + uses: ggml-org/ccache-action@v1.2.21 + with: + key: ggml-ci-arm64-cpu-high-perf-sve + evict-old-files: 1d + save: ${{ github.event_name == 'push' && github.ref == 'refs/heads/master' }} + + - name: Dependencies + id: depends + run: | + sudo apt-get update + sudo apt-get install build-essential + + - name: Test + id: ggml-ci + run: | + LLAMA_ARG_THREADS=$(nproc) GG_BUILD_NO_BF16=1 GG_BUILD_EXTRA_TESTS_0=1 bash ./ci/run.sh ./tmp/results ./tmp/mnt + + ggml-ci-arm64-cpu-kleidiai: + runs-on: ubuntu-22.04-arm + + steps: + - name: Clone + id: checkout + uses: actions/checkout@v6 + + - name: ccache + uses: ggml-org/ccache-action@v1.2.21 + with: + key: ggml-ci-arm64-cpu-kleidiai + evict-old-files: 1d + save: ${{ github.event_name == 'push' && github.ref == 'refs/heads/master' }} + + - name: Dependencies + id: depends + run: | + sudo apt-get update + sudo apt-get install -y build-essential + + - name: Test + id: ggml-ci + run: | + GG_BUILD_KLEIDIAI=1 GG_BUILD_EXTRA_TESTS_0=1 bash ./ci/run.sh ./tmp/results ./tmp/mnt + + ggml-ci-arm64-cpu-kleidiai-graviton4: + runs-on: ah-ubuntu_22_04-c8g_8x + + steps: + - name: Clone + id: checkout + uses: actions/checkout@v6 + + - name: Dependencies + id: depends + run: | + set -euxo pipefail + sudo apt-get update + sudo DEBIAN_FRONTEND=noninteractive NEEDRESTART_MODE=a \ + apt-get install -y \ + build-essential \ + python3-venv \ + gpg \ + wget \ + time \ + git-lfs + + git lfs install + + # install the latest cmake + sudo install -d /usr/share/keyrings + wget -O - https://apt.kitware.com/keys/kitware-archive-latest.asc \ + | gpg --dearmor \ + | sudo tee /usr/share/keyrings/kitware-archive-keyring.gpg >/dev/null + echo 'deb [signed-by=/usr/share/keyrings/kitware-archive-keyring.gpg] https://apt.kitware.com/ubuntu/ jammy main' \ + | sudo tee /etc/apt/sources.list.d/kitware.list + sudo apt-get update + sudo apt-get install -y cmake + + - name: ccache + uses: ggml-org/ccache-action@v1.2.21 + with: + key: ggml-ci-arm64-cpu-kleidiai-graviton4 + evict-old-files: 1d + save: ${{ github.event_name == 'push' && github.ref == 'refs/heads/master' }} + + - name: Test + id: ggml-ci + run: | + GG_BUILD_KLEIDIAI=1 \ + GG_BUILD_EXTRA_TESTS_0=1 \ + bash ./ci/run.sh ./tmp/results ./tmp/mnt From cfb7b69c7e2839ee2d7e5af056cf782911caae3f Mon Sep 17 00:00:00 2001 From: Nikhil Jain Date: Thu, 19 Mar 2026 17:55:02 -0700 Subject: [PATCH 9/9] Update after merge --- ggml/src/ggml-webgpu/ggml-webgpu.cpp | 10 +++++----- 1 file changed, 5 insertions(+), 5 deletions(-) diff --git a/ggml/src/ggml-webgpu/ggml-webgpu.cpp b/ggml/src/ggml-webgpu/ggml-webgpu.cpp index b19e2c57efd..1020ff260a7 100644 --- a/ggml/src/ggml-webgpu/ggml-webgpu.cpp +++ b/ggml/src/ggml-webgpu/ggml-webgpu.cpp @@ -1051,7 +1051,7 @@ static webgpu_command ggml_webgpu_set(webgpu_context & ctx, ggml_tensor * src0, .size = ggml_webgpu_tensor_binding_size(ctx, dst) }); uint32_t wg_x = CEIL_DIV(ne, decisions->wg_size); - return ggml_backend_webgpu_build(ctx->global_ctx, ctx->param_buf_pool, pipeline, params, entries, wg_x); + return ggml_backend_webgpu_build(ctx->global_ctx, ctx, ctx->param_buf_pool, pipeline, params, entries, wg_x); } static webgpu_command ggml_webgpu_pad(webgpu_context & ctx, ggml_tensor * src, ggml_tensor * dst) { @@ -1167,7 +1167,7 @@ static webgpu_command ggml_webgpu_solve_tri(webgpu_context & ctx, const uint32_t wg_x = CEIL_DIV((uint32_t) src1->ne[0], decisions->wg_size); const uint32_t wg_y = (uint32_t) (dst->ne[2] * dst->ne[3]); - return ggml_backend_webgpu_build(ctx->global_ctx, ctx->param_buf_pool, pipeline, params, entries, wg_x, wg_y); + return ggml_backend_webgpu_build(ctx->global_ctx, ctx, ctx->param_buf_pool, pipeline, params, entries, wg_x, wg_y); } static webgpu_command ggml_webgpu_ssm_conv(webgpu_context & ctx, @@ -1223,7 +1223,7 @@ static webgpu_command ggml_webgpu_ssm_conv(webgpu_context & ctx, const uint32_t wg_x = CEIL_DIV((uint32_t) src0->ne[1], decisions->block_size); const uint32_t wg_y = token_tiles * (uint32_t) dst->ne[2]; - return ggml_backend_webgpu_build(ctx->global_ctx, ctx->param_buf_pool, pipeline, params, entries, wg_x, wg_y); + return ggml_backend_webgpu_build(ctx->global_ctx, ctx, ctx->param_buf_pool, pipeline, params, entries, wg_x, wg_y); } static webgpu_command ggml_webgpu_gated_delta_net(webgpu_context & ctx, @@ -1308,7 +1308,7 @@ static webgpu_command ggml_webgpu_gated_delta_net(webgpu_context & ctx, .size = ggml_webgpu_tensor_binding_size(ctx, dst) } }; - return ggml_backend_webgpu_build(ctx->global_ctx, ctx->param_buf_pool, pipeline, params, entries, h, n_seqs); + return ggml_backend_webgpu_build(ctx->global_ctx, ctx, ctx->param_buf_pool, pipeline, params, entries, h, n_seqs); } static std::optional ggml_webgpu_set_rows(webgpu_context & ctx, @@ -2018,7 +2018,7 @@ static webgpu_command ggml_webgpu_row_norm(webgpu_context & ctx, ggml_tensor * s }; webgpu_pipeline pipeline = ctx->shader_lib->get_row_norm_pipeline(shader_lib_ctx); - return ggml_backend_webgpu_build(ctx->global_ctx, ctx->param_buf_pool, pipeline, params, entries, ggml_nrows(src)); + return ggml_backend_webgpu_build(ctx->global_ctx, ctx, ctx->param_buf_pool, pipeline, params, entries, ggml_nrows(src)); } static webgpu_command ggml_webgpu_rope(webgpu_context & ctx,