[tests] Fix for DEV builds. Fix testing on 16GB GPUs. #1729

atamazov · 2022-09-01T00:07:37Z

A pack of changes that expected to fix problems with testing of DEV builds and problems with testing on 16GB GPUs (e.g. Radeon VII, MI25, Radeon RX 6800/6900). May improve CI stability.

🔴 [tests][Pooling] Remove unused var from MIOpenPoolingBwd.cl to avoid warning
- Resolves LRN testing issues with DEV builds, see [gfx1030][FP16][ROCm5.2] test_lrn_test failure due to hipRTC issue #1674 (comment)
🔴 [tests] default_allocator/deallocator: do not fail if GetAvailableMemory reports insufficient memory but use hipHostMalloc instead.
- This fixes testing failures of test_lrn_test, test_activation, test_pooling2d (and possibly more) with 16GB GPU cards, see [gfx1030][FP16][ROCm5.2] test_lrn_test failure due to hipRTC issue #1674 (comment). Also prevents similar problems that may happen with GPUs that have more memory onboard.
🟡 [tests] handle_test: Better diags in test_warnings
- Makes error messages more clear, partially resolving issue mentioned at [gfx1030][FP16][ROCm5.2] test_lrn_test failure due to hipRTC issue #1674 (comment).
🟢 [tests] Improve logging related to GPU memory allocation/deallocation.

@shurale-nkn @junliume @muralinr Please review. I would recommend urgency_high. Thanks.

Local `make check` passed on Radeon VII with both FP16 and FP32.

[Informative] COMMANDS TO BUILD AND RUN TESTS:

$ CXX=/opt/rocm/llvm/bin/clang++ \
CXXFLAGS='-Werror' \
cmake \
-DMIOPEN_TEST_FLAGS=' --disable-verification-cache ' \
-DCMAKE_BUILD_TYPE=release \
-DBUILD_DEV=On \
-DMIOPEN_GPU_SYNC=Off \
-DMIOPEN_TEST_ALL=On \
-DMIOPEN_TEST_HALF=On \
../..
$ CTEST_PARALLEL_LEVEL=4 make -j `nproc --ignore=4` check

… HIP version >= 5.3.22305 (Mainline build 10659)

…cl to avoid warning

… if GetAvailableMemory reports insufficient memory but use hipHostMalloc instead. Improve logging.

…e-1674-2

atamazov · 2022-09-01T23:13:43Z

@junliume @shurale-nkn Locally tested. Description updated. Ready for review.

atamazov · 2022-09-01T23:17:15Z

@junliume If some memory-consuming tests were disabled due to instability, then we can try enabling them after this PR is merged in.

muralinr · 2022-09-02T03:19:58Z

src/kernels/MIOpenPoolingBwd.cl

