From 751eeb6fb02598af6c323f1635e0d07b0ee5d65a Mon Sep 17 00:00:00 2001
From: Li Zhang <lzhang329@gmail.com>
Date: Thu, 30 Nov 2023 04:23:48 +0000
Subject: [PATCH] fix msvc build

---
 builder/windows/README.md                   |  2 +-
 builder/windows/generate.ps1                |  2 +-
 src/turbomind/models/llama/llama_kernels.cu | 17 ++++++++---------
 3 files changed, 10 insertions(+), 11 deletions(-)

diff --git a/builder/windows/README.md b/builder/windows/README.md
index e22efda685..6933506782 100644
--- a/builder/windows/README.md
+++ b/builder/windows/README.md
@@ -3,7 +3,7 @@
 ## Requirements
 
 - [CMake 3.17+](https://github.com/Kitware/CMake/releases)
-- [Visual Studio 2022+](https://visualstudio.microsoft.com/downloads/)
+- [Visual Studio 2019+](https://visualstudio.microsoft.com/downloads/)
 - [CUDA Toolkit 11.8+](https://developer.nvidia.com/cuda-toolkit-archive)
 
 ## Build lmdeploy wheel
diff --git a/builder/windows/generate.ps1 b/builder/windows/generate.ps1
index a0b7faf83c..7c793dc382 100644
--- a/builder/windows/generate.ps1
+++ b/builder/windows/generate.ps1
@@ -1,4 +1,4 @@
-cmake .. -A x64 -T v143,cuda="$env:CUDA_PATH" `
+cmake .. -A x64 -T v142,cuda="$env:CUDA_PATH" `
     -DCMAKE_BUILD_TYPE=Release `
     -DCMAKE_INSTALL_PREFIX=install `
     -DBUILD_PY_FFI=ON `
diff --git a/src/turbomind/models/llama/llama_kernels.cu b/src/turbomind/models/llama/llama_kernels.cu
index 7fc328ac4b..af8bf76157 100644
--- a/src/turbomind/models/llama/llama_kernels.cu
+++ b/src/turbomind/models/llama/llama_kernels.cu
@@ -821,15 +821,14 @@ void invokeBatchedCopy(void** src_ptr, void** dst_ptr, int* size, int count, cud
                     params.size[i] = size[c + i] / sizeof(T);
                 }
                 const int max_size = *std::max_element(params.size.begin(), params.size.end());
-                dispatch(
-                    std::integer_sequence<int, 1, 2, 4, 8, 16, 32, 64, 128>{},
-                    [&](auto S) { return max_size <= S; },
-                    [&](auto S) {
-                        constexpr int threads         = 128;
-                        constexpr int items_per_block = threads / S;
-                        const int     blocks          = (count + items_per_block - 1) / items_per_block;
-                        batchedCopy<S><<<blocks, threads, 0, st>>>(params);
-                    });
+                auto      pred     = [&](auto S) { return max_size <= S; };
+                auto      func     = [&](auto S) {
+                    constexpr int threads         = 128;
+                    constexpr int items_per_block = threads / S;
+                    const int     blocks          = (count + items_per_block - 1) / items_per_block;
+                    batchedCopy<S><<<blocks, threads, 0, st>>>(params);
+                };
+                dispatch(std::integer_sequence<int, 1, 2, 4, 8, 16, 32, 64, 128>{}, pred, func);
             }
         });
 }