Move autogenerated files to 'autogen' folder

ghstack-source-id: e8fd3486722f3e201a67e3db5771eef79c4e4ae2 Pull Request resolved: https://github.com/fairinternal/xformers/pull/456 __original_commit__ = fairinternal/xformers@dfbacbc69253f36c3cf3a13df9b2eb66348447c3
P2Enjoy · Feb 2, 2023 · 48a77cc · 48a77cc
1 parent 82d5881
commit 48a77cc
Show file tree

Hide file tree

Showing 53 changed files with 110 additions and 156 deletions.
diff --git a/.github/workflows/linters_reusable.yml b/.github/workflows/linters_reusable.yml
@@ -46,4 +46,4 @@ jobs:
           clang-format --version
 
           # apply to our files - excluding autogenerated files
-          ./.circleci/run-clang-format.py -e "*fmha/kernels" -r xformers/csrc
+          ./.circleci/run-clang-format.py -e "*fmha/autogen" -r xformers/csrc
diff --git a/xformers/csrc/attention/cuda/fmha/attention_backward_generic.cu b/xformers/csrc/attention/cuda/fmha/attention_backward_generic.cu
@@ -10,9 +10,9 @@
 #include <torch/library.h>
 #include "ATen/ops/empty_like.h"
 
+#include "autogen/cutlassB.h"
 #include "gemm_kernel_utils.h"
 #include "kernel_backward.h"
