huggingface/candle
1
0
Fork 0
mirror of https://github.com/huggingface/candle.git synced 2025-06-16 02:38:10 +00:00
Code Issues Packages Projects Releases Wiki Activity

Attempt at fixing M1/M2 metal async copy bug

Browse Source
This commit is contained in:
Ivar Flakstad
2024-09-06 15:59:35 +02:00
parent 8712ceb84f
commit 7ec4f64d38
1 changed files with 3 additions and 1 deletions
Download Patch File Download Diff File Expand all files Collapse all files

4
candle-metal-kernels/src/gemm.metal
View File

@ -607,10 +607,12 @@ METAL_FUNC void async_access_accumulator(threadgroup T *C_block, device T *C,
min(uint(M_group), M - C_offset.y)); min(uint(M_group), M - C_offset.y));
auto C_src = simdgroup_matrix_storage<T>::apply_offset(C, N, C_offset); auto C_src = simdgroup_matrix_storage<T>::apply_offset(C, N, C_offset);
simdgroup_event event;
if (is_store) { if (is_store) {
simdgroup_event event;
event.async_copy(C_src, N, C_tile, C_block, N_group, C_tile); event.async_copy(C_src, N, C_tile, C_block, N_group, C_tile);
simdgroup_event::wait(1, &event);
} else { } else {
simdgroup_event event;
event.async_copy(C_block, N_group, C_tile, C_src, N, C_tile); event.async_copy(C_block, N_group, C_tile, C_src, N, C_tile);
simdgroup_event::wait(1, &event); simdgroup_event::wait(1, &event);
} }