From 4772f2c7a003f2fb3783a21072a37acca3d9379e Mon Sep 17 00:00:00 2001
From: Vyacheslav Zhdanovskiy <zeronsix@gmail.com>
Date: Thu, 20 Feb 2025 12:52:48 +0300
Subject: [PATCH] Fix debug Turing switch

---
 inference_lib/src/aqlm/inference_kernels/cuda_kernel.cu | 4 ++--
 1 file changed, 2 insertions(+), 2 deletions(-)

diff --git a/inference_lib/src/aqlm/inference_kernels/cuda_kernel.cu b/inference_lib/src/aqlm/inference_kernels/cuda_kernel.cu
index 7e021ee..4e7173f 100644
--- a/inference_lib/src/aqlm/inference_kernels/cuda_kernel.cu
+++ b/inference_lib/src/aqlm/inference_kernels/cuda_kernel.cu
@@ -589,7 +589,7 @@ void  code2x8_matvec_cuda(
   int threads = 32 * thread_m;
   cudaStream_t stream = at::cuda::getCurrentCUDAStream().stream();
   const bool is_turing = cc_major == 7 && cc_minor == 5;
-  if (is_turing) {
+  if (!is_turing) {
     int shared = 16 * (2 * 256 * 8 + 32 * 9);
     cudaFuncSetAttribute(
       Code2x8MatVec<use_bfloat16>, cudaFuncAttributeMaxDynamicSharedMemorySize, shared
@@ -655,7 +655,7 @@ void  code2x8_dequant_cuda(
   int threads = 32 * thread_m;
   cudaStream_t stream = at::cuda::getCurrentCUDAStream().stream();
   const bool is_turing = cc_major == 7 && cc_minor == 5;
-  if (is_turing) {
+  if (!is_turing) {
     int shared = 16 * (2 * 256 * 8 + 32 * 9);
     if (use_bfloat16) {
       cudaFuncSetAttribute(