Skip to content

GitLab

Commit f6653ed9 authored by zhuwenwen's avatar zhuwenwen
Browse files

__syncwarp isn't defined

parent 3280b0a0
Show whitespace changes
Inline Side-by-side
Showing with 9 additions and 9 deletions
csrc/fused_qknorm_rope_kernel.cu
View file @ f6653ed9
...@@ -38,15 +38,15 @@ ...@@ -38,15 +38,15 @@
#ifdef USE_ROCM #ifdef USE_ROCM
#define FINAL_MASK 0xffffffffffffffffULL #define FINAL_MASK 0xffffffffffffffffULL
#if defined(HIP_VERSION) && HIP_VERSION < 70000000 // #if defined(HIP_VERSION) && HIP_VERSION < 70000000
// On ROCm versions before 7.0, __syncwarp isn't defined. The below // // On ROCm versions before 7.0, __syncwarp isn't defined. The below
// implementation is copy/pasted from the implementation in ROCm 7.0 // // implementation is copy/pasted from the implementation in ROCm 7.0
__device__ inline void __syncwarp() { // __device__ inline void __syncwarp() {
__builtin_amdgcn_fence(__ATOMIC_RELEASE, "wavefront"); // __builtin_amdgcn_fence(__ATOMIC_RELEASE, "wavefront");
__builtin_amdgcn_wave_barrier(); // __builtin_amdgcn_wave_barrier();
__builtin_amdgcn_fence(__ATOMIC_ACQUIRE, "wavefront"); // __builtin_amdgcn_fence(__ATOMIC_ACQUIRE, "wavefront");
} // }
#endif // #endif
#else #else
#define FINAL_MASK 0xffffffff #define FINAL_MASK 0xffffffff
#endif #endif
......
Markdown is supported
0% or .
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment