Group all archs in the same file

facebookresearch · Sep 6, 2022 · 6018600 · 6018600
1 parent 02062f7
commit 6018600
Show file tree

Hide file tree

Showing 21 changed files with 37 additions and 61 deletions.
diff --git a/xformers/components/attention/csrc/cuda/mem_eff_attention/kernels/backward_f16.cu b/xformers/components/attention/csrc/cuda/mem_eff_attention/kernels/backward_f16.cu
@@ -0,0 +1,6 @@
+// This file is auto-generated. See "generate_kernels.sh"
+#include "../kernel_backward.h"
+INSTANTIATE_ATTENTION_KERNEL_BACKWARD_SM50(cutlass::half_t, false);
+INSTANTIATE_ATTENTION_KERNEL_BACKWARD_SM70(cutlass::half_t, false);
+INSTANTIATE_ATTENTION_KERNEL_BACKWARD_SM75(cutlass::half_t, false);
+INSTANTIATE_ATTENTION_KERNEL_BACKWARD_SM80(cutlass::half_t, false);
diff --git a/xformers/components/attention/csrc/cuda/mem_eff_attention/kernels/backward_f16_aligned.cu b/xformers/components/attention/csrc/cuda/mem_eff_attention/kernels/backward_f16_aligned.cu
@@ -0,0 +1,6 @@
+// This file is auto-generated. See "generate_kernels.sh"
+#include "../kernel_backward.h"
+INSTANTIATE_ATTENTION_KERNEL_BACKWARD_SM50(cutlass::half_t, true);
+INSTANTIATE_ATTENTION_KERNEL_BACKWARD_SM70(cutlass::half_t, true);
+INSTANTIATE_ATTENTION_KERNEL_BACKWARD_SM75(cutlass::half_t, true);
+INSTANTIATE_ATTENTION_KERNEL_BACKWARD_SM80(cutlass::half_t, true);
diff --git a/xformers/components/attention/csrc/cuda/mem_eff_attention/kernels/backward_f16_sm50.cu b/xformers/components/attention/csrc/cuda/mem_eff_attention/kernels/backward_f16_sm50.cu
diff --git a/...ers/components/attention/csrc/cuda/mem_eff_attention/kernels/backward_f16_sm50_aligned.cu b/...ers/components/attention/csrc/cuda/mem_eff_attention/kernels/backward_f16_sm50_aligned.cu
diff --git a/xformers/components/attention/csrc/cuda/mem_eff_attention/kernels/backward_f16_sm70.cu b/xformers/components/attention/csrc/cuda/mem_eff_attention/kernels/backward_f16_sm70.cu
diff --git a/...ers/components/attention/csrc/cuda/mem_eff_attention/kernels/backward_f16_sm70_aligned.cu b/...ers/components/attention/csrc/cuda/mem_eff_attention/kernels/backward_f16_sm70_aligned.cu
diff --git a/xformers/components/attention/csrc/cuda/mem_eff_attention/kernels/backward_f16_sm75.cu b/xformers/components/attention/csrc/cuda/mem_eff_attention/kernels/backward_f16_sm75.cu
diff --git a/...ers/components/attention/csrc/cuda/mem_eff_attention/kernels/backward_f16_sm75_aligned.cu b/...ers/components/attention/csrc/cuda/mem_eff_attention/kernels/backward_f16_sm75_aligned.cu
diff --git a/xformers/components/attention/csrc/cuda/mem_eff_attention/kernels/backward_f16_sm80.cu b/xformers/components/attention/csrc/cuda/mem_eff_attention/kernels/backward_f16_sm80.cu
diff --git a/...ers/components/attention/csrc/cuda/mem_eff_attention/kernels/backward_f16_sm80_aligned.cu b/...ers/components/attention/csrc/cuda/mem_eff_attention/kernels/backward_f16_sm80_aligned.cu
diff --git a/xformers/components/attention/csrc/cuda/mem_eff_attention/kernels/backward_f32.cu b/xformers/components/attention/csrc/cuda/mem_eff_attention/kernels/backward_f32.cu
@@ -0,0 +1,6 @@
+// This file is auto-generated. See "generate_kernels.sh"
+#include "../kernel_backward.h"
+INSTANTIATE_ATTENTION_KERNEL_BACKWARD_SM50(float, false);
+INSTANTIATE_ATTENTION_KERNEL_BACKWARD_SM70(float, false);
+INSTANTIATE_ATTENTION_KERNEL_BACKWARD_SM75(float, false);
+INSTANTIATE_ATTENTION_KERNEL_BACKWARD_SM80(float, false);
diff --git a/xformers/components/attention/csrc/cuda/mem_eff_attention/kernels/backward_f32_aligned.cu b/xformers/components/attention/csrc/cuda/mem_eff_attention/kernels/backward_f32_aligned.cu
@@ -0,0 +1,6 @@
+// This file is auto-generated. See "generate_kernels.sh"
+#include "../kernel_backward.h"
+INSTANTIATE_ATTENTION_KERNEL_BACKWARD_SM50(float, true);
+INSTANTIATE_ATTENTION_KERNEL_BACKWARD_SM70(float, true);
+INSTANTIATE_ATTENTION_KERNEL_BACKWARD_SM75(float, true);
+INSTANTIATE_ATTENTION_KERNEL_BACKWARD_SM80(float, true);
diff --git a/xformers/components/attention/csrc/cuda/mem_eff_attention/kernels/backward_f32_sm50.cu b/xformers/components/attention/csrc/cuda/mem_eff_attention/kernels/backward_f32_sm50.cu
diff --git a/...ers/components/attention/csrc/cuda/mem_eff_attention/kernels/backward_f32_sm50_aligned.cu b/...ers/components/attention/csrc/cuda/mem_eff_attention/kernels/backward_f32_sm50_aligned.cu
diff --git a/xformers/components/attention/csrc/cuda/mem_eff_attention/kernels/backward_f32_sm70.cu b/xformers/components/attention/csrc/cuda/mem_eff_attention/kernels/backward_f32_sm70.cu
diff --git a/...ers/components/attention/csrc/cuda/mem_eff_attention/kernels/backward_f32_sm70_aligned.cu b/...ers/components/attention/csrc/cuda/mem_eff_attention/kernels/backward_f32_sm70_aligned.cu
diff --git a/xformers/components/attention/csrc/cuda/mem_eff_attention/kernels/backward_f32_sm75.cu b/xformers/components/attention/csrc/cuda/mem_eff_attention/kernels/backward_f32_sm75.cu
diff --git a/...ers/components/attention/csrc/cuda/mem_eff_attention/kernels/backward_f32_sm75_aligned.cu b/...ers/components/attention/csrc/cuda/mem_eff_attention/kernels/backward_f32_sm75_aligned.cu
diff --git a/xformers/components/attention/csrc/cuda/mem_eff_attention/kernels/backward_f32_sm80.cu b/xformers/components/attention/csrc/cuda/mem_eff_attention/kernels/backward_f32_sm80.cu
diff --git a/...ers/components/attention/csrc/cuda/mem_eff_attention/kernels/backward_f32_sm80_aligned.cu b/...ers/components/attention/csrc/cuda/mem_eff_attention/kernels/backward_f32_sm80_aligned.cu
diff --git a/xformers/components/attention/csrc/cuda/mem_eff_attention/kernels/generate_kernels.sh b/xformers/components/attention/csrc/cuda/mem_eff_attention/kernels/generate_kernels.sh
@@ -4,22 +4,22 @@ rm -f *.cu
 IFS=","
 for kernel in "BACKWARD"; do
     kernel_lower=`echo "\$kernel" | awk '{print tolower($0)}'`
