From f4a59590550c1a663524f7b77f320c6c9044f857 Mon Sep 17 00:00:00 2001 From: Angelos Katharopoulos Date: Tue, 4 Mar 2025 21:26:31 -0800 Subject: [PATCH] Fix update gpu --- mlx/backend/metal/kernels/fence.metal | 5 ++++- 1 file changed, 4 insertions(+), 1 deletion(-) diff --git a/mlx/backend/metal/kernels/fence.metal b/mlx/backend/metal/kernels/fence.metal index f736b33b0..2f87018b5 100644 --- a/mlx/backend/metal/kernels/fence.metal +++ b/mlx/backend/metal/kernels/fence.metal @@ -29,7 +29,10 @@ constexpr constant metal::thread_scope thread_scope_system = [[kernel]] void fence_update( volatile coherent(system) device uint* timestamp [[buffer(0)]], constant uint& value [[buffer(1)]]) { - timestamp[0] = value; + uint tmp = timestamp[0]; + if (tmp < value) { + timestamp[0] = value; + } metal::atomic_thread_fence( metal::mem_flags::mem_device, metal::memory_order_seq_cst,