gemini-code-assist[bot] commented on PR #18877: URL: https://github.com/apache/tvm/pull/18877#issuecomment-4008791360
## Summary of Changes Hello, I'm Gemini Code Assist[^1]! I'm currently reviewing this pull request and will post my feedback shortly. In the meantime, here's a summary to help you and other reviewers quickly get up to speed! This pull request significantly enhances Metal runtime performance by optimizing command dispatch and memory copy operations. By batching compute kernels and inlining blit operations into a single command buffer, it reduces overhead associated with frequent command buffer creation and submission. The introduction of a staging buffer pool further streamlines CPU-to-GPU transfers, while conditional synchronization prevents unnecessary GPU stalls. These changes result in substantial speedups, particularly for smaller models that are dispatch-bound, as demonstrated by the provided benchmarks. ### Highlights * **Batched Compute Dispatch**: Kernel dispatches are now accumulated in a single `MTLCommandBuffer` via a shared `MTLComputeCommandEncoder`, instead of each dispatch creating and committing its own command buffer. The command buffer is flushed on GPU-to-CPU readback, buffer deallocation, or stream sync. * **Inline Blit Encoders for Copies**: CPU-to-GPU and GPU-to-GPU copies now utilize blit encoders on the same pending command buffer, eliminating the need for separate command buffers per copy. Metal's sequential ordering guarantee within a command buffer ensures correct execution without explicit synchronization. * **Staging Buffer Pool for CPU-to-GPU Copies**: A per-device `StagingBufferPool` has been introduced to manage shared-mode buffers for inlined CPU-to-GPU copies. These buffers are recycled after a flush or sync, allowing multiple CPU-to-GPU copies to be inlined efficiently. * **Conditional Sync in FreeDataSpace**: The `FreeDataSpace` function now checks for pending work (`HasPendingWork()`) before performing a stream synchronization. This avoids redundant syncs when the GPU-to-CPU readback path has already flushed and waited, improving efficiency. 🧠**New Feature in Public Preview:** You can now enable **Memory** to help **Gemini Code Assist** learn from your team's feedback. This makes future code reviews more consistent and personalized to your project's style. **Click [here](https://codeassist.google/code-review/login) to enable Memory in your admin console.** <details> <summary><b>Changelog</b></summary> * **src/runtime/metal/metal_common.h** * Refactored `Stream` class to support batched command dispatch and blit operations, including new methods like `FlushCommandBuffer`, `Synchronize`, `GetPendingComputeEncoder`, and `GetBlitEncoderOnPendingBuffer`. * Added `ProfileCounters` struct to `Stream` for tracking various GPU operation metrics (dispatches, flushes, syncs, blits, copies). * Introduced private helper methods `GetOrCreatePendingCommandBuffer` and `PauseComputeEncoder` within the `Stream` class. * Defined `StagingBufferPool` struct and integrated it into `MetalThreadEntry` to manage buffers for CPU-to-GPU copies. * Added `staging_pools_` vector and `GetOrCreateStagingBuffer` method to `MetalThreadEntry`. * **src/runtime/metal/metal_device_api.mm** * Modified `FreeDataSpace` to conditionally synchronize the stream only if pending work exists, preventing unnecessary stalls. * Updated `CopyDataFromTo` to use batched blit encoders for GPU-to-GPU and CPU-to-GPU copies, and to explicitly flush the command buffer for GPU-to-CPU copies. * Integrated the `StagingBufferPool` for efficient CPU-to-GPU copies. * Replaced direct command buffer creation with `Stream::Synchronize` in `StreamSync` and added staging buffer pool index reset. * Adjusted `GetTempBuffer` logic to conditionally sync before releasing old temporary buffers. * Exposed `metal.GetProfileCounters` and `metal.ResetProfileCounters` via FFI for runtime profiling. * **src/runtime/metal/metal_module.mm** * Modified `MetalFunc::operator()` to use `Stream::GetPendingComputeEncoder` for kernel dispatches, enabling batching of compute operations. * Removed immediate command buffer commit and associated error handling from kernel dispatch, deferring these responsibilities to the `Stream` class. </details> <details> <summary><b>Using Gemini Code Assist</b></summary> <br> The full guide for Gemini Code Assist can be found on our [documentation page](https://developers.google.com/gemini-code-assist/docs/review-github-code), here are some quick tips. <b>Invoking Gemini</b> You can request assistance from Gemini at any point by creating a comment using either `/gemini <command>` or `@gemini-code-assist <command>`. Below is a summary of the supported commands on the current page. Feature | Command | Description --- | --- | --- Code Review | `/gemini review` | Performs a code review for the current pull request in its current state. Pull Request Summary | `/gemini summary` | Provides a summary of the current pull request in its current state. Comment | @gemini-code-assist | Responds in comments when explicitly tagged, both in pull request comments and review comments. Help | `/gemini help` | Displays a list of available commands. <b>Customization</b> To customize Gemini Code Assist for GitHub experience, repository maintainers can create a configuration file and/or provide a custom code review style guide (such as PEP-8 for Python) by creating and adding files to a `.gemini/` folder in the base of the repository. Detailed instructions can be found [here](https://developers.google.com/gemini-code-assist/docs/customize-gemini-behavior-github). <b>Limitations & Feedback</b> Gemini Code Assist may make mistakes. Please leave feedback on any instances where its feedback is incorrect or counter productive. You can react with :thumbsup: and :thumbsdown: on @gemini-code-assist comments. If you're interested in giving your feedback about your experience with Gemini Code Assist for Github and other Google products, sign up [here](https://google.qualtrics.com/jfe/form/SV_2cyuGuTWsEw84yG). <b>You can also get AI-powered code generation, chat, as well as code reviews directly in the IDE at no cost with the [Gemini Code Assist IDE Extension](https://cloud.google.com/products/gemini/code-assist).</b> </details> [^1]: Review the [Privacy Notices](https://policies.google.com/privacy), [Generative AI Prohibited Use Policy](https://policies.google.com/terms/generative-ai/use-policy), [Terms of Service](https://policies.google.com/terms), and learn how to configure Gemini Code Assist in GitHub [here](https://developers.google.com/gemini-code-assist/docs/customize-gemini-behavior-github). Gemini can make mistakes, so double check it and [use code with caution](https://support.google.com/legal/answer/13505487). -- This is an automated message from the Apache Git Service. To respond to the message, please log on to GitHub and use the URL above to go to the specific comment. To unsubscribe, e-mail: [email protected] For queries about this service, please contact Infrastructure at: [email protected] --------------------------------------------------------------------- To unsubscribe, e-mail: [email protected] For additional commands, e-mail: [email protected]
