(vllm) Add vllm-project/vllm#8776

Author: sasha0552

patches/vllm_pascal/main/8776.patch

@@ -0,0 +1,50 @@
+diff --git a/csrc/custom_all_reduce.cuh b/csrc/custom_all_reduce.cuh
+index 632b579c55af..a2f7e4330000 100644
+--- a/csrc/custom_all_reduce.cuh
++++ b/csrc/custom_all_reduce.cuh
+@@ -131,15 +131,26 @@ DINLINE O downcast(array_t<float, O::size> val) {
+ }
+ 
+ static DINLINE void st_flag_release(FlagType* flag_addr, FlagType flag) {
++#if defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 700
+   asm volatile("st.release.sys.global.u32 [%1], %0;" ::"r"(flag),
+                "l"(flag_addr));
++#else
++  asm volatile("membar.sys; st.volatile.global.u32 [%1], %0;" ::"r"(flag),
++               "l"(flag_addr));
++#endif
+ }
+ 
+ static DINLINE FlagType ld_flag_acquire(FlagType* flag_addr) {
+   FlagType flag;
++#if defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 700
+   asm volatile("ld.acquire.sys.global.u32 %0, [%1];"
+                : "=r"(flag)
+                : "l"(flag_addr));
++#else
++  asm volatile("ld.volatile.global.u32 %0, [%1]; membar.gl;"
++               : "=r"(flag)
++               : "l"(flag_addr));
++#endif
+   return flag;
+ }
+ 
+diff --git a/csrc/custom_all_reduce_test.cu b/csrc/custom_all_reduce_test.cu
+index c8b5d0a013f6..9d83a001c77a 100644
+--- a/csrc/custom_all_reduce_test.cu
++++ b/csrc/custom_all_reduce_test.cu
+@@ -44,7 +44,14 @@
+   } while (0)
+ 
+ __global__ void dummy_kernel() {
++#if defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 700
+   for (int i = 0; i < 100; i++) __nanosleep(1000000);  // 100ms
++#else
++  for (int i = 0; i < 100; i++) {
++    long long int start = clock64();
++    while (clock64() - start < 1000000);  // something like 100ms
++  }
++#endif
+ }
+ 
+ template <typename T>