SwiGLU: more info when kernel cant run

ghstack-source-id: 3c4ef8dd55c19fca9117ba5d43ee1b484e42d02a Pull Request resolved: https://github.com/fairinternal/xformers/pull/457 __original_commit__ = fairinternal/xformers@5529f67e955cb0801c02f656bf87128013d36b3d
P2Enjoy · Feb 9, 2023 · 36e23c5 · 36e23c5
1 parent 44bc216
commit 36e23c5
Show file tree

Hide file tree

Showing 2 changed files with 6 additions and 2 deletions.
diff --git a/xformers/csrc/swiglu/cuda/dual_gemm_silu_identity_mul.cu b/xformers/csrc/swiglu/cuda/dual_gemm_silu_identity_mul.cu
@@ -144,7 +144,9 @@ std::tuple<at::Tensor, at::Tensor, at::Tensor> dual_gemm_silu_identity_mul_(
       x.options().dtype(at::ScalarType::Byte));
   cutlass::Status status = dual_gemm.can_implement(arguments);
   TORCH_CHECK(
-      status == cutlass::Status::kSuccess, "not supported by this kernel");
+      status == cutlass::Status::kSuccess,
+      "`dual_gemm_silu_identity_mul` does not support this input: ",
+      cutlass::cutlassGetStatusString(status));
   status = dual_gemm.initialize(arguments, (uint8_t*)workspace.data_ptr());
   TORCH_CHECK(status == cutlass::Status::kSuccess, "kernel initialize failed");
   status = dual_gemm(stream);

diff --git a/xformers/csrc/swiglu/cuda/gemm_fused_operand_sum.cu b/xformers/csrc/swiglu/cuda/gemm_fused_operand_sum.cu
@@ -186,7 +186,9 @@ void gemm_fused_operand_sum_(
       a.options().dtype(at::ScalarType::Byte));
   cutlass::Status status = gemm_op.can_implement(arguments);
   TORCH_CHECK(
-      status == cutlass::Status::kSuccess, "not supported by this kernel");
+      status == cutlass::Status::kSuccess,
+      "`gemm_fused_operand_sum` does not support this input: ",
+      cutlass::cutlassGetStatusString(status));
   status = gemm_op.initialize(arguments, (uint8_t*)workspace.data_ptr());
   TORCH_CHECK(status == cutlass::Status::kSuccess, "kernel initialize failed");
   status = gemm_op(stream);