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..5be6c577a 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", @@ -229,11 +220,7 @@ "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..d7cc153d4 100644 --- a/catch/unit/graph/hipGraphAddChildGraphNode.cc +++ b/catch/unit/graph/hipGraphAddChildGraphNode.cc @@ -126,23 +126,23 @@ 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/catch/unit/graph/hipGraphInstantiateWithParams.cc b/catch/unit/graph/hipGraphInstantiateWithParams.cc index d9ad3690c..9ce949594 100644 --- a/catch/unit/graph/hipGraphInstantiateWithParams.cc +++ b/catch/unit/graph/hipGraphInstantiateWithParams.cc @@ -56,23 +56,22 @@ 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)); } SECTION("Passing invalid flag") { @@ -81,9 +80,9 @@ 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)); } } 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..79a0c4f4e 100644 --- a/catch/unit/memory/hipMallocMipmappedArray.cc +++ b/catch/unit/memory/hipMallocMipmappedArray.cc @@ -390,10 +390,15 @@ TEST_CASE("Unit_hipMallocMipmappedArray_Negative_NumLevels") { unsigned int numLevels = floor(log2(size)) + 2; hipChannelFormatDesc desc = hipCreateChannelDesc(); - const auto flag = GENERATE(from_range(std::begin(validFlags), std::end(validFlags))); + const auto flag = hipArrayDefault; +#if HT_AMD HIP_CHECK_ERRORS( hipMallocMipmappedArray(&array, &desc, makeMipmappedExtent(flag, size), numLevels, flag), hipErrorInvalidValue, hipErrorNotSupported); +#else + HIP_CHECK( + hipMallocMipmappedArray(&array, &desc, makeMipmappedExtent(flag, size), numLevels, flag)); +#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/module/hipModuleGetTexRef.cc b/catch/unit/module/hipModuleGetTexRef.cc index 651856bf3..fb712733a 100644 --- a/catch/unit/module/hipModuleGetTexRef.cc +++ b/catch/unit/module/hipModuleGetTexRef.cc @@ -24,6 +24,8 @@ THE SOFTWARE. #include #include +#if defined(__HIP_PLATFORM_AMD__) || CUDA_VERSION < CUDA_12000 + static hipModule_t GetModule() { HIP_CHECK(hipFree(nullptr)); static const auto mg = ModuleGuard::LoadModule("get_tex_ref_module.code"); @@ -68,4 +70,6 @@ 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 +} + +#endif