@@ -284,9 +284,8 @@ mloPoolingMaxBwd(const __global _FLOAT* top_df,

    barrier(CLK_LOCAL_MEM_FENCE);
    _FLOAT add_val;
-    int bt_y  = (y + lcl_id1 * MLO_POOLBWD_N_VERT_OUT_PIX);
-    int bt_x  = (x + lcl_id0 * MLO_POOLBWD_N_HORIZ_OUT_PIX);
-    int b_idx = bt_y * mlo_bot_width + bt_x;


is "b_idx" not used anywhere and redundant?

@muralinr it seems that we have met the same issue:
#1732

Yes Jun. I looked at code. 'b_idx" is computed and it was not used inside code anywhere. It's should be removed it to avoid compilation issues.

muralinr · 2022-09-02T03:21:40Z

src/hip/hiperrors.cpp

@@ -30,7 +30,7 @@ namespace miopen {

 std::string HIPErrorMessage(int error, const std::string& msg)
 {
-    return msg + " " + hipGetErrorString(static_cast<hipError_t>(error));
+    return msg + ": " + hipGetErrorString(static_cast<hipError_t>(error));


is ":" change mandatory?

Yes, if we want to automatically separate error message from the error code, see https://github.com/ROCmSoftwarePlatform/MIOpen/pull/1729/files#r962057473.

muralinr · 2022-09-02T03:22:49Z

test/handle_test.cpp

@@ -211,6 +211,7 @@ void test_warnings(kernel_type_t kern_type)
    if(kern_type == miopenOpenCLKernelType)
        EXPECT(throws([&] {
            h.AddKernel("GEMM", "", WriteNop(kern_type), "write", {1, 1, 1}, {1, 1, 1}, "");
+            MIOPEN_LOG_E("FAILED: Build of the OpenCL kernel should produce warnings");


This change is very good.

muralinr · 2022-09-02T03:23:07Z

test/handle_test.cpp

@@ -223,7 +224,8 @@ void test_warnings(kernel_type_t kern_type)
                        "",
                        0,
                        false,
-                        WriteNop(miopenHIPKernelType));
+                        WriteNop(kern_type));
+            MIOPEN_LOG_E("FAILED: Build of the HIP kernel 'nop_hip.cpp' should produce warnings");


Same here. This change is good.

shurale-nkn · 2022-09-02T19:46:08Z

src/hip/handlehip.cpp

+    auto status = hipMemPtrGetInfo(mem, &size);
+    if(status != hipSuccess)
+        MIOPEN_LOG_W("hipMemPtrGetInfo at " << mem << " status: " << status);
+    status = hipFree(mem);


is it ok to use hipFree for device and for host memory?

src/hip/handlehip.cpp

atamazov · 2022-09-02T23:05:37Z

src/hip/handlehip.cpp

-            MIOPEN_THROW_HIP_STATUS(status,
-                                    "Hip error creating buffer " + std::to_string(sz) + ": ");


@muralinr I do not think we want to append + ": " to each MIOPEN_THROW_HIP_STATUS.

atamazov added 9 commits August 17, 2022 23:07

fix-issue-1674(01) W/A for issue 1674: Disable -Wgnu-line-marker when…

a7110f4

… HIP version >= 5.3.22305 (Mainline build 10659)

fix-issue-1674(02) [tests] handle_test: Better diags in test_warnings

1b966c5

fix-issue-1674(03) [Pooling] Remove unused var from MIOpenPoolingBwd.…

ccecacb

…cl to avoid warning

fix-issue-1674(04) [tests] default_allocator/deallocator: do not fail…

dd59a7f

… if GetAvailableMemory reports insufficient memory but use hipHostMalloc instead. Improve logging.

fix-issue-1674(05) Formatting

4b846d3

Merge commit '1b966c5bc64a9e2d8986b411f9512775ee20cd61' into fix-issu…

288384d

…e-1674-2

Merge commit 'ccecacbc027aa2e724faa7b85501714e98e7631e' into fix-issu…

3520dad

…e-1674-2

Merge commit 'dd59a7fecf4af77a526321117da7d65ad6773f64' into fix-issu…

c5e5c8d

…e-1674-2

Merge commit '4b846d3b5944f256b5018ae3d48014c010e03518' into fix-issu…

93ac5e2

…e-1674-2

atamazov changed the title ~~Fix issue 1674 2~~ [tests] Fix for DEV builds. Fix testing on 16GB GPUs. Sep 1, 2022

atamazov marked this pull request as ready for review September 1, 2022 23:13

muralinr reviewed Sep 2, 2022

View reviewed changes

shurale-nkn reviewed Sep 2, 2022

View reviewed changes

junliume reviewed Sep 2, 2022

View reviewed changes

src/hip/handlehip.cpp Outdated Show resolved Hide resolved

junliume added the urgency_high label Sep 2, 2022

junliume reviewed Sep 2, 2022

View reviewed changes

src/hip/handlehip.cpp Outdated Show resolved Hide resolved

atamazov mentioned this pull request Sep 2, 2022

[ROCm 5.3] OpenCL kernels compilation issues [-Werror,-Wunused-but-set-variable] #1732

Closed

atamazov commented Sep 2, 2022

View reviewed changes

Update src/hip/handlehip.cpp

802b37c

junliume previously approved these changes Sep 4, 2022

View reviewed changes

Merge branch 'develop' into fix-issue-1674-2

6fdb01f

junliume dismissed their stale review via 6fdb01f September 4, 2022 05:15

junliume approved these changes Sep 4, 2022

View reviewed changes

junliume merged commit 601fdf5 into ROCm:develop Sep 4, 2022

junliume mentioned this pull request Dec 13, 2023

[Issue]: [Tests] test_handle_test will fail under DEV builds #2600

Closed

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

[tests] Fix for DEV builds. Fix testing on 16GB GPUs. #1729

[tests] Fix for DEV builds. Fix testing on 16GB GPUs. #1729

atamazov commented Sep 1, 2022 •

edited

Loading

atamazov commented Sep 1, 2022

atamazov commented Sep 1, 2022

muralinr Sep 2, 2022

junliume Sep 2, 2022

muralinr Sep 2, 2022

muralinr Sep 2, 2022

atamazov Sep 2, 2022 •

edited

Loading

muralinr Sep 3, 2022

muralinr Sep 2, 2022

muralinr Sep 2, 2022

shurale-nkn Sep 2, 2022

atamazov Sep 2, 2022

		MIOPEN_THROW_HIP_STATUS(status,
		"Hip error creating buffer " + std::to_string(sz) + ": ");

[tests] Fix for DEV builds. Fix testing on 16GB GPUs. #1729

[tests] Fix for DEV builds. Fix testing on 16GB GPUs. #1729

Conversation

atamazov commented Sep 1, 2022 • edited Loading

Local make check passed on Radeon VII with both FP16 and FP32.

atamazov commented Sep 1, 2022

atamazov commented Sep 1, 2022

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

atamazov Sep 2, 2022 • edited Loading

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

atamazov commented Sep 1, 2022 •

edited

Loading

Local `make check` passed on Radeon VII with both FP16 and FP32.

atamazov Sep 2, 2022 •

edited

Loading