From 6018600ec16246e074b672fd4e25718562b16245 Mon Sep 17 00:00:00 2001 From: danthe3rd Date: Mon, 5 Sep 2022 14:58:44 +0000 Subject: [PATCH] Group all archs in the same file --- .../mem_eff_attention/kernels/backward_f16.cu | 6 +++++ .../kernels/backward_f16_aligned.cu | 6 +++++ .../kernels/backward_f16_sm50.cu | 3 --- .../kernels/backward_f16_sm50_aligned.cu | 3 --- .../kernels/backward_f16_sm70.cu | 3 --- .../kernels/backward_f16_sm70_aligned.cu | 3 --- .../kernels/backward_f16_sm75.cu | 3 --- .../kernels/backward_f16_sm75_aligned.cu | 3 --- .../kernels/backward_f16_sm80.cu | 3 --- .../kernels/backward_f16_sm80_aligned.cu | 3 --- .../mem_eff_attention/kernels/backward_f32.cu | 6 +++++ .../kernels/backward_f32_aligned.cu | 6 +++++ .../kernels/backward_f32_sm50.cu | 3 --- .../kernels/backward_f32_sm50_aligned.cu | 3 --- .../kernels/backward_f32_sm70.cu | 3 --- .../kernels/backward_f32_sm70_aligned.cu | 3 --- .../kernels/backward_f32_sm75.cu | 3 --- .../kernels/backward_f32_sm75_aligned.cu | 3 --- .../kernels/backward_f32_sm80.cu | 3 --- .../kernels/backward_f32_sm80_aligned.cu | 3 --- .../kernels/generate_kernels.sh | 26 +++++++++---------- 21 files changed, 37 insertions(+), 61 deletions(-) create mode 100644 xformers/components/attention/csrc/cuda/mem_eff_attention/kernels/backward_f16.cu create mode 100644 xformers/components/attention/csrc/cuda/mem_eff_attention/kernels/backward_f16_aligned.cu delete mode 100644 xformers/components/attention/csrc/cuda/mem_eff_attention/kernels/backward_f16_sm50.cu delete mode 100644 xformers/components/attention/csrc/cuda/mem_eff_attention/kernels/backward_f16_sm50_aligned.cu delete mode 100644 xformers/components/attention/csrc/cuda/mem_eff_attention/kernels/backward_f16_sm70.cu delete mode 100644 xformers/components/attention/csrc/cuda/mem_eff_attention/kernels/backward_f16_sm70_aligned.cu delete mode 100644 xformers/components/attention/csrc/cuda/mem_eff_attention/kernels/backward_f16_sm75.cu delete mode 100644 xformers/components/attention/csrc/cuda/mem_eff_attention/kernels/backward_f16_sm75_aligned.cu delete mode 100644 xformers/components/attention/csrc/cuda/mem_eff_attention/kernels/backward_f16_sm80.cu delete mode 100644 xformers/components/attention/csrc/cuda/mem_eff_attention/kernels/backward_f16_sm80_aligned.cu create mode 100644 xformers/components/attention/csrc/cuda/mem_eff_attention/kernels/backward_f32.cu create mode 100644 xformers/components/attention/csrc/cuda/mem_eff_attention/kernels/backward_f32_aligned.cu delete mode 100644 xformers/components/attention/csrc/cuda/mem_eff_attention/kernels/backward_f32_sm50.cu delete mode 100644 xformers/components/attention/csrc/cuda/mem_eff_attention/kernels/backward_f32_sm50_aligned.cu delete mode 100644 xformers/components/attention/csrc/cuda/mem_eff_attention/kernels/backward_f32_sm70.cu delete mode 100644 xformers/components/attention/csrc/cuda/mem_eff_attention/kernels/backward_f32_sm70_aligned.cu delete mode 100644 xformers/components/attention/csrc/cuda/mem_eff_attention/kernels/backward_f32_sm75.cu delete mode 100644 xformers/components/attention/csrc/cuda/mem_eff_attention/kernels/backward_f32_sm75_aligned.cu delete mode 100644 xformers/components/attention/csrc/cuda/mem_eff_attention/kernels/backward_f32_sm80.cu delete mode 100644 xformers/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/backward_f16.cu b/xformers/components/attention/csrc/cuda/mem_eff_attention/kernels/backward_f16.cu new file mode 100644 index 0000000000..a04bdbd6e0 --- /dev/null +++ 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 new file mode 100644 index 0000000000..d7aebba93d --- /dev/null +++ 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 deleted file mode 100644 index d9d7cccdf7..0000000000 --- a/xformers/components/attention/csrc/cuda/mem_eff_attention/kernels/backward_f16_sm50.cu +++ /dev/null @@ -1,3 +0,0 @@ -// This file is auto-generated. See "generate_kernels.sh" -#include "../kernel_backward.h" -INSTANTIATE_ATTENTION_KERNEL_BACKWARD_SM50(cutlass::half_t, false); diff --git a/xformers/components/attention/csrc/cuda/mem_eff_attention/kernels/backward_f16_sm50_aligned.cu b/xformers/components/attention/csrc/cuda/mem_eff_attention/kernels/backward_f16_sm50_aligned.cu deleted file mode 100644 index 89aaba232a..0000000000 --- a/xformers/components/attention/csrc/cuda/mem_eff_attention/kernels/backward_f16_sm50_aligned.cu +++ /dev/null @@ -1,3 +0,0 @@ -// This file is auto-generated. See "generate_kernels.sh" -#include "../kernel_backward.h" -INSTANTIATE_ATTENTION_KERNEL_BACKWARD_SM50(cutlass::half_t, true); 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 deleted file mode 100644 index 4ae107effe..0000000000 --- a/xformers/components/attention/csrc/cuda/mem_eff_attention/kernels/backward_f16_sm70.cu +++ /dev/null @@ -1,3 +0,0 @@ -// This file is auto-generated. See "generate_kernels.sh" -#include "../kernel_backward.h" -INSTANTIATE_ATTENTION_KERNEL_BACKWARD_SM70(cutlass::half_t, false); diff --git a/xformers/components/attention/csrc/cuda/mem_eff_attention/kernels/backward_f16_sm70_aligned.cu b/xformers/components/attention/csrc/cuda/mem_eff_attention/kernels/backward_f16_sm70_aligned.cu deleted file mode 100644 index 588fd994e9..0000000000 --- a/xformers/components/attention/csrc/cuda/mem_eff_attention/kernels/backward_f16_sm70_aligned.cu +++ /dev/null @@ -1,3 +0,0 @@ -// This file is auto-generated. See "generate_kernels.sh" -#include "../kernel_backward.h" -INSTANTIATE_ATTENTION_KERNEL_BACKWARD_SM70(cutlass::half_t, true); 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 deleted file mode 100644 index cbac52f3c7..0000000000 --- a/xformers/components/attention/csrc/cuda/mem_eff_attention/kernels/backward_f16_sm75.cu +++ /dev/null @@ -1,3 +0,0 @@ -// This file is auto-generated. See "generate_kernels.sh" -#include "../kernel_backward.h" -INSTANTIATE_ATTENTION_KERNEL_BACKWARD_SM75(cutlass::half_t, false); diff --git a/xformers/components/attention/csrc/cuda/mem_eff_attention/kernels/backward_f16_sm75_aligned.cu b/xformers/components/attention/csrc/cuda/mem_eff_attention/kernels/backward_f16_sm75_aligned.cu deleted file mode 100644 index 760e983bee..0000000000 --- a/xformers/components/attention/csrc/cuda/mem_eff_attention/kernels/backward_f16_sm75_aligned.cu +++ /dev/null @@ -1,3 +0,0 @@ -// This file is auto-generated. See "generate_kernels.sh" -#include "../kernel_backward.h" -INSTANTIATE_ATTENTION_KERNEL_BACKWARD_SM75(cutlass::half_t, true); 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 deleted file mode 100644 index 9f2ea62c15..0000000000 --- a/xformers/components/attention/csrc/cuda/mem_eff_attention/kernels/backward_f16_sm80.cu +++ /dev/null @@ -1,3 +0,0 @@ -// This file is auto-generated. See "generate_kernels.sh" -#include "../kernel_backward.h" -INSTANTIATE_ATTENTION_KERNEL_BACKWARD_SM80(cutlass::half_t, false); diff --git a/xformers/components/attention/csrc/cuda/mem_eff_attention/kernels/backward_f16_sm80_aligned.cu b/xformers/components/attention/csrc/cuda/mem_eff_attention/kernels/backward_f16_sm80_aligned.cu deleted file mode 100644 index 95c5d06742..0000000000 --- a/xformers/components/attention/csrc/cuda/mem_eff_attention/kernels/backward_f16_sm80_aligned.cu +++ /dev/null @@ -1,3 +0,0 @@ -// This file is auto-generated. See "generate_kernels.sh" -#include "../kernel_backward.h" -INSTANTIATE_ATTENTION_KERNEL_BACKWARD_SM80(cutlass::half_t, true); 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 new file mode 100644 index 0000000000..d14d2434cf --- /dev/null +++ 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 new file mode 100644 index 0000000000..dc4a824ff2 --- /dev/null +++ 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 deleted file mode 100644 index 93e4fd5891..0000000000 --- a/xformers/components/attention/csrc/cuda/mem_eff_attention/kernels/backward_f32_sm50.cu +++ /dev/null @@ -1,3 +0,0 @@ -// This file is auto-generated. See "generate_kernels.sh" -#include "../kernel_backward.h" -INSTANTIATE_ATTENTION_KERNEL_BACKWARD_SM50(float, false); diff --git a/xformers/components/attention/csrc/cuda/mem_eff_attention/kernels/backward_f32_sm50_aligned.cu b/xformers/components/attention/csrc/cuda/mem_eff_attention/kernels/backward_f32_sm50_aligned.cu deleted file mode 100644 index 372d7f1e49..0000000000 --- a/xformers/components/attention/csrc/cuda/mem_eff_attention/kernels/backward_f32_sm50_aligned.cu +++ /dev/null @@ -1,3 +0,0 @@ -// This file is auto-generated. See "generate_kernels.sh" -#include "../kernel_backward.h" -INSTANTIATE_ATTENTION_KERNEL_BACKWARD_SM50(float, true); 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 deleted file mode 100644 index c260439c43..0000000000 --- a/xformers/components/attention/csrc/cuda/mem_eff_attention/kernels/backward_f32_sm70.cu +++ /dev/null @@ -1,3 +0,0 @@ -// This file is auto-generated. See "generate_kernels.sh" -#include "../kernel_backward.h" -INSTANTIATE_ATTENTION_KERNEL_BACKWARD_SM70(float, false); diff --git a/xformers/components/attention/csrc/cuda/mem_eff_attention/kernels/backward_f32_sm70_aligned.cu b/xformers/components/attention/csrc/cuda/mem_eff_attention/kernels/backward_f32_sm70_aligned.cu deleted file mode 100644 index dc85033aa6..0000000000 --- a/xformers/components/attention/csrc/cuda/mem_eff_attention/kernels/backward_f32_sm70_aligned.cu +++ /dev/null @@ -1,3 +0,0 @@ -// This file is auto-generated. See "generate_kernels.sh" -#include "../kernel_backward.h" -INSTANTIATE_ATTENTION_KERNEL_BACKWARD_SM70(float, true); 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 deleted file mode 100644 index f2bfb046b6..0000000000 --- a/xformers/components/attention/csrc/cuda/mem_eff_attention/kernels/backward_f32_sm75.cu +++ /dev/null @@ -1,3 +0,0 @@ -// This file is auto-generated. See "generate_kernels.sh" -#include "../kernel_backward.h" -INSTANTIATE_ATTENTION_KERNEL_BACKWARD_SM75(float, false); diff --git a/xformers/components/attention/csrc/cuda/mem_eff_attention/kernels/backward_f32_sm75_aligned.cu b/xformers/components/attention/csrc/cuda/mem_eff_attention/kernels/backward_f32_sm75_aligned.cu deleted file mode 100644 index 85fb6ab493..0000000000 --- a/xformers/components/attention/csrc/cuda/mem_eff_attention/kernels/backward_f32_sm75_aligned.cu +++ /dev/null @@ -1,3 +0,0 @@ -// This file is auto-generated. See "generate_kernels.sh" -#include "../kernel_backward.h" -INSTANTIATE_ATTENTION_KERNEL_BACKWARD_SM75(float, true); 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 deleted file mode 100644 index 16ad339d9b..0000000000 --- a/xformers/components/attention/csrc/cuda/mem_eff_attention/kernels/backward_f32_sm80.cu +++ /dev/null @@ -1,3 +0,0 @@ -// This file is auto-generated. See "generate_kernels.sh" -#include "../kernel_backward.h" -INSTANTIATE_ATTENTION_KERNEL_BACKWARD_SM80(float, false); diff --git a/xformers/components/attention/csrc/cuda/mem_eff_attention/kernels/backward_f32_sm80_aligned.cu b/xformers/components/attention/csrc/cuda/mem_eff_attention/kernels/backward_f32_sm80_aligned.cu deleted file mode 100644 index 72db5e3c1b..0000000000 --- a/xformers/components/attention/csrc/cuda/mem_eff_attention/kernels/backward_f32_sm80_aligned.cu +++ /dev/null @@ -1,3 +0,0 @@ -// This file is auto-generated. See "generate_kernels.sh" -#include "../kernel_backward.h" -INSTANTIATE_ATTENTION_KERNEL_BACKWARD_SM80(float, true); 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 index 8c2f658f5b..b5bf6a0114 100755 --- 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 < $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 < $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;