Hi Xuanteng, On 2024-05-05 08:30, Xuanteng Huang wrote:
I’m going to package MIOpen [1], ROCm’s counterpart to CUDA cuBLAS which provides several deep learning primitives (Conv, Pooling, etc.) for PyTorch and other DL frameworks.
Minor correction. MIOpen is more akin to cuDNN. The ROCm counterpart to cuBLAS is rocBLAS/hipBLAS.
There exists some large binary files (*.kdb.bz2) stored as Git LFS[2]. These files are served as kernel performance database [3] for different GPU generations, which will be loaded at runtime by MIOpen [4] to obtain the kernel code objects based on the queried problem size and corresponding solution.
I don't have an answer for you, unfortunately.When looking through this, I also became distracted by the assembly files in the directory you linked [2]. Did someone really write three thousand lines of uncommented assembly code, or is this file a disassembly of some other source [9]? Maybe they started with a disassembly of the compiled reference kernel, then tweaked it? I do wonder how these files were created and how they're maintained.
I suppose I have the same question about the performance database files. How are they created and how are they maintained?
Sincerely, Cory Bloor
[9]: https://salsa.debian.org/rocm-team/miopen/-/blob/upstream/5.5.1/src/kernels/Conv_Winograd_v14_3_3_fp16dot_stride2_dec.inc?ref_type=tags[1]: https://rocm.docs.amd.com/projects/MIOpen/en/latest/ [2]: https://salsa.debian.org/rocm-team/miopen/-/tree/master/src/kernels?ref_type=heads [3]: https://rocm.docs.amd.com/projects/MIOpen/en/develop/conceptual/perfdb.html [4]: https://salsa.debian.org/rocm-team/miopen/-/blob/master/src/binary_cache.cpp?ref_type=heads#L114-129 [5]: https://salsa.debian.org/rocm-team/miopen/-/blob/master/CMakeLists.txt?ref_type=heads#L501-521 [6]: https://salsa.debian.org/rocm-team/miopen/-/blob/master/src/binary_cache.cpp?ref_type=heads#L162-180 [7]: https://salsa.debian.org/rocm-team/miopen/-/blob/master/src/hip/handlehip.cpp?ref_type=heads#L403-413 [8]: https://salsa.debian.org/rocm-team/miopen/-/blob/master/src/hip/handlehip.cpp?ref_type=heads#L415-424