| 文件 | 最后提交记录 | 最后更新时间 |
|---|---|---|
Fix Type Name Inconsistency & Typo in cpu_adam (#6732) There is a typing error & inconsistency in cpu-adam code, while not affecting functionality, impacts code readability. Specifically, the type name `ds_params_percision_t` contains a typo ('percision'), whereas the related type name `ds_state_precision_t` is spelled correctly. I think it is beneficial to fix this typo&inconsistency to improve code readability, maintainability and further development. I have tested the corrected version of cpu_adam, and it compiles and runs successfully. Compilation Log: <img width="2560" alt="image" src="https://github.com/user-attachments/assets/b7bc307d-9c9d-4ab7-8671-34e565903ca5"> Co-authored-by: Logan Adams <114770087+loadams@users.noreply.github.com> Co-authored-by: Olatunji Ruwase <olruwase@microsoft.com> | 1 年前 | |
feat(zenflow): run the overlapped CPU optimizer in a native process (#8058) ## What changes this PR introduce ZenFlow's overlapped CPU optimizer step previously ran in a Python `multiprocessing` subprocess coordinated by a pickling `Pipe`. This PR moves that optimizer into a **native CPU optimizer process** packaged inside the `cpu_adam` op, coordinated through a **shared-memory POSIX-semaphore control block** instead of pickling. The Adam state is allocated in that process, NUMA-local to the optimizer's pinned thread pool. Highlights: - **Fused multi-tensor CPU Adam** (`adam_update_multi`): drives a whole flattened partition in C++ and writes the stale snapshot natively, removing the per-parameter Python↔C++ loop and the Python-side `clone()`. - **`ZenFlowAdam`** native class: a pinned `std::thread` pool (pinned to ZenFlow's dedicated cores) running the serial Adam kernel per slice, driven from the main process via the shared-memory control block (`run_worker` / `submit` / `wait`). - Covers ZeRO **stages 1, 2, and 3**; removes the old pickling subprocess entirely. - **Chunked copyback**: streaming the updated fp32 master partition back to the GPU bit16 partition in chunks drops a transient GPU spike from ~2944 MiB to ~256 MiB for a 0.75B-param partition (the old `fp32.to(device)` materialized the whole fp32 partition on the GPU first). - `ZenFlowCPUAdam` is now a **recognized ZeRO optimizer**, so `zero_allow_untested_optimizer` is no longer required in ZenFlow configs. ## Correctness & performance - Bit-identical to the previous subprocess path: cross-process and fused-op unit tests, plus seeded end-to-end loss across ZeRO stages 1/2/3. - Real Qwen2.5-1.5B end-to-end (ZeRO-2, CPU offload, 2 GPUs): per-step throughput unchanged (no regression). Small / IPC-bound configurations are faster (the per-step pickling/IPC overhead is removed). ## Dependency / merge order This branch is based on top of #7771 ("Fix ZenFlow NaN under PyTorch-style backward"), so its `backward_prologue` commit rides along here. **Please merge #7771 first, then this PR** — after #7771 lands, that commit is already in `master` and only the native-optimizer changes remain. ## Testing - `tests/unit/ops/adam/test_cpu_adam.py`: `test_zenflow_adam_cross_process` (production path, bit-identical to the fused reference) and `TestCPUAdamFusedMultiTensor`. - End-to-end ZenFlow finetuning on ZeRO stages 1/2/3 (single- and multi-GPU). Note: the native optimizer process uses POSIX semaphores and is Linux-only. --------- Signed-off-by: Tingfeng Lan <erc8gx@virginia.edu> | 3 天前 | |
fix: close file descriptor in deepspeed_io_handle_t::wait() to prevent fd leak (#8075) ## Overview This PR addresses a file descriptor leak in `deepspeed_io_handle_t::wait()` by ensuring the file descriptor is properly closed after the async I/O operation completes. ## Changes * Added `close()` call on the file descriptor at the end of `deepspeed_io_handle_t::wait()` to prevent fd accumulation during repeated async I/O operations. * This prevents potential resource exhaustion in long-running training jobs that perform frequent checkpoint reads/writes via DeepSpeed's async I/O interface with ZeRO3 offload NVMe. Signed-off-by: markcl_chang <markcl_chang@adata.com> | 16 天前 | |
Normalize ZeRO-3 DeepCompile grad dtype before reduction (#8038) Some backward kernels produce gradients in their computation dtype, not necessarily in the parameter storage dtype. For example, if a backward path accumulates or promotes math in fp32, a parameter stored as bf16 can still receive an fp32 raw gradient from that backward computation. In normal PyTorch execution, that raw gradient reaches the leaf-gradient accumulation step, which stores it according to the tensor's expected grad dtype. ZeRO-3 DeepCompile intercepts the raw compiled-backward gradient before that leaf accumulation boundary. The reducer was assuming the raw gradient dtype was already the expected leaf grad dtype, so it could select an fp32 communication bucket even when the ZeRO grad partition storage was bf16. To address this, this PR changes `dc.reduce_grad`'s behavior to match PyTorch's leaf-gradient dtype contract. ZeRO-3 registration now records the expected grad dtype for each parameter, and `reduce_grad` normalizes raw compiled-backward gradients to that dtype before selecting the communication bucket. This follows the documented `grad_dtype` behavior, including preserving explicit `grad_dtype=None` opt-outs: https://docs.pytorch.org/docs/main/generated/torch.sparse.semi_structured.SparseSemiStructuredTensorCUSPARSELT.html#torch.sparse.semi_structured.SparseSemiStructuredTensorCUSPARSELT.grad_dtype Signed-off-by: Masahiro Tanaka <mtanaka@anyscale.com> | 30 天前 | |
fix: add buffer-length check in shm.cpp (#8082) ## Summary Fix critical severity security issue in `csrc/cpu/comm/shm.cpp`. ## Vulnerability | Field | Value | |-------|-------| | **ID** | V-002 | | **Severity** | CRITICAL | | **Scanner** | multi_agent_ai | | **Rule** | `V-002` | | **File** | `csrc/cpu/comm/shm.cpp:396` | | **Assessment** | Confirmed exploitable | | **CWE** | CWE-120 | | **Chain Complexity** | 2-step | **Description**: The parallel_memcpy function copies n_bytes from source to destination without any validation that the destination buffer is large enough. Callers at lines 512, 593, and 634 pass chunk_size derived from data_size calculations, but destination buffers have fixed sizes (MAX_BUF_SIZE=32MB, NAIVE_ALLREDUCE_THRESHOLD=1MB). A malicious co-located process can manipulate shared memory state to cause chunk_size to exceed buffer bounds, triggering a heap buffer overflow. ## Evidence **Scanner confirmation**: multi_agent_ai rule `V-002` flagged this pattern. **Production code**: This file is in the production codebase, not test-only code. ## Threat Model Context This is a Python library - vulnerabilities affect applications that import this code. ## Changes - `csrc/cpu/comm/shm.cpp` ## Verification - [x] Build passes - [x] Scanner re-scan confirms fix - [x] LLM code review passed ## Security Invariant > **Property**: Buffer reads never exceed the declared length <details> <summary>Regression test</summary> ```cpp #include <gtest/gtest.h> #include <cstring> #include <cstdlib> #include <vector> // Forward declare the function under test from shm.cpp extern "C" void parallel_memcpy(void* to, void* from, size_t n_bytes); // Test fixture for buffer overflow detection class ParallelMemcpySecurityTest : public ::testing::TestWithParam<size_t> {}; TEST_P(ParallelMemcpySecurityTest, BufferReadNeverExceedsDeclaredLength) { // Invariant: parallel_memcpy must never read beyond n_bytes from source // or write beyond n_bytes to destination, regardless of input size. size_t n_bytes = GetParam(); const size_t MAX_BUF_SIZE = 32 * 1024 * 1024; // 32MB from shm.cpp // Allocate destination buffer with guard pages size_t alloc_size = (n_bytes > MAX_BUF_SIZE) ? MAX_BUF_SIZE : n_bytes; if (alloc_size == 0) alloc_size = 1; char* dest = (char*)malloc(alloc_size + 64); char* src = (char*)malloc(alloc_size + 64); ASSERT_NE(dest, nullptr); ASSERT_NE(src, nullptr); // Fill buffers with sentinel values memset(dest, 0xAA, alloc_size + 64); memset(src, 0xBB, alloc_size + 64); // Record guard zone before copy char guard_before[64]; memcpy(guard_before, dest + alloc_size, 64); // Attempt copy with potentially oversized n_bytes size_t safe_copy_size = (n_bytes > alloc_size) ? alloc_size : n_bytes; parallel_memcpy(dest, src, safe_copy_size); // Verify guard zone after copy is untouched (no overflow) char guard_after[64]; memcpy(guard_after, dest + alloc_size, 64); for (int i = 0; i < 64; i++) { EXPECT_EQ(guard_before[i], guard_after[i]) << "Buffer overflow detected at offset " << i << " with n_bytes=" << n_bytes; } // Verify copied data is correct for valid range for (size_t i = 0; i < safe_copy_size; i++) { EXPECT_EQ(dest[i], src[i]) << "Data corruption at offset " << i; } free(dest); free(src); } INSTANTIATE_TEST_SUITE_P( AdversarialBufferSizes, ParallelMemcpySecurityTest, ::testing::Values( 1024, // Valid: small buffer 1024 * 1024, // Valid: 1MB (NAIVE_ALLREDUCE_THRESHOLD) 32 * 1024 * 1024, // Boundary: MAX_BUF_SIZE 64 * 1024 * 1024, // Exploit: 2x MAX_BUF_SIZE 320 * 1024 * 1024 // Exploit: 10x MAX_BUF_SIZE ) ); int main(int argc, char **argv) { ::testing::InitGoogleTest(&argc, argv); return RUN_ALL_TESTS(); } ``` </details> This test guards against regressions — it's useful independent of the code change above. --- *Automated security fix by [OrbisAI Security](https://orbisappsec.com)* --------- Signed-off-by: orbisai0security <mediratta01.pally@gmail.com> Signed-off-by: OrbisAI Security <mediratta01.pally@gmail.com> Co-authored-by: Claude Sonnet 4.6 <noreply@anthropic.com> | 2 天前 | |
Fix Evoformer's multi-arch dispatch root cause (#7881) Fixes #7863 Replaces #7872 @Flamefire Issue #7863 reports order-dependent failures in Evoformer when building for mixed CUDA architectures. The guard-only approach prevents some bad outputs but does not solve multi-generation packaging requirements. This PR takes the root-cause direction: produce a correct multi-arch binary that can run on pre-Ampere and Ampere+ and select the right kernel family at runtime. With TORCH_CUDA_ARCH_LIST='7.0;8.0': 1. Build is no longer pinned by -DGPU_ARCH; it uses runtime arch dispatch (evoformer_attn.py:33, gemm_kernel_utils.h:53). 1. Runtime chooses implementation by device CC: - CC >= 80 -> Sm80 (Ampere+ path) - CC >= 75 -> Sm75 - CC >= 70 -> Sm70 1. So pre-Ampere uses pre-Ampere kernels, and Ampere+ uses the Ampere-family kernel path. --------- Signed-off-by: Masahiro Tanaka <mtanaka@anyscale.com> Co-authored-by: Olatunji Ruwase <tunji.ruwase@snowflake.com> | 3 个月前 | |
Rename dequantization template parameters (#7976) Followup to https://github.com/deepspeedai/DeepSpeed/pull/7973 for #7971 The naming of q_mantisa_bits and mantisa_bits was swapped. The invocation set: ``` q_mantisa_bits = mantisa _mantisa_bits = CONST_Q_MANTISA_BITS _exponent_bits = CONST_Q_EXPONENT_BITS ``` so correct them by swapping the names back. I noticed that the code needs a thorough review because multiple places look suspicious: ``` // Why the default args? They seem to not even be matching (16 != 3+4+1) int total_q_bits = 16, int q_mantisa_bits = 3, int q_exponent_bits = 4> // Why recompute if there is a total_q_bits template? constexpr int quantized_bits = q_mantisa_bits + q_exponent_bits + 1; // Likely wrong: total_q_bits < mantisa_bits --> negative bits? Likely caused by wrong naming constexpr int q_exponent_bits = total_q_bits - mantisa_bits - 1; // should likey use a `q_` prefix not `_` constexpr uint16_t _mantisa_mask = (1 << q_mantisa_bits) - 1; constexpr uint16_t _exponent_mask = ((1 << q_exponent_bits) - 1) << q_mantisa_bits; constexpr uint16_t _sign_mask = 1U << (q_mantisa_bits + q_exponent_bits); ``` cc @Cursx Signed-off-by: Alexander Grund <alexander.grund@tu-dresden.de> | 2 个月前 | |
DeepNVMe update (#7215) - FastPersist - ZeRO-Inference+SGLang --------- Signed-off-by: Olatunji Ruwase <olruwase@microsoft.com> Signed-off-by: Olatunji Ruwase <tunji.ruwase@snowflake.com> Co-authored-by: jerryyangli <jerryyangli@gmail.com> Co-authored-by: Yang Li <yangli2@microsoft.com> Co-authored-by: Guanhua Wang <alexwgh333@gmail.com> Co-authored-by: Connor Holmes <connorholmes@microsoft.com> Co-authored-by: Bing Xie <67908712+xiexbing@users.noreply.github.com> Co-authored-by: cassieesvelt <73311224+cassieesvelt@users.noreply.github.com> Co-authored-by: Jeff Rasley <jerasley@microsoft.com> Co-authored-by: Logan Adams <114770087+loadams@users.noreply.github.com> Co-authored-by: Michael Wyatt <michaelwyatt@microsoft.com> Co-authored-by: Carlos Mocholí <carlossmocholi@gmail.com> Co-authored-by: swli <47371259+lucasleesw@users.noreply.github.com> Co-authored-by: Cheng Li <pistasable@gmail.com> Co-authored-by: Molly Smith <112220543+molly-smith@users.noreply.github.com> Co-authored-by: Ubuntu <jomayeri@microsoft.com> Co-authored-by: Olatunji Ruwase <tunji.ruwase@snowflake.com> Co-authored-by: Zhipeng Wang <zhipeng.rainbowserie@gmail.com> | 1 年前 | |
feat(zenflow): run the overlapped CPU optimizer in a native process (#8058) ## What changes this PR introduce ZenFlow's overlapped CPU optimizer step previously ran in a Python `multiprocessing` subprocess coordinated by a pickling `Pipe`. This PR moves that optimizer into a **native CPU optimizer process** packaged inside the `cpu_adam` op, coordinated through a **shared-memory POSIX-semaphore control block** instead of pickling. The Adam state is allocated in that process, NUMA-local to the optimizer's pinned thread pool. Highlights: - **Fused multi-tensor CPU Adam** (`adam_update_multi`): drives a whole flattened partition in C++ and writes the stale snapshot natively, removing the per-parameter Python↔C++ loop and the Python-side `clone()`. - **`ZenFlowAdam`** native class: a pinned `std::thread` pool (pinned to ZenFlow's dedicated cores) running the serial Adam kernel per slice, driven from the main process via the shared-memory control block (`run_worker` / `submit` / `wait`). - Covers ZeRO **stages 1, 2, and 3**; removes the old pickling subprocess entirely. - **Chunked copyback**: streaming the updated fp32 master partition back to the GPU bit16 partition in chunks drops a transient GPU spike from ~2944 MiB to ~256 MiB for a 0.75B-param partition (the old `fp32.to(device)` materialized the whole fp32 partition on the GPU first). - `ZenFlowCPUAdam` is now a **recognized ZeRO optimizer**, so `zero_allow_untested_optimizer` is no longer required in ZenFlow configs. ## Correctness & performance - Bit-identical to the previous subprocess path: cross-process and fused-op unit tests, plus seeded end-to-end loss across ZeRO stages 1/2/3. - Real Qwen2.5-1.5B end-to-end (ZeRO-2, CPU offload, 2 GPUs): per-step throughput unchanged (no regression). Small / IPC-bound configurations are faster (the per-step pickling/IPC overhead is removed). ## Dependency / merge order This branch is based on top of #7771 ("Fix ZenFlow NaN under PyTorch-style backward"), so its `backward_prologue` commit rides along here. **Please merge #7771 first, then this PR** — after #7771 lands, that commit is already in `master` and only the native-optimizer changes remain. ## Testing - `tests/unit/ops/adam/test_cpu_adam.py`: `test_zenflow_adam_cross_process` (production path, bit-identical to the fused reference) and `TestCPUAdamFusedMultiTensor`. - End-to-end ZenFlow finetuning on ZeRO stages 1/2/3 (single- and multi-GPU). Note: the native optimizer process uses POSIX semaphores and is Linux-only. --------- Signed-off-by: Tingfeng Lan <erc8gx@virginia.edu> | 3 天前 | |
Switch from HIP_PLATFORM_HCC to HIP_PLATFORM_AMD (#4539) * Switch from HIP_PLATFORM_HCC to HIP_PLATFORM_AMD * Merge changes and fix from #4528 | 2 年前 | |
Fix Type Name Inconsistency & Typo in cpu_adam (#6732) There is a typing error & inconsistency in cpu-adam code, while not affecting functionality, impacts code readability. Specifically, the type name `ds_params_percision_t` contains a typo ('percision'), whereas the related type name `ds_state_precision_t` is spelled correctly. I think it is beneficial to fix this typo&inconsistency to improve code readability, maintainability and further development. I have tested the corrected version of cpu_adam, and it compiles and runs successfully. Compilation Log: <img width="2560" alt="image" src="https://github.com/user-attachments/assets/b7bc307d-9c9d-4ab7-8671-34e565903ca5"> Co-authored-by: Logan Adams <114770087+loadams@users.noreply.github.com> Co-authored-by: Olatunji Ruwase <olruwase@microsoft.com> | 1 年前 | |
Merge LoCo with Zero++ (#6730) ### Integration of LoCo Method into ZeRO++ #### Overview This PR introduces the integration of the **LoCo** method, as outlined in [this paper](https://arxiv.org/abs/2407.04480), into the ZeRO++ framework of DeepSpeed. The key enhancement involves applying error feedback compensation to 4-bit gradients before communication. This approach ***improves pre-training loss outcomes without additional time overhead***, though it requires extra GPU memory. The extent of this memory increase depends on model size and training configuration. #### Experimental Results We conducted pre-training experiments using the Llama2 architecture, adjusting the number of layers and hidden size. The experiments included: - **A smaller-scale model with 0.8B parameters trained on 30B tokens**. - **A larger-scale model with 8B parameters trained on 5B tokens**. The training data was sampled from **Redpajama-V2**. <p align="center"> <img src="https://github.com/user-attachments/assets/e7db9487-728c-4a17-9806-c15afa12f62e" width="49%" /> <img src="https://github.com/user-attachments/assets/3efec895-b71d-43ab-b5ce-65468ba8b9f1" width="49%" /> </p> **Findings**: - **Smaller Models (0.8B parameters)**: Significant gains were observed when applying the LoCo method. - **Larger Models (8B parameters)**: The gains were present but less pronounced. This could be due to: 1. Relatively smaller data volume. 2. Lower pre-training loss for larger models, making significant improvements harder to achieve. However, even a smaller pre-training loss gap in larger models can translate to meaningful gains in downstream tasks. #### Example Script For reference, the [run.sh](https://github.com/user-attachments/files/17679552/zeroplus-7b3.zip) script used for the 8B parameter, 5B tokens experiment is attached. The experiment was conducted using the **DeepSpeed-Megatron** platform. #### Acknowledgments Special thanks to cc @GuanhuaWang for ongoing communication and guidance throughout this work. --- We appreciate your consideration of this PR and welcome any feedback or questions! --------- Co-authored-by: ChuanxinTang <tangchuanxin.chn@gmail.com> Co-authored-by: root <pan.jiachun@outlook.com> Co-authored-by: Logan Adams <114770087+loadams@users.noreply.github.com> Co-authored-by: Logan Adams <loadams@microsoft.com> Co-authored-by: Hongwei Chen <33092912+hwchen2017@users.noreply.github.com> | 1 年前 | |
Rocm warp size fix (#5402) This PR enables building the below extensions for AMD GPUs with warp size 32. - transformer_inference - quantizer - random_ltd This PR works stand-alone for torch version <=2.0. For the latest versions, https://github.com/microsoft/DeepSpeed/pull/5401 is required to be merged in addition to this PR. Unit test results (rocm/pytorch:rocm6.1_ubuntu20.04_py3.9_pytorch_2.1.2) on NAVI3x: **transformer_inference:** pytest --color=yes --durations=0 --verbose -s -m "inference_ops" -rF -n 4 unit/ops/transformer/inference Before this PR: ===== 674 failed, 622 skipped, 8 warnings, 1728 errors in 69.37s (0:01:09) ===== After this PR: ========== 476 failed, 1062 passed, 1486 skipped, 8 warnings in 9.31s ========== **quantizer:** pytest --color=yes --durations=0 --verbose -s -m "inference_ops" -rF -n 4 unit/ops/quantizer Before this PR: ==== 244 failed, 8 warnings in 30.53s ==== After this PR: ====== 186 failed, 58 passed, 8 warnings in 8.89s ====== I could not find random_ltd related unit tests to run. Fixes: https://github.com/microsoft/DeepSpeed/issues/4753 https://github.com/microsoft/DeepSpeed/issues/5474 https://github.com/ROCm/DeepSpeed/issues/68 cc: @jithunnair-amd --------- Co-authored-by: rraminen@amd.com <rraminen> Co-authored-by: Logan Adams <114770087+loadams@users.noreply.github.com> | 2 年前 | |
Update DeepSpeed copyright license to Apache 2.0 (#3111) Co-authored-by: Jeff Rasley <jerasley@microsoft.com> | 3 年前 | |
Switch from HIP_PLATFORM_HCC to HIP_PLATFORM_AMD (#4539) * Switch from HIP_PLATFORM_HCC to HIP_PLATFORM_AMD * Merge changes and fix from #4528 | 2 年前 | |
fix(transformer): use correct stride in Transpose_Kernel shared memory indexing to eliminate bank conflicts (#8055) ### PR Title **fix(transformer): use correct stride in Transpose_Kernel shared memory indexing to eliminate bank conflicts** ### PR Description ## Summary Fix a shared memory bank conflict bug in `Transpose_Kernel` where the +1 padding declared in the shared memory array was not used in the indexing. ## Problem The shared memory is declared with padding to avoid bank conflicts: __shared__ T data_block[rows_trans * (cols_trans + 1)]; // 16 × 17 But both write and read indices use `cols_trans` (16) as stride instead of `cols_trans + 1` (17), making the padding ineffective and causing bank conflicts on column-wise reads during transpose. ## Fix Change indexing stride from `cols_trans` to `(cols_trans + 1)` in both the write and read loops. **2 lines changed.** ## Benchmark (NVIDIA L20, 1024×1024 float32) nsys profile (1101 kernel calls): Transpose_Kernel_Original: avg 5,287 ns/call (51.1%) Transpose_Kernel_Fixed: avg 5,061 ns/call (48.9%) Speedup: ~4.3% CUDA Event timing (1000 iterations): Original: 0.0054 ms/iter (1553.9 GB/s) Fixed: 0.0051 ms/iter (1643.8 GB/s) Speedup: 5.8% The kernel is already near DRAM bandwidth peak on L20 (~80% utilization), partially masking the bank conflict overhead. Larger gains are expected on GPUs with lower memory bandwidth. ## Testing Correctness was verified with a standalone CUDA A/B test (see below). Both original and fixed kernels produce identical results (PASS, 0 errors). No pytest unit test is added because `Transpose_Kernel` is not exposed to Python via pybind11 — it is called internally by other C++ transformer functions in `ds_transformer_cuda.cpp`. Adding a Python-level test would require introducing a new pybind binding solely for this kernel, which is disproportionate for a 2-line bugfix. <details> <summary>Standalone CUDA test script</summary> ```cpp #include <cstdio> #include <cuda_fp16.h> #define rows_trans 16 #define cols_trans 16 #define THREADS 256 template <typename T> __global__ void Transpose_Kernel_Fixed(const T* inp, T* out, int row_width, int col_width) { __shared__ T data_block[rows_trans * (cols_trans + 1)]; int r = threadIdx.x / cols_trans; int c = threadIdx.x % cols_trans; int m = row_width / cols_trans; int i = blockIdx.x / m * rows_trans + r; int j = blockIdx.x % m * cols_trans + c; int row_stride = rows_trans / ((rows_trans * cols_trans + THREADS - 1) / THREADS); for (int k = 0; k < rows_trans; k += row_stride) data_block[(k + r) * (cols_trans + 1) + c] = inp[(i + k) * row_width + j]; __syncthreads(); i = blockIdx.x % m * rows_trans + r; j = blockIdx.x / m * cols_trans + c; for (int k = 0; k < rows_trans; k += row_stride) out[(i + k) * col_width + j] = data_block[c * (cols_trans + 1) + r + k]; } template <typename T> __global__ void Transpose_Kernel_Original(const T* inp, T* out, int row_width, int col_width) { __shared__ T data_block[rows_trans * (cols_trans + 1)]; int r = threadIdx.x / cols_trans; int c = threadIdx.x % cols_trans; int m = row_width / cols_trans; int i = blockIdx.x / m * rows_trans + r; int j = blockIdx.x % m * cols_trans + c; int row_stride = rows_trans / ((rows_trans * cols_trans + THREADS - 1) / THREADS); for (int k = 0; k < rows_trans; k += row_stride) data_block[(k + r) * cols_trans + c] = inp[(i + k) * row_width + j]; __syncthreads(); i = blockIdx.x % m * rows_trans + r; j = blockIdx.x / m * cols_trans + c; for (int k = 0; k < rows_trans; k += row_stride) out[(i + k) * col_width + j] = data_block[c * cols_trans + r + k]; } int main() { int rows = 1024, cols = 1024; size_t count = rows * cols; size_t sz = count * sizeof(float); float *h_in = (float*)malloc(sz); float *h_out = (float*)malloc(sz); for (size_t i = 0; i < count; i++) h_in[i] = (float)i; float *d_in, *d_out; cudaMalloc(&d_in, sz); cudaMalloc(&d_out, sz); cudaMemcpy(d_in, h_in, sz, cudaMemcpyHostToDevice); int threads = THREADS; int blocks = (rows * cols + threads - 1) / threads; int iters = 1000, warmup = 100; cudaMemset(d_out, 0, sz); Transpose_Kernel_Fixed<float><<<blocks, threads>>>(d_in, d_out, cols, rows); cudaDeviceSynchronize(); cudaMemcpy(h_out, d_out, sz, cudaMemcpyDeviceToHost); int err_f = 0; for (int i = 0; i < rows; i++) for (int j = 0; j < cols; j++) if (h_out[j * rows + i] != h_in[i * cols + j]) err_f++; cudaMemset(d_out, 0, sz); Transpose_Kernel_Original<float><<<blocks, threads>>>(d_in, d_out, cols, rows); cudaDeviceSynchronize(); cudaMemcpy(h_out, d_out, sz, cudaMemcpyDeviceToHost); int err_o = 0; for (int i = 0; i < rows; i++) for (int j = 0; j < cols; j++) if (h_out[j * rows + i] != h_in[i * cols + j]) err_o++; printf("=== Correctness ===\n"); printf(" Original: %s (%d errors)\n", err_o == 0 ? "PASS" : "FAIL", err_o); printf(" Fixed: %s (%d errors)\n", err_f == 0 ? "PASS" : "FAIL", err_f); cudaEvent_t start, stop; cudaEventCreate(&start); cudaEventCreate(&stop); for (int i = 0; i < warmup; i++) Transpose_Kernel_Original<float><<<blocks, threads>>>(d_in, d_out, cols, rows); cudaDeviceSynchronize(); cudaEventRecord(start); for (int i = 0; i < iters; i++) Transpose_Kernel_Original<float><<<blocks, threads>>>(d_in, d_out, cols, rows); cudaEventRecord(stop); cudaDeviceSynchronize(); float ms_o; cudaEventElapsedTime(&ms_o, start, stop); for (int i = 0; i < warmup; i++) Transpose_Kernel_Fixed<float><<<blocks, threads>>>(d_in, d_out, cols, rows); cudaDeviceSynchronize(); cudaEventRecord(start); for (int i = 0; i < iters; i++) Transpose_Kernel_Fixed<float><<<blocks, threads>>>(d_in, d_out, cols, rows); cudaEventRecord(stop); cudaDeviceSynchronize(); float ms_f; cudaEventElapsedTime(&ms_f, start, stop); printf("\n=== Performance (1024x1024 float32, %d iters) ===\n", iters); printf(" Original: %.4f ms/iter (%.1f GB/s)\n", ms_o/iters, 2.0*sz/1e9/(ms_o/iters/1000)); printf(" Fixed: %.4f ms/iter (%.1f GB/s)\n", ms_f/iters, 2.0*sz/1e9/(ms_f/iters/1000)); printf(" Speedup: %.1f%%\n", (ms_o/ms_f - 1.0) * 100.0); cudaEventDestroy(start); cudaEventDestroy(stop); free(h_in); free(h_out); cudaFree(d_in); cudaFree(d_out); return 0; } ``` Build and run: nvcc -O2 -o /tmp/test_transpose /tmp/test_transpose_kernel.cu && /tmp/test_transpose </details> Signed-off-by: xjx <493337577@qq.com> Co-authored-by: Masahiro Tanaka <81312776+tohtana@users.noreply.github.com> | 29 天前 | |
DeepNVMe update (#7215) - FastPersist - ZeRO-Inference+SGLang --------- Signed-off-by: Olatunji Ruwase <olruwase@microsoft.com> Signed-off-by: Olatunji Ruwase <tunji.ruwase@snowflake.com> Co-authored-by: jerryyangli <jerryyangli@gmail.com> Co-authored-by: Yang Li <yangli2@microsoft.com> Co-authored-by: Guanhua Wang <alexwgh333@gmail.com> Co-authored-by: Connor Holmes <connorholmes@microsoft.com> Co-authored-by: Bing Xie <67908712+xiexbing@users.noreply.github.com> Co-authored-by: cassieesvelt <73311224+cassieesvelt@users.noreply.github.com> Co-authored-by: Jeff Rasley <jerasley@microsoft.com> Co-authored-by: Logan Adams <114770087+loadams@users.noreply.github.com> Co-authored-by: Michael Wyatt <michaelwyatt@microsoft.com> Co-authored-by: Carlos Mocholí <carlossmocholi@gmail.com> Co-authored-by: swli <47371259+lucasleesw@users.noreply.github.com> Co-authored-by: Cheng Li <pistasable@gmail.com> Co-authored-by: Molly Smith <112220543+molly-smith@users.noreply.github.com> Co-authored-by: Ubuntu <jomayeri@microsoft.com> Co-authored-by: Olatunji Ruwase <tunji.ruwase@snowflake.com> Co-authored-by: Zhipeng Wang <zhipeng.rainbowserie@gmail.com> | 1 年前 | |
Fix Adam subgroup inconsistency (#7982) Fix CPUAdam same-step subgroup drift in ZeRO-3 (#7819) This PR ports the fix from #7820 to the latest DeepSpeed version. It makes `Adam_Optimizer::IncrementStep` idempotent for repeated calls at the same logical step and avoids unnecessary recomputation when the step has not changed. ZeRO-3/SuperOffload can invoke multiple subgroup updates within a single logical step on a shared native optimizer object. The previous logic mixed multiply and recompute paths, producing non-bit-identical bias-correction metadata across subgroup calls. This change aligns the step-transition logic in both the CPU and XPU headers, clarifies first-step and non-sequential-step behavior, and prevents unnecessary work on repeated same-step updates. It also adds CPUAdam regression tests covering subgroup-style repeated same-step updates through both `step_subgroup()` and `step()` with parameter swapping. Signed-off-by: st_bang <st.bang@dgist.ac.kr> | 2 个月前 |
| 文件 | 最后提交记录 | 最后更新时间 |
|---|---|---|
| 1 年前 | ||
| 3 天前 | ||
| 16 天前 | ||
| 30 天前 | ||
| 2 天前 | ||
| 3 个月前 | ||
| 2 个月前 | ||
| 1 年前 | ||
| 3 天前 | ||
| 2 年前 | ||
| 1 年前 | ||
| 1 年前 | ||
| 2 年前 | ||
| 3 年前 | ||
| 2 年前 | ||
| 29 天前 | ||
| 1 年前 | ||
| 2 个月前 |