diff --git a/test/AllocTests.cpp b/test/AllocTests.cpp index 12800ae443..15178c044f 100644 --- a/test/AllocTests.cpp +++ b/test/AllocTests.cpp @@ -1,5 +1,5 @@ /************************************************************************* - * Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2025 Advanced Micro Devices, Inc. All rights reserved. * * See LICENSE.txt for license information ************************************************************************/ @@ -7,7 +7,7 @@ #include #include #include - + #include "TestBed.hpp" template ncclResult_t ncclCudaMemcpy(float*, float*, size_t); @@ -16,44 +16,47 @@ namespace RcclUnitTesting TEST(Alloc, ncclIbMallocDebugNonZero) { void* ptr = nullptr; size_t size = 4096; - + ncclResult_t result = ncclIbMalloc(&ptr, size); - + EXPECT_EQ(result, ncclSuccess); ASSERT_NE(ptr, nullptr); - + char* char_ptr = static_cast(ptr); for (size_t i = 0; i < size; ++i) { ASSERT_EQ(char_ptr[i], 0); } - + free(ptr); } - + TEST(Alloc, ncclIbMallocDebugZeroSize) { void* ptr = (void*)0xdeadbeef; ncclResult_t result = ncclIbMalloc(&ptr, 0); - + EXPECT_EQ(result, ncclSuccess); EXPECT_EQ(ptr, nullptr); } - - + + TEST(Alloc, ncclCuMemHostAlloc) { void* ptr = NULL; void* handle = NULL; size_t size = 1024; - ncclResult_t result = ncclCuMemHostAlloc(&ptr, handle, size); + ncclResult_t result = ncclCuMemHostAlloc(&ptr, handle, size); ASSERT_EQ(result, ncclInternalError); } TEST(Alloc, ncclCuMemHostFree) { void* dummyPtr = reinterpret_cast(0x1234); // any dummy address - ncclResult_t result = ncclCuMemHostFree(dummyPtr); + ncclResult_t result = ncclCuMemHostFree(dummyPtr); ASSERT_EQ(result, ncclInternalError); } - + +#if ROCM_VERSION < 70000 + // This test is only valid for ROCm versions < 7.0.0 + // In ROCm 7.0.0+, the ncclCuMemAlloc signature changed TEST(Alloc, ncclCuMemAlloc) { void* ptr = reinterpret_cast(0x1234); // dummy non-null input @@ -61,14 +64,14 @@ namespace RcclUnitTesting size_t size = 1024; hipMemAllocationHandleType type = hipMemHandleTypeNone; ncclResult_t result = ncclCuMemAlloc(&ptr, &handle, type, size); - EXPECT_EQ(result, ncclInternalError); + EXPECT_EQ(result, ncclInternalError); } TEST(Alloc, ncclCuMemFree) { void* dummyPtr = reinterpret_cast(0xdeadbeef); // arbitrary non-null ncclResult_t result = ncclCuMemFree(dummyPtr); - EXPECT_EQ(result, ncclInternalError); + EXPECT_EQ(result, ncclInternalError); } TEST(Alloc, ncclCuMemAllocAddr) @@ -76,39 +79,40 @@ namespace RcclUnitTesting void* ptr = reinterpret_cast(0x1111); // Dummy non-null input hipMemGenericAllocationHandle_t handle = reinterpret_cast(0x1234); size_t size = 4096; - ncclResult_t result = ncclCuMemAllocAddr(&ptr, &handle, size); + ncclResult_t result = ncclCuMemAllocAddr(&ptr, &handle, size); ASSERT_EQ(result, ncclInternalError); } - + TEST(Alloc, ncclCuMemFreeAddr) { void* testPtr = reinterpret_cast(0xbeefcafe); // Arbitrary non-null pointer ncclResult_t result = ncclCuMemFreeAddr(testPtr); - ASSERT_EQ(result, ncclInternalError); + ASSERT_EQ(result, ncclInternalError); } - +#endif // ROCM_VERSION < 70000 + TEST(Alloc, NcclCudaMemcpy) { constexpr size_t N = 128; float *d_src = nullptr, *d_dst = nullptr; float h_src[N], h_dst[N]; - + for (size_t i = 0; i < N; ++i) h_src[i] = static_cast(i + 1); // Allocate device memory - + ASSERT_EQ(hipMalloc(&d_src, N * sizeof(float)), hipSuccess); ASSERT_EQ(hipMalloc(&d_dst, N * sizeof(float)), hipSuccess); - + // Copy from host to device (source buffer) ASSERT_EQ(hipMemcpy(d_src, h_src, N * sizeof(float), hipMemcpyHostToDevice), hipSuccess); - + // Perform the tested function ncclResult_t result = ncclCudaMemcpy(d_dst, d_src, N); - + ASSERT_EQ(result, ncclSuccess); // Fixed typo: was ncclSsuccess - + // Copy result back to host ASSERT_EQ(hipMemcpy(h_dst, d_dst, N * sizeof(float), hipMemcpyDeviceToHost), hipSuccess); - + // Check correctness for (size_t i = 0; i < N; ++i) { EXPECT_EQ(h_src[i], h_dst[i]) << "Mismatch at index " << i; @@ -116,7 +120,7 @@ namespace RcclUnitTesting // Free memory hipFree(d_src); hipFree(d_dst); - + } TEST(Alloc, ZeroElementMemcpy) {