mirror of
https://github.com/NVIDIA/TensorRT-LLM.git
synced 2026-02-18 16:55:08 +08:00
[https://nvbugs/5829830][fix] Declare the var in the correct scope (#11066)
Signed-off-by: ziyixiong-nv <219238287+ziyixiong-nv@users.noreply.github.com> Signed-off-by: Wangshanshan <30051912+dominicshanshan@users.noreply.github.com>
This commit is contained in:
parent
d9f787a8d2
commit
5d73194ffb
@ -1138,6 +1138,10 @@ CUBIN_EXPORT __global__
|
||||
auto& xBar = smem.xBar[idxXBuf];
|
||||
auto& vBar = smem.vBar[idxVBuf];
|
||||
auto const& vBuf = smem.vBuf(idxVBuf);
|
||||
#if !SWAP_AB
|
||||
CtaBarrierPair& vtBar = smem.vtBar[idxVBuf];
|
||||
auto& vtBuf = smem.vtBuf(idxVBuf);
|
||||
#endif
|
||||
xBar.produced.arrive_and_wait();
|
||||
#if SKIP_SOFTMAX_ATTN
|
||||
bool shouldSkipSoftmaxAttn = smem.skipSoftmaxVotesGemm0ToGemm1[idxXBuf]; // guarded by xBar
|
||||
@ -1153,8 +1157,6 @@ CUBIN_EXPORT __global__
|
||||
{
|
||||
arrive_tx_and_wait(vBar.produced, exactDiv(sizeof(SharedMem::VBuffer), gemm1NbThrds));
|
||||
#if !SWAP_AB
|
||||
CtaBarrierPair& vtBar = smem.vtBar[idxVBuf];
|
||||
auto& vtBuf = smem.vtBuf(idxVBuf);
|
||||
vtBar.consumed.arrive_and_wait();
|
||||
transposeVTile(warpRank, laneId(), vtBuf, vBuf);
|
||||
vBar.consumed.arrive();
|
||||
|
||||
Loading…
Reference in New Issue
Block a user