From c87014a54cab1333fc7fff80c8d678d2ee97445e Mon Sep 17 00:00:00 2001 From: Gerardo Hernandez Date: Tue, 25 Nov 2025 06:37:47 +0000 Subject: [PATCH] SWDEV-534207 fix order of kernel launch parameters when calling notifiedKernel in some tests: kernel<<>> instead of kernel<<>>. This was causing out of bounds accesses (#1860) --- .../catch/unit/memory/hipMemPoolApi.cc | 24 +++++++++---------- .../unit/memory/hipMemPoolSetGetAttribute.cc | 24 +++++++++---------- .../catch/unit/memory/hipMemPoolTrimTo.cc | 2 +- .../catch/unit/memory/mempool_common.hh | 8 +++---- 4 files changed, 29 insertions(+), 29 deletions(-) diff --git a/projects/hip-tests/catch/unit/memory/hipMemPoolApi.cc b/projects/hip-tests/catch/unit/memory/hipMemPoolApi.cc index dd781cfce6..002aea92a4 100644 --- a/projects/hip-tests/catch/unit/memory/hipMemPoolApi.cc +++ b/projects/hip-tests/catch/unit/memory/hipMemPoolApi.cc @@ -127,7 +127,7 @@ TEST_CASE("Unit_hipMemPoolApi_BasicAlloc") { int blocks = 1024; hipMemPoolAttr attr; - notifiedKernel<<<32, blocks, 0, stream>>>(B, notified); + notifiedKernel<<>>(B, notified); HIP_CHECK(hipFreeAsync(reinterpret_cast(B), stream)); @@ -215,7 +215,7 @@ TEST_CASE("Unit_hipMemPoolApi_BasicTrim") { mem_pool, stream)); int blocks = 2; - notifiedKernel<<<32, blocks, 0, stream>>>(B, notified); + notifiedKernel<<>>(B, notified); hipMemPoolAttr attr; attr = hipMemPoolAttrReleaseThreshold; @@ -303,7 +303,7 @@ TEST_CASE("Unit_hipMemPoolApi_BasicReuse") { mem_pool, stream)); int blocks = 2; - notifiedKernel<<<32, blocks, 0, stream>>>(A, notified); + notifiedKernel<<>>(A, notified); hipMemPoolAttr attr; // Not a real free, since kernel isn't done @@ -322,7 +322,7 @@ TEST_CASE("Unit_hipMemPoolApi_BasicReuse") { // Second kernel launch with new memory *notified = 0; - notifiedKernel<<<32, blocks, 0, stream>>>(B, notified); + notifiedKernel<<>>(B, notified); std::this_thread::sleep_for(std::chrono::milliseconds(500)); *notified = 1; // Notify kernel loop to exit @@ -396,7 +396,7 @@ TEST_CASE("Unit_hipMemPoolApi_Opportunistic") { HIP_CHECK(hipMemPoolSetAttribute(mem_pool, attr, &value)); // Run kernel in the first stream - notifiedKernel<<<32, blocks, 0, stream1>>>(A, notified1); + notifiedKernel<<>>(A, notified1); // Not a real free, since kernel isn't done HIP_CHECK(hipFreeAsync(reinterpret_cast(A), stream1)); @@ -414,7 +414,7 @@ TEST_CASE("Unit_hipMemPoolApi_Opportunistic") { REQUIRE(A != B); // Run kernel with the new memory in the second streamn - notifiedKernel<<<32, blocks, 0, stream2>>>(B, notified2); + notifiedKernel<<>>(B, notified2); std::this_thread::sleep_for(std::chrono::milliseconds(500)); *notified2 = 1; // Notify kernel loop to exit @@ -435,7 +435,7 @@ TEST_CASE("Unit_hipMemPoolApi_Opportunistic") { HIP_CHECK(hipMemPoolSetAttribute(mem_pool, attr, &value)); // Run kernel in the first stream - notifiedKernel<<<32, blocks, 0, stream1>>>(A, notified1); + notifiedKernel<<>>(A, notified1); // Not a real free, since kernel isn't done HIP_CHECK(hipFreeAsync(reinterpret_cast(A), stream1)); @@ -453,7 +453,7 @@ TEST_CASE("Unit_hipMemPoolApi_Opportunistic") { REQUIRE(A == B); // Run kernel with the new memory in the second stream - notifiedKernel<<<32, blocks, 0, stream2>>>(B, notified2); + notifiedKernel<<>>(B, notified2); std::this_thread::sleep_for(std::chrono::milliseconds(500)); *notified2 = 1; // Notify kernel loop to exit @@ -475,7 +475,7 @@ TEST_CASE("Unit_hipMemPoolApi_Opportunistic") { HIP_CHECK(hipMemPoolSetAttribute(mem_pool, attr, &value)); // Run kernel in the first stream - notifiedKernel<<<32, blocks, 0, stream1>>>(A, notified1); + notifiedKernel<<>>(A, notified1); // Not a real free, since kernel isn't done HIP_CHECK(hipFreeAsync(reinterpret_cast(A), stream1)); @@ -489,7 +489,7 @@ TEST_CASE("Unit_hipMemPoolApi_Opportunistic") { REQUIRE(A != B); // Run kernel with the new memory in the second stream - notifiedKernel<<<32, blocks, 0, stream2>>>(B, notified2); + notifiedKernel<<>>(B, notified2); std::this_thread::sleep_for(std::chrono::milliseconds(500)); *notified1 = 1; // Notify kernel loop to exit @@ -532,7 +532,7 @@ TEST_CASE("Unit_hipMemPoolApi_Default") { HIP_CHECK(hipMallocAsync(reinterpret_cast(&C), numElements * sizeof(float), stream)); int blocks = 2; - notifiedKernel<<<32, blocks, 0, stream>>>(A, notified); + notifiedKernel<<>>(A, notified); hipMemPoolAttr attr; // Not a real free, since kernel isn't done @@ -550,7 +550,7 @@ TEST_CASE("Unit_hipMemPoolApi_Default") { // Second kernel launch with new memory *notified = 0; - notifiedKernel<<<32, blocks, 0, stream>>>(B, notified); + notifiedKernel<<>>(B, notified); HIP_CHECK(hipFreeAsync(reinterpret_cast(B), stream)); diff --git a/projects/hip-tests/catch/unit/memory/hipMemPoolSetGetAttribute.cc b/projects/hip-tests/catch/unit/memory/hipMemPoolSetGetAttribute.cc index c45b837b12..93667cabb0 100644 --- a/projects/hip-tests/catch/unit/memory/hipMemPoolSetGetAttribute.cc +++ b/projects/hip-tests/catch/unit/memory/hipMemPoolSetGetAttribute.cc @@ -156,7 +156,7 @@ TEST_CASE("Unit_hipMemPoolSetAttribute_Opportunistic") { HIP_CHECK(hipMemPoolSetAttribute(mempool.mempool(), attr, &value)); // Run kernel in the first stream - notifiedKernel<<<32, blocks, 0, stream1.stream()>>>(alloc_mem1, notified1); + notifiedKernel<<>>(alloc_mem1, notified1); // Not a real free, since kernel isn't done HIP_CHECK(hipFreeAsync(reinterpret_cast(alloc_mem1), stream1.stream())); std::this_thread::sleep_for(std::chrono::milliseconds(500)); @@ -171,7 +171,7 @@ TEST_CASE("Unit_hipMemPoolSetAttribute_Opportunistic") { REQUIRE(alloc_mem1 != alloc_mem2); // Run kernel with the new memory in the second stream - notifiedKernel<<<32, blocks, 0, stream2.stream()>>>(alloc_mem2, notified2); + notifiedKernel<<>>(alloc_mem2, notified2); std::this_thread::sleep_for(std::chrono::milliseconds(500)); *notified2 = 1; @@ -195,7 +195,7 @@ TEST_CASE("Unit_hipMemPoolSetAttribute_Opportunistic") { HIP_CHECK(hipMemPoolSetAttribute(mempool.mempool(), attr, &value)); // Run kernel in the first stream - notifiedKernel<<<32, blocks, 0, stream1.stream()>>>(alloc_mem1, notified1); + notifiedKernel<<>>(alloc_mem1, notified1); // Not a real free, since kernel isn't done HIP_CHECK(hipFreeAsync(reinterpret_cast(alloc_mem1), stream1.stream())); @@ -211,7 +211,7 @@ TEST_CASE("Unit_hipMemPoolSetAttribute_Opportunistic") { REQUIRE(alloc_mem1 == alloc_mem2); // Run kernel with the new memory in the first stream - notifiedKernel<<<32, blocks, 0, stream1.stream()>>>(alloc_mem2, notified2); + notifiedKernel<<>>(alloc_mem2, notified2); std::this_thread::sleep_for(std::chrono::milliseconds(500)); *notified2 = 1; @@ -230,7 +230,7 @@ TEST_CASE("Unit_hipMemPoolSetAttribute_Opportunistic") { // Enable Opportunistic HIP_CHECK(hipMemPoolSetAttribute(mempool.mempool(), attr, &value)); // Run kernel in the first stream - notifiedKernel<<<32, blocks, 0, stream1.stream()>>>(alloc_mem1, notified1); + notifiedKernel<<>>(alloc_mem1, notified1); // Not a real free, since kernel isn't done HIP_CHECK(hipFreeAsync(reinterpret_cast(alloc_mem1), stream1.stream())); @@ -248,7 +248,7 @@ TEST_CASE("Unit_hipMemPoolSetAttribute_Opportunistic") { REQUIRE(alloc_mem1 == alloc_mem2); // Run kernel with the new memory in the second stream - notifiedKernel<<<32, blocks, 0, stream2.stream()>>>(alloc_mem2, notified2); + notifiedKernel<<>>(alloc_mem2, notified2); std::this_thread::sleep_for(std::chrono::milliseconds(500)); *notified2 = 1; // Notifiy kernel to exit after 500 ms @@ -270,7 +270,7 @@ TEST_CASE("Unit_hipMemPoolSetAttribute_Opportunistic") { HIP_CHECK(hipMemPoolSetAttribute(mempool.mempool(), attr, &value)); // Run kernel in the first stream - notifiedKernel<<<32, blocks, 0, stream1.stream()>>>(alloc_mem1, notified1); + notifiedKernel<<>>(alloc_mem1, notified1); // Not a real free, since kernel isn't done HIP_CHECK(hipFreeAsync(reinterpret_cast(alloc_mem1), stream1.stream())); @@ -283,7 +283,7 @@ TEST_CASE("Unit_hipMemPoolSetAttribute_Opportunistic") { REQUIRE(alloc_mem1 != alloc_mem2); // Run kernel with the new memory in the second stream - notifiedKernel<<<32, blocks, 0, stream2.stream()>>>(alloc_mem2, notified2); + notifiedKernel<<>>(alloc_mem2, notified2); std::this_thread::sleep_for(std::chrono::milliseconds(500)); *notified1 = 1; @@ -352,7 +352,7 @@ TEST_CASE("Unit_hipMemPoolSetAttribute_EventDependencies") { HIP_CHECK(hipMemPoolSetAttribute(mempool.mempool(), attr, &value)); // Run kernel in the first stream - notifiedKernel<<<32, blocks, 0, stream1.stream()>>>(alloc_mem1, notified1); + notifiedKernel<<>>(alloc_mem1, notified1); // Not a real free, since kernel isn't done HIP_CHECK(hipFreeAsync(reinterpret_cast(alloc_mem1), stream1.stream())); @@ -367,7 +367,7 @@ TEST_CASE("Unit_hipMemPoolSetAttribute_EventDependencies") { REQUIRE(alloc_mem1 == alloc_mem2); // Run kernel with the new memory in the second stream - notifiedKernel<<<32, blocks, 0, stream2.stream()>>>(alloc_mem2, notified2); + notifiedKernel<<>>(alloc_mem2, notified2); std::this_thread::sleep_for(std::chrono::milliseconds(500)); *notified1 = 1; @@ -389,7 +389,7 @@ TEST_CASE("Unit_hipMemPoolSetAttribute_EventDependencies") { HIP_CHECK(hipMemPoolSetAttribute(mempool.mempool(), attr, &value)); // Run kernel in the first stream - notifiedKernel<<<32, blocks, 0, stream1.stream()>>>(alloc_mem1, notified1); + notifiedKernel<<>>(alloc_mem1, notified1); // Not a real free, since kernel isn't done HIP_CHECK(hipFreeAsync(reinterpret_cast(alloc_mem1), stream1.stream())); HIP_CHECK(hipEventRecord(event, stream1.stream())); @@ -403,7 +403,7 @@ TEST_CASE("Unit_hipMemPoolSetAttribute_EventDependencies") { REQUIRE(alloc_mem1 != alloc_mem2); // Run kernel with the new memory in the second stream - notifiedKernel<<<32, blocks, 0, stream2.stream()>>>(alloc_mem2, notified2); + notifiedKernel<<>>(alloc_mem2, notified2); std::this_thread::sleep_for(std::chrono::milliseconds(500)); *notified1 = 1; diff --git a/projects/hip-tests/catch/unit/memory/hipMemPoolTrimTo.cc b/projects/hip-tests/catch/unit/memory/hipMemPoolTrimTo.cc index 622839477c..dde9bd4aaa 100644 --- a/projects/hip-tests/catch/unit/memory/hipMemPoolTrimTo.cc +++ b/projects/hip-tests/catch/unit/memory/hipMemPoolTrimTo.cc @@ -87,7 +87,7 @@ TEST_CASE("Unit_hipMemPoolTrimTo_Positive_Basic") { mempool.mempool(), stream.stream())); int blocks = 2; - notifiedKernel<<<32, blocks, 0, stream.stream()>>>(alloc_mem1, notified); + notifiedKernel<<>>(alloc_mem1, notified); hipMemPoolAttr attr; attr = hipMemPoolAttrReleaseThreshold; diff --git a/projects/hip-tests/catch/unit/memory/mempool_common.hh b/projects/hip-tests/catch/unit/memory/mempool_common.hh index 8720a386ed..8baf1a2690 100644 --- a/projects/hip-tests/catch/unit/memory/mempool_common.hh +++ b/projects/hip-tests/catch/unit/memory/mempool_common.hh @@ -124,7 +124,7 @@ template void MallocMemPoolAsync_OneAlloc(F malloc_func, const MemP int blocks = 16; hipMemPoolAttr attr; - notifiedKernel<<<32, blocks, 0, stream.stream()>>>(alloc_mem, notified); + notifiedKernel<<>>(alloc_mem, notified); const auto element_count = allocation_size / sizeof(int); constexpr auto thread_count = 1024; @@ -187,7 +187,7 @@ void MallocMemPoolAsync_TwoAllocs(F malloc_func, const MemPools mempool_type) { int blocks = 16; hipMemPoolAttr attr; - notifiedKernel<<<32, blocks, 0, stream.stream()>>>(alloc_mem1, notified); + notifiedKernel<<>>(alloc_mem1, notified); const auto element_count = allocation_size / sizeof(int); constexpr auto thread_count = 1024; @@ -267,7 +267,7 @@ template void MallocMemPoolAsync_Reuse(F malloc_func, const MemPool int blocks = 2; - notifiedKernel<<<32, blocks, 0, stream.stream()>>>(alloc_mem1, notified); + notifiedKernel<<>>(alloc_mem1, notified); hipMemPoolAttr attr; // Not a real free, since kernel isn't done @@ -283,7 +283,7 @@ template void MallocMemPoolAsync_Reuse(F malloc_func, const MemPool HIP_CHECK(hipStreamSynchronize(stream.stream())); *notified = 0; // Second kernel launch with new memory - notifiedKernel<<<32, blocks, 0, stream.stream()>>>(alloc_mem2, notified); + notifiedKernel<<>>(alloc_mem2, notified); *notified = 1; HIP_CHECK(hipStreamSynchronize(stream.stream()));