hipblaslt: add support for gfx1250-strict build

projects/hipblaslt/CMakeLists.txt

@@ -592,17 +592,6 @@ if(HIPBLASLT_ENABLE_DEVICE OR HIPBLASLT_ENABLE_ROCROLLER)
         DESTINATION ${HIPBLASLT_TENSILE_LIBRARY_DIR}
         COMPONENT runtime
     )
-    # The gfx1250 v0 GEMM tree builds into its own directory (see
-    # device-library/CMakeLists.txt) but installs as one more subdir of the same
-    # library/ root. The property is only set when that tree is built.
-    get_property(_gfx1250v0_library_dir GLOBAL PROPERTY HIPBLASLT_GFX1250V0_LIBRARY_DIR)
-    if(_gfx1250v0_library_dir)
-        rocm_install(
-            DIRECTORY "${_gfx1250v0_library_dir}"
-            DESTINATION "${HIPBLASLT_TENSILE_LIBRARY_DIR}/library"
-            COMPONENT runtime
-        )
-    endif()
 endif()
 
 set(CPACK_RESOURCE_FILE_LICENSE "${CMAKE_CURRENT_SOURCE_DIR}/LICENSE.md")
@@ -766,11 +755,12 @@ if(HIPBLASLT_INSTALL_TENSILELITE_TEST_ARTIFACTS AND NOT WIN32)
         COMPONENT tests
     )
 
