[rocm-libraries] ROCm/rocm-libraries#1561 (commit c28c222)

[MIOpen] Asm code fix and test update

## Motivation

With last compiler version kernel build failed:
```
/tmp/comgr-xxx/input/conv1x1wrw.s:1606:21: error: redefinition of 'data_prefetch'
MIOpen Error: /xxx/src/MIOpen/src/hipoc/hipoc_program.cpp:299: Code object build failed. Source: conv1x1wrw.s
```
In latest versions of clang , looks like assembler has become stricter
about symbol definitions and treating that the data_prefetch is defined
twice once with default and the other with the = assignment.

## Technical Details

Added test with affected kernel configuration.
Redefinition replaced by new variable.

## Test Plan

mi2xx
./bin/miopen_gtest --gtest_filter=*GPU_Conv2d*AsmBwdWrw*

<!-- Explain any relevant testing done to verify this PR. -->

PASSED

<!-- Briefly summarize test outcomes. -->

## Submission Checklist

- [x] Look over the contributing guidelines at
https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests.

Author: shurale-nkn

src/env_debug.cpp

@@ -65,6 +65,7 @@ MIOPEN_DECLARE_ENV_VAR_STR(MIOPEN_FIND_ENFORCE)
 MIOPEN_DECLARE_ENV_VAR_STR(MIOPEN_FIND_MODE)
 MIOPEN_DECLARE_ENV_VAR_STR(MIOPEN_USER_DB_PATH)
 MIOPEN_DECLARE_ENV_VAR_UINT64(MIOPEN_LOG_LEVEL)
