ROCm · JDas-AMD · Jun 12, 2026 · Aug 27, 2026 · Aug 28, 2026
@@ -81,7 +81,7 @@ def create_test_definition(
     for entry in disabled:
         tags_str += f"[exclude_{entry}]"
 
-    return f'#define {case_name} "{case_name}", "{tags_str}"'
+    return f'#define {case_name} "{tags_str}"'
 
 
 def generate_parameter_header(cmd_options, output_path):

@@ -75,11 +75,11 @@ per-API notes on why each gap exists.
        tags: [<domain>]
    ```
 
-   This is **required, not optional**. `HIP_TEST_CASE(name)` expands through
-   `GET_TAGS(name)` → `SECOND_ARG(name)`, which needs a `name` macro that the
-   build **generates from this YAML** into `hip_test_config.hh`. A test-case name
+   This is **required, not optional**. `HIP_TEST_CASE(name)` expands to
+   `TEST_CASE(#name, name)`, which needs a `name` macro that the build
+   **generates from this YAML** into `hip_test_config.hh`. A test-case name
    with no YAML entry fails to compile with
-   *"too few arguments to function-like macro SECOND_ARG"*. After editing the
+   *"use of undeclared identifier '<name>'"*. After editing the
    YAML you must re-run CMake configure so the header regenerates.
 
 6. **Regenerate the test plan and update the docs**: run

@@ -26,19 +26,16 @@
 #ifdef ENABLE_YAML_TAGS
 #include "hip_test_config.hh"
 
-#define SECOND_ARG(a, b, ...) b
-#define GET_TAGS(...) SECOND_ARG(__VA_ARGS__)
-#define HIP_TEST_CASE(name) TEST_CASE(#name, GET_TAGS(name))
-#define HIP_TEMPLATE_TEST_CASE(name, ...) TEMPLATE_TEST_CASE(#name, GET_TAGS(name), __VA_ARGS__)
+#define HIP_TEST_CASE(name) TEST_CASE(#name, name)
+#define HIP_TEMPLATE_TEST_CASE(name, ...) TEMPLATE_TEST_CASE(#name, name, __VA_ARGS__)
 
 #else
-#define GET_TAGS(...)
 #define HIP_TEST_CASE(name) TEST_CASE(#name, "")
 #define HIP_TEMPLATE_TEST_CASE(name, ...) TEMPLATE_TEST_CASE(#name, "", __VA_ARGS__)
 #endif
 
 /**
- * @brief Check if running at quick level (level_0).
+ * @brief Check if r1ning at quick level (level_0).
  * Use this to reduce test parameters for faster execution.
  */
 inline bool isQuickLevel() {

@@ -177,7 +177,7 @@ static void runUnsafeAtomicAddHalfAndBfloatTest() {
   HIP_CHECK(hipFree(out));
 }
 
-TEST_CASE(Unit_unsafe_atomic_add_half_and_bfloat) {
+HIP_TEST_CASE(Unit_unsafe_atomic_add_half_and_bfloat) {
   SECTION("__half2") { runUnsafeAtomicAddHalfAndBfloatTest<__half2>(); }
   SECTION("__hip_bfloat162") { runUnsafeAtomicAddHalfAndBfloatTest<__hip_bfloat162>(); }
   SECTION("__half") { runUnsafeAtomicAddHalfAndBfloatTest<__half>(); }

@@ -992,7 +992,7 @@ HIP_TEST_CASE(Unit_Thread_Block_Tile_Reduce_Trivially_Copyable_Parameters)
   }
 }
 
-TEST_CASE(Unit_Thread_Block_Tile_Scan_Trivially_Copyable_Parameters)
+HIP_TEST_CASE(Unit_Thread_Block_Tile_Scan_Trivially_Copyable_Parameters)
 {
   CHECK_COOPERATIVE_LAUNCH_SUPPORT
 
@@ -1278,7 +1278,7 @@ HIP_TEST_CASE(Unit_Thread_Block_Tile_Reduce_All_Parameter_Sizes)
   }
 }
 
-TEST_CASE(Unit_Thread_Block_Tile_Scan_All_Parameter_Sizes)
+HIP_TEST_CASE(Unit_Thread_Block_Tile_Scan_All_Parameter_Sizes)
 {
   CHECK_COOPERATIVE_LAUNCH_SUPPORT
 
@@ -1475,7 +1475,7 @@ void testScanForTileSize()
   }
 }
 
-TEST_CASE(Unit_Thread_Block_Tile_Inclusive_Scan_Basic)
+HIP_TEST_CASE(Unit_Thread_Block_Tile_Inclusive_Scan_Basic)
 {
   CHECK_COOPERATIVE_LAUNCH_SUPPORT
 
@@ -1495,7 +1495,7 @@ TEST_CASE(Unit_Thread_Block_Tile_Inclusive_Scan_Basic)
   }
 }
 
