Fix testSolutionStructsUtilities Unit Test

projects/hipblaslt/clients/CMakeLists.txt

@@ -28,19 +28,21 @@ target_link_libraries(hipblaslt-clients-common
         hip::device
 )
 
-if(HIPBLASLT_ENABLE_ROCROLLER)
-    if(NOT ROCM_LIBS_SUPERBUILD)
-        if(HIPBLASLT_ENABLE_THEROCK)
-            find_package(mxDataGenerator REQUIRED)
-        else()
-            add_subdirectory("${CMAKE_CURRENT_SOURCE_DIR}/../../../shared/mxdatagenerator" "${CMAKE_CURRENT_BINARY_DIR}/mxdatagenerator")
-        endif()
+if(NOT ROCM_LIBS_SUPERBUILD)
+    if(HIPBLASLT_ENABLE_THEROCK)
+        find_package(mxDataGenerator REQUIRED)
+    else()
+        add_subdirectory("${CMAKE_CURRENT_SOURCE_DIR}/../../../shared/mxdatagenerator" "${CMAKE_CURRENT_BINARY_DIR}/mxdatagenerator")
     endif()
+endif()
+target_compile_features(hipblaslt-clients-common PRIVATE cxx_std_20)
+target_link_libraries(hipblaslt-clients-common PRIVATE roc::mxDataGenerator)
+
+if(HIPBLASLT_ENABLE_ROCROLLER)
     target_compile_definitions(hipblaslt-clients-common PRIVATE HIPBLASLT_USE_ROCROLLER)
-    target_link_libraries(hipblaslt-clients-common PRIVATE roc::mxDataGenerator)
-    target_compile_features(hipblaslt-clients-common PRIVATE cxx_std_20)
 endif()
 
+
 if(HIPBLASLT_ENABLE_ASAN)
     hipblaslt_target_configure_sanitizers(hipblaslt-clients-common PUBLIC)
 endif()

projects/hipblaslt/clients/common/include/mxDataGen.hpp

@@ -26,10 +26,10 @@
 
 #pragma once
 
-#include <hipblaslt/hipblaslt.h>
-#include <stdint.h>
+#include <hip/hip_runtime.h>
+
+#include <vector>
 
