AI/rocm-systems
2
0
Förgrening 0
Kod Ärenden Pull-förfrågningar Actions Packages Projekt Släpp Wiki Aktiviteter

SWDEV-534207 fix order of kernel launch parameters when calling notifiedKernel in some tests: kernel<<<gridDim, blockDim>>> instead of kernel<<<blockDim, gridDim>>>. This was causing out of bounds accesses (#1860)

Bläddra i källkod
This commit is contained in:
Gerardo Hernandez
2025-11-25 06:37:47 +00:00
committad av GitHub
förälder 6c31785eaf
incheckning c87014a54c
4 ändrade filer med 29 tillägg och 29 borttagningar
Download Patch File Download Diff File Expand all files Collapse all files
projects/hip-tests/catch/unit/memory/hipMemPoolApi.cc
+12 -12
Visa fil
@@ -127,7 +127,7 @@ TEST_CASE("Unit_hipMemPoolApi_BasicAlloc") {
int blocks = 1024;
hipMemPoolAttr attr;
notifiedKernel<<<32, blocks, 0, stream>>>(B, notified);
notifiedKernel<<<blocks, 32, 0, stream>>>(B, notified);
HIP_CHECK(hipFreeAsync(reinterpret_cast<void*>(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<<<blocks, 32, 0, stream>>>(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<<<blocks, 32, 0, stream>>>(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<<<blocks, 32, 0, stream>>>(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<<<blocks, 32, 0, stream1>>>(A, notified1);
// Not a real free, since kernel isn't done
HIP_CHECK(hipFreeAsync(reinterpret_cast<void*>(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<<<blocks, 32, 0, stream2>>>(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<<<blocks, 32, 0, stream1>>>(A, notified1);
// Not a real free, since kernel isn't done
HIP_CHECK(hipFreeAsync(reinterpret_cast<void*>(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<<<blocks, 32, 0, stream2>>>(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<<<blocks, 32, 0, stream1>>>(A, notified1);
// Not a real free, since kernel isn't done
HIP_CHECK(hipFreeAsync(reinterpret_cast<void*>(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<<<blocks, 32, 0, stream2>>>(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<void**>(&C), numElements * sizeof(float), stream));
int blocks = 2;
notifiedKernel<<<32, blocks, 0, stream>>>(A, notified);
notifiedKernel<<<blocks, 32, 0, stream>>>(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<<<blocks, 32, 0, stream>>>(B, notified);
HIP_CHECK(hipFreeAsync(reinterpret_cast<void*>(B), stream));
projects/hip-tests/catch/unit/memory/hipMemPoolSetGetAttribute.cc
+12 -12
Visa fil
@@ -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<<<blocks, 32, 0, stream1.stream()>>>(alloc_mem1, notified1);
// Not a real free, since kernel isn't done
HIP_CHECK(hipFreeAsync(reinterpret_cast<void*>(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<<<blocks, 32, 0, stream2.stream()>>>(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<<<blocks, 32, 0, stream1.stream()>>>(alloc_mem1, notified1);
// Not a real free, since kernel isn't done
HIP_CHECK(hipFreeAsync(reinterpret_cast<void*>(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<<<blocks, 32, 0, stream1.stream()>>>(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<<<blocks, 32, 0, stream1.stream()>>>(alloc_mem1, notified1);
// Not a real free, since kernel isn't done
HIP_CHECK(hipFreeAsync(reinterpret_cast<void*>(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<<<blocks, 32, 0, stream2.stream()>>>(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<<<blocks, 32, 0, stream1.stream()>>>(alloc_mem1, notified1);
// Not a real free, since kernel isn't done
HIP_CHECK(hipFreeAsync(reinterpret_cast<void*>(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<<<blocks, 32, 0, stream2.stream()>>>(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<<<blocks, 32, 0, stream1.stream()>>>(alloc_mem1, notified1);
// Not a real free, since kernel isn't done
HIP_CHECK(hipFreeAsync(reinterpret_cast<void*>(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<<<blocks, 32, 0, stream2.stream()>>>(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<<<blocks, 32, 0, stream1.stream()>>>(alloc_mem1, notified1);
// Not a real free, since kernel isn't done
HIP_CHECK(hipFreeAsync(reinterpret_cast<void*>(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<<<blocks, 32, 0, stream2.stream()>>>(alloc_mem2, notified2);
std::this_thread::sleep_for(std::chrono::milliseconds(500));
*notified1 = 1;
projects/hip-tests/catch/unit/memory/hipMemPoolTrimTo.cc
+1 -1
Visa fil
@@ -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<<<blocks, 32, 0, stream.stream()>>>(alloc_mem1, notified);
hipMemPoolAttr attr;
attr = hipMemPoolAttrReleaseThreshold;
projects/hip-tests/catch/unit/memory/mempool_common.hh
+4 -4
Visa fil
@@ -124,7 +124,7 @@ template <typename F> void MallocMemPoolAsync_OneAlloc(F malloc_func, const MemP
int blocks = 16;
hipMemPoolAttr attr;
notifiedKernel<<<32, blocks, 0, stream.stream()>>>(alloc_mem, notified);
notifiedKernel<<<blocks, 32, 0, stream.stream()>>>(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<<<blocks, 32, 0, stream.stream()>>>(alloc_mem1, notified);
const auto element_count = allocation_size / sizeof(int);
constexpr auto thread_count = 1024;
@@ -267,7 +267,7 @@ template <typename F> void MallocMemPoolAsync_Reuse(F malloc_func, const MemPool
int blocks = 2;
notifiedKernel<<<32, blocks, 0, stream.stream()>>>(alloc_mem1, notified);
notifiedKernel<<<blocks, 32, 0, stream.stream()>>>(alloc_mem1, notified);
hipMemPoolAttr attr;
// Not a real free, since kernel isn't done
@@ -283,7 +283,7 @@ template <typename F> 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<<<blocks, 32, 0, stream.stream()>>>(alloc_mem2, notified);
*notified = 1;
HIP_CHECK(hipStreamSynchronize(stream.stream()));