-TEMPLATE_TEST_CASE(Unit_Thread_Block_Tile_Exclusive_Scan_Basic, int, half)
+HIP_TEMPLATE_TEST_CASE(Unit_Thread_Block_Tile_Exclusive_Scan_Basic, int, half)
 {
   CHECK_COOPERATIVE_LAUNCH_SUPPORT
 
@@ -1533,7 +1533,7 @@ TEMPLATE_TEST_CASE(Unit_Thread_Block_Tile_Exclusive_Scan_Basic, int, half)
 
 // for all the tile sizes and all input types, using random input values, calculates the scan
 // values. Additionally, randomly make some threads not participate for the coalesced_threads case
-TEMPLATE_TEST_CASE(Unit_Thread_Block_Tile_Scan_Random_arithmetic,  int, unsigned int, long long,
+HIP_TEMPLATE_TEST_CASE(Unit_Thread_Block_Tile_Scan_Random_arithmetic,  int, unsigned int, long long,
                    unsigned long long, float, half, double)
 {
   CHECK_COOPERATIVE_LAUNCH_SUPPORT
@@ -1559,7 +1559,7 @@ TEMPLATE_TEST_CASE(Unit_Thread_Block_Tile_Scan_Random_arithmetic,  int, unsigned
   }
 }
 
-TEMPLATE_TEST_CASE(Unit_Thread_Block_Tile_Scan_Random_boolean, int, unsigned int, long long,
+HIP_TEMPLATE_TEST_CASE(Unit_Thread_Block_Tile_Scan_Random_boolean, int, unsigned int, long long,
                    unsigned long long)
 {
   CHECK_COOPERATIVE_LAUNCH_SUPPORT
@@ -1586,7 +1586,7 @@ TEMPLATE_TEST_CASE(Unit_Thread_Block_Tile_Scan_Random_boolean, int, unsigned int
 }
 
 // make sures that tiled blocks that use the y or z dimension work correctly
-TEST_CASE(Unit_Thread_Block_Tile_2D_3D_Blocks)
+HIP_TEST_CASE(Unit_Thread_Block_Tile_2D_3D_Blocks)
 {
   CHECK_COOPERATIVE_LAUNCH_SUPPORT
 
@@ -1621,7 +1621,7 @@ TEST_CASE(Unit_Thread_Block_Tile_2D_3D_Blocks)
   }
 }
 
-TEMPLATE_TEST_CASE(Unit_Thread_Block_Coalesced_Scan_arithmetic, int, unsigned int, long long,
+HIP_TEMPLATE_TEST_CASE(Unit_Thread_Block_Coalesced_Scan_arithmetic, int, unsigned int, long long,
                    unsigned long long, float, half, double)
 {
   CHECK_COOPERATIVE_LAUNCH_SUPPORT
@@ -1647,7 +1647,7 @@ TEMPLATE_TEST_CASE(Unit_Thread_Block_Coalesced_Scan_arithmetic, int, unsigned in
   }
 }
 
-TEMPLATE_TEST_CASE(Unit_Thread_Block_Coalesced_Scan_boolean, int, unsigned int, long long,
+HIP_TEMPLATE_TEST_CASE(Unit_Thread_Block_Coalesced_Scan_boolean, int, unsigned int, long long,
                    unsigned long long)
 {
   CHECK_COOPERATIVE_LAUNCH_SUPPORT
@@ -1706,7 +1706,7 @@ void __global__ binaryPartitionTiled(int* out, int* ranks)
   }
 }
 
-TEST_CASE(Unit_Thread_Block_Scan_partition)
+HIP_TEST_CASE(Unit_Thread_Block_Scan_partition)
 {
   CHECK_COOPERATIVE_LAUNCH_SUPPORT
 

@@ -140,7 +140,7 @@ static void runAtomicAddCoherentUnsafeFlagTest(const std::string& gfxName) {
   HIP_CHECK(hipHostFree(result));
 }
 
-TEST_CASE(Unit_AtomicAdd_Coherent) {
+HIP_TEST_CASE(Unit_AtomicAdd_Coherent) {
   hipDeviceProp_t prop;
   int device;
   HIP_CHECK(hipGetDevice(&device));

@@ -134,7 +134,7 @@ static void runAtomicAddNonCoherentUnsafeFlagTest() {
   HIP_CHECK(hipHostFree(result));
 }
 
-TEST_CASE(Unit_AtomicAdd_NonCoherent) {
+HIP_TEST_CASE(Unit_AtomicAdd_NonCoherent) {
   hipDeviceProp_t prop;
   int device;
   HIP_CHECK(hipGetDevice(&device));

@@ -167,7 +167,7 @@ static void runUnsafeAtomicAddCoherentUnsafeFlagTest(const std::string& gfxName)
   HIP_CHECK(hipHostFree(result));
 }
 
-TEST_CASE(Unit_unsafeAtomicAdd_Coherent) {
+HIP_TEST_CASE(Unit_unsafeAtomicAdd_Coherent) {
   hipDeviceProp_t prop;
   int device;
   HIP_CHECK(hipGetDevice(&device));

@@ -146,7 +146,7 @@ static void runUnsafeAtomicAddNonCoherentUnsafeFlagTest() {
   HIP_CHECK(hipHostFree(result));
 }
 
-TEST_CASE(Unit_unsafeAtomicAdd_NonCoherent) {
+HIP_TEST_CASE(Unit_unsafeAtomicAdd_NonCoherent) {
   hipDeviceProp_t prop;
   int device;
   HIP_CHECK(hipGetDevice(&device));

@@ -261,7 +261,7 @@ HIP_TEST_CASE(Unit_Rtc_CoopReduce)
   }
 }
 
-TEST_CASE(Unit_Rtc_CoopScan)
+HIP_TEST_CASE(Unit_Rtc_CoopScan)
 {
   const std::tuple<int, unsigned int, long long, unsigned long long, float, half, double> allTypes;
   const std::tuple<int, unsigned int, long long, unsigned long long> integralTypes;