-#ifdef HIPBLASLT_USE_ROCROLLER
 std::vector<float> generateMXInput(hipDataType                dataType,
                                    void*                      data,
                                    void*                      scale,
@@ -45,4 +45,3 @@ std::vector<float> generateMXInput(hipDataType                dataType,
                                    std::string_view const     initMethod = "Bounded",
                                    float                      min_val    = -1.0f,
                                    float                      max_val    = 1.0f);
-#endif

projects/hipblaslt/clients/common/include/testing_matmul.hpp

@@ -1918,7 +1918,6 @@ void testing_matmul_with_bias(const Arguments& arg,
 
         hipblaslt_seedrand();
 
-#ifdef HIPBLASLT_USE_ROCROLLER
         if(isBlockScaling(arg.scaleA))
         {
             if(arg.initialization != hipblaslt_initialization::hpl
@@ -1963,7 +1962,6 @@ void testing_matmul_with_bias(const Arguments& arg,
         }
         else
         {
-#endif
             hipblaslt_init_device(ABC_dims::A,
                                   arg.initialization,
                                   alpha_isnan_type(arg, Talpha),
@@ -1975,7 +1973,6 @@ void testing_matmul_with_bias(const Arguments& arg,
                                   (do_swizzle_a && stride_a[i] != 0) ? A_row[i] * A_col[i]
                                                                      : stride_a[i],
                                   num_batches[i]);
-#ifdef HIPBLASLT_USE_ROCROLLER
         }
         if(isBlockScaling(arg.scaleB))
         {
@@ -2019,7 +2016,6 @@ void testing_matmul_with_bias(const Arguments& arg,
         }
         else
         {
-#endif
             hipblaslt_init_device(ABC_dims::B,
                                   arg.initialization,
                                   alpha_isnan_type(arg, Talpha),
@@ -2031,9 +2027,7 @@ void testing_matmul_with_bias(const Arguments& arg,
                                   (do_swizzle_b && stride_b[i] != 0) ? B_row[i] * B_col[i]
                                                                      : stride_b[i],
                                   num_batches[i]);
-#ifdef HIPBLASLT_USE_ROCROLLER
         }
-#endif
         hipblaslt_init_device(ABC_dims::C,
                               arg.initialization,
                               beta_isnan_type(arg, Talpha),

projects/hipblaslt/clients/common/src/CMakeLists.txt

@@ -14,9 +14,7 @@ target_sources(hipblaslt-clients-common
         "${CMAKE_CURRENT_SOURCE_DIR}/hipblaslt_init_device.cpp"
 )
 
-if(HIPBLASLT_ENABLE_ROCROLLER)
-    target_sources(hipblaslt-clients-common PRIVATE "${CMAKE_CURRENT_SOURCE_DIR}/mxDataGen.cpp")
-endif()
+target_sources(hipblaslt-clients-common PRIVATE "${CMAKE_CURRENT_SOURCE_DIR}/mxDataGen.cpp")
 
 if(HIPBLASLT_ENABLE_BLIS)
     target_sources(hipblaslt-clients-common

projects/hipblaslt/clients/common/src/mxDataGen.cpp

@@ -234,7 +234,6 @@ std::vector<float> generateData(T                           dgen,
 
     std::vector<uint8_t> scaleBytes = dgen.getScaleBytes();
 
-#ifdef HIPBLASLT_USE_ROCROLLER
     // Apply pre-swizzle to scale data
     size_t scaleRows = sizes[0] / elementsPerMXBlock;
     size_t scaleCols = sizes[1];
@@ -244,7 +243,6 @@ std::vector<float> generateData(T                           dgen,
         scaleBytes = DGen::preSwizzleScalesGFX950(scaleBytes, {scaleCols, scaleRows});
 
     }
-#endif
 
     std::memcpy(scale, scaleBytes.data(), scaleBytes.size() * sizeof(uint8_t));
 
@@ -290,7 +288,6 @@ std::vector<float> generateData(T                           dgen,
     }
 }
 
-#ifdef HIPBLASLT_USE_ROCROLLER
 /**
  * @brief Generate random data for OCP (MX) F8/F6/F4 types
  *
@@ -388,7 +385,7 @@ std::vector<float> generateMXInput(hipDataType                dataType,
                                                                   preSwizzleTile,
                                                                   preTile);
     }
-    else if(static_cast<hipDataType>(dataType) == HIP_R_6F_E2M3_EXT)
+    else if(static_cast<hipDataType>(dataType) == HIP_R_6F_E2M3)
     {
         DGen::DataGenerator<DGen::ocp_e2m3_mxfp6> dgen;
         return generateData<decltype(dgen), DGen::ocp_e2m3_mxfp6>(dgen,
@@ -404,7 +401,7 @@ std::vector<float> generateMXInput(hipDataType                dataType,
                                                                   preSwizzleTile,
                                                                   preTile);
     }
-    else if(static_cast<hipDataType>(dataType) == HIP_R_6F_E3M2_EXT)
+    else if(static_cast<hipDataType>(dataType) == HIP_R_6F_E3M2)
     {
         DGen::DataGenerator<DGen::ocp_e3m2_mxfp6> dgen;
         return generateData<decltype(dgen), DGen::ocp_e3m2_mxfp6>(dgen,
@@ -420,7 +417,7 @@ std::vector<float> generateMXInput(hipDataType                dataType,
                                                                   preSwizzleTile,
                                                                   preTile);
     }
-    else if(static_cast<hipDataType>(dataType) == HIP_R_4F_E2M1_EXT)
+    else if(static_cast<hipDataType>(dataType) == HIP_R_4F_E2M1)
     {
         DGen::DataGenerator<DGen::ocp_e2m1_mxfp4> dgen;
         return generateData<decltype(dgen), DGen::ocp_e2m1_mxfp4>(dgen,
@@ -441,4 +438,3 @@ std::vector<float> generateMXInput(hipDataType                dataType,
         throw std::runtime_error("Unsupported data types in MX data generation!");
     }
 }
-#endif

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

@@ -2560,4 +2560,23 @@ Tests:
   beta: [ 0.0, 1.0 ]
   gpu_arch: '90a'
 
+# This is for testing MX FP4 kernel using Tensile
+- name: matmul_tensile_fp4
+  category: quick
+  function:
+    matmul:
+    - { a_type: f4_r, b_type: f4_r, c_type: f32_r, d_type: f32_r, compute_type: c_f32_r, scaleA: 3, scaleB: 3, scale_type: f32_r}
+  M: [2048]
+  N: [2048]
+  K: [4096]
+  transA: T
+  transB: N
+  alpha: 1.0
+  beta: 0.0
+  initialization: hpl
+  unit_check: 0
+  norm_check: 1
+  requested_solution_num: 1
+  gpu_arch: '950'
+
 ...