-    # Imported directly by unit tests, so these generators must ship at the
+    # Imported or executed directly by unit tests, so these must ship at the
     # tensilelite root, not under scripts/.
     install(
         FILES
             "${_tensilelite_src}/AMaxGenerator.py"
+            "${_tensilelite_src}/SoftmaxGenerator.py"
             "${_tensilelite_src}/LayerNormGenerator.py"
         DESTINATION "${_tensilelite_install_dest}"
         COMPONENT tests

projects/hipblaslt/README.md

@@ -277,7 +277,6 @@ inv build --architecture gfx1100 --clean
 * `TENSILELITE_EXPERIMENTAL` Process experimental logic files (default: see below)
 * `HIPBLASLT_LIBLOGIC_PATH` Path to library logic files (will use 'library' if unset) (default: `Off`)
 * `HIPBLASLT_TENSILE_LIBPATH` Path to output the device gemm libraries (default: `build/Tensile`)
-* `HIPBLASLT_ASIC_REVISION` gfx1250 ASIC revision to build for, `v0` or `v1`; `invoke build` probes the local GPU and sets it when unset (default: empty, treated as `v1`)
 
 > [!NOTE]
 > To determine defaults for the `TensileCreateLibrary` command generated when building the device

projects/hipblaslt/clients/tests/data/known_bugs.yaml

@@ -10,15 +10,18 @@
 
 Known bugs:
   # ROCM-1545: TF32 matmul with INF inputs returns NaN on gfx950 (MI355)
-  - { name: matmul_tf32_inf_ROCM1545, function: matmul, initialization: inf, known_bug_platforms: "gfx950, gfx1250" }
+  # gfx1250-strict is listed alongside gfx1250 because the match is an exact
+  # token comparison against the name the runtime reports, and a v0 part reports
+  # the revision; the defect is in silicon the revision shares.
+  - { name: matmul_tf32_inf_ROCM1545, function: matmul, initialization: inf, known_bug_platforms: "gfx950, gfx1250, gfx1250-strict" }
   # ROCM-1545 (same root cause): the special-FP sweep hits it on the TF32 (xf32)
   # compute type when the input is +Inf or -Inf -- the GPU returns -nan where the
   # CPU reference is +Inf, failing check_special_value_consistency(). Only the xf32
   # Inf/-Inf cases fail; the f16/bf16/f32/f64 Inf cases and the xf32 NaN case pass,
   # so scope the quarantine to compute_type xf32_r + inf/neg_inf. REMOVE once the
   # TF32 Inf fix lands so these re-activate as an enforced consistency gate.
-  - { name: matmul_special_fp_64x64, function: matmul, compute_type: c_xf32_r, initialization: inf, known_bug_platforms: "gfx950, gfx1250" }
-  - { name: matmul_special_fp_64x64, function: matmul, compute_type: c_xf32_r, initialization: neg_inf, known_bug_platforms: "gfx950, gfx1250" }
+  - { name: matmul_special_fp_64x64, function: matmul, compute_type: c_xf32_r, initialization: inf, known_bug_platforms: "gfx950, gfx1250, gfx1250-strict" }
+  - { name: matmul_special_fp_64x64, function: matmul, compute_type: c_xf32_r, initialization: neg_inf, known_bug_platforms: "gfx950, gfx1250, gfx1250-strict" }
   # ROCM-26455: bf16 matmul output-store address overflow -- an output element
   # offset > 2^31 (from an oversized output leading dimension ldc/ldd) wraps in the
   # 32-bit global-write address path and produces silent wrong results. The gate

projects/hipblaslt/clients/tests/src/arch_revision_gtest.cpp

@@ -12,33 +12,37 @@
 
 namespace
 {
-    TEST(ArchRevisionSmoke, Gfx1250Revision0IsTheV0Subtree)
+    TEST(arch_revision_smoke, gfx1250_revision0_is_the_v0_subtree)
     {
-        // The one case that diverges: pre-production v0 loads its own tree.
+        // An A0 part in non-strict mode reports gfx1250 and loads its own tree.
         EXPECT_EQ(rocblaslt_revisioned_arch_name("gfx1250", 0), "gfx1250v0");
     }
 
-    TEST(ArchRevisionSmoke, Gfx1250V1RevisionKeepsTheBaseName)
+    TEST(arch_revision_smoke, gfx1250_strict_keeps_its_reported_name)
     {
-        // v1 is revision 1; it must map to the plain gfx1250 tree.
+        // An A0 part in strict mode reports gfx1250-strict (asicRevision is still 0)
+        // and keeps loading library/gfx1250-strict/.
+        EXPECT_EQ(rocblaslt_revisioned_arch_name("gfx1250-strict", 0), "gfx1250-strict");
+    }
+
+    TEST(arch_revision_smoke, gfx1250_nonzero_revision_keeps_the_base_name)
+    {
+        // Assumes B0 reports a non-zero revision; it maps to the plain gfx1250 tree.
         EXPECT_EQ(rocblaslt_revisioned_arch_name("gfx1250", 1), "gfx1250");
+        EXPECT_EQ(rocblaslt_revisioned_arch_name("gfx1250", 2), "gfx1250");
     }
 
-    TEST(ArchRevisionSmoke, Gfx1250UnknownOrFutureRevisionDefaultsToV1)
+    TEST(arch_revision_smoke, gfx1250_unknown_revision_keeps_the_base_name)
     {
-        // -1 is what HIP reports when it is too old to expose the field; any
-        // unseen value must default to the v1 tree rather than invent a
-        // subtree that was never built.
+        // Only a revision known to be A0 selects the A0 tree.
         EXPECT_EQ(rocblaslt_revisioned_arch_name("gfx1250", -1), "gfx1250");
-        EXPECT_EQ(rocblaslt_revisioned_arch_name("gfx1250", 2), "gfx1250");
     }
 
-    TEST(ArchRevisionSmoke, OtherArchesAreUnaffectedByRevision)
+    TEST(arch_revision_smoke, other_arches_are_unaffected_by_revision)
     {
-        // Only gfx1250 has a revision split; every other arch is returned
-        // unchanged regardless of what asicRevision happens to be.
         EXPECT_EQ(rocblaslt_revisioned_arch_name("gfx942", 0), "gfx942");
         EXPECT_EQ(rocblaslt_revisioned_arch_name("gfx950", 0), "gfx950");
         EXPECT_EQ(rocblaslt_revisioned_arch_name("gfx1250v0", 0), "gfx1250v0");
+        EXPECT_EQ(rocblaslt_revisioned_arch_name("gfx1250-strict", 1), "gfx1250-strict");
     }
 }

projects/hipblaslt/clients/tests/src/row_uniformity_gtest.cpp

@@ -37,6 +37,8 @@
 #include <origami/hardware.hpp>
 #include <origami/streamk.hpp>
 
+#include "../../../library/src/amd_detail/rocblaslt/src/include/rocblaslt_arch_revision.hpp"
+
 #ifdef _WIN32
 #ifndef WIN32_LEAN_AND_MEAN
 #define WIN32_LEAN_AND_MEAN
@@ -169,14 +171,23 @@ namespace
             // that never appear in the library filename.
             const std::string arch      = gpuArchName();
             const std::string processor = arch.substr(0, arch.find(':'));
+            // The subtree the runtime loads, which for an A0 part reporting
+            // gfx1250 is not the one named for processor.
+            hipDeviceProp_t props{};
+            int             device = 0;
+            const int       rev    = hipGetDevice(&device) == hipSuccess
+                                    && hipGetDeviceProperties(&props, device) == hipSuccess
+                                ? props.asicRevision
+                                : -1;
+            const std::string libArch = rocblaslt_revisioned_arch_name(processor, rev);
 
             std::string tried;
             for(const auto& root : libraryRootCandidates())
             {
                 // Always the logical single-extension name: the loader resolves
                 // the shipped ".dat.zlib" by appending the suffix itself.
                 const std::filesystem::path logical
-                    = root / processor / ("TensileLibrary_lazy_" + processor + ".dat");
+                    = root / libArch / ("TensileLibrary_lazy_" + processor + ".dat");
                 tried += (tried.empty() ? "" : ", ") + logical.string();
 
                 if(!std::filesystem::exists(logical)

projects/hipblaslt/cmake/hipblaslt_codegen.cmake

@@ -130,17 +130,21 @@ function(hipblaslt_detect_sanitizer_runtime out_options out_lib_dirs)
     set(${out_lib_dirs} "${_lib_dirs}" PARENT_SCOPE)
 endfunction()
 
+# When ON, a build that requests gfx1250 also produces library/gfx1250v0/
+# (gfx1250-strict's kernels built for gfx1250). Never add gfx1250v0 to GPU_TARGETS.
+option(HIPBLASLT_BUILD_GFX1250V0 "Build library/gfx1250v0/ alongside gfx1250 for A0 parts reporting gfx1250." ON)
+
 function(create_device_library)
     set(_opts HOST_ASAN HOST_TSAN)
     set(_one
         TARGET LOGIC_PATH OUTPUT_DIR CODEGEN_ROOT PYTHON_EXECUTABLE CXX_COMPILER OFFLOAD_BUNDLER JOBS LOGIC_FILTER
         ASAN YAML_FORMAT NO_COMPRESS EXPERIMENTAL GEMM_A2A_FUSION LAZY_LOAD ASM_COMMENTS
-        KEEP_BUILD_TMP ASM_DEBUG REQUIRE_GFX1250V0_OVERLAY)
+        KEEP_BUILD_TMP ASM_DEBUG)
     set(_multi ARCHES)
     cmake_parse_arguments(_cdl "${_opts}" "${_one}" "${_multi}" ${ARGN})
 
     if(_cdl_UNPARSED_ARGUMENTS)
-        message(FATAL_ERROR "create_device_library: unexpected arguments: ${_cdl_UNPARSED_ARGUMENTS} (permitted options: HOST_ASAN, HOST_TSAN; single-value keywords: TARGET, LOGIC_PATH, OUTPUT_DIR, CODEGEN_ROOT, PYTHON_EXECUTABLE, CXX_COMPILER, OFFLOAD_BUNDLER, JOBS, LOGIC_FILTER, ASAN, YAML_FORMAT, NO_COMPRESS, EXPERIMENTAL, GEMM_A2A_FUSION, LAZY_LOAD, ASM_COMMENTS, KEEP_BUILD_TMP, ASM_DEBUG, REQUIRE_GFX1250V0_OVERLAY; multi-value keyword: ARCHES)")
+        message(FATAL_ERROR "create_device_library: unexpected arguments: ${_cdl_UNPARSED_ARGUMENTS} (permitted options: HOST_ASAN, HOST_TSAN; single-value keywords: TARGET, LOGIC_PATH, OUTPUT_DIR, CODEGEN_ROOT, PYTHON_EXECUTABLE, CXX_COMPILER, OFFLOAD_BUNDLER, JOBS, LOGIC_FILTER, ASAN, YAML_FORMAT, NO_COMPRESS, EXPERIMENTAL, GEMM_A2A_FUSION, LAZY_LOAD, ASM_COMMENTS, KEEP_BUILD_TMP, ASM_DEBUG; multi-value keyword: ARCHES)")
     endif()
     if(NOT _cdl_LOGIC_PATH)
         message(FATAL_ERROR "create_device_library: LOGIC_PATH is required")
@@ -293,17 +297,23 @@ function(create_device_library)
     if(NOT _cdl_ASM_COMMENTS)
         list(APPEND _opts_list "--disable-asm-comments")
     endif()
+    set(_logic_arches ${_cdl_ARCHES})
+    if(HIPBLASLT_BUILD_GFX1250V0)
+        list(APPEND _opts_list "--gfx1250v0")
+        # gfx1250v0 is built from gfx1250-strict's logic, so validate that too.
+        if("gfx1250" IN_LIST _cdl_ARCHES AND NOT "gfx1250-strict" IN_LIST _cdl_ARCHES)
+            list(APPEND _logic_arches "gfx1250-strict")
+        endif()
+    endif()
+    list(JOIN _logic_arches "$<SEMICOLON>" _logic_arches_semi)
 
     set(_tensile_logic_args
         "${_cdl_LOGIC_PATH}"
         --architecture
-        "${_arches_semi}"
+        "${_logic_arches_semi}"
         --use-bundled-known-bugs
         --check-all
     )
-    if(_cdl_REQUIRE_GFX1250V0_OVERLAY)
-        list(APPEND _tensile_logic_args --require-gfx1250v0-overlay)
-    endif()
     set(_codegen_dependencies "${_known_bugs_resource}")
     if(TARGET _rocisa)
         list(APPEND _codegen_dependencies _rocisa)
@@ -323,6 +333,11 @@ function(create_device_library)
         USES_TERMINAL
     )
 
+    # All architectures go to this one command, a stepping included. Covering a
+    # stepping and the architecture it steps from takes more than one run, since
+    # a run names its target by the ISA and the two spell one ISA -- but that
+    # split is TensileCreateLibrary's: it cannot happen at configure time, when
+    # Tensile is not yet importable.
     set(_output_stamp "${CMAKE_CURRENT_BINARY_DIR}/${_cdl_TARGET}.stamp")
     set(_tcl_command
         ${_python_command} -m Tensile.TensileCreateLibrary

projects/hipblaslt/cmake/tensilelite_supported_architectures.cmake

@@ -18,6 +18,7 @@ set(SUPPORTED_ARCHITECTURES
     "gfx1200"
     "gfx1201"
     "gfx1250"
+    "gfx1250-strict"
     "gfx908:xnack+"
     "gfx908:xnack-"
     "gfx90a:xnack+"
@@ -44,6 +45,7 @@ if(HIPBLASLT_ENABLE_ASAN OR THEROCK_SANITIZER STREQUAL "ASAN" OR THEROCK_SANITIZ
         "gfx942:xnack+"
         "gfx950:xnack+"
         "gfx1250"
+        "gfx1250-strict"
         )
 else()
     # For non address sanitizer builds, "all" is non-xnack architectures.
@@ -61,7 +63,8 @@ else()
         "gfx1153"
         "gfx1200"
         "gfx1201"
-        "gfx1250")
+        "gfx1250"
+        "gfx1250-strict")
 endif()
 
 # Validate that all BASE_ARCHITECTURES are in the SUPPORTED_ARCHITECTURES list

projects/hipblaslt/device-library/CMakeLists.txt

@@ -25,22 +25,6 @@ mark_as_advanced(
 set(HIPBLASLT_LIBLOGIC_PATH "" CACHE STRING "Path to library logic files (will use 'library' if unset).")
 set(HIPBLASLT_TENSILE_LIBPATH "${hipblaslt_BINARY_DIR}/Tensile" CACHE STRING "Path to output the device gemm libraries.")
 
-set(HIPBLASLT_ASIC_REVISION "" CACHE STRING "gfx1250 ASIC revisions to build ('' = both v0 and v1, 'v1' = v1 only, 'v0' = v0 only).")
-
-# This policy belongs to the caller that owns the logic corpus, not to the
-# gfx1250v0 architecture name itself. hipSPARSELt reuses this device-library
-# CMakeLists with a corpus that deliberately has no v0 overlay, while
-# hipBLASLt's own corpus must retain one. The embedding hipSPARSELt CMakeLists
-# overrides this to OFF before adding hipBLASLt as a subdirectory.
-if(NOT DEFINED HIPBLASLT_REQUIRE_GFX1250V0_OVERLAY)
-    set(HIPBLASLT_REQUIRE_GFX1250V0_OVERLAY ON)
-endif()
-
-# Reject an unrecognized value rather than silently building the wrong revisions.
-if(NOT HIPBLASLT_ASIC_REVISION MATCHES "^(v0|v1)?$")
-    message(FATAL_ERROR "HIPBLASLT_ASIC_REVISION must be '', 'v0' or 'v1', got '${HIPBLASLT_ASIC_REVISION}'")
-endif()
-
 if(NOT HIPBLASLT_LIBLOGIC_PATH)
     set(HIPBLASLT_LIBLOGIC_PATH "${hipblaslt_SOURCE_DIR}/library")
 endif()
@@ -71,63 +55,15 @@ if(HIPBLASLT_ENABLE_GEMM_A2A_FUSION)
     list(APPEND _codegen_product_args GEMM_A2A_FUSION "${HIPBLASLT_ENABLE_GEMM_A2A_FUSION}")
 endif()
 
-# gfx1250's two ASIC revisions (v0/v1) need different tuned kernels; the runtime
-# picks library/gfx1250/ (v1) vs library/gfx1250v0/ by asicRevision. The default
-# builds both (v1 here, v0 by the second command below); HIPBLASLT_ASIC_REVISION
-# only prunes which. GPU_TARGETS is untouched (it also feeds extops/transform/
-# validation on the base name); only the bare gfx1250 reaches here.
-set(_main_arches ${GPU_TARGETS})
-set(_build_gfx1250v0 OFF)
-set(_main_requires_gfx1250v0_overlay OFF)
-if("gfx1250" IN_LIST GPU_TARGETS)
-    if(HIPBLASLT_ASIC_REVISION STREQUAL "v0")
-        list(TRANSFORM _main_arches REPLACE "^gfx1250$" "gfx1250v0")  # v0 in place of v1
-        set(_main_requires_gfx1250v0_overlay ${HIPBLASLT_REQUIRE_GFX1250V0_OVERLAY})
-        message(STATUS "hipBLASLt: gfx1250 ASIC revision v0 only")
-    elseif(HIPBLASLT_ASIC_REVISION STREQUAL "v1")
-        message(STATUS "hipBLASLt: gfx1250 ASIC revision v1 only")
-    else()
-        set(_build_gfx1250v0 ON)  # both
-        message(STATUS "hipBLASLt: gfx1250 ASIC revisions v0 and v1 (runtime selects by asicRevision)")
-    endif()
-elseif(HIPBLASLT_ASIC_REVISION STREQUAL "v0")
-    message(WARNING "HIPBLASLT_ASIC_REVISION=v0 but GPU_TARGETS names no gfx1250 "
-                    "(${GPU_TARGETS}); nothing in this build is for the v0 revision.")
-endif()
-
 create_device_library(
+    TARGET tensilelite-device-libraries
     LOGIC_PATH "${HIPBLASLT_LIBLOGIC_PATH}"
     OUTPUT_DIR "${HIPBLASLT_TENSILE_LIBPATH}"
-    ARCHES ${_main_arches}
-    REQUIRE_GFX1250V0_OVERLAY ${_main_requires_gfx1250v0_overlay}
+    ARCHES ${GPU_TARGETS}
     ${_codegen_host_sanitizer_args}
     ${_codegen_product_args}
 )
 
-if(_build_gfx1250v0)
-    # Needs its own OUTPUT_DIR: the two runs share a compiler target and each
-    # writes Kernels.cpp + a build_tmp tree under its output dir (reusing a .o
-    # that is already there), so a shared dir would serialize them or cross-reuse
-    # objects. Installed by the top-level CMakeLists.txt via the property below.
-    set(_gfx1250v0_output_dir "${HIPBLASLT_TENSILE_LIBPATH}-gfx1250v0")
-    file(MAKE_DIRECTORY "${_gfx1250v0_output_dir}/library/gfx1250v0")
-    set_property(GLOBAL PROPERTY HIPBLASLT_GFX1250V0_LIBRARY_DIR
-                 "${_gfx1250v0_output_dir}/library/gfx1250v0")
-
-    create_device_library(
-        TARGET tensilelite-device-libraries-gfx1250v0
-        LOGIC_PATH "${HIPBLASLT_LIBLOGIC_PATH}"
-        OUTPUT_DIR "${_gfx1250v0_output_dir}"
-        ARCHES gfx1250v0
-        # hipBLASLt owns the gfx1250/gfx1250v0 split for its own corpus, but
-        # hipSPARSELt reuses this CMakeLists with a no-split corpus. Let the
-        # embedding caller choose whether this target requires the directory.
-        REQUIRE_GFX1250V0_OVERLAY ${HIPBLASLT_REQUIRE_GFX1250V0_OVERLAY}
-        ${_codegen_host_sanitizer_args}
-        ${_codegen_product_args}
-    )
-endif()
-
 if(HIPBLASLT_ENABLE_EXTOPS)
     add_subdirectory(extops)
 endif()

projects/hipblaslt/device-library/extops/CMakeLists.txt

@@ -3,7 +3,10 @@
 
 set(ops_dir "${CMAKE_CURRENT_SOURCE_DIR}/../../tensilelite")
 set(output_dir "${CMAKE_CURRENT_BINARY_DIR}")
-string(REGEX MATCHALL "gfx[a-z0-9]+" archs "${GPU_TARGETS}")
+# Keeps a "-strict" suffix while still dropping ":xnack+" (re-added below by
+# tensilelite_offload_target). Truncating it would collapse gfx1250-strict onto
+# gfx1250, leaving gfx1250-strict with code objects it cannot load.
+string(REGEX MATCHALL "gfx[a-z0-9]+(-[a-z0-9]+)*" archs "${GPU_TARGETS}")
 list(REMOVE_DUPLICATES archs)
 message(STATUS "Architectures to process: ${archs}")
 
@@ -17,13 +20,15 @@ foreach(arch IN LISTS archs)
         set(extop_xnack_args --xnack)
     endif()
     set(WAVEFRONT "-mwavefrontsize64")
-    if(arch MATCHES "^gfx1[12][0-9][0-9]$")
+    if(arch MATCHES "^gfx1[12][0-9][0-9](-[a-z0-9]+)*$")
         set(WAVEFRONT "-mno-wavefrontsize64")
     endif()
     # Enable true16 on all gfx11*/gfx12* (NoSDWA). "SHELL:" keeps the flag group
-    # intact so CMake de-duplication does not drop the paired +real-true16.
+    # intact so CMake de-duplication does not drop the paired +real-true16. The
+    # stepping suffix is matched like it is for the wavefront size above: true16
+    # belongs to the ISA, which a stepping shares with the architecture it steps.
     set(TRUE16_FEATURE "")
-    if(arch MATCHES "^gfx1[12][0-9][0-9]$")
+    if(arch MATCHES "^gfx1[12][0-9][0-9](-[a-z0-9]+)*$")
         set(TRUE16_FEATURE "SHELL:-Xclangas -target-feature -Xclangas +real-true16")
     endif()
     add_library(extop-obj-${arch} OBJECT)

projects/hipblaslt/device-library/matrix-transform/CMakeLists.txt

@@ -23,7 +23,10 @@
 
 # Need to support Release/Debug
 
-string(REGEX MATCHALL "gfx[a-z0-9]+" archs "${GPU_TARGETS}")
+# Keeps a "-strict" suffix while still dropping ":xnack+" (re-added below by
+# tensilelite_offload_target). Truncating it would collapse gfx1250-strict onto
+# gfx1250, leaving gfx1250-strict with code objects it cannot load.
+string(REGEX MATCHALL "gfx[a-z0-9]+(-[a-z0-9]+)*" archs "${GPU_TARGETS}")
 list(REMOVE_DUPLICATES archs)
 
 set(matrix_transform_cpp "${CMAKE_CURRENT_SOURCE_DIR}/matrix_transform.cpp")

projects/hipblaslt/install.sh

@@ -70,7 +70,6 @@ while [[ $# -gt 0 ]]; do
             echo "  -g, --debug               Debug build"
             echo "  -k, --relwithdebinfo      RelWithDebInfo build"
             echo "  -a, --architecture <arch>  GPU target(s)"
-            echo "  --asic-revision <rev>  build only one gfx1250 ASIC-revision tree, 'v0' or 'v1' (default: build both; runtime selects by asicRevision)"
             echo "  -n, --client-only         Build without Tensile"
             echo "  -j <N>                    Parallel jobs (caps both cmake and Tensile kernel generation)"
             echo "  -l, --logic <path>        Tensile logic path"
@@ -92,7 +91,6 @@ while [[ $# -gt 0 ]]; do
         -g|--debug)         INVOKE_ARGS="$INVOKE_ARGS --debug" ;;
         -k|--relwithdebinfo) INVOKE_ARGS="$INVOKE_ARGS --relwithdebinfo" ;;
         -a|--architecture)  shift; INVOKE_ARGS="$INVOKE_ARGS --architecture='$1'" ;;
-        --asic-revision) shift; INVOKE_ARGS="$INVOKE_ARGS --asic-revision='$1'" ;;
         -n|--client-only)   INVOKE_ARGS="$INVOKE_ARGS --no-tensile" ;;
         -j)                 shift; INVOKE_ARGS="$INVOKE_ARGS --jobs $1 --tensile-threads $1" ;;
         -l|--logic)         shift; INVOKE_ARGS="$INVOKE_ARGS --tensile-logic=$1" ;;