Skip to content

Commit

Permalink
Need one more syncthreads before syncgrid for Mixtral
Browse files Browse the repository at this point in the history
Again not clear why any of this is necessary yet. One possibility is
that the writes are buffered or cached, and we could write directly to
L2 instead but that requires careful code modifications.
  • Loading branch information
zeux committed Jun 7, 2024
1 parent f442e4a commit 74eec37
Showing 1 changed file with 1 addition and 0 deletions.
1 change: 1 addition & 0 deletions src/infer.cu
Original file line number Diff line number Diff line change
Expand Up @@ -483,6 +483,7 @@ __global__ __launch_bounds__(1024, 1) static void kernel_forward(const __grid_co
}
}

__syncthreads(); // TODO: unclear why this is needed for determinism
syncgrid();
coopstage(args.perfstats, 0);

Expand Down

0 comments on commit 74eec37

Please sign in to comment.