9/18/2024
AMD Distributed Training Overview
9/17/2024
What is IREE turbine
IREE-Turbine is a package or toolset that combines PyTorch, Torch-MLIR, IREE, and additional tools to provide a comprehensive solution for compiling, optimizing, and executing PyTorch models using IREE's infrastructure. Based on the information in the image, IREE-Turbine offers the following key features:
1. AOT Export: This allows for Ahead-Of-Time compilation of PyTorch modules (nn.Modules) into deployment-ready artifacts. These compiled artifacts can then take full advantage of IREE's runtime features.
2. Eager Execution: It provides a torch.compile backend and a Turbine Tensor/Device for interactive PyTorch sessions. This enables users to work with PyTorch in a familiar environment while leveraging IREE's optimization capabilities.
3. Custom Ops: IREE-Turbine offers integration for defining custom PyTorch operations and implementing them using either IREE's backend IR or the Pythonic kernel language. This allows for extending PyTorch's functionality while maintaining compatibility with IREE's optimization pipeline.
In essence, IREE-Turbine acts as a bridge between PyTorch and IREE, allowing PyTorch users to benefit from IREE's advanced compilation and runtime features while maintaining a familiar PyTorch-based workflow. It aims to provide a seamless experience for compiling PyTorch models to run efficiently on various hardware targets supported by IREE.
HIP kernel for matrix multiplication that can leverage Matrix Cores
Key points about this example:
1. It uses `half` precision for input matrices A and B, which can potentially benefit from Matrix Core acceleration.
2. The kernel is designed for 16x16 matrices, which is a common size for Matrix Core operations.
3. Shared memory is used to improve performance by reducing global memory accesses.
4. The main computation loop uses `__half2float` conversions. On GPUs with native FP16 support, these conversions might be optimized out.
5. The kernel uses a tiled approach, which is generally efficient for matrix multiplication.
6. Error checking is included for HIP calls.
Important considerations:
1. This kernel doesn't guarantee the use of Matrix Cores. The actual use of Matrix Cores depends on the GPU architecture and the HIP compiler's optimizations.
2. For larger matrices, you'd need to implement a more sophisticated tiling strategy.
3. Performance tuning is crucial. You might need to experiment with different block sizes and memory access patterns for optimal performance.
4. The HIP runtime and compiler will attempt to optimize this code for the target GPU, potentially leveraging Matrix Cores if available.
5. For production use, you should implement proper error handling and potentially use more sophisticated synchronization methods.
To fully leverage Matrix Cores, you might need to use specific intrinsics or rely on compiler optimizations. The exact method can vary depending on the GPU architecture and HIP version. Always profile your code to ensure you're getting the expected performance benefits.
Creating a custom CUDA kernel that directly utilizes tensor cores
Creating a custom CUDA kernel that directly utilizes tensor cores is an advanced topic, as tensor cores are typically accessed through higher-level libraries like cuBLAS or cuDNN. However, NVIDIA does provide a way to use tensor cores in custom kernels through their CUDA Core library, specifically with Warp Matrix Multiply-Accumulate (WMMA) API. Here's an overview of how to create a kernel that works on tensor cores:
Key points about this example:
1. We're using 16x16 matrices as this is a common size for tensor core operations.
2. The kernel uses `nvcuda::wmma::fragment` to define matrix fragments.
3. `load_matrix_sync`, `mma_sync`, and `store_matrix_sync` are used to load data, perform matrix multiplication, and store results using tensor cores.
4. The kernel operates on half-precision input (`half`) and produces single-precision output (`float`).
To use this kernel:
1. Compile with a CUDA compiler that supports tensor cores (CUDA 9.0 or later).
2. Use appropriate GPU architecture flags (e.g., `-arch=sm_70` for Volta, `-arch=sm_75` for Turing).
3. Allocate memory and copy data to the GPU before calling `launch_wmma_kernel`.
Important considerations:
1. Error checking is omitted for brevity but should be included in production code.
2. This is a basic example. Real-world usage often involves tiling and more complex memory access patterns for larger matrices.
3. Performance tuning is crucial. The exact dimensions and data types should be chosen based on your specific use case and target GPU architecture.
4. Not all operations can be efficiently mapped to tensor cores. They're most beneficial for large matrix multiplications common in deep learning workloads.
Remember, while this approach gives you direct control over tensor core usage, in many cases, using higher-level libraries like cuBLAS or cuDNN is more practical and can automatically leverage tensor cores when appropriate.
9/16/2024
Pytorch model to mlir -> llvm -> executable file on Mac book m1
-
This article explain how to access the thread index when you make block and thread with two dimensions. please refer to this page about me...
-
As you can see in the following video, I created a class that stitching n cameras in real time. https://www.youtube.com/user/feelmare/sear...
-
* Introduction - The solution shows panorama image from multi images. The panorama images is processing by real-time stitching algorithm...
-
The latent SVM tells the learning method used in this paper -> "Discriminatively trained deformable part models". The authors s...
-
EMD(earth mover distance) method is very good method to compare image similarity. But processing time is slow. For using the EMD compare, ...
-
This is GPU version of this page -> http://feelmare.blogspot.kr/2013/12/finding-largest-subset-images-that-is.html Please refer detail ...
-
This is dithering example, it make image like a stippling effect. I referenced to blew website. wiki page: https://en.wikipedia.org/wik...
-
My Environment : MS VS 2008 & MFC(Dialog Based) Joy Stick : Logitech Extreme 3D pro (XBox Type) Cteated Date : 2012. 03 [source code]...
-
If you meet this error message "afxcontrolbars.h - No such file or directory...", while complie on the VS C++. You can solve this...
-
Created Date : 2011.10 Language : C/C++ Tool : Microsoft Visual C++ 2008 Library & Utilized : OpenCV 2.3 Reference : SIFT referenc...