-#include "kernels/cutlassB.h"
 #include "pytorch_utils.h"
 
 namespace {

diff --git a/xformers/csrc/attention/cuda/fmha/attention_forward_generic.cu b/xformers/csrc/attention/cuda/fmha/attention_forward_generic.cu
@@ -12,8 +12,8 @@
 #include <torch/library.h>
 #include <ATen/cuda/CUDAGraphsUtils.cuh>
 
+#include "autogen/cutlassF.h"
 #include "kernel_forward.h"
-#include "kernels/cutlassF.h"
 #include "pytorch_utils.h"
 
 namespace {

diff --git a/...rc/attention/cuda/fmha/kernels/cutlassB.h → ...rc/attention/cuda/fmha/autogen/cutlassB.h b/...rc/attention/cuda/fmha/kernels/cutlassB.h → ...rc/attention/cuda/fmha/autogen/cutlassB.h
@@ -1,8 +1,7 @@
 // This file is auto-generated. See "generate_kernels.py"
-#include "../kernel_backward.h"
-
 #pragma once
 #ifndef XFORMERS_MEM_EFF_ATTENTION_DISABLE_BACKWARD
+#include "../kernel_backward.h"
 // ======== f16 / sm50 ========
 __global__ void __launch_bounds__(
     AttentionBackwardKernel<cutlass::arch::Sm50, cutlass::half_t, true, false, false, 64, 64, 32>::kNumThreads,

diff --git a/...rc/attention/cuda/fmha/kernels/cutlassF.h → ...rc/attention/cuda/fmha/autogen/cutlassF.h b/...rc/attention/cuda/fmha/kernels/cutlassF.h → ...rc/attention/cuda/fmha/autogen/cutlassF.h
@@ -1,8 +1,7 @@
 // This file is auto-generated. See "generate_kernels.py"
-#include "../kernel_forward.h"
-
 #pragma once
 #ifndef XFORMERS_MEM_EFF_ATTENTION_DISABLE_FORWARD
+#include "../kernel_forward.h"
 // ======== bf16 / sm80 ========
 __global__ void __launch_bounds__(
     AttentionKernel<cutlass::bfloat16_t, cutlass::arch::Sm80, true, 64, 64, true, true, true>::kNumThreads,

diff --git a/...mha/kernels/cutlassB_bf16_aligned_k128.cu → ...utogen/impl/cutlassB_bf16_aligned_k128.cu b/...mha/kernels/cutlassB_bf16_aligned_k128.cu → ...utogen/impl/cutlassB_bf16_aligned_k128.cu
@@ -1,7 +1,6 @@
-#ifndef XFORMERS_MEM_EFF_ATTENTION_DISABLE_BACKWARD
 // This file is auto-generated. See "generate_kernels.py"
-#include "../kernel_backward.h"
-
+#ifndef XFORMERS_MEM_EFF_ATTENTION_DISABLE_BACKWARD
+#include "../../kernel_backward.h"
 __global__ void __launch_bounds__(
     AttentionBackwardKernel<cutlass::arch::Sm80, cutlass::bfloat16_t, true, false, true, 128, 128, 128>::kNumThreads,
     AttentionBackwardKernel<cutlass::arch::Sm80, cutlass::bfloat16_t, true, false, true, 128, 128, 128>::kMinBlocksPerSm)

diff --git a/...els/cutlassB_bf16_aligned_k128_dropout.cu → ...mpl/cutlassB_bf16_aligned_k128_dropout.cu b/...els/cutlassB_bf16_aligned_k128_dropout.cu → ...mpl/cutlassB_bf16_aligned_k128_dropout.cu
@@ -1,7 +1,6 @@
-#ifndef XFORMERS_MEM_EFF_ATTENTION_DISABLE_BACKWARD
 // This file is auto-generated. See "generate_kernels.py"
-#include "../kernel_backward.h"
-
+#ifndef XFORMERS_MEM_EFF_ATTENTION_DISABLE_BACKWARD
+#include "../../kernel_backward.h"
 __global__ void __launch_bounds__(
     AttentionBackwardKernel<cutlass::arch::Sm80, cutlass::bfloat16_t, true, true, true, 128, 128, 128>::kNumThreads,
     AttentionBackwardKernel<cutlass::arch::Sm80, cutlass::bfloat16_t, true, true, true, 128, 128, 128>::kMinBlocksPerSm)

diff --git a/...fmha/kernels/cutlassB_bf16_aligned_k32.cu → ...autogen/impl/cutlassB_bf16_aligned_k32.cu b/...fmha/kernels/cutlassB_bf16_aligned_k32.cu → ...autogen/impl/cutlassB_bf16_aligned_k32.cu
@@ -1,7 +1,6 @@
-#ifndef XFORMERS_MEM_EFF_ATTENTION_DISABLE_BACKWARD
 // This file is auto-generated. See "generate_kernels.py"
-#include "../kernel_backward.h"
-
+#ifndef XFORMERS_MEM_EFF_ATTENTION_DISABLE_BACKWARD
+#include "../../kernel_backward.h"
 __global__ void __launch_bounds__(
     AttentionBackwardKernel<cutlass::arch::Sm80, cutlass::bfloat16_t, true, false, true, 64, 64, 32>::kNumThreads,
     AttentionBackwardKernel<cutlass::arch::Sm80, cutlass::bfloat16_t, true, false, true, 64, 64, 32>::kMinBlocksPerSm)

diff --git a/...nels/cutlassB_bf16_aligned_k32_dropout.cu → ...impl/cutlassB_bf16_aligned_k32_dropout.cu b/...nels/cutlassB_bf16_aligned_k32_dropout.cu → ...impl/cutlassB_bf16_aligned_k32_dropout.cu
@@ -1,7 +1,6 @@
-#ifndef XFORMERS_MEM_EFF_ATTENTION_DISABLE_BACKWARD
 // This file is auto-generated. See "generate_kernels.py"
-#include "../kernel_backward.h"
-
+#ifndef XFORMERS_MEM_EFF_ATTENTION_DISABLE_BACKWARD
+#include "../../kernel_backward.h"
 __global__ void __launch_bounds__(
     AttentionBackwardKernel<cutlass::arch::Sm80, cutlass::bfloat16_t, true, true, true, 64, 64, 32>::kNumThreads,
     AttentionBackwardKernel<cutlass::arch::Sm80, cutlass::bfloat16_t, true, true, true, 64, 64, 32>::kMinBlocksPerSm)

diff --git a/...fmha/kernels/cutlassB_bf16_aligned_k64.cu → ...autogen/impl/cutlassB_bf16_aligned_k64.cu b/...fmha/kernels/cutlassB_bf16_aligned_k64.cu → ...autogen/impl/cutlassB_bf16_aligned_k64.cu
@@ -1,7 +1,6 @@
-#ifndef XFORMERS_MEM_EFF_ATTENTION_DISABLE_BACKWARD
 // This file is auto-generated. See "generate_kernels.py"
-#include "../kernel_backward.h"
-
+#ifndef XFORMERS_MEM_EFF_ATTENTION_DISABLE_BACKWARD
+#include "../../kernel_backward.h"
 __global__ void __launch_bounds__(
     AttentionBackwardKernel<cutlass::arch::Sm80, cutlass::bfloat16_t, true, false, true, 64, 64, 64>::kNumThreads,
     AttentionBackwardKernel<cutlass::arch::Sm80, cutlass::bfloat16_t, true, false, true, 64, 64, 64>::kMinBlocksPerSm)

diff --git a/...nels/cutlassB_bf16_aligned_k64_dropout.cu → ...impl/cutlassB_bf16_aligned_k64_dropout.cu b/...nels/cutlassB_bf16_aligned_k64_dropout.cu → ...impl/cutlassB_bf16_aligned_k64_dropout.cu
@@ -1,7 +1,6 @@
-#ifndef XFORMERS_MEM_EFF_ATTENTION_DISABLE_BACKWARD
 // This file is auto-generated. See "generate_kernels.py"
-#include "../kernel_backward.h"
-
+#ifndef XFORMERS_MEM_EFF_ATTENTION_DISABLE_BACKWARD
+#include "../../kernel_backward.h"
 __global__ void __launch_bounds__(
     AttentionBackwardKernel<cutlass::arch::Sm80, cutlass::bfloat16_t, true, true, true, 64, 64, 64>::kNumThreads,
     AttentionBackwardKernel<cutlass::arch::Sm80, cutlass::bfloat16_t, true, true, true, 64, 64, 64>::kMinBlocksPerSm)

diff --git a/...a/kernels/cutlassB_bf16_aligned_k65536.cu → ...ogen/impl/cutlassB_bf16_aligned_k65536.cu b/...a/kernels/cutlassB_bf16_aligned_k65536.cu → ...ogen/impl/cutlassB_bf16_aligned_k65536.cu
@@ -1,7 +1,6 @@
-#ifndef XFORMERS_MEM_EFF_ATTENTION_DISABLE_BACKWARD
 // This file is auto-generated. See "generate_kernels.py"
-#include "../kernel_backward.h"
-
+#ifndef XFORMERS_MEM_EFF_ATTENTION_DISABLE_BACKWARD
+#include "../../kernel_backward.h"
 __global__ void __launch_bounds__(
     AttentionBackwardKernel<cutlass::arch::Sm80, cutlass::bfloat16_t, true, false, false, 128, 64, 65536>::kNumThreads,
     AttentionBackwardKernel<cutlass::arch::Sm80, cutlass::bfloat16_t, true, false, false, 128, 64, 65536>::kMinBlocksPerSm)

diff --git a/...s/cutlassB_bf16_aligned_k65536_dropout.cu → ...l/cutlassB_bf16_aligned_k65536_dropout.cu b/...s/cutlassB_bf16_aligned_k65536_dropout.cu → ...l/cutlassB_bf16_aligned_k65536_dropout.cu
@@ -1,7 +1,6 @@
-#ifndef XFORMERS_MEM_EFF_ATTENTION_DISABLE_BACKWARD
 // This file is auto-generated. See "generate_kernels.py"
-#include "../kernel_backward.h"
-
+#ifndef XFORMERS_MEM_EFF_ATTENTION_DISABLE_BACKWARD
+#include "../../kernel_backward.h"
 __global__ void __launch_bounds__(
     AttentionBackwardKernel<cutlass::arch::Sm80, cutlass::bfloat16_t, true, true, false, 128, 64, 65536>::kNumThreads,
     AttentionBackwardKernel<cutlass::arch::Sm80, cutlass::bfloat16_t, true, true, false, 128, 64, 65536>::kMinBlocksPerSm)

diff --git a/...fmha/kernels/cutlassB_bf16_aligned_k96.cu → ...autogen/impl/cutlassB_bf16_aligned_k96.cu b/...fmha/kernels/cutlassB_bf16_aligned_k96.cu → ...autogen/impl/cutlassB_bf16_aligned_k96.cu
@@ -1,7 +1,6 @@
-#ifndef XFORMERS_MEM_EFF_ATTENTION_DISABLE_BACKWARD
 // This file is auto-generated. See "generate_kernels.py"
-#include "../kernel_backward.h"
-
+#ifndef XFORMERS_MEM_EFF_ATTENTION_DISABLE_BACKWARD
+#include "../../kernel_backward.h"
 __global__ void __launch_bounds__(
     AttentionBackwardKernel<cutlass::arch::Sm80, cutlass::bfloat16_t, true, false, true, 128, 64, 96>::kNumThreads,
     AttentionBackwardKernel<cutlass::arch::Sm80, cutlass::bfloat16_t, true, false, true, 128, 64, 96>::kMinBlocksPerSm)

diff --git a/...fmha/kernels/cutlassB_f16_aligned_k128.cu → ...autogen/impl/cutlassB_f16_aligned_k128.cu b/...fmha/kernels/cutlassB_f16_aligned_k128.cu → ...autogen/impl/cutlassB_f16_aligned_k128.cu
@@ -1,7 +1,6 @@
-#ifndef XFORMERS_MEM_EFF_ATTENTION_DISABLE_BACKWARD
 // This file is auto-generated. See "generate_kernels.py"
-#include "../kernel_backward.h"
-
+#ifndef XFORMERS_MEM_EFF_ATTENTION_DISABLE_BACKWARD
+#include "../../kernel_backward.h"
 __global__ void __launch_bounds__(
     AttentionBackwardKernel<cutlass::arch::Sm70, cutlass::half_t, true, false, false, 128, 64, 128>::kNumThreads,
     AttentionBackwardKernel<cutlass::arch::Sm70, cutlass::half_t, true, false, false, 128, 64, 128>::kMinBlocksPerSm)

diff --git a/...nels/cutlassB_f16_aligned_k128_dropout.cu → ...impl/cutlassB_f16_aligned_k128_dropout.cu b/...nels/cutlassB_f16_aligned_k128_dropout.cu → ...impl/cutlassB_f16_aligned_k128_dropout.cu
@@ -1,7 +1,6 @@
-#ifndef XFORMERS_MEM_EFF_ATTENTION_DISABLE_BACKWARD
 // This file is auto-generated. See "generate_kernels.py"
-#include "../kernel_backward.h"
-
+#ifndef XFORMERS_MEM_EFF_ATTENTION_DISABLE_BACKWARD
+#include "../../kernel_backward.h"
 __global__ void __launch_bounds__(
     AttentionBackwardKernel<cutlass::arch::Sm70, cutlass::half_t, true, true, false, 128, 64, 128>::kNumThreads,
     AttentionBackwardKernel<cutlass::arch::Sm70, cutlass::half_t, true, true, false, 128, 64, 128>::kMinBlocksPerSm)

diff --git a/.../fmha/kernels/cutlassB_f16_aligned_k32.cu → .../autogen/impl/cutlassB_f16_aligned_k32.cu b/.../fmha/kernels/cutlassB_f16_aligned_k32.cu → .../autogen/impl/cutlassB_f16_aligned_k32.cu
@@ -1,7 +1,6 @@
-#ifndef XFORMERS_MEM_EFF_ATTENTION_DISABLE_BACKWARD
 // This file is auto-generated. See "generate_kernels.py"
-#include "../kernel_backward.h"
-
+#ifndef XFORMERS_MEM_EFF_ATTENTION_DISABLE_BACKWARD
+#include "../../kernel_backward.h"
 __global__ void __launch_bounds__(
     AttentionBackwardKernel<cutlass::arch::Sm50, cutlass::half_t, true, false, false, 64, 64, 32>::kNumThreads,
     AttentionBackwardKernel<cutlass::arch::Sm50, cutlass::half_t, true, false, false, 64, 64, 32>::kMinBlocksPerSm)

diff --git a/...rnels/cutlassB_f16_aligned_k32_dropout.cu → .../impl/cutlassB_f16_aligned_k32_dropout.cu b/...rnels/cutlassB_f16_aligned_k32_dropout.cu → .../impl/cutlassB_f16_aligned_k32_dropout.cu
@@ -1,7 +1,6 @@
-#ifndef XFORMERS_MEM_EFF_ATTENTION_DISABLE_BACKWARD
 // This file is auto-generated. See "generate_kernels.py"
-#include "../kernel_backward.h"
-
+#ifndef XFORMERS_MEM_EFF_ATTENTION_DISABLE_BACKWARD
+#include "../../kernel_backward.h"
 __global__ void __launch_bounds__(
     AttentionBackwardKernel<cutlass::arch::Sm50, cutlass::half_t, true, true, false, 64, 64, 32>::kNumThreads,
     AttentionBackwardKernel<cutlass::arch::Sm50, cutlass::half_t, true, true, false, 64, 64, 32>::kMinBlocksPerSm)

diff --git a/.../fmha/kernels/cutlassB_f16_aligned_k64.cu → .../autogen/impl/cutlassB_f16_aligned_k64.cu b/.../fmha/kernels/cutlassB_f16_aligned_k64.cu → .../autogen/impl/cutlassB_f16_aligned_k64.cu
@@ -1,7 +1,6 @@
-#ifndef XFORMERS_MEM_EFF_ATTENTION_DISABLE_BACKWARD
 // This file is auto-generated. See "generate_kernels.py"
-#include "../kernel_backward.h"
-
+#ifndef XFORMERS_MEM_EFF_ATTENTION_DISABLE_BACKWARD
+#include "../../kernel_backward.h"
 __global__ void __launch_bounds__(
     AttentionBackwardKernel<cutlass::arch::Sm50, cutlass::half_t, true, false, false, 64, 64, 64>::kNumThreads,
     AttentionBackwardKernel<cutlass::arch::Sm50, cutlass::half_t, true, false, false, 64, 64, 64>::kMinBlocksPerSm)

diff --git a/...rnels/cutlassB_f16_aligned_k64_dropout.cu → .../impl/cutlassB_f16_aligned_k64_dropout.cu b/...rnels/cutlassB_f16_aligned_k64_dropout.cu → .../impl/cutlassB_f16_aligned_k64_dropout.cu
@@ -1,7 +1,6 @@
-#ifndef XFORMERS_MEM_EFF_ATTENTION_DISABLE_BACKWARD
 // This file is auto-generated. See "generate_kernels.py"
-#include "../kernel_backward.h"
-
+#ifndef XFORMERS_MEM_EFF_ATTENTION_DISABLE_BACKWARD
+#include "../../kernel_backward.h"
 __global__ void __launch_bounds__(
     AttentionBackwardKernel<cutlass::arch::Sm50, cutlass::half_t, true, true, false, 64, 64, 64>::kNumThreads,
     AttentionBackwardKernel<cutlass::arch::Sm50, cutlass::half_t, true, true, false, 64, 64, 64>::kMinBlocksPerSm)

diff --git a/...ha/kernels/cutlassB_f16_aligned_k65536.cu → ...togen/impl/cutlassB_f16_aligned_k65536.cu b/...ha/kernels/cutlassB_f16_aligned_k65536.cu → ...togen/impl/cutlassB_f16_aligned_k65536.cu
@@ -1,7 +1,6 @@
-#ifndef XFORMERS_MEM_EFF_ATTENTION_DISABLE_BACKWARD
 // This file is auto-generated. See "generate_kernels.py"
-#include "../kernel_backward.h"
-
+#ifndef XFORMERS_MEM_EFF_ATTENTION_DISABLE_BACKWARD
+#include "../../kernel_backward.h"
 __global__ void __launch_bounds__(
     AttentionBackwardKernel<cutlass::arch::Sm70, cutlass::half_t, true, false, false, 128, 64, 65536>::kNumThreads,
     AttentionBackwardKernel<cutlass::arch::Sm70, cutlass::half_t, true, false, false, 128, 64, 65536>::kMinBlocksPerSm)

diff --git a/...ls/cutlassB_f16_aligned_k65536_dropout.cu → ...pl/cutlassB_f16_aligned_k65536_dropout.cu b/...ls/cutlassB_f16_aligned_k65536_dropout.cu → ...pl/cutlassB_f16_aligned_k65536_dropout.cu
@@ -1,7 +1,6 @@
-#ifndef XFORMERS_MEM_EFF_ATTENTION_DISABLE_BACKWARD
 // This file is auto-generated. See "generate_kernels.py"
-#include "../kernel_backward.h"
-
+#ifndef XFORMERS_MEM_EFF_ATTENTION_DISABLE_BACKWARD
+#include "../../kernel_backward.h"
 __global__ void __launch_bounds__(
     AttentionBackwardKernel<cutlass::arch::Sm70, cutlass::half_t, true, true, false, 128, 64, 65536>::kNumThreads,
     AttentionBackwardKernel<cutlass::arch::Sm70, cutlass::half_t, true, true, false, 128, 64, 65536>::kMinBlocksPerSm)

diff --git a/.../fmha/kernels/cutlassB_f16_aligned_k96.cu → .../autogen/impl/cutlassB_f16_aligned_k96.cu b/.../fmha/kernels/cutlassB_f16_aligned_k96.cu → .../autogen/impl/cutlassB_f16_aligned_k96.cu
@@ -1,7 +1,6 @@
-#ifndef XFORMERS_MEM_EFF_ATTENTION_DISABLE_BACKWARD
 // This file is auto-generated. See "generate_kernels.py"
-#include "../kernel_backward.h"
-
+#ifndef XFORMERS_MEM_EFF_ATTENTION_DISABLE_BACKWARD
+#include "../../kernel_backward.h"
 __global__ void __launch_bounds__(
     AttentionBackwardKernel<cutlass::arch::Sm80, cutlass::half_t, true, false, true, 128, 64, 96>::kNumThreads,
     AttentionBackwardKernel<cutlass::arch::Sm80, cutlass::half_t, true, false, true, 128, 64, 96>::kMinBlocksPerSm)

diff --git a/...a/kernels/cutlassB_f16_notaligned_k128.cu → ...ogen/impl/cutlassB_f16_notaligned_k128.cu b/...a/kernels/cutlassB_f16_notaligned_k128.cu → ...ogen/impl/cutlassB_f16_notaligned_k128.cu
@@ -1,7 +1,6 @@
-#ifndef XFORMERS_MEM_EFF_ATTENTION_DISABLE_BACKWARD
 // This file is auto-generated. See "generate_kernels.py"
-#include "../kernel_backward.h"
-
+#ifndef XFORMERS_MEM_EFF_ATTENTION_DISABLE_BACKWARD
+#include "../../kernel_backward.h"
 __global__ void __launch_bounds__(
     AttentionBackwardKernel<cutlass::arch::Sm70, cutlass::half_t, false, false, false, 128, 64, 128>::kNumThreads,
     AttentionBackwardKernel<cutlass::arch::Sm70, cutlass::half_t, false, false, false, 128, 64, 128>::kMinBlocksPerSm)

diff --git a/...s/cutlassB_f16_notaligned_k128_dropout.cu → ...l/cutlassB_f16_notaligned_k128_dropout.cu b/...s/cutlassB_f16_notaligned_k128_dropout.cu → ...l/cutlassB_f16_notaligned_k128_dropout.cu
@@ -1,7 +1,6 @@
-#ifndef XFORMERS_MEM_EFF_ATTENTION_DISABLE_BACKWARD
 // This file is auto-generated. See "generate_kernels.py"
-#include "../kernel_backward.h"
-
+#ifndef XFORMERS_MEM_EFF_ATTENTION_DISABLE_BACKWARD
+#include "../../kernel_backward.h"
 __global__ void __launch_bounds__(
     AttentionBackwardKernel<cutlass::arch::Sm70, cutlass::half_t, false, true, false, 128, 64, 128>::kNumThreads,
     AttentionBackwardKernel<cutlass::arch::Sm70, cutlass::half_t, false, true, false, 128, 64, 128>::kMinBlocksPerSm)

diff --git a/...ha/kernels/cutlassB_f16_notaligned_k32.cu → ...togen/impl/cutlassB_f16_notaligned_k32.cu b/...ha/kernels/cutlassB_f16_notaligned_k32.cu → ...togen/impl/cutlassB_f16_notaligned_k32.cu
@@ -1,7 +1,6 @@
-#ifndef XFORMERS_MEM_EFF_ATTENTION_DISABLE_BACKWARD
 // This file is auto-generated. See "generate_kernels.py"
-#include "../kernel_backward.h"
-
+#ifndef XFORMERS_MEM_EFF_ATTENTION_DISABLE_BACKWARD
+#include "../../kernel_backward.h"
 __global__ void __launch_bounds__(
     AttentionBackwardKernel<cutlass::arch::Sm50, cutlass::half_t, false, false, false, 64, 64, 32>::kNumThreads,
     AttentionBackwardKernel<cutlass::arch::Sm50, cutlass::half_t, false, false, false, 64, 64, 32>::kMinBlocksPerSm)

diff --git a/...ls/cutlassB_f16_notaligned_k32_dropout.cu → ...pl/cutlassB_f16_notaligned_k32_dropout.cu b/...ls/cutlassB_f16_notaligned_k32_dropout.cu → ...pl/cutlassB_f16_notaligned_k32_dropout.cu
@@ -1,7 +1,6 @@
-#ifndef XFORMERS_MEM_EFF_ATTENTION_DISABLE_BACKWARD
 // This file is auto-generated. See "generate_kernels.py"
-#include "../kernel_backward.h"
-
+#ifndef XFORMERS_MEM_EFF_ATTENTION_DISABLE_BACKWARD
+#include "../../kernel_backward.h"
 __global__ void __launch_bounds__(
     AttentionBackwardKernel<cutlass::arch::Sm50, cutlass::half_t, false, true, false, 64, 64, 32>::kNumThreads,
     AttentionBackwardKernel<cutlass::arch::Sm50, cutlass::half_t, false, true, false, 64, 64, 32>::kMinBlocksPerSm)

diff --git a/...ha/kernels/cutlassB_f16_notaligned_k64.cu → ...togen/impl/cutlassB_f16_notaligned_k64.cu b/...ha/kernels/cutlassB_f16_notaligned_k64.cu → ...togen/impl/cutlassB_f16_notaligned_k64.cu
@@ -1,7 +1,6 @@
-#ifndef XFORMERS_MEM_EFF_ATTENTION_DISABLE_BACKWARD
 // This file is auto-generated. See "generate_kernels.py"
-#include "../kernel_backward.h"
-
+#ifndef XFORMERS_MEM_EFF_ATTENTION_DISABLE_BACKWARD
+#include "../../kernel_backward.h"
 __global__ void __launch_bounds__(
     AttentionBackwardKernel<cutlass::arch::Sm50, cutlass::half_t, false, false, false, 64, 64, 64>::kNumThreads,
     AttentionBackwardKernel<cutlass::arch::Sm50, cutlass::half_t, false, false, false, 64, 64, 64>::kMinBlocksPerSm)

diff --git a/...ls/cutlassB_f16_notaligned_k64_dropout.cu → ...pl/cutlassB_f16_notaligned_k64_dropout.cu b/...ls/cutlassB_f16_notaligned_k64_dropout.cu → ...pl/cutlassB_f16_notaligned_k64_dropout.cu
@@ -1,7 +1,6 @@
-#ifndef XFORMERS_MEM_EFF_ATTENTION_DISABLE_BACKWARD
 // This file is auto-generated. See "generate_kernels.py"
-#include "../kernel_backward.h"
-
+#ifndef XFORMERS_MEM_EFF_ATTENTION_DISABLE_BACKWARD
+#include "../../kernel_backward.h"
 __global__ void __launch_bounds__(
     AttentionBackwardKernel<cutlass::arch::Sm50, cutlass::half_t, false, true, false, 64, 64, 64>::kNumThreads,
     AttentionBackwardKernel<cutlass::arch::Sm50, cutlass::half_t, false, true, false, 64, 64, 64>::kMinBlocksPerSm)

diff --git a/...kernels/cutlassB_f16_notaligned_k65536.cu → ...en/impl/cutlassB_f16_notaligned_k65536.cu b/...kernels/cutlassB_f16_notaligned_k65536.cu → ...en/impl/cutlassB_f16_notaligned_k65536.cu
@@ -1,7 +1,6 @@
-#ifndef XFORMERS_MEM_EFF_ATTENTION_DISABLE_BACKWARD
 // This file is auto-generated. See "generate_kernels.py"
-#include "../kernel_backward.h"
-
+#ifndef XFORMERS_MEM_EFF_ATTENTION_DISABLE_BACKWARD
+#include "../../kernel_backward.h"
 __global__ void __launch_bounds__(
     AttentionBackwardKernel<cutlass::arch::Sm70, cutlass::half_t, false, false, false, 128, 64, 65536>::kNumThreads,
     AttentionBackwardKernel<cutlass::arch::Sm70, cutlass::half_t, false, false, false, 128, 64, 65536>::kMinBlocksPerSm)

diff --git a/...cutlassB_f16_notaligned_k65536_dropout.cu → ...cutlassB_f16_notaligned_k65536_dropout.cu b/...cutlassB_f16_notaligned_k65536_dropout.cu → ...cutlassB_f16_notaligned_k65536_dropout.cu
@@ -1,7 +1,6 @@
-#ifndef XFORMERS_MEM_EFF_ATTENTION_DISABLE_BACKWARD
 // This file is auto-generated. See "generate_kernels.py"
-#include "../kernel_backward.h"
-
+#ifndef XFORMERS_MEM_EFF_ATTENTION_DISABLE_BACKWARD
+#include "../../kernel_backward.h"
 __global__ void __launch_bounds__(
     AttentionBackwardKernel<cutlass::arch::Sm70, cutlass::half_t, false, true, false, 128, 64, 65536>::kNumThreads,
     AttentionBackwardKernel<cutlass::arch::Sm70, cutlass::half_t, false, true, false, 128, 64, 65536>::kMinBlocksPerSm)

diff --git a/...fmha/kernels/cutlassB_f32_aligned_k128.cu → ...autogen/impl/cutlassB_f32_aligned_k128.cu b/...fmha/kernels/cutlassB_f32_aligned_k128.cu → ...autogen/impl/cutlassB_f32_aligned_k128.cu
@@ -1,7 +1,6 @@
-#ifndef XFORMERS_MEM_EFF_ATTENTION_DISABLE_BACKWARD
 // This file is auto-generated. See "generate_kernels.py"
-#include "../kernel_backward.h"
-
+#ifndef XFORMERS_MEM_EFF_ATTENTION_DISABLE_BACKWARD
+#include "../../kernel_backward.h"
 __global__ void __launch_bounds__(
     AttentionBackwardKernel<cutlass::arch::Sm80, float, true, false, false, 128, 64, 128>::kNumThreads,
     AttentionBackwardKernel<cutlass::arch::Sm80, float, true, false, false, 128, 64, 128>::kMinBlocksPerSm)

diff --git a/...nels/cutlassB_f32_aligned_k128_dropout.cu → ...impl/cutlassB_f32_aligned_k128_dropout.cu b/...nels/cutlassB_f32_aligned_k128_dropout.cu → ...impl/cutlassB_f32_aligned_k128_dropout.cu
@@ -1,7 +1,6 @@
-#ifndef XFORMERS_MEM_EFF_ATTENTION_DISABLE_BACKWARD
 // This file is auto-generated. See "generate_kernels.py"
-#include "../kernel_backward.h"
-
+#ifndef XFORMERS_MEM_EFF_ATTENTION_DISABLE_BACKWARD
+#include "../../kernel_backward.h"
 __global__ void __launch_bounds__(
     AttentionBackwardKernel<cutlass::arch::Sm80, float, true, true, false, 128, 64, 128>::kNumThreads,
     AttentionBackwardKernel<cutlass::arch::Sm80, float, true, true, false, 128, 64, 128>::kMinBlocksPerSm)

diff --git a/.../fmha/kernels/cutlassB_f32_aligned_k32.cu → .../autogen/impl/cutlassB_f32_aligned_k32.cu b/.../fmha/kernels/cutlassB_f32_aligned_k32.cu → .../autogen/impl/cutlassB_f32_aligned_k32.cu
@@ -1,7 +1,6 @@
-#ifndef XFORMERS_MEM_EFF_ATTENTION_DISABLE_BACKWARD
 // This file is auto-generated. See "generate_kernels.py"
-#include "../kernel_backward.h"
-
+#ifndef XFORMERS_MEM_EFF_ATTENTION_DISABLE_BACKWARD
+#include "../../kernel_backward.h"
 __global__ void __launch_bounds__(
     AttentionBackwardKernel<cutlass::arch::Sm50, float, true, false, false, 64, 64, 32>::kNumThreads,
     AttentionBackwardKernel<cutlass::arch::Sm50, float, true, false, false, 64, 64, 32>::kMinBlocksPerSm)

diff --git a/...rnels/cutlassB_f32_aligned_k32_dropout.cu → .../impl/cutlassB_f32_aligned_k32_dropout.cu b/...rnels/cutlassB_f32_aligned_k32_dropout.cu → .../impl/cutlassB_f32_aligned_k32_dropout.cu
@@ -1,7 +1,6 @@
-#ifndef XFORMERS_MEM_EFF_ATTENTION_DISABLE_BACKWARD
 // This file is auto-generated. See "generate_kernels.py"
-#include "../kernel_backward.h"
-
+#ifndef XFORMERS_MEM_EFF_ATTENTION_DISABLE_BACKWARD
+#include "../../kernel_backward.h"
 __global__ void __launch_bounds__(
     AttentionBackwardKernel<cutlass::arch::Sm50, float, true, true, false, 64, 64, 32>::kNumThreads,
     AttentionBackwardKernel<cutlass::arch::Sm50, float, true, true, false, 64, 64, 32>::kMinBlocksPerSm)

diff --git a/.../fmha/kernels/cutlassB_f32_aligned_k64.cu → .../autogen/impl/cutlassB_f32_aligned_k64.cu b/.../fmha/kernels/cutlassB_f32_aligned_k64.cu → .../autogen/impl/cutlassB_f32_aligned_k64.cu
@@ -1,7 +1,6 @@
-#ifndef XFORMERS_MEM_EFF_ATTENTION_DISABLE_BACKWARD
 // This file is auto-generated. See "generate_kernels.py"
-#include "../kernel_backward.h"
-
+#ifndef XFORMERS_MEM_EFF_ATTENTION_DISABLE_BACKWARD
+#include "../../kernel_backward.h"
 __global__ void __launch_bounds__(
     AttentionBackwardKernel<cutlass::arch::Sm50, float, true, false, false, 64, 64, 64>::kNumThreads,
     AttentionBackwardKernel<cutlass::arch::Sm50, float, true, false, false, 64, 64, 64>::kMinBlocksPerSm)

diff --git a/...rnels/cutlassB_f32_aligned_k64_dropout.cu → .../impl/cutlassB_f32_aligned_k64_dropout.cu b/...rnels/cutlassB_f32_aligned_k64_dropout.cu → .../impl/cutlassB_f32_aligned_k64_dropout.cu
@@ -1,7 +1,6 @@
-#ifndef XFORMERS_MEM_EFF_ATTENTION_DISABLE_BACKWARD
 // This file is auto-generated. See "generate_kernels.py"
-#include "../kernel_backward.h"
-
+#ifndef XFORMERS_MEM_EFF_ATTENTION_DISABLE_BACKWARD
+#include "../../kernel_backward.h"
 __global__ void __launch_bounds__(
     AttentionBackwardKernel<cutlass::arch::Sm50, float, true, true, false, 64, 64, 64>::kNumThreads,
     AttentionBackwardKernel<cutlass::arch::Sm50, float, true, true, false, 64, 64, 64>::kMinBlocksPerSm)

diff --git a/...ha/kernels/cutlassB_f32_aligned_k65536.cu → ...togen/impl/cutlassB_f32_aligned_k65536.cu b/...ha/kernels/cutlassB_f32_aligned_k65536.cu → ...togen/impl/cutlassB_f32_aligned_k65536.cu
@@ -1,7 +1,6 @@
-#ifndef XFORMERS_MEM_EFF_ATTENTION_DISABLE_BACKWARD
 // This file is auto-generated. See "generate_kernels.py"
-#include "../kernel_backward.h"
-
+#ifndef XFORMERS_MEM_EFF_ATTENTION_DISABLE_BACKWARD
+#include "../../kernel_backward.h"
 __global__ void __launch_bounds__(
     AttentionBackwardKernel<cutlass::arch::Sm80, float, true, false, false, 128, 64, 65536>::kNumThreads,
     AttentionBackwardKernel<cutlass::arch::Sm80, float, true, false, false, 128, 64, 65536>::kMinBlocksPerSm)

diff --git a/...ls/cutlassB_f32_aligned_k65536_dropout.cu → ...pl/cutlassB_f32_aligned_k65536_dropout.cu b/...ls/cutlassB_f32_aligned_k65536_dropout.cu → ...pl/cutlassB_f32_aligned_k65536_dropout.cu
@@ -1,7 +1,6 @@
-#ifndef XFORMERS_MEM_EFF_ATTENTION_DISABLE_BACKWARD
 // This file is auto-generated. See "generate_kernels.py"
-#include "../kernel_backward.h"
-
+#ifndef XFORMERS_MEM_EFF_ATTENTION_DISABLE_BACKWARD
+#include "../../kernel_backward.h"
 __global__ void __launch_bounds__(
     AttentionBackwardKernel<cutlass::arch::Sm80, float, true, true, false, 128, 64, 65536>::kNumThreads,
     AttentionBackwardKernel<cutlass::arch::Sm80, float, true, true, false, 128, 64, 65536>::kMinBlocksPerSm)

diff --git a/...a/kernels/cutlassB_f32_notaligned_k128.cu → ...ogen/impl/cutlassB_f32_notaligned_k128.cu b/...a/kernels/cutlassB_f32_notaligned_k128.cu → ...ogen/impl/cutlassB_f32_notaligned_k128.cu
@@ -1,7 +1,6 @@
-#ifndef XFORMERS_MEM_EFF_ATTENTION_DISABLE_BACKWARD
 // This file is auto-generated. See "generate_kernels.py"
-#include "../kernel_backward.h"
-
+#ifndef XFORMERS_MEM_EFF_ATTENTION_DISABLE_BACKWARD
+#include "../../kernel_backward.h"
 __global__ void __launch_bounds__(
     AttentionBackwardKernel<cutlass::arch::Sm50, float, false, false, false, 64, 64, 128>::kNumThreads,
     AttentionBackwardKernel<cutlass::arch::Sm50, float, false, false, false, 64, 64, 128>::kMinBlocksPerSm)