diff --git a/csrc/cache_kernels.cu b/csrc/cache_kernels.cu index cf26ae544deaa..a02fcb617910f 100644 --- a/csrc/cache_kernels.cu +++ b/csrc/cache_kernels.cu @@ -451,9 +451,6 @@ __global__ void indexer_k_quant_and_cache_kernel( for (int i = 0; i < VEC_SIZE; i++) { amax = fmaxf(amax, fabsf(float(k_val_ptr[i]))); } -#ifndef USE_ROCM - __syncwarp(); -#endif // Reduced amax for (int mask = 16; mask > 0; mask /= 2) { @@ -463,9 +460,7 @@ __global__ void indexer_k_quant_and_cache_kernel( amax = fmaxf(amax, __shfl_xor_sync(unsigned(-1), amax, mask)); #endif } -#ifndef USE_ROCM - __syncwarp(); -#endif + #if defined(__gfx942__) float scale = fmaxf(amax, 1e-4) / 224.0f; #else