zhangyiss/mlx
1
0
Fork 0
mirror of https://github.com/ml-explore/mlx.git synced 2025-10-31 16:21:27 +08:00
Code Issues Actions 1 Packages Projects Releases Wiki Activity

Fix update gpu

Browse Source
This commit is contained in:
Angelos Katharopoulos
2025-03-04 21:26:31 -08:00
parent 90801467d8
commit f4a5959055
1 changed files with 4 additions and 1 deletions
Download Patch File Download Diff File Expand all files Collapse all files

3
mlx/backend/metal/kernels/fence.metal
View File

@@ -29,7 +29,10 @@ constexpr constant metal::thread_scope thread_scope_system =
[[kernel]] void fence_update( [[kernel]] void fence_update(
volatile coherent(system) device uint* timestamp [[buffer(0)]], volatile coherent(system) device uint* timestamp [[buffer(0)]],
constant uint& value [[buffer(1)]]) { constant uint& value [[buffer(1)]]) {
uint tmp = timestamp[0];
if (tmp < value) {
timestamp[0] = value; timestamp[0] = value;
}
metal::atomic_thread_fence( metal::atomic_thread_fence(
metal::mem_flags::mem_device, metal::mem_flags::mem_device,
metal::memory_order_seq_cst, metal::memory_order_seq_cst,