minor udpate

Author: tinebp

ci/blackbox.sh

@@ -74,6 +74,7 @@ parse_args() {
             --log=*)    LOGFILE=${i#*=} ;;
             --nohup)    TEMPBUILD=1 ;;
             --help)     show_help; exit 0 ;;
+            --*)        echo "Invalid argument: $i"; show_usage; exit 1 ;;
             *)          show_usage; exit 1 ;;
         esac
     done
@@ -140,6 +141,11 @@ main() {
     set_driver_path
     set_app_path
 
+    if [ $SAIF -eq 1 ] && [ "$DRIVER" = "simx" ]; then
+        echo "Error: SAIF is not supported with the simx driver"
+        exit 1
+    fi
+
     # execute on default installed GPU
     if [ "$DRIVER" = "gpu" ]; then
         run_app

kernel/include/vx_tensor.h

@@ -790,7 +790,8 @@ struct wgmma_context {
 
   // Load A fragment (NRA=4 config)
   template <mem_layout src_layout = row_major, typename Frag>
-  static __attribute__((always_inline)) void load_a_sync(Frag &dst, const void *src, size_t ldm) {
+  static __attribute__((always_inline)) void load_matrix_sync(Frag &dst, const void *src, size_t ldm) {
+    stactic_assert(Frag::Use == matrix_a, "only matrix_a fragment can be loaded from registers in wgmma_context");
     ctx_a::template load_matrix_sync<src_layout>(dst, src, ldm);
   }
 

tests/regression/sgemm_tcu_wg/kernel.cpp

@@ -57,16 +57,16 @@ __kernel void kernel_main(kernel_arg_t* __UNIFORM__ arg) {
     auto A_warp = A_smem + warp_rank * ctx::xtileM * ctx::tileK;
     auto desc_b = vt::vx_make_smem_desc(B_smem, ctx::xtileN * sizeof(ctx::input_t));
 
-#if defined(WGMMA_RS)
+  #if defined(WGMMA_RS)
     // RS: A from registers, B from smem
     ctx::fragment_a fragA;
-    ctx::load_a_sync(fragA, A_warp, ctx::tileK);
+    ctx::load_matrix_sync(fragA, A_warp, ctx::tileK);
     ctx::wgmma_sync(fragC, fragA, desc_b, fragC);
-#else
+  #else
     // SS: both from smem
     auto desc_a = vt::vx_make_smem_desc(A_warp, ctx::tileK * sizeof(ctx::input_t));
     ctx::wgmma_sync(fragC, desc_a, desc_b, fragC);
-#endif
+  #endif
 
     __syncthreads();
   }