-    for sm in 50 70 75 80; do
-        for aligned in "false" "true"; do
-            [[ $aligned = "true" ]] && aligned_suffix="_aligned" || aligned_suffix=""
-            for dtype_name in "f32" "f16"; do
-                case "$dtype_name" in
-                    "f32") dtype="float" ;;
-                    "f16") dtype="cutlass::half_t" ;;
-                esac
-                FNAME="${kernel_lower}_${dtype_name}_sm${sm}${aligned_suffix}.cu"
-                echo $FNAME
-                cat <<EOF > $FNAME
+    for aligned in "false" "true"; do
+        [[ $aligned = "true" ]] && aligned_suffix="_aligned" || aligned_suffix=""
+        for dtype_name in "f32" "f16"; do
+            case "$dtype_name" in
+                "f32") dtype="float" ;;
+                "f16") dtype="cutlass::half_t" ;;
+            esac
+            FNAME="${kernel_lower}_${dtype_name}${aligned_suffix}.cu"
+            echo $FNAME
+            cat <<EOF > $FNAME
 // This file is auto-generated. See "generate_kernels.sh"
 #include "../kernel_backward.h"
-INSTANTIATE_ATTENTION_KERNEL_${kernel}_SM${sm}($dtype, $aligned);
 EOF
+            for sm in 50 70 75 80; do
+                echo "INSTANTIATE_ATTENTION_KERNEL_${kernel}_SM${sm}($dtype, $aligned);" >> $FNAME
             done;
-        done
+        done;
     done
 done;