+MIOPEN_DECLARE_ENV_VAR_STR(MIOPEN_DEBUG_CONV_DIRECT_ASM_WRW1X1_PERF_VALS)
 
 namespace miopen {
 namespace debug {
@@ -200,6 +201,7 @@ const LibEnvVar& FindEnvVariable(std::string_view name)
         {MIOPEN_FIND_MODE.GetName(), MIOPEN_FIND_MODE},
         {MIOPEN_LOG_LEVEL.GetName(), MIOPEN_LOG_LEVEL},
         {MIOPEN_USER_DB_PATH.GetName(), MIOPEN_USER_DB_PATH},
+        {MIOPEN_DEBUG_CONV_DIRECT_ASM_WRW1X1_PERF_VALS.GetName(), MIOPEN_DEBUG_CONV_DIRECT_ASM_WRW1X1_PERF_VALS}
         // clang-format on
     };
 

src/kernels/conv1x1wrw.s

@@ -811,14 +811,15 @@ miopenGcnAsmConv1x1WrW:
         .endr
     .endm
 
+data_prefetch_limited = data_prefetch
 .if(full_reads > 0 && full_reads < data_prefetch + 1)
-    data_prefetch = full_reads - 1
+    data_prefetch_limited = full_reads - 1
 .endif
 
 LD_PARTIAL_CHUNKS = 1
 LD_FULL_CHUNKS = 0
 LD_PART_A_ID = 0
-LD_PART_B_ID = data_prefetch
+LD_PART_B_ID = data_prefetch_limited
 last_free_ld_part = LD_PART_B_ID
 
 
@@ -843,14 +844,14 @@ last_free_ld_part = LD_PART_B_ID
 
 .macro data_prefetch_init_q q_wait_cnt, singl_wait_cnt
     q_id = LD_PART_A_ID
-    .rept data_prefetch
+    .rept data_prefetch_limited
         ld_buffers_inc_pointers_rept_waitcnt LD_FULL_CHUNKS, q_id, \singl_wait_cnt
         \q_wait_cnt = \q_wait_cnt + \singl_wait_cnt
         q_id = (q_id + 1)
     .endr
 .endm
 
-.macro data_ld_prefetch_active_loop q_wait_cnt, loop_cnt=data_prefetch+1
+.macro data_ld_prefetch_active_loop q_wait_cnt, loop_cnt=data_prefetch_limited+1
     q_id_conv = LD_PART_A_ID
     q_id_data_ld = LD_PART_B_ID
     .rept \loop_cnt
@@ -860,19 +861,19 @@ last_free_ld_part = LD_PART_B_ID
 
         m_conv_accums elements_per_lane, q_id_conv
 
-        q_id_conv = ((q_id_conv + 1) % (data_prefetch + 1))
-        q_id_data_ld = ((q_id_data_ld + 1) % (data_prefetch + 1))
+        q_id_conv = ((q_id_conv + 1) % (data_prefetch_limited + 1))
+        q_id_data_ld = ((q_id_data_ld + 1) % (data_prefetch_limited + 1))
     .endr
 .endm
 
 .macro data_prefetch_conv_finalizing q_wait_cnt, singl_wait_cnt, q_id_conv_off=0
-    q_id_conv = ((LD_PART_A_ID + \q_id_conv_off) % (data_prefetch + 1))
+    q_id_conv = ((LD_PART_A_ID + \q_id_conv_off) % (data_prefetch_limited + 1))
 
-    .rept data_prefetch
+    .rept data_prefetch_limited
         \q_wait_cnt = (\q_wait_cnt - \singl_wait_cnt)
         s_wait \q_wait_cnt
         m_conv_accums elements_per_lane, q_id_conv
-        q_id_conv = ((q_id_conv + 1) % (data_prefetch + 1))
+        q_id_conv = ((q_id_conv + 1) % (data_prefetch_limited + 1))
     .endr
 .endm
 
@@ -889,21 +890,21 @@ loop_n_begin: // loop over batch (n)
 
         data_prefetch_init_q q_wait_vec_load_full, single_wait_vec_load_full
 
-        .if(full_reads >= 2 * data_prefetch + 1)
+        .if(full_reads >= 2 * data_prefetch_limited + 1)
             loop_hw_begin:
                 data_ld_prefetch_active_loop q_wait_vec_load_full
 
             loop_hw_end:
-                s_addk_i32 s[loop_hw_cnt], 1 + data_prefetch
-                s_cmpk_gt_u32 s[loop_hw_cnt], 0 + full_reads - (2 * data_prefetch + 1)
+                s_addk_i32 s[loop_hw_cnt], 1 + data_prefetch_limited
+                s_cmpk_gt_u32 s[loop_hw_cnt], 0 + full_reads - (2 * data_prefetch_limited + 1)
                 s_cbranch_scc0 loop_hw_begin
         .endif
-        .if( (full_reads - data_prefetch) % (data_prefetch + 1) != 0)
-            loop_resi = ((full_reads - data_prefetch) % (data_prefetch + 1))
+        .if( (full_reads - data_prefetch_limited) % (data_prefetch_limited + 1) != 0)
+            loop_resi = ((full_reads - data_prefetch_limited) % (data_prefetch_limited + 1))
 
             data_ld_prefetch_active_loop q_wait_vec_load_full, loop_resi
 
-            last_free_ld_part = ((LD_PART_B_ID + loop_resi) % (data_prefetch + 1))
+            last_free_ld_part = ((LD_PART_B_ID + loop_resi) % (data_prefetch_limited + 1))
         .endif
     .endif
 
@@ -919,7 +920,7 @@ loop_n_begin: // loop over batch (n)
         q_wait_vec_load_full = q_wait_vec_load_full + wait_vec_load_part
     .endif
 
-    .if(full_reads != 0 && data_prefetch != 0)
+    .if(full_reads != 0 && data_prefetch_limited != 0)
         data_prefetch_conv_finalizing q_wait_vec_load_full, single_wait_vec_load_full, loop_resi
     .endif
 

test/gtest/gtest_common.hpp

@@ -120,6 +120,14 @@ inline void tuning_check(const std::string& err)
     default_check(err);
 }
 
+inline void compiler_check(const std::string& err)
+{
+    // the test should fail if kernel build failed.
+    EXPECT_FALSE(err.find("Error") != std::string::npos ||
+                 err.find("Code object build failed") != std::string::npos);
+    default_check(err);
+}
+
 inline void db_check(const std::string& err)
 {
     EXPECT_FALSE(err.find("Perf Db: record not found") != std::string::npos);

test/gtest/smoke_solver_convasmbwdwrw.cpp

@@ -32,6 +32,26 @@
 
 namespace {
 
+MIOPEN_LIB_ENV_VAR(MIOPEN_DEBUG_CONV_DIRECT_ASM_WRW1X1_PERF_VALS)
+
+auto GetTestCompilerRegressionCases()
+{
+    const auto env_regression_check =
+        std::tuple{std::pair{MIOPEN_FIND_ENFORCE, "SEARCH"},
+                   std::pair{wa::MIOPEN_DEBUG_TUNING_ITERATIONS_MAX, 1},
+                   std::pair{MIOPEN_FIND_MODE, "normal"},
+                   std::pair{MIOPEN_DEBUG_CONV_DIRECT_ASM_WRW1X1_PERF_VALS, "2,8,4,2,4,2,2,4,0,2"},
+                   std::pair{MIOPEN_DEBUG_FIND_ONLY_SOLVER, "ConvAsmBwdWrW1x1"}};
+
+    const std::string vw = " --verbose --disable-forward --disable-backward-data";
+
+    return std::vector{
+        // clang-format off
+    std::pair{env_regression_check, vw + " --input 1 4 6 6 --weights 4 4 1 1 --pads_strides_dilations 0 0 1 1 1 1"}
+        // clang-format on
+    };
+}
+
 auto GetTestCases()
 {
     const auto env_w1 = std::tuple{std::pair{MIOPEN_FIND_ENFORCE, "SEARCH_DB_UPDATE"},
@@ -48,7 +68,8 @@ auto GetTestCases()
     };
 }
 
-using TestCase = decltype(GetTestCases())::value_type;
+using RegrTestCase = decltype(GetTestCompilerRegressionCases())::value_type;
+using TestCase     = decltype(GetTestCases())::value_type;
 
 bool SkipTest() { return get_handle_xnack(); }
 
@@ -61,6 +82,10 @@ bool IsTestSupportedForDevice()
 
 } // namespace
 
+class GPU_Conv2dSingleAsmBwdWrw_FP32 : public FloatTestCase<std::vector<RegrTestCase>>
+{
+};
+
 class GPU_Conv2dTuningAsmBwdWrw_FP32 : public FloatTestCase<std::vector<TestCase>>
 {
 };
@@ -85,6 +110,18 @@ TEST_P(GPU_Conv2dTuningAsmBwdWrw_FP32, FloatTest_smoke_solver_convasmbwdwrw)
     }
 };
 
+TEST_P(GPU_Conv2dSingleAsmBwdWrw_FP32, FloatTest_smoke_solver_convasmbwdwrw)
+{
+    if(IsTestSupportedForDevice() && !SkipTest())
+    {
+        invoke_with_params<conv2d_driver, GPU_Conv2dSingleAsmBwdWrw_FP32>(compiler_check);
+    }
+    else
+    {
+        GTEST_SKIP();
+    }
+};
+
 TEST_P(GPU_Conv2dTuningAsmBwdWrw_FP16, HalfTest_smoke_solver_convasmbwdwrw)
 {
     if(IsTestSupportedForDevice() && !SkipTest())
@@ -109,6 +146,9 @@ TEST_P(GPU_Conv2dTuningAsmBwdWrw_BFP16, Bf16Test_smoke_solver_convasmbwdwrw)
     }
 };
 
+INSTANTIATE_TEST_SUITE_P(Smoke,
+                         GPU_Conv2dSingleAsmBwdWrw_FP32,
+                         testing::Values(GetTestCompilerRegressionCases()));
 INSTANTIATE_TEST_SUITE_P(Smoke, GPU_Conv2dTuningAsmBwdWrw_FP32, testing::Values(GetTestCases()));
 INSTANTIATE_TEST_SUITE_P(Smoke, GPU_Conv2dTuningAsmBwdWrw_FP16, testing::Values(GetTestCases()));
 INSTANTIATE_TEST_SUITE_P(Smoke, GPU_Conv2dTuningAsmBwdWrw_BFP16, testing::Values(GetTestCases()));