Replace pragma once with header guard

Author: WoosukKwon

csrc/attention/attention_dtypes.cuh

@@ -1,5 +1,8 @@
-#pragma once
+#ifndef CACHEFLOW_ATTENTION_ATTENTION_DTYPES_H_
+#define CACHEFLOW_ATTENTION_ATTENTION_DTYPES_H_
 
 #include "attention_generic.cuh"
 #include "dtype_float16.cuh"
 #include "dtype_float32.cuh"
+
+#endif  // CACHEFLOW_ATTENTION_ATTENTION_DTYPES_H_

csrc/attention/attention_generic.cuh

@@ -1,4 +1,5 @@
-#pragma once
+#ifndef CACHEFLOW_ATTENTION_ATTENTION_GENERIC_H_
+#define CACHEFLOW_ATTENTION_ATTENTION_GENERIC_H_
 
 #include <stdint.h>
 
@@ -45,3 +46,5 @@ inline __device__ void zero(T& dst) {
 }
 
 } // namespace cacheflow
+
+#endif  // CACHEFLOW_ATTENTION_ATTENTION_GENERIC_H_

csrc/attention/attention_utils.cuh

@@ -1,4 +1,5 @@
-#pragma once
+#ifndef CACHEFLOW_ATTENTION_ATTENTION_UTILS_H_
+#define CACHEFLOW_ATTENTION_ATTENTION_UTILS_H_
 
 #include "attention_dtypes.cuh"
 
@@ -36,3 +37,5 @@ struct Qk_dot {
 };
 
 } // namespace cacheflow
+
+#endif  // CACHEFLOW_ATTENTION_ATTENTION_UTILS_H_

csrc/attention/dtype_float16.cuh

@@ -1,4 +1,5 @@
-#pragma once
+#ifndef CACHEFLOW_ATTENTION_DTYPE_FLOAT16_H_
+#define CACHEFLOW_ATTENTION_DTYPE_FLOAT16_H_
 
 #include "attention_generic.cuh"
 #include "dtype_float32.cuh"
@@ -424,3 +425,5 @@ inline __device__ Float8_ to_float(uint4 u) {
 }
 
 } // namespace cacheflow
+
+#endif  // CACHEFLOW_ATTENTION_DTYPE_FLOAT16_H_

csrc/attention/dtype_float32.cuh

@@ -1,4 +1,5 @@
-#pragma once
+#ifndef CACHEFLOW_ATTENTION_DTYPE_FLOAT32_H_
+#define CACHEFLOW_ATTENTION_DTYPE_FLOAT32_H_
 
 #include "attention_generic.cuh"
 
@@ -248,3 +249,5 @@ inline __device__ Float8_ to_float(Float8_ u) {
 }
 
 } // namespace cacheflow
+
+#endif  // CACHEFLOW_ATTENTION_DTYPE_FLOAT32_H_

csrc/reduction_utils.cuh

@@ -1,4 +1,5 @@
-#pragma once
+#ifndef CACHEFLOW_REDUCTION_UTILS_H_
+#define CACHEFLOW_REDUCTION_UTILS_H_
 
 namespace cacheflow {
 
@@ -32,3 +33,5 @@ __inline__ __device__ T blockReduceSum(T val) {
 }
 
 } // namespace cacheflow
+
+#endif  // CACHEFLOW_REDUCTION_UTILS_H_