AI/rocm-systems
2
0
Rozštěpit 0
Zdrojový kód Úkoly Pull requesty Akce Balíčky Projekty Vydání Wiki Aktivita

SWDEV-454895 - Fix for Atomic Memory Order testcase failures

Procházet zdrojové kódy 
Change-Id: I66f92b57527c364b18a695bc9475f4c3432e742b
Tento commit je obsažen v:
Rahul Manocha
2024-04-03 13:44:03 -07:00
odevzdal Rahul Manocha
rodič ef9558a160
revize 5672b0d380
4 změnil soubory, kde provedl 79 přidání a 85 odebrání
Stáhněte soubor záplaty Stáhněte rozdílový soubor Rozbalit všechny soubory Sbalit všechny soubory
catch/hipTestMain/config/config_amd_linux
-20
Zobrazit soubor
@@ -911,26 +911,6 @@
"Unit_Device___longlong_as_double_Negative_RTC",
"Unit_Device___hiloint2double_Positive",
"Unit_Device___hiloint2double_Negative_RTC",
"Unit___hip_atomic_load_store_Positive_Acquire_Release",
"Unit___hip_atomic_exchange_Positive_Acquire_Release",
"Unit___hip_atomic_compare_exchange_strong_Positive_Acquire_Release",
"Unit___hip_atomic_compare_exchange_weak_Positive_Acquire_Release",
"Unit___hip_atomic_fetch_add_Positive_Acquire_Release",
"Unit___hip_atomic_fetch_and_Positive_Acquire_Release",
"Unit___hip_atomic_fetch_or_Positive_Acquire_Release",
"Unit___hip_atomic_fetch_xor_Positive_Acquire_Release",
"Unit___hip_atomic_fetch_min_Positive_Acquire_Release",
"Unit___hip_atomic_fetch_max_Positive_Acquire_Release",
"Unit___hip_atomic_load_store_Positive_Sequential_Consistency",
"Unit___hip_atomic_exchange_Positive_Sequential_Consistency",
"Unit___hip_atomic_compare_exchange_strong_Positive_Sequential_Consistency",
"Unit___hip_atomic_compare_exchange_weak_Positive_Sequential_Consistency",
"Unit___hip_atomic_fetch_add_Positive_Sequential_Consistency",
"Unit___hip_atomic_fetch_and_Positive_Sequential_Consistency",
"Unit___hip_atomic_fetch_or_Positive_Sequential_Consistency",
"Unit___hip_atomic_fetch_xor_Positive_Sequential_Consistency",
"Unit___hip_atomic_fetch_min_Positive_Sequential_Consistency",
"Unit___hip_atomic_fetch_max_Positive_Sequential_Consistency",
"Unit_atomicAdd_Negative_Parameters_RTC",
"Unit_atomicSub_Negative_Parameters_RTC",
"Unit_atomicInc_Negative_Parameters_RTC",
catch/hipTestMain/config/config_amd_windows
-20
Zobrazit soubor
@@ -1094,26 +1094,6 @@
"Unit_Device___dmul_rn_Accuracy_Positive",
"Unit_Device___ddiv_rn_Accuracy_Positive",
"Unit_Device___fma_rn_Accuracy_Positive",
"Unit___hip_atomic_load_store_Positive_Acquire_Release",
"Unit___hip_atomic_exchange_Positive_Acquire_Release",
"Unit___hip_atomic_compare_exchange_strong_Positive_Acquire_Release",
"Unit___hip_atomic_compare_exchange_weak_Positive_Acquire_Release",
"Unit___hip_atomic_fetch_add_Positive_Acquire_Release",
"Unit___hip_atomic_fetch_and_Positive_Acquire_Release",
"Unit___hip_atomic_fetch_or_Positive_Acquire_Release",
"Unit___hip_atomic_fetch_xor_Positive_Acquire_Release",
"Unit___hip_atomic_fetch_min_Positive_Acquire_Release",
"Unit___hip_atomic_fetch_max_Positive_Acquire_Release",
"Unit___hip_atomic_load_store_Positive_Sequential_Consistency",
"Unit___hip_atomic_exchange_Positive_Sequential_Consistency",
"Unit___hip_atomic_compare_exchange_strong_Positive_Sequential_Consistency",
"Unit___hip_atomic_compare_exchange_weak_Positive_Sequential_Consistency",
"Unit___hip_atomic_fetch_add_Positive_Sequential_Consistency",
"Unit___hip_atomic_fetch_and_Positive_Sequential_Consistency",
"Unit___hip_atomic_fetch_or_Positive_Sequential_Consistency",
"Unit___hip_atomic_fetch_xor_Positive_Sequential_Consistency",
"Unit___hip_atomic_fetch_min_Positive_Sequential_Consistency",
"Unit___hip_atomic_fetch_max_Positive_Sequential_Consistency",
"Unit_atomicAdd_Negative_Parameters_RTC",
"Unit_atomicSub_Negative_Parameters_RTC",
"Unit_atomicInc_Negative_Parameters_RTC",
catch/unit/atomics/memory_order_common.hh
+55 -37
Zobrazit soubor
@@ -129,7 +129,6 @@ __host__ __device__ void Producer(int* const flag, int* const data) {
memory_order == __ATOMIC_ACQUIRE ? __ATOMIC_RELEASE : memory_order;
data[0] = kTestValue;
SetFlag<operation, actual_memory_order, memory_scope>(flag);
}
@@ -170,12 +169,10 @@ __global__ void TestKernel(int* const flag, int* data, int* const ret) {
if (producer) {
Producer<operation, memory_order, memory_scope>(flag, data);
return;
}
if (consumer) {
Consumer<operation, memory_order, memory_scope>(flag, data, ret);
return;
}
}
@@ -216,7 +213,7 @@ template <BuiltinAtomicOperation operation, int memory_order, int memory_scope>
LinearAllocGuard<int> ret(LinearAllocs::hipMallocManaged, sizeof(int));
SECTION("Global memory") {
const auto alloc_type = GENERATE(LinearAllocs::hipMalloc, LinearAllocs::hipMallocManaged);
const auto alloc_type = LinearAllocs::hipMalloc;
LinearAllocGuard<int> data(alloc_type, sizeof(int));
TestKernel<operation, memory_order, memory_scope>
<<<blocks, threads>>>(flag.ptr(), data.ptr(), ret.ptr());
@@ -235,27 +232,33 @@ template <BuiltinAtomicOperation operation, int memory_order, int memory_scope>
}
template <BuiltinAtomicOperation operation, int memory_order> void SystemTest() {
HipTest::HIP_SKIP_TEST("Skip system scope tests due to random failures!!");
return;
std::thread host_thread;
LinearAllocGuard<int> flag(LinearAllocs::hipMallocManaged, sizeof(int));
LinearAllocGuard<int> ret(LinearAllocs::hipMallocManaged, sizeof(int));
SECTION("Global memory") {
const auto alloc_type = GENERATE(LinearAllocs::hipHostMalloc, LinearAllocs::hipMallocManaged);
const auto alloc_type = GENERATE(LinearAllocs::hipHostMalloc , LinearAllocs::hipMallocManaged);
LinearAllocGuard<int> data(alloc_type, sizeof(int));
if constexpr(operation == BuiltinAtomicOperation::kAnd) {
flag.ptr()[0] = 1;
}
SECTION("Host producer - Device consumer") {
host_thread = std::thread([&] {
Producer<operation, memory_order, __HIP_MEMORY_SCOPE_SYSTEM>(flag.host_ptr(), data.host_ptr());
});
ConsumerKernel<operation, memory_order, __HIP_MEMORY_SCOPE_SYSTEM>
<<<1, 1>>>(flag.ptr(), data.ptr(), ret.ptr());
host_thread = std::thread([&] {
Producer<operation, memory_order, __HIP_MEMORY_SCOPE_SYSTEM>(flag.ptr(), data.ptr());
});
}
SECTION("Device producer - Host consumer") {
host_thread = std::thread([&] {
Consumer<operation, memory_order, __HIP_MEMORY_SCOPE_SYSTEM>(flag.ptr(), data.ptr(),
ret.ptr());
Consumer<operation, memory_order, __HIP_MEMORY_SCOPE_SYSTEM>(flag.host_ptr(), data.host_ptr(),
ret.host_ptr());
});
ProducerKernel<operation, memory_order, __HIP_MEMORY_SCOPE_SYSTEM>
<<<1, 1>>>(flag.ptr(), data.ptr());
@@ -274,24 +277,27 @@ namespace SequentialConsistency {
template <BuiltinAtomicOperation operation, int memory_scope>
__host__ __device__ void Producer(int* const flag) {
__atomic_store_n(flag, 1, __ATOMIC_SEQ_CST);
if constexpr (operation == BuiltinAtomicOperation::kAnd) {
__atomic_store_n(flag, 0, __ATOMIC_SEQ_CST);
}
else {
__atomic_store_n(flag, 1, __ATOMIC_SEQ_CST);
}
}
template <BuiltinAtomicOperation operation, int memory_scope>
__host__ __device__ void Consumer(int* const flag1, int* const flag2, int* const counter) {
while (!FetchFlag<operation, __ATOMIC_SEQ_CST, memory_scope>(flag1))
;
if (FetchFlag<operation, __ATOMIC_SEQ_CST, memory_scope>(flag2)) {
__host__ __device__ void Consumer(int* const flag1, int* const counter) {
while (!FetchFlag<operation, __ATOMIC_SEQ_CST, memory_scope>(flag1)) {};
#ifdef __HIP_DEVICE_COMPILE__
__hip_atomic_fetch_add(counter, 1, __ATOMIC_SEQ_CST, memory_scope);
#else
__atomic_fetch_add(counter, 1, __ATOMIC_SEQ_CST);
#endif
}
}
template <BuiltinAtomicOperation operation, int memory_scope>
__global__ void TestKernel(int* flag1, int* flag2, int* const counter) {
__global__ void TestKernel(int* flag1, int* flag2, int* const counter1, int* const counter2) {
__shared__ int shared_mem[2];
if (flag1 == nullptr) flag1 = &shared_mem[0];
@@ -329,22 +335,18 @@ __global__ void TestKernel(int* flag1, int* flag2, int* const counter) {
if (producer1) {
Producer<operation, memory_scope>(flag1);
return;
}
if (consumer1) {
Consumer<operation, memory_scope>(flag1, flag2, counter);
return;
Consumer<operation, memory_scope>(flag1, counter1);
}
if (producer2) {
Producer<operation, memory_scope>(flag2);
return;
}
if (consumer2) {
Consumer<operation, memory_scope>(flag2, flag1, counter);
return;
Consumer<operation, memory_scope>(flag2, counter2);
}
}
@@ -358,12 +360,12 @@ __global__ void ProducerKernel(int* const flag) {
}
template <BuiltinAtomicOperation operation, int memory_scope>
__global__ void ConsumerKernel(int* const flag1, int* const flag2, int* const counter) {
__global__ void ConsumerKernel(int* const flag1, int* const counter) {
if (!(blockIdx.x == 0 && threadIdx.x == 0)) {
return;
}
Consumer<operation, memory_scope>(flag1, flag2, counter);
Consumer<operation, memory_scope>(flag1, counter);
}
template <BuiltinAtomicOperation operation, int memory_scope> void Test() {
@@ -381,53 +383,69 @@ template <BuiltinAtomicOperation operation, int memory_scope> void Test() {
threads = 1;
}
LinearAllocGuard<int> counter(LinearAllocs::hipMallocManaged, sizeof(int));
LinearAllocGuard<int> counter1(LinearAllocs::hipMallocManaged, sizeof(int));
LinearAllocGuard<int> counter2(LinearAllocs::hipMallocManaged, sizeof(int));
SECTION("Global memory") {
const auto alloc_type = GENERATE(LinearAllocs::hipMalloc);
const auto alloc_type = LinearAllocs::hipMalloc;
LinearAllocGuard<int> flag1(alloc_type, sizeof(int));
LinearAllocGuard<int> flag2(alloc_type, sizeof(int));
TestKernel<operation, memory_scope>
<<<blocks, threads>>>(flag1.ptr(), flag2.ptr(), counter.ptr());
<<<blocks, threads>>>(flag1.ptr(), flag2.ptr(), counter1.ptr(), counter2.ptr());
}
if (memory_scope != __HIP_MEMORY_SCOPE_AGENT && memory_scope != __HIP_MEMORY_SCOPE_SYSTEM) {
SECTION("Shared memory") {
TestKernel<operation, memory_scope><<<blocks, threads>>>(nullptr, nullptr, counter.ptr());
TestKernel<operation, memory_scope><<<blocks, threads>>>(nullptr, nullptr, counter1.ptr(), counter2.ptr());
}
}
HIP_CHECK(hipDeviceSynchronize());
REQUIRE(counter.ptr()[0] != 0);
REQUIRE(counter1.ptr()[0] != 0);
REQUIRE(counter2.ptr()[0] != 0);
}
template <BuiltinAtomicOperation operation> void SystemTest() {
HipTest::HIP_SKIP_TEST("Skip system scope tests due to random failures!!");
return;
std::thread host_producer, host_consumer;
LinearAllocGuard<int> counter(LinearAllocs::hipMallocManaged, sizeof(int));
LinearAllocGuard<int> counter1(LinearAllocs::hipMallocManaged, sizeof(int));
LinearAllocGuard<int> counter2(LinearAllocs::hipMallocManaged, sizeof(int));
std::vector<StreamGuard> streams;
for (auto j = 0; j < 2; ++j) {
streams.emplace_back(Streams::created);
}
SECTION("Global memory") {
const auto alloc_type = GENERATE(LinearAllocs::hipMallocManaged);
const auto alloc_type = LinearAllocs::hipMallocManaged;
LinearAllocGuard<int> flag1(alloc_type, sizeof(int));
LinearAllocGuard<int> flag2(alloc_type, sizeof(int));
const auto &stream1 = streams[0].stream();
ConsumerKernel<operation, __HIP_MEMORY_SCOPE_SYSTEM>
<<<1, 1>>>(flag1.ptr(), flag2.ptr(), counter.ptr());
<<<1, 1, 0, stream1>>>(flag1.ptr(), counter1.ptr());
host_consumer = std::thread([&] {
Consumer<operation, __HIP_MEMORY_SCOPE_SYSTEM>(flag2.ptr(), flag1.ptr(), counter.ptr());
Consumer<operation, __HIP_MEMORY_SCOPE_SYSTEM>(flag2.ptr(), counter2.ptr());
});
ProducerKernel<operation, __HIP_MEMORY_SCOPE_SYSTEM><<<1, 1>>>(flag1.ptr());
const auto &stream2 = streams[1].stream();
ProducerKernel<operation, __HIP_MEMORY_SCOPE_SYSTEM><<<1, 1, 0 , stream2>>>(flag2.ptr());
host_producer =
std::thread([&] { Producer<operation, __HIP_MEMORY_SCOPE_SYSTEM>(flag2.ptr()); });
std::thread([&] { Producer<operation, __HIP_MEMORY_SCOPE_SYSTEM>(flag1.ptr()); });
}
HIP_CHECK(hipDeviceSynchronize());
host_producer.join();
host_consumer.join();
REQUIRE(counter.ptr()[0] != 0);
REQUIRE(counter1.ptr()[0] != 0);
REQUIRE(counter2.ptr()[0] != 0);
}
} // namespace SequentialConsistency
catch/unit/atomics/sequential_consistency.cc
+24 -8
Zobrazit soubor
@@ -34,7 +34,9 @@ TEST_CASE("Unit___hip_atomic_load_store_Positive_Sequential_Consistency") {
SECTION("AGENT") {
SequentialConsistency::Test<BuiltinAtomicOperation::kLoadStore, __HIP_MEMORY_SCOPE_AGENT>();
}
SECTION("SYSTEM") { SequentialConsistency::SystemTest<BuiltinAtomicOperation::kLoadStore>(); }
SECTION("SYSTEM") {
SequentialConsistency::SystemTest<BuiltinAtomicOperation::kLoadStore>();
}
}
TEST_CASE("Unit___hip_atomic_exchange_Positive_Sequential_Consistency") {
@@ -47,7 +49,9 @@ TEST_CASE("Unit___hip_atomic_exchange_Positive_Sequential_Consistency") {
SECTION("AGENT") {
SequentialConsistency::Test<BuiltinAtomicOperation::kExchange, __HIP_MEMORY_SCOPE_AGENT>();
}
SECTION("SYSTEM") { SequentialConsistency::SystemTest<BuiltinAtomicOperation::kExchange>(); }
SECTION("SYSTEM") {
SequentialConsistency::SystemTest<BuiltinAtomicOperation::kExchange>();
}
}
TEST_CASE("Unit___hip_atomic_compare_exchange_strong_Positive_Sequential_Consistency") {
@@ -96,7 +100,9 @@ TEST_CASE("Unit___hip_atomic_fetch_add_Positive_Sequential_Consistency") {
SECTION("AGENT") {
SequentialConsistency::Test<BuiltinAtomicOperation::kAdd, __HIP_MEMORY_SCOPE_AGENT>();
}
SECTION("SYSTEM") { SequentialConsistency::SystemTest<BuiltinAtomicOperation::kAdd>(); }
SECTION("SYSTEM") {
SequentialConsistency::SystemTest<BuiltinAtomicOperation::kAdd>();
}
}
TEST_CASE("Unit___hip_atomic_fetch_and_Positive_Sequential_Consistency") {
@@ -109,7 +115,9 @@ TEST_CASE("Unit___hip_atomic_fetch_and_Positive_Sequential_Consistency") {
SECTION("AGENT") {
SequentialConsistency::Test<BuiltinAtomicOperation::kAnd, __HIP_MEMORY_SCOPE_AGENT>();
}
SECTION("SYSTEM") { SequentialConsistency::SystemTest<BuiltinAtomicOperation::kAnd>(); }
SECTION("SYSTEM") {
SequentialConsistency::SystemTest<BuiltinAtomicOperation::kAnd>();
}
}
TEST_CASE("Unit___hip_atomic_fetch_or_Positive_Sequential_Consistency") {
@@ -122,7 +130,9 @@ TEST_CASE("Unit___hip_atomic_fetch_or_Positive_Sequential_Consistency") {
SECTION("AGENT") {
SequentialConsistency::Test<BuiltinAtomicOperation::kOr, __HIP_MEMORY_SCOPE_AGENT>();
}
SECTION("SYSTEM") { SequentialConsistency::SystemTest<BuiltinAtomicOperation::kOr>(); }
SECTION("SYSTEM") {
SequentialConsistency::SystemTest<BuiltinAtomicOperation::kOr>();
}
}
TEST_CASE("Unit___hip_atomic_fetch_xor_Positive_Sequential_Consistency") {
@@ -135,7 +145,9 @@ TEST_CASE("Unit___hip_atomic_fetch_xor_Positive_Sequential_Consistency") {
SECTION("AGENT") {
SequentialConsistency::Test<BuiltinAtomicOperation::kXor, __HIP_MEMORY_SCOPE_AGENT>();
}
SECTION("SYSTEM") { SequentialConsistency::SystemTest<BuiltinAtomicOperation::kXor>(); }
SECTION("SYSTEM") {
SequentialConsistency::SystemTest<BuiltinAtomicOperation::kXor>();
}
}
TEST_CASE("Unit___hip_atomic_fetch_min_Positive_Sequential_Consistency") {
@@ -148,7 +160,9 @@ TEST_CASE("Unit___hip_atomic_fetch_min_Positive_Sequential_Consistency") {
SECTION("AGENT") {
SequentialConsistency::Test<BuiltinAtomicOperation::kMin, __HIP_MEMORY_SCOPE_AGENT>();
}
SECTION("SYSTEM") { SequentialConsistency::SystemTest<BuiltinAtomicOperation::kMin>(); }
SECTION("SYSTEM") {
SequentialConsistency::SystemTest<BuiltinAtomicOperation::kMin>();
}
}
TEST_CASE("Unit___hip_atomic_fetch_max_Positive_Sequential_Consistency") {
@@ -161,5 +175,7 @@ TEST_CASE("Unit___hip_atomic_fetch_max_Positive_Sequential_Consistency") {
SECTION("AGENT") {
SequentialConsistency::Test<BuiltinAtomicOperation::kMax, __HIP_MEMORY_SCOPE_AGENT>();
}
SECTION("SYSTEM") { SequentialConsistency::SystemTest<BuiltinAtomicOperation::kMax>(); }
SECTION("SYSTEM") {
SequentialConsistency::SystemTest<BuiltinAtomicOperation::kMax>();
}
}