From ae0312c61ce2db78d16ba3b399bb9487043651d4 Mon Sep 17 00:00:00 2001 From: Marko Arandjelovic Date: Mon, 14 Apr 2025 14:35:59 +0530 Subject: [PATCH 1/6] SWDEV-523137 - Enable and fix failing tests on NV --- catch/hipTestMain/config/config_amd_windows | 1 - .../config/config_nvidia_linux.json | 15 ------------ .../config/config_nvidia_windows.json | 1 - catch/unit/device/hipDeviceSetGetLimit.cc | 8 +++++++ catch/unit/device/hipDeviceSynchronize.cc | 12 ---------- catch/unit/device/hipGetSetDeviceFlags.cc | 14 ++++------- catch/unit/graph/hipGraphAddChildGraphNode.cc | 5 ++-- .../graph/hipGraphInstantiateWithParams.cc | 3 +++ .../graph/hipGraphMemAllocNodeGetParams.cc | 17 +++++++++++-- catch/unit/memory/hipFreeMipmappedArray.cc | 24 ++----------------- catch/unit/memory/hipMallocMipmappedArray.cc | 15 ++++++++---- catch/unit/memory/hipMemPoolSetGetAccess.cc | 15 +++++++++--- catch/unit/memory/hipMemcpyWithStream.cc | 17 ------------- catch/unit/module/hipModuleGetTexRef.cc | 4 +++- 14 files changed, 61 insertions(+), 90 deletions(-) diff --git a/catch/hipTestMain/config/config_amd_windows b/catch/hipTestMain/config/config_amd_windows index 91a6107dc..858faadc8 100644 --- a/catch/hipTestMain/config/config_amd_windows +++ b/catch/hipTestMain/config/config_amd_windows @@ -344,7 +344,6 @@ "Unit_hipMallocMipmappedArray_MultiThread", "Unit_hipMallocMipmappedArray_Negative_InvalidFlags", "Unit_hipGetMipmappedArrayLevel_Negative", - "Unit_hipFreeMipmappedArray_Negative_DoubleFree", "Unit_hipFreeMipmappedArrayMultiTArray - int", "Unit_hipMemcpy3D_Positive_Synchronization_Behavior", "Unit_hipMemcpyParam2D_Positive_Synchronization_Behavior", diff --git a/catch/hipTestMain/config/config_nvidia_linux.json b/catch/hipTestMain/config/config_nvidia_linux.json index fc8a35dd1..29bedf69f 100644 --- a/catch/hipTestMain/config/config_nvidia_linux.json +++ b/catch/hipTestMain/config/config_nvidia_linux.json @@ -31,7 +31,6 @@ "Unit_hipMemcpyParam2D_multiDevice-D2D - long double", "Unit_hipMemcpyParam2DAsync_multiDevice-StreamOnDiffDevice - char", "Unit_hipMemsetFunctional_ZeroValue_hipMemsetD16", - "Unit_hipStreamAttachMemAsync_Negative_Parameters", "hipStreamPerThread_CoopLaunch", "hipCGMultiGridGroupType", "Grid_Group_Getters_Positive_Basic", @@ -54,14 +53,7 @@ "Unit_Device_Complex_hipCfma_Negative_Parameters_RTC", "Unit_Device_make_Complex_Negative_Parameters_RTC", "Unit_Device_Complex_Cast_Negative_Parameters_RTC", - "=== Below 2 tests are disabled due to defect EXSWHTEC-342 ===", - "Unit_hipDeviceSetLimit_Negative_Parameters", - "Unit_hipDeviceGetLimit_Negative_Parameters", "=== Below tests are failing PSDB ===", - "Unit_hipGraphMem_Alloc_Free_NodeGetParams_Functional_3", - "Unit_hipMemPoolSetAccess_Negative_Parameters", - "Unit_hipMallocMipmappedArray_Negative_NumLevels", - "Unit_hipFreeMipmappedArray_Negative_Nullptr", "Unit_hipFreeMipmappedArrayMultiTArray - int", "Unit_hipFreeMipmappedArray_Negative_Parameters", "Unit_hipCreateSurfaceObject_Negative_Parameters", @@ -69,7 +61,6 @@ "Unit_hipMemcpy3D_Positive_Synchronization_Behavior", "Unit_hipMemcpy2D_Positive_Synchronization_Behavior", "Unit_hipDrvMemcpy3D_Positive_Synchronization_Behavior", - "Unit_hipFreeMipmappedArray_Negative_DoubleFree", "Unit_hipModuleLoad_Positive_Basic", "Unit_hipModuleLoad_Negative_Load_From_A_File_That_Is_Not_A_Module", "Unit_hipModuleLoadData_Positive_Basic", @@ -81,7 +72,6 @@ "Performance_hipMemsetD32", "Performance_hipMemsetD32Async", "Unit_hipMemcpyParam2D_Positive_Synchronization_Behavior", - "Unit_hipMemcpy_Positive_Synchronization_Behavior", "Unit_tex1Dfetch_Positive_ReadModeElementType - char", "Unit_tex1Dfetch_Positive_ReadModeElementType - unsigned char", "Unit_tex1Dfetch_Positive_ReadModeElementType - short", @@ -228,12 +218,7 @@ "Unit_hipDeviceSetLimit_Negative_PrintfFifoSize", "Unit_hipDeviceSetLimit_Negative_MallocHeapSize", "=== Disabling tests which no longer behave the same on nvidia platform ===", - "Unit_hipGraphInstantiateWithParams_Negative", - "Unit_hipGraphAddChildGraphNode_OrgGraphAsChildGraph", - "Unit_hipDeviceSynchronize_Positive_Nullstream", - "Unit_hipDeviceSynchronize_Functional", "Unit_hipDeviceReset_Positive_Basic", "Unit_hipDeviceReset_Positive_Threaded", - "Unit_hipModuleGetTexRef_Positive_Basic" ] } diff --git a/catch/hipTestMain/config/config_nvidia_windows.json b/catch/hipTestMain/config/config_nvidia_windows.json index 8c18107be..a291f0eb6 100644 --- a/catch/hipTestMain/config/config_nvidia_windows.json +++ b/catch/hipTestMain/config/config_nvidia_windows.json @@ -26,7 +26,6 @@ "Unit_hipMemcpy3D_Positive_Synchronization_Behavior", "Unit_hipMemcpy2D_Positive_Synchronization_Behavior", "Unit_hipDrvMemcpy3D_Positive_Synchronization_Behavior", - "Unit_hipFreeMipmappedArray_Negative_DoubleFree", "Unit_hipModuleLoad_Positive_Basic", "Unit_hipModuleLoad_Negative_Load_From_A_File_That_Is_Not_A_Module", "Unit_hipModuleLoadData_Positive_Basic", diff --git a/catch/unit/device/hipDeviceSetGetLimit.cc b/catch/unit/device/hipDeviceSetGetLimit.cc index 7071f979e..0e676f89b 100644 --- a/catch/unit/device/hipDeviceSetGetLimit.cc +++ b/catch/unit/device/hipDeviceSetGetLimit.cc @@ -152,7 +152,11 @@ TEST_CASE("Unit_hipDeviceSetLimit_Negative_MallocHeapSize") { * - HIP_VERSION >= 5.3 */ TEST_CASE("Unit_hipDeviceSetLimit_Negative_Parameters") { +#if HT_AMD HIP_CHECK_ERROR(hipDeviceSetLimit(static_cast(-1), 1024), hipErrorUnsupportedLimit); +#else + HIP_CHECK_ERROR(hipDeviceSetLimit(static_cast(-1), 1024), hipErrorInvalidValue); +#endif } /** @@ -186,7 +190,11 @@ TEST_CASE("Unit_hipDeviceGetLimit_Negative_Parameters") { SECTION("unsupported limit") { size_t val; +#if HT_AMD HIP_CHECK_ERROR(hipDeviceGetLimit(&val, static_cast(-1)), hipErrorUnsupportedLimit); +#else + HIP_CHECK_ERROR(hipDeviceGetLimit(&val, static_cast(-1)), hipErrorInvalidValue); +#endif } } diff --git a/catch/unit/device/hipDeviceSynchronize.cc b/catch/unit/device/hipDeviceSynchronize.cc index 221df92f7..fdd3eba73 100644 --- a/catch/unit/device/hipDeviceSynchronize.cc +++ b/catch/unit/device/hipDeviceSynchronize.cc @@ -89,19 +89,15 @@ TEST_CASE("Unit_hipDeviceSynchronize_Positive_Nullstream") { INFO("Current device: " << device); int *A_h = nullptr, *A_d = nullptr; - HipTest::BlockingContext b_context{nullptr}; HIP_CHECK(hipHostMalloc(reinterpret_cast(&A_h), _SIZE, hipHostMallocDefault)); A_h[0] = 1; HIP_CHECK(hipMalloc(reinterpret_cast(&A_d), _SIZE)); HIP_CHECK(hipMemcpyAsync(A_d, A_h, _SIZE, hipMemcpyHostToDevice, NULL)); - b_context.block_stream(); - REQUIRE(b_context.is_blocked()); hipLaunchKernelGGL(HIP_KERNEL_NAME(Iter), dim3(1), dim3(1), 0, NULL, A_d, 1 << 30); HIP_CHECK(hipMemcpyAsync(A_h, A_d, _SIZE, hipMemcpyDeviceToHost, NULL)); REQUIRE(1 << 30 != A_h[0] - 1); - b_context.unblock_stream(); HIP_CHECK(hipDeviceSynchronize()); REQUIRE(1 << 30 == A_h[0] - 1); HIP_CHECK(hipHostFree(A_h)); @@ -124,22 +120,17 @@ TEST_CASE("Unit_hipDeviceSynchronize_Functional") { int* A[NUM_STREAMS]; int* Ad[NUM_STREAMS]; hipStream_t stream[NUM_STREAMS]; - std::vector b_context; - b_context.reserve(NUM_STREAMS); for (int i = 0; i < NUM_STREAMS; i++) { HIP_CHECK(hipHostMalloc(reinterpret_cast(&A[i]), _SIZE, hipHostMallocDefault)); A[i][0] = 1; HIP_CHECK(hipMalloc(reinterpret_cast(&Ad[i]), _SIZE)); HIP_CHECK(hipStreamCreate(&stream[i])); - b_context.emplace_back(HipTest::BlockingContext(stream[i])); } for (int i = 0; i < NUM_STREAMS; i++) { HIP_CHECK(hipMemcpyAsync(Ad[i], A[i], _SIZE, hipMemcpyHostToDevice, stream[i])); } for (int i = 0; i < NUM_STREAMS; i++) { - b_context[i].block_stream(); - REQUIRE(b_context[i].is_blocked()); hipLaunchKernelGGL(HIP_KERNEL_NAME(Iter), dim3(1), dim3(1), 0, stream[i], Ad[i], NUM_ITERS); } for (int i = 0; i < NUM_STREAMS; i++) { @@ -154,9 +145,6 @@ TEST_CASE("Unit_hipDeviceSynchronize_Functional") { // fail, ie if HIP_LAUNCH_BLOCKING=true. REQUIRE(NUM_ITERS != A[NUM_STREAMS - 1][0] - 1); - for (int i = 0; i < NUM_STREAMS; i++) { - b_context[i].unblock_stream(); - } HIP_CHECK(hipDeviceSynchronize()); REQUIRE(NUM_ITERS == A[NUM_STREAMS - 1][0] - 1); for (int i = 0; i < NUM_STREAMS; i++) { diff --git a/catch/unit/device/hipGetSetDeviceFlags.cc b/catch/unit/device/hipGetSetDeviceFlags.cc index b35cc88c8..4e605d46d 100644 --- a/catch/unit/device/hipGetSetDeviceFlags.cc +++ b/catch/unit/device/hipGetSetDeviceFlags.cc @@ -233,15 +233,11 @@ TEST_CASE("Unit_hipGetDeviceFlags_Positive_Context") { * - HIP_VERSION >= 5.2 */ TEST_CASE("Unit_hipGetSetDeviceFlags_InvalidFlag") { -#if HT_AMD - HipTest::HIP_SKIP_TEST("EXSWCPHIPT-115"); - return; -#endif - const unsigned int invalidFlag = GENERATE(0b011, // schedule flags should not overlap - 0b101, // schedule flags should not overlap - 0b110, // schedule flags should not overlap - 0b111, // schedule flags should not overlap - //0b100000, // out of bounds is no longer invalid + const unsigned int invalidFlag = GENERATE(0xb011, // schedule flags should not overlap + 0xb101, // schedule flags should not overlap + 0xb110, // schedule flags should not overlap + 0xb111, // schedule flags should not overlap + 0xb100000, // out of bounds 0xFFFF); CAPTURE(invalidFlag); HIP_CHECK_ERROR(hipSetDeviceFlags(invalidFlag), hipErrorInvalidValue); diff --git a/catch/unit/graph/hipGraphAddChildGraphNode.cc b/catch/unit/graph/hipGraphAddChildGraphNode.cc index 9c8bf4c09..cded1d62d 100644 --- a/catch/unit/graph/hipGraphAddChildGraphNode.cc +++ b/catch/unit/graph/hipGraphAddChildGraphNode.cc @@ -126,13 +126,14 @@ and verify the number of the nodes in the original graph TEST_CASE("Unit_hipGraphAddChildGraphNode_OrgGraphAsChildGraph") { constexpr size_t N = 1024; constexpr size_t Nbytes = N * sizeof(int); - hipGraph_t graph; + hipGraph_t graph, childGraph; hipGraphExec_t graphExec; int *A_d{nullptr}, *B_d{nullptr}; int *A_h{nullptr}, *B_h{nullptr}; HipTest::initArrays(&A_d, &B_d, nullptr, &A_h, &B_h, nullptr, N, false); HIP_CHECK(hipGraphCreate(&graph, 0)); + HIP_CHECK(hipGraphCreate(&childGraph, 0)); hipGraphNode_t memcpyH2D_A, memcpyH2D_B, childGraphNode1; size_t numNodes; hipStream_t streamForGraph; @@ -142,7 +143,7 @@ TEST_CASE("Unit_hipGraphAddChildGraphNode_OrgGraphAsChildGraph") { HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_A, graph, nullptr, 0, A_h, B_d, Nbytes, hipMemcpyDeviceToHost)); HIP_CHECK(hipGraphAddChildGraphNode(&childGraphNode1, graph, - nullptr, 0, graph)); + nullptr, 0, childGraph)); HIP_CHECK(hipGraphAddDependencies(graph, &memcpyH2D_B, &memcpyH2D_A, 1)); diff --git a/catch/unit/graph/hipGraphInstantiateWithParams.cc b/catch/unit/graph/hipGraphInstantiateWithParams.cc index d9ad3690c..50bbe1b9c 100644 --- a/catch/unit/graph/hipGraphInstantiateWithParams.cc +++ b/catch/unit/graph/hipGraphInstantiateWithParams.cc @@ -81,9 +81,12 @@ TEST_CASE("Unit_hipGraphInstantiateWithParams_Negative") { hipGraphExec_t graphExec; hipGraphInstantiateParams params; params.flags = 10; + // Cuda segfaults here! +#if HT_AMD REQUIRE(hipGraphInstantiateWithParams(&graphExec, graph, ¶ms) == hipErrorInvalidValue); REQUIRE(params.result_out == hipGraphInstantiateError); +#endif } } diff --git a/catch/unit/graph/hipGraphMemAllocNodeGetParams.cc b/catch/unit/graph/hipGraphMemAllocNodeGetParams.cc index 47f167a0d..0ec9a16d9 100644 --- a/catch/unit/graph/hipGraphMemAllocNodeGetParams.cc +++ b/catch/unit/graph/hipGraphMemAllocNodeGetParams.cc @@ -51,12 +51,25 @@ THE SOFTWARE. *  - HIP_VERSION >= 6.0 */ -static bool validateAllocParam(hipMemAllocNodeParams in, hipMemAllocNodeParams out) { +static bool validateAllocParam(hipMemAllocNodeParams in, hipMemAllocNodeParams out, + bool accessDesc = false) { if (in.bytesize != out.bytesize) return false; if (in.poolProps.allocType != out.poolProps.allocType) return false; if (in.poolProps.location.id != out.poolProps.location.id) return false; if (in.poolProps.location.type != out.poolProps.location.type) return false; + if (accessDesc) { + if (in.accessDescs->location.type != out.accessDescs->location.type) { + return false; + } + if (in.accessDescs->location.id != out.accessDescs->location.id) { + return false; + } + if (in.accessDescs->flags != out.accessDescs->flags) { + return false; + } + } + return true; } @@ -299,7 +312,7 @@ TEST_CASE("Unit_hipGraphMem_Alloc_Free_NodeGetParams_Functional_3") { hipMemAllocNodeParams get_alloc_params; HIP_CHECK(hipGraphMemAllocNodeGetParams(alloc_node, &get_alloc_params)); - REQUIRE(memcmp(&alloc_param, &get_alloc_params, sizeof(hipMemAllocNodeParams)) == 0); + REQUIRE(validateAllocParam(alloc_param, get_alloc_params, true) == true); constexpr int fill_value = 11; hipGraphNode_t memset_node; diff --git a/catch/unit/memory/hipFreeMipmappedArray.cc b/catch/unit/memory/hipFreeMipmappedArray.cc index fc87b6546..4f7e6d1fc 100644 --- a/catch/unit/memory/hipFreeMipmappedArray.cc +++ b/catch/unit/memory/hipFreeMipmappedArray.cc @@ -80,31 +80,11 @@ TEMPLATE_TEST_CASE("Unit_hipFreeMipmappedArrayImplicitSyncArray", "", char, floa } TEST_CASE("Unit_hipFreeMipmappedArray_Negative_Nullptr") { - HIP_CHECK_ERROR(hipFreeMipmappedArray(nullptr), hipErrorInvalidValue); -} - -TEST_CASE("Unit_hipFreeMipmappedArray_Negative_DoubleFree") { - hipMipmappedArray_t arrayPtr{}; - hipExtent extent{}; - hipChannelFormatDesc desc = hipCreateChannelDesc(); - #if HT_AMD - const unsigned int flags = hipArrayDefault; + HIP_CHECK_ERROR(hipFreeMipmappedArray(nullptr), hipErrorInvalidValue); #else - const unsigned int flags = GENERATE(hipArrayDefault, hipArraySurfaceLoadStore); + HIP_CHECK(hipFreeMipmappedArray(nullptr)); #endif - - extent.width = GENERATE(64, 512, 1024); - extent.height = GENERATE(64, 512, 1024); - extent.depth = GENERATE(0, 64, 512, 1024); - - const unsigned int numLevels = GENERATE(1, 5, 7); - - HIP_CHECK_IGNORED_RETURN(hipMallocMipmappedArray(&arrayPtr, &desc, extent, numLevels, flags), - hipErrorNotSupported); - - HIP_CHECK(hipFreeMipmappedArray(arrayPtr)); - HIP_CHECK_ERROR(hipFreeMipmappedArray(arrayPtr), hipErrorContextIsDestroyed); } TEMPLATE_TEST_CASE("Unit_hipFreeMipmappedArrayMultiTArray", "", char, int) { diff --git a/catch/unit/memory/hipMallocMipmappedArray.cc b/catch/unit/memory/hipMallocMipmappedArray.cc index e68550923..e86ee38e1 100644 --- a/catch/unit/memory/hipMallocMipmappedArray.cc +++ b/catch/unit/memory/hipMallocMipmappedArray.cc @@ -387,13 +387,18 @@ TEMPLATE_TEST_CASE("Unit_hipMallocMipmappedArray_Negative_Non2DTextureGather", " TEST_CASE("Unit_hipMallocMipmappedArray_Negative_NumLevels") { hipMipmappedArray_t array; constexpr size_t size = 6; - unsigned int numLevels = floor(log2(size)) + 2; + unsigned int numLevels = -1; hipChannelFormatDesc desc = hipCreateChannelDesc(); - const auto flag = GENERATE(from_range(std::begin(validFlags), std::end(validFlags))); - HIP_CHECK_ERRORS( - hipMallocMipmappedArray(&array, &desc, makeMipmappedExtent(flag, size), numLevels, flag), - hipErrorInvalidValue, hipErrorNotSupported); + const auto flag = hipArrayDefault; +#if HT_AMD + HIP_CHECK_ERROR(hipMallocMipmappedArray(&array, &desc, makeMipmappedExtent(flag, size), numLevels, + hipArrayDefault), + hipErrorNotSupported); +#else + HIP_CHECK(hipMallocMipmappedArray(&array, &desc, makeMipmappedExtent(flag, size), numLevels, + hipArrayDefault)); +#endif } TEST_CASE("Unit_hipGetMipmappedArrayLevel_Negative") { diff --git a/catch/unit/memory/hipMemPoolSetGetAccess.cc b/catch/unit/memory/hipMemPoolSetGetAccess.cc index 116c297a6..a6d51aa11 100644 --- a/catch/unit/memory/hipMemPoolSetGetAccess.cc +++ b/catch/unit/memory/hipMemPoolSetGetAccess.cc @@ -261,32 +261,41 @@ TEST_CASE("Unit_hipMemPoolSetAccess_Negative_Parameters") { SECTION("Mempool is nullptr") { HIP_CHECK_ERROR(hipMemPoolSetAccess(nullptr, &desc, 1), hipErrorInvalidValue); } + + // Cuda segfaults here! #if HT_AMD SECTION("Desc is nullptr and count is > 0") { HIP_CHECK_ERROR(hipMemPoolSetAccess(mempool.mempool(), nullptr, 1), hipErrorInvalidValue); } #endif + SECTION("Count > num_device") { +#if HT_AMD HIP_CHECK_ERROR(hipMemPoolSetAccess(mempool.mempool(), &desc, (num_dev + 1)), hipErrorInvalidDevice); +#else + HIP_CHECK_ERROR(hipMemPoolSetAccess(mempool.mempool(), &desc, (num_dev + 1)), + hipErrorNotSupported); +#endif } SECTION("Passing invalid desc location type") { desc.location.type = hipMemLocationTypeInvalid; +#if HT_AMD HIP_CHECK_ERROR(hipMemPoolSetAccess(mempool.mempool(), &desc, 1), hipErrorInvalidValue); - desc.location.type = hipMemLocationTypeDevice; +#else + HIP_CHECK_ERROR(hipMemPoolSetAccess(mempool.mempool(), &desc, 1), hipErrorNotSupported); +#endif } SECTION("Passing invalid desc location id") { desc.location.id = num_dev; HIP_CHECK_ERROR(hipMemPoolSetAccess(mempool.mempool(), &desc, 1), hipErrorInvalidDevice); - desc.location.id = device_id; } SECTION("Revoking access to own memory pool") { desc.flags = hipMemAccessFlagsProtNone; HIP_CHECK_ERROR(hipMemPoolSetAccess(mempool.mempool(), &desc, 1), hipErrorInvalidDevice); - desc.flags = hipMemAccessFlagsProtReadWrite; } } diff --git a/catch/unit/memory/hipMemcpyWithStream.cc b/catch/unit/memory/hipMemcpyWithStream.cc index 36f5c9316..8a8fd7eed 100644 --- a/catch/unit/memory/hipMemcpyWithStream.cc +++ b/catch/unit/memory/hipMemcpyWithStream.cc @@ -31,16 +31,6 @@ TEST_CASE("Unit_hipMemcpy_Positive_Synchronization_Behavior") { using namespace std::placeholders; HIP_CHECK(hipDeviceSynchronize()); - // For transfers from pageable host memory to device memory, a stream sync is performed before - // the copy is initiated. The function will return once the pageable buffer has been copied to - // the staging memory for DMA transfer to device memory, but the DMA to final destination may - // not have completed. - // For transfers from pinned host memory to device memory, the function is synchronous with - // respect to the host - SECTION("Host memory to device memory") { - MemcpyHPageabletoDSyncBehavior(std::bind(hipMemcpy, _1, _2, _3, hipMemcpyHostToDevice), true); - } - // For transfers from device to either pageable or pinned host memory, the function returns only // once the copy has completed SECTION("Device memory to host memory") { @@ -51,13 +41,6 @@ TEST_CASE("Unit_hipMemcpy_Positive_Synchronization_Behavior") { // For transfers from device memory to device memory, no host-side synchronization is performed. SECTION("Device memory to device memory") { - // This behavior differs on NVIDIA and AMD, on AMD the hipMemcpy calls is synchronous with - // respect to the host -#if HT_AMD - HipTest::HIP_SKIP_TEST( - "EXSWCPHIPT-127 - Memcpy from device to device memory behavior differs on AMD and Nvidia"); - return; -#endif MemcpyDtoDSyncBehavior(std::bind(hipMemcpy, _1, _2, _3, hipMemcpyDeviceToDevice), false); } diff --git a/catch/unit/module/hipModuleGetTexRef.cc b/catch/unit/module/hipModuleGetTexRef.cc index 651856bf3..ddeb1b34b 100644 --- a/catch/unit/module/hipModuleGetTexRef.cc +++ b/catch/unit/module/hipModuleGetTexRef.cc @@ -32,9 +32,11 @@ static hipModule_t GetModule() { TEST_CASE("Unit_hipModuleGetTexRef_Positive_Basic") { CHECK_IMAGE_SUPPORT + CTX_CREATE(); hipTexRef tex_ref = nullptr; HIP_CHECK(hipModuleGetTexRef(&tex_ref, GetModule(), "tex")); REQUIRE(tex_ref != nullptr); + CTX_DESTROY(); } TEST_CASE("Unit_hipModuleGetTexRef_Negative_Parameters") { @@ -68,4 +70,4 @@ TEST_CASE("Unit_hipModuleGetTexRef_Negative_Name_Is_Empty_String") { hipTexRef tex_ref = nullptr; HIP_CHECK_ERROR(hipModuleGetTexRef(&tex_ref, module, ""), hipErrorInvalidValue); -} \ No newline at end of file +} From 1ff3a6c2e3f665691001dae884465940fb0e2ceb Mon Sep 17 00:00:00 2001 From: Marko Arandjelovic Date: Mon, 14 Apr 2025 19:26:32 +0530 Subject: [PATCH 2/6] SWDEV-523137 - hipModuleGetTexRef is deprecated and is using deprecated texture template --- catch/unit/module/hipModuleGetTexRef.cc | 6 ++++-- 1 file changed, 4 insertions(+), 2 deletions(-) diff --git a/catch/unit/module/hipModuleGetTexRef.cc b/catch/unit/module/hipModuleGetTexRef.cc index ddeb1b34b..89a2419e2 100644 --- a/catch/unit/module/hipModuleGetTexRef.cc +++ b/catch/unit/module/hipModuleGetTexRef.cc @@ -24,6 +24,8 @@ THE SOFTWARE. #include #include +#if CUDA_VERSION < CUDA_12000 + static hipModule_t GetModule() { HIP_CHECK(hipFree(nullptr)); static const auto mg = ModuleGuard::LoadModule("get_tex_ref_module.code"); @@ -32,11 +34,9 @@ static hipModule_t GetModule() { TEST_CASE("Unit_hipModuleGetTexRef_Positive_Basic") { CHECK_IMAGE_SUPPORT - CTX_CREATE(); hipTexRef tex_ref = nullptr; HIP_CHECK(hipModuleGetTexRef(&tex_ref, GetModule(), "tex")); REQUIRE(tex_ref != nullptr); - CTX_DESTROY(); } TEST_CASE("Unit_hipModuleGetTexRef_Negative_Parameters") { @@ -71,3 +71,5 @@ TEST_CASE("Unit_hipModuleGetTexRef_Negative_Name_Is_Empty_String") { HIP_CHECK_ERROR(hipModuleGetTexRef(&tex_ref, module, ""), hipErrorInvalidValue); } + +#endif From 9f321cebca1e3b343801d66fdead594cdf3af857 Mon Sep 17 00:00:00 2001 From: Marko Arandjelovic Date: Tue, 15 Apr 2025 19:04:44 +0530 Subject: [PATCH 3/6] SWDEV-523137 - Fix negative num_levels test --- catch/unit/memory/hipMallocMipmappedArray.cc | 12 ++++++------ 1 file changed, 6 insertions(+), 6 deletions(-) diff --git a/catch/unit/memory/hipMallocMipmappedArray.cc b/catch/unit/memory/hipMallocMipmappedArray.cc index e86ee38e1..79a0c4f4e 100644 --- a/catch/unit/memory/hipMallocMipmappedArray.cc +++ b/catch/unit/memory/hipMallocMipmappedArray.cc @@ -387,17 +387,17 @@ TEMPLATE_TEST_CASE("Unit_hipMallocMipmappedArray_Negative_Non2DTextureGather", " TEST_CASE("Unit_hipMallocMipmappedArray_Negative_NumLevels") { hipMipmappedArray_t array; constexpr size_t size = 6; - unsigned int numLevels = -1; + unsigned int numLevels = floor(log2(size)) + 2; hipChannelFormatDesc desc = hipCreateChannelDesc(); const auto flag = hipArrayDefault; #if HT_AMD - HIP_CHECK_ERROR(hipMallocMipmappedArray(&array, &desc, makeMipmappedExtent(flag, size), numLevels, - hipArrayDefault), - hipErrorNotSupported); + HIP_CHECK_ERRORS( + hipMallocMipmappedArray(&array, &desc, makeMipmappedExtent(flag, size), numLevels, flag), + hipErrorInvalidValue, hipErrorNotSupported); #else - HIP_CHECK(hipMallocMipmappedArray(&array, &desc, makeMipmappedExtent(flag, size), numLevels, - hipArrayDefault)); + HIP_CHECK( + hipMallocMipmappedArray(&array, &desc, makeMipmappedExtent(flag, size), numLevels, flag)); #endif } From 700de8aa84d8d5e3027acd29e5d70300c7c743c0 Mon Sep 17 00:00:00 2001 From: Marko Arandjelovic Date: Thu, 17 Apr 2025 09:43:07 +0000 Subject: [PATCH 4/6] SWDEV-523137 - Small fixes --- .../hipTestMain/config/config_nvidia_linux.json | 1 + catch/unit/memory/hipMemcpyWithStream.cc | 17 +++++++++++++++++ catch/unit/module/hipModuleGetTexRef.cc | 2 +- 3 files changed, 19 insertions(+), 1 deletion(-) diff --git a/catch/hipTestMain/config/config_nvidia_linux.json b/catch/hipTestMain/config/config_nvidia_linux.json index 29bedf69f..f09df7110 100644 --- a/catch/hipTestMain/config/config_nvidia_linux.json +++ b/catch/hipTestMain/config/config_nvidia_linux.json @@ -72,6 +72,7 @@ "Performance_hipMemsetD32", "Performance_hipMemsetD32Async", "Unit_hipMemcpyParam2D_Positive_Synchronization_Behavior", + "Unit_hipMemcpy_Positive_Synchronization_Behavior", "Unit_tex1Dfetch_Positive_ReadModeElementType - char", "Unit_tex1Dfetch_Positive_ReadModeElementType - unsigned char", "Unit_tex1Dfetch_Positive_ReadModeElementType - short", diff --git a/catch/unit/memory/hipMemcpyWithStream.cc b/catch/unit/memory/hipMemcpyWithStream.cc index 8a8fd7eed..36f5c9316 100644 --- a/catch/unit/memory/hipMemcpyWithStream.cc +++ b/catch/unit/memory/hipMemcpyWithStream.cc @@ -31,6 +31,16 @@ TEST_CASE("Unit_hipMemcpy_Positive_Synchronization_Behavior") { using namespace std::placeholders; HIP_CHECK(hipDeviceSynchronize()); + // For transfers from pageable host memory to device memory, a stream sync is performed before + // the copy is initiated. The function will return once the pageable buffer has been copied to + // the staging memory for DMA transfer to device memory, but the DMA to final destination may + // not have completed. + // For transfers from pinned host memory to device memory, the function is synchronous with + // respect to the host + SECTION("Host memory to device memory") { + MemcpyHPageabletoDSyncBehavior(std::bind(hipMemcpy, _1, _2, _3, hipMemcpyHostToDevice), true); + } + // For transfers from device to either pageable or pinned host memory, the function returns only // once the copy has completed SECTION("Device memory to host memory") { @@ -41,6 +51,13 @@ TEST_CASE("Unit_hipMemcpy_Positive_Synchronization_Behavior") { // For transfers from device memory to device memory, no host-side synchronization is performed. SECTION("Device memory to device memory") { + // This behavior differs on NVIDIA and AMD, on AMD the hipMemcpy calls is synchronous with + // respect to the host +#if HT_AMD + HipTest::HIP_SKIP_TEST( + "EXSWCPHIPT-127 - Memcpy from device to device memory behavior differs on AMD and Nvidia"); + return; +#endif MemcpyDtoDSyncBehavior(std::bind(hipMemcpy, _1, _2, _3, hipMemcpyDeviceToDevice), false); } diff --git a/catch/unit/module/hipModuleGetTexRef.cc b/catch/unit/module/hipModuleGetTexRef.cc index 89a2419e2..fb712733a 100644 --- a/catch/unit/module/hipModuleGetTexRef.cc +++ b/catch/unit/module/hipModuleGetTexRef.cc @@ -24,7 +24,7 @@ THE SOFTWARE. #include #include -#if CUDA_VERSION < CUDA_12000 +#if defined(__HIP_PLATFORM_AMD__) || CUDA_VERSION < CUDA_12000 static hipModule_t GetModule() { HIP_CHECK(hipFree(nullptr)); From 3de3fc9125aad73cbac9cec504082a62bc2959d9 Mon Sep 17 00:00:00 2001 From: Marko Arandjelovic Date: Mon, 23 Jun 2025 15:03:09 +0000 Subject: [PATCH 5/6] SWDEV-523137 - Avoid using guards on faulting behavior --- catch/hipTestMain/config/config_nvidia_linux.json | 1 + catch/unit/graph/hipGraphInstantiateWithParams.cc | 6 +++--- 2 files changed, 4 insertions(+), 3 deletions(-) diff --git a/catch/hipTestMain/config/config_nvidia_linux.json b/catch/hipTestMain/config/config_nvidia_linux.json index f09df7110..5be6c577a 100644 --- a/catch/hipTestMain/config/config_nvidia_linux.json +++ b/catch/hipTestMain/config/config_nvidia_linux.json @@ -219,6 +219,7 @@ "Unit_hipDeviceSetLimit_Negative_PrintfFifoSize", "Unit_hipDeviceSetLimit_Negative_MallocHeapSize", "=== Disabling tests which no longer behave the same on nvidia platform ===", + "Unit_hipGraphInstantiateWithParams_Negative", "Unit_hipDeviceReset_Positive_Basic", "Unit_hipDeviceReset_Positive_Threaded", ] diff --git a/catch/unit/graph/hipGraphInstantiateWithParams.cc b/catch/unit/graph/hipGraphInstantiateWithParams.cc index 50bbe1b9c..74c5a796d 100644 --- a/catch/unit/graph/hipGraphInstantiateWithParams.cc +++ b/catch/unit/graph/hipGraphInstantiateWithParams.cc @@ -58,6 +58,7 @@ TEST_CASE("Unit_hipGraphInstantiateWithParams_Negative") { HIP_CHECK(hipGraphCreate(&graph, 0)); REQUIRE(hipGraphInstantiateWithParams(nullptr, graph, ¶ms) == hipErrorInvalidValue); + HIP_CHECK(hipGraphDestroy(graph)); } SECTION("Passing nullptr to graph") { @@ -73,6 +74,7 @@ TEST_CASE("Unit_hipGraphInstantiateWithParams_Negative") { hipGraphExec_t graphExec; REQUIRE(hipGraphInstantiateWithParams(&graphExec, graph, nullptr) == hipErrorInvalidValue); + HIP_CHECK(hipGraphDestroy(graph)); } SECTION("Passing invalid flag") { @@ -81,12 +83,10 @@ TEST_CASE("Unit_hipGraphInstantiateWithParams_Negative") { hipGraphExec_t graphExec; hipGraphInstantiateParams params; params.flags = 10; - // Cuda segfaults here! -#if HT_AMD REQUIRE(hipGraphInstantiateWithParams(&graphExec, graph, ¶ms) == hipErrorInvalidValue); REQUIRE(params.result_out == hipGraphInstantiateError); -#endif + HIP_CHECK(hipGraphDestroy(graph)); } } From 90749dd58ac6fb11201270e11c852565a355429d Mon Sep 17 00:00:00 2001 From: Dragoslav Sicarov Date: Thu, 17 Jul 2025 13:08:41 +0000 Subject: [PATCH 6/6] SWDEV-523137 - Run clang-format --- catch/unit/graph/hipGraphAddChildGraphNode.cc | 11 +++++------ catch/unit/graph/hipGraphInstantiateWithParams.cc | 12 ++++-------- 2 files changed, 9 insertions(+), 14 deletions(-) diff --git a/catch/unit/graph/hipGraphAddChildGraphNode.cc b/catch/unit/graph/hipGraphAddChildGraphNode.cc index cded1d62d..d7cc153d4 100644 --- a/catch/unit/graph/hipGraphAddChildGraphNode.cc +++ b/catch/unit/graph/hipGraphAddChildGraphNode.cc @@ -138,12 +138,11 @@ TEST_CASE("Unit_hipGraphAddChildGraphNode_OrgGraphAsChildGraph") { size_t numNodes; hipStream_t streamForGraph; HIP_CHECK(hipStreamCreate(&streamForGraph)); - HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_B, graph, nullptr, 0, B_d, B_h, - Nbytes, hipMemcpyHostToDevice)); - HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_A, graph, nullptr, 0, A_h, B_d, - Nbytes, hipMemcpyDeviceToHost)); - HIP_CHECK(hipGraphAddChildGraphNode(&childGraphNode1, graph, - nullptr, 0, childGraph)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_B, graph, nullptr, 0, B_d, B_h, Nbytes, + hipMemcpyHostToDevice)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_A, graph, nullptr, 0, A_h, B_d, Nbytes, + hipMemcpyDeviceToHost)); + HIP_CHECK(hipGraphAddChildGraphNode(&childGraphNode1, graph, nullptr, 0, childGraph)); HIP_CHECK(hipGraphAddDependencies(graph, &memcpyH2D_B, &memcpyH2D_A, 1)); diff --git a/catch/unit/graph/hipGraphInstantiateWithParams.cc b/catch/unit/graph/hipGraphInstantiateWithParams.cc index 74c5a796d..9ce949594 100644 --- a/catch/unit/graph/hipGraphInstantiateWithParams.cc +++ b/catch/unit/graph/hipGraphInstantiateWithParams.cc @@ -56,24 +56,21 @@ TEST_CASE("Unit_hipGraphInstantiateWithParams_Negative") { hipGraph_t graph; hipGraphInstantiateParams params; HIP_CHECK(hipGraphCreate(&graph, 0)); - REQUIRE(hipGraphInstantiateWithParams(nullptr, - graph, ¶ms) == hipErrorInvalidValue); + REQUIRE(hipGraphInstantiateWithParams(nullptr, graph, ¶ms) == hipErrorInvalidValue); HIP_CHECK(hipGraphDestroy(graph)); } SECTION("Passing nullptr to graph") { hipGraphExec_t graphExec; hipGraphInstantiateParams params; - REQUIRE(hipGraphInstantiateWithParams(&graphExec, - nullptr, ¶ms) == hipErrorInvalidValue); + REQUIRE(hipGraphInstantiateWithParams(&graphExec, nullptr, ¶ms) == hipErrorInvalidValue); } SECTION("Passing nullptr to params") { hipGraph_t graph; HIP_CHECK(hipGraphCreate(&graph, 0)); hipGraphExec_t graphExec; - REQUIRE(hipGraphInstantiateWithParams(&graphExec, - graph, nullptr) == hipErrorInvalidValue); + REQUIRE(hipGraphInstantiateWithParams(&graphExec, graph, nullptr) == hipErrorInvalidValue); HIP_CHECK(hipGraphDestroy(graph)); } @@ -83,8 +80,7 @@ TEST_CASE("Unit_hipGraphInstantiateWithParams_Negative") { hipGraphExec_t graphExec; hipGraphInstantiateParams params; params.flags = 10; - REQUIRE(hipGraphInstantiateWithParams(&graphExec, - graph, ¶ms) == hipErrorInvalidValue); + REQUIRE(hipGraphInstantiateWithParams(&graphExec, graph, ¶ms) == hipErrorInvalidValue); REQUIRE(params.result_out == hipGraphInstantiateError); HIP_CHECK(hipGraphDestroy(graph)); }