Expose resolved GEMM schedules in library descriptions

media/docs/cpp/code_organization.md

@@ -121,6 +121,48 @@ python/
 When CMake is executed, the CUTLASS Instance Library generator scripts are executed to construct a set of
 instantiations in `build/tools/library/generated/`.
 
+### Runtime GEMM schedule metadata
+
+The instance library exposes the selected CUTLASS 3.x mainloop and epilogue strategies
+through `GemmDescription`. These fields describe the resolved collective policies,
+including when the kernel was constructed with `KernelScheduleAuto` or
+`EpilogueScheduleAuto`:
+
+| Field | Meaning |
+| --- | --- |
+| `mainloop_schedule` | Multistage, TMA, warp-specialized, ping-pong, cooperative, or Blackwell 1-SM / 2-SM execution |
+| `mainloop_load` | `cp.async`, TMA, or mixed TMA / `cp.async` operand loads |
+| `epilogue_schedule` | Direct stores without shared memory, or a TMA epilogue |
+| `accumulation_kind` | Default accumulation or the opt-in Hopper FP8 fast-accumulation policy |
+
+For example, a client can select cooperative Hopper kernels without parsing their names:
+
+```cpp
+using namespace cutlass::library;
+auto const& description = static_cast<GemmDescription const&>(operation.description());
+if (description.mainloop_schedule == MainloopScheduleKind::kWarpSpecializedCooperative &&
+    description.accumulation_kind == AccumulationKind::kDefault) {
+  // Consider this kernel for the application's dispatch table.
+}
+```
+
+Use the description type matching the operation kind. `BlockScaledGemmDescription`
+and `BlockwiseGemmDescription` expose the same fields; grouped operations expose them
+through `GroupedGemmDescription::gemm`. The 1-SM / 2-SM distinction comes from the
+selected MMA atom's CTA group, rather than the cluster size.
+
+Unrecognized custom policies and descriptions populated by older operation wrappers
+retain `kUnknown` fields. Clients should handle `kUnknown` explicitly. In particular,
+the pre-Hopper `KernelMultistage` tag alone does not distinguish synchronous loads
+from `cp.async`, so its load kind is unknown. Existing constructor arguments are
+unchanged; rebuild the instance library and its clients together when updating the
+public description structures.
+
+The `cutlass_test_unit_library` target checks this metadata using real collective
+builders and operation descriptions. It requires a CUDA toolkit for compilation but
+runs on the host without launching kernels or requiring a GPU. Architecture-specific
+cases are enabled according to the toolkit's support for the respective builders.
+
 ### CUTLASS Profiler
 
 The CUTLASS Profiler is designed to load the CUTLASS Instance Library and execute all operations contained therein.

test/unit/CMakeLists.txt

@@ -173,6 +173,7 @@ add_custom_target(cutlass_test_unit)
 add_custom_target(test_unit)
 
 set(SUBDIRS
+  library
   core
   cute
   gemm

test/unit/library/CMakeLists.txt

@@ -0,0 +1,40 @@
+# Copyright (c) 2017 - 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+# SPDX-License-Identifier: BSD-3-Clause
+#
+# Redistribution and use in source and binary forms, with or without
+# modification, are permitted provided that the following conditions are met:
+#
+# 1. Redistributions of source code must retain the above copyright notice, this
+# list of conditions and the following disclaimer.
+#
+# 2. Redistributions in binary form must reproduce the above copyright notice,
+# this list of conditions and the following disclaimer in the documentation
+# and/or other materials provided with the distribution.
+#
+# 3. Neither the name of the copyright holder nor the names of its
+# contributors may be used to endorse or promote products derived from
+# this software without specific prior written permission.
+#
+# THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
+# AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
+# IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
+# DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE
+# FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL
+# DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR
+# SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
+# CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY,
+# OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
+# OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+
+
+cutlass_test_unit_add_executable(
+  cutlass_test_unit_library
+  WITHOUT_CUDA
+  gemm_description.cu
+  EXTRA_INCLUDE_DIRS
+    ${CUTLASS_SOURCE_DIR}/tools/library/include
+    ${CUTLASS_SOURCE_DIR}/tools/library/src
+)
+
+target_link_libraries(cutlass_test_unit_library PRIVATE
+  CUTLASS cutlass_tools_util_includes GTest::gtest_main cuda_driver)