[rocWMMA] Cherry picks for ROCm 7.2 (#2842)

Author: bsyrowik

CHANGELOG.md

@@ -3,12 +3,30 @@
 Documentation for rocWMMA is available at
 [https://rocm.docs.amd.com/projects/rocWMMA/en/latest](https://rocm.docs.amd.com/projects/rocWMMA/en/latest).
 
-## (Unreleased) rocWMMA 2.1.0 for ROCm 7.1.0
+## (Unreleased) rocWMMA 2.2.0 for ROCm 7.2.0
+
+### Added
+
+* Added sample `perf_i8gemm` to demonstrate `int8_t` as matrix input data type
+* Added support for the gfx1150 target
+
+### Changed
+
+* Removed unnecessary const keyword to avoid compiler warnings
+* rocWMMA has been moved into the new rocm-libraries "monorepo" repository (https://github.com/ROCm/rocm-libraries). This repository consolidates a number of separate ROCm libraries and shared components.
+  * The repository migration requires a few changes to the CMake configuration of rocWMMA
+  * The repository migration rquired the GTest dependency to be updated to v1.16.0
+
+### Resolved issues
+
+* Skip invalid test configurations when using 'register file' LDS mapping
+* Ensured transform functions in samples are only available on the device
+
+## rocWMMA 2.1.0 for ROCm 7.1.1
 
 ### Added
 
 * Added more unit tests to increase the code coverage.
-* Added sample `perf_i8gemm` to demonstrate `int8_t` as matrix input data type.
 
 ### Changed
 

CMakeLists.txt

@@ -30,6 +30,7 @@ message(STATUS "CMake version: ${CMAKE_VERSION}")
 if(NOT WIN32)
   if(CMAKE_GENERATOR STREQUAL "Ninja")
     message(STATUS "The CMake generator is Ninja.")
+    set(CMAKE_BUILD_WITH_INSTALL_RPATH ON)
   else()
     message(WARNING "The CMake generator is not Ninja. Ninja is preferred over Make.")
   endif()
@@ -87,7 +88,7 @@ include(ROCMCheckTargetIds)
 include(ROCMClients)
 
 # Versioning via rocm-cmake
-set ( VERSION_STRING "2.1.0" )
+set ( VERSION_STRING "2.2.0" )
 rocm_setup_version( VERSION ${VERSION_STRING} )
 
 # configure a header file to pass the CMake version settings to the source

CONTRIBUTING.md

@@ -10,7 +10,7 @@ External contributions and feedback are welcome for rocWMMA. Please see the foll
 
 ## Issue Discussion ##
 
-Please use the [GitHub Issues](https://github.com/ROCm/rocWMMA/issues) tab to notify us of issues.
+Please use the [GitHub Issues](https://github.com/ROCm/rocm-libraries/issues) tab to notify us of issues.
 
 * Use your best judgement for issue creation. If your issue is already listed, upvote the existing issue and
   comment or post to provide additional details, such as how you reproduced this issue.
@@ -37,13 +37,17 @@ New issues should use the following templates:
             - Expected Behavior:
             - Actual Behavior:
         3. Any additional information:
+        4. Add the `project: rocWMMA` tag
+        5. Specify 'rocWMMA' in the 'ROCm Component' field, if available.
 * **Enhancement Requests**
     + Any proposed enhancements to rocWMMA should be thoroughly documented with the template below.
         1. Description: ***Please be clear and concise***
         2. Value and Motivation
             - Feature and/or Functionalities Enabled:
             - Any Alternatives
         3. Any additional information:
+        4. Add the `project: rocWMMA` tag
+        5. Specify 'rocWMMA' in the 'ROCm Component' field, if available.
 
 ## Acceptance Criteria ##
 
@@ -65,7 +69,7 @@ Contributors wishing to submit new features for rocWMMA should follow the guidel
 - WMMA Porting
     * Developers wishing to implement gap closures with nvcuda::wmma may suggest additional features to do so
 
-All new features and fixes should also tie into the rocWMMA [GitHub Issues](https://github.com/ROCm/rocWMMA/issues) tab:
+All new features and fixes should also tie into the rocWMMA [GitHub Issues](https://github.com/ROCm/rocm-libraries/issues) tab:
 
 - **Enhancements**
     + Any implementations of pre-filed enhancement requests should clearly link to the original issue.
@@ -80,7 +84,7 @@ Exceptions to these criteria will be handled on a case-by-case basis, and should
 
 ## Code Structure ##
 
-The organization of the rocWMMA library is explained in detail in the [Programmers Guide](https://github.com/ROCm/rocWMMA/blob/develop/docs/programmers-guide.rst)
+The organization of the rocWMMA library is explained in detail in the [Programmers Guide](https://github.com/ROCm/rocm-libraries/blob/develop/projects/rocwmma/docs/conceptual/programmers-guide.rst).
 
 ## Coding Style ##
 
@@ -191,7 +195,7 @@ rocWMMA has a set of required deliverables for every pull request that are as fo
     kernels.
 
 2. **API Documentation**:
-    - Any new outward facing rocWMMA API functions must be properly documented and included in the [API Reference Guide](https://github.com/ROCm/rocWMMA/blob/develop/docs/api-reference-guide.rst).
+    - Any new outward facing rocWMMA API functions must be properly documented and included in the [API Reference Guide](https://github.com/ROCm/rocm-libraries/blob/develop/projects/rocwmma/docs/api-reference/api-reference-guide.rst).
 
 3. **Type Support**:
     - All features introduced to rocWMMA must maintain support for the following types:

README.md

@@ -14,7 +14,7 @@ performant GEMMs, DLRM, GEMV and hipRTC integration.
 The test suite includes validation and benchmarking projects that focus on unit testing, GEMMs and DLRM.
 
 > [!NOTE]
-> The published rocWMMA documentation is available at [rocWMMA](https://rocm.docs.amd.com/projects/rocWMMA/en/latest/index.html) in an organized, easy-to-read format, with search and a table of contents. The documentation source files reside in the rocWMMA/docs folder of this repository. As with all ROCm projects, the documentation is open source. For more information, see [Contribute to ROCm documentation](https://rocm.docs.amd.com/en/latest/contribute/contributing.html).
+> The published rocWMMA documentation is available at [rocWMMA](https://rocm.docs.amd.com/projects/rocWMMA/en/latest/index.html) in an organized, easy-to-read format, with search and a table of contents. The documentation source files reside in the `projects/rocwmma/docs` folder of this repository. As with all ROCm projects, the documentation is open source. For more information, see [Contribute to ROCm documentation](https://rocm.docs.amd.com/en/latest/contribute/contributing.html).
 
 
 ## Requirements
@@ -90,7 +90,7 @@ pip3 install -r sphinx/requirements.txt
 python3 -m sphinx -T -E -b html -d _build/doctrees -D language=en . _build/html
 ```
 
-The HTML documentation can be viewed in your browser by opening docs/_build/html/index.html result.
+The HTML documentation can be viewed in your browser by opening the `docs/_build/html/index.html` result.
 
 ### Pdf
 
@@ -118,4 +118,4 @@ The latest official documentation for rocWMMA is available at:
 
 ## Contributing to the rocWMMA Library
 
-Community collaboration is encouraged! If you are considering contributing, please follow the [rocWMMA Contribution Guide](https://github.com/ROCm/rocWMMA/blob/develop/CONTRIBUTING.md) to get started.
+Community collaboration is encouraged! If you are considering contributing, please follow the [rocWMMA Contribution Guide](https://github.com/ROCm/rocm-libraries/blob/develop/projects/rocwmma/CONTRIBUTING.md) to get started.

docs/api-reference/api-reference-guide.rst

@@ -294,31 +294,31 @@ col_major
 
 
 default_schedule
-^^^^^^^^^
+^^^^^^^^^^^^^^^^
 
 .. doxygentypedef:: rocwmma::fragment_scheduler::default_schedule
 
 
 coop_row_major_2d
-^^^^^^^^^
+^^^^^^^^^^^^^^^^^
 
 .. doxygentypedef:: rocwmma::fragment_scheduler::coop_row_major_2d
 
 
 coop_col_major_2d
-^^^^^^^^^
+^^^^^^^^^^^^^^^^^
 
 .. doxygentypedef:: rocwmma::fragment_scheduler::coop_col_major_2d
 
 
 coop_row_slice_2d
-^^^^^^^^^
+^^^^^^^^^^^^^^^^^
 
 .. doxygentypedef:: rocwmma::fragment_scheduler::coop_row_slice_2d
 
 
 coop_col_slice_2d
-^^^^^^^^^
+^^^^^^^^^^^^^^^^^
 
 .. doxygentypedef:: rocwmma::fragment_scheduler::coop_col_slice_2d
 
@@ -379,8 +379,8 @@ rocWMMA transforms API functions
 Sample programs
 ----------------
 
-A sample demonstrating the use of rocWMMA functions ``load_matrix_sync``, ``store_matrix_sync``, ``fill_fragment``, and ``mma_sync`` is available `here <https://github.com/ROCm/rocWMMA/blob/develop/samples/simple_hgemm.cpp>`_.
-For more sample programs, refer to the `samples directory <https://github.com/ROCm/rocWMMA/tree/develop/samples>`_.
+A sample demonstrating the use of rocWMMA functions ``load_matrix_sync``, ``store_matrix_sync``, ``fill_fragment``, and ``mma_sync`` is available `here <https://github.com/ROCm/rocm-libraries/blob/develop/projects/rocwmma/samples/simple_hgemm.cpp>`_.
+For more sample programs, refer to the `samples directory <https://github.com/ROCm/rocm-libraries/tree/develop/projects/rocwmma/samples>`_.
 
 Emulation tests
 ---------------

docs/conceptual/programmers-guide.rst

@@ -17,7 +17,12 @@ Infrastructure
 
 - Doxygen and Sphinx are used to generate the project's documentation.
 - Jenkins is used to automate Continuous Integration (CI) testing, with configurations stored in the ``.jenkins`` folder.
-- rocWMMA is hosted and maintained by AMD on `GitHub  <https://github.com/ROCm/rocWMMA>`_.
+- rocWMMA is hosted and maintained by AMD on `GitHub  <https://github.com/ROCm/rocm-libraries/tree/develop/projects/rocwmma>`_.
+
+  .. note::
+
+    The rocWMMA repository for ROCm 7.1.1 and earlier is located at `<https://github.com/ROCm/rocWMMA>`_.
+
 - The rocWMMA project is organized and configured using ``CMake``, with ``CMakeLists.txt`` files in the root of each directory.
 - ``clang-format`` is used to format C++ code. ``.githooks/install`` ensures that a clang-format pass will run on each committed file.
 - ``GTest`` is used to implement test suite organization and execution.
@@ -31,7 +36,7 @@ hipRTC support
 The HIP runtime compilation (hipRTC) environment enables on-the-fly runtime compilation, loading, and execution of device code on AMD GPUs. The rocWMMA library is compatible with hipRTC, so it can be used for runtime-generated kernels.
 A simple GEMM sample is included to demonstrate compatibility.
 
-For more information, refer to the `HIP API Reference  <https://rocm.docs.amd.com/projects/HIP/en/latest/doxygen/html/index.html>`_
+For more information, refer to the `HIP API Reference  <https://rocm.docs.amd.com/projects/HIP/en/latest/index.html>`_
 
 --------------------------------
 Design concepts
@@ -228,4 +233,4 @@ The ``test`` directory contains the following test code support:
 Contributing
 ^^^^^^^^^^^^
 
-To contribute to the project, see `Contributing to rocWMMA  <https://github.com/ROCm/rocWMMA/blob/develop/CONTRIBUTING.md>`_.
+To contribute to the project, see `Contributing to rocWMMA  <https://github.com/ROCm/rocm-libraries/blob/develop/projects/rocwmma/CONTRIBUTING.md>`_.

docs/index.rst

@@ -12,7 +12,11 @@ rocWMMA is a C++ header library for accelerating mixed-precision matrix multiply
 It leverages specialized GPU matrix cores on the latest AMD discrete GPUs.
 For more information, see :doc:`What is rocWMMA? <./what-is-rocwmma>`
 
-The rocWMMA public repository is located at `<https://github.com/ROCm/rocWMMA>`_.
+The rocWMMA public repository is located at `<https://github.com/ROCm/rocm-libraries/tree/develop/projects/rocwmma>`_.
+
+.. note::
+
+   The rocWMMA repository for ROCm 7.1.1 and earlier is located at `<https://github.com/ROCm/rocWMMA>`_.
 
 .. grid:: 2
   :gutter: 3
@@ -28,12 +32,12 @@ The rocWMMA public repository is located at `<https://github.com/ROCm/rocWMMA>`_
 
   .. grid-item-card:: Examples
 
-    * `Samples <https://github.com/ROCm/rocWMMA/tree/develop/samples>`_ 
+    * `Samples <https://github.com/ROCm/rocm-libraries/tree/develop/projects/rocwmma/samples>`_ 
 
   .. grid-item-card:: API reference
 
     * :doc:`API reference guide <./api-reference/api-reference-guide>`
 
 To contribute to the documentation, see `Contributing to ROCm <https://rocm.docs.amd.com/en/latest/contribute/contributing.html>`_.
 
-You can find licensing information on the `Licensing <https://rocm.docs.amd.com/en/latest/about/license.html>`_ page.
+You can find licensing information on the `Licensing <https://rocm.docs.amd.com/en/latest/about/license.html>`_ page.

docs/install/installation.rst

@@ -111,9 +111,9 @@ The following dependencies are required:
 *  `ROCm <https://github.com/ROCm/ROCm>`_ (Version 6.4 or later)
 *  `CMake <https://cmake.org/>`_ (Version 3.14 or later)
 *  `rocm-cmake <https://github.com/ROCm/rocm-cmake>`_ (Version 0.8.0 or later)
-*  `HIP runtime <https://github.com/ROCm/hip>`_ (Version 6.4.0 or later) (Or the ROCm hip-runtime-amd package)
+*  `HIP runtime <https://github.com/ROCm/rocm-systems/tree/develop/projects/hip>`_ (Version 6.4.0 or later) (Or the ROCm hip-runtime-amd package)
 *  `LLVM OpenMP <https://openmp.llvm.org/>`_ runtime dev package (Version 10.0 or later) (Also available as the ROCm rocm-llvm-dev package)
-*  (Optional, only required to use rocBLAS for validation) `rocBLAS <https://github.com/ROCm/rocBLAS>`_ (Version 4.3.0 for ROCm 6.3 or later) (Or the ROCm rocblas and rocblas-dev packages).
+*  (Optional, only required to use rocBLAS for validation) `rocBLAS <https://github.com/ROCm/rocm-libraries/tree/develop/projects/rocblas>`_ (Version 4.3.0 for ROCm 6.3 or later) (Or the ROCm rocblas and rocblas-dev packages).
 
 .. <!-- spellcheck-enable -->
 
@@ -124,9 +124,13 @@ The following dependencies are required:
 Downloading rocWMMA
 -------------------------------------------
 
-The rocWMMA source code is available from the `rocWMMA GitHub <https://github.com/ROCm/rocWMMA>`_.
+The rocWMMA source code is available from the `rocWMMA GitHub <https://github.com/ROCm/rocm-libraries/tree/develop/projects/rocwmma>`_.
 ROCm version 6.4 or later is required.
 
+.. note::
+
+   The rocWMMA repository for ROCm 7.1.1 and earlier is located at `<https://github.com/ROCm/rocWMMA>`_.
+
 To verify the ROCm version installed on an Ubuntu system, use this command:
 
 .. code-block:: shell
@@ -146,10 +150,22 @@ The rocWMMA GitHub repository has branches with names like ``rocm-major.minor.x`
 where ``major`` and ``minor`` are the same as for the ROCm version.
 To download rocWMMA on ROCm version ``x.y``, use this command:
 
-.. code-block:: shell
+.. tab-set::
+    .. tab-item:: ROCm 7.2.0 and later
+
+        .. code-block:: shell
+
+           git clone -b release/rocm-rel-x.y https://github.com/ROCm/rocm-libraries.git
+           cd projects/rocwmma
+
+        Alternatively, you can use ``sparse-checkout`` to clone only the rocWMMA project from the ``rocm-libraries`` monorepo. For more information, see `Contributing to the ROCm Libraries <https://github.com/ROCm/rocm-libraries/blob/develop/CONTRIBUTING.md#option-b-clone-the-monorepo-with-sparse-checkout>`_.
+
+    .. tab-item:: ROCm 7.1.1 and prior
+
+        .. code-block:: shell
 
-   git clone -b release/rocm-rel-x.y https://github.com/ROCm/rocWMMA.git
-   cd rocWMMA
+           git clone -b release/rocm-rel-x.y https://github.com/ROCm/rocWMMA.git
+           cd rocWMMA
 
 Replace ``x.y`` in the above command with the version of ROCm installed on your machine.
 For example, if you have ROCm 6.0 installed, then replace ``release/rocm-rel-x.y`` with ``release/rocm-rel-6.0``.

docs/sphinx/_toc.yml.in

@@ -21,7 +21,7 @@ subtrees:
 
 - caption: Examples
   entries:
-  - url: https://github.com/ROCm/rocWMMA/tree/develop/samples
+  - url: https://github.com/ROCm/rocm-libraries/tree/develop/projects/rocwmma/samples
     title: Samples
 
 - caption: API reference

samples/CMakeLists.txt

@@ -42,7 +42,8 @@ function(add_rocwmma_sample TEST_TARGET TEST_SOURCE)
   add_executable(${TEST_TARGET} ${TEST_SOURCE})
   target_link_libraries(${TEST_TARGET} OpenMP::OpenMP_CXX "-L${HIP_CLANG_ROOT}/lib")
   if(UNIX)
-    target_link_libraries(${TEST_TARGET} "-Wl,-rpath=$ORIGIN/../llvm/lib"  "-fno-rtlib-add-rpath")
+    set_target_properties(${TEST_TARGET} PROPERTIES INSTALL_RPATH "\$ORIGIN/../llvm/lib")
+    target_link_libraries(${TEST_TARGET} "-fno-rtlib-add-rpath")
   endif()
   target_link_libraries(${TEST_TARGET} rocwmma hiprtc::hiprtc)
   target_include_directories(${TEST_TARGET} PRIVATE