20231088/vllm
1
0
Fork 0
Code Issues Pull Requests Actions 1 Packages Projects Releases Wiki Activity

[Bugfix] Fix marlin kernel crash on H100 (#4218)

Browse Source 
This PR addresses the Marlin kernel H100 crash that was reported here: neuralmagic#187.
The reason for the crash was the inline PTX assembly that introduced the async_copy with streaming behavior. The solution is to use the more standard PTX for async_copy (without the fractional L2 policy for "evict_first"). There is no performance difference between standard async_copy PTX and the previous one.
This commit is contained in:
alexm-nm 2024-04-24 13:35:01 -04:00 committed by GitHub
parent 7923dcad12
commit aae08249ac
No known key found for this signature in database
GPG Key ID: B5690EEEBB952194
1 changed files with 8 additions and 15 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

23
csrc/quantization/marlin/marlin_cuda_kernel.cu
View File

@ -67,20 +67,13 @@ __device__ inline void cp_async4_pred(void *smem_ptr, const void *glob_ptr,
"r"(smem), "l"(glob_ptr), "n"(BYTES)); "r"(smem), "l"(glob_ptr), "n"(BYTES));
} }
// Asynchronous global->shared copy with a cache hint indicating that the values // Asynchronous global->shared copy
// may be evicted immediately; used for quantized weights B, which are only __device__ inline void cp_async4(void *smem_ptr, const void *glob_ptr) {
// accessed precisely once and should thus not pollute the L2 cache which we
// need for inputs A and outputs C.
__device__ inline void cp_async4_stream(void *smem_ptr, const void *glob_ptr) {
const int BYTES = 16; const int BYTES = 16;
uint32_t smem = static_cast<uint32_t>(__cvta_generic_to_shared(smem_ptr)); uint32_t smem = static_cast<uint32_t>(__cvta_generic_to_shared(smem_ptr));
asm volatile( asm volatile("{\n"
"{\n" " cp.async.cg.shared.global [%0], [%1], %2;\n"
" .reg .b64 p;\n" "}\n" :: "r"(smem), "l"(glob_ptr), "n"(BYTES));
" createpolicy.fractional.L2::evict_first.b64 p, 1.0;"
" cp.async.cg.shared.global.L2::cache_hint [%0], [%1], %2, p;\n"
"}\n" ::"r"(smem),
"l"(glob_ptr), "n"(BYTES));
} }
// Async copy fence. // Async copy fence.
@ -448,14 +441,14 @@ Marlin(const int4 *__restrict__ A, // fp16 input matrix of shape mxk
int4 *sh_b_stage = sh_b + b_sh_stage * pipe; int4 *sh_b_stage = sh_b + b_sh_stage * pipe;
#pragma unroll #pragma unroll
for (int i = 0; i < b_sh_wr_iters; i++) { for (int i = 0; i < b_sh_wr_iters; i++) {
cp_async4_stream(&sh_b_stage[b_sh_wr_delta * i + b_sh_wr], B_ptr[i]); cp_async4(&sh_b_stage[b_sh_wr_delta * i + b_sh_wr], B_ptr[i]);
B_ptr[i] += b_gl_rd_delta_o; B_ptr[i] += b_gl_rd_delta_o;
} }
// Only fetch scales if this tile starts a new group // Only fetch scales if this tile starts a new group
if (group_blocks != -1 && pipe % (group_blocks / thread_k_blocks) == 0) { if (group_blocks != -1 && pipe % (group_blocks / thread_k_blocks) == 0) {
int4 *sh_s_stage = sh_s + s_sh_stage * pipe; int4 *sh_s_stage = sh_s + s_sh_stage * pipe;
if (s_sh_wr_pred) if (s_sh_wr_pred)
cp_async4_stream(&sh_s_stage[s_sh_wr], &s[s_gl_rd]); cp_async4(&sh_s_stage[s_sh_wr], &s[s_gl_rd]);
s_gl_rd += s_gl_rd_delta; s_gl_rd += s_gl_rd_delta;
} }
} }
@ -750,7 +743,7 @@ Marlin(const int4 *__restrict__ A, // fp16 input matrix of shape mxk
// write-out // write-out
if (group_blocks == -1 && last) { if (group_blocks == -1 && last) {
if (s_sh_wr_pred) if (s_sh_wr_pred)
cp_async4_stream(&sh_s[s_sh_wr], &s[s_gl_rd]); cp_async4(&sh_s[s_sh_wr], &s[s_gl_rd]);
cp_async_fence(); cp_async_fence();
} }
thread_block_reduce(); thread_block_reduce();