diff --git a/.jenkins/Jenkinsfile b/.jenkins/Jenkinsfile index d64775c026..a4c767be38 100644 --- a/.jenkins/Jenkinsfile +++ b/.jenkins/Jenkinsfile @@ -58,8 +58,8 @@ def hipBuildTest(String backendLabel) { set -x # Check if backend label contains string "amd" or backend host is a server with amd gpu if [[ $backendLabel =~ amd ]]; then - LLVM_PATH=/opt/rocm/llvm ctest -E 'cooperative_streams_least_capacity.tst|cooperative_streams_half_capacity.tst|cooperative_streams_full_capacity.tst|grid_group_data_sharing.tst|hipIpcMemAccessTest.tst|p2p_copy_coherency.tst' sleep 120 + LLVM_PATH=/opt/rocm/llvm ctest -E 'cooperative_streams_least_capacity.tst|cooperative_streams_half_capacity.tst|cooperative_streams_full_capacity.tst|grid_group_data_sharing.tst|hipIpcMemAccessTest.tst|p2p_copy_coherency.tst' else make test fi @@ -100,9 +100,9 @@ def hipBuildTest(String backendLabel) { set -x # Check if backend label contains string "amd" or backend host is a server with amd gpu if [[ $backendLabel =~ amd ]]; then + sleep 120 export HT_CONFIG_FILE="$HIP_DIR/tests/catch/hipTestMain/config/config_amd_linux.json" LLVM_PATH=/opt/rocm/llvm ctest -E 'Unit_hipGraphChildGraphNodeGetGraph_Functional|Unit_hipGraphExecMemcpyNodeSetParamsFromSymbol_Negative|Unit_hipPtrGetAttribute_Simple|Unit_hipStreamPerThread_DeviceReset_2' - sleep 120 else make test fi diff --git a/docs/markdown/hip_build.md b/docs/markdown/hip_build.md index 17bc1af955..54d85fa540 100755 --- a/docs/markdown/hip_build.md +++ b/docs/markdown/hip_build.md @@ -141,7 +141,7 @@ After build and install HIP commands, catch tests can be built via the following ``` cd "$HIP_DIR" mkdir -p build; cd build -export HIP_PATH=$HIPAMD_DIR/build +export HIP_PATH=$HIPAMD_DIR/build/install cmake ../tests/catch/ -DHIP_PLATFORM=amd make -j$(nproc) build_tests ctest # run tests diff --git a/docs/markdown/hip_deprecated_api_list.md b/docs/markdown/hip_deprecated_api_list.md index 783ea4a390..b26aceecee 100644 --- a/docs/markdown/hip_deprecated_api_list.md +++ b/docs/markdown/hip_deprecated_api_list.md @@ -1,9 +1,11 @@ # HIP Deprecated APIs + ## HIP Context Management APIs CUDA supports cuCtx API, the Driver API that defines "Context" and "Devices" as separate entities. Contexts contain a single device, and a device can theoretically have multiple contexts. HIP initially added limited support for these API to facilitate easy porting from existing driver codes. These API are marked as deprecated now since there are better alternate interface (such as hipSetDevice or the stream API) to achieve the required functions. - +### hipCtxCreate +### hipCtxDestroy ### hipCtxPopCurrent ### hipCtxPushCurrent ### hipCtxSetCurrent @@ -19,6 +21,7 @@ CUDA supports cuCtx API, the Driver API that defines "Context" and "Devices" as ### hipCtxEnablePeerAccess ### hipCtxDisablePeerAccess + ## HIP Memory Management APIs ### hipMallocHost @@ -31,4 +34,49 @@ Should use "hipHostMalloc" instead. Should use "hipHostMalloc" instead. ### hipFreeHost -Should use "hipHostFree" instead. \ No newline at end of file +Should use "hipHostFree" instead. + +### hipMemcpyToArray +### hipMemcpyFromArray + + +## HIP Profiler Control APIs + +### hipProfilerStart +Should use roctracer/rocTX instead + +### hipProfilerStop +Should use roctracer/rocTX instead + + +## HIP Texture Management APIs + +###hipGetTextureReference +###hipTexRefSetAddressMode +###hipTexRefSetArray +###hipTexRefSetFilterMode +###hipTexRefSetFlags +###hipTexRefSetFormat +###hipBindTexture +###hipBindTexture2D +###hipBindTextureToArray +###hipGetTextureAlignmentOffset +###hipUnbindTexture +###hipTexRefGetAddress +###hipTexRefGetAddressMode +###hipTexRefGetFilterMode +###hipTexRefGetFlags +###hipTexRefGetFormat +###hipTexRefGetMaxAnisotropy +###hipTexRefGetMipmapFilterMode +###hipTexRefGetMipmapLevelBias +###hipTexRefGetMipmapLevelClamp +###hipTexRefGetMipMappedArray +###hipTexRefSetAddress +###hipTexRefSetAddress2D +###hipTexRefSetMaxAnisotropy +###hipTexRefSetBorderColor +###hipTexRefSetMipmapFilterMode +###hipTexRefSetMipmapLevelBias +###hipTexRefSetMipmapLevelClamp +###hipTexRefSetMipmappedArray \ No newline at end of file diff --git a/docs/markdown/hip_programming_guide.md b/docs/markdown/hip_programming_guide.md index e9d25b06bc..9f48870300 100644 --- a/docs/markdown/hip_programming_guide.md +++ b/docs/markdown/hip_programming_guide.md @@ -129,6 +129,9 @@ For more details on hipRTC APIs, refer to HIP-API.pdf in GitHub (https://github. The link here(https://github.com/ROCm-Developer-Tools/HIP/blob/main/tests/src/hiprtc/saxpy.cpp) shows an example how to program HIP application using runtime compilation mechanism, and detail hipRTC programming guide is also available in Github (https://github.com/ROCm-Developer-Tools/HIP/blob/main/docs/markdown/hip_rtc.md). +## HIP Graph +HIP graph is supported. For more details, refer to the HIP API Guide. + ## Device-Side Malloc HIP-Clang now supports device-side malloc and free. diff --git a/include/hip/hip_runtime_api.h b/include/hip/hip_runtime_api.h index f8a244d046..868f0de44a 100644 --- a/include/hip/hip_runtime_api.h +++ b/include/hip/hip_runtime_api.h @@ -657,7 +657,7 @@ enum hipLimit_t { #define hipDeviceScheduleBlockingSync 0x4 #define hipDeviceScheduleMask 0x7 #define hipDeviceMapHost 0x8 -#define hipDeviceLmemResizeToMax 0x16 +#define hipDeviceLmemResizeToMax 0x10 /** Default HIP array allocation flag.*/ #define hipArrayDefault 0x00 #define hipArrayLayered 0x01 diff --git a/samples/0_Intro/square/square.hipref.cpp b/samples/0_Intro/square/square.hipref.cpp index b8213e0074..80ff2a00bf 100644 --- a/samples/0_Intro/square/square.hipref.cpp +++ b/samples/0_Intro/square/square.hipref.cpp @@ -38,7 +38,7 @@ THE SOFTWARE. */ template __global__ void vector_square(T* C_d, const T* A_d, size_t N) { - size_t offset = (blockIdx.x * blockDim_x + threadIdx.x); + size_t offset = (blockIdx.x * blockDim.x + threadIdx.x); size_t stride = blockDim.x * gridDim.x; for (size_t i = offset; i < N; i += stride) { diff --git a/tests/catch/CMakeLists.txt b/tests/catch/CMakeLists.txt index ff12be86e3..a8a488002d 100644 --- a/tests/catch/CMakeLists.txt +++ b/tests/catch/CMakeLists.txt @@ -41,6 +41,7 @@ if(NOT DEFINED ROCM_PATH) endif() endif() endif() +file(TO_CMAKE_PATH "${ROCM_PATH}" ROCM_PATH) message(STATUS "ROCM Path: ${ROCM_PATH}") @@ -179,6 +180,8 @@ if(NOT DEFINED OFFLOAD_ARCH_STR else() message(STATUS "ROCm Agent Enumurator found no valid architectures") endif() +elseif(DEFINED OFFLOAD_ARCH_STR) + string(REPLACE "--offload-arch=" "" HIP_GPU_ARCH_LIST ${OFFLOAD_ARCH_STR}) endif() if(DEFINED OFFLOAD_ARCH_STR) diff --git a/tests/catch/README.md b/tests/catch/README.md index 7d1ff6265b..a0c3c3adc2 100644 --- a/tests/catch/README.md +++ b/tests/catch/README.md @@ -144,7 +144,7 @@ If there arises a condition where certain flag is disabled and due to which a te ```cpp TEST_CASE("TestOnlyOnXnack") { if(!XNACKEnabled) { - HIP_SKIP_TEST("Test only runs on system with XNACK enabled"); + HipTest::HIP_SKIP_TEST("Test only runs on system with XNACK enabled"); return; } // Rest of test functionality @@ -162,11 +162,11 @@ These macros are to be called in multi process tests, inside a process which get There is a special interface available for process isolation. ```hip::SpawnProc``` in ```hip_test_process.hh```. Using this interface test can spawn a process and place passing conditions on its return value or its output to stdout. This can be useful for testing printf output. Sample Usage: ```cpp -hip::SpawnProc proc(, ); +hip::SpawnProc proc(, ); REQUIRE(0 == proc.run()); // Test of return value of the proc REQUIRE(exepctedOutput == proc.getOutput()); // Test on expected output of the process ``` -The process can be a standalone exe (see tests/catch/unit/printfExe for more information). +The process must be a standalone exe inside the same folder as other tests. ## Enabling New Tests Initially, the new tests can be enabled via using ```-DHIP_CATCH_TEST=1```. After porting existing tests, this will be turned on by default. diff --git a/tests/catch/hipTestMain/config/config_amd_linux.json b/tests/catch/hipTestMain/config/config_amd_linux.json index fc2c1331ed..91c89ae7dc 100644 --- a/tests/catch/hipTestMain/config/config_amd_linux.json +++ b/tests/catch/hipTestMain/config/config_amd_linux.json @@ -3,7 +3,8 @@ [ "# Following test is related to ticket EXSWCPHIPT-41", "Unit_hipStreamGetPriority_happy", - "Unit_hipStreamPerThread_DeviceReset_1" + "Unit_hipStreamPerThread_DeviceReset_1", + "Unit_hipGraphAddKernelNode_Negative" ] } diff --git a/tests/catch/hipTestMain/config/config_amd_windows.json b/tests/catch/hipTestMain/config/config_amd_windows.json index 2d92fd5752..547d92a4bf 100644 --- a/tests/catch/hipTestMain/config/config_amd_windows.json +++ b/tests/catch/hipTestMain/config/config_amd_windows.json @@ -29,6 +29,23 @@ "Unit_hipDeviceGetUuid", "Unit_hipGraphMemcpyNodeSetParamsFromSymbol_Functional", "Unit_hipGraphExecEventWaitNodeSetEvent_Negative", - "Unit_hipGraphExecEventWaitNodeSetEvent_SetAndVerifyMemory" + "Unit_hipGraphExecEventWaitNodeSetEvent_SetAndVerifyMemory", + "Unit_hipGraphMemcpyNodeSetParams_Functional", + "Unit_hipMalloc3D_ValidatePitch", + "Unit_hipArrayCreate_happy", + "Unit_hipGraphAddKernelNode_Negative", + "Unit_hipHostRegister_Negative - int", + "Unit_hipHostRegister_Negative - float", + "Unit_hipHostRegister_Negative - double", + "Unit_hipMemAllocPitch_ValidatePitch", + "Unit_hipArrayCreate_happy - int", + "Unit_hipArrayCreate_happy - int4", + "Unit_hipArrayCreate_happy - short2", + "Unit_hipArrayCreate_happy - char", + "Unit_hipArrayCreate_happy - char4", + "Unit_hipArrayCreate_happy - float", + "Unit_hipArrayCreate_happy - float2", + "Unit_hipArrayCreate_happy - float4", + "Unit_hipMemVmm_Basic" ] } diff --git a/tests/catch/include/hip_test_helper.hh b/tests/catch/include/hip_test_helper.hh index fbbcb6cb06..9d4cbcd73f 100644 --- a/tests/catch/include/hip_test_helper.hh +++ b/tests/catch/include/hip_test_helper.hh @@ -38,7 +38,7 @@ static inline int getGeviceCount() { } // Get Free Memory from the system -static size_t getMemoryAmount() { +static inline size_t getMemoryAmount() { #ifdef __linux__ struct sysinfo info{}; sysinfo(&info); diff --git a/tests/catch/include/hip_test_process.hh b/tests/catch/include/hip_test_process.hh index 665f5d6b97..335e6c7bba 100644 --- a/tests/catch/include/hip_test_process.hh +++ b/tests/catch/include/hip_test_process.hh @@ -1,5 +1,5 @@ /* -Copyright (c) 2021 Advanced Micro Devices, Inc. All rights reserved. +Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved. Permission is hereby granted, free of charge, to any person obtaining a copy of this software and associated documentation files (the "Software"), to deal @@ -21,7 +21,9 @@ THE SOFTWARE. */ #pragma once + #include "hip_test_common.hh" +#include "hip_test_filesystem.hh" #include #include @@ -29,9 +31,17 @@ THE SOFTWARE. #include #include #include -#include "hip_test_filesystem.hh" namespace hip { +/* +Class to spawn a process in isolation and test its standard output and return status +Good for printf tests and environment variable tests + +How to use: +Have the stand alone exe in the same folder +Init a class using hip::SpawnProc proc("ExeName", yes_or_no_to_capture_output); +proc.run("Optional command line args"); +*/ class SpawnProc { std::string exeName; std::string resultStr; @@ -53,18 +63,31 @@ class SpawnProc { public: SpawnProc(std::string exeName_, bool captureOutput_ = false) : exeName(exeName_), captureOutput(captureOutput_) { - auto dir = fs::path(TestContext::get().currentPath()).parent_path(); + auto dir = fs::path(TestContext::get().currentPath()); dir /= exeName; exeName = dir.string(); + + INFO("Testing that exe exists: " << exeName); + REQUIRE(fs::exists(exeName)); + if (captureOutput) { auto path = fs::temp_directory_path(); path /= getRandomString(); tmpFileName = path.string(); + INFO("Testing that capture file does not exist already: " << tmpFileName); + REQUIRE(!fs::exists(tmpFileName)); } } - int run() { + int run(std::string commandLineArgs = "") { std::string execCmd = exeName; + + // Append command line args + if (commandLineArgs.size() > 0) { + execCmd += " "; // Add space for command line args + execCmd += commandLineArgs; + } + if (captureOutput) { execCmd += " > "; execCmd += tmpFileName; @@ -77,7 +100,11 @@ class SpawnProc { resultStr = std::string((std::istreambuf_iterator(t)), std::istreambuf_iterator()); } +#if HT_LINUX + return WEXITSTATUS(res); +#else return res; +#endif } std::string getOutput() { return resultStr; } diff --git a/tests/catch/packaging/hip-tests.txt b/tests/catch/packaging/hip-tests.txt index bbf3ec0ade..77b3a0e46c 100644 --- a/tests/catch/packaging/hip-tests.txt +++ b/tests/catch/packaging/hip-tests.txt @@ -30,7 +30,7 @@ install(FILES @PROJECT_BINARY_DIR@/CTestTestfile.cmake # Packaging steps ############################# set(CPACK_SET_DESTDIR TRUE) -set(CPACK_INSTALL_PREFIX @ROCM_PATH@/test/hip/) +set(CPACK_INSTALL_PREFIX "@ROCM_PATH@/test/hip/") set(PKG_NAME catch-@HIP_PLATFORM@) set(CPACK_PACKAGE_NAME ${PKG_NAME}) set(CPACK_PACKAGE_DESCRIPTION_SUMMARY "HIP: Heterogenous-computing Interface for Portability [CATCH TESTS]") diff --git a/tests/catch/unit/CMakeLists.txt b/tests/catch/unit/CMakeLists.txt index 9ebff9b9de..28c6fcb6aa 100644 --- a/tests/catch/unit/CMakeLists.txt +++ b/tests/catch/unit/CMakeLists.txt @@ -27,7 +27,6 @@ add_subdirectory(event) add_subdirectory(occupancy) add_subdirectory(device) add_subdirectory(printf) -add_subdirectory(printfExe) add_subdirectory(texture) add_subdirectory(streamperthread) add_subdirectory(kernel) diff --git a/tests/catch/unit/device/CMakeLists.txt b/tests/catch/unit/device/CMakeLists.txt index bebee8af71..247867bf5c 100644 --- a/tests/catch/unit/device/CMakeLists.txt +++ b/tests/catch/unit/device/CMakeLists.txt @@ -18,7 +18,13 @@ set(TEST_SRC hipDeviceGetUuid.cc ) +set_source_files_properties(hipGetDeviceCount.cc PROPERTIES COMPILE_FLAGS -std=c++17) + +add_executable(getDeviceCount EXCLUDE_FROM_ALL getDeviceCount_exe.cc) + hip_add_exe_to_target(NAME DeviceTest TEST_SRC ${TEST_SRC} TEST_TARGET_NAME build_tests COMPILE_OPTIONS -std=c++14) + +add_dependencies(DeviceTest getDeviceCount) diff --git a/tests/catch/unit/device/getDeviceCount_exe.cc b/tests/catch/unit/device/getDeviceCount_exe.cc new file mode 100644 index 0000000000..bf9951b6d4 --- /dev/null +++ b/tests/catch/unit/device/getDeviceCount_exe.cc @@ -0,0 +1,89 @@ +/* +Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved. + +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#include +#include +#include + +bool UNSETENV(std::string var) { + int result = -1; + #ifdef __unix__ + result = unsetenv(var.c_str()); + #else + result = _putenv((var + '=').c_str()); + #endif + return (result == 0) ? true: false; +} + +bool SETENV(std::string var, std::string value, int overwrite) { + int result = -1; + #ifdef __unix__ + result = setenv(var.c_str(), value.c_str(), overwrite); + #else + result = _putenv((var + '=' + value).c_str()); + #endif + return (result == 0) ? true: false; +} + +// Expects 1 command line arg, which is the Device Visible String +int main(int argc, char** argv) { + if (argc != 2) { + std::cerr << "Invalid number of args passed.\n" + << "argc : " << argc << std::endl; + for (int i = 0; i < argc; i++) { + std::cerr << " argv[" << i << "] : " << argv[0] << std::endl; + } + std::cerr << "The program expects device visibility string i.e. 0,1,2" << std::endl; + return -1; + } + + // disable visible_devices env from shell +#ifdef __HIP_PLATFORM_NVCC__ + UNSETENV("CUDA_VISIBLE_DEVICES"); + SETENV("CUDA_VISIBLE_DEVICES", argv[1], 1); + auto init_res = hipInit(0); + if (hipSuccess != init_res) { + std::cerr << "CUDA INIT API returned : " << hipGetErrorString(init_res) << std::endl; + return -1; + } +#else + UNSETENV("ROCR_VISIBLE_DEVICES"); + UNSETENV("HIP_VISIBLE_DEVICES"); + SETENV("ROCR_VISIBLE_DEVICES", argv[1], 1); + SETENV("HIP_VISIBLE_DEVICES", argv[1], 1); +#endif + + int count = 0; + auto res = hipGetDeviceCount(&count); + if (hipSuccess != res) { + std::cerr << "HIP API returned : " << hipGetErrorString(res) << std::endl; + return -1; + } + +#ifdef __HIP_PLATFORM_NVCC__ + UNSETENV("CUDA_VISIBLE_DEVICES"); +#else + UNSETENV("ROCR_VISIBLE_DEVICES"); + UNSETENV("HIP_VISIBLE_DEVICES"); +#endif + return count; +} \ No newline at end of file diff --git a/tests/catch/unit/device/hipGetDeviceCount.cc b/tests/catch/unit/device/hipGetDeviceCount.cc index 2179bdad05..2facbf9515 100644 --- a/tests/catch/unit/device/hipGetDeviceCount.cc +++ b/tests/catch/unit/device/hipGetDeviceCount.cc @@ -23,6 +23,7 @@ THE SOFTWARE. */ #include +#include /** * hipGetDeviceCount tests @@ -32,3 +33,25 @@ TEST_CASE("Unit_hipGetDeviceCount_NegTst") { // Scenario1 REQUIRE_FALSE(hipGetDeviceCount(nullptr) == hipSuccess); } + +TEST_CASE("Unit_hipGetDeviceCount_HideDevices") { + int deviceCount = HipTest::getDeviceCount(); + if (deviceCount < 2) { + HipTest::HIP_SKIP_TEST("This test requires more than 2 GPUs. Skipping."); + return; + } + + for (int i = deviceCount; i >= 1; i--) { + std::string visibleStr; + for (int j = 0; j < i; j++) { // Generate a string which has first i devices + visibleStr += std::to_string(j); + if (j != (i - 1)) { + visibleStr += ","; + } + } + + hip::SpawnProc proc("getDeviceCount", true); + INFO("Output from process : " << proc.getOutput()); + REQUIRE(proc.run(visibleStr) == i); + } +} diff --git a/tests/catch/unit/event/Unit_hipEventRecord.cc b/tests/catch/unit/event/Unit_hipEventRecord.cc index 408f802d41..41c793d23e 100644 --- a/tests/catch/unit/event/Unit_hipEventRecord.cc +++ b/tests/catch/unit/event/Unit_hipEventRecord.cc @@ -19,72 +19,102 @@ LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. */ + // Test hipEventRecord serialization behavior. -// Through manual inspection of the reported timestamps, can determine if recording a NULL event -// forces synchronization : set -#include -#include -#include + #include +#include +#include +#include + TEST_CASE("Unit_hipEventRecord") { - size_t N = 4 * 1024 * 1024; - unsigned threadsPerBlock = 256; - int iterations = 1; + constexpr size_t N = 1024; + constexpr int iterations = 1; - unsigned blocks = (N + threadsPerBlock - 1) / threadsPerBlock; - if (blocks > 1024) blocks = 1024; - if (blocks == 0) blocks = 1; + constexpr int blocks = 1024; - printf("N=%zu (A+B+C= %6.1f MB total) blocks=%u threadsPerBlock=%u iterations=%d\n", N, - ((double)3 * N * sizeof(float)) / 1024 / 1024, blocks, threadsPerBlock, iterations); - printf("iterations=%d\n", iterations); + constexpr size_t Nbytes = N * sizeof(float); - size_t Nbytes = N * sizeof(float); + float *A_h, *B_h, *C_h; + float *A_d, *B_d, *C_d; + HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N); - float *A_h, *B_h, *C_h; - float *A_d, *B_d, *C_d; - HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N); + enum TestType { + WithFlags_Default = hipEventDefault, + WithFlags_Blocking = hipEventBlockingSync, + WithFlags_DisableTiming = hipEventDisableTiming, +#if HT_AMD + WithFlags_ReleaseToDevice = hipEventReleaseToDevice, + WithFlags_ReleaseToSystem = hipEventReleaseToSystem, +#endif + WithoutFlags + }; - hipEvent_t start, stop; +#if HT_AMD + auto flags = GENERATE(WithFlags_Default, WithFlags_Blocking, WithFlags_DisableTiming, + WithFlags_ReleaseToDevice, WithFlags_ReleaseToSystem, WithoutFlags); +#endif - // NULL stream check: +#if HT_NVIDIA + auto flags = + GENERATE(WithFlags_Default, WithFlags_Blocking, WithFlags_DisableTiming, WithoutFlags); +#endif + + + hipEvent_t start{}, stop{}; + + if (flags == WithoutFlags) { HIP_CHECK(hipEventCreate(&start)); HIP_CHECK(hipEventCreate(&stop)); + } else { + HIP_CHECK(hipEventCreateWithFlags(&start, flags)); + HIP_CHECK(hipEventCreateWithFlags(&stop, flags)); + } - HIP_CHECK(hipMemcpy(A_d, A_h, Nbytes, hipMemcpyHostToDevice)); - HIP_CHECK(hipMemcpy(B_d, B_h, Nbytes, hipMemcpyHostToDevice)); + HIP_CHECK(hipMemcpy(A_d, A_h, Nbytes, hipMemcpyHostToDevice)); + HIP_CHECK(hipMemcpy(B_d, B_h, Nbytes, hipMemcpyHostToDevice)); - for (int i = 0; i < iterations; i++) { - //--- START TIMED REGION - long long hostStart = HipTest::get_time(); - // Record the start event - HIP_CHECK(hipEventRecord(start, NULL)); + for (int i = 0; i < iterations; i++) { + //--- START TIMED REGION + long long hostStart = HipTest::get_time(); + // Record the start event + HIP_CHECK(hipEventRecord(start, NULL)); - HipTest::launchKernel(HipTest::vectorADD, blocks, threadsPerBlock, 0, 0, -static_cast(A_d), static_cast(B_d), C_d, N); + HipTest::launchKernel(HipTest::vectorADD, blocks, 1, 0, 0, + static_cast(A_d), static_cast(B_d), + C_d, N); - HIP_CHECK(hipEventRecord(stop, NULL)); - HIP_CHECK(hipEventSynchronize(stop)); - long long hostStop = HipTest::get_time(); - //--- STOP TIMED REGION + HIP_CHECK(hipEventRecord(stop, NULL)); + HIP_CHECK(hipEventSynchronize(stop)); + long long hostStop = HipTest::get_time(); + //--- STOP TIMED REGION - float eventMs = 1.0f; - HIP_CHECK(hipEventElapsedTime(&eventMs, start, stop)); - float hostMs = HipTest::elapsed_time(hostStart, hostStop); + float hostMs = HipTest::elapsed_time(hostStart, hostStop); - printf("host_time (chrono) =%6.3fms\n", hostMs); - printf("kernel_time (hipEventElapsedTime) =%6.3fms\n", eventMs); - printf("\n"); + INFO("host_time (chrono) = " << hostMs); - // Make sure timer is timing something... - REQUIRE(eventMs > 0.0f); + // Make sure timer is timing something... + if (flags != WithFlags_DisableTiming) { + float eventMs = 1.0f; + HIP_CHECK(hipEventElapsedTime(&eventMs, start, stop)); + INFO("kernel_time (hipEventElapsedTime) = " << eventMs); + REQUIRE(eventMs > 0.0f); } + } - HIP_CHECK(hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost)); + HIP_CHECK(hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost)); - HIP_CHECK(hipEventDestroy(start)); - HIP_CHECK(hipEventDestroy(stop)); + HIP_CHECK(hipEventDestroy(start)); + HIP_CHECK(hipEventDestroy(stop)); - HipTest::checkVectorADD(A_h, B_h, C_h, N, true); + HipTest::checkVectorADD(A_h, B_h, C_h, N, true); + HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false); + TestContext::get().cleanContext(); } + +TEST_CASE("Unit_hipEventRecord_Negative") { + SECTION("Nullptr event") { + HIP_CHECK_ERROR(hipEventRecord(nullptr, nullptr), hipErrorInvalidResourceHandle); + } +} \ No newline at end of file diff --git a/tests/catch/unit/graph/CMakeLists.txt b/tests/catch/unit/graph/CMakeLists.txt index a8a37733d4..33a6ce3cd5 100644 --- a/tests/catch/unit/graph/CMakeLists.txt +++ b/tests/catch/unit/graph/CMakeLists.txt @@ -64,6 +64,10 @@ set(TEST_SRC hipStreamEndCapture.cc hipGraphMemcpyNodeSetParamsFromSymbol.cc hipGraphExecEventWaitNodeSetEvent.cc + hipGraphAddMemsetNode.cc + hipGraphAddKernelNode.cc + hipGraphMemcpyNodeGetParams.cc + hipGraphMemcpyNodeSetParams.cc ) hip_add_exe_to_target(NAME GraphsTest diff --git a/tests/catch/unit/graph/hipGraphAddKernelNode.cc b/tests/catch/unit/graph/hipGraphAddKernelNode.cc new file mode 100644 index 0000000000..79ed3c76ea --- /dev/null +++ b/tests/catch/unit/graph/hipGraphAddKernelNode.cc @@ -0,0 +1,101 @@ +/* +Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved. +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#include +#include + +/* Test verifies hipGraphAddKernelNode API Negative scenarios. + */ + +TEST_CASE("Unit_hipGraphAddKernelNode_Negative") { + constexpr int N = 1024; + size_t NElem{N}; + constexpr auto blocksPerCU = 6; // to hide latency + constexpr auto threadsPerBlock = 256; + unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N); + int *A_d, *B_d, *C_d; + hipGraph_t graph; + hipError_t ret; + hipGraphNode_t kNode; + hipKernelNodeParams kNodeParams{}; + std::vector dependencies; + + HIP_CHECK(hipMalloc(&A_d, sizeof(int) * N)); + HIP_CHECK(hipMalloc(&B_d, sizeof(int) * N)); + HIP_CHECK(hipMalloc(&C_d, sizeof(int) * N)); + HIP_CHECK(hipGraphCreate(&graph, 0)); + + void* kernelArgs[] = {&A_d, &B_d, &C_d, reinterpret_cast(&NElem)}; + kNodeParams.func = reinterpret_cast(HipTest::vectorADD); + kNodeParams.gridDim = dim3(blocks); + kNodeParams.blockDim = dim3(threadsPerBlock); + kNodeParams.sharedMemBytes = 0; + kNodeParams.kernelParams = reinterpret_cast(kernelArgs); + kNodeParams.extra = nullptr; + + SECTION("Pass pGraphNode as nullptr") { + ret = hipGraphAddKernelNode(nullptr, graph, nullptr, 0, &kNodeParams); + REQUIRE(hipErrorInvalidValue == ret); + } + SECTION("Pass Graph as nullptr") { + ret = hipGraphAddKernelNode(&kNode, nullptr, nullptr, 0, &kNodeParams); + REQUIRE(hipErrorInvalidValue == ret); + } + SECTION("Pass invalid numDependencies") { + ret = hipGraphAddKernelNode(&kNode, graph, nullptr, 11, &kNodeParams); + REQUIRE(hipErrorInvalidValue == ret); + } + SECTION("Pass invalid numDependencies and valid list for dependencies") { + HIP_CHECK(hipGraphAddKernelNode(&kNode, graph, nullptr, 0, &kNodeParams)); + dependencies.push_back(kNode); + ret = hipGraphAddKernelNode(&kNode, graph, + dependencies.data(), dependencies.size()+1, &kNodeParams); + REQUIRE(hipErrorInvalidValue == ret); + } + SECTION("Pass NodeParams as nullptr") { + ret = hipGraphAddKernelNode(&kNode, graph, + dependencies.data(), dependencies.size(), nullptr); + REQUIRE(hipErrorInvalidValue == ret); + } + SECTION("Pass NodeParams func datamember as nullptr") { + kNodeParams.func = nullptr; + ret = hipGraphAddKernelNode(&kNode, graph, nullptr, 0, &kNodeParams); + REQUIRE(hipSuccess != ret); + } + SECTION("Pass kernelParams datamember as nullptr") { + kNodeParams.func = reinterpret_cast(HipTest::vectorADD); + kNodeParams.kernelParams = nullptr; + ret = hipGraphAddKernelNode(&kNode, graph, nullptr, 0, &kNodeParams); + REQUIRE(hipErrorInvalidValue == ret); + } +#if HT_AMD +// On Cuda setup this test case getting failed + SECTION("Try adding kernel node after destroy the already created graph") { + kNodeParams.kernelParams = reinterpret_cast(kernelArgs); + HIP_CHECK(hipGraphDestroy(graph)); + ret = hipGraphAddKernelNode(&kNode, graph, nullptr, 0, &kNodeParams); + REQUIRE(hipErrorInvalidValue == ret); + } +#endif + + HIP_CHECK(hipFree(A_d)); + HIP_CHECK(hipFree(B_d)); + HIP_CHECK(hipFree(C_d)); +} + diff --git a/tests/catch/unit/graph/hipGraphAddMemsetNode.cc b/tests/catch/unit/graph/hipGraphAddMemsetNode.cc new file mode 100644 index 0000000000..0984e194c8 --- /dev/null +++ b/tests/catch/unit/graph/hipGraphAddMemsetNode.cc @@ -0,0 +1,123 @@ +/* +Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved. +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANNTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +/** +Negative Testcase Scenarios for api hipGraphAddMemsetNode : +1) Pass pGraphNode as nullptr and check if api returns error. +2) Pass pGraphNode as un-initialize object and check. +3) Pass Graph as nullptr and check if api returns error. +4) Pass Graph as empty object(skipping graph creation), api should return error code. +5) Pass pDependencies as nullptr, api should return success. +6) Pass numDependencies is max(size_t) and pDependencies is not valid ptr, api expected to return error code. +7) Pass pDependencies is nullptr, but numDependencies is non-zero, api expected to return error. +8) Pass pMemsetParams as nullptr and check if api returns error code. +9) Pass pMemsetParams as un-initialize object and check if api returns error code. +10) Pass hipMemsetParams::dst as nullptr should return error code. +11) Pass hipMemsetParams::element size other than 1, 2, or 4 and check api should return error code. +12) Pass hipMemsetParams::height as zero and check api should return error code. +*/ + +#include + +/** + * Negative Test for API hipGraphAddMemsetNode + */ + +TEST_CASE("Unit_hipGraphAddMemsetNode_Negative") { + hipError_t ret; + hipGraph_t graph; + hipGraphNode_t memsetNode; + char *devData; + + HIP_CHECK(hipMalloc(&devData, 1024)); + HIP_CHECK(hipGraphCreate(&graph, 0)); + + hipMemsetParams memsetParams{}; + memset(&memsetParams, 0, sizeof(memsetParams)); + memsetParams.dst = reinterpret_cast(devData); + memsetParams.value = 0; + memsetParams.pitch = 0; + memsetParams.elementSize = sizeof(char); + memsetParams.width = 1024; + memsetParams.height = 1; + + SECTION("Pass pGraphNode as nullptr") { + ret = hipGraphAddMemsetNode(nullptr, graph, nullptr, 0, &memsetParams); + REQUIRE(hipErrorInvalidValue == ret); + } + SECTION("Pass pGraphNode as un-initialize object") { + hipGraphNode_t memsetNode_1; + ret = hipGraphAddMemsetNode(&memsetNode_1, graph, + nullptr, 0, &memsetParams); + REQUIRE(hipSuccess == ret); + } + SECTION("Pass graph as nullptr") { + ret = hipGraphAddMemsetNode(&memsetNode, nullptr, + nullptr, 0, &memsetParams); + REQUIRE(hipErrorInvalidValue == ret); + } + SECTION("Pass Graph as empty object") { + hipGraph_t graph_1{}; + ret = hipGraphAddMemsetNode(&memsetNode, graph_1, + nullptr, 0, &memsetParams); + REQUIRE(hipErrorInvalidValue == ret); + } + SECTION("Pass pDependencies as nullptr") { + ret = hipGraphAddMemsetNode(&memsetNode, graph, nullptr, 0, &memsetParams); + REQUIRE(hipSuccess == ret); + } + SECTION("Pass numDependencies is max and pDependencies is not valid ptr") { + ret = hipGraphAddMemsetNode(&memsetNode, graph, + nullptr, INT_MAX, &memsetParams); + REQUIRE(hipErrorInvalidValue == ret); + } + SECTION("Pass pDependencies as nullptr, but numDependencies is non-zero") { + ret = hipGraphAddMemsetNode(&memsetNode, graph, nullptr, 9, &memsetParams); + REQUIRE(hipErrorInvalidValue == ret); + } + SECTION("Pass pMemsetParams as nullptr") { + ret = hipGraphAddMemsetNode(&memsetNode, graph, nullptr, 0, nullptr); + REQUIRE(hipErrorInvalidValue == ret); + } + SECTION("Pass pMemsetParams as un-initialize object") { + hipMemsetParams memsetParams1; + ret = hipGraphAddMemsetNode(&memsetNode, graph, nullptr, 0, + &memsetParams1); + REQUIRE(hipErrorInvalidValue == ret); + } + SECTION("Pass hipMemsetParams::dst as nullptr") { + memsetParams.dst = nullptr; + ret = hipGraphAddMemsetNode(&memsetNode, graph, nullptr, 0, &memsetParams); + REQUIRE(hipErrorInvalidValue == ret); + } + SECTION("Pass hipMemsetParams::element size other than 1, 2, or 4") { + memsetParams.dst = reinterpret_cast(devData); + memsetParams.elementSize = 9; + ret = hipGraphAddMemsetNode(&memsetNode, graph, nullptr, 0, &memsetParams); + REQUIRE(hipErrorInvalidValue == ret); + } + SECTION("Pass hipMemsetParams::height as zero") { + memsetParams.elementSize = sizeof(char); + memsetParams.height = 0; + ret = hipGraphAddMemsetNode(&memsetNode, graph, nullptr, 0, &memsetParams); + REQUIRE(hipErrorInvalidValue == ret); + } + HIP_CHECK(hipFree(devData)); + HIP_CHECK(hipGraphDestroy(graph)); +} diff --git a/tests/catch/unit/graph/hipGraphMemcpyNodeGetParams.cc b/tests/catch/unit/graph/hipGraphMemcpyNodeGetParams.cc new file mode 100644 index 0000000000..6393358bb2 --- /dev/null +++ b/tests/catch/unit/graph/hipGraphMemcpyNodeGetParams.cc @@ -0,0 +1,232 @@ +/* +Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved. +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +/** +Testcase Scenarios : +Negative - +1) Pass node as nullptr and verify api returns error code. +2) Pass un-initialize node and verify api returns error code. +3) Pass pNodeParams as nullptr and verify api returns error code. +Functional - +1) Create a graph, add Memcpy node to graph with desired node params. + Verify api fetches the node params mentioned while adding Memcpy node. +2) Set Memcpy node params with hipGraphMemcpyNodeSetParams, + now get the params and verify both are same. +*/ + +#include +#include + +#define SIZE 10 +#define UPDATESIZE 8 + +/* Test verifies hipGraphMemcpyNodeGetParams API Negative scenarios. + */ +TEST_CASE("Unit_hipGraphMemcpyNodeGetParams_Negative") { + constexpr int width{SIZE}, height{SIZE}, depth{SIZE}; + hipArray *devArray; + hipChannelFormatKind formatKind = hipChannelFormatKindSigned; + hipMemcpy3DParms myparms; + int* hData; + uint32_t size = width * height * depth * sizeof(int); + hData = reinterpret_cast(malloc(size)); + REQUIRE(hData != nullptr); + memset(hData, 0, size); + for (int i = 0; i < depth; i++) { + for (int j = 0; j < height; j++) { + for (int k = 0; k < width; k++) { + hData[i*width*height + j*width + k] = i*width*height + j*width + k; + } + } + } + hipChannelFormatDesc channelDesc = hipCreateChannelDesc(sizeof(int)*8, + 0, 0, 0, formatKind); + HIP_CHECK(hipMalloc3DArray(&devArray, &channelDesc, make_hipExtent(width, + height, depth), hipArrayDefault)); + memset(&myparms, 0x0, sizeof(hipMemcpy3DParms)); + myparms.srcPos = make_hipPos(0, 0, 0); + myparms.dstPos = make_hipPos(0, 0, 0); + myparms.extent = make_hipExtent(width , height, depth); + myparms.srcPtr = make_hipPitchedPtr(hData, width * sizeof(int), + width, height); + myparms.dstArray = devArray; + myparms.kind = hipMemcpyHostToDevice; + + hipGraph_t graph; + hipError_t ret; + hipGraphNode_t memcpyNode; + HIP_CHECK(hipGraphCreate(&graph, 0)); + HIP_CHECK(hipGraphAddMemcpyNode(&memcpyNode, graph, NULL, 0, &myparms)); + + SECTION("Pass node as nullptr") { + ret = hipGraphMemcpyNodeGetParams(nullptr, &myparms); + REQUIRE(hipErrorInvalidValue == ret); + } + SECTION("Pass un-initilize node") { + hipGraphNode_t memcpyNode_uninit{}; + ret = hipGraphMemcpyNodeGetParams(memcpyNode_uninit, &myparms); + REQUIRE(hipErrorInvalidValue == ret); + } + SECTION("Pass GetNodeParams as nullptr") { + ret = hipGraphMemcpyNodeGetParams(memcpyNode, nullptr); + REQUIRE(hipErrorInvalidValue == ret); + } + HIP_CHECK(hipFreeArray(devArray)); + free(hData); + HIP_CHECK(hipGraphDestroy(graph)); +} + +/* Test verifies hipGraphMemcpyNodeGetParams API Functional scenarios. + */ + +static bool compareHipPos(hipPos hPos1, hipPos hPos2) { + if ((hPos1.x == hPos2.x) && (hPos1.y == hPos2.y) && (hPos1.z == hPos2.z)) + return true; + else + return false; +} +static bool compareHipExtent(hipExtent hExt1, hipExtent hExt2) { + if ((hExt1.width == hExt2.width) && (hExt1.height == hExt2.height) && + (hExt1.depth == hExt2.depth)) + return true; + else + return false; +} +static bool compareHipPitchedPtr(hipPitchedPtr hpPtr1, hipPitchedPtr hpPtr2) { + if ((reinterpret_cast(hpPtr1.ptr) == + reinterpret_cast(hpPtr2.ptr)) + && (hpPtr1.pitch == hpPtr2.pitch) + #if HT_AMD + && (hpPtr1.xsize == hpPtr2.xsize) + /* xsize check below is disabled on nvidia as xsize value + * is not being updated properly due to issue with CUDA api */ + #endif + && (hpPtr1.ysize == hpPtr2.ysize)) + return true; + else + return false; +} + +static bool memcpyNodeCompare(hipMemcpy3DParms *mNode1, + hipMemcpy3DParms *mNode2) { + if (mNode1->srcArray != mNode2->srcArray) + return false; + if (!compareHipPos(mNode1->srcPos, mNode2->srcPos)) + return false; + if (!compareHipPitchedPtr(mNode1->srcPtr, mNode2->srcPtr)) + return false; + if (mNode1->dstArray != mNode2->dstArray) + return false; + if (!compareHipPos(mNode1->dstPos, mNode2->dstPos)) + return false; + if (!compareHipPitchedPtr(mNode1->dstPtr, mNode2->dstPtr)) + return false; + if (!compareHipExtent(mNode1->extent, mNode2->extent)) + return false; + if (mNode1->kind != mNode2->kind) + return false; + return true; +} + +TEST_CASE("Unit_hipGraphMemcpyNodeGetParams_Functional") { + constexpr int width{SIZE}, height{SIZE}, depth{SIZE}; + hipArray *devArray; + hipChannelFormatKind formatKind = hipChannelFormatKindSigned; + hipMemcpy3DParms myparms; + int* hData; + uint32_t size = width * height * depth * sizeof(int); + hData = reinterpret_cast(malloc(size)); + REQUIRE(hData != nullptr); + memset(hData, 0, size); + for (int i = 0; i < depth; i++) { + for (int j = 0; j < height; j++) { + for (int k = 0; k < width; k++) { + hData[i*width*height + j*width + k] = i*width*height + j*width + k; + } + } + } + hipChannelFormatDesc channelDesc = hipCreateChannelDesc(sizeof(int)*8, + 0, 0, 0, formatKind); + HIP_CHECK(hipMalloc3DArray(&devArray, &channelDesc, make_hipExtent(width, + height, depth), hipArrayDefault)); + memset(&myparms, 0x0, sizeof(hipMemcpy3DParms)); + myparms.srcPos = make_hipPos(0, 0, 0); + myparms.dstPos = make_hipPos(0, 0, 0); + myparms.extent = make_hipExtent(width , height, depth); + myparms.srcPtr = make_hipPitchedPtr(hData, width * sizeof(int), + width, height); + myparms.dstArray = devArray; + myparms.kind = hipMemcpyHostToDevice; + + hipGraph_t graph; + hipGraphNode_t memcpyNode; + HIP_CHECK(hipGraphCreate(&graph, 0)); + HIP_CHECK(hipGraphAddMemcpyNode(&memcpyNode, graph, NULL, 0, &myparms)); + + SECTION("Get Memcpy Param and verify.") { + hipMemcpy3DParms m3DGetParams; + REQUIRE(hipSuccess == hipGraphMemcpyNodeGetParams(memcpyNode, + &m3DGetParams)); + // Validating the result + REQUIRE(true == memcpyNodeCompare(&myparms, &m3DGetParams)); + } + + SECTION("Set memcpy params and Get param and verify.") { + hipMemcpy3DParms myparms1, m3DGetParams1; + constexpr int width1{UPDATESIZE}, height1{UPDATESIZE}, depth1{UPDATESIZE}; + hipArray *devArray1; + hipChannelFormatKind formatKind1 = hipChannelFormatKindSigned; + int* hData1; + uint32_t size1 = width1 * height1 * depth1 * sizeof(int); + hData1 = reinterpret_cast(malloc(size1)); + REQUIRE(hData1 != nullptr); + memset(hData1, 0, size1); + for (int i = 0; i < depth1; i++) { + for (int j = 0; j < height1; j++) { + for (int k = 0; k < width1; k++) { + hData1[i*width1*height1 + j*width1 + k] = i*width1*height1 + + j*width1 + k; + } + } + } + hipChannelFormatDesc channelDesc1 = hipCreateChannelDesc(sizeof(int)*8, + 0, 0, 0, formatKind1); + HIP_CHECK(hipMalloc3DArray(&devArray1, &channelDesc1, + make_hipExtent(width1, height1, depth1), hipArrayDefault)); + memset(&myparms1, 0x0, sizeof(hipMemcpy3DParms)); + myparms1.srcPos = make_hipPos(0, 0, 0); + myparms1.dstPos = make_hipPos(0, 0, 0); + myparms1.extent = make_hipExtent(width1 , height1, depth1); + myparms1.srcPtr = make_hipPitchedPtr(hData1, width1 * sizeof(int), + width1, height1); + myparms1.dstArray = devArray1; + myparms1.kind = hipMemcpyHostToDevice; + + REQUIRE(hipSuccess == hipGraphMemcpyNodeSetParams(memcpyNode, &myparms1)); + REQUIRE(hipSuccess == hipGraphMemcpyNodeGetParams(memcpyNode, + &m3DGetParams1)); + REQUIRE(true == memcpyNodeCompare(&myparms1, &m3DGetParams1)); + + HIP_CHECK(hipFreeArray(devArray1)); + free(hData1); + } + HIP_CHECK(hipFreeArray(devArray)); + free(hData); + HIP_CHECK(hipGraphDestroy(graph)); +} diff --git a/tests/catch/unit/graph/hipGraphMemcpyNodeSetParams.cc b/tests/catch/unit/graph/hipGraphMemcpyNodeSetParams.cc new file mode 100644 index 0000000000..12cf7e22eb --- /dev/null +++ b/tests/catch/unit/graph/hipGraphMemcpyNodeSetParams.cc @@ -0,0 +1,215 @@ +/* +Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved. +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +/** +Testcase Scenarios : +Negative - +1) Pass node as nullptr and verify api returns error code. +2) Pass un-initialize node and verify api returns error code. +3) Pass pNodeParams as nullptr and verify api returns error code. +Functional - +1) Add Memcpy node to graph, update the Memcpy node params with set and + launch the graph and check updated params are taking effect. +2) Add Memcpy node to graph, launch graph, then update the Memcpy node params + with set and launch the graph and check updated params are taking effect. +*/ + +#include +#include + +#define SIZE 10 + +/* Test verifies hipGraphMemcpyNodeSetParams API Negative scenarios. + */ +TEST_CASE("Unit_hipGraphMemcpyNodeSetParams_Negative") { + constexpr int width{SIZE}, height{SIZE}, depth{SIZE}; + hipArray *devArray; + hipChannelFormatKind formatKind = hipChannelFormatKindSigned; + hipMemcpy3DParms myparms; + int* hData; + uint32_t size = width * height * depth * sizeof(int); + hData = reinterpret_cast(malloc(size)); + REQUIRE(hData != nullptr); + memset(hData, 0, size); + for (int i = 0; i < depth; i++) { + for (int j = 0; j < height; j++) { + for (int k = 0; k < width; k++) { + hData[i*width*height + j*width + k] = i*width*height + j*width + k; + } + } + } + hipChannelFormatDesc channelDesc = hipCreateChannelDesc(sizeof(int)*8, + 0, 0, 0, formatKind); + HIP_CHECK(hipMalloc3DArray(&devArray, &channelDesc, make_hipExtent(width, + height, depth), hipArrayDefault)); + memset(&myparms, 0x0, sizeof(hipMemcpy3DParms)); + myparms.srcPos = make_hipPos(0, 0, 0); + myparms.dstPos = make_hipPos(0, 0, 0); + myparms.extent = make_hipExtent(width , height, depth); + myparms.srcPtr = make_hipPitchedPtr(hData, width * sizeof(int), + width, height); + myparms.dstArray = devArray; + myparms.kind = hipMemcpyHostToDevice; + + hipGraph_t graph; + hipError_t ret; + hipGraphNode_t memcpyNode; + HIP_CHECK(hipGraphCreate(&graph, 0)); + HIP_CHECK(hipGraphAddMemcpyNode(&memcpyNode, graph, NULL, 0, &myparms)); + + SECTION("Pass node as nullptr") { + ret = hipGraphMemcpyNodeSetParams(nullptr, &myparms); + REQUIRE(hipErrorInvalidValue == ret); + } + SECTION("Pass un-initialize node") { + hipGraphNode_t memcpyNode_uninit{}; + ret = hipGraphMemcpyNodeSetParams(memcpyNode_uninit, &myparms); + REQUIRE(hipErrorInvalidValue == ret); + } + SECTION("Pass SetNodeParams as nullptr") { + ret = hipGraphMemcpyNodeSetParams(memcpyNode, nullptr); + REQUIRE(hipErrorInvalidValue == ret); + } + HIP_CHECK(hipFreeArray(devArray)); + free(hData); + HIP_CHECK(hipGraphDestroy(graph)); +} + +/* Test verifies hipGraphMemcpyNodeSetParams API Functional scenarios. + */ +TEST_CASE("Unit_hipGraphMemcpyNodeSetParams_Functional") { + constexpr int width{SIZE}, height{SIZE}, depth{SIZE}; + hipArray *devArray; + hipChannelFormatKind formatKind = hipChannelFormatKindSigned; + hipMemcpy3DParms myparms, myparms1; + uint32_t size = width * height * depth * sizeof(int); + + int *hData = reinterpret_cast(malloc(size)); + REQUIRE(hData != nullptr); + memset(hData, 0, size); + int *hDataTemp = reinterpret_cast(malloc(size)); + REQUIRE(hDataTemp != nullptr); + memset(hDataTemp, 0, size); + int *hOutputData = reinterpret_cast(malloc(size)); + REQUIRE(hOutputData != nullptr); + memset(hOutputData, 0, size); + int *hOutputData1 = reinterpret_cast(malloc(size)); + REQUIRE(hOutputData1 != nullptr); + memset(hOutputData1, 0, size); + + for (int i = 0; i < depth; i++) { + for (int j = 0; j < height; j++) { + for (int k = 0; k < width; k++) { + hData[i*width*height + j*width + k] = i*width*height + j*width + k; + } + } + } + hipChannelFormatDesc channelDesc = hipCreateChannelDesc(sizeof(int)*8, + 0, 0, 0, formatKind); + HIP_CHECK(hipMalloc3DArray(&devArray, &channelDesc, make_hipExtent(width, + height, depth), hipArrayDefault)); + memset(&myparms, 0x0, sizeof(hipMemcpy3DParms)); + + // Host to Device + myparms.srcPos = make_hipPos(0, 0, 0); + myparms.dstPos = make_hipPos(0, 0, 0); + myparms.extent = make_hipExtent(width , height, depth); + myparms.srcPtr = make_hipPitchedPtr(hData, width * sizeof(int), + width, height); + myparms.dstArray = devArray; + myparms.kind = hipMemcpyHostToDevice; + + hipGraph_t graph; + hipGraphNode_t memcpyNode; + std::vector dependencies; + hipStream_t streamForGraph; + hipGraphExec_t graphExec; + + HIP_CHECK(hipStreamCreate(&streamForGraph)); + HIP_CHECK(hipGraphCreate(&graph, 0)); + HIP_CHECK(hipGraphAddMemcpyNode(&memcpyNode, graph, NULL, 0, &myparms)); + dependencies.push_back(memcpyNode); + + // Device to host + memset(&myparms1, 0x0, sizeof(hipMemcpy3DParms)); + myparms1.srcPos = make_hipPos(0, 0, 0); + myparms1.dstPos = make_hipPos(0, 0, 0); + myparms1.dstPtr = make_hipPitchedPtr(hDataTemp, width * sizeof(int), + width, height); + myparms1.srcArray = devArray; + myparms1.extent = make_hipExtent(width, height, depth); + myparms1.kind = hipMemcpyDeviceToHost; + + HIP_CHECK(hipGraphAddMemcpyNode(&memcpyNode, graph, dependencies.data(), + dependencies.size(), &myparms1)); + + SECTION("Update the memcpyNode and check") { + // Device to host with updated host ptr hDataTemp -> hOutputData + memset(&myparms1, 0x0, sizeof(hipMemcpy3DParms)); + myparms1.srcPos = make_hipPos(0, 0, 0); + myparms1.dstPos = make_hipPos(0, 0, 0); + myparms1.dstPtr = make_hipPitchedPtr(hOutputData, width * sizeof(int), + width, height); + myparms1.srcArray = devArray; + myparms1.extent = make_hipExtent(width, height, depth); + myparms1.kind = hipMemcpyDeviceToHost; + + HIP_CHECK(hipGraphMemcpyNodeSetParams(memcpyNode, &myparms1)); + + // Instantiate and launch the graph + HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); + HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph)); + HIP_CHECK(hipStreamSynchronize(streamForGraph)); + + // Check result + HipTest::checkArray(hData, hOutputData, width, height, depth); + } + + SECTION("Update the memcpyNode again and check") { + // Device to host with updated host ptr hOutputData -> hOutputData1 + memset(&myparms1, 0x0, sizeof(hipMemcpy3DParms)); + myparms1.srcPos = make_hipPos(0, 0, 0); + myparms1.dstPos = make_hipPos(0, 0, 0); + myparms1.dstPtr = make_hipPitchedPtr(hOutputData1, width * sizeof(int), + width, height); + myparms1.srcArray = devArray; + myparms1.extent = make_hipExtent(width, height, depth); + myparms1.kind = hipMemcpyDeviceToHost; + + HIP_CHECK(hipGraphAddMemcpyNode(&memcpyNode, graph, dependencies.data(), + dependencies.size(), &myparms1)); + HIP_CHECK(hipGraphMemcpyNodeSetParams(memcpyNode, &myparms1)); + + // Instantiate and launch the graph + HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); + HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph)); + HIP_CHECK(hipStreamSynchronize(streamForGraph)); + + // Check result + HipTest::checkArray(hData, hOutputData1, width, height, depth); + } + HIP_CHECK(hipGraphExecDestroy(graphExec)); + HIP_CHECK(hipGraphDestroy(graph)); + HIP_CHECK(hipStreamDestroy(streamForGraph)); + HIP_CHECK(hipFreeArray(devArray)); + free(hData); + free(hDataTemp); + free(hOutputData); + free(hOutputData1); +} diff --git a/tests/catch/unit/memory/CMakeLists.txt b/tests/catch/unit/memory/CMakeLists.txt index af72155549..997572ee6d 100644 --- a/tests/catch/unit/memory/CMakeLists.txt +++ b/tests/catch/unit/memory/CMakeLists.txt @@ -1,4 +1,4 @@ -# Copyright (c) 2021 Advanced Micro Devices, Inc. All Rights Reserved. +# Copyright (c) 2022 Advanced Micro Devices, Inc. All Rights Reserved. # # Permission is hereby granted, free of charge, to any person obtaining a copy # of this software and associated documentation files (the "Software"), to deal @@ -21,6 +21,7 @@ # Common Tests - Test independent of all platforms if(HIP_PLATFORM MATCHES "amd") set(TEST_SRC + DriverContext.cc memset.cc malloc.cc hipMemcpy2DToArray.cc @@ -38,6 +39,8 @@ set(TEST_SRC hipMemcpyAllApiNegative.cc hipMemcpy_MultiThread.cc hipHostRegister.cc + hipHostUnregister.cc + hipMallocPitch.cc hipMemPtrGetInfo.cc hipPointerGetAttributes.cc hipHostGetFlags.cc @@ -73,11 +76,13 @@ set(TEST_SRC hipHostMalloc.cc hipMemcpy.cc hipMemcpyAsync.cc + hipMemsetFunctional.cc hipMallocPitch.cc hipMallocArray.cc hipMalloc3D.cc hipMalloc3DArray.cc hipArrayCreate.cc + hipArray3DCreate.cc hipDrvMemcpy3D.cc hipDrvMemcpy3DAsync.cc hipPointerGetAttribute.cc @@ -85,9 +90,11 @@ set(TEST_SRC hipMallocMngdMultiThread.cc hipMemPrefetchAsync.cc hipArray.cc + hipMemVmm.cc ) else() set(TEST_SRC + DriverContext.cc memset.cc malloc.cc hipMemcpy2DToArray.cc @@ -105,6 +112,8 @@ set(TEST_SRC hipMemcpyAllApiNegative.cc hipMemcpy_MultiThread.cc hipHostRegister.cc + hipHostUnregister.cc + hipMallocPitch.cc hipHostGetFlags.cc hipHostGetDevicePointer.cc hipMallocManaged_MultiScenario.cc @@ -136,11 +145,13 @@ set(TEST_SRC hipHostMalloc.cc hipMemcpy.cc hipMemcpyAsync.cc + hipMemsetFunctional.cc hipMallocPitch.cc hipMallocArray.cc hipMalloc3D.cc hipMalloc3DArray.cc hipArrayCreate.cc + hipArray3DCreate.cc hipDrvMemcpy3D.cc hipDrvMemcpy3DAsync.cc hipPointerGetAttribute.cc @@ -159,4 +170,5 @@ endif() hip_add_exe_to_target(NAME MemoryTest TEST_SRC ${TEST_SRC} - TEST_TARGET_NAME build_tests) + TEST_TARGET_NAME build_tests + COMPILE_OPTIONS -std=c++14) diff --git a/tests/catch/unit/memory/DriverContext.cc b/tests/catch/unit/memory/DriverContext.cc new file mode 100644 index 0000000000..dc27a62000 --- /dev/null +++ b/tests/catch/unit/memory/DriverContext.cc @@ -0,0 +1,36 @@ +/* +Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved. + +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#include "DriverContext.hh" +#include + +DriverContext::DriverContext() { + HIP_CHECK(hipInit(0)); + HIP_CHECK(hipDeviceGet(&device, 0)); + HIP_CHECK(hipDevicePrimaryCtxRetain(&ctx, device)); + HIP_CHECK(hipCtxPushCurrent(ctx)); +} + +DriverContext::~DriverContext() { + HIP_CHECK(hipCtxPopCurrent(&ctx)); + HIP_CHECK(hipDevicePrimaryCtxRelease(device)); +} diff --git a/tests/catch/unit/memory/DriverContext.hh b/tests/catch/unit/memory/DriverContext.hh new file mode 100644 index 0000000000..5593c512d4 --- /dev/null +++ b/tests/catch/unit/memory/DriverContext.hh @@ -0,0 +1,39 @@ +/* +Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved. + +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#pragma once + +#include + +class DriverContext { + private: + hipCtx_t ctx; + hipDevice_t device; + + public: + DriverContext(); + ~DriverContext(); + + // Rule of three + DriverContext(const DriverContext& other) = delete; + DriverContext(DriverContext&& other) noexcept = delete; +}; diff --git a/tests/catch/unit/memory/hipArray3DCreate.cc b/tests/catch/unit/memory/hipArray3DCreate.cc new file mode 100644 index 0000000000..faa3822b9b --- /dev/null +++ b/tests/catch/unit/memory/hipArray3DCreate.cc @@ -0,0 +1,321 @@ +/* +Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved. +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#include +#include "DriverContext.hh" +#include "hipArrayCommon.hh" +#include "hip_test_common.hh" + +namespace { +void checkArrayIsExpected(const hiparray array, const HIP_ARRAY3D_DESCRIPTOR& expected_desc) { +// hipArray3DGetDescriptor doesn't currently exist (EXSWCPHIPT-87) +#if HT_AMD + std::ignore = array; + std::ignore = expected_desc; +#else + CUDA_ARRAY3D_DESCRIPTOR queried_desc; + cuArray3DGetDescriptor(&queried_desc, array); + + REQUIRE(queried_desc.Width == expected_desc.Width); + REQUIRE(queried_desc.Height == expected_desc.Height); + REQUIRE(queried_desc.Depth == expected_desc.Depth); + REQUIRE(queried_desc.Format == expected_desc.Format); + REQUIRE(queried_desc.NumChannels == expected_desc.NumChannels); + REQUIRE(queried_desc.Flags == expected_desc.Flags); +#endif +} + +void testInvalidDescription(HIP_ARRAY3D_DESCRIPTOR desc) { + hiparray array; + HIP_CHECK_ERROR(hipArray3DCreate(&array, &desc), hipErrorInvalidValue); +} +} // namespace + +TEMPLATE_TEST_CASE("Unit_hipArray3DCreate_happy", "", char, uchar2, uint2, int4, short4, float, + float2, float4) { + using vec_info = vector_info; + DriverContext ctx; + + HIP_ARRAY3D_DESCRIPTOR desc{}; + desc.Format = vec_info::format; + desc.NumChannels = vec_info::size; +#if HT_AMD + desc.Flags = 0; +#else + desc.Flags = GENERATE(0, CUDA_ARRAY3D_SURFACE_LDST); +#endif + + constexpr size_t size = 64; + + std::vector extents{ + {size, 0, 0}, // 1D array + {size, size, 0}, // 2D array + {size, size, size} // 3D array + }; + + for (auto& extent : extents) { + desc.Width = extent.width; + desc.Height = extent.height; + desc.Depth = extent.depth; + + CAPTURE(desc.Width, desc.Height, desc.Depth); + + hiparray array; + HIP_CHECK(hipArray3DCreate(&array, &desc)); + checkArrayIsExpected(array, desc); + HIP_CHECK(hipArrayDestroy(array)); + } +} + +TEMPLATE_TEST_CASE("Unit_hipArray3DCreate_MaxTexture", "", int, uint4, short, ushort2, + unsigned char, float, float4) { +#if HT_AMD + HipTest::HIP_SKIP_TEST("EXSWCPHIPT-97"); + return; +#endif + + using vec_info = vector_info; + DriverContext ctx; + + hiparray array; + HIP_ARRAY3D_DESCRIPTOR desc{}; + desc.Format = vec_info::format; + desc.NumChannels = vec_info::size; +#if HT_AMD + desc.Flags = 0; +#else + desc.Flags = GENERATE(0, CUDA_ARRAY3D_SURFACE_LDST); + if (desc.Flags == CUDA_ARRAY3D_SURFACE_LDST) { + HipTest::HIP_SKIP_TEST("EXSWCPHIPT-58"); + return; + } +#endif + CAPTURE(desc.Flags); + + const Sizes sizes(desc.Flags); + CAPTURE(sizes.max1D, sizes.max2D, sizes.max3D); + + const size_t s = 64; + SECTION("Happy") { + // stored in a vector so some values can be ifdef'd out + std::vector extentsToTest{ + make_hipExtent(sizes.max1D, 0, 0), // 1D max + make_hipExtent(sizes.max2D[0], s, 0), // 2D max width + make_hipExtent(s, sizes.max2D[1], 0), // 2D max height + make_hipExtent(sizes.max2D[0], sizes.max2D[1], 0), // 2D max + make_hipExtent(sizes.max3D[0], s, s), // 3D max width + make_hipExtent(s, sizes.max3D[1], s), // 3D max height + make_hipExtent(s, s, sizes.max3D[2]), // 3D max depth + make_hipExtent(s, sizes.max3D[1], sizes.max3D[2]), // 3D max height and depth + make_hipExtent(sizes.max3D[0], s, sizes.max3D[2]), // 3D max width and depth + make_hipExtent(sizes.max3D[0], sizes.max3D[1], s), // 3D max width and height + make_hipExtent(sizes.max3D[0], sizes.max3D[1], sizes.max3D[2]) // 3D max + }; + const auto extent = + GENERATE_COPY(from_range(std::begin(extentsToTest), std::end(extentsToTest))); + + desc.Width = extent.width; + desc.Height = extent.height; + desc.Depth = extent.depth; + + CAPTURE(desc.Width, desc.Height, desc.Depth); + + auto maxArrayCreateError = hipArray3DCreate(&array, &desc); + // this can try to alloc many GB of memory, so out of memory is acceptable + if (maxArrayCreateError == hipErrorOutOfMemory) return; + HIP_CHECK(maxArrayCreateError); + checkArrayIsExpected(array, desc); + HIP_CHECK(hipArrayDestroy(array)); + } + SECTION("Negative") { + std::vector extentsToTest { + make_hipExtent(sizes.max1D + 1, 0, 0), // 1D max + make_hipExtent(sizes.max2D[0] + 1, s, 0), // 2D max width + make_hipExtent(s, sizes.max2D[1] + 1, 0), // 2D max height + make_hipExtent(sizes.max2D[0] + 1, sizes.max2D[1] + 1, 0), // 2D max + make_hipExtent(sizes.max3D[0] + 1, s, s), // 3D max width + make_hipExtent(s, sizes.max3D[1] + 1, s), // 3D max height +#if !HT_NVIDIA // leads to hipSuccess on NVIDIA + make_hipExtent(s, s, sizes.max3D[2] + 1), // 3D max depth +#endif + make_hipExtent(s, sizes.max3D[1] + 1, sizes.max3D[2] + 1), // 3D max height and depth + make_hipExtent(sizes.max3D[0] + 1, s, sizes.max3D[2] + 1), // 3D max width and depth + make_hipExtent(sizes.max3D[0] + 1, sizes.max3D[1] + 1, s), // 3D max width and height + make_hipExtent(sizes.max3D[0] + 1, sizes.max3D[1] + 1, sizes.max3D[2] + 1) // 3D max + }; + const auto extent = + GENERATE_COPY(from_range(std::begin(extentsToTest), std::end(extentsToTest))); + + desc.Width = extent.width; + desc.Height = extent.height; + desc.Depth = extent.depth; + + CAPTURE(desc.Width, desc.Height, desc.Depth); + + HIP_CHECK_ERROR(hipArray3DCreate(&array, &desc), hipErrorInvalidValue); + } +} + +#if HT_NVIDIA +constexpr std::array validFlags{ + 0, + hipArraySurfaceLoadStore, + hipArrayLayered, + hipArrayLayered | hipArraySurfaceLoadStore, + hipArrayCubemap, + hipArrayCubemap | hipArrayLayered, + hipArrayCubemap | hipArraySurfaceLoadStore, + hipArrayCubemap | hipArrayLayered | hipArraySurfaceLoadStore, + hipArrayTextureGather}; +#else +constexpr std::array validFlags{ + 0, hipArrayCubemap, hipArrayCubemap | hipArrayLayered, + hipArrayCubemap | hipArraySurfaceLoadStore, + hipArrayCubemap | hipArrayLayered | hipArraySurfaceLoadStore}; +#endif + +constexpr HIP_ARRAY3D_DESCRIPTOR defaultDescriptor(unsigned int flags, size_t size) { + HIP_ARRAY3D_DESCRIPTOR desc{}; + desc.Format = HIP_AD_FORMAT_FLOAT; + desc.NumChannels = 4; + desc.Flags = flags; + desc.Width = size; + desc.Height = size; + desc.Depth = size; + +#if HT_NVIDIA + if (flags == CUDA_ARRAY3D_TEXTURE_GATHER) { + desc.Depth = 0; + } +#endif + return desc; +} + +// Providing the array pointer as nullptr should return an error +TEST_CASE("Unit_hipArray3DCreate_Negative_NullArrayPtr") { + auto desc = defaultDescriptor(0, 64); + + DriverContext ctx; + HIP_CHECK_ERROR(hipArray3DCreate(nullptr, &desc), hipErrorInvalidValue); +} + +// Providing the description pointer as nullptr should return an error +TEST_CASE("Unit_hipArray3DCreate_Negative_NullDescPtr") { +#if HT_AMD + HipTest::HIP_SKIP_TEST("EXSWCPHIPT-130"); + return; +#endif + + DriverContext ctx; + hiparray array; + HIP_CHECK_ERROR(hipArray3DCreate(&array, nullptr), hipErrorInvalidValue); +} + + +// Zero width arrays are not allowed +TEST_CASE("Unit_hipArray3DCreate_Negative_ZeroWidth") { + DriverContext ctx; + + unsigned int flags = GENERATE(from_range(std::begin(validFlags), std::end(validFlags))); + auto desc = defaultDescriptor(flags, 6); + desc.Width = 0; + CAPTURE(desc.Flags); + + testInvalidDescription(desc); +} + +// Zero height arrays are only allowed for 1D arrays and layered arrays +TEST_CASE("Unit_hipArray3DCreate_Negative_ZeroHeight") { + DriverContext ctx; + + unsigned int flags = GENERATE(from_range(std::begin(validFlags), std::end(validFlags))); + auto desc = defaultDescriptor(flags, 6); +#if HT_NVIDIA + std::array exceptions{CUDA_ARRAY3D_LAYERED, + CUDA_ARRAY3D_LAYERED | CUDA_ARRAY3D_SURFACE_LDST}; +#else + std::array exceptions{}; +#endif + desc.Height = 0; + + if (std::find(std::begin(exceptions), std::end(exceptions), desc.Flags) == std::end(exceptions)) { + // flag is not in list of exceptions + testInvalidDescription(desc); + } +} + +// Arrays must be created with a valid data format +TEST_CASE("Unit_hipArray3DCreate_Negative_InvalidFormat") { + DriverContext ctx; + + unsigned int flags = GENERATE(from_range(std::begin(validFlags), std::end(validFlags))); + auto desc = defaultDescriptor(flags, 6); + + desc.Format = static_cast(0xDEADBEEF); + REQUIRE(std::find(std::begin(driverFormats), std::end(driverFormats), desc.Format) == + std::end(driverFormats)); + + testInvalidDescription(desc); +} + +// An array must have either 1,2, or 4 channels +TEST_CASE("Unit_hipArray3DCreate_Negative_NumChannels") { + DriverContext ctx; + unsigned int flags = GENERATE(from_range(std::begin(validFlags), std::end(validFlags))); + auto desc = defaultDescriptor(flags, 6); + desc.NumChannels = GENERATE(0, 3, 5); + + testInvalidDescription(desc); +} + +// Using invalid flags should result in an error +TEST_CASE("Unit_hipArray3DCreate_Negative_InvalidFlags") { + DriverContext ctx; + + // FIXME: use the same flags for both tests when the values exist for hip +#if HT_NVIDIA + unsigned int flags = + GENERATE(0xDEADBEEF, CUDA_ARRAY3D_TEXTURE_GATHER | CUDA_ARRAY3D_SURFACE_LDST, + CUDA_ARRAY3D_TEXTURE_GATHER | CUDA_ARRAY3D_CUBEMAP, + CUDA_ARRAY3D_TEXTURE_GATHER | CUDA_ARRAY3D_SURFACE_LDST | CUDA_ARRAY3D_CUBEMAP); +#else + unsigned int flags = 0xDEADBEEF; +#endif + + CAPTURE(flags); + + auto desc = defaultDescriptor(flags, 6); + + + REQUIRE(std::find(std::begin(validFlags), std::end(validFlags), desc.Flags) == + std::end(validFlags)); + + testInvalidDescription(desc); +} + + +// hipArray3DCreate should handle the max numeric value gracefully. +TEST_CASE("Unit_hipArray3DCreate_Negative_NumericLimit") { + DriverContext ctx; + + unsigned int flags = GENERATE(from_range(std::begin(validFlags), std::end(validFlags))); + auto desc = defaultDescriptor(flags, std::numeric_limits::max()); + + testInvalidDescription(desc); +} diff --git a/tests/catch/unit/memory/hipArrayCommon.hh b/tests/catch/unit/memory/hipArrayCommon.hh new file mode 100644 index 0000000000..8502ac218d --- /dev/null +++ b/tests/catch/unit/memory/hipArrayCommon.hh @@ -0,0 +1,211 @@ +/* +Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved. + +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#pragma once + +#include + +constexpr size_t BlockSize = 16; + +template struct type_and_size_and_format { + using type = T; + static constexpr size_t size = N; + static constexpr hipArray_Format format = Format; +}; + +// Create a map of type to scalar type, vector size and scalar type format enum. +// This is useful for creating simpler function that depend on the vector size. +template struct vector_info; +template <> +struct vector_info : type_and_size_and_format {}; +template <> struct vector_info : type_and_size_and_format {}; +template <> +struct vector_info : type_and_size_and_format {}; +template <> +struct vector_info : type_and_size_and_format {}; +template <> +struct vector_info + : type_and_size_and_format {}; +template <> +struct vector_info + : type_and_size_and_format {}; +template <> +struct vector_info + : type_and_size_and_format {}; + +template <> +struct vector_info : type_and_size_and_format {}; +template <> struct vector_info : type_and_size_and_format {}; +template <> +struct vector_info : type_and_size_and_format {}; +template <> +struct vector_info : type_and_size_and_format {}; +template <> +struct vector_info + : type_and_size_and_format {}; +template <> +struct vector_info + : type_and_size_and_format {}; +template <> +struct vector_info + : type_and_size_and_format {}; + +template <> +struct vector_info : type_and_size_and_format {}; +template <> struct vector_info : type_and_size_and_format {}; +template <> +struct vector_info : type_and_size_and_format {}; +template <> +struct vector_info : type_and_size_and_format {}; +template <> +struct vector_info + : type_and_size_and_format {}; +template <> +struct vector_info + : type_and_size_and_format {}; +template <> +struct vector_info + : type_and_size_and_format {}; + +// read from a texture using normalized coordinates +constexpr size_t ChannelToRead = 1; +template +__global__ void readFromTexture(T* output, hipTextureObject_t texObj, size_t width, size_t height, + bool textureGather) { + // Calculate normalized texture coordinates + const unsigned int x = blockIdx.x * blockDim.x + threadIdx.x; + const unsigned int y = blockIdx.y * blockDim.y + threadIdx.y; + const float u = x / (float)width; + + // Read from texture and write to global memory + if (height == 0) { + output[x] = tex1D(texObj, u); + } else { + const float v = y / (float)height; + output[y * width + x] = + textureGather ? tex2Dgather(texObj, u, v, ChannelToRead) : tex2D(texObj, u, v); + } +} + +template void checkDataIsAscending(const std::vector& hostData) { + bool allMatch = true; + size_t i = 0; + for (; i < hostData.size(); ++i) { + allMatch = allMatch && hostData[i] == static_cast(i); + if (!allMatch) break; + } + INFO("hostData[" << i << "] == " << static_cast(hostData[i])); + REQUIRE(allMatch); +} + +inline size_t getFreeMem() { + size_t free = 0, total = 0; + HIP_CHECK(hipMemGetInfo(&free, &total)); + return free; +} + +struct Sizes { + int max1D; + std::array max2D; + std::array max3D; + + Sizes(unsigned int flag) { + int device; + HIP_CHECK(hipGetDevice(&device)); + static_assert( + hipArrayDefault == 0, + "hipArrayDefault is assumed to be equivalent to 0 for the following switch statment"); +#if HT_NVIDIA + static_assert(hipArraySurfaceLoadStore == CUDA_ARRAY3D_SURFACE_LDST, + "hipArraySurface is assumed to be equivalent to CUDA_ARRAY3D_SURFACE_LDST for " + "the following switch statment"); +#endif + switch (flag) { + case hipArrayDefault: { // 0 + hipDeviceProp_t prop; + HIP_CHECK(hipGetDeviceProperties(&prop, device)); + max1D = prop.maxTexture1D; + max2D = {prop.maxTexture2D[0], prop.maxTexture2D[1]}; + max3D = {prop.maxTexture3D[0], prop.maxTexture3D[1], prop.maxTexture3D[2]}; + return; + } + case hipArraySurfaceLoadStore: { // CUDA_ARRAY3D_SURFACE_LDST + int value; + HIP_CHECK(hipDeviceGetAttribute(&value, hipDeviceAttributeMaxSurface1D, device)); + max1D = value; + HIP_CHECK(hipDeviceGetAttribute(&value, hipDeviceAttributeMaxSurface2D, device)); + max2D = {value, value}; + HIP_CHECK(hipDeviceGetAttribute(&value, hipDeviceAttributeMaxSurface3D, device)); + max3D = {value, value, value}; + return; + } + default: { + INFO("Array flag not supported"); + REQUIRE(false); + return; + } + } + } +}; + +inline const char* channelFormatString(hipChannelFormatKind formatKind) noexcept { + switch (formatKind) { + case hipChannelFormatKindFloat: + return "float"; + case hipChannelFormatKindSigned: + return "signed"; + case hipChannelFormatKindUnsigned: + return "unsigned"; + default: + return "error"; + } +} + +// All the possible formats for channel data in an array. +static const std::vector driverFormats{ + HIP_AD_FORMAT_UNSIGNED_INT8, HIP_AD_FORMAT_UNSIGNED_INT16, HIP_AD_FORMAT_UNSIGNED_INT32, + HIP_AD_FORMAT_SIGNED_INT8, HIP_AD_FORMAT_SIGNED_INT16, HIP_AD_FORMAT_SIGNED_INT32, + HIP_AD_FORMAT_HALF, HIP_AD_FORMAT_FLOAT}; + +// Helpful for printing errors +inline const char* formatToString(hipArray_Format f) { + switch (f) { + case HIP_AD_FORMAT_UNSIGNED_INT8: + return "Unsigned Int 8"; + case HIP_AD_FORMAT_UNSIGNED_INT16: + return "Unsigned Int 16"; + case HIP_AD_FORMAT_UNSIGNED_INT32: + return "Unsigned Int 32"; + case HIP_AD_FORMAT_SIGNED_INT8: + return "Signed Int 8"; + case HIP_AD_FORMAT_SIGNED_INT16: + return "Signed Int 16"; + case HIP_AD_FORMAT_SIGNED_INT32: + return "Signed Int 32"; + case HIP_AD_FORMAT_HALF: + return "Float 16"; + case HIP_AD_FORMAT_FLOAT: + return "Float 32"; + default: + return "not found"; + } +} diff --git a/tests/catch/unit/memory/hipArrayCreate.cc b/tests/catch/unit/memory/hipArrayCreate.cc index 3bbf6cbd54..836c0d1c45 100644 --- a/tests/catch/unit/memory/hipArrayCreate.cc +++ b/tests/catch/unit/memory/hipArrayCreate.cc @@ -24,12 +24,16 @@ hipArrayCreate API test scenarios 3. Multithreaded scenario */ +#include +#include #include +#include "hipArrayCommon.hh" +#include "DriverContext.hh" -static constexpr auto NUM_W{4}; -static constexpr auto BIGNUM_W{100}; -static constexpr auto NUM_H{4}; -static constexpr auto BIGNUM_H{100}; +static constexpr size_t NUM_W{4}; +static constexpr size_t BIGNUM_W{100}; +static constexpr size_t NUM_H{4}; +static constexpr size_t BIGNUM_H{100}; static constexpr auto ARRAY_LOOP{100}; /* @@ -47,77 +51,35 @@ static constexpr auto ARRAY_LOOP{100}; */ static void ArrayCreate_DiffSizes(int gpu) { - HIP_CHECK(hipSetDevice(gpu)); - std::vector array_size; - array_size.push_back(NUM_W); - array_size.push_back(BIGNUM_W); - for (auto &size : array_size) { - HIP_ARRAY array[ARRAY_LOOP]; - size_t tot, avail, ptot, pavail; - HIP_CHECK(hipMemGetInfo(&pavail, &ptot)); - for (int i = 0; i < ARRAY_LOOP; i++) { - HIP_ARRAY_DESCRIPTOR desc; - desc.NumChannels = 1; - if (size == NUM_W) { - desc.Width = NUM_W; - desc.Height = NUM_H; - } else { - desc.Width = BIGNUM_W; - desc.Height = BIGNUM_H; - } - desc.Format = HIP_AD_FORMAT_FLOAT; - HIP_CHECK(hipArrayCreate(&array[i], &desc)); - } - for (int i = 0; i < ARRAY_LOOP; i++) { - ARRAY_DESTROY(array[i]); - } - HIP_CHECK(hipMemGetInfo(&avail, &tot)); - if ((pavail != avail)) { - HIPASSERT(false); - } - } -} + HIP_CHECK_THREAD(hipSetDevice(gpu)); + std::pair size = + GENERATE(std::make_pair(NUM_W, NUM_H), std::make_pair(BIGNUM_W, BIGNUM_H)); + std::array array; + size_t pavail, avail; + HIP_CHECK_THREAD(hipMemGetInfo(&pavail, nullptr)); + HIP_ARRAY_DESCRIPTOR desc; + desc.NumChannels = 1; + desc.Width = std::get<0>(size); + desc.Height = std::get<1>(size); + desc.Format = HIP_AD_FORMAT_FLOAT; -/*Thread function*/ -static void ArrayCreateThreadFunc(int gpu) { - ArrayCreate_DiffSizes(gpu); + for (int i = 0; i < ARRAY_LOOP; i++) { + HIP_CHECK_THREAD(hipArrayCreate(&array[i], &desc)); + } + for (int i = 0; i < ARRAY_LOOP; i++) { + HIP_CHECK_THREAD(hipArrayDestroy(array[i])); + } + + HIP_CHECK_THREAD(hipMemGetInfo(&avail, nullptr)); + REQUIRE_THREAD(pavail == avail); } /* This testcase verifies hipArrayCreate API for small and big chunks data*/ TEST_CASE("Unit_hipArrayCreate_DiffSizes") { ArrayCreate_DiffSizes(0); + HIP_CHECK_THREAD_FINALIZE(); } - -/* This testcase verifies the negative scenarios of - * hipArrayCreate API - */ -TEST_CASE("Unit_hipArrayCreate_Negative") { - HIP_ARRAY_DESCRIPTOR desc; - HIP_ARRAY array; - desc.Format = HIP_AD_FORMAT_FLOAT; - desc.NumChannels = 1; - desc.Width = NUM_W; - desc.Height = NUM_H; -#if HT_NVIDIA - SECTION("NullPointer to Array") { - REQUIRE(hipArrayCreate(nullptr, &desc) != hipSuccess); - } - - SECTION("NullPointer to Channel Descriptor") { - REQUIRE(hipArrayCreate(&array, nullptr) != hipSuccess); - } -#endif - SECTION("Width 0 for Array Descriptor") { - desc.Width = 0; - REQUIRE(hipArrayCreate(&array, &desc) != hipSuccess); - } - - SECTION("Invalid NumChannels") { - desc.NumChannels = 3; - REQUIRE(hipArrayCreate(&array, &desc) != hipSuccess); - } -} /* This testcase verifies the hipArrayCreate API in multithreaded scenario by launching threads in parallel on multiple GPUs @@ -129,16 +91,16 @@ TEST_CASE("Unit_hipArrayCreate_MultiThread") { devCnt = HipTest::getDeviceCount(); - size_t tot, avail, ptot, pavail; - HIP_CHECK(hipMemGetInfo(&pavail, &ptot)); + const size_t pavail = getFreeMem(); for (int i = 0; i < devCnt; i++) { - threadlist.push_back(std::thread(ArrayCreateThreadFunc, i)); + threadlist.push_back(std::thread(ArrayCreate_DiffSizes, i)); } - for (auto &t : threadlist) { + for (auto& t : threadlist) { t.join(); } - HIP_CHECK(hipMemGetInfo(&avail, &tot)); + HIP_CHECK_THREAD_FINALIZE(); + const size_t avail = getFreeMem(); if (pavail != avail) { WARN("Memory leak of hipMalloc3D API in multithreaded scenario"); @@ -146,3 +108,273 @@ TEST_CASE("Unit_hipArrayCreate_MultiThread") { } } + +// Tests ///////////////////////////////////////// + +#if HT_AMD +constexpr auto MemoryTypeHost = hipMemoryTypeHost; +constexpr auto MemoryTypeArray = hipMemoryTypeArray; +constexpr auto NORMALIZED_COORDINATES = HIP_TRSF_NORMALIZED_COORDINATES; +constexpr auto READ_AS_INTEGER = HIP_TRSF_READ_AS_INTEGER; +#else +constexpr auto MemoryTypeHost = CU_MEMORYTYPE_HOST; +constexpr auto MemoryTypeArray = CU_MEMORYTYPE_ARRAY; +// (EXSWCPHIPT-92) HIP equivalents not defined for CUDA backend. +constexpr auto NORMALIZED_COORDINATES = CU_TRSF_NORMALIZED_COORDINATES; +constexpr auto READ_AS_INTEGER = CU_TRSF_READ_AS_INTEGER; +#endif + +// Copy data from host to the hiparray, accounting 1D or 2D arrays +template +void copyToArray(hiparray dst, const std::vector& src, const size_t height) { + const auto sizeInBytes = src.size() * sizeof(T); + if (height == 0) { + // FIXME(EXSWCPHIPT-64) remove cast when API is fixed (will require major version change) + HIP_CHECK(hipMemcpyHtoA(reinterpret_cast(dst), 0, src.data(), sizeInBytes)); + } else { + const auto pitch = sizeInBytes / height; + hip_Memcpy2D copyParams{}; + copyParams.srcMemoryType = MemoryTypeHost; + copyParams.srcXInBytes = 0; // x offset + copyParams.srcY = 0; // y offset + copyParams.srcHost = src.data(); + copyParams.srcPitch = pitch; + + + copyParams.dstMemoryType = MemoryTypeArray; + copyParams.dstXInBytes = 0; // x offset + copyParams.dstY = 0; // y offset + copyParams.dstArray = dst; + + copyParams.WidthInBytes = pitch; + copyParams.Height = height; + + HIP_CHECK(hipMemcpyParam2D(©Params)); + } +} + +// Test the allocated array by generating a texture from it then reading from that texture. +// Textures are read-only, so write to the array then copy that into normal device memory. +template +void testArrayAsTexture(hiparray array, const size_t width, const size_t height) { + using vec_info = vector_info; + using scalar_type = typename vec_info::type; + const auto h = height ? height : 1; + const auto size = sizeof(T) * width * h; + + // set hip array + std::vector hostData(width * h * vec_info::size); + // assigned ascending values to the data array to show indexing is working + std::iota(std::begin(hostData), std::end(hostData), 0); + + copyToArray(array, hostData, height); + + // create texture + hipTextureObject_t textObj{}; + + HIP_RESOURCE_DESC resDesc{}; + memset(&resDesc, 0, sizeof(HIP_RESOURCE_DESC)); + resDesc.resType = HIP_RESOURCE_TYPE_ARRAY; + resDesc.res.array.hArray = array; + resDesc.flags = 0; + + HIP_TEXTURE_DESC texDesc{}; + memset(&texDesc, 0, sizeof(HIP_TEXTURE_DESC)); + // use the actual values in the texture, not normalized data + texDesc.filterMode = HIP_TR_FILTER_MODE_POINT; + // Use normalized coordinates and also read the data in the original data type + texDesc.flags |= NORMALIZED_COORDINATES | READ_AS_INTEGER; + + HIP_CHECK(hipTexObjectCreate(&textObj, &resDesc, &texDesc, nullptr)); + + // run kernel + T* device_data{}; + HIP_CHECK(hipMalloc(&device_data, size)); + readFromTexture<<>>(device_data, textObj, width, + height, false); + HIP_CHECK(hipGetLastError()); // check for errors when running the kernel + + // copy data back and then test it + std::fill(std::begin(hostData), std::end(hostData), 0); + HIP_CHECK(hipMemcpy(hostData.data(), device_data, size, hipMemcpyDeviceToHost)); + + checkDataIsAscending(hostData); + + // clean up + HIP_CHECK(hipTexObjectDestroy(textObj)); + HIP_CHECK(hipFree(device_data)); +} + +// Selection of types chosen since trying all types would be slow to compile +// Test the happy path of the hipArrayCreate +TEMPLATE_TEST_CASE("Unit_hipArrayCreate_happy", "", uint, int, int4, ushort, short2, char, uchar2, + char4, float, float2, float4) { +#if HT_AMD + if (std::is_same::value || std::is_same::value || + std::is_same::value) { + HipTest::HIP_SKIP_TEST("Probably EXSWCPHIPT-62"); + return; + } +#endif + using vec_info = vector_info; + DriverContext ctx; + + HIP_ARRAY_DESCRIPTOR desc; + desc.Format = vec_info::format; + desc.NumChannels = vec_info::size; + desc.Width = 1024; + desc.Height = GENERATE(0, 1024); + + size_t initFree = getFreeMem(); + + // pointer to the array in device memory + hiparray array{}; + + HIP_CHECK(hipArrayCreate(&array, &desc)); + + testArrayAsTexture(array, desc.Width, desc.Height); + + size_t finalFree = getFreeMem(); + + const size_t allocSize = sizeof(TestType) * desc.Width * (desc.Height ? desc.Height : 1); + // will be aligned to some size, so this is not exact + REQUIRE(initFree - finalFree >= allocSize); + + HIP_CHECK(hipArrayDestroy(array)); +} + + +// Only widths and Heights up to the maxTexture size is supported +TEMPLATE_TEST_CASE("Unit_hipArrayCreate_maxTexture", "", uint, int, int4, ushort, short2, char, + uchar2, char4, float, float2, float4) { + using vec_info = vector_info; + DriverContext ctx; + + HIP_ARRAY_DESCRIPTOR desc; + desc.Format = vec_info::format; + desc.NumChannels = vec_info::size; + + const Sizes sizes(hipArrayDefault); + const size_t s = 64; + + hiparray array{}; + SECTION("Happy") { + SECTION("1D - Max") { + desc.Width = sizes.max1D; + desc.Height = 0; + } + SECTION("2D - Max Width") { + desc.Width = sizes.max2D[0]; + desc.Height = s; + } + SECTION("2D - Max Height") { + desc.Width = s; + desc.Height = sizes.max2D[1]; + } + SECTION("2D - Max Width and Height") { + desc.Width = sizes.max2D[0]; + desc.Height = sizes.max2D[1]; + } + auto maxArrayCreateError = hipArrayCreate(&array, &desc); + // this can try to alloc many GB of memory, so out of memory is acceptable + // return to avoid destroy + if (maxArrayCreateError == hipErrorOutOfMemory) return; + HIP_CHECK(maxArrayCreateError); + HIP_CHECK(hipArrayDestroy(array)); + } + SECTION("Negative") { + SECTION("1D - More Than Max") { + desc.Width = sizes.max1D + 1; + desc.Height = 0; + } + SECTION("2D - More Than Max Width") { + desc.Width = sizes.max2D[0] + 1; + desc.Height = s; + } + SECTION("2D - More Than Max Height") { + desc.Width = s; + desc.Height = sizes.max2D[1] + 1; + } + SECTION("2D - More Than Max Width and Height") { + desc.Width = sizes.max2D[0] + 1; + desc.Height = sizes.max2D[1] + 1; + } + HIP_CHECK_ERROR(hipArrayCreate(&array, &desc), hipErrorInvalidValue); + } +} + +// zero-width array is not supported +TEST_CASE("Unit_hipArrayCreate_ZeroWidth") { + DriverContext ctx; + HIP_ARRAY_DESCRIPTOR desc; + desc.Format = driverFormats[0]; + desc.NumChannels = 4; + desc.Width = 0; + desc.Height = GENERATE(0, 1024); + + // pointer to the array in device memory + hiparray array; + HIP_CHECK_ERROR(hipArrayCreate(&array, &desc), hipErrorInvalidValue); +} + +// HipArrayCreate will return an error when nullptr is used as the array argument +TEST_CASE("Unit_hipArrayCreate_Nullptr") { +#if HT_AMD + HipTest::HIP_SKIP_TEST("EXSWCPHIPT-130"); + return; +#endif + DriverContext ctx; + SECTION("Null array") { + HIP_ARRAY_DESCRIPTOR desc; + desc.Format = driverFormats[0]; + desc.NumChannels = 4; + desc.Width = 1024; + desc.Height = 1024; + + HIP_CHECK_ERROR(hipArrayCreate(nullptr, &desc), hipErrorInvalidValue); + } + SECTION("Null Description") { + hiparray array; + HIP_CHECK_ERROR(hipArrayCreate(&array, nullptr), hipErrorInvalidValue); + } +} + +// Only elements with 1,2, or 4 channels is supported +TEST_CASE("Unit_hipArrayCreate_BadNumberChannelElement") { + DriverContext ctx; + HIP_ARRAY_DESCRIPTOR desc; + desc.Format = GENERATE(from_range(std::begin(driverFormats), std::end(driverFormats))); + desc.NumChannels = GENERATE(-1, 0, 3, 5, 8); + desc.Width = 1024; + desc.Height = GENERATE(0, 1024); + + hiparray array; + + INFO("Format: " << formatToString(desc.Format) << " NumChannels: " << desc.NumChannels + << " Height: " << desc.Height) + HIP_CHECK_ERROR(hipArrayCreate(&array, &desc), hipErrorInvalidValue); +} + +// Only certain channel formats are acceptable. +TEST_CASE("Unit_hipArrayCreate_BadChannelFormat") { + DriverContext ctx; + HIP_ARRAY_DESCRIPTOR desc; + + // create a bad format + desc.Format = + std::accumulate(std::begin(driverFormats), std::end(driverFormats), driverFormats[0], + [](auto i, auto f) { return static_cast(i + f); }); + for (auto&& format : driverFormats) { + REQUIRE(desc.Format != format); + } + + desc.NumChannels = 4; + desc.Width = 1024; + desc.Height = GENERATE(0, 1024); + + hiparray array; + + INFO("Format: " << formatToString(desc.Format) << " Height: " << desc.Height) + HIP_CHECK_ERROR(hipArrayCreate(&array, &desc), hipErrorInvalidValue); +} diff --git a/tests/catch/unit/memory/hipHostGetFlags.cc b/tests/catch/unit/memory/hipHostGetFlags.cc index f150aaa5a8..25fb5b2734 100644 --- a/tests/catch/unit/memory/hipHostGetFlags.cc +++ b/tests/catch/unit/memory/hipHostGetFlags.cc @@ -1,5 +1,5 @@ /* -Copyright (c) 2021 Advanced Micro Devices, Inc. All rights reserved. +Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved. Permission is hereby granted, free of charge, to any person obtaining a copy of this software and associated documentation files (the "Software"), to deal @@ -27,8 +27,34 @@ This testcase verifies the basic scenario of hipHostGetFlags API #include #include #include +#include +#include -static constexpr auto LEN{1024*1024}; +std::vector FlagPart1Vec{hipHostMallocDefault, + hipHostMallocDefault | hipHostMallocPortable, + hipHostMallocDefault | hipHostMallocMapped, + hipHostMallocDefault | hipHostMallocWriteCombined, + hipHostMallocPortable, + hipHostMallocPortable | hipHostMallocMapped, + hipHostMallocPortable | hipHostMallocWriteCombined, + hipHostMallocMapped, + hipHostMallocMapped | hipHostMallocWriteCombined, + hipHostMallocWriteCombined}; +#if HT_AMD +// For cases where flags from FlagPart1Vec are not used, +// hipHostMallocDefault is the default on AMD +// and hipHostMallocMapped on Nvidia +std::vector FlagPart2Vec{0x0, + hipHostMallocNumaUser, + hipHostMallocNumaUser | hipHostMallocCoherent, + hipHostMallocNumaUser | hipHostMallocNonCoherent, + hipHostMallocCoherent, + hipHostMallocNonCoherent}; +#else +std::vector FlagPart2Vec{0x0}; +#endif + +static constexpr auto LEN{1024 * 1024}; /* This testcase verifies hipHostGetFlags API basic scenario @@ -38,57 +64,160 @@ This testcase verifies hipHostGetFlags API basic scenario 3. Validates it with the initial flags used while allocating memory */ -TEMPLATE_TEST_CASE("Unit_hipHostGetFlags_Basic", "", int, - float, double) { - constexpr auto SIZE{LEN * sizeof(TestType)}; - TestType *A_h{nullptr}, *B_h{nullptr}, *C_h{nullptr}; - TestType *A_d{nullptr}, *B_d{nullptr}, *C_d{nullptr}; - unsigned int FlagA, FlagB, FlagC; - FlagA = hipHostMallocWriteCombined | hipHostMallocMapped; - FlagB = hipHostMallocWriteCombined | hipHostMallocMapped; - FlagC = hipHostMallocMapped; +/* Possible host flags + * hipHostMallocDefault 0x0 + * hipHostMallocPortable 0x1 + * hipHostMallocMapped 0x2 + * hipHostMallocWriteCombined 0x4 + * NOT on Nvidia + * hipHostMallocNumaUser 0x20000000 + * hipHostMallocCoherent 0x40000000 + * hipHostMallocNonCoherent 0x80000000 + */ + +inline void checkFlags(unsigned int expected, unsigned int obtained) { + // Account for cases where flags from FlagPart1Vec do not include hipHostMallocMapped, + // on Nvidia devices it is added by default +#if HT_NVIDIA + expected = expected | hipHostMallocMapped; +#endif + REQUIRE(expected == obtained); +} + +TEST_CASE("Unit_hipHostGetFlags_flagCombos") { + + constexpr auto SIZE{LEN * sizeof(int)}; + int* A_h{nullptr}; + + const unsigned int FlagPart1 = GENERATE(from_range(FlagPart1Vec.begin(), FlagPart1Vec.end())); + const unsigned int FlagPart2 = GENERATE(from_range(FlagPart2Vec.begin(), FlagPart2Vec.end())); + + unsigned int FlagComp = FlagPart1 | FlagPart2; hipDeviceProp_t prop; - int device; + int device{}; + HIP_CHECK(hipGetDevice(&device)); + HIP_CHECK(hipGetDeviceProperties(&prop, device)); + + // Skip test if device does not support the property canMapHostMemory + if (prop.canMapHostMemory != 1) { + HipTest::HIP_SKIP_TEST("Device Property canMapHostMemory is not set"); + return; + } else { + // Allocate using the generated flags combos + INFO("Flag passed when allocating: 0x" << std::hex << FlagComp << "\n"); + HIP_CHECK(hipHostMalloc(reinterpret_cast(&A_h), SIZE, FlagComp)); + unsigned int flagA{}; + + // get the flags from allocations and check if they are the same as the one set + HIP_CHECK(hipHostGetFlags(&flagA, A_h)); + + checkFlags(FlagComp, flagA); + HIP_CHECK(hipHostFree(A_h)); + } +} + +// Test Allocation with flags and getting flags in another thread +TEST_CASE("Unit_hipHostGetFlags_DifferentThreads") { + constexpr auto SIZE{LEN * sizeof(int)}; + int* A_h{nullptr}; + + const unsigned int FlagPart1 = GENERATE(from_range(FlagPart1Vec.begin(), FlagPart1Vec.end())); + const unsigned int FlagPart2 = GENERATE(from_range(FlagPart2Vec.begin(), FlagPart2Vec.end())); + + + unsigned int FlagComp = FlagPart1 | FlagPart2; + + hipDeviceProp_t prop; + int device{}; HIP_CHECK(hipGetDevice(&device)); HIP_CHECK(hipGetDeviceProperties(&prop, device)); if (prop.canMapHostMemory != 1) { - SUCCEED("Device Property canMapHostMemory is not set"); + HipTest::HIP_SKIP_TEST("Device Property canMapHostMemory is not set"); + return; } else { - HIP_CHECK(hipHostMalloc(reinterpret_cast(&A_h), SIZE, - hipHostMallocWriteCombined | hipHostMallocMapped)); - HIP_CHECK(hipHostMalloc(reinterpret_cast(&B_h), SIZE, - hipHostMallocWriteCombined | hipHostMallocMapped)); - HIP_CHECK(hipHostMalloc(reinterpret_cast(&C_h), SIZE, - hipHostMallocMapped)); - - unsigned int flagA, flagB, flagC; - HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast(&A_d), A_h, 0)); - HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast(&B_d), B_h, 0)); - HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast(&C_d), C_h, 0)); + // Make sure we allocate before trying to get the flags + std::thread malloc_thread( + [&]() { HIP_CHECK_THREAD(hipHostMalloc(reinterpret_cast(&A_h), SIZE, FlagComp)); }); + malloc_thread.join(); + HIP_CHECK_THREAD_FINALIZE(); + unsigned int flagA{}; HIP_CHECK(hipHostGetFlags(&flagA, A_h)); - HIP_CHECK(hipHostGetFlags(&flagB, B_h)); - HIP_CHECK(hipHostGetFlags(&flagC, C_h)); - HipTest::setDefaultData(LEN, A_h, B_h, C_h); + checkFlags(FlagComp, flagA); - dim3 dimGrid(LEN / 512, 1, 1); - dim3 dimBlock(512, 1, 1); - hipLaunchKernelGGL(HipTest::vectorADD, dimGrid, dimBlock, - 0, 0, static_cast(A_d), - static_cast(B_d), C_d, LEN); - - HIP_CHECK(hipMemcpy(C_h, C_d, SIZE, hipMemcpyDeviceToHost)); - // Note this really HostToHost not - // DeviceToHost, since memory is mapped... - HipTest::checkVectorADD(A_h, B_h, C_h, LEN); - - REQUIRE(flagA == FlagA); - REQUIRE(flagB == FlagB); - REQUIRE(flagC == FlagC); HIP_CHECK(hipHostFree(A_h)); - HIP_CHECK(hipHostFree(B_h)); - HIP_CHECK(hipHostFree(C_h)); + } +} + +// Test behaviour of hipHostGetFlags with invalid args +TEST_CASE("Unit_hipHostGetFlags_InvalidArgs") { + constexpr auto SIZE{LEN * sizeof(int)}; + int* A_h{nullptr}; + + hipDeviceProp_t prop; + int device{}; + HIP_CHECK(hipGetDevice(&device)); + HIP_CHECK(hipGetDeviceProperties(&prop, device)); + + // Skip test if device does not support the property canMapHostMemory + if (prop.canMapHostMemory != 1) { + HipTest::HIP_SKIP_TEST("Device Property canMapHostMemory is not set"); + return; + } else { + SECTION("Invalid flag ptr being passed to hipHostGetFlags") { + // Use default flag + unsigned int FlagComp = 0x0; + + // Allocate using the generated flags combos + HIP_CHECK(hipHostMalloc(reinterpret_cast(&A_h), SIZE, FlagComp)); + + // use a nullptr to return flags to + unsigned int* flagA = nullptr; + + // get the flags from allocations and check if they are the same as the one set + HIP_CHECK_ERROR(hipHostGetFlags(flagA, A_h), hipErrorInvalidValue); + + HIP_CHECK(hipHostFree(A_h)); + } + + SECTION("Device ptr allocated with hipMalloc passed to hipHostGetFlags") { + unsigned int FlagComp = 0x4; + + // Allocate memory on device + HIP_CHECK(hipMalloc(reinterpret_cast(&A_h), SIZE)); + + unsigned int flagA{}; + + // get the flags from allocations and check if they are the same as the one set + HIP_CHECK_ERROR(hipHostGetFlags(&flagA, A_h), hipErrorInvalidValue); + INFO("Flag passed when allocating: " << std::hex << FlagComp << " Returned flag: " << std::hex + << flagA << "\n"); + + HIP_CHECK(hipFree(A_h)); + } + + SECTION("Ptr from hipHostGetDevicePointer passed to hipHostGetFlags") { + unsigned int FlagComp = 0x4; + + int* A_d{nullptr}; + // Allocate memory on device + HIP_CHECK(hipHostMalloc(reinterpret_cast(&A_h), SIZE, FlagComp)); + HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast(&A_d), A_h, 0)); + + unsigned int flagA; + + // get the flags from allocations and check if they are the same as the one set + HIP_CHECK(hipHostGetFlags(&flagA, A_d)); + INFO("Flag passed when allocating: " << std::hex << FlagComp << " Returned flag: " << std::hex + << flagA << "\n"); +#if HT_NVIDIA + // on Nvidia adjust for cudaHostAllocMapped being set by default + FlagComp = FlagComp | hipHostMallocMapped; +#endif + REQUIRE(flagA == FlagComp); + HIP_CHECK(hipHostFree(A_h)); + } } } diff --git a/tests/catch/unit/memory/hipHostRegister.cc b/tests/catch/unit/memory/hipHostRegister.cc index 8312cf7f50..417ab0857a 100644 --- a/tests/catch/unit/memory/hipHostRegister.cc +++ b/tests/catch/unit/memory/hipHostRegister.cc @@ -1,5 +1,5 @@ /* -Copyright (c) 2021 Advanced Micro Devices, Inc. All rights reserved. +Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved. Permission is hereby granted, free of charge, to any person obtaining a copy of this software and associated documentation files (the "Software"), to deal @@ -28,51 +28,49 @@ This testfile verifies the following scenarios of hipHostRegister API */ #include -#include -#include +#include #define OFFSET 128 -static constexpr auto LEN{1024*1024}; +static constexpr auto LEN{1024 * 1024}; -template -__global__ void Inc(T* Ad) { - int tx = threadIdx.x + blockIdx.x * blockDim.x; - Ad[tx] = Ad[tx] + static_cast(1); +template __global__ void Inc(T* Ad) { + int tx = threadIdx.x + blockIdx.x * blockDim.x; + Ad[tx] = Ad[tx] + static_cast(1); } template -void doMemCopy(size_t numElements, int offset, T* A, T* Bh, T* Bd, - bool internalRegister) { - constexpr auto memsetval = 13.0f; - A = A + offset; - numElements -= offset; +void doMemCopy(size_t numElements, int offset, T* A, T* Bh, T* Bd, bool internalRegister) { + constexpr auto memsetval = 13.0f; + A = A + offset; + numElements -= offset; - size_t sizeBytes = numElements * sizeof(T); + size_t sizeBytes = numElements * sizeof(T); - if (internalRegister) { - HIP_CHECK(hipHostRegister(A, sizeBytes, 0)); - } + if (internalRegister) { + HIP_CHECK(hipHostRegister(A, sizeBytes, 0)); + } - // Reset - for (size_t i = 0; i < numElements; i++) { - A[i] = static_cast(i); - Bh[i] = 0.0f; - } + // Reset + for (size_t i = 0; i < numElements; i++) { + A[i] = static_cast(i); + Bh[i] = 0.0f; + } - HIP_CHECK(hipMemset(Bd, memsetval, sizeBytes)); + HIP_CHECK(hipMemset(Bd, memsetval, sizeBytes)); - HIP_CHECK(hipMemcpy(Bd, A, sizeBytes, hipMemcpyHostToDevice)); - HIP_CHECK(hipMemcpy(Bh, Bd, sizeBytes, hipMemcpyDeviceToHost)); + HIP_CHECK(hipMemcpy(Bd, A, sizeBytes, hipMemcpyHostToDevice)); + HIP_CHECK(hipMemcpy(Bh, Bd, sizeBytes, hipMemcpyDeviceToHost)); - // Make sure the copy worked - for (size_t i = 0; i < numElements; i++) { - REQUIRE(Bh[i] == A[i]); - } + // Make sure the copy worked + for (size_t i = 0; i < numElements; i++) { + REQUIRE(Bh[i] == A[i]); + } - if (internalRegister) { - HIP_CHECK(hipHostUnregister(A)); - } + if (internalRegister) { + HIP_CHECK(hipHostUnregister(A)); + } } + /* This testcase verifies the hipHostRegister API by 1. Allocating the memory using malloc @@ -81,9 +79,7 @@ This testcase verifies the hipHostRegister API by 4. Launching kernel and access the device pointer variable 5. performing hipMemset on the device pointer variable */ -TEMPLATE_TEST_CASE("Unit_hipHostRegister_ReferenceFromKernelandhipMemset", - "", int, - float, double) { +TEMPLATE_TEST_CASE("Unit_hipHostRegister_ReferenceFromKernelandhipMemset", "", int, float, double) { size_t sizeBytes{LEN * sizeof(TestType)}; TestType *A, **Ad; int num_devices; @@ -118,14 +114,14 @@ TEMPLATE_TEST_CASE("Unit_hipHostRegister_ReferenceFromKernelandhipMemset", HIP_CHECK(hipHostUnregister(A)); free(A); - delete [] Ad; + delete[] Ad; } + /* This testcase verifies hipHostRegister API by performing memcpy on the hipHostRegistered variable. */ -TEMPLATE_TEST_CASE("Unit_hipHostRegister_Memcpy", "", - int, float, double) { +TEMPLATE_TEST_CASE("Unit_hipHostRegister_Memcpy", "", int, float, double) { // 1 refers to hipHostRegister // 0 refers to malloc auto mem_type = GENERATE(0, 1); @@ -156,5 +152,49 @@ TEMPLATE_TEST_CASE("Unit_hipHostRegister_Memcpy", "", free(A); free(Bh); - hipFree(Bd); + HIP_CHECK(hipFree(Bd)); +} + +template __global__ void fill_kernel(T* dataPtr, T value) { + size_t tid{blockIdx.x * blockDim.x + threadIdx.x}; + dataPtr[tid] = value; +} + +TEMPLATE_TEST_CASE("Unit_hipHostRegister_Negative", "", int, float, double) { + TestType* hostPtr = nullptr; + + size_t sizeBytes = 1 * sizeof(TestType); + SECTION("hipHostRegister Negative Test - nullptr") { + HIP_CHECK_ERROR(hipHostRegister(hostPtr, 1, 0), hipErrorInvalidValue); + } + + hostPtr = reinterpret_cast(malloc(sizeBytes)); + SECTION("hipHostRegister Negative Test - zero size") { + HIP_CHECK_ERROR(hipHostRegister(hostPtr, 0, 0), hipErrorInvalidValue); + } + +#if HT_NVIDIA + // Flags aren't used for AMD devices currently + SECTION("hipHostRegister Negative Test - invalid flag") { + HIP_CHECK_ERROR(hipHostRegister(hostPtr, sizeBytes, 0b11111111), hipErrorInvalidValue); + } +#endif + + size_t devMemAvail{0}, devMemFree{0}; + HIP_CHECK(hipMemGetInfo(&devMemFree, &devMemAvail)); + auto hostMemFree = HipTest::getMemoryAmount() /* In MB */ * 1024 * 1024; // In bytes + REQUIRE(devMemFree > 0); + REQUIRE(devMemAvail > 0); + REQUIRE(hostMemFree > 0); + + size_t memFree = (std::min)(devMemFree, hostMemFree); // which is the limiter cpu or gpu + + SECTION("hipHostRegister Negative Test - invalid memory size") { + HIP_CHECK_ERROR(hipHostRegister(hostPtr, memFree, 0), hipErrorInvalidValue); + } + + free(hostPtr); + SECTION("hipHostRegister Negative Test - freed memory") { + HIP_CHECK_ERROR(hipHostRegister(hostPtr, 0, 0), hipErrorInvalidValue); + } } diff --git a/tests/catch/unit/memory/hipHostUnregister.cc b/tests/catch/unit/memory/hipHostUnregister.cc new file mode 100644 index 0000000000..68a34122b9 --- /dev/null +++ b/tests/catch/unit/memory/hipHostUnregister.cc @@ -0,0 +1,103 @@ +/* +Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved. + +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#include + +namespace hipHostUnregisterTests { +constexpr unsigned int allFlags = hipHostRegisterDefault & // 0 + hipHostRegisterPortable & // 1 + hipHostRegisterMapped & // 2 + hipHostRegisterIoMemory // 4 +#if HT_NVIDIA + & cudaHostRegisterReadOnly; // 8 +#else + ; +#endif + +inline bool hipHostRegisterSupported() { +#if HT_NVIDIA + // unable to query for cudaDevAttrHostRegisterSupported equivalent + HipTest::HIP_SKIP_TEST("EXSWCPHIPT-40"); + HipTest::HIP_SKIP_TEST("hipHostRegister is not supported on this device"); + return false; +#else + return true; +#endif +} + + +TEST_CASE("Unit_hipHostUnregister_MemoryNotAccessableAfterUnregister") { + if (!hipHostRegisterSupported()) { + return; + } + // try all combinations of flags + for (unsigned int flag = 0; flag <= allFlags; ++flag) { + DYNAMIC_SECTION("Using flag: " << flag) { + auto x = std::unique_ptr(new int); + HIP_CHECK(hipHostRegister(x.get(), sizeof(int), flag)); + + void* device_memory; + HIP_CHECK(hipHostGetDevicePointer(&device_memory, x.get(), 0)); + + HIP_CHECK(hipHostUnregister(x.get())); + HIP_CHECK_ERROR(hipHostGetDevicePointer(&device_memory, x.get(), 0), hipErrorInvalidValue); + } + } +} + +TEST_CASE("Unit_hipHostUnregister_NullPtr") { +#if HT_AMD + HipTest::HIP_SKIP_TEST("EXSWCPHIPT-90"); + return; +#endif + HIP_CHECK_ERROR(hipHostUnregister(nullptr), hipErrorInvalidValue); +} + +TEST_CASE("Unit_hipHostUnregister_NotRegisteredPointer") { +#if HT_AMD + HipTest::HIP_SKIP_TEST("TODO-MATCH-ERRORCODE"); + return; +#endif + auto x = std::unique_ptr(new int); + HIP_CHECK_ERROR(hipHostUnregister(x.get()), hipErrorHostMemoryNotRegistered); +} + +TEST_CASE("Unit_hipHostUnregister_AlreadyUnregisteredPointer") { +#if HT_AMD + HipTest::HIP_SKIP_TEST("TODO-MATCH-ERRORCODE"); + return; +#endif + if (!hipHostRegisterSupported()) { + return; + } + // try all combinations of flags + for (unsigned int flag = 0; flag <= allFlags; ++flag) { + DYNAMIC_SECTION("Using flag: " << flag) { + auto x = std::unique_ptr(new int); + HIP_CHECK(hipHostRegister(x.get(), sizeof(int), flag)); + HIP_CHECK(hipHostUnregister(x.get())); + HIP_CHECK_ERROR(hipHostUnregister(x.get()), hipErrorHostMemoryNotRegistered); + } + } +} + +} // namespace hipHostUnregisterTests diff --git a/tests/catch/unit/memory/hipMalloc3D.cc b/tests/catch/unit/memory/hipMalloc3D.cc index 2efc9ff7db..386eff0a51 100644 --- a/tests/catch/unit/memory/hipMalloc3D.cc +++ b/tests/catch/unit/memory/hipMalloc3D.cc @@ -64,26 +64,6 @@ static void Malloc3DThreadFunc(int gpu) { MemoryAlloc3DDiffSizes(gpu); } -/* - * This verifies the negative scenarios of hipMalloc3D API - */ -TEST_CASE("Unit_hipMalloc3D_Negative") { - size_t width = SMALL_SIZE * sizeof(char); - size_t height{SMALL_SIZE}, depth{SMALL_SIZE}; - hipPitchedPtr devPitchedPtr; - - SECTION("Passing nullptr to device pitched pointer") { - hipExtent extent = make_hipExtent(width, height, depth); - REQUIRE(hipMalloc3D(nullptr, extent) != hipSuccess); - } - - SECTION("Passing Max values to extent") { - hipExtent extent = make_hipExtent(std::numeric_limits::max(), - std::numeric_limits::max(), - std::numeric_limits::max()); - REQUIRE(hipMalloc3D(&devPitchedPtr, extent) != hipSuccess); - } -} /* * This verifies the hipMalloc3D API by * assigning width,height and depth as 10 diff --git a/tests/catch/unit/memory/hipMalloc3DArray.cc b/tests/catch/unit/memory/hipMalloc3DArray.cc index fe8ba3f8b0..1a8ec2f0d0 100644 --- a/tests/catch/unit/memory/hipMalloc3DArray.cc +++ b/tests/catch/unit/memory/hipMalloc3DArray.cc @@ -25,14 +25,15 @@ hipMalloc3DArray API test scenarios 4. Multithreaded scenario */ - - +#include #include +#include "hipArrayCommon.hh" static constexpr auto ARRAY_SIZE{4}; static constexpr auto BIG_ARRAY_SIZE{100}; static constexpr auto ARRAY_LOOP{100}; + /* * This API verifies memory allocations for small and * bigger chunks of data. @@ -47,126 +48,31 @@ static constexpr auto ARRAY_LOOP{100}; * */ static void Malloc3DArray_DiffSizes(int gpu) { - HIP_CHECK(hipSetDevice(gpu)); - std::vector array_size; - array_size.push_back(ARRAY_SIZE); - array_size.push_back(BIG_ARRAY_SIZE); - for (auto &size : array_size) { - int width{size}, height{size}, depth{size}; - hipChannelFormatDesc channelDesc = hipCreateChannelDesc(sizeof(float)*8, 0, - 0, 0, hipChannelFormatKindFloat); - hipArray *arr[ARRAY_LOOP]; - size_t tot, avail, ptot, pavail; - HIP_CHECK(hipMemGetInfo(&pavail, &ptot)); - for (int i = 0; i < ARRAY_LOOP; i++) { - HIP_CHECK(hipMalloc3DArray(&arr[i], &channelDesc, make_hipExtent(width, - height, depth), hipArrayDefault)); - } - for (int i = 0; i < ARRAY_LOOP; i++) { - HIP_CHECK(hipFreeArray(arr[i])); - } - HIP_CHECK(hipMemGetInfo(&avail, &tot)); - if ((pavail != avail)) { - HIPASSERT(false); - } + HIP_CHECK_THREAD(hipSetDevice(gpu)); + const int size = GENERATE(ARRAY_SIZE, BIG_ARRAY_SIZE); + int width{size}, height{size}, depth{size}; + hipChannelFormatDesc channelDesc = hipCreateChannelDesc(); + std::array arr; + size_t pavail, avail; + HIP_CHECK_THREAD(hipMemGetInfo(&pavail, nullptr)); + + for (int i = 0; i < ARRAY_LOOP; i++) { + HIP_CHECK_THREAD(hipMalloc3DArray(&arr[i], &channelDesc, make_hipExtent(width, height, depth), + hipArrayDefault)); } -} - -/* Thread Function */ -static void Malloc3DArrayThreadFunc(int gpu) { - Malloc3DArray_DiffSizes(gpu); -} - -/* - * Verifies the negative scenarios of hipMalloc3DArray API - */ -TEST_CASE("Unit_hipMalloc3DArray_Negative") { - constexpr int width{ARRAY_SIZE}, height{ARRAY_SIZE}, depth{ARRAY_SIZE}; - hipChannelFormatDesc channelDesc = hipCreateChannelDesc(sizeof(float)*8, 0, - 0, 0, hipChannelFormatKindFloat); - hipArray *arr; -#if HT_NVIDIA - SECTION("NullPointer to Array") { - REQUIRE(hipMalloc3DArray(nullptr, &channelDesc, make_hipExtent(width, - height, depth), hipArrayDefault) != hipSuccess); + for (int i = 0; i < ARRAY_LOOP; i++) { + HIP_CHECK_THREAD(hipFreeArray(arr[i])); } - SECTION("NullPointer to Channel Descriptor") { - REQUIRE(hipMalloc3DArray(&arr, nullptr, make_hipExtent(width, - height, depth), hipArrayDefault) != hipSuccess); - } -#endif - SECTION("Width 0 in hipExtent") { - REQUIRE(hipMalloc3DArray(&arr, &channelDesc, make_hipExtent(0, - height, width), hipArrayDefault) != hipSuccess); - } - - SECTION("Height 0 in hipExtent") { - REQUIRE(hipMalloc3DArray(&arr, &channelDesc, make_hipExtent(width, - 0, width), hipArrayDefault) != hipSuccess); - } - - SECTION("Invalid Flag") { - REQUIRE(hipMalloc3DArray(&arr, &channelDesc, make_hipExtent(width, - height, depth), 100) != hipSuccess); - } - - SECTION("Width,Height & Depth 0 in hipExtent") { - REQUIRE(hipMalloc3DArray(&arr, &channelDesc, make_hipExtent(0, - 0, 0), hipArrayDefault) != hipSuccess); - } - - SECTION("Max int values to extent") { - REQUIRE(hipMalloc3DArray(&arr, &channelDesc, - make_hipExtent(std::numeric_limits::max(), - std::numeric_limits::max(), - std::numeric_limits::max()), - hipArrayDefault) != hipSuccess); - } -} -/* - * Verifies the extent validation scenarios - * 1. Passing depth as 0 would create 2D array - * 2. Passing height and depth as 0 would create 1D array - * from hipMalloc3DArray API - */ -TEST_CASE("Unit_hipMalloc3DArray_ExtentValidation") { - constexpr int width{ARRAY_SIZE}, height{ARRAY_SIZE}; - hipChannelFormatDesc channelDesc = hipCreateChannelDesc(sizeof(float)*8, 0, - 0, 0, hipChannelFormatKindFloat); - hipArray *arr; - - SECTION("Depth 0 in hipExtent") { - REQUIRE(hipMalloc3DArray(&arr, &channelDesc, make_hipExtent(width, - height, 0), hipArrayDefault) == hipSuccess); - HIP_CHECK(hipFreeArray(arr)); - } - - SECTION("Height & Depth 0 in hipExtent") { - REQUIRE(hipMalloc3DArray(&arr, &channelDesc, make_hipExtent(width, - 0, 0), hipArrayDefault) == hipSuccess); - HIP_CHECK(hipFreeArray(arr)); - } -} - -/* - * Verifies hipMalloc3DArray API by passing width,height - * and depth as 10 - */ -TEST_CASE("Unit_hipMalloc3DArray_Basic") { - constexpr int width{ARRAY_SIZE}, height{ARRAY_SIZE}, depth{ARRAY_SIZE}; - hipChannelFormatDesc channelDesc = hipCreateChannelDesc(sizeof(float)*8, 0, - 0, 0, hipChannelFormatKindFloat); - hipArray *arr; - - REQUIRE(hipMalloc3DArray(&arr, &channelDesc, make_hipExtent(width, - height, depth), hipArrayDefault) == hipSuccess); - HIP_CHECK(hipFreeArray(arr)); + HIP_CHECK_THREAD(hipMemGetInfo(&avail, nullptr)); + REQUIRE_THREAD(pavail == avail); } TEST_CASE("Unit_hipMalloc3DArray_DiffSizes") { Malloc3DArray_DiffSizes(0); + HIP_CHECK_THREAD_FINALIZE(); } + /* This testcase verifies the hipMalloc3DArray API in multithreaded scenario by launching threads in parallel on multiple GPUs @@ -176,16 +82,16 @@ TEST_CASE("Unit_hipMalloc3DArray_MultiThread") { std::vector threadlist; int devCnt = 0; devCnt = HipTest::getDeviceCount(); - size_t tot, avail, ptot, pavail; - HIP_CHECK(hipMemGetInfo(&pavail, &ptot)); + const auto pavail = getFreeMem(); for (int i = 0; i < devCnt; i++) { - threadlist.push_back(std::thread(Malloc3DArrayThreadFunc, i)); + threadlist.push_back(std::thread(Malloc3DArray_DiffSizes, i)); } - for (auto &t : threadlist) { + for (auto& t : threadlist) { t.join(); } - HIP_CHECK(hipMemGetInfo(&avail, &tot)); + HIP_CHECK_THREAD_FINALIZE(); + const auto avail = getFreeMem(); if (pavail != avail) { WARN("Memory leak of hipMalloc3D API in multithreaded scenario"); @@ -193,3 +99,324 @@ TEST_CASE("Unit_hipMalloc3DArray_MultiThread") { } } +namespace { +void checkArrayIsExpected(hipArray_t array, const hipChannelFormatDesc& expected_desc, + const hipExtent& expected_extent, const unsigned int expected_flags) { +// hipArrayGetInfo doesn't currently exist (EXSWCPHIPT-87) +#if HT_AMD + std::ignore = array; + std::ignore = expected_desc; + std::ignore = expected_extent; + std::ignore = expected_flags; +#else + cudaChannelFormatDesc queried_desc; + cudaExtent queried_extent; + unsigned int queried_flags; + + cudaArrayGetInfo(&queried_desc, &queried_extent, &queried_flags, array); + + REQUIRE(expected_desc.x == queried_desc.x); + REQUIRE(expected_desc.y == queried_desc.y); + REQUIRE(expected_desc.z == queried_desc.z); + REQUIRE(expected_desc.f == queried_desc.f); + + REQUIRE(expected_extent.width == queried_extent.width); + REQUIRE(expected_extent.height == queried_extent.height); + REQUIRE(expected_extent.depth == queried_extent.depth); + + REQUIRE(expected_flags == queried_flags); +#endif +} +} // namespace + +TEMPLATE_TEST_CASE("Unit_hipMalloc3DArray_happy", "", char, uchar2, uint2, int4, short4, float, + float2, float4) { + hipArray_t array; + const auto desc = hipCreateChannelDesc(); +#if HT_AMD + const unsigned int flags = hipArrayDefault; +#else + const unsigned int flags = GENERATE(hipArrayDefault, hipArraySurfaceLoadStore); +#endif + constexpr size_t size = 64; + hipExtent extent; + + SECTION("1D Array") { extent = make_hipExtent(size, 0, 0); } + SECTION("2D Array") { extent = make_hipExtent(size, size, 0); } + SECTION("3D Array") { extent = make_hipExtent(size, size, size); } + + HIP_CHECK(hipMalloc3DArray(&array, &desc, extent, flags)); + checkArrayIsExpected(array, desc, extent, flags); + HIP_CHECK(hipFreeArray(array)); +} + +TEMPLATE_TEST_CASE("Unit_hipMalloc3DArray_MaxTexture", "", int, uint4, short, ushort2, + unsigned char, float, float4) { +#if HT_AMD + HipTest::HIP_SKIP_TEST("EXSWCPHIPT-97"); + return; +#endif + + hipArray_t array; + const hipChannelFormatDesc desc = hipCreateChannelDesc(); +#if HT_AMD + const unsigned int flag = hipArrayDefault; +#else + const unsigned int flag = GENERATE(hipArrayDefault, hipArraySurfaceLoadStore); +#endif + if (flag == hipArraySurfaceLoadStore) { + HipTest::HIP_SKIP_TEST("EXSWCPHIPT-58"); + return; + } + CAPTURE(flag); + const Sizes sizes(flag); + CAPTURE(sizes.max1D, sizes.max2D, sizes.max3D); + + const size_t s = 64; + SECTION("Happy") { + // stored in a vector so some values can be ifdef'd out + std::vector extentsToTest{ + make_hipExtent(sizes.max1D, 0, 0), // 1D max + make_hipExtent(sizes.max2D[0], s, 0), // 2D max width + make_hipExtent(s, sizes.max2D[1], 0), // 2D max height + make_hipExtent(sizes.max2D[0], sizes.max2D[1], 0), // 2D max + make_hipExtent(sizes.max3D[0], s, s), // 3D max width + make_hipExtent(s, sizes.max3D[1], s), // 3D max height + make_hipExtent(s, s, sizes.max3D[2]), // 3D max depth + make_hipExtent(s, sizes.max3D[1], sizes.max3D[2]), // 3D max height and depth + make_hipExtent(sizes.max3D[0], s, sizes.max3D[2]), // 3D max width and depth + make_hipExtent(sizes.max3D[0], sizes.max3D[1], s), // 3D max width and height + make_hipExtent(sizes.max3D[0], sizes.max3D[1], sizes.max3D[2]) // 3D max + }; + const auto extent = + GENERATE_COPY(from_range(std::begin(extentsToTest), std::end(extentsToTest))); + CAPTURE(extent.width, extent.height, extent.depth); + auto maxArrayCreateError = hipMalloc3DArray(&array, &desc, extent, flag); + // this can try to alloc many GB of memory, so out of memory is acceptable + if (maxArrayCreateError == hipErrorOutOfMemory) return; + HIP_CHECK(maxArrayCreateError); + checkArrayIsExpected(array, desc, extent, flag); + HIP_CHECK(hipFreeArray(array)); + } + SECTION("Negative") { + std::vector extentsToTest { + make_hipExtent(sizes.max1D + 1, 0, 0), // 1D max + make_hipExtent(sizes.max2D[0] + 1, s, 0), // 2D max width + make_hipExtent(s, sizes.max2D[1] + 1, 0), // 2D max height + make_hipExtent(sizes.max2D[0] + 1, sizes.max2D[1] + 1, 0), // 2D max + make_hipExtent(sizes.max3D[0] + 1, s, s), // 3D max width + make_hipExtent(s, sizes.max3D[1] + 1, s), // 3D max height +#if !HT_NVIDIA // leads to hipSuccess on NVIDIA + make_hipExtent(s, s, sizes.max3D[2] + 1), // 3D max depth +#endif + make_hipExtent(s, sizes.max3D[1] + 1, sizes.max3D[2] + 1), // 3D max height and depth + make_hipExtent(sizes.max3D[0] + 1, s, sizes.max3D[2] + 1), // 3D max width and depth + make_hipExtent(sizes.max3D[0] + 1, sizes.max3D[1] + 1, s), // 3D max width and height + make_hipExtent(sizes.max3D[0] + 1, sizes.max3D[1] + 1, sizes.max3D[2] + 1) // 3D max + }; + const auto extent = + GENERATE_COPY(from_range(std::begin(extentsToTest), std::end(extentsToTest))); + CAPTURE(extent.width, extent.height, extent.depth); + HIP_CHECK_ERROR(hipMalloc3DArray(&array, &desc, extent, flag), hipErrorInvalidValue); + } +} + + +#if HT_AMD +constexpr std::array validFlags{hipArrayDefault}; +#else +constexpr std::array validFlags{ + hipArrayDefault, + hipArrayDefault | hipArraySurfaceLoadStore, + hipArrayLayered, + hipArrayLayered | hipArraySurfaceLoadStore, + hipArrayCubemap, + hipArrayCubemap | hipArrayLayered, + hipArrayCubemap | hipArraySurfaceLoadStore, + hipArrayCubemap | hipArrayLayered | hipArraySurfaceLoadStore, + hipArrayTextureGather}; +#endif + +hipExtent makeExtent(unsigned int flag, size_t s) { + if (flag == hipArrayTextureGather) { + return make_hipExtent(s, s, 0); + } + return make_hipExtent(s, s, s); +} + + +// Providing the array pointer as nullptr should return an error +TEST_CASE("Unit_hipMalloc3DArray_Negative_NullArrayPtr") { + hipChannelFormatDesc desc = hipCreateChannelDesc(); + constexpr size_t s = 6; + + const auto flag = GENERATE(from_range(std::begin(validFlags), std::end(validFlags))); + HIP_CHECK_ERROR(hipMalloc3DArray(nullptr, &desc, makeExtent(flag, s), flag), + hipErrorInvalidValue); +} + +// Providing the description pointer as nullptr should return an error +TEST_CASE("Unit_hipMalloc3DArray_Negative_NullDescPtr") { + constexpr size_t s = 6; // 6 to keep cubemap happy + hipArray_t array; + + const auto flag = GENERATE(from_range(std::begin(validFlags), std::end(validFlags))); + + HIP_CHECK_ERROR(hipMalloc3DArray(&array, nullptr, makeExtent(flag, s), flag), + hipErrorInvalidValue); +} + +// Zero width arrays are not allowed +TEST_CASE("Unit_hipMalloc3DArray_Negative_ZeroWidth") { + constexpr size_t s = 6; // 6 to keep cubemap happy + hipArray_t array; + hipChannelFormatDesc desc = hipCreateChannelDesc(); + + const auto flag = GENERATE(from_range(std::begin(validFlags), std::end(validFlags))); + + HIP_CHECK_ERROR(hipMalloc3DArray(&array, &desc, make_hipExtent(0, s, s), flag), + hipErrorInvalidValue); +} + +// Zero height arrays are only allowed for 1D arrays and layered arrays +TEST_CASE("Unit_hipMalloc3DArray_Negative_ZeroHeight") { + constexpr size_t s = 6; // 6 to keep cubemap happy + hipArray_t array; + hipChannelFormatDesc desc = hipCreateChannelDesc(); + std::array exceptions{hipArrayLayered, + hipArrayLayered | hipArraySurfaceLoadStore}; + + const auto flag = GENERATE(from_range(std::begin(validFlags), std::end(validFlags))); + + if (std::find(std::begin(exceptions), std::end(exceptions), flag) == std::end(exceptions)) { + // flag is not in list of exceptions + HIP_CHECK_ERROR(hipMalloc3DArray(&array, &desc, make_hipExtent(s, 0, s), flag), + hipErrorInvalidValue); + } +} + +TEST_CASE("Unit_hipMalloc3DArray_Negative_InvalidFlags") { + constexpr size_t s = 6; // 6 to keep cubemap happy + hipArray_t array; + hipChannelFormatDesc desc = hipCreateChannelDesc(); + +#if HT_AMD + const unsigned int flag = 0xDEADBEEF; +#else + const unsigned int flag = + GENERATE(0xDEADBEEF, hipArrayTextureGather | hipArraySurfaceLoadStore, + hipArrayTextureGather | hipArrayCubemap, + hipArrayTextureGather | hipArraySurfaceLoadStore | hipArrayCubemap); +#endif + + CAPTURE(flag); + + REQUIRE(std::find(std::begin(validFlags), std::end(validFlags), flag) == std::end(validFlags)); + + HIP_CHECK_ERROR(hipMalloc3DArray(&array, &desc, makeExtent(flag, s), flag), hipErrorInvalidValue); +} + +void testInvalidDescription(hipChannelFormatDesc desc){ + constexpr size_t s = 6; // 6 to keep cubemap happy + hipArray_t array; + +#if HT_NVIDIA + hipError_t expectedError = hipErrorUnknown; +#else + hipError_t expectedError = hipErrorInvalidValue; +#endif + + const auto flag = GENERATE(from_range(std::begin(validFlags), std::end(validFlags))); + HIP_CHECK_ERROR(hipMalloc3DArray(&array, &desc, makeExtent(flag, s), flag), expectedError); +} + +TEST_CASE("Unit_hipMalloc3DArray_Negative_InvalidFormat") { + hipChannelFormatDesc desc = hipCreateChannelDesc(); + desc.f = GENERATE(hipChannelFormatKindNone, 0xBEEF); + testInvalidDescription(desc); +} + +TEST_CASE("Unit_hipMalloc3DArray_Negative_BadChannelLayout") { +#if HT_AMD + HipTest::HIP_SKIP_TEST("EXSWCPHIPT-129"); + return; +#endif + + const int bits = GENERATE(8, 16, 32); + const hipChannelFormatKind formatKind = + GENERATE(hipChannelFormatKindSigned, hipChannelFormatKindUnsigned, hipChannelFormatKindFloat); + if (bits == 8 && formatKind == hipChannelFormatKindFloat) return; + + + hipChannelFormatDesc desc = GENERATE_COPY(hipCreateChannelDesc(bits, bits, bits, 0, formatKind), + hipCreateChannelDesc(0, bits, bits, 0, formatKind), + hipCreateChannelDesc(0, bits, bits, bits, formatKind), + hipCreateChannelDesc(bits, 0, bits, 0, formatKind), + hipCreateChannelDesc(bits, bits, 0, bits, formatKind), + hipCreateChannelDesc(0, 0, bits, 0, formatKind), + hipCreateChannelDesc(0, 0, bits, bits, formatKind)); + + INFO("kind: " << channelFormatString(formatKind)); + INFO("x: " << desc.x << ", y: " << desc.y << ", z: " << desc.z << ", w: " << desc.w); + + testInvalidDescription(desc); +} + +TEST_CASE("Unit_hipMalloc3DArray_Negative_8BitFloat") { + hipChannelFormatDesc desc = GENERATE(hipCreateChannelDesc(8, 0, 0, 0, hipChannelFormatKindFloat), + hipCreateChannelDesc(8, 8, 0, 0, hipChannelFormatKindFloat), + hipCreateChannelDesc(8, 8, 8, 8, hipChannelFormatKindFloat)); + + testInvalidDescription(desc); +} + +TEST_CASE("Unit_hipMalloc3DArray_Negative_DifferentChannelSizes") { +#if HT_AMD + HipTest::HIP_SKIP_TEST("EXSWCPHIPT-129"); + return; +#endif + + const int bitsX = GENERATE(8, 16, 32); + const int bitsY = GENERATE(8, 16, 32); + const int bitsZ = GENERATE(8, 16, 32); + const int bitsW = GENERATE(8, 16, 32); + if (bitsX == bitsY && bitsY == bitsZ && bitsZ == bitsW) return; // skip when they are equal + + const hipChannelFormatKind channelFormat = + GENERATE(hipChannelFormatKindSigned, hipChannelFormatKindUnsigned, hipChannelFormatKindFloat); + + if (channelFormat == hipChannelFormatKindFloat && + (bitsX == 8 || bitsY == 8 || bitsZ == 8 || bitsW == 8)) + return; // 8 bit floats aren't allowed + + hipChannelFormatDesc desc = hipCreateChannelDesc(bitsX, bitsY, bitsZ, bitsW, channelFormat); + + INFO("format: " << channelFormatString(channelFormat) << ", x bits: " << bitsX + << ", y bits: " << bitsY << ", z bits: " << bitsZ << ", w bits: " << bitsW); + + + testInvalidDescription(desc); +} + +TEST_CASE("Unit_hipMalloc3DArray_Negative_BadChannelSize") { + const int badBits = GENERATE(-1, 0, 10, 100); + const hipChannelFormatKind formatKind = + GENERATE(hipChannelFormatKindSigned, hipChannelFormatKindUnsigned, hipChannelFormatKindFloat); + hipChannelFormatDesc desc = hipCreateChannelDesc(badBits, badBits, badBits, badBits, formatKind); + + INFO("Number of bits: " << badBits); + + testInvalidDescription(desc); +} + + +// hipMalloc3DArray should handle the max numeric value gracefully. +TEST_CASE("Unit_hipMalloc3DArray_Negative_NumericLimit") { + hipArray_t arrayPtr; + hipChannelFormatDesc desc = hipCreateChannelDesc(); + + size_t size = std::numeric_limits::max(); + const auto flag = GENERATE(from_range(std::begin(validFlags), std::end(validFlags))); + HIP_CHECK_ERROR(hipMalloc3DArray(&arrayPtr, &desc, makeExtent(flag, size), flag), + hipErrorInvalidValue); +} diff --git a/tests/catch/unit/memory/hipMallocArray.cc b/tests/catch/unit/memory/hipMallocArray.cc index 22e1ce9864..f1f6ebc9d2 100644 --- a/tests/catch/unit/memory/hipMallocArray.cc +++ b/tests/catch/unit/memory/hipMallocArray.cc @@ -27,15 +27,14 @@ hipMallocArray API test scenarios #include #include -#if defined(_WIN32) || defined(_WIN64) #include -#endif +#include "hipArrayCommon.hh" -static constexpr auto NUM_W{4}; -static constexpr auto BIGNUM_W{100}; -static constexpr auto BIGNUM_H{100}; -static constexpr auto NUM_H{4}; -static constexpr auto ARRAY_LOOP{100}; +static constexpr size_t NUM_W{4}; +static constexpr size_t NUM_H{4}; +static constexpr size_t BIGNUM_W{100}; +static constexpr size_t BIGNUM_H{100}; +static constexpr int ARRAY_LOOP{100}; /* * This API verifies memory allocations for small and @@ -51,28 +50,30 @@ static constexpr auto ARRAY_LOOP{100}; * */ static void MallocArray_DiffSizes(int gpu) { - HIP_CHECK(hipSetDevice(gpu)); - std::vector> array_size{{NUM_W, NUM_H}, {BIGNUM_W, BIGNUM_H}}; - for (auto& size : array_size) { - std::array A_d; - size_t tot, avail, ptot, pavail; - hipChannelFormatDesc desc = hipCreateChannelDesc(); - HIP_CHECK(hipMemGetInfo(&pavail, &ptot)); - for (int i = 0; i < ARRAY_LOOP; i++) { - HIP_CHECK( - hipMallocArray(&A_d[i], &desc, std::get<0>(size), std::get<1>(size), hipArrayDefault)); - } - for (int i = 0; i < ARRAY_LOOP; i++) { - HIP_CHECK(hipFreeArray(A_d[i])); - } - HIP_CHECK(hipMemGetInfo(&avail, &tot)); - if ((pavail != avail)) { - HIPASSERT(false); - } + HIP_CHECK_THREAD(hipSetDevice(gpu)); + std::pair size = + GENERATE(std::make_pair(NUM_W, NUM_H), std::make_pair(BIGNUM_W, BIGNUM_H)); + hipChannelFormatDesc desc = hipCreateChannelDesc(); + std::array A_d; + size_t pavail, avail; + HIP_CHECK_THREAD(hipMemGetInfo(&pavail, nullptr)); + + for (int i = 0; i < ARRAY_LOOP; i++) { + HIP_CHECK_THREAD( + hipMallocArray(&A_d[i], &desc, std::get<0>(size), std::get<1>(size), hipArrayDefault)); } + for (int i = 0; i < ARRAY_LOOP; i++) { + HIP_CHECK_THREAD(hipFreeArray(A_d[i])); + } + + HIP_CHECK_THREAD(hipMemGetInfo(&avail, nullptr)); + REQUIRE_THREAD(pavail == avail); } -TEST_CASE("Unit_hipMallocArray_DiffSizes") { MallocArray_DiffSizes(0); } +TEST_CASE("Unit_hipMallocArray_DiffSizes") { + MallocArray_DiffSizes(0); + HIP_CHECK_THREAD_FINALIZE(); +} /* This testcase verifies the hipMallocArray API in multithreaded @@ -83,81 +84,25 @@ TEST_CASE("Unit_hipMallocArray_MultiThread") { std::vector threadlist; int devCnt = 0; devCnt = HipTest::getDeviceCount(); - size_t tot, avail, ptot, pavail; - HIP_CHECK(hipMemGetInfo(&pavail, &ptot)); + const auto pavail = getFreeMem(); for (int i = 0; i < devCnt; i++) { - // TODO the HIP_CHECK and HIPASSERT are not threadsafe so this test is broken. threadlist.push_back(std::thread(MallocArray_DiffSizes, i)); } for (auto& t : threadlist) { t.join(); } - HIP_CHECK(hipMemGetInfo(&avail, &tot)); + HIP_CHECK_THREAD_FINALIZE(); + const auto avail = getFreeMem(); if (pavail != avail) { WARN("Memory leak of hipMalloc3D API in multithreaded scenario"); REQUIRE(false); } } - -constexpr size_t BlockSize = 16; - -template struct type_and_size { - using type = T; - static constexpr size_t size = N; -}; - -// scalars are interpreted as a vector of 1 length. -// template using int_constant = std::integral_constant; -template struct vector_info; -template <> struct vector_info : type_and_size {}; -template <> struct vector_info : type_and_size {}; -template <> struct vector_info : type_and_size {}; -template <> struct vector_info : type_and_size {}; -template <> struct vector_info : type_and_size {}; -template <> struct vector_info : type_and_size {}; -template <> struct vector_info : type_and_size {}; - -template <> struct vector_info : type_and_size {}; -template <> struct vector_info : type_and_size {}; -template <> struct vector_info : type_and_size {}; -template <> struct vector_info : type_and_size {}; -template <> struct vector_info : type_and_size {}; -template <> struct vector_info : type_and_size {}; -template <> struct vector_info : type_and_size {}; - -template <> struct vector_info : type_and_size {}; -template <> struct vector_info : type_and_size {}; -template <> struct vector_info : type_and_size {}; -template <> struct vector_info : type_and_size {}; -template <> struct vector_info : type_and_size {}; -template <> struct vector_info : type_and_size {}; -template <> struct vector_info : type_and_size {}; - // Kernels /////////////////////////////////////// -// read from a texture using normalized coordinates -constexpr size_t ChannelToRead = 1; -template -__global__ void readFromTexture(T* output, hipTextureObject_t texObj, size_t width, size_t height, - bool textureGather) { - // Calculate normalized texture coordinates - const unsigned int x = blockIdx.x * blockDim.x + threadIdx.x; - const unsigned int y = blockIdx.y * blockDim.y + threadIdx.y; - const float u = x / (float)width; - - // Read from texture and write to global memory - if (height == 0) { - output[x] = tex1D(texObj, u); - } else { - const float v = y / (float)height; - output[y * width + x] = - textureGather ? tex2Dgather(texObj, u, v, ChannelToRead) : tex2D(texObj, u, v); - } -} - template __device__ void addOne(T* a) { using scalar_type = typename vector_info::type; auto as = reinterpret_cast(a); @@ -190,30 +135,6 @@ template size_t getAllocSize(const size_t width, const size_t heigh return sizeof(T) * width * (height ? height : 1); } -template void checkDataIsAscending(const std::vector& hostData) { - bool allMatch = true; - size_t i = 0; - for (; i < hostData.size(); ++i) { - allMatch = allMatch && hostData[i] == static_cast(i); - if (!allMatch) break; - } - INFO("hostData[" << i << "] == " << static_cast(hostData[i])); - REQUIRE(allMatch); -} - -const char* channelFormatString(hipChannelFormatKind formatKind) noexcept { - switch (formatKind) { - case hipChannelFormatKindFloat: - return "float"; - case hipChannelFormatKindSigned: - return "signed"; - case hipChannelFormatKindUnsigned: - return "unsigned"; - default: - return "error"; - } -} - // Tests ///////////////////////////////////////// // Test the default array by generating a texture from it then reading from that texture. @@ -458,12 +379,6 @@ void testArrayAsSurface(hipArray_t arrayPtr, const size_t width, const size_t he HIP_CHECK(hipFree(device_data)); } -size_t getFreeMem() { - size_t free = 0, total = 0; - HIP_CHECK(hipMemGetInfo(&free, &total)); - return free; -} - // The happy path of a default array and a SurfaceLoadStore array should work // Selection of types chosen to reduce compile times TEMPLATE_TEST_CASE("Unit_hipMallocArray_happy", "", uint, int, int4, ushort, short2, char, uchar2, @@ -480,7 +395,7 @@ TEMPLATE_TEST_CASE("Unit_hipMallocArray_happy", "", uint, int, int4, ushort, sho // pointer to the array in device memory hipArray_t arrayPtr{}; size_t width = 1024; - size_t height; + size_t height{}; SECTION("hipArrayDefault") { height = GENERATE(0, 1024); @@ -522,54 +437,55 @@ TEMPLATE_TEST_CASE("Unit_hipMallocArray_happy", "", uint, int, int4, ushort, sho // EXSWCPHIPT-71 - no equivalent value for maxSurface and maxTexture2DGather. TEMPLATE_TEST_CASE("Unit_hipMallocArray_MaxTexture_Default", "", uint, int4, ushort, short2, char, char4, float2, float4) { - int device; - HIP_CHECK(hipGetDevice(&device)); - hipDeviceProp_t prop; - HIP_CHECK(hipGetDeviceProperties(&prop, device)); size_t width, height; hipArray_t array{}; hipChannelFormatDesc desc = hipCreateChannelDesc(); const unsigned int flag = hipArrayDefault; + const Sizes sizes(flag); + CAPTURE(sizes.max1D, sizes.max2D, sizes.max3D); + + const size_t s = 64; + SECTION("Happy") { SECTION("1D - Max") { - width = prop.maxTexture1D; + width = sizes.max1D; height = 0; } SECTION("2D - Max Width") { - width = prop.maxTexture2D[0]; - height = 64; + width = sizes.max2D[0]; + height = s; } SECTION("2D - Max Height") { - width = 64; - height = prop.maxTexture2D[1]; + width = s; + height = sizes.max2D[1]; } SECTION("2D - Max Width and Height") { - width = prop.maxTexture2D[0]; - height = prop.maxTexture2D[1]; + width = sizes.max2D[0]; + height = sizes.max2D[1]; } auto maxArrayCreateError = hipMallocArray(&array, &desc, width, height, flag); - // this can try to alloc many GB of memory, so out of memory is fair + // this can try to alloc many GB of memory, so out of memory is acceptable if (maxArrayCreateError == hipErrorOutOfMemory) return; HIP_CHECK(maxArrayCreateError); HIP_CHECK(hipFreeArray(array)); } SECTION("Negative") { SECTION("1D - More Than Max") { - width = prop.maxTexture1D + 1; + width = sizes.max1D + 1; height = 0; } SECTION("2D - More Than Max Width") { - width = prop.maxTexture2D[0] + 1; - height = 64; + width = sizes.max2D[0] + 1; + height = s; } SECTION("2D - More Than Max Height") { - width = 64; - height = prop.maxTexture2D[1] + 1; + width = s; + height = sizes.max2D[1] + 1; } SECTION("2D - More Than Max Width and Height") { - width = prop.maxTexture2D[0] + 1; - height = prop.maxTexture2D[1] + 1; + width = sizes.max2D[0] + 1; + height = sizes.max2D[1] + 1; } HIP_CHECK_ERROR(hipMallocArray(&array, &desc, width, height, flag), hipErrorInvalidValue); } @@ -579,7 +495,7 @@ TEMPLATE_TEST_CASE("Unit_hipMallocArray_MaxTexture_Default", "", uint, int4, ush // Arrays with channels of different size are not allowed. TEST_CASE("Unit_hipMallocArray_Negative_DifferentChannelSizes") { #if HT_AMD - HipTest::HIP_SKIP_TEST("EXSWCPHIPT-59"); + HipTest::HIP_SKIP_TEST("EXSWCPHIPT-129"); return; #endif const int bitsX = GENERATE(8, 16, 32); @@ -634,10 +550,6 @@ TEST_CASE("Unit_hipMallocArray_Negative_ZeroWidth") { // Providing the array pointer as nullptr should return an error TEST_CASE("Unit_hipMallocArray_Negative_NullArrayPtr") { -#if HT_AMD - HipTest::HIP_SKIP_TEST("EXSWCPHIPT-45"); - return; -#endif hipChannelFormatDesc desc = hipCreateChannelDesc(); HIP_CHECK_ERROR(hipMallocArray(nullptr, &desc, 1024, 0, hipArrayDefault), hipErrorInvalidValue); @@ -645,10 +557,6 @@ TEST_CASE("Unit_hipMallocArray_Negative_NullArrayPtr") { // Providing the desc pointer as nullptr should return an error TEST_CASE("Unit_hipMallocArray_Negative_NullDescPtr") { -#if HT_AMD - HipTest::HIP_SKIP_TEST("EXSWCPHIPT-83"); - return; -#endif hipArray_t arrayPtr; HIP_CHECK_ERROR(hipMallocArray(&arrayPtr, nullptr, 1024, 0, hipArrayDefault), hipErrorInvalidValue); @@ -656,10 +564,6 @@ TEST_CASE("Unit_hipMallocArray_Negative_NullDescPtr") { // Inappropriate but related flags should still return an error TEST_CASE("Unit_hipMallocArray_Negative_BadFlags") { -#if HT_AMD - HipTest::HIP_SKIP_TEST("EXSWCPHIPT-72"); - return; -#endif hipChannelFormatDesc desc = hipCreateChannelDesc(); hipArray_t arrayPtr; @@ -757,7 +661,7 @@ TEST_CASE("Unit_hipMallocArray_Negative_3ChannelElement") { // The bit channel description should not allow any channels after a zero channel TEST_CASE("Unit_hipMallocArray_Negative_ChannelAfterZeroChannel") { #if HT_AMD - HipTest::HIP_SKIP_TEST("EXSWCPHIPT-59"); + HipTest::HIP_SKIP_TEST("EXSWCPHIPT-129"); return; #endif const int bits = GENERATE(8, 16, 32); diff --git a/tests/catch/unit/memory/hipMallocPitch.cc b/tests/catch/unit/memory/hipMallocPitch.cc index 08691b2045..33133f4a52 100644 --- a/tests/catch/unit/memory/hipMallocPitch.cc +++ b/tests/catch/unit/memory/hipMallocPitch.cc @@ -1,13 +1,16 @@ /* Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved. + Permission is hereby granted, free of charge, to any person obtaining a copy of this software and associated documentation files (the "Software"), to deal in the Software without restriction, including without limitation the rights to use, copy, modify, merge, publish, distribute, sublicense, and/or sell copies of the Software, and to permit persons to whom the Software is furnished to do so, subject to the following conditions: + The above copyright notice and this permission notice shall be included in all copies or substantial portions of the Software. + THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE @@ -17,20 +20,294 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. */ -/* -Test Scenarios of hipMallocPitch API -1. Negative Scenarios -2. Basic Functionality Scenario -3. Allocate memory using hipMallocPitch API, Launch Kernel validate result. -4. Allocate Memory in small chunks and large chunks and check for possible memory leaks -5. Allocate Memory using hipMallocPitch API, Memcpy2D on the allocated variables. -6. Multithreaded scenario -*/ - #include +#include +#include +#include "hip/driver_types.h" +#include +#include +#include #include #include +/** + * @brief Test hipMalloc3D, hipMallocPitch and hipMemAllocPitch with multiple input values. + * Checks that the memory has been allocated with the specified pitch and extent sizes. + */ + +struct MemoryInfo { + size_t freeMem; + size_t totalMem; +}; + +inline static MemoryInfo createMemoryInfo() { + MemoryInfo memoryInfo{}; + HIP_CHECK(hipMemGetInfo(&memoryInfo.freeMem, &memoryInfo.totalMem)); + return memoryInfo; +} + +static void validateMemory(void* devPtr, hipExtent extent, size_t pitch, + MemoryInfo memBeforeAllocation) { + INFO("Width: " << extent.width << " Height: " << extent.height << " Depth: " << extent.depth); + + MemoryInfo memAfterAllocation{createMemoryInfo()}; + const size_t theoreticalAllocatedMemory{pitch * extent.height * extent.depth}; + const size_t allocatedMemory = memBeforeAllocation.freeMem - memAfterAllocation.freeMem; + + if (theoreticalAllocatedMemory == 0) { + REQUIRE(theoreticalAllocatedMemory == allocatedMemory); + return; /* If there was no memory allocated then we don't need to do further checks. */ + } else { + REQUIRE(theoreticalAllocatedMemory <= allocatedMemory); + } + + std::unique_ptr hostPtr{new char[theoreticalAllocatedMemory]}; + std::memset(hostPtr.get(), 2, theoreticalAllocatedMemory); + + hipPitchedPtr devPitchedPtr{devPtr, pitch, extent.width, extent.height}; + hipPitchedPtr hostPitchedPtr{hostPtr.get(), pitch, extent.width, extent.height}; + + HIP_CHECK(hipMemset3D(devPitchedPtr, 1, extent)); + + hipMemcpy3DParms params{}; + params.srcPtr = devPitchedPtr; + params.kind = hipMemcpyKind::hipMemcpyDeviceToHost; + params.dstPtr = hostPitchedPtr; + params.extent = extent; + HIP_CHECK(hipMemcpy3D(¶ms)) + + bool mismatch = false; + for (size_t width = 0; width < extent.width; ++width) { + for (size_t height = 0; height < extent.height; ++height) { + for (size_t depth = 0; depth < extent.depth; ++depth) { + char* reinterpretedPtr = reinterpret_cast(hostPtr.get()); + size_t index = (pitch * extent.height * depth) + (pitch * height) + width; + if (*(reinterpretedPtr + index) != 1) { + mismatch = true; + } + } + } + } + REQUIRE(!mismatch); +} + +class ExtentGenerator { + public: + static constexpr size_t totalRandomValues{20}; + static constexpr size_t seed{1337}; + + std::uniform_int_distribution width_distribution{1, 1024}; + std::uniform_int_distribution height_distribution{1, 100}; + std::uniform_int_distribution depth_distribution{1, 100}; + + std::vector extents2D{}; + std::vector extents3D{}; + + static ExtentGenerator& getInstance() { + static ExtentGenerator instance; + return instance; + } + + private: + ExtentGenerator() { + std::mt19937 randomGenerator{seed}; + extents3D = std::vector{hipExtent{0, 0, 0}, hipExtent{1, 0, 0}, hipExtent{0, 1, 0}, + hipExtent{0, 0, 1}}; + + for (size_t i = 0; i < totalRandomValues; ++i) { + extents3D.push_back(hipExtent{width_distribution(randomGenerator), + height_distribution(randomGenerator), + depth_distribution(randomGenerator)}); + } + + extents2D = std::vector{hipExtent{0, 0, 1}, hipExtent{1, 0, 1}, hipExtent{0, 1, 1}}; + + for (size_t i = 0; i < totalRandomValues; ++i) { + extents2D.push_back( + hipExtent{width_distribution(randomGenerator), height_distribution(randomGenerator), 1}); + } + } +}; + +enum class AllocationApi { hipMalloc3D, hipMallocPitch, hipMemAllocPitch }; + +hipExtent generateExtent(AllocationApi api) { + hipExtent extent; + if (api == AllocationApi::hipMalloc3D) { + auto& extents3D = ExtentGenerator::getInstance().extents3D; + extent = GENERATE_REF(from_range(extents3D.begin(), extents3D.end())); + } else { + auto& extents2D = ExtentGenerator::getInstance().extents2D; + extent = GENERATE_REF(from_range(extents2D.begin(), extents2D.end())); + } + + return extent; +} + + +TEST_CASE("Unit_hipMalloc3D_ValidatePitch") { + hipPitchedPtr hipPitchedPtr; + hipExtent validExtent{generateExtent(AllocationApi::hipMalloc3D)}; + + MemoryInfo memBeforeAllocation{createMemoryInfo()}; + HIP_CHECK(hipMalloc3D(&hipPitchedPtr, validExtent)); + validateMemory(hipPitchedPtr.ptr, validExtent, hipPitchedPtr.pitch, memBeforeAllocation); + HIP_CHECK(hipFree(hipPitchedPtr.ptr)); +} + +TEST_CASE("Unit_hipMemAllocPitch_ValidatePitch") { + size_t pitch; + hipDeviceptr_t ptr; + hipExtent validExtent{generateExtent(AllocationApi::hipMemAllocPitch)}; + MemoryInfo memBeforeAllocation{createMemoryInfo()}; + unsigned int elementSizeBytes = GENERATE(4, 8, 16); + + if (validExtent.width == 0 || validExtent.height == 0) { + return; + } + + HIP_CHECK( + hipMemAllocPitch(&ptr, &pitch, validExtent.width, validExtent.height, elementSizeBytes)); + validateMemory(reinterpret_cast(ptr), validExtent, pitch, memBeforeAllocation); + HIP_CHECK(hipFree(reinterpret_cast(ptr))); +} + +TEST_CASE("Unit_hipMallocPitch_ValidatePitch") { +#if HT_AMD + HipTest::HIP_SKIP_TEST("TODO-FIX-EXTENT-GENERATOR"); + return; +#endif + size_t pitch; + void* ptr; + hipExtent validExtent{generateExtent(AllocationApi::hipMemAllocPitch)}; + MemoryInfo memBeforeAllocation{createMemoryInfo()}; + HIP_CHECK(hipMallocPitch(&ptr, &pitch, validExtent.width, validExtent.height)); + validateMemory(ptr, validExtent, pitch, memBeforeAllocation); + HIP_CHECK(hipFree(ptr)); +} + +TEST_CASE("Unit_hipMalloc3D_Negative") { + SECTION("Invalid ptr") { + hipExtent validExtent{1, 1, 1}; + HIP_CHECK_ERROR(hipMalloc3D(nullptr, validExtent), hipErrorInvalidValue); + } + + hipPitchedPtr ptr; + constexpr size_t maxSizeT = std::numeric_limits::max(); + +#if HT_NVIDIA + //TODO-MATCH-ERRORCODE + SECTION("Max size_t width") { + hipExtent validExtent{maxSizeT, 1, 1}; + HIP_CHECK_ERROR(hipMalloc3D(&ptr, validExtent), hipErrorInvalidValue); + } +#endif + + SECTION("Max size_t height") { + hipExtent validExtent{1, maxSizeT, 1}; + HIP_CHECK_ERROR(hipMalloc3D(&ptr, validExtent), hipErrorOutOfMemory); + } + + SECTION("Max size_t depth") { + hipExtent validExtent{1, 1, maxSizeT}; + HIP_CHECK_ERROR(hipMalloc3D(&ptr, validExtent), hipErrorOutOfMemory); + } + +#if HT_NVIDIA + //TODO-MATCH-ERRORCODE + SECTION("Max size_t all dimensions") { + hipExtent validExtent{maxSizeT, maxSizeT, maxSizeT}; + HIP_CHECK_ERROR(hipMalloc3D(&ptr, validExtent), hipErrorInvalidValue); + } +#endif +} + +TEST_CASE("Unit_hipMallocPitch_Negative") { + size_t pitch; + void* ptr; + constexpr size_t maxSizeT = std::numeric_limits::max(); + + SECTION("Invalid ptr") { + HIP_CHECK_ERROR(hipMallocPitch(nullptr, &pitch, 1, 1), hipErrorInvalidValue); + } + + SECTION("Invalid pitch") { + HIP_CHECK_ERROR(hipMallocPitch(&ptr, nullptr, 1, 1), hipErrorInvalidValue); + } + +#if HT_NVIDIA + //TODO-MATCH-ERRORCODE + SECTION("Max size_t width") { + HIP_CHECK_ERROR(hipMallocPitch(&ptr, &pitch, maxSizeT, 1), hipErrorInvalidValue); + } +#endif + + SECTION("Max size_t height") { + HIP_CHECK_ERROR(hipMallocPitch(&ptr, &pitch, 1, maxSizeT), hipErrorOutOfMemory); + } +} + +TEST_CASE("Unit_hipMemAllocPitch_Negative") { + size_t pitch; + hipDeviceptr_t ptr{}; + unsigned int validElementSizeBytes{4}; + constexpr size_t maxSizeT = std::numeric_limits::max(); + +#if HT_NVIDIA + /* Device synchronize is used here to initialize the device. + * Nvidia does not implicitly do it for this Api. And hipInit(0) does not work either. + */ + HIP_CHECK(hipDeviceSynchronize()); + + SECTION("Invalid elementSizeBytes") { + unsigned int invalidElementSizeBytes = GENERATE(0, 7, 12, 17); + HIP_CHECK_ERROR(hipMemAllocPitch(&ptr, &pitch, 1, 1, invalidElementSizeBytes), + hipErrorInvalidValue); + } + + SECTION("Zero width") { + HIP_CHECK_ERROR(hipMemAllocPitch(&ptr, &pitch, 0, 1, validElementSizeBytes), + hipErrorInvalidValue); + } + SECTION("Zero height") { + HIP_CHECK_ERROR(hipMemAllocPitch(&ptr, &pitch, 1, 0, validElementSizeBytes), + hipErrorInvalidValue); + } +#endif + + SECTION("Invalid dptr") { + HIP_CHECK_ERROR(hipMemAllocPitch(nullptr, &pitch, 1, 1, validElementSizeBytes), + hipErrorInvalidValue); + } + + SECTION("Invalid pitch") { + HIP_CHECK_ERROR(hipMemAllocPitch(&ptr, nullptr, 1, 1, validElementSizeBytes), + hipErrorInvalidValue); + } + +#if HT_NVIDIA + //TODO-MATCH-ERRORCODE + SECTION("Max size_t width") { + HIP_CHECK_ERROR(hipMemAllocPitch(&ptr, &pitch, maxSizeT, 1, validElementSizeBytes), + hipErrorInvalidValue); + } +#endif + + SECTION("Max size_t height") { + HIP_CHECK_ERROR(hipMemAllocPitch(&ptr, &pitch, 1, maxSizeT, validElementSizeBytes), + hipErrorOutOfMemory); + } +} + +/* +Test Scenarios of hipMallocPitch API +1. Basic Functionality Scenario +2. Allocate memory using hipMallocPitch API, Launch Kernel validate result. +3. Allocate Memory in small chunks and large chunks and check for possible memory leaks +4. Allocate Memory using hipMallocPitch API, Memcpy2D on the allocated variables. +5. Multithreaded scenario +*/ + static constexpr auto SMALLCHUNK_NUMW{4}; static constexpr auto SMALLCHUNK_NUMH{4}; static constexpr auto LARGECHUNK_NUMW{1025}; @@ -295,4 +572,3 @@ TEMPLATE_TEST_CASE("Unit_hipMallocPitch_KernelLaunch", "" A_h, B_h, C_h, false); } - diff --git a/tests/catch/unit/memory/hipMemVmm.cc b/tests/catch/unit/memory/hipMemVmm.cc new file mode 100644 index 0000000000..1d3041eae5 --- /dev/null +++ b/tests/catch/unit/memory/hipMemVmm.cc @@ -0,0 +1,85 @@ +/* + Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved. + Permission is hereby granted, free of charge, to any person obtaining a copy + of this software and associated documentation files (the "Software"), to deal + in the Software without restriction, including without limitation the rights + to use, copy, modify, merge, publish, distribute, sublicense, and/or sell + copies of the Software, and to permit persons to whom the Software is + furnished to do so, subject to the following conditions: + The above copyright notice and this permission notice shall be included in + all copies or substantial portions of the Software. + THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANNTY OF ANY KIND, EXPRESS OR + IMPLIED, INNCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANNY CLAIM, DAMAGES OR OTHER + LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN + THE SOFTWARE. + */ + +/* Test Case Description: + 1) This testcase verifies the basic scenario - supported on + all devices +*/ + +#include +#include +#include +#include +#include +#include +#include +#include +#include + +/* + This testcase verifies HIP Mem VMM API basic scenario - supported on all devices + */ + +TEST_CASE("Unit_hipMemVmm_Basic") { + int vmm = 0; + HIP_CHECK(hipDeviceGetAttribute(&vmm, hipDeviceAttributeVirtualMemoryManagementSupported, 0)); + INFO("hipDeviceAttributeVirtualMemoryManagementSupported: " << vmm); + + if (vmm == 0) { + SUCCEED("GPU 0 doesn't support hipDeviceAttributeVirtualMemoryManagement " + "attribute. Hence skipping the testing with Pass result.\n"); + return; + } + + size_t granularity = 0; + + hipMemAllocationProp memAllocationProp; + memAllocationProp.type = hipMemAllocationTypePinned; + memAllocationProp.location.id = 0; + memAllocationProp.location.type = hipMemLocationTypeDevice; + + HIP_CHECK(hipMemGetAllocationGranularity(&granularity, &memAllocationProp, hipMemAllocationGranularityRecommended)); + + size_t size = 4 * 1024; + void* reservedAddress{nullptr}; + HIP_CHECK(hipMemAddressReserve(&reservedAddress, size, granularity, nullptr, 0)); + + hipMemGenericAllocationHandle_t gaHandle{nullptr}; + HIP_CHECK(hipMemCreate(&gaHandle, size, &memAllocationProp, 0)); + + HIP_CHECK(hipMemMap(reservedAddress, size, 0, gaHandle, 0)); + + hipMemAccessDesc desc; + std::vector values(size); + const char value = 1; + + HIP_CHECK(hipMemSetAccess(reservedAddress, size, &desc, 1)); + HIP_CHECK(hipMemset(reservedAddress, value, size)); + HIP_CHECK(hipMemcpy(&values[0], reservedAddress, size, hipMemcpyDeviceToHost)); + + for (size_t i=0; i < size; ++i) { + REQUIRE(values[i] == value); + } + + HIP_CHECK(hipMemUnmap(reservedAddress, size)); + + HIP_CHECK(hipMemRelease(gaHandle)); + HIP_CHECK(hipMemAddressFree(reservedAddress, size)); +} + diff --git a/tests/catch/unit/memory/hipMemsetFunctional.cc b/tests/catch/unit/memory/hipMemsetFunctional.cc new file mode 100644 index 0000000000..274074266d --- /dev/null +++ b/tests/catch/unit/memory/hipMemsetFunctional.cc @@ -0,0 +1,560 @@ +/* + * Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved. + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to deal + * in the Software without restriction, including without limitation the rights + * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell + * copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * The above copyright notice and this permission notice shall be included in + * all copies or substantial portions of the Software. + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANNTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN + * THE SOFTWARE. + */ + +/** + * Testcase Scenarios: + * For hipMemset, hipMemsetD8, hipMemsetD16, hipMemsetD32, hipMemset2D, hipMemset3D and all async + * counterparts + * 1) (ZeroValue) - Test setting a specified range to zero. + * 2) (SmallSize) - Test setting a unique memset value for small sizes. + * 3) (ZeroSize) - Test that trying to set memory with a zero dimension does not fail and doesn't + * affect the memory. + * 4) (PartialSet) - Test setting a partial range of total allocated memory and + * ensure the full range isn't affected. + */ +#include + +constexpr size_t FULL_DIM = 10; + +// Enum used to determine which 1D memset function to use. +enum MemsetType { + hipMemsetTypeDefault = 1, + hipMemsetTypeD8 = 2, + hipMemsetTypeD16 = 3, + hipMemsetTypeD32 = 4 +}; + +// Macro used to assert all elements in a flat vector range is equal to a specified value +#define HIP_ASSERT_VEC_EQ(ptr, value, N) \ + for (size_t i = 0; i < N; i++) { \ + CAPTURE(N, i, ptr[i], value); \ + HIP_ASSERT(ptr[i] == value); \ + } + +// Copies device data to host and checks that each element is equal to the +// specified value +template void check_device_data(T* devPtr, T value, size_t numElems) { + std::unique_ptr hostPtr(new T[numElems]); + HIP_CHECK(hipMemcpy(hostPtr.get(), devPtr, numElems * sizeof(T), hipMemcpyDeviceToHost)); + HIP_ASSERT_VEC_EQ(hostPtr.get(), value, numElems); +} + +// Macro to assist calling and then checking the result of the 1D memset API with the necessary +// manipulation to the arguments. +#define HIP_MEMSET_CHECK(hipMemsetFunc, devPtr, value, count, async) \ + using scalar_t = decltype(value); \ + size_t sizeBytes = count * sizeof(scalar_t); \ + HIP_CHECK(hipMemsetFunc(devPtr, value, sizeBytes)); \ + if (async) { \ + HIP_CHECK(hipStreamSynchronize(stream)); \ + } \ + check_device_data(devPtr, value, count); + +#define HIP_MEMSET_CHECK_DTYPE(hipMemsetFunc, devPtr, value, count, async) \ + HIP_CHECK(hipMemsetFunc(reinterpret_cast(devPtr), value, count)); \ + if (async) { \ + HIP_CHECK(hipStreamSynchronize(stream)); \ + } \ + check_device_data(devPtr, value, count); + +// Enum for specifying wether to allocate the data using hipMalloc, hipHostMalloc or not at all. +enum MemsetMallocType { hipDeviceMalloc_t = 1, hipHostMalloc_t = 2, hipNoMalloc_t }; + +// Helper function for allocating memory, setting data with the specified 1D memset API and then +// checking result of operation. +template +void checkMemset(T value, size_t count, MemsetType memsetType, bool async = false, + MemsetMallocType mallocType = hipDeviceMalloc_t, T* devPtr = nullptr) { + hipStream_t stream{nullptr}; + if (async) { + hipStreamCreate(&stream); + } + + // Allocate Memory + if (mallocType == hipDeviceMalloc_t) { + HIP_CHECK(hipMalloc(&devPtr, count * sizeof(T))); + } else if (mallocType == hipHostMalloc_t) { + HIP_CHECK(hipHostMalloc(&devPtr, count * sizeof(T))); + } + + // memset API calls + switch (memsetType) { + case hipMemsetTypeDefault: + if (!async) { + INFO("Testing hipMemset call") + HIP_MEMSET_CHECK(hipMemset, devPtr, value, count, false); + } else { + INFO("Testing hipMemsetAsync call") + HIP_MEMSET_CHECK(hipMemsetAsync, devPtr, value, count, true); + } + break; + case hipMemsetTypeD8: + if (!async) { + INFO("Testing hipMemsetD8 call") + HIP_MEMSET_CHECK_DTYPE(hipMemsetD8, devPtr, value, count, false); + } else { + INFO("Testing hipMemsetD8Async call") + HIP_MEMSET_CHECK_DTYPE(hipMemsetD8Async, devPtr, value, count, true); + } + break; + case hipMemsetTypeD16: + if (!async) { + INFO("Testing hipMemsetD16 call") + HIP_MEMSET_CHECK_DTYPE(hipMemsetD16, devPtr, value, count, false); + } else { + INFO("Testing hipMemsetD16Async call") + HIP_MEMSET_CHECK_DTYPE(hipMemsetD16Async, devPtr, value, count, true); + } + break; + case hipMemsetTypeD32: + if (!async) { + INFO("Testing hipMemsetD32 call") + HIP_MEMSET_CHECK_DTYPE(hipMemsetD32, devPtr, value, count, false); + } else { + INFO("Testing hipMemsetD32Async call") + HIP_MEMSET_CHECK_DTYPE(hipMemsetD32Async, devPtr, value, count, true); + } + break; + } + + // Cleanup + if (async) { + HIP_CHECK(hipStreamDestroy(stream)); + } + + // Free memory + if (mallocType == hipDeviceMalloc_t) { + HIP_CHECK(hipFree(devPtr)); + } else if (mallocType == hipHostMalloc_t) { + HIP_CHECK(hipHostFree(devPtr)); + } +} + +// Macro which defines a TEST_CASE which calls and then checks the result of the 1D memset macros +// for all combinations of sync/async and hipMalloc/hipHostMalloc, given the value and memory range. +#define DEFINE_1D_BASIC_TEST_CASE(suffix, memsetType, T, value, count) \ + TEST_CASE("Unit_hipMemsetFunctional_" + std::string(suffix)) { \ + const std::string memsetStr = std::string(suffix); \ + SECTION(memsetStr + " - Device Malloc") { \ + checkMemset(static_cast(value), count, memsetType, false, hipDeviceMalloc_t); \ + } \ + SECTION(memsetStr + " - Host Malloc") { \ + checkMemset(static_cast(value), count, memsetType, false, hipHostMalloc_t); \ + } \ + SECTION(memsetStr + "Async - Device Malloc") { \ + checkMemset(static_cast(value), count, memsetType, true, hipDeviceMalloc_t); \ + } \ + SECTION(memsetStr + "Async - Host Malloc") { \ + checkMemset(static_cast(value), count, memsetType, true, hipHostMalloc_t); \ + } \ + } + +DEFINE_1D_BASIC_TEST_CASE("ZeroValue_hipMemset", hipMemsetTypeDefault, float, 0, 1024) +DEFINE_1D_BASIC_TEST_CASE("ZeroValue_hipMemsetD32", hipMemsetTypeD32, uint32_t, 0, 1024) +DEFINE_1D_BASIC_TEST_CASE("ZeroValue_hipMemsetD16", hipMemsetTypeD16, int16_t, 0, 1024) +DEFINE_1D_BASIC_TEST_CASE("ZeroValue_hipMemsetD8", hipMemsetTypeD8, int8_t, 0, 1024) + +DEFINE_1D_BASIC_TEST_CASE("SmallSize_hipMemset", hipMemsetTypeDefault, char, 0x42, 1) +DEFINE_1D_BASIC_TEST_CASE("SmallSize_hipMemsetD32", hipMemsetTypeD32, uint32_t, 0x101, 1) +DEFINE_1D_BASIC_TEST_CASE("SmallSize_hipMemsetD16", hipMemsetTypeD16, int16_t, 0x10, 1) +DEFINE_1D_BASIC_TEST_CASE("SmallSize_hipMemsetD8", hipMemsetTypeD8, int8_t, 0x1, 1) + +DEFINE_1D_BASIC_TEST_CASE("ZeroSize_hipMemset", hipMemsetTypeDefault, char, 0x42, 0) +DEFINE_1D_BASIC_TEST_CASE("ZeroSize_hipMemsetD32", hipMemsetTypeD32, uint32_t, 0x101, 0) +DEFINE_1D_BASIC_TEST_CASE("ZeroSize_hipMemsetD16", hipMemsetTypeD16, int16_t, 0x10, 0) +DEFINE_1D_BASIC_TEST_CASE("ZeroSize_hipMemsetD8", hipMemsetTypeD8, int8_t, 0x1, 0) + +// Helper function that sets a full region of memory with an initial value, sets a smaller subregion +// with another value and check that the memset API do not write outside of the subregion of data. +template +void partialMemsetTest(T valA, T valB, size_t count, size_t offset, MemsetType memsetType, + bool async) { + T* devPtr; + size_t subSize{count - offset}; + HIP_CHECK(hipMalloc(&devPtr, count * sizeof(T))); + + // Set entire region to be first value. + INFO("Setting full region"); + checkMemset(valA, count, memsetType, async, hipNoMalloc_t, devPtr); + + // Set partial region to be second value. + INFO("Setting partial region"); + checkMemset(valB, subSize, memsetType, async, hipNoMalloc_t, devPtr + offset); + + // Ensure the first section remains unchanged + check_device_data(devPtr, valA, offset); + HIP_CHECK(hipFree(devPtr)); +} + +TEST_CASE("Unit_hipMemsetFunctional_PartialSet_1D") { + for (auto widthOffset = 8; widthOffset <= 8; widthOffset *= 2) { + SECTION("hipMemset - Partial Set") { + partialMemsetTest(0x1, 0x42, 1024, widthOffset, hipMemsetTypeDefault, false); + } + SECTION("hipMemsetAsync - Partial Set") { + partialMemsetTest(0x1, 0x42, 1024, widthOffset, hipMemsetTypeDefault, true); + } + SECTION("hipMemsetD8 - Partial Set") { + partialMemsetTest(0x1, 0xDE, 1024, widthOffset, hipMemsetTypeD8, false); + } + SECTION("hipMemsetD8Async - Partial Set") { + partialMemsetTest(0x1, 0xDE, 1024, widthOffset, hipMemsetTypeD8, true); + } + SECTION("hipMemsetD16 - Partial Set") { + partialMemsetTest(0x1, 0xDEAD, 1024, widthOffset, hipMemsetTypeD16, false); + } + SECTION("hipMemsetD16Async - Partial Set") { + partialMemsetTest(0x1, 0xDEAD, 1024, widthOffset, hipMemsetTypeD16, true); + } + SECTION("hipMemsetD32 - Partial Set") { + partialMemsetTest(0x1, 0xDEADBEEF, 1024, widthOffset, hipMemsetTypeD32, false); + } + SECTION("hipMemsetD32Async - Partial Set") { + partialMemsetTest(0x1, 0xDEADBEEF, 1024, widthOffset, hipMemsetTypeD32, true); + } + } +} + +// Helper function that copies the device data to the host and returns a unique_ptr to that data. +template +std::unique_ptr get_device_data_2D(T* devPtr, size_t pitch, size_t width, size_t height) { + std::unique_ptr hostPtr(new T[width * height]); + constexpr size_t elementSize = sizeof(T); + HIP_CHECK(hipMemcpy2D(hostPtr.get(), width * elementSize, devPtr, pitch, width, height, + hipMemcpyDeviceToHost)); + return hostPtr; +} + +// Copies device data to host and checks that each element is equal to the +// specified value +template +void check_device_data_2D(T* devPtr, T value, size_t pitch, size_t width, size_t height) { + auto hostPtr = get_device_data_2D(devPtr, pitch, width, height); + HIP_ASSERT_VEC_EQ(hostPtr.get(), value, width * height); +} + +// Helper function for allocating memory, setting data with the specified 2D memset API and then +// checking result of operation. +template +void checkMemset2D(T value, size_t width, size_t height, bool async = false, size_t pitch = 0, + T* devPtr = nullptr) { + hipStream_t stream{nullptr}; + hipStreamCreate(&stream); + constexpr size_t elementSize = sizeof(T); + bool freeDevPtr = false; + if (devPtr == nullptr) { + freeDevPtr = true; + HIP_CHECK( + hipMallocPitch(reinterpret_cast(&devPtr), &pitch, width * elementSize, height)); + } + + if (!async) { + INFO("Testing hipMemset2D call") + HIP_CHECK(hipMemset2D(devPtr, pitch, value, width * elementSize, height)); + } else { + INFO("Testing hipMemset2DAsync call") + HIP_CHECK(hipMemset2DAsync(devPtr, pitch, value, width * elementSize, height, stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + } + if (width * height > 0) { + check_device_data_2D(devPtr, value, pitch, width, height); + } + if (freeDevPtr) { + HIP_CHECK(hipFree(devPtr)); + } + hipStreamDestroy(stream); +} + +TEST_CASE("Unit_hipMemsetFunctional_ZeroValue_2D") { + constexpr size_t width{128}; + constexpr size_t height{128}; + constexpr char memsetVal = 0; + SECTION("hipMemset2D - Zero Value") { checkMemset2D(memsetVal, width, height, false); } + SECTION("hipMemset2DAsync - Zero Value") { checkMemset2D(memsetVal, width, height, true); } +} + +TEST_CASE("Unit_hipMemsetFunctional_SmallSize_2D") { + constexpr char memsetVal = 0x42; + SECTION("hipMemset2D - Small Size") { checkMemset2D(memsetVal, 1, 1, false); } + SECTION("hipMemset2DAsync - Small Size") { checkMemset2D(memsetVal, 1, 1, true); } +} + +TEST_CASE("Unit_hipMemsetFunctional_ZeroSize_2D") { + size_t pitch{0}; + size_t width{10}; + size_t height{10}; + char* devPtr{nullptr}; + HIP_CHECK( + hipMallocPitch(reinterpret_cast(&devPtr), &pitch, width * sizeof(char), height)); + + const char initValue = 0x1; + const char testValue = 0x11; + // Set full region to initial value + checkMemset2D(initValue, width, height, false, pitch, devPtr); + + SECTION("hipMemset2D - Zero Width") { + checkMemset2D(testValue, 0, height, false, pitch, devPtr); + check_device_data_2D(devPtr, initValue, pitch, width, height); + } + SECTION("hipMemset2DAsync - Zero Width") { + checkMemset2D(testValue, 0, height, true, pitch, devPtr); + check_device_data_2D(devPtr, initValue, pitch, width, height); + } + SECTION("hipMemset2D - Zero Height") { + checkMemset2D(testValue, width, 0, false, pitch, devPtr); + check_device_data_2D(devPtr, initValue, pitch, width, height); + } + SECTION("hipMemset2DAsync - Zero Height") { + checkMemset2D(testValue, width, 0, true, pitch, devPtr); + check_device_data_2D(devPtr, initValue, pitch, width, height); + } + SECTION("hipMemset2D - Zero Width and Height") { + checkMemset2D(testValue, 0, 0, false, pitch, devPtr); + check_device_data_2D(devPtr, initValue, pitch, width, height); + } + SECTION("hipMemset2DAsync - Zero Width and Height") { + checkMemset2D(testValue, 0, 0, true, pitch, devPtr); + check_device_data_2D(devPtr, initValue, pitch, width, height); + } + HIP_CHECK(hipFree(devPtr)); +} + +// Helper function that sets a full region of memory with an initial value, sets a smaller subregion +// with another value and check that the memset API do not write outside of the subregion of data. +template +void partialMemsetTest2D(T valA, T valB, size_t width, size_t height, size_t widthOffset, + size_t heightOffset, bool async) { + T* devPtr{nullptr}; + size_t pitch{0}; + size_t subWidth{width - widthOffset}; + size_t subHeight{height - heightOffset}; + constexpr size_t elementSize = sizeof(T); + HIP_CHECK(hipMallocPitch(reinterpret_cast(&devPtr), &pitch, width * elementSize, height)); + + // Set entire region to be first value. + INFO("Setting full square region"); + checkMemset2D(valA, width, height, async, pitch, devPtr); + + // Set partial region to be second value. + INFO("Setting partial square region") + checkMemset2D(valB, subWidth, subHeight, async, pitch, devPtr); + + auto hostPtr = get_device_data_2D(devPtr, pitch, width, height); + T comparVal{0}; + size_t idx{0}; + for (size_t i = 0; i < width; i++) { + for (size_t j = 0; j < height; j++) { + if (i < subWidth && j < subHeight) { + // Compare subregion value + comparVal = valB; + } else { + // Compare full region value + comparVal = valA; + } + idx = i * height + j; + CAPTURE(width, height, subWidth, subHeight, i, j, idx, hostPtr[idx], comparVal); + HIP_ASSERT(hostPtr[idx] == comparVal); + } + } + HIP_CHECK(hipFree(devPtr)); +} + +TEST_CASE("Unit_hipMemsetFunctional_PartialSet_2D") { + for (auto widthOffset = 8; widthOffset <= 128; widthOffset *= 2) { + for (auto heightOffset = 8; heightOffset <= 128; heightOffset *= 2) { + SECTION("hipMemset2D - Partial Set") { + partialMemsetTest2D('a', 'b', 200, 200, widthOffset, heightOffset, false); + } + SECTION("hipMemset2DAsync - Partial Set") { + partialMemsetTest2D('a', 'b', 200, 200, widthOffset, heightOffset, true); + } + } + } +} + +// Helper function that copies the device data to the host and returns a unique_ptr to that data. +template +std::unique_ptr get_device_data_3D(hipPitchedPtr& devPitchedPtr, hipExtent extent) { + constexpr size_t elementSize = sizeof(T); + std::unique_ptr hostPtr( + new T[devPitchedPtr.pitch * extent.width * extent.height / elementSize]); + hipMemcpy3DParms myparms{}; + myparms.srcPos = make_hipPos(0, 0, 0); + myparms.dstPos = make_hipPos(0, 0, 0); + myparms.dstPtr = make_hipPitchedPtr(hostPtr.get(), devPitchedPtr.pitch, + extent.width / elementSize, extent.height); + myparms.srcPtr = devPitchedPtr; + myparms.extent = extent; + myparms.kind = hipMemcpyDeviceToHost; + HIP_CHECK(hipMemcpy3D(&myparms)); + return hostPtr; +} + +// Copies device data to host and checks that each element is equal to the +// specified value +template +void check_device_data_3D(hipPitchedPtr& devPitchedPtr, T value, hipExtent extent) { + auto hostPtr = get_device_data_3D(devPitchedPtr, extent); + size_t width = extent.width / sizeof(T); + size_t height = extent.height; + size_t depth = extent.depth; + size_t idx; + for (size_t k = 0; k < depth; k++) { + for (size_t j = 0; j < height; j++) { + for (size_t i = 0; i < width; i++) { + idx = devPitchedPtr.pitch * height * k + devPitchedPtr.pitch * j + i; + INFO("idx=" << idx << " hostPtr[idx]=" << hostPtr[idx] << " value=" << value) + HIP_ASSERT(hostPtr[idx] == value); + } + } + } +} + +// Helper function for allocating memory, setting data with the specified 3D memset API and then +// checking result of operation. +template +void checkMemset3D(hipPitchedPtr& devPitchedPtr, T value, hipExtent extent, bool async = false) { + hipStream_t stream{nullptr}; + hipStreamCreate(&stream); + if (devPitchedPtr.ptr == nullptr) { + HIP_CHECK(hipMalloc3D(&devPitchedPtr, extent)); + } + if (!async) { + INFO("Testing hipMemset3D call") + HIP_CHECK(hipMemset3D(devPitchedPtr, value, extent)); + } else { + INFO("Testing hipMemset3DAsync call") + HIP_CHECK(hipMemset3DAsync(devPitchedPtr, value, extent, stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + } + if (extent.width * extent.height * extent.depth > 0) { + check_device_data_3D(devPitchedPtr, value, extent); + } + hipStreamDestroy(stream); +} + +void check_memset_3D(std::string sectionStr, size_t width, size_t height, size_t depth, + char value) { + hipPitchedPtr devPitchedPtr; + hipExtent fullExtent; + constexpr char fullVal = 0x21; + hipExtent extent = make_hipExtent(width, height, depth); + // Check if any of the dimensions are zero + bool anyZero = width * height * depth == 0; + if (anyZero) { + // If they are zero then set a full region with memset value to later check if it's changed. + devPitchedPtr.ptr = nullptr; + fullExtent = make_hipExtent(FULL_DIM, FULL_DIM, FULL_DIM); + checkMemset3D(devPitchedPtr, fullVal, fullExtent, false); + } + SECTION("hipMemset3D - " + sectionStr) { + if (!anyZero) { + devPitchedPtr.ptr = nullptr; + } + checkMemset3D(devPitchedPtr, value, extent, false); + if (anyZero) { + // Check to make sure memsets with a zero dimension did not affect above set region. + check_device_data_3D(devPitchedPtr, fullVal, fullExtent); + } + HIP_CHECK(hipFree(devPitchedPtr.ptr)); + } + SECTION("hipMemset3DAsync - " + sectionStr) { + if (!anyZero) { + devPitchedPtr.ptr = nullptr; + } + checkMemset3D(devPitchedPtr, value, extent, true); + if (anyZero) { + // Check to make sure memsets with a zero dimension did not affect above set region. + check_device_data_3D(devPitchedPtr, fullVal, fullExtent); + } + HIP_CHECK(hipFree(devPitchedPtr.ptr)); + } +} + +TEST_CASE("Unit_hipMemsetFunctional_ZeroValue_3D") { + check_memset_3D("Zero Value", 128, 128, 10, 0); +} + +TEST_CASE("Unit_hipMemsetFunctional_SmallSize_3D") { check_memset_3D("Small Size", 1, 1, 1, 0x42); } + +TEST_CASE("Unit_hipMemsetFunctional_ZeroSize_3D") { + constexpr size_t elementSize = sizeof(char); + check_memset_3D("Zero Width", 0, FULL_DIM, FULL_DIM, 0x23); + check_memset_3D("Zero Height", FULL_DIM * elementSize, 0, FULL_DIM, 0x23); + check_memset_3D("Zero Depth", FULL_DIM * elementSize, FULL_DIM, 0, 0x23); + check_memset_3D("Zero Width and Height", 0 * elementSize, 0, FULL_DIM, 0x23); + check_memset_3D("Zero Width and Depth", 0 * elementSize, FULL_DIM, 0, 0x23); + check_memset_3D("Zero Height and Depth", FULL_DIM * elementSize, 0, 0, 0x23); + check_memset_3D("Zero Width, Height and Depth", 0 * elementSize, 0, 0, 0x23); +} + +// Helper function that sets a full region of memory with an initial value, sets a smaller subregion +// with another value and check that the memset API do not write outside of the subregion of data. +template +void partialMemsetTest3D(T valA, T valB, size_t width, size_t height, size_t depth, + size_t widthOffset, size_t heightOffset, size_t depthOffset, bool async) { + size_t subWidth{width - widthOffset}; + size_t subHeight{height - heightOffset}; + size_t subDepth{depth - depthOffset}; + hipPitchedPtr devPitchedPtr; + devPitchedPtr.ptr = nullptr; + hipExtent extent = make_hipExtent(width * sizeof(T), height, depth); + hipExtent subExtent = make_hipExtent(subWidth * sizeof(T), subHeight, subDepth); + + // Set entire region to be first value. + INFO("Setting full cuboid region") { checkMemset3D(devPitchedPtr, valA, extent, async); } + // Set partial region to be second value. + INFO("Setting partial cuboid region") { checkMemset3D(devPitchedPtr, valB, subExtent, async); } + auto pitch = devPitchedPtr.pitch; + auto hostPtr = get_device_data_3D(devPitchedPtr, extent); + T comparVal{0}; + size_t idx{0}; + for (size_t k = 0; k < depth; k++) { + for (size_t j = 0; j < height; j++) { + for (size_t i = 0; i < width; i++) { + if (i < subWidth && j < subHeight && k < subDepth) { + comparVal = valB; + } else { + comparVal = valA; + } + idx = devPitchedPtr.pitch * height * k + devPitchedPtr.pitch * j + i; + CAPTURE(width, height, depth, pitch, subWidth, subHeight, subDepth, i, j, k, idx, + hostPtr[idx], comparVal); + HIP_ASSERT(hostPtr[idx] == comparVal); + } + } + } + HIP_CHECK(hipFree(devPitchedPtr.ptr)); +} + +TEST_CASE("Unit_hipMemsetFunctional_PartialSet_3D") { + for (auto widthOffset = 8; widthOffset <= 128; widthOffset *= 2) { + for (auto heightOffset = 8; heightOffset <= 128; heightOffset *= 2) { + for (auto depthOffset = 2; depthOffset <= 5; depthOffset++) { + SECTION("hipMemset3D - Partial Set") { + partialMemsetTest3D('a', 'b', 200, 200, 10, widthOffset, heightOffset, depthOffset, + false); + } + SECTION("hipMemset3DAsync - Partial Set") { + partialMemsetTest3D('a', 'b', 200, 200, 10, widthOffset, heightOffset, depthOffset, true); + } + } + } + } +} diff --git a/tests/catch/unit/printf/CMakeLists.txt b/tests/catch/unit/printf/CMakeLists.txt index b911796d66..d2a9d6eced 100644 --- a/tests/catch/unit/printf/CMakeLists.txt +++ b/tests/catch/unit/printf/CMakeLists.txt @@ -16,3 +16,10 @@ elseif (HIP_PLATFORM MATCHES "nvidia") TEST_TARGET_NAME build_tests COMPILE_OPTIONS -std=c++17) endif() + +# Standalone exes +add_executable(printfFlags EXCLUDE_FROM_ALL printfFlags_exe.cc) +add_executable(printfSpecifiers EXCLUDE_FROM_ALL printfSpecifiers_exe.cc) + +add_dependencies(printfTests printfFlags) +add_dependencies(printfTests printfSpecifiers) diff --git a/tests/catch/unit/printf/printfFlags.cc b/tests/catch/unit/printf/printfFlags.cc index 8a9aa75a3f..b9955b0b60 100644 --- a/tests/catch/unit/printf/printfFlags.cc +++ b/tests/catch/unit/printf/printfFlags.cc @@ -1,5 +1,5 @@ /* -Copyright (c) 2021 Advanced Micro Devices, Inc. All rights reserved. +Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved. Permission is hereby granted, free of charge, to any person obtaining a copy of this software and associated documentation files (the "Software"), to deal @@ -23,20 +23,6 @@ THE SOFTWARE. #include #include -__global__ void test_kernel() { - printf("%08d\n", 42); - printf("%08i\n", -42); - printf("%08u\n", 42); - printf("%08g\n", 123.456); - printf("%0+8d\n", 42); - printf("%+d\n", -42); - printf("%+08d\n", 42); - printf("%-8s\n", "xyzzy"); - printf("% i\n", -42); - printf("%-16.8d\n", 42); - printf("%16.8d\n", 42); -} - TEST_CASE("Unit_printf_flags") { std::string reference(R"here(00000042 -0000042 @@ -51,7 +37,7 @@ xyzzy 00000042 )here"); - hip::SpawnProc proc("printfExe/printfFlags", true); + hip::SpawnProc proc("printfFlags", true); REQUIRE(proc.run() == 0); REQUIRE(proc.getOutput() == reference); } diff --git a/tests/catch/unit/printfExe/printfFlags.cc b/tests/catch/unit/printf/printfFlags_exe.cc similarity index 95% rename from tests/catch/unit/printfExe/printfFlags.cc rename to tests/catch/unit/printf/printfFlags_exe.cc index a96c099340..6ce9691b17 100644 --- a/tests/catch/unit/printfExe/printfFlags.cc +++ b/tests/catch/unit/printf/printfFlags_exe.cc @@ -1,5 +1,5 @@ /* -Copyright (c) 2021 Advanced Micro Devices, Inc. All rights reserved. +Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved. Permission is hereby granted, free of charge, to any person obtaining a copy of this software and associated documentation files (the "Software"), to deal diff --git a/tests/catch/unit/printf/printfSpecifiers.cc b/tests/catch/unit/printf/printfSpecifiers.cc index 37e8584943..223a3c338b 100644 --- a/tests/catch/unit/printf/printfSpecifiers.cc +++ b/tests/catch/unit/printf/printfSpecifiers.cc @@ -1,5 +1,5 @@ /* -Copyright (c) 2021 Advanced Micro Devices, Inc. All rights reserved. +Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved. Permission is hereby granted, free of charge, to any person obtaining a copy of this software and associated documentation files (the "Software"), to deal @@ -89,7 +89,7 @@ x )here"); #endif - hip::SpawnProc proc("printfExe/printfSepcifiers", true); + hip::SpawnProc proc("printfSpecifiers", true); REQUIRE(0 == proc.run()); REQUIRE(proc.getOutput() == reference); } diff --git a/tests/catch/unit/printfExe/printfSepcifiers.cc b/tests/catch/unit/printf/printfSpecifiers_exe.cc similarity index 97% rename from tests/catch/unit/printfExe/printfSepcifiers.cc rename to tests/catch/unit/printf/printfSpecifiers_exe.cc index 364376fb0a..5dcabdc4d4 100644 --- a/tests/catch/unit/printfExe/printfSepcifiers.cc +++ b/tests/catch/unit/printf/printfSpecifiers_exe.cc @@ -1,5 +1,5 @@ /* -Copyright (c) 2021 Advanced Micro Devices, Inc. All rights reserved. +Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved. Permission is hereby granted, free of charge, to any person obtaining a copy of this software and associated documentation files (the "Software"), to deal diff --git a/tests/catch/unit/printfExe/CMakeLists.txt b/tests/catch/unit/printfExe/CMakeLists.txt deleted file mode 100644 index 3b73810b35..0000000000 --- a/tests/catch/unit/printfExe/CMakeLists.txt +++ /dev/null @@ -1,5 +0,0 @@ -add_executable(printfFlags EXCLUDE_FROM_ALL printfFlags.cc) -add_executable(printfSepcifiers EXCLUDE_FROM_ALL printfSepcifiers.cc) - -add_dependencies(build_tests printfFlags) -add_dependencies(build_tests printfSepcifiers) diff --git a/tests/catch/unit/stream/CMakeLists.txt b/tests/catch/unit/stream/CMakeLists.txt index 932e76f718..229df157c1 100644 --- a/tests/catch/unit/stream/CMakeLists.txt +++ b/tests/catch/unit/stream/CMakeLists.txt @@ -25,9 +25,13 @@ set(TEST_SRC hipStreamCreateWithFlags.cc hipStreamCreateWithPriority.cc hipAPIStreamDisable.cc + # hipStreamAttachMemAsync.cc # Disabling it on nvidia due to issue in function definition of hipStreamAttachMemAsync + # Fixing would break ABI, to be re-enabled when the fix is made. streamCommon.cc hipStreamValue.cc ) + + set_source_files_properties(hipStreamAttachMemAsync.cc PROPERTIES COMPILE_FLAGS -std=c++17) endif() hip_add_exe_to_target(NAME StreamTest diff --git a/tests/catch/unit/stream/hipStreamAttachMemAsync.cc b/tests/catch/unit/stream/hipStreamAttachMemAsync.cc new file mode 100644 index 0000000000..e02754dfbb --- /dev/null +++ b/tests/catch/unit/stream/hipStreamAttachMemAsync.cc @@ -0,0 +1,143 @@ +/* +Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved. +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +// TODO Enable it after hipStreamAttachMemAsync is feature complete on HIP + +#include +#include + +__device__ __managed__ int var = 0; + +enum class StreamAttachTestType { NullStream = 0, StreamPerThread, CreatedStream }; + +TEST_CASE("Unit_hipStreamAttachMemAsync_Negative") { + hipStream_t stream{nullptr}; + + auto streamType = + GENERATE(StreamAttachTestType::NullStream, StreamAttachTestType::StreamPerThread, + StreamAttachTestType::CreatedStream); + + if (streamType == StreamAttachTestType::StreamPerThread) { + stream = hipStreamPerThread; + } else if (streamType == StreamAttachTestType::CreatedStream) { + HIP_CHECK(hipStreamCreate(&stream)); + REQUIRE(stream != nullptr); + } + + SECTION("Invalid Resource Handle") { + int definitelyNotAManagedVariable = 0; + HIP_CHECK_ERROR( + hipStreamAttachMemAsync(stream, reinterpret_cast(&definitelyNotAManagedVariable), + sizeof(int), hipMemAttachSingle), + hipErrorInvalidValue); + } + + SECTION("Invalid devptr") { + HIP_CHECK_ERROR(hipStreamAttachMemAsync(stream, nullptr, sizeof(int), hipMemAttachSingle), + hipErrorInvalidValue); + } + + SECTION("Invalid Resource Size") { + HIP_CHECK_ERROR(hipStreamAttachMemAsync(stream, reinterpret_cast(&var), sizeof(int) - 1, + hipMemAttachSingle), + hipErrorInvalidValue); + } + + SECTION("Invalid Flags") { + HIP_CHECK_ERROR( + hipStreamAttachMemAsync(stream, reinterpret_cast(&var), sizeof(int) - 1, + hipMemAttachSingle | hipMemAttachHost | hipMemAttachGlobal), + hipErrorInvalidValue); + } + + if (streamType == StreamAttachTestType::CreatedStream) { + HIP_CHECK(hipStreamDestroy(stream)); + } +} + +__global__ void kernel(int* ptr, size_t size) { + auto i = threadIdx.x; + if (i < size) { + ptr[i] = 1024; + } +} + +constexpr size_t size = 1024; +__device__ __managed__ int m_memory[size]; + +TEST_CASE("Unit_hipStreamAttachMemAsync_UseCase") { + hipStream_t stream{nullptr}; + + auto streamType = + GENERATE(StreamAttachTestType::NullStream, StreamAttachTestType::StreamPerThread, + StreamAttachTestType::CreatedStream); + + if (streamType == StreamAttachTestType::CreatedStream) { + HIP_CHECK(hipStreamCreate(&stream)); + REQUIRE(stream != nullptr); + } + + SECTION("Size zero is valid") { + int* d_memory{nullptr}; + HIP_CHECK(hipMallocManaged(&d_memory, sizeof(int) * size, hipMemAttachHost)); + HIP_CHECK( + hipStreamAttachMemAsync(stream, reinterpret_cast(d_memory), 0, hipMemAttachHost)); + HIP_CHECK(hipStreamSynchronize(stream)); // Wait for command to complete + HIP_CHECK(hipFree(d_memory)); + } + + SECTION("Access from device and host") { + int* d_memory{nullptr}; + + HIP_CHECK(hipMallocManaged(&d_memory, sizeof(int) * size, hipMemAttachHost)); + HIP_CHECK(hipMemset(d_memory, 0, sizeof(int) * size)); + HIP_CHECK( + hipStreamAttachMemAsync(stream, reinterpret_cast(d_memory), 0, hipMemAttachHost)); + HIP_CHECK(hipStreamSynchronize(stream)); // Wait for the command to complete + + kernel<<<1, size, 0, stream>>>(d_memory, size); + HIP_CHECK(hipStreamSynchronize(stream)); // Wait for the kernel to complete + + auto ptr = std::make_unique(size); + std::copy(d_memory, d_memory + size, ptr.get()); + + HIP_CHECK(hipFree(d_memory)); + + REQUIRE(std::all_of(ptr.get(), ptr.get() + size, [](int n) { return n == size; })); + } + + SECTION("Access ManagedMemory") { + HIP_CHECK(hipMemset(m_memory, 0, sizeof(int) * size)); + HIP_CHECK( + hipStreamAttachMemAsync(stream, reinterpret_cast(m_memory), 0, hipMemAttachHost)); + HIP_CHECK(hipStreamSynchronize(stream)); // Wait for the command to complete + + kernel<<<1, size, 0, stream>>>(m_memory, size); + HIP_CHECK(hipStreamSynchronize(stream)); // Wait for the kernel to complete + + auto ptr = std::make_unique(size); + std::copy(m_memory, m_memory + size, ptr.get()); + + REQUIRE(std::all_of(ptr.get(), ptr.get() + size, [](int n) { return n == size; })); + } + + if (streamType == StreamAttachTestType::CreatedStream) { + HIP_CHECK(hipStreamDestroy(stream)); + } +} \ No newline at end of file diff --git a/tests/catch/unit/streamperthread/CMakeLists.txt b/tests/catch/unit/streamperthread/CMakeLists.txt index a47e1e8c69..dd5b54dabc 100644 --- a/tests/catch/unit/streamperthread/CMakeLists.txt +++ b/tests/catch/unit/streamperthread/CMakeLists.txt @@ -4,6 +4,7 @@ set(TEST_SRC hipStreamPerThread_Event.cc hipStreamPerThread_MultiThread.cc hipStreamPerThread_DeviceReset.cc + hipStreamPerThrdTsts.cc ) hip_add_exe_to_target(NAME StreamPerThreadTest diff --git a/tests/catch/unit/streamperthread/hipStreamPerThrdTsts.cc b/tests/catch/unit/streamperthread/hipStreamPerThrdTsts.cc new file mode 100644 index 0000000000..8aba0fa0f5 --- /dev/null +++ b/tests/catch/unit/streamperthread/hipStreamPerThrdTsts.cc @@ -0,0 +1,597 @@ +/* +Copyright (c) 2021 Advanced Micro Devices, Inc. All rights reserved. +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + + +/* Test Description: + Scenario-1: Launch a kernel in hipStreamPerThread, while it is in flight + check for hipStreamQuery(hipStreamPerThread) it should return + hipErrorNotReady. + Scenario-2: Testing hipStreamPerThread stream object with hipMallocManaged() + memory + Scenario-3: To check the working of hipStreamPerThread in forked process + Scenario-4: The following test case tests the working of hipEventSynchronize + in multiple threads which are launched in quick succession + Scenario-5: The following test case checks the working of + hipStreamWaitEvent() with hipStreamWaitEvent() + Scenario-6: Testing hipLaunchCooperativeKernel() api with hipStreamPerThread + Scenario-7: Testing hipLaunchCooperativeKernelMultiDevice() with + hipStreamPerThread +*/ +#include +#include +#include +#ifdef _WIN32 + #include + #define sleep(x) _sleep(x) +#endif +#ifdef __linux__ + #include + #include + #include +#endif + +#include +#ifdef HT_AMD + #include "hip/hip_cooperative_groups.h" +#endif +using namespace std::chrono; +using namespace cooperative_groups; +#if HT_AMD +#define HIPRT_CB +#endif + + +static bool IfTestPassed = false; +// kernel +__global__ void StreamPerThrd(int *Ad, int *Ad1, size_t n, int Pk_Clk, + int Wait, int WaitEvnt = 0) { + size_t index = blockIdx.x * blockDim.x + threadIdx.x; + if (index < n) { + Ad[index] = Ad[index] + 10; + } + if (Wait) { + int64_t GpuFrq = (Pk_Clk * 1000); + int64_t StrtTck = clock64(); + if (index == 0) { + // The following while loop checks the value in ptr for around 4 seconds + while ((clock64() - StrtTck) <= (6 * GpuFrq)) { + } + if (WaitEvnt == 1) { + *Ad1 = 1; + } + } + } +} + + +__global__ void StreamPerThrd1(int *A, int Pk_Clk) { + int64_t GpuFrq = (Pk_Clk * 1000); + int64_t StrtTck = clock64(); + // The following while loop checks the value in ptr for around 3-4 seconds + while ((clock64() - StrtTck) <= (3 * GpuFrq)) { + } + *A = 1; +} + +__global__ void MiniKernel(int *A) { + if (*A == 0) { + *A = 2; // Fail condition + } else if (*A == 1) { + *A = 3; // Pass condition + } else { + *A = 4; // Garbage value found in A + } +} + +__global__ void StreamPerThrdCoopKrnl(int *Ad, int *n) { + int NumElms = (*n); + int index = blockIdx.x * blockDim.x + threadIdx.x; + if (index < NumElms) { + Ad[index] = Ad[index] + 10; + } +} + +#if HT_AMD +__global__ void test_gwsPerThrd(uint* buf, uint bufSize, int64_t* tmpBuf, + int64_t* result) { + extern __shared__ int64_t tmp[]; + uint groups = gridDim.x; + uint group_id = blockIdx.x; + uint local_id = threadIdx.x; + uint chunk = gridDim.x * blockDim.x; + + uint i = group_id * blockDim.x + local_id; + int64_t sum = 0; + while (i < bufSize) { + sum += buf[i]; + i += chunk; + } + tmp[local_id] = sum; + __syncthreads(); + i = 0; + if (local_id == 0) { + sum = 0; + while (i < blockDim.x) { + sum += tmp[i]; + i++; + } + tmpBuf[group_id] = sum; + } + + // wait + cooperative_groups::this_grid().sync(); + + if (((blockIdx.x * blockDim.x) + threadIdx.x) == 0) { + for (uint i = 1; i < groups; ++i) { + sum += tmpBuf[i]; + } + // *result = sum; + result[1 + cooperative_groups::this_multi_grid().grid_rank()] = sum; + } + cooperative_groups::this_multi_grid().sync(); + if (cooperative_groups::this_multi_grid().grid_rank() == 0) { + sum = 0; + for (uint i = 1; i <= cooperative_groups::this_multi_grid().num_grids(); + ++i) { + sum += result[i]; + } + *result = sum; + } +} +#endif +static const uint BufferSizeInDwords = 256 * 1024 * 1024; +static constexpr uint NumKernelArgs = 4; +static constexpr uint MaxGPUs = 8; +// callback function +static void HIPRT_CB CallBackFunctn(hipStream_t strm, hipError_t err, + void *ChkVal) { + // The following HIPASSERT() is just to satisfy catch2 framework. + // As it ensures the use of all the variables. + HIPASSERT(strm); + HIPCHECK(err); + if (*(reinterpret_cast(ChkVal)) == 1) { + IfTestPassed = true; + } else { + IfTestPassed = false; + } +} + +static void EventSync() { + int *Ad = nullptr, *Ah = nullptr, NumElms = 4096, CONST_NUM = 123; + int blockSize = 32, peak_clk; + HIP_CHECK(hipMalloc(&Ad, NumElms * sizeof(int))); + Ah = new int[NumElms]; + for (int i = 0; i < NumElms; ++i) { + Ah[i] = CONST_NUM; + } + // creating event objects + hipEvent_t start, end; + HIP_CHECK(hipEventCreate(&start)); + HIP_CHECK(hipEventCreate(&end)); + HIP_CHECK(hipMemcpy(Ad, Ah, NumElms * sizeof(int), hipMemcpyHostToDevice)); + HIP_CHECK(hipDeviceGetAttribute(&peak_clk, hipDeviceAttributeClockRate, 0)); + dim3 dimBlock(blockSize, 1, 1); + dim3 dimGrid((NumElms + blockSize -1)/blockSize, 1, 1); + HIP_CHECK(hipEventRecord(start, hipStreamPerThread)); + StreamPerThrd<<>>(Ad, NULL, NumElms, + peak_clk, 0); + HIP_CHECK(hipEventRecord(end, hipStreamPerThread)); + HIP_CHECK(hipEventSynchronize(end)); + HIP_CHECK(hipMemcpy(Ah, Ad, NumElms * sizeof(int), hipMemcpyDeviceToHost)); + int MisMatch = 0; + for (int i = 0; i < NumElms; ++i) { + if (Ah[i] != (CONST_NUM + 10)) { + MisMatch++; + } + } + delete[] Ah; + HIP_CHECK(hipFree(Ad)); + if (MisMatch) { + WARN("Data Mismatch observed!!\n"); + IfTestPassed = false; + } else { + IfTestPassed = true; + } +} + +/* Launch a kernel in hipStreamPerThread, while it is in flight check for + hipStreamQuery(hipStreamPerThread) it should return hipErrorNotReady.*/ +TEST_CASE("Unit_hipStreamPerThreadTst_StrmQuery") { + int *Ad = nullptr, *Ah = nullptr, NumElms = 4096, CONST_NUM = 123; + int blockSize = 32, peak_clk; + hipError_t err; + HIP_CHECK(hipMalloc(&Ad, NumElms * sizeof(int))); + Ah = new int[NumElms]; + for (int i = 0; i < NumElms; ++i) { + Ah[i] = CONST_NUM; + } + HIP_CHECK(hipMemcpy(Ad, Ah, NumElms * sizeof(int), hipMemcpyHostToDevice)); + HIP_CHECK(hipDeviceGetAttribute(&peak_clk, hipDeviceAttributeClockRate, 0)); + dim3 dimBlock(blockSize, 1, 1); + dim3 dimGrid((NumElms + blockSize -1)/blockSize, 1, 1); + SECTION("Test working of hipStreamQuery") { + StreamPerThrd<<>>(Ad, NULL, + NumElms, peak_clk, 1); + err = hipStreamQuery(hipStreamPerThread); + if (err != hipErrorNotReady) { + WARN("hipStreamQuery on hipStreamPerThread didnt return expected error!"); + IfTestPassed = false; + } else { + IfTestPassed = true; + } + } + SECTION("check working of hipStreamAddCallback() with hipStreamPerThread") { + int *Hptr = nullptr, *A_d = nullptr; + HIP_CHECK(hipHostMalloc(&Hptr, sizeof(int))); + *Hptr = 0; + HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast(&A_d), Hptr, 0)); + StreamPerThrd1<<<1, 1, 0, hipStreamPerThread>>>(A_d, peak_clk); + HIP_CHECK(hipStreamAddCallback(hipStreamPerThread, CallBackFunctn, A_d, 0)); + HIP_CHECK(hipStreamSynchronize(hipStreamPerThread)); + HIP_CHECK(hipHostFree(Hptr)); + } + HIP_CHECK(hipFree(Ad)); + delete[] Ah; + REQUIRE(IfTestPassed); +} + +/* Testing hipStreamPerThread stream object with hipMallocManaged() memory*/ +TEST_CASE("Unit_hipStreamPerThread_MangdMem") { + int managed = 0; + HIP_CHECK(hipDeviceGetAttribute(&managed, hipDeviceAttributeManagedMemory, + 0)); + if (managed == 1) { + int *Hmm = nullptr, NumElms = 4096, CONST_NUM = 123, blockSize = 32; + SECTION("Using Managed memory") { + HIP_CHECK(hipMallocManaged(&Hmm, NumElms * sizeof(int))); + for (int i = 0; i < NumElms; ++i) { + Hmm[i] = CONST_NUM; + } + } + SECTION("Prefetching Managed memory to device") { + HIP_CHECK(hipMallocManaged(&Hmm, NumElms * sizeof(int))); + for (int i = 0; i < NumElms; ++i) { + Hmm[i] = CONST_NUM; + } + HIP_CHECK(hipMemPrefetchAsync(Hmm, NumElms * sizeof(int), 0, + hipStreamPerThread)); + } + int peak_clk; + HIP_CHECK(hipDeviceGetAttribute(&peak_clk, hipDeviceAttributeClockRate, 0)); + dim3 dimBlock(blockSize, 1, 1); + dim3 dimGrid((NumElms + blockSize -1)/blockSize, 1, 1); + StreamPerThrd<<>>(Hmm, NULL, + NumElms, peak_clk, 0); + HIP_CHECK(hipStreamSynchronize(hipStreamPerThread)); + // Validating the result + int MisMatch = 0; + for (int i = 0; i < NumElms; ++i) { + if (Hmm[i] != (CONST_NUM + 10)) { + MisMatch++; + } + } + HIP_CHECK(hipFree(Hmm)); + if (MisMatch) { + WARN("Data mismatch observed!!\n"); + REQUIRE(false); + } + } else { + SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory " + "attribute. Hence skipping the testing with Pass result.\n"); + } +} + +/* To check the working of hipStreamPerThread in forked process*/ +#ifdef __linux__ +TEST_CASE("Unit_hipStreamPerThread_ChildProc") { + if (fork() == 0) { // child process + int *Ad = nullptr, *Ah = nullptr, NumElms = 4096, CONST_NUM = 123; + int blockSize = 32, peak_clk; + HIP_CHECK(hipMalloc(&Ad, NumElms * sizeof(int))); + Ah = new int[NumElms]; + for (int i = 0; i < NumElms; ++i) { + Ah[i] = CONST_NUM; + } + HIP_CHECK(hipMemcpy(Ad, Ah, NumElms * sizeof(int), hipMemcpyHostToDevice)); + HIP_CHECK(hipDeviceGetAttribute(&peak_clk, hipDeviceAttributeClockRate, 0)); + dim3 dimBlock(blockSize, 1, 1); + dim3 dimGrid((NumElms + blockSize -1)/blockSize, 1, 1); + StreamPerThrd<<>>(Ad, NULL, + NumElms, peak_clk, 0); + HIP_CHECK(hipStreamSynchronize(hipStreamPerThread)); + HIP_CHECK(hipMemcpy(Ah, Ad, NumElms * sizeof(int), hipMemcpyDeviceToHost)); + int MisMatch = 0; + for (int i = 0; i < NumElms; ++i) { + if (Ah[i] != (CONST_NUM + 10)) { + MisMatch++; + } + } + delete[] Ah; + HIP_CHECK(hipFree(Ad)); + if (MisMatch) { + WARN("Data Mismatch observed!!\n"); + exit(9); + } else { + exit(10); + } + } else { // Parent process + int stat; + wait(&stat); + int Result = WEXITSTATUS(stat); + if (Result != 10) { + REQUIRE(false); + } + } +} +#endif + +/* The following test case tests the working of hipEventSynchronize in + multiple threads which are launched in quick succession*/ +TEST_CASE("Unit_hipStreamPerThread_EvtRcrdMThrd") { + IfTestPassed = true; + int MAX_THREAD_CNT = 20; + std::vector threads(MAX_THREAD_CNT); + for (auto &th : threads) { + th = std::thread(EventSync); + } + for (auto& th : threads) { + th.join(); + } + REQUIRE(IfTestPassed); +} + +/* The following test case checks the working of hipStreamWaitEvent() with + hipStreamWaitEvent()*/ +TEST_CASE("Unit_hipStreamPerThread_StrmWaitEvt") { + IfTestPassed = true; + int *Ad = nullptr, NumElms = 4096, CONST_NUM = 123, blockSize = 32, *Ah = nullptr; + int *Ad1 = nullptr, *Ah1 = nullptr; + Ah = new int[NumElms]; + Ah1 = new int; + hipStream_t Strm; + HIP_CHECK(hipStreamCreate(&Strm)); + for (int i = 0; i < NumElms; ++i) { + Ah[i] = CONST_NUM; + } + Ah1[0] = 0; + HIP_CHECK(hipMalloc(&Ad, NumElms * sizeof(int))); + HIP_CHECK(hipMemcpy(Ad, Ah, NumElms * sizeof(int), hipMemcpyHostToDevice)); + memset(Ah, 0, NumElms * sizeof(int)); + HIP_CHECK(hipMalloc(&Ad1, sizeof(int))); + HIP_CHECK(hipMemset(Ad1, 0, sizeof(int))); + int peak_clk; + HIP_CHECK(hipDeviceGetAttribute(&peak_clk, hipDeviceAttributeClockRate, 0)); + dim3 dimBlock(blockSize, 1, 1); + dim3 dimGrid((NumElms + blockSize -1)/blockSize, 1, 1); + hipEvent_t e1; + HIPCHECK(hipEventCreate(&e1)); + StreamPerThrd<<>>(Ad, Ad1, NumElms, + peak_clk, 1, 1); + HIP_CHECK(hipEventRecord(e1, Strm)); + HIP_CHECK(hipStreamWaitEvent(hipStreamPerThread, e1, 0 /*flags*/)); + MiniKernel<<<1, 1, 0, hipStreamPerThread>>>(Ad1); + sleep(1); + HIP_CHECK(hipMemcpy(Ah1, Ad1, sizeof(int), hipMemcpyDeviceToHost)); + if (*Ah1 != 3) { + IfTestPassed = false; + if (*Ah1 == 2) { + WARN("hipStreamPerThread didn't honour hipStreamWaitEvent()"); + } else if (*Ah1 == 4) { + WARN("Unexpected behavior observed with hipStreamPerThread"); + } + } + // Validating the result + HIP_CHECK(hipMemcpy(Ah, Ad, NumElms * sizeof(int), hipMemcpyDeviceToHost)); + int MisMatch = 0; + for (int i = 0; i < NumElms; ++i) { + if (Ah[i] != (CONST_NUM + 10)) { + MisMatch++; + } + } + HIP_CHECK(hipFree(Ad)); + HIP_CHECK(hipFree(Ad1)); + HIP_CHECK(hipStreamDestroy(Strm)); + delete[] Ah; + delete Ah1; + if (MisMatch) { + WARN("Data mismatch observed!!\n"); + IfTestPassed = false; + } + REQUIRE(IfTestPassed); +} + + +/* Testing hipLaunchCooperativeKernel() api with hipStreamPerThread*/ +TEST_CASE("Unit_hipStreamPerThread_CoopLaunch") { + hipDeviceProp_t device_properties; + HIPCHECK(hipGetDeviceProperties(&device_properties, 0)); + /* Test whether target device supports cooperative groups ****************/ + if (device_properties.cooperativeLaunch == 0) { + SUCCEED("Cooperative group support not available..."); + } else { + /* We will launch enough waves to fill up all of the GPU *****************/ + int warp_size = device_properties.warpSize; + int num_sms = device_properties.multiProcessorCount; + // long long totalTicks = device_properties.clockRate ; + int max_blocks_per_sm = 0; + // Calculate the device occupancy to know how many blocks can be run. + HIPCHECK(hipOccupancyMaxActiveBlocksPerMultiprocessor(&max_blocks_per_sm, + StreamPerThrdCoopKrnl, + warp_size, 0)); + int max_active_blocks = max_blocks_per_sm * num_sms; + int *Ad = nullptr, *Ah = nullptr, *DNumElms = nullptr, NumElms = 4096; + int Const = 123; + Ah = new int[NumElms]; + for (int i = 0; i < NumElms; ++i) { + Ah[i] = Const; + } + HIP_CHECK(hipMalloc(&Ad, sizeof(int) * NumElms)); + HIP_CHECK(hipMalloc(&DNumElms, sizeof(int))); + HIP_CHECK(hipMemcpyAsync(Ad, Ah, sizeof(int) * NumElms, + hipMemcpyHostToDevice, hipStreamPerThread)); + HIP_CHECK(hipMemcpyAsync(DNumElms, &NumElms, sizeof(int), + hipMemcpyHostToDevice, hipStreamPerThread)); + HIP_CHECK(hipStreamSynchronize(hipStreamPerThread)); + + void *coop_params[2]; + coop_params[0] = reinterpret_cast(&Ad); + coop_params[1] = reinterpret_cast(&DNumElms); + HIP_CHECK(hipLaunchCooperativeKernel( + reinterpret_cast(StreamPerThrdCoopKrnl), + max_active_blocks, warp_size, + coop_params, 0, hipStreamPerThread)); + HIP_CHECK(hipMemcpy(Ah, Ad, sizeof(int) * NumElms, hipMemcpyDeviceToHost)); + // Verifying the result + int DataMismatch = 0; + for (int i = 0; i < NumElms; ++i) { + if (Ah[i] != (Const + 10)) { + DataMismatch++; + } + } + HIP_CHECK(hipFree(Ad)); + HIP_CHECK(hipFree(DNumElms)); + delete[] Ah; + if (DataMismatch > 0) { + REQUIRE(false); + } + } +} + +/* Testing hipLaunchCooperativeKernelMultiDevice() with hipStreamPerThread*/ +#if HT_AMD +TEST_CASE("Unit_hipStreamPerThread_CoopLaunchMDev") { + uint* dA[MaxGPUs]; + int64_t* dB[MaxGPUs]; + int64_t* dC; + + uint32_t* init = new uint32_t[BufferSizeInDwords]; + for (uint32_t i = 0; i < BufferSizeInDwords; ++i) { + init[i] = i; + } + + int nGpu = 0; + HIPCHECK(hipGetDeviceCount(&nGpu)); + size_t copySizeInDwords = BufferSizeInDwords / nGpu; + hipDeviceProp_t deviceProp[MaxGPUs]; + + for (int i = 0; i < nGpu; i++) { + HIPCHECK(hipSetDevice(i)); + + // Calculate the device occupancy to know how many blocks can be + // run concurrently + hipGetDeviceProperties(&deviceProp[i], 0); + if (!deviceProp[i].cooperativeMultiDeviceLaunch) { + WARN("Device doesn't support cooperative launch!"); + SUCCEED(""); + } + size_t SIZE = copySizeInDwords * sizeof(uint); + + HIPCHECK(hipMalloc(reinterpret_cast(&dA[i]), SIZE)); + HIPCHECK(hipMalloc(reinterpret_cast(&dB[i]), + 64 * deviceProp[i].multiProcessorCount * sizeof(int64_t))); + if (i == 0) { + HIPCHECK(hipHostMalloc(reinterpret_cast(&dC), + (nGpu + 1) * sizeof(int64_t))); + } + HIPCHECK(hipMemcpy(dA[i], &init[i * copySizeInDwords] , SIZE, + hipMemcpyHostToDevice)); + hipDeviceSynchronize(); + } + + dim3 dimBlock; + dim3 dimGrid; + dimGrid.x = 1; + dimGrid.y = 1; + dimGrid.z = 1; + dimBlock.x = 64; + dimBlock.y = 1; + dimBlock.z = 1; + + int numBlocks = 0; + uint workgroups[3] = {64, 128, 256}; + + hipLaunchParams* launchParamsList = new hipLaunchParams[nGpu]; + std::time_t end_time; + double time = 0; + for (uint set = 0; set < 3; ++set) { + void* args[MaxGPUs * NumKernelArgs]; + WARN("---------- Test#" << set << ", size: "<< BufferSizeInDwords << + " dwords ---------------\n"); + for (int i = 0; i < nGpu; i++) { + HIPCHECK(hipSetDevice(i)); + dimBlock.x = workgroups[set]; + HIPCHECK(hipOccupancyMaxActiveBlocksPerMultiprocessor(&numBlocks, + test_gwsPerThrd, dimBlock.x * dimBlock.y * dimBlock.z, + dimBlock.x * sizeof(int64_t))); + + WARN("GPU(" << i << ") Block size: " << dimBlock.x << + " Num blocks per CU: " << numBlocks << "\n"); + + dimGrid.x = deviceProp[i].multiProcessorCount * (std::min)(numBlocks, 32); + + args[i * NumKernelArgs] = reinterpret_cast(&dA[i]); + args[i * NumKernelArgs + 1] = reinterpret_cast(©SizeInDwords); + args[i * NumKernelArgs + 2] = reinterpret_cast(&dB[i]); + args[i * NumKernelArgs + 3] = reinterpret_cast(&dC); + + launchParamsList[i].func = reinterpret_cast(test_gwsPerThrd); + launchParamsList[i].gridDim = dimGrid; + launchParamsList[i].blockDim = dimBlock; + launchParamsList[i].sharedMem = dimBlock.x * sizeof(int64_t); + + launchParamsList[i].stream = hipStreamPerThread; + launchParamsList[i].args = &args[i * NumKernelArgs]; + } + + system_clock::time_point start = system_clock::now(); + hipLaunchCooperativeKernelMultiDevice(launchParamsList, nGpu, 0); + for (int i = 0; i < nGpu; i++) { + HIP_CHECK(hipSetDevice(i)); + HIP_CHECK(hipDeviceSynchronize()); + } + system_clock::time_point end = system_clock::now(); + std::chrono::duration elapsed_seconds = end - start; + end_time = std::chrono::system_clock::to_time_t(end); + + time += elapsed_seconds.count(); + + size_t processedDwords = copySizeInDwords * nGpu; + if (*dC != (((int64_t)(processedDwords) * (processedDwords - 1)) / 2)) { + WARN("Data validation failed ("<< *dC << " != " << + (((int64_t)(BufferSizeInDwords) * (BufferSizeInDwords - 1)) / 2) << + ") for grid size = " << dimGrid.x << " and block size = " << + dimBlock.x << "\n"); + WARN("Test failed!"); + } + } + + delete [] launchParamsList; + + WARN("finished computation at " << std::ctime(&end_time)); + WARN("elapsed time: " << time << "s\n"); + + hipSetDevice(0); + hipFree(dC); + for (int i = 0; i < nGpu; i++) { + hipFree(dA[i]); + hipFree(dB[i]); + } + delete [] init; +} +#endif diff --git a/tests/performance/compute/hipPerfDotProduct.cpp b/tests/performance/compute/hipPerfDotProduct.cpp index 3b2e0c72dd..e30d5ab039 100644 --- a/tests/performance/compute/hipPerfDotProduct.cpp +++ b/tests/performance/compute/hipPerfDotProduct.cpp @@ -38,7 +38,7 @@ __global__ void vectors_not_equal(int n, const double* __restrict__ x, const double* __restrict__ y, double* __restrict__ workspace) { - int gid = hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x; + int gid = blockIdx.x * blockDim.x + threadIdx.x; double sum = 0.0; for(int idx = gid; idx < n; idx += hipGridDim_x * hipBlockDim_x) { @@ -46,51 +46,51 @@ __global__ void vectors_not_equal(int n, } __shared__ double sdata[BLOCKSIZE]; - sdata[hipThreadIdx_x] = sum; + sdata[threadIdx.x] = sum; __syncthreads(); - if(hipThreadIdx_x < 128) { - sdata[hipThreadIdx_x] += sdata[hipThreadIdx_x + 128]; + if(threadIdx.x < 128) { + sdata[threadIdx.x] += sdata[threadIdx.x + 128]; } __syncthreads(); - if(hipThreadIdx_x < 64){ - sdata[hipThreadIdx_x] += sdata[hipThreadIdx_x + 64]; + if(threadIdx.x < 64){ + sdata[threadIdx.x] += sdata[threadIdx.x + 64]; } __syncthreads(); - if(hipThreadIdx_x < 32){ - sdata[hipThreadIdx_x] += sdata[hipThreadIdx_x + 32]; + if(threadIdx.x < 32){ + sdata[threadIdx.x] += sdata[threadIdx.x + 32]; } __syncthreads(); - if(hipThreadIdx_x < 16) { - sdata[hipThreadIdx_x] += sdata[hipThreadIdx_x + 16]; + if(threadIdx.x < 16) { + sdata[threadIdx.x] += sdata[threadIdx.x + 16]; } __syncthreads(); - if(hipThreadIdx_x < 8) { - sdata[hipThreadIdx_x] += sdata[hipThreadIdx_x + 8]; + if(threadIdx.x < 8) { + sdata[threadIdx.x] += sdata[threadIdx.x + 8]; } __syncthreads(); - if(hipThreadIdx_x < 4) { - sdata[hipThreadIdx_x] += sdata[hipThreadIdx_x + 4]; + if(threadIdx.x < 4) { + sdata[threadIdx.x] += sdata[threadIdx.x + 4]; } __syncthreads(); - if(hipThreadIdx_x < 2) { - sdata[hipThreadIdx_x] += sdata[hipThreadIdx_x + 2]; + if(threadIdx.x < 2) { + sdata[threadIdx.x] += sdata[threadIdx.x + 2]; } __syncthreads(); - if(hipThreadIdx_x < 1) { - sdata[hipThreadIdx_x] += sdata[hipThreadIdx_x + 1]; + if(threadIdx.x < 1) { + sdata[threadIdx.x] += sdata[threadIdx.x + 1]; } - if(hipThreadIdx_x == 0) { - workspace[hipBlockIdx_x] = sdata[0]; + if(threadIdx.x == 0) { + workspace[blockIdx.x] = sdata[0]; } } @@ -99,59 +99,59 @@ template __launch_bounds__(BLOCKSIZE) __global__ void vectors_equal(int n, const double* __restrict__ x, double* __restrict__ workspace) { - int gid = hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x; + int gid = blockIdx.x * blockDim.x + threadIdx.x; double sum = 0.0; - for(int idx = gid; idx < n; idx += hipGridDim_x * hipBlockDim_x) { + for(int idx = gid; idx < n; idx += hipGridDim_x * blockDim.x) { sum = fma(x[idx], x[idx], sum); } __shared__ double sdata[BLOCKSIZE]; - sdata[hipThreadIdx_x] = sum; + sdata[threadIdx.x] = sum; __syncthreads(); - if(hipThreadIdx_x < 128) { - sdata[hipThreadIdx_x] += sdata[hipThreadIdx_x + 128]; + if(threadIdx.x < 128) { + sdata[threadIdx.x] += sdata[threadIdx.x + 128]; } __syncthreads(); - if(hipThreadIdx_x < 64) { - sdata[hipThreadIdx_x] += sdata[hipThreadIdx_x + 64]; + if(threadIdx.x < 64) { + sdata[threadIdx.x] += sdata[threadIdx.x + 64]; } __syncthreads(); - if(hipThreadIdx_x < 32) { - sdata[hipThreadIdx_x] += sdata[hipThreadIdx_x + 32]; + if(threadIdx.x < 32) { + sdata[threadIdx.x] += sdata[threadIdx.x + 32]; } __syncthreads(); - if(hipThreadIdx_x < 16) { - sdata[hipThreadIdx_x] += sdata[hipThreadIdx_x + 16]; + if(threadIdx.x < 16) { + sdata[threadIdx.x] += sdata[threadIdx.x + 16]; } __syncthreads(); - if(hipThreadIdx_x < 8) { - sdata[hipThreadIdx_x] += sdata[hipThreadIdx_x + 8]; + if(threadIdx.x < 8) { + sdata[threadIdx.x] += sdata[threadIdx.x + 8]; } __syncthreads(); - if(hipThreadIdx_x < 4) { - sdata[hipThreadIdx_x] += sdata[hipThreadIdx_x + 4]; + if(threadIdx.x < 4) { + sdata[threadIdx.x] += sdata[threadIdx.x + 4]; } __syncthreads(); - if(hipThreadIdx_x < 2) { - sdata[hipThreadIdx_x] += sdata[hipThreadIdx_x + 2]; + if(threadIdx.x < 2) { + sdata[threadIdx.x] += sdata[threadIdx.x + 2]; } __syncthreads(); - if(hipThreadIdx_x < 1) { - sdata[hipThreadIdx_x] += sdata[hipThreadIdx_x + 1]; + if(threadIdx.x < 1) { + sdata[threadIdx.x] += sdata[threadIdx.x + 1]; } - if(hipThreadIdx_x == 0) { - workspace[hipBlockIdx_x] = sdata[0]; + if(threadIdx.x == 0) { + workspace[blockIdx.x] = sdata[0]; } } @@ -161,49 +161,49 @@ __global__ void dot_reduction(double* __restrict__ workspace) { __shared__ double sdata[BLOCKSIZE]; - sdata[hipThreadIdx_x] = workspace[hipThreadIdx_x]; + sdata[threadIdx.x] = workspace[threadIdx.x]; __syncthreads(); - if(hipThreadIdx_x < 128) { - sdata[hipThreadIdx_x] += sdata[hipThreadIdx_x + 128]; + if(threadIdx.x < 128) { + sdata[threadIdx.x] += sdata[threadIdx.x + 128]; } __syncthreads(); - if(hipThreadIdx_x < 64) { - sdata[hipThreadIdx_x] += sdata[hipThreadIdx_x + 64]; + if(threadIdx.x < 64) { + sdata[threadIdx.x] += sdata[threadIdx.x + 64]; } __syncthreads(); - if(hipThreadIdx_x < 32) { - sdata[hipThreadIdx_x] += sdata[hipThreadIdx_x + 32]; + if(threadIdx.x < 32) { + sdata[threadIdx.x] += sdata[threadIdx.x + 32]; } __syncthreads(); - if(hipThreadIdx_x < 16) { - sdata[hipThreadIdx_x] += sdata[hipThreadIdx_x + 16]; + if(threadIdx.x < 16) { + sdata[threadIdx.x] += sdata[threadIdx.x + 16]; } __syncthreads(); - if(hipThreadIdx_x < 8) { - sdata[hipThreadIdx_x] += sdata[hipThreadIdx_x + 8]; + if(threadIdx.x < 8) { + sdata[threadIdx.x] += sdata[threadIdx.x + 8]; } __syncthreads(); - if(hipThreadIdx_x < 4) { - sdata[hipThreadIdx_x] += sdata[hipThreadIdx_x + 4]; + if(threadIdx.x < 4) { + sdata[threadIdx.x] += sdata[threadIdx.x + 4]; } __syncthreads(); - if(hipThreadIdx_x < 2) { - sdata[hipThreadIdx_x] += sdata[hipThreadIdx_x + 2]; + if(threadIdx.x < 2) { + sdata[threadIdx.x] += sdata[threadIdx.x + 2]; } __syncthreads(); - if(hipThreadIdx_x < 1) { - sdata[hipThreadIdx_x] += sdata[hipThreadIdx_x + 1]; + if(threadIdx.x < 1) { + sdata[threadIdx.x] += sdata[threadIdx.x + 1]; } - if(hipThreadIdx_x == 0) { + if(threadIdx.x == 0) { workspace[0] = sdata[0]; } diff --git a/tests/src/runtimeApi/stream/hipStreamGetCUMask.cpp b/tests/src/runtimeApi/stream/hipStreamGetCUMask.cpp index 1ce5c52bff..3468a1fa53 100644 --- a/tests/src/runtimeApi/stream/hipStreamGetCUMask.cpp +++ b/tests/src/runtimeApi/stream/hipStreamGetCUMask.cpp @@ -83,12 +83,12 @@ int main(int argc, char* argv[]) { defaultCUMask.push_back(temp); } - str_out = hipGetErrorString(hipExtStreamGetCUMask(0, cuMask.size(), 0)); + str_out = hipGetErrorName(hipExtStreamGetCUMask(0, cuMask.size(), 0)); if ((str_err.compare(str_out)) != 0) { failed("hipExtStreamGetCUMask returned wrong error code!"); } - str_out = hipGetErrorString(hipExtStreamGetCUMask(0, 0, &cuMask[0])); + str_out = hipGetErrorName(hipExtStreamGetCUMask(0, 0, &cuMask[0])); if ((str_err.compare(str_out)) != 0) { failed("hipExtStreamGetCUMask returned wrong error code!"); }