Skip to content

Commit

Permalink
fix msvc build
Browse files Browse the repository at this point in the history
  • Loading branch information
lzhangzz committed Nov 30, 2023
1 parent 6a3fe75 commit 751eeb6
Show file tree
Hide file tree
Showing 3 changed files with 10 additions and 11 deletions.
2 changes: 1 addition & 1 deletion builder/windows/README.md
Original file line number Diff line number Diff line change
Expand Up @@ -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
Expand Down
2 changes: 1 addition & 1 deletion builder/windows/generate.ps1
Original file line number Diff line number Diff line change
@@ -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 `
Expand Down
17 changes: 8 additions & 9 deletions src/turbomind/models/llama/llama_kernels.cu
Original file line number Diff line number Diff line change
Expand Up @@ -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);
}
});
}
Expand Down

0 comments on commit 751eeb6

Please sign in to comment.