test: re-enable all TMEM ops (alloc, zero, dealloc)
This commit is contained in:
@@ -71,14 +71,14 @@ test_umma_qk_hd16(
|
||||
__syncthreads();
|
||||
uint32_t tmem_base = *sTmemBase;
|
||||
|
||||
// SKIP zeroing TMEM — just test alloc
|
||||
// if (wid == 0) {
|
||||
// for (int col = 0; col < 128; col++) {
|
||||
// tmem_store(tmem_base + col, 0, 0, 0, 0);
|
||||
// }
|
||||
// tmem_fence_store();
|
||||
// }
|
||||
// __syncthreads();
|
||||
// Zero TMEM
|
||||
if (wid == 0) {
|
||||
for (int col = 0; col < 128; col++) {
|
||||
tmem_store(tmem_base + col, 0, 0, 0, 0);
|
||||
}
|
||||
tmem_fence_store();
|
||||
}
|
||||
__syncthreads();
|
||||
|
||||
// ================================================================
|
||||
// Load Q and K into SMEM in canonical layout
|
||||
@@ -149,10 +149,10 @@ test_umma_qk_hd16(
|
||||
}
|
||||
__syncthreads();
|
||||
|
||||
// TMEM dealloc — skipped
|
||||
// if (wid == 0) {
|
||||
// tmem_dealloc(tmem_base, 128);
|
||||
// }
|
||||
// TMEM dealloc
|
||||
if (wid == 0) {
|
||||
tmem_dealloc(tmem_base, 128);
|
||||
}
|
||||
}
|
||||
|
||||
// ==================================================================
|
||||
|
||||
Reference in New Issue
Block a user