From dbaa53dac3c389223853c5864fc48d3d6cd49bea 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 --- .../hipTestMain/config/config_amd_windows | 1 - .../config/config_nvidia_linux.json | 14 ----------- .../config/config_nvidia_windows.json | 1 - .../catch/unit/device/hipDeviceSetGetLimit.cc | 8 +++++++ .../catch/unit/device/hipDeviceSynchronize.cc | 12 ---------- .../catch/unit/device/hipGetSetDeviceFlags.cc | 14 ++++------- .../unit/graph/hipGraphAddChildGraphNode.cc | 14 ++++++----- .../graph/hipGraphInstantiateWithParams.cc | 6 ++++- .../graph/hipGraphMemAllocNodeGetParams.cc | 17 +++++++++++-- .../unit/memory/hipFreeMipmappedArray.cc | 24 ++----------------- .../unit/memory/hipMallocMipmappedArray.cc | 15 ++++++++---- .../unit/memory/hipMemPoolSetGetAccess.cc | 15 +++++++++--- .../catch/unit/module/hipModuleGetTexRef.cc | 4 +++- 13 files changed, 68 insertions(+), 77 deletions(-) diff --git a/projects/hip-tests/catch/hipTestMain/config/config_amd_windows b/projects/hip-tests/catch/hipTestMain/config/config_amd_windows index 6f61bb0a9be..4017b037fa9 100644 --- a/projects/hip-tests/catch/hipTestMain/config/config_amd_windows +++ b/projects/hip-tests/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_Thread_Block_Tile_Dynamic_Getters_Positive_Basic", "Performance_hipMemcpy2D_HostToHost", diff --git a/projects/hip-tests/catch/hipTestMain/config/config_nvidia_linux.json b/projects/hip-tests/catch/hipTestMain/config/config_nvidia_linux.json index 636f1d006a5..52cb73c2965 100644 --- a/projects/hip-tests/catch/hipTestMain/config/config_nvidia_linux.json +++ b/projects/hip-tests/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,20 +53,12 @@ "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", "Unit_hipDestroySurfaceObject_Negative_Parameters", "Unit_hipMemcpy2D_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", @@ -224,12 +215,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/projects/hip-tests/catch/hipTestMain/config/config_nvidia_windows.json b/projects/hip-tests/catch/hipTestMain/config/config_nvidia_windows.json index 8c18107be88..a291f0eb638 100644 --- a/projects/hip-tests/catch/hipTestMain/config/config_nvidia_windows.json +++ b/projects/hip-tests/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/projects/hip-tests/catch/unit/device/hipDeviceSetGetLimit.cc b/projects/hip-tests/catch/unit/device/hipDeviceSetGetLimit.cc index aed5fd8bc27..5f70c186665 100644 --- a/projects/hip-tests/catch/unit/device/hipDeviceSetGetLimit.cc +++ b/projects/hip-tests/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/projects/hip-tests/catch/unit/device/hipDeviceSynchronize.cc b/projects/hip-tests/catch/unit/device/hipDeviceSynchronize.cc index 26e661330a3..399e8765f93 100644 --- a/projects/hip-tests/catch/unit/device/hipDeviceSynchronize.cc +++ b/projects/hip-tests/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/projects/hip-tests/catch/unit/device/hipGetSetDeviceFlags.cc b/projects/hip-tests/catch/unit/device/hipGetSetDeviceFlags.cc index 67cbd32d1e8..4e605d46d5e 100644 --- a/projects/hip-tests/catch/unit/device/hipGetSetDeviceFlags.cc +++ b/projects/hip-tests/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/projects/hip-tests/catch/unit/graph/hipGraphAddChildGraphNode.cc b/projects/hip-tests/catch/unit/graph/hipGraphAddChildGraphNode.cc index bdd0ab7f605..df774408a75 100644 --- a/projects/hip-tests/catch/unit/graph/hipGraphAddChildGraphNode.cc +++ b/projects/hip-tests/catch/unit/graph/hipGraphAddChildGraphNode.cc @@ -119,22 +119,24 @@ 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; 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, graph)); + 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/projects/hip-tests/catch/unit/graph/hipGraphInstantiateWithParams.cc b/projects/hip-tests/catch/unit/graph/hipGraphInstantiateWithParams.cc index ea6ee1b8dc4..20071b2563c 100644 --- a/projects/hip-tests/catch/unit/graph/hipGraphInstantiateWithParams.cc +++ b/projects/hip-tests/catch/unit/graph/hipGraphInstantiateWithParams.cc @@ -78,8 +78,12 @@ TEST_CASE("Unit_hipGraphInstantiateWithParams_Negative") { hipGraphExec_t graphExec; hipGraphInstantiateParams params; params.flags = 10; - REQUIRE(hipGraphInstantiateWithParams(&graphExec, graph, ¶ms) == hipErrorInvalidValue); + // Cuda segfaults here! +#if HT_AMD + REQUIRE(hipGraphInstantiateWithParams(&graphExec, + graph, ¶ms) == hipErrorInvalidValue); REQUIRE(params.result_out == hipGraphInstantiateError); +#endif } } diff --git a/projects/hip-tests/catch/unit/graph/hipGraphMemAllocNodeGetParams.cc b/projects/hip-tests/catch/unit/graph/hipGraphMemAllocNodeGetParams.cc index 618381ce388..ef0a596b6e5 100644 --- a/projects/hip-tests/catch/unit/graph/hipGraphMemAllocNodeGetParams.cc +++ b/projects/hip-tests/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/projects/hip-tests/catch/unit/memory/hipFreeMipmappedArray.cc b/projects/hip-tests/catch/unit/memory/hipFreeMipmappedArray.cc index fc87b6546c6..4f7e6d1fc51 100644 --- a/projects/hip-tests/catch/unit/memory/hipFreeMipmappedArray.cc +++ b/projects/hip-tests/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/projects/hip-tests/catch/unit/memory/hipMallocMipmappedArray.cc b/projects/hip-tests/catch/unit/memory/hipMallocMipmappedArray.cc index 7fdbee81384..d7c1fce9c8f 100644 --- a/projects/hip-tests/catch/unit/memory/hipMallocMipmappedArray.cc +++ b/projects/hip-tests/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/projects/hip-tests/catch/unit/memory/hipMemPoolSetGetAccess.cc b/projects/hip-tests/catch/unit/memory/hipMemPoolSetGetAccess.cc index abfe0ddcb28..bbad79748ee 100644 --- a/projects/hip-tests/catch/unit/memory/hipMemPoolSetGetAccess.cc +++ b/projects/hip-tests/catch/unit/memory/hipMemPoolSetGetAccess.cc @@ -260,32 +260,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/projects/hip-tests/catch/unit/module/hipModuleGetTexRef.cc b/projects/hip-tests/catch/unit/module/hipModuleGetTexRef.cc index 651856bf31f..ddeb1b34bf7 100644 --- a/projects/hip-tests/catch/unit/module/hipModuleGetTexRef.cc +++ b/projects/hip-tests/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 70662e88637b89d3c81652e8d734bbd8375f8848 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 --- projects/hip-tests/catch/unit/module/hipModuleGetTexRef.cc | 6 ++++-- 1 file changed, 4 insertions(+), 2 deletions(-) diff --git a/projects/hip-tests/catch/unit/module/hipModuleGetTexRef.cc b/projects/hip-tests/catch/unit/module/hipModuleGetTexRef.cc index ddeb1b34bf7..89a2419e295 100644 --- a/projects/hip-tests/catch/unit/module/hipModuleGetTexRef.cc +++ b/projects/hip-tests/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 234a67f7ca36e7627a435480d1108a22a06a75a2 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/projects/hip-tests/catch/unit/memory/hipMallocMipmappedArray.cc b/projects/hip-tests/catch/unit/memory/hipMallocMipmappedArray.cc index d7c1fce9c8f..aebd4bee169 100644 --- a/projects/hip-tests/catch/unit/memory/hipMallocMipmappedArray.cc +++ b/projects/hip-tests/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 0f9cbee1aea526e50b4d2af5d47d71689f1fd20f 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 --- .../catch/unit/memory/hipMemcpyWithStream.cc | 17 +++++++++++++++++ .../catch/unit/module/hipModuleGetTexRef.cc | 2 +- 2 files changed, 18 insertions(+), 1 deletion(-) diff --git a/projects/hip-tests/catch/unit/memory/hipMemcpyWithStream.cc b/projects/hip-tests/catch/unit/memory/hipMemcpyWithStream.cc index 8a8fd7eed2e..36f5c931681 100644 --- a/projects/hip-tests/catch/unit/memory/hipMemcpyWithStream.cc +++ b/projects/hip-tests/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/projects/hip-tests/catch/unit/module/hipModuleGetTexRef.cc b/projects/hip-tests/catch/unit/module/hipModuleGetTexRef.cc index 89a2419e295..fb712733acf 100644 --- a/projects/hip-tests/catch/unit/module/hipModuleGetTexRef.cc +++ b/projects/hip-tests/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 681d404f2fe2e10fc3bceaaef86b2c1aee91ef65 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 --- .../hipTestMain/config/config_nvidia_linux.json | 1 + .../unit/graph/hipGraphInstantiateWithParams.cc | 12 +++++++----- 2 files changed, 8 insertions(+), 5 deletions(-) diff --git a/projects/hip-tests/catch/hipTestMain/config/config_nvidia_linux.json b/projects/hip-tests/catch/hipTestMain/config/config_nvidia_linux.json index 52cb73c2965..cc97f9e59b4 100644 --- a/projects/hip-tests/catch/hipTestMain/config/config_nvidia_linux.json +++ b/projects/hip-tests/catch/hipTestMain/config/config_nvidia_linux.json @@ -215,6 +215,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/projects/hip-tests/catch/unit/graph/hipGraphInstantiateWithParams.cc b/projects/hip-tests/catch/unit/graph/hipGraphInstantiateWithParams.cc index 20071b2563c..18627a41d50 100644 --- a/projects/hip-tests/catch/unit/graph/hipGraphInstantiateWithParams.cc +++ b/projects/hip-tests/catch/unit/graph/hipGraphInstantiateWithParams.cc @@ -56,7 +56,9 @@ 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") { @@ -69,7 +71,9 @@ TEST_CASE("Unit_hipGraphInstantiateWithParams_Negative") { 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)); } SECTION("Passing invalid flag") { @@ -78,12 +82,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 2ca40bc3ca8c3dc5adc502faf4bceca40a9a5f35 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 | 9 +++------ 2 files changed, 8 insertions(+), 12 deletions(-) diff --git a/projects/hip-tests/catch/unit/graph/hipGraphAddChildGraphNode.cc b/projects/hip-tests/catch/unit/graph/hipGraphAddChildGraphNode.cc index df774408a75..adc6588d6b2 100644 --- a/projects/hip-tests/catch/unit/graph/hipGraphAddChildGraphNode.cc +++ b/projects/hip-tests/catch/unit/graph/hipGraphAddChildGraphNode.cc @@ -131,12 +131,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/projects/hip-tests/catch/unit/graph/hipGraphInstantiateWithParams.cc b/projects/hip-tests/catch/unit/graph/hipGraphInstantiateWithParams.cc index 18627a41d50..cf4fb1c0820 100644 --- a/projects/hip-tests/catch/unit/graph/hipGraphInstantiateWithParams.cc +++ b/projects/hip-tests/catch/unit/graph/hipGraphInstantiateWithParams.cc @@ -56,8 +56,7 @@ 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)); } @@ -71,8 +70,7 @@ TEST_CASE("Unit_hipGraphInstantiateWithParams_Negative") { 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)); } @@ -82,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)); }