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..cc97f9e59b4 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", @@ -225,11 +216,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/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..adc6588d6b2 100644 --- a/projects/hip-tests/catch/unit/graph/hipGraphAddChildGraphNode.cc +++ b/projects/hip-tests/catch/unit/graph/hipGraphAddChildGraphNode.cc @@ -119,13 +119,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; @@ -134,7 +135,7 @@ TEST_CASE("Unit_hipGraphAddChildGraphNode_OrgGraphAsChildGraph") { 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(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..cf4fb1c0820 100644 --- a/projects/hip-tests/catch/unit/graph/hipGraphInstantiateWithParams.cc +++ b/projects/hip-tests/catch/unit/graph/hipGraphInstantiateWithParams.cc @@ -57,6 +57,7 @@ TEST_CASE("Unit_hipGraphInstantiateWithParams_Negative") { hipGraphInstantiateParams params; HIP_CHECK(hipGraphCreate(&graph, 0)); REQUIRE(hipGraphInstantiateWithParams(nullptr, graph, ¶ms) == hipErrorInvalidValue); + HIP_CHECK(hipGraphDestroy(graph)); } SECTION("Passing nullptr to graph") { @@ -70,6 +71,7 @@ TEST_CASE("Unit_hipGraphInstantiateWithParams_Negative") { HIP_CHECK(hipGraphCreate(&graph, 0)); hipGraphExec_t graphExec; REQUIRE(hipGraphInstantiateWithParams(&graphExec, graph, nullptr) == hipErrorInvalidValue); + HIP_CHECK(hipGraphDestroy(graph)); } SECTION("Passing invalid flag") { @@ -80,6 +82,7 @@ TEST_CASE("Unit_hipGraphInstantiateWithParams_Negative") { params.flags = 10; REQUIRE(hipGraphInstantiateWithParams(&graphExec, graph, ¶ms) == hipErrorInvalidValue); REQUIRE(params.result_out == hipGraphInstantiateError); + HIP_CHECK(hipGraphDestroy(graph)); } } 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..aebd4bee169 100644 --- a/projects/hip-tests/catch/unit/memory/hipMallocMipmappedArray.cc +++ b/projects/hip-tests/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/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/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 651856bf31f..fb712733acf 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 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