diff --git a/hipamd/CMakeLists.txt b/hipamd/CMakeLists.txt index f64ca45559..d1a2b133c5 100644 --- a/hipamd/CMakeLists.txt +++ b/hipamd/CMakeLists.txt @@ -377,6 +377,7 @@ if(HIP_PLATFORM STREQUAL "hcc") target_link_libraries(hip_hcc_static PRIVATE hc_am) add_library(hiprtc SHARED src/hiprtc.cpp src/code_object_bundle.cpp) + target_compile_options(hiprtc PRIVATE -DDISABLE_REDUCED_GPU_BLOB_COPY) set_property ( TARGET hiprtc PROPERTY VERSION "${HIP_LIB_VERSION_STRING}" ) set_property ( TARGET hiprtc PROPERTY SOVERSION "${HIP_LIB_VERSION_MAJOR}" ) diff --git a/hipamd/Jenkinsfile b/hipamd/Jenkinsfile index b8bd24cd74..e38f7824d2 100644 --- a/hipamd/Jenkinsfile +++ b/hipamd/Jenkinsfile @@ -177,7 +177,7 @@ def docker_build_inside_image( def build_image, String inside_args, String platf cd ${build_dir_rel} make install -j\$(nproc) make build_tests -i -j\$(nproc) - ctest -E "(hipMultiThreadDevice-pyramid|hipMemoryAllocateCoherentDriver)" + ctest --output-on-failure -E "(hipMultiThreadDevice-pyramid|hipMemoryAllocateCoherentDriver)" """ // If unit tests output a junit or xunit file in the future, jenkins can parse that file // to display test results on the dashboard @@ -295,13 +295,13 @@ def docker_upload_dockerhub( String local_org, String image_name, String remote_ String build_config = 'Release' String job_name = env.JOB_NAME.toLowerCase( ) -// The following launches 3 builds in parallel: rocm-head, rocm-3.0.x and cuda-10.x -parallel rocm_3_0: +// The following launches 3 builds in parallel: rocm-head, rocm-3.1.x and cuda-10.x +parallel rocm_3_1: { node('hip-rocm') { - String hcc_ver = 'rocm-3.0.x' - String from_image = 'ci_test_nodes/rocm-3.0.x/ubuntu-16.04:latest' + String hcc_ver = 'rocm-3.1.x' + String from_image = 'ci_test_nodes/rocm-3.1.x/ubuntu-16.04:latest' String inside_args = '--device=/dev/kfd --device=/dev/dri --group-add=video' // Checkout source code, dependencies and version files diff --git a/hipamd/bin/hipcc b/hipamd/bin/hipcc index 7eacc8f6b6..33a07c39dc 100755 --- a/hipamd/bin/hipcc +++ b/hipamd/bin/hipcc @@ -726,6 +726,12 @@ if ($HIP_PLATFORM eq "clang") { $HIPCXXFLAGS .= " -O3"; $HIPLDFLAGS .= " -O3"; } + if ($optArg ne "-O0") { + $HIPCXXFLAGS .= " -mllvm -amdgpu-early-inline-all=true -mllvm -amdgpu-function-calls=false"; + if ($needLDFLAGS and not $needCXXFLAGS) { + $HIPLDFLAGS .= " -mllvm -amdgpu-early-inline-all=true -mllvm -amdgpu-function-calls=false"; + } + } $HIP_DEVLIB_FLAGS = " --hip-device-lib-path=$DEVICE_LIB_PATH"; $HIPCXXFLAGS .= " $HIP_DEVLIB_FLAGS"; if (not $isWindows) { diff --git a/hipamd/bin/hipconfig b/hipamd/bin/hipconfig index a73e8af8b9..c56b56ecd8 100755 --- a/hipamd/bin/hipconfig +++ b/hipamd/bin/hipconfig @@ -1,7 +1,7 @@ #!/usr/bin/perl -w $HIP_BASE_VERSION_MAJOR = "3"; -$HIP_BASE_VERSION_MINOR = "1"; +$HIP_BASE_VERSION_MINOR = "2"; # Need perl > 5.10 to use logic-defined or use 5.006; use v5.10.1; diff --git a/hipamd/bin/hipify-perl b/hipamd/bin/hipify-perl index 4fe9b80349..f6de5abae4 100755 --- a/hipamd/bin/hipify-perl +++ b/hipamd/bin/hipify-perl @@ -1172,6 +1172,7 @@ sub simpleSubstitutions { $ft{'numeric_literal'} += s/\bCUDA_ERROR_ASSERT\b/hipErrorAssert/g; $ft{'numeric_literal'} += s/\bCUDA_ERROR_CONTEXT_ALREADY_CURRENT\b/hipErrorContextAlreadyCurrent/g; $ft{'numeric_literal'} += s/\bCUDA_ERROR_CONTEXT_ALREADY_IN_USE\b/hipErrorContextAlreadyInUse/g; + $ft{'numeric_literal'} += s/\bCUDA_ERROR_COOPERATIVE_LAUNCH_TOO_LARGE\b/hipErrorCooperativeLaunchTooLarge/g; $ft{'numeric_literal'} += s/\bCUDA_ERROR_DEINITIALIZED\b/hipErrorDeinitialized/g; $ft{'numeric_literal'} += s/\bCUDA_ERROR_ECC_UNCORRECTABLE\b/hipErrorECCNotCorrectable/g; $ft{'numeric_literal'} += s/\bCUDA_ERROR_FILE_NOT_FOUND\b/hipErrorFileNotFound/g; @@ -1528,6 +1529,7 @@ sub simpleSubstitutions { $ft{'numeric_literal'} += s/\bCU_MEMORYTYPE_UNIFIED\b/hipMemoryTypeUnified/g; $ft{'numeric_literal'} += s/\bCU_MEM_ATTACH_GLOBAL\b/hipMemAttachGlobal/g; $ft{'numeric_literal'} += s/\bCU_MEM_ATTACH_HOST\b/hipMemAttachHost/g; + $ft{'numeric_literal'} += s/\bCU_OCCUPANCY_DEFAULT\b/hipOccupancyDefault/g; $ft{'numeric_literal'} += s/\bCU_RESOURCE_TYPE_ARRAY\b/hipResourceTypeArray/g; $ft{'numeric_literal'} += s/\bCU_RESOURCE_TYPE_LINEAR\b/hipResourceTypeLinear/g; $ft{'numeric_literal'} += s/\bCU_RESOURCE_TYPE_MIPMAPPED_ARRAY\b/hipResourceTypeMipmappedArray/g; @@ -1636,6 +1638,7 @@ sub simpleSubstitutions { $ft{'numeric_literal'} += s/\bcudaErrorAlreadyMapped\b/hipErrorAlreadyMapped/g; $ft{'numeric_literal'} += s/\bcudaErrorArrayIsMapped\b/hipErrorArrayIsMapped/g; $ft{'numeric_literal'} += s/\bcudaErrorAssert\b/hipErrorAssert/g; + $ft{'numeric_literal'} += s/\bcudaErrorCooperativeLaunchTooLarge\b/hipErrorCooperativeLaunchTooLarge/g; $ft{'numeric_literal'} += s/\bcudaErrorCudartUnloading\b/hipErrorDeinitialized/g; $ft{'numeric_literal'} += s/\bcudaErrorDeviceAlreadyInUse\b/hipErrorContextAlreadyInUse/g; $ft{'numeric_literal'} += s/\bcudaErrorDeviceUninitialized\b/hipErrorInvalidContext/g; @@ -1749,6 +1752,8 @@ sub simpleSubstitutions { $ft{'define'} += s/\bCUDA_ARRAY3D_LAYERED\b/hipArrayLayered/g; $ft{'define'} += s/\bCUDA_ARRAY3D_SURFACE_LDST\b/hipArraySurfaceLoadStore/g; $ft{'define'} += s/\bCUDA_ARRAY3D_TEXTURE_GATHER\b/hipArrayTextureGather/g; + $ft{'define'} += s/\bCUDA_COOPERATIVE_LAUNCH_MULTI_DEVICE_NO_POST_LAUNCH_SYNC\b/hipCooperativeLaunchMultiDeviceNoPostSync/g; + $ft{'define'} += s/\bCUDA_COOPERATIVE_LAUNCH_MULTI_DEVICE_NO_PRE_LAUNCH_SYNC\b/hipCooperativeLaunchMultiDeviceNoPreSync/g; $ft{'define'} += s/\bCU_LAUNCH_PARAM_BUFFER_POINTER\b/HIP_LAUNCH_PARAM_BUFFER_POINTER/g; $ft{'define'} += s/\bCU_LAUNCH_PARAM_BUFFER_SIZE\b/HIP_LAUNCH_PARAM_BUFFER_SIZE/g; $ft{'define'} += s/\bCU_LAUNCH_PARAM_END\b/HIP_LAUNCH_PARAM_END/g; @@ -1769,6 +1774,8 @@ sub simpleSubstitutions { $ft{'define'} += s/\bcudaArrayLayered\b/hipArrayLayered/g; $ft{'define'} += s/\bcudaArraySurfaceLoadStore\b/hipArraySurfaceLoadStore/g; $ft{'define'} += s/\bcudaArrayTextureGather\b/hipArrayTextureGather/g; + $ft{'define'} += s/\bcudaCooperativeLaunchMultiDeviceNoPostSync\b/hipCooperativeLaunchMultiDeviceNoPostSync/g; + $ft{'define'} += s/\bcudaCooperativeLaunchMultiDeviceNoPreSync\b/hipCooperativeLaunchMultiDeviceNoPreSync/g; $ft{'define'} += s/\bcudaDeviceBlockingSync\b/hipDeviceScheduleBlockingSync/g; $ft{'define'} += s/\bcudaDeviceLmemResizeToMax\b/hipDeviceLmemResizeToMax/g; $ft{'define'} += s/\bcudaDeviceMapHost\b/hipDeviceMapHost/g; @@ -1792,6 +1799,7 @@ sub simpleSubstitutions { $ft{'define'} += s/\bcudaIpcMemLazyEnablePeerAccess\b/hipIpcMemLazyEnablePeerAccess/g; $ft{'define'} += s/\bcudaMemAttachGlobal\b/hipMemAttachGlobal/g; $ft{'define'} += s/\bcudaMemAttachHost\b/hipMemAttachHost/g; + $ft{'define'} += s/\bcudaOccupancyDefault\b/hipOccupancyDefault/g; $ft{'define'} += s/\bcudaStreamDefault\b/hipStreamDefault/g; $ft{'define'} += s/\bcudaStreamNonBlocking\b/hipStreamNonBlocking/g; $ft{'define'} += s/\bcudaTextureType1D\b/hipTextureType1D/g; diff --git a/hipamd/docs/markdown/CUDA_Driver_API_functions_supported_by_HIP.md b/hipamd/docs/markdown/CUDA_Driver_API_functions_supported_by_HIP.md index ff14cf9bc1..e77997b0b6 100644 --- a/hipamd/docs/markdown/CUDA_Driver_API_functions_supported_by_HIP.md +++ b/hipamd/docs/markdown/CUDA_Driver_API_functions_supported_by_HIP.md @@ -447,6 +447,7 @@ | 717 |*`CUDA_ERROR_INVALID_ADDRESS_SPACE`* | | | 718 |*`CUDA_ERROR_INVALID_PC`* | | | 719 |*`CUDA_ERROR_LAUNCH_FAILED`* |*`hipErrorLaunchFailure`* | +| 720 |*`CUDA_ERROR_COOPERATIVE_LAUNCH_TOO_LARGE`* |*`hipErrorCooperativeLaunchTooLarge`* | | 800 |*`CUDA_ERROR_NOT_PERMITTED`* | | | 801 |*`CUDA_ERROR_NOT_SUPPORTED`* |*`hipErrorNotSupported`* | | 802 |*`CUDA_ERROR_SYSTEM_NOT_READY`* | | 10.0 | @@ -1125,10 +1126,10 @@ | **CUDA** | **HIP** |**CUDA version\***| |-----------------------------------------------------------|---------------------------------------------------------|------------------| -| `cuOccupancyMaxActiveBlocksPerMultiprocessor` |`hipOccupancyMaxActiveBlocksPerMultiprocessor` | -| `cuOccupancyMaxActiveBlocksPerMultiprocessorWithFlags` |`hipOccupancyMaxActiveBlocksPerMultiprocessorWithFlags` | -| `cuOccupancyMaxPotentialBlockSize` |`hipOccupancyMaxPotentialBlockSize` | -| `cuOccupancyMaxPotentialBlockSizeWithFlags` | | +| `cuOccupancyMaxActiveBlocksPerMultiprocessor` |`hipDrvOccupancyMaxActiveBlocksPerMultiprocessor` | +| `cuOccupancyMaxActiveBlocksPerMultiprocessorWithFlags` |`hipDrvOccupancyMaxActiveBlocksPerMultiprocessorWithFlags` | +| `cuOccupancyMaxPotentialBlockSize` |`hipOccupancyMaxPotentialBlockSize` | +| `cuOccupancyMaxPotentialBlockSizeWithFlags` | | ## **22. Texture Reference Management [DEPRECATED]** diff --git a/hipamd/docs/markdown/CUDA_Runtime_API_functions_supported_by_HIP.md b/hipamd/docs/markdown/CUDA_Runtime_API_functions_supported_by_HIP.md index 201d2aad0f..54e0c89e06 100644 --- a/hipamd/docs/markdown/CUDA_Runtime_API_functions_supported_by_HIP.md +++ b/hipamd/docs/markdown/CUDA_Runtime_API_functions_supported_by_HIP.md @@ -800,7 +800,7 @@ | 717 |*`cudaErrorInvalidAddressSpace`* | | | | 718 |*`cudaErrorInvalidPc`* | | | | 719 |*`cudaErrorLaunchFailure`* | |*`hipErrorLaunchFailure`* | -| 720 |*`cudaErrorCooperativeLaunchTooLarge`* | 9.0 | | +| 720 |*`cudaErrorCooperativeLaunchTooLarge`* | 9.0 |*`hipErrorCooperativeLaunchTooLarge`* | | 800 |*`cudaErrorNotPermitted`* | | | | 801 |*`cudaErrorNotSupported`* | |*`hipErrorNotSupported`* | | 802 |*`cudaErrorSystemNotReady`* | 10.0 | | @@ -1045,7 +1045,7 @@ | define |`cudaMemAttachGlobal` | |`hipMemAttachGlobal` | | define |`cudaMemAttachHost` | |`hipMemAttachHost` | | define |`cudaMemAttachSingle` | | | -| define |`cudaOccupancyDefault` | | | +| define |`cudaOccupancyDefault` | |`hipOccupancyDefault` | | define |`cudaOccupancyDisableCachingOverride` | | | | define |`cudaPeerAccessDefault` | | | | define |`cudaStreamDefault` | |`hipStreamDefault` | diff --git a/hipamd/docs/markdown/hip_tracing.md b/hipamd/docs/markdown/hip_tracing.md new file mode 100644 index 0000000000..40513f4e3c --- /dev/null +++ b/hipamd/docs/markdown/hip_tracing.md @@ -0,0 +1,72 @@ +# Profiling HIP Code + +This section describes the tracing and debugging capabilities that HIP provides. + + +- [Tracing and Debug](#tracing-and-debug) + * [Tracing HIP APIs](#tracing-hip-apis) + + [Color](#color) + + + +## Tracing and Debug + +### Tracing HIP APIs +The HIP runtime can print the HIP function strings to stderr using HIP_TRACE_API environment variable. +The trace prints two messages for each API - one at the beginning of the API call (line starts with "<<") and one at the end of the API call (line ends with ">>"). +Here's an example for one API followed by a description for the sections of the trace: + +``` +<> +``` + +- `<> +info: running on device gfx803 +info: allocate host mem ( 7.63 MB) +info: allocate device mem ( 7.63 MB) +<> +<> +info: copy Host2Device +<> +info: launch 'vector_square' kernel +1.5 hipLaunchKernel 'HIP_KERNEL_NAME(vector_square)' gridDim:{512,1,1} groupDim:{256,1,1} sharedMem:+0 stream#0.0 +info: copy Device2Host +<> +info: check result +PASSED! +``` + +HIP_TRACE_API supports multiple levels of debug information: + - 0x1 = print all HIP APIs. This is the most verbose setting; the flags below allow selecting a subset. + - 0x2 = print HIP APIs which initiate GPU kernel commands. Includes hipLaunchKernel, hipLaunchModuleKernel + - 0x4 = print HIP APIs which initiate GPU memory commands. Includes hipMemcpy*, hipMemset*. + - 0x8 = print HIP APIs which allocate or free memory. Includes hipMalloc, hipHostMalloc, hipFree, hipHostFree. + +These can be combined. For example, HIP_TRACE_API=6 shows a concise view of the HIP commands (both kernel and memory) that are sent to the GPU. + + +#### Color +Note this trace mode uses colors. "less -r" can handle raw control characters and will display the debug output in proper colors. +You can change the color used for the trace mode with the HIP_TRACE_API_COLOR environment variable. Possible values are None/Red/Green/Yellow/Blue/Magenta/Cyan/White. +None will disable use of color control codes for both the opening and closing and may be useful when saving the trace file or when a pure text trace is desired. + + + diff --git a/hipamd/hip_prof_gen.py b/hipamd/hip_prof_gen.py index 9e90c1558c..d2da7cd4df 100755 --- a/hipamd/hip_prof_gen.py +++ b/hipamd/hip_prof_gen.py @@ -447,7 +447,8 @@ if len(sys.argv) > 3: OUTPUT = sys.argv[3] # API declaration map api_map = { - 'hipHccModuleLaunchKernel': '' + 'hipHccModuleLaunchKernel': '', + 'hipExtModuleLaunchKernel': '' } # API options map opts_map = {} diff --git a/hipamd/hipify-clang/README.md b/hipamd/hipify-clang/README.md index 86a97dbe29..88d7a72ccd 100644 --- a/hipamd/hipify-clang/README.md +++ b/hipamd/hipify-clang/README.md @@ -42,10 +42,10 @@ After applying all the matchers, the output HIP source is produced. `hipify-clang` requires: -1. [**LLVM+CLANG**](http://releases.llvm.org) of at least version [3.8.0](http://releases.llvm.org/download.html#3.8.0); the latest stable and recommended release: [**9.0.1**](http://releases.llvm.org/download.html#9.0.1), the latest release candidate: [10.0.0-rc1](https://github.com/llvm/llvm-project/releases/tag/llvmorg-10.0.0-rc1). +1. [**LLVM+CLANG**](http://releases.llvm.org) of at least version [3.8.0](http://releases.llvm.org/download.html#3.8.0); the latest stable and recommended release: [**9.0.1**](http://releases.llvm.org/download.html#9.0.1), the latest release candidate: [10.0.0-rc3](https://github.com/llvm/llvm-project/releases/tag/llvmorg-10.0.0-rc3). 2. [**CUDA**](https://developer.nvidia.com/cuda-downloads) of at least version [7.0](https://developer.nvidia.com/cuda-toolkit-70), the latest supported version is [**10.1 Update 2**](https://developer.nvidia.com/cuda-10.1-download-archive-base). -To use the latest CUDA version [10.2](https://developer.nvidia.com/cuda-downloads) please use the latest `LLVM` release candidate: [10.0.0-rc1](https://github.com/llvm/llvm-project/releases/tag/llvmorg-10.0.0-rc1). +To use the latest CUDA version [10.2](https://developer.nvidia.com/cuda-downloads) please use the latest `LLVM` release candidate: [10.0.0-rc3](https://github.com/llvm/llvm-project/releases/tag/llvmorg-10.0.0-rc3). | **LLVM release version** | **CUDA latest supported version** | **Windows** | **Linux** | |:----------------------------------------------------------:|:------------------------------------------------------------------------:|:-----------:|:---------:| @@ -67,7 +67,7 @@ To use the latest CUDA version [10.2](https://developer.nvidia.com/cuda-download | [8.0.1](http://releases.llvm.org/download.html#8.0.1) | [10.0](https://developer.nvidia.com/cuda-10.0-download-archive) | -
not working due to
the clang's bug [38811](https://bugs.llvm.org/show_bug.cgi?id=38811)
+
[patch](patches/patch_for_clang_8.0.1_bug_38811.zip)*
| + | | [9.0.0](http://releases.llvm.org/download.html#9.0.0) | [10.1](https://developer.nvidia.com/cuda-10.1-download-archive-base) | + | + | | [**9.0.1**](http://releases.llvm.org/download.html#9.0.1) | [**10.1**](https://developer.nvidia.com/cuda-10.1-download-archive-base) | +
**LATEST STABLE RELEASE** | +
**LATEST STABLE RELEASE** | -| [10.0.0-rc1](https://github.com/llvm/llvm-project/releases/tag/llvmorg-10.0.0-rc1) | [10.2](https://developer.nvidia.com/cuda-downloads) | + | + | +| [10.0.0-rc3](https://github.com/llvm/llvm-project/releases/tag/llvmorg-10.0.0-rc3) | [10.2](https://developer.nvidia.com/cuda-downloads) | + | + | `*` Download the patch and unpack it into your `LLVM` distributive directory; a few header files will be overwritten; rebuilding of `LLVM` is not needed. @@ -158,7 +158,7 @@ Run `Visual Studio 16 2019`, open the generated `LLVM.sln`, build all, build pro **LLVM 10.0.0 or newer:** -1. download [`LLVM project`](https://github.com/llvm/llvm-project/archive/llvmorg-10.0.0-rc1.tar.gz) sources; +1. download [`LLVM project`](https://github.com/llvm/llvm-project/archive/llvmorg-10.0.0-rc3.tar.gz) sources; 2. build [`LLVM project`](http://llvm.org/docs/CMake.html): **Linux**: @@ -168,6 +168,7 @@ Run `Visual Studio 16 2019`, open the generated `LLVM.sln`, build all, build pro -DLLVM_SOURCE_DIR=../llvm-project \ -DLLVM_TARGETS_TO_BUILD="X86;NVPTX" \ -DLLVM_ENABLE_PROJECTS="clang" \ + -DLLVM_TEMPORARILY_ALLOW_OLD_TOOLCHAIN=ON -DCMAKE_BUILD_TYPE=Release \ ../llvm-project/llvm make -j install @@ -181,6 +182,7 @@ Run `Visual Studio 16 2019`, open the generated `LLVM.sln`, build all, build pro -DLLVM_SOURCE_DIR=../llvm-project \ -DLLVM_TARGETS_TO_BUILD="NVPTX" \ -DLLVM_ENABLE_PROJECTS="clang" \ + -DLLVM_TEMPORARILY_ALLOW_OLD_TOOLCHAIN=ON -DCMAKE_BUILD_TYPE=Release \ -Thost=x64 \ ../llvm-project/llvm @@ -247,7 +249,7 @@ On Linux the following configurations are tested: Ubuntu 14: LLVM 5.0.0 - 6.0.1, CUDA 7.0 - 9.0, cudnn-5.0.5 - cudnn-7.6.5.32 -Ubuntu 16-18: LLVM 8.0.0 - 10.0.0-rc1, CUDA 8.0 - 10.2, cudnn-5.1.10 - cudnn-7.6.5.32 +Ubuntu 16-18: LLVM 8.0.0 - 10.0.0-rc3, CUDA 8.0 - 10.2, cudnn-5.1.10 - cudnn-7.6.5.32 Minimum build system requirements for the above configurations: @@ -393,19 +395,19 @@ Testing Time: 3.07s ``` ### hipify-clang: Windows -On Windows 10 the following configurations are tested: +*Tested configurations:* -LLVM 5.0.0 - 5.0.2, CUDA 8.0, cudnn 5.1.10 - 7.1.4.18 +| **LLVM** | **CUDA** | **cuDNN** | **Visual Studio** | **cmake** | **Python** | +|:--------------:|---------:|--------------------:|--------------------------:|----------:|-----------:| +| 5.0.0 - 5.0.2 | 8.0 | 5.1.10 - 7.1.4.18 | 2017.15.5.2 | 3.5.1 | 3.6.4 | +| 6.0.0 - 6.0.1 | 9.0 | 7.0.5.15 - 7.6.5.32 | 2017.15.5.5 | 3.6.0 | 3.7.2 | +| 7.0.0 - 7.1.0 | 9.2 | 7.6.5.32 | 2017.15.9.11 | 3.13.3 | 3.7.3 | +| 8.0.0 - 8.0.1 | 10.0 | 7.6.5.32 | 2017.15.9.15 | 3.14.2 | 3.7.4 | +| 9.0.0 - 9.0.1 | 10.1 | 7.6.5.32 | 2017.15.9.20, 2019.16.4.5 | 3.16.4 | 3.8.0 | +| 10.0.0-rc1-rc3 | 10.2 | 7.6.5.32 | 2017.15.9.20, 2019.16.4.5 | 3.16.4 | 3.8.1 | +| 11.0.0git | 10.2 | 7.6.5.32 | 2017.15.9.20, 2019.16.4.5 | 3.16.5 | 3.8.2 | -LLVM 6.0.0 - 6.0.1, CUDA 9.0, cudnn 7.0.5.15 - 7.6.5.32 - -LLVM 7.0.0 - 10.0.0-rc1, CUDA 7.5 - 10.2, cudnn 7.0.5.15 - 7.6.5.32 - -Build system requirements for the latest stable configuration LLVM 9.0.1/CUDA 10.1 Update 2: - -Python 3.6.0 - 3.8.1, cmake 3.5.1 - 3.16.3, Visual Studio 2017 (15.5.2) - 2019 (16.4.4). - -Here is an example of building `hipify-clang` with testing support on `Windows 10` by `Visual Studio 16 2019`: +*Building with testing support on `Windows 10` by `Visual Studio 16 2019`:* ```shell cmake @@ -429,7 +431,7 @@ cmake -- - CMake module path: F:/LLVM/9.0.1/dist/lib/cmake/llvm -- - Include path : F:/LLVM/9.0.1/dist/include -- - Binary path : F:/LLVM/9.0.1/dist/bin --- Found PythonInterp: C:/Program Files/Python38/python.exe (found suitable version "3.8.1", minimum required is "3.6") +-- Found PythonInterp: C:/Program Files/Python38/python.exe (found suitable version "3.8.2", minimum required is "3.6") -- Found lit: C:/Program Files/Python38/Scripts/lit.exe -- Found FileCheck: F:/LLVM/9.0.1/dist/bin/FileCheck.exe -- Found CUDA: C:/Program Files/NVIDIA GPU Computing Toolkit/CUDA/v10.1 (found version "10.1") diff --git a/hipamd/hipify-clang/src/CUDA2HIP_Driver_API_functions.cpp b/hipamd/hipify-clang/src/CUDA2HIP_Driver_API_functions.cpp index 5d700631c4..ab07a10e93 100644 --- a/hipamd/hipify-clang/src/CUDA2HIP_Driver_API_functions.cpp +++ b/hipamd/hipify-clang/src/CUDA2HIP_Driver_API_functions.cpp @@ -545,9 +545,9 @@ const std::map CUDA_DRIVER_FUNCTION_MAP{ // 5.21. Occupancy // cudaOccupancyMaxActiveBlocksPerMultiprocessor - {"cuOccupancyMaxActiveBlocksPerMultiprocessor", {"hipOccupancyMaxActiveBlocksPerMultiprocessor", "", CONV_OCCUPANCY, API_DRIVER}}, + {"cuOccupancyMaxActiveBlocksPerMultiprocessor", {"hipDrvOccupancyMaxActiveBlocksPerMultiprocessor", "", CONV_OCCUPANCY, API_DRIVER}}, // cudaOccupancyMaxActiveBlocksPerMultiprocessorWithFlags - {"cuOccupancyMaxActiveBlocksPerMultiprocessorWithFlags", {"hipOccupancyMaxActiveBlocksPerMultiprocessorWithFlags", "", CONV_OCCUPANCY, API_DRIVER}}, + {"cuOccupancyMaxActiveBlocksPerMultiprocessorWithFlags", {"hipDrvOccupancyMaxActiveBlocksPerMultiprocessorWithFlags", "", CONV_OCCUPANCY, API_DRIVER}}, // cudaOccupancyMaxPotentialBlockSize {"cuOccupancyMaxPotentialBlockSize", {"hipOccupancyMaxPotentialBlockSize", "", CONV_OCCUPANCY, API_DRIVER}}, // cudaOccupancyMaxPotentialBlockSizeWithFlags diff --git a/hipamd/hipify-clang/src/CUDA2HIP_Driver_API_types.cpp b/hipamd/hipify-clang/src/CUDA2HIP_Driver_API_types.cpp index 1c3c2634f7..054de19800 100644 --- a/hipamd/hipify-clang/src/CUDA2HIP_Driver_API_types.cpp +++ b/hipamd/hipify-clang/src/CUDA2HIP_Driver_API_types.cpp @@ -1047,7 +1047,7 @@ const std::map CUDA_DRIVER_TYPE_NAME_MAP{ {"CUoccupancy_flags_enum", {"hipOccupancyFlags", "", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}}, // CUoccupancy_flags enum values // cudaOccupancyDefault - {"CU_OCCUPANCY_DEFAULT", {"hipOccupancyDefault", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 0x00 + {"CU_OCCUPANCY_DEFAULT", {"hipOccupancyDefault", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 0x00 // cudaOccupancyDisableCachingOverride {"CU_OCCUPANCY_DISABLE_CACHING_OVERRIDE", {"hipOccupancyDisableCachingOverride", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 0x01 @@ -1284,6 +1284,8 @@ const std::map CUDA_DRIVER_TYPE_NAME_MAP{ {"CUDA_ERROR_INVALID_PC", {"hipErrorInvalidPc", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 718 // cudaErrorLaunchFailure {"CUDA_ERROR_LAUNCH_FAILED", {"hipErrorLaunchFailure", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 719 + // cudaErrorCooperativeLaunchTooLarge + {"CUDA_ERROR_COOPERATIVE_LAUNCH_TOO_LARGE", {"hipErrorCooperativeLaunchTooLarge", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 720 // cudaErrorNotPermitted {"CUDA_ERROR_NOT_PERMITTED", {"hipErrorNotPermitted", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 800 // cudaErrorNotSupported @@ -1598,9 +1600,9 @@ const std::map CUDA_DRIVER_TYPE_NAME_MAP{ // cudaArrayColorAttachment {"CUDA_ARRAY3D_COLOR_ATTACHMENT", {"hipArrayColorAttachment", "", CONV_DEFINE, API_DRIVER, HIP_UNSUPPORTED}}, // 0x20 // cudaCooperativeLaunchMultiDeviceNoPreSync - {"CUDA_COOPERATIVE_LAUNCH_MULTI_DEVICE_NO_PRE_LAUNCH_SYNC", {"hipCooperativeLaunchMultiDeviceNoPreSync", "", CONV_DEFINE, API_DRIVER, HIP_UNSUPPORTED}}, // 0x01 + {"CUDA_COOPERATIVE_LAUNCH_MULTI_DEVICE_NO_PRE_LAUNCH_SYNC", {"hipCooperativeLaunchMultiDeviceNoPreSync", "", CONV_DEFINE, API_DRIVER}}, // 0x01 // cudaCooperativeLaunchMultiDeviceNoPostSync - {"CUDA_COOPERATIVE_LAUNCH_MULTI_DEVICE_NO_POST_LAUNCH_SYNC", {"hipCooperativeLaunchMultiDeviceNoPostSync", "", CONV_DEFINE, API_DRIVER, HIP_UNSUPPORTED}}, // 0x02 + {"CUDA_COOPERATIVE_LAUNCH_MULTI_DEVICE_NO_POST_LAUNCH_SYNC", {"hipCooperativeLaunchMultiDeviceNoPostSync", "", CONV_DEFINE, API_DRIVER}}, // 0x02 // cudaExternalMemoryDedicated {"CUDA_EXTERNAL_MEMORY_DEDICATED", {"hipExternalMemoryDedicated", "", CONV_DEFINE, API_DRIVER, HIP_UNSUPPORTED}}, // 0x1 // cudaExternalSemaphoreSignalSkipNvSciBufMemSync diff --git a/hipamd/hipify-clang/src/CUDA2HIP_Runtime_API_types.cpp b/hipamd/hipify-clang/src/CUDA2HIP_Runtime_API_types.cpp index 5993f5d770..6eb9bfb2be 100644 --- a/hipamd/hipify-clang/src/CUDA2HIP_Runtime_API_types.cpp +++ b/hipamd/hipify-clang/src/CUDA2HIP_Runtime_API_types.cpp @@ -787,8 +787,8 @@ const std::map CUDA_RUNTIME_TYPE_NAME_MAP { {"cudaErrorInvalidPc", {"hipErrorInvalidPc", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 718 // CUDA_ERROR_LAUNCH_FAILED {"cudaErrorLaunchFailure", {"hipErrorLaunchFailure", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 719 - // no analogue - {"cudaErrorCooperativeLaunchTooLarge", {"hipErrorCooperativeLaunchTooLarge", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 720 + // CUDA_ERROR_COOPERATIVE_LAUNCH_TOO_LARGE + {"cudaErrorCooperativeLaunchTooLarge", {"hipErrorCooperativeLaunchTooLarge", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 720 // CUDA_ERROR_NOT_PERMITTED {"cudaErrorNotPermitted", {"hipErrorNotPermitted", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 800 // CUDA_ERROR_NOT_SUPPORTED @@ -1327,9 +1327,9 @@ const std::map CUDA_RUNTIME_TYPE_NAME_MAP { // CUDA_ARRAY3D_COLOR_ATTACHMENT {"cudaArrayColorAttachment", {"hipArrayColorAttachment", "", CONV_DEFINE, API_RUNTIME, HIP_UNSUPPORTED}}, // 0x20 // CUDA_COOPERATIVE_LAUNCH_MULTI_DEVICE_NO_PRE_LAUNCH_SYNC - {"cudaCooperativeLaunchMultiDeviceNoPreSync", {"hipCooperativeLaunchMultiDeviceNoPreSync", "", CONV_DEFINE, API_RUNTIME, HIP_UNSUPPORTED}}, // 0x01 + {"cudaCooperativeLaunchMultiDeviceNoPreSync", {"hipCooperativeLaunchMultiDeviceNoPreSync", "", CONV_DEFINE, API_RUNTIME}}, // 0x01 // CUDA_COOPERATIVE_LAUNCH_MULTI_DEVICE_NO_POST_LAUNCH_SYNC - {"cudaCooperativeLaunchMultiDeviceNoPostSync", {"hipCooperativeLaunchMultiDeviceNoPostSync", "", CONV_DEFINE, API_RUNTIME, HIP_UNSUPPORTED}}, // 0x02 + {"cudaCooperativeLaunchMultiDeviceNoPostSync", {"hipCooperativeLaunchMultiDeviceNoPostSync", "", CONV_DEFINE, API_RUNTIME}}, // 0x02 // CU_DEVICE_CPU ((CUdevice)-1) {"cudaCpuDeviceId", {"hipCpuDeviceId", "", CONV_DEFINE, API_RUNTIME, HIP_UNSUPPORTED}}, // ((int)-1) // CU_DEVICE_INVALID ((CUdevice)-2) @@ -1412,7 +1412,7 @@ const std::map CUDA_RUNTIME_TYPE_NAME_MAP { // no analogue {"cudaTextureTypeCubemapLayered", {"hipTextureTypeCubemapLayered", "", CONV_DEFINE, API_RUNTIME}}, // 0xFC // CU_OCCUPANCY_DEFAULT - {"cudaOccupancyDefault", {"hipOccupancyDefault", "", CONV_DEFINE, API_RUNTIME, HIP_UNSUPPORTED}}, // 0x00 + {"cudaOccupancyDefault", {"hipOccupancyDefault", "", CONV_DEFINE, API_RUNTIME}}, // 0x00 // CU_OCCUPANCY_DISABLE_CACHING_OVERRIDE {"cudaOccupancyDisableCachingOverride", {"hipOccupancyDisableCachingOverride", "", CONV_DEFINE, API_RUNTIME, HIP_UNSUPPORTED}}, // 0x01 // CU_STREAM_DEFAULT diff --git a/hipamd/include/hip/hcc_detail/code_object_bundle.hpp b/hipamd/include/hip/hcc_detail/code_object_bundle.hpp index f312d2e79b..77e0d706d6 100644 --- a/hipamd/include/hip/hcc_detail/code_object_bundle.hpp +++ b/hipamd/include/hip/hcc_detail/code_object_bundle.hpp @@ -31,9 +31,11 @@ THE SOFTWARE. #include #include #include - +#include namespace hip_impl { - +#if !defined(DISABLE_REDUCED_GPU_BLOB_COPY) +std::unordered_set& get_all_gpuarch(); +#endif inline std::string transmogrify_triple(const std::string& triple) { @@ -43,7 +45,6 @@ std::string transmogrify_triple(const std::string& triple) if (triple.find(old_prefix) == 0) { return new_prefix + triple.substr(sizeof(old_prefix) - 1); } - return (triple.find(new_prefix) == 0) ? triple : ""; } @@ -114,9 +115,7 @@ class Bundled_code_header { friend inline bool read(RandomAccessIterator f, RandomAccessIterator l, Bundled_code_header& x) { if (f == l) return false; - std::copy_n(f, sizeof(x.header_.cbuf_), x.header_.cbuf_); - if (valid(x)) { x.bundles_.resize(x.header_.bundle_cnt_); @@ -126,11 +125,16 @@ class Bundled_code_header { it += sizeof(y.header.cbuf); y.triple.assign(it, it + y.header.triple_sz); - + #ifdef DISABLE_REDUCED_GPU_BLOB_COPY std::copy_n(f + y.header.offset, y.header.bundle_sz, std::back_inserter(y.blob)); - + #else + auto& gpuArch = get_all_gpuarch(); + auto itgpuArch = std::find(gpuArch.begin(),gpuArch.end(),y.triple); + if (itgpuArch != gpuArch.end()){ + std::copy_n(f + y.header.offset, y.header.bundle_sz, std::back_inserter(y.blob)); + } + #endif it += y.header.triple_sz; - x.bundled_code_size = std::max(x.bundled_code_size, y.header.offset + y.header.bundle_sz); } diff --git a/hipamd/include/hip/hcc_detail/device_functions.h b/hipamd/include/hip/hcc_detail/device_functions.h index 68e3277270..7096841da8 100644 --- a/hipamd/include/hip/hcc_detail/device_functions.h +++ b/hipamd/include/hip/hcc_detail/device_functions.h @@ -319,6 +319,41 @@ double __shfl(double var, int src_lane, int width = warpSize) { double tmp1; __builtin_memcpy(&tmp1, &tmp0, sizeof(tmp0)); return tmp1; } +__device__ +inline +long __shfl(long var, int src_lane, int width = warpSize) +{ + #ifndef _MSC_VER + static_assert(sizeof(long) == 2 * sizeof(int), ""); + static_assert(sizeof(long) == sizeof(uint64_t), ""); + + int tmp[2]; __builtin_memcpy(tmp, &var, sizeof(tmp)); + tmp[0] = __shfl(tmp[0], src_lane, width); + tmp[1] = __shfl(tmp[1], src_lane, width); + + uint64_t tmp0 = (static_cast(tmp[1]) << 32ull) | static_cast(tmp[0]); + long tmp1; __builtin_memcpy(&tmp1, &tmp0, sizeof(tmp0)); + return tmp1; + #else + static_assert(sizeof(long) == sizeof(int), ""); + return static_cast(__shfl(static_cast(var), src_lane, width)); + #endif +} +__device__ +inline +long long __shfl(long long var, int src_lane, int width = warpSize) +{ + static_assert(sizeof(long long) == 2 * sizeof(int), ""); + static_assert(sizeof(long long) == sizeof(uint64_t), ""); + + int tmp[2]; __builtin_memcpy(tmp, &var, sizeof(tmp)); + tmp[0] = __shfl(tmp[0], src_lane, width); + tmp[1] = __shfl(tmp[1], src_lane, width); + + uint64_t tmp0 = (static_cast(tmp[1]) << 32ull) | static_cast(tmp[0]); + long long tmp1; __builtin_memcpy(&tmp1, &tmp0, sizeof(tmp0)); + return tmp1; +} __device__ inline @@ -356,6 +391,39 @@ double __shfl_up(double var, unsigned int lane_delta, int width = warpSize) { double tmp1; __builtin_memcpy(&tmp1, &tmp0, sizeof(tmp0)); return tmp1; } +__device__ +inline +long __shfl_up(long var, unsigned int lane_delta, int width = warpSize) +{ + #ifndef _MSC_VER + static_assert(sizeof(long) == 2 * sizeof(int), ""); + static_assert(sizeof(long) == sizeof(uint64_t), ""); + + int tmp[2]; __builtin_memcpy(tmp, &var, sizeof(tmp)); + tmp[0] = __shfl_up(tmp[0], lane_delta, width); + tmp[1] = __shfl_up(tmp[1], lane_delta, width); + + uint64_t tmp0 = (static_cast(tmp[1]) << 32ull) | static_cast(tmp[0]); + long tmp1; __builtin_memcpy(&tmp1, &tmp0, sizeof(tmp0)); + return tmp1; + #else + static_assert(sizeof(long) == sizeof(int), ""); + return static_cast(__shfl_up(static_cast(var), lane_delta, width)); + #endif +} +__device__ +inline +long long __shfl_up(long long var, unsigned int lane_delta, int width = warpSize) +{ + static_assert(sizeof(long long) == 2 * sizeof(int), ""); + static_assert(sizeof(long long) == sizeof(uint64_t), ""); + int tmp[2]; __builtin_memcpy(tmp, &var, sizeof(tmp)); + tmp[0] = __shfl_up(tmp[0], lane_delta, width); + tmp[1] = __shfl_up(tmp[1], lane_delta, width); + uint64_t tmp0 = (static_cast(tmp[1]) << 32ull) | static_cast(tmp[0]); + long long tmp1; __builtin_memcpy(&tmp1, &tmp0, sizeof(tmp0)); + return tmp1; +} __device__ inline @@ -393,6 +461,39 @@ double __shfl_down(double var, unsigned int lane_delta, int width = warpSize) { double tmp1; __builtin_memcpy(&tmp1, &tmp0, sizeof(tmp0)); return tmp1; } +__device__ +inline +long __shfl_down(long var, unsigned int lane_delta, int width = warpSize) +{ + #ifndef _MSC_VER + static_assert(sizeof(long) == 2 * sizeof(int), ""); + static_assert(sizeof(long) == sizeof(uint64_t), ""); + + int tmp[2]; __builtin_memcpy(tmp, &var, sizeof(tmp)); + tmp[0] = __shfl_down(tmp[0], lane_delta, width); + tmp[1] = __shfl_down(tmp[1], lane_delta, width); + + uint64_t tmp0 = (static_cast(tmp[1]) << 32ull) | static_cast(tmp[0]); + long tmp1; __builtin_memcpy(&tmp1, &tmp0, sizeof(tmp0)); + return tmp1; + #else + static_assert(sizeof(long) == sizeof(int), ""); + return static_cast(__shfl_down(static_cast(var), lane_delta, width)); + #endif +} +__device__ +inline +long long __shfl_down(long long var, unsigned int lane_delta, int width = warpSize) +{ + static_assert(sizeof(long long) == 2 * sizeof(int), ""); + static_assert(sizeof(long long) == sizeof(uint64_t), ""); + int tmp[2]; __builtin_memcpy(tmp, &var, sizeof(tmp)); + tmp[0] = __shfl_down(tmp[0], lane_delta, width); + tmp[1] = __shfl_down(tmp[1], lane_delta, width); + uint64_t tmp0 = (static_cast(tmp[1]) << 32ull) | static_cast(tmp[0]); + long long tmp1; __builtin_memcpy(&tmp1, &tmp0, sizeof(tmp0)); + return tmp1; +} __device__ inline @@ -430,6 +531,39 @@ double __shfl_xor(double var, int lane_mask, int width = warpSize) { double tmp1; __builtin_memcpy(&tmp1, &tmp0, sizeof(tmp0)); return tmp1; } +__device__ +inline +long __shfl_xor(long var, int lane_mask, int width = warpSize) +{ + #ifndef _MSC_VER + static_assert(sizeof(long) == 2 * sizeof(int), ""); + static_assert(sizeof(long) == sizeof(uint64_t), ""); + + int tmp[2]; __builtin_memcpy(tmp, &var, sizeof(tmp)); + tmp[0] = __shfl_xor(tmp[0], lane_mask, width); + tmp[1] = __shfl_xor(tmp[1], lane_mask, width); + + uint64_t tmp0 = (static_cast(tmp[1]) << 32ull) | static_cast(tmp[0]); + long tmp1; __builtin_memcpy(&tmp1, &tmp0, sizeof(tmp0)); + return tmp1; + #else + static_assert(sizeof(long) == sizeof(int), ""); + return static_cast(__shfl_down(static_cast(var), lane_delta, width)); + #endif +} +__device__ +inline +long long __shfl_xor(long long var, int lane_mask, int width = warpSize) +{ + static_assert(sizeof(long long) == 2 * sizeof(int), ""); + static_assert(sizeof(long long) == sizeof(uint64_t), ""); + int tmp[2]; __builtin_memcpy(tmp, &var, sizeof(tmp)); + tmp[0] = __shfl_xor(tmp[0], lane_mask, width); + tmp[1] = __shfl_xor(tmp[1], lane_mask, width); + uint64_t tmp0 = (static_cast(tmp[1]) << 32ull) | static_cast(tmp[0]); + long long tmp1; __builtin_memcpy(&tmp1, &tmp0, sizeof(tmp0)); + return tmp1; +} #define MASK1 0x00ff00ff #define MASK2 0xff00ff00 diff --git a/hipamd/include/hip/hcc_detail/driver_types.h b/hipamd/include/hip/hcc_detail/driver_types.h index 0c29542c7e..1941f44617 100644 --- a/hipamd/include/hip/hcc_detail/driver_types.h +++ b/hipamd/include/hip/hcc_detail/driver_types.h @@ -255,14 +255,14 @@ typedef struct hipMemcpy3DParms { hipArray_t srcArray; struct hipPos srcPos; struct hipPitchedPtr srcPtr; - hipArray_t dstArray; struct hipPos dstPos; struct hipPitchedPtr dstPtr; - struct hipExtent extent; enum hipMemcpyKind kind; +} hipMemcpy3DParms; +typedef struct HIP_MEMCPY3D { size_t Depth; size_t Height; size_t WidthInBytes; @@ -283,10 +283,7 @@ typedef struct hipMemcpy3DParms { size_t srcLOD; hipMemoryType srcMemoryType; size_t srcPitch; - size_t srcXInBytes; - size_t srcY; - size_t srcZ; -}hipMemcpy3DParms; +} HIP_MEMCPY3D; static inline struct hipPitchedPtr make_hipPitchedPtr(void* d, size_t p, size_t xsz, size_t ysz) { diff --git a/hipamd/include/hip/hcc_detail/hip_runtime_api.h b/hipamd/include/hip/hcc_detail/hip_runtime_api.h index 936170b1fd..1f0b474863 100644 --- a/hipamd/include/hip/hcc_detail/hip_runtime_api.h +++ b/hipamd/include/hip/hcc_detail/hip_runtime_api.h @@ -212,6 +212,11 @@ enum hipLimit_t { #define hipArrayCubemap 0x04 #define hipArrayTextureGather 0x08 +#define hipOccupancyDefault 0x00 + +#define hipCooperativeLaunchMultiDeviceNoPreSync 0x01 +#define hipCooperativeLaunchMultiDeviceNoPostSync 0x02 + /* * @brief hipJitOption * @enum @@ -2903,7 +2908,7 @@ hipError_t hipModuleLaunchKernel(hipFunction_t f, unsigned int gridDimX, unsigne * @param [in] stream Stream where the kernel should be dispatched. May be 0, in which case th * default stream is used with associated synchronization rules. * - * @returns hipSuccess, hipInvalidDevice, hipErrorNotInitialized, hipErrorInvalidValue + * @returns hipSuccess, hipInvalidDevice, hipErrorNotInitialized, hipErrorInvalidValue, hipErrorCooperativeLaunchTooLarge */ hipError_t hipLaunchCooperativeKernel(const void* f, dim3 gridDim, dim3 blockDimX, void** kernelParams, unsigned int sharedMemBytes, @@ -2917,7 +2922,7 @@ hipError_t hipLaunchCooperativeKernel(const void* f, dim3 gridDim, dim3 blockDim * @param [in] numDevices Size of the launchParamsList array. * @param [in] flags Flags to control launch behavior. * - * @returns hipSuccess, hipInvalidDevice, hipErrorNotInitialized, hipErrorInvalidValue + * @returns hipSuccess, hipInvalidDevice, hipErrorNotInitialized, hipErrorInvalidValue, hipErrorCooperativeLaunchTooLarge */ hipError_t hipLaunchCooperativeKernelMultiDevice(hipLaunchParams* launchParamsList, int numDevices, unsigned int flags); @@ -2954,13 +2959,36 @@ hipError_t hipOccupancyMaxActiveBlocksPerMultiprocessor( * @brief Returns occupancy for a device function. * * @param [out] numBlocks Returned occupancy - * @param [in] func Kernel function for which occupancy is calulated + * @param [in] func Kernel function (hipFunction) for which occupancy is calulated + * @param [in] blockSize Block size the kernel is intended to be launched with + * @param [in] dynSharedMemPerBlk dynamic shared memory usage (in bytes) intended for each block + */ +hipError_t hipDrvOccupancyMaxActiveBlocksPerMultiprocessor( + int* numBlocks, hipFunction_t f, int blockSize, size_t dynSharedMemPerBlk); + +/** + * @brief Returns occupancy for a device function. + * + * @param [out] numBlocks Returned occupancy + * @param [in] f Kernel function for which occupancy is calulated * @param [in] blockSize Block size the kernel is intended to be launched with * @param [in] dynSharedMemPerBlk dynamic shared memory usage (in bytes) intended for each block * @param [in] flags Extra flags for occupancy calculation (currently ignored) */ hipError_t hipOccupancyMaxActiveBlocksPerMultiprocessorWithFlags( - uint32_t* numBlocks, hipFunction_t f, uint32_t blockSize, size_t dynSharedMemPerBlk, unsigned int flags); + uint32_t* numBlocks, hipFunction_t f, uint32_t blockSize, size_t dynSharedMemPerBlk, unsigned int flags __dparm(hipOccupancyDefault)); + +/** + * @brief Returns occupancy for a device function. + * + * @param [out] numBlocks Returned occupancy + * @param [in] f Kernel function(hipFunction_t) for which occupancy is calulated + * @param [in] blockSize Block size the kernel is intended to be launched with + * @param [in] dynSharedMemPerBlk dynamic shared memory usage (in bytes) intended for each block + * @param [in] flags Extra flags for occupancy calculation (currently ignored) + */ +hipError_t hipDrvOccupancyMaxActiveBlocksPerMultiprocessorWithFlags( + int* numBlocks, hipFunction_t f, int blockSize, size_t dynSharedMemPerBlk, unsigned int flags); /** * @brief Launches kernels on multiple devices and guarantees all specified kernels are dispatched diff --git a/hipamd/include/hip/hcc_detail/hiprtc.h b/hipamd/include/hip/hcc_detail/hiprtc.h index 26d3129dbc..624f1ea157 100644 --- a/hipamd/include/hip/hcc_detail/hiprtc.h +++ b/hipamd/include/hip/hcc_detail/hiprtc.h @@ -19,10 +19,14 @@ 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 +#ifndef HIPRTC_H +#define HIPRTC_H -#include -#include +#ifdef __cplusplus +extern "C" { +#endif /* __cplusplus */ + +#include enum hiprtcResult { HIPRTC_SUCCESS = 0, @@ -41,29 +45,22 @@ enum hiprtcResult { const char* hiprtcGetErrorString(hiprtcResult result); -inline -hiprtcResult hiprtcVersion(int* major, int* minor) noexcept -{ // TODO: NVRTC versioning is somewhat unclear. - if (!major || !minor) return HIPRTC_ERROR_INVALID_INPUT; - // TODO: this should be generic / set by the build infrastructure. - *major = 9; - *minor = 0; +hiprtcResult hiprtcVersion(int* major, int* minor); - return HIPRTC_SUCCESS; -} - -struct _hiprtcProgram; -using hiprtcProgram = _hiprtcProgram*; +typedef struct _hiprtcProgram* hiprtcProgram; hiprtcResult hiprtcAddNameExpression(hiprtcProgram prog, const char* name_expression); -hiprtcResult hiprtcCompileProgram(hiprtcProgram prog, int numOptions, +hiprtcResult hiprtcCompileProgram(hiprtcProgram prog, + int numOptions, const char** options); -hiprtcResult hiprtcCreateProgram(hiprtcProgram* prog, const char* src, - const char* name, int numHeaders, +hiprtcResult hiprtcCreateProgram(hiprtcProgram* prog, + const char* src, + const char* name, + int numHeaders, const char** headers, const char** includeNames); @@ -76,37 +73,14 @@ hiprtcResult hiprtcGetLoweredName(hiprtcProgram prog, hiprtcResult hiprtcGetProgramLog(hiprtcProgram prog, char* log); hiprtcResult hiprtcGetProgramLogSize(hiprtcProgram prog, - std::size_t* logSizeRet); + size_t* logSizeRet); hiprtcResult hiprtcGetCode(hiprtcProgram prog, char* code); -hiprtcResult hiprtcGetCodeSize(hiprtcProgram prog, std::size_t* codeSizeRet); +hiprtcResult hiprtcGetCodeSize(hiprtcProgram prog, size_t* codeSizeRet); -namespace hip_impl -{ - char* demangle(const char* mangled_expression); +#ifdef __cplusplus } +#endif /* __cplusplus */ -#if defined(HIPRTC_GET_TYPE_NAME) - #include - - #if defined(_WIN32) - #include - - template - hiprtcResult hiprtcGetTypeName(std::string*) = delete; - #else - template - inline - hiprtcResult hiprtcGetTypeName(std::string* result) - { - if (!result) return HIPRTC_ERROR_INVALID_INPUT; - - char * res= hip_impl::demangle(typeid(T).name()); - result->assign(res == nullptr ? "" : res); - std::free(res); - return (result->empty()) ? HIPRTC_ERROR_INTERNAL_ERROR : - HIPRTC_SUCCESS; - } - #endif -#endif +#endif //HIPRTC_H diff --git a/hipamd/include/hip/hip_profile.h b/hipamd/include/hip/hip_profile.h index 95224af4a3..ff18239e44 100644 --- a/hipamd/include/hip/hip_profile.h +++ b/hipamd/include/hip/hip_profile.h @@ -1,16 +1,13 @@ /* Copyright (c) 2015 - present 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 @@ -23,23 +20,8 @@ THE SOFTWARE. #ifndef HIP_INCLUDE_HIP_HIP_PROFILE_H #define HIP_INCLUDE_HIP_HIP_PROFILE_H -#warning "HIP Profiling through markers is deprecated, please check roctrace/rocTX support." - -#if not defined(ENABLE_HIP_PROFILE) -#define ENABLE_HIP_PROFILE 1 -#endif - -#if defined(__HIP_PLATFORM_HCC__) and (ENABLE_HIP_PROFILE == 1) -#warning "HIP Markers are deprecated and would be removed soon." -#include -#define HIP_SCOPED_MARKER(markerName, group) \ - amdtScopedMarker __scopedMarker(markerName, group, nullptr); -#define HIP_BEGIN_MARKER(markerName, group) amdtBeginMarker(markerName, group, nullptr); -#define HIP_END_MARKER() amdtEndMarker(); -#else #define HIP_SCOPED_MARKER(markerName, group) #define HIP_BEGIN_MARKER(markerName, group) #define HIP_END_MARKER() -#endif #endif diff --git a/hipamd/include/hip/hip_runtime_api.h b/hipamd/include/hip/hip_runtime_api.h index 64b2a85d8a..025688e98c 100644 --- a/hipamd/include/hip/hip_runtime_api.h +++ b/hipamd/include/hip/hip_runtime_api.h @@ -255,6 +255,10 @@ typedef enum __HIP_NODISCARD hipError_t { 713, ///< Produced when trying to unlock a non-page-locked memory. hipErrorLaunchFailure = 719, ///< An exception occurred on the device while executing a kernel. + hipErrorCooperativeLaunchTooLarge = + 720, ///< This error indicates that the number of blocks launched per grid for a kernel + ///< that was launched via cooperative launch APIs exceeds the maximum number of + ///< allowed blocks for the current device hipErrorNotSupported = 801, ///< Produced when the hip API is not supported/implemented hipErrorUnknown = 999, //< Unknown error. // HSA Runtime Error Codes start here. diff --git a/hipamd/include/hip/nvcc_detail/hip_runtime_api.h b/hipamd/include/hip/nvcc_detail/hip_runtime_api.h index 05a162478a..6e0d02d0c0 100644 --- a/hipamd/include/hip/nvcc_detail/hip_runtime_api.h +++ b/hipamd/include/hip/nvcc_detail/hip_runtime_api.h @@ -140,6 +140,14 @@ typedef enum cudaChannelFormatKind hipChannelFormatKind; #define hipLimitMallocHeapSize cudaLimitMallocHeapSize #define hipIpcMemLazyEnablePeerAccess cudaIpcMemLazyEnablePeerAccess +#define hipOccupancyDefault cudaOccupancyDefault + +#define hipCooperativeLaunchMultiDeviceNoPreSync \ + cudaCooperativeLaunchMultiDeviceNoPreSync +#define hipCooperativeLaunchMultiDeviceNoPostSync \ + cudaCooperativeLaunchMultiDeviceNoPostSync + + // enum CUjit_option redefines #define hipJitOptionMaxRegisters CU_JIT_MAX_REGISTERS #define hipJitOptionThreadsPerBlock CU_JIT_THREADS_PER_BLOCK @@ -267,6 +275,8 @@ inline static hipError_t hipCUDAErrorTohipError(cudaError_t cuError) { return hipErrorNotInitialized; case cudaErrorLaunchFailure: return hipErrorLaunchFailure; + case cudaErrorCooperativeLaunchTooLarge: + return hipErrorCooperativeLaunchTooLarge; case cudaErrorPriorLaunchFailure: return hipErrorPriorLaunchFailure; case cudaErrorLaunchOutOfResources: @@ -445,6 +455,8 @@ inline static hipError_t hipCUResultTohipError(CUresult cuError) { return hipErrorHostMemoryNotRegistered; case CUDA_ERROR_LAUNCH_FAILED: return hipErrorLaunchFailure; + case CUDA_ERROR_COOPERATIVE_LAUNCH_TOO_LARGE: + return hipErrorCooperativeLaunchTooLarge; case CUDA_ERROR_NOT_SUPPORTED: return hipErrorNotSupported; case CUDA_ERROR_UNKNOWN: @@ -601,6 +613,8 @@ inline static cudaError_t hipErrorToCudaError(hipError_t hError) { return cudaErrorSetOnActiveProcess; case hipErrorLaunchFailure: return cudaErrorLaunchFailure; + case hipErrorCooperativeLaunchTooLarge: + return cudaErrorCooperativeLaunchTooLarge; case hipErrorNotSupported: return cudaErrorNotSupported; // HSA: does not exist in CUDA diff --git a/hipamd/lpl_ca/CMakeLists.txt b/hipamd/lpl_ca/CMakeLists.txt index b36d73bbcb..ac01a6a0ab 100644 --- a/hipamd/lpl_ca/CMakeLists.txt +++ b/hipamd/lpl_ca/CMakeLists.txt @@ -26,7 +26,7 @@ target_include_directories(ca PUBLIC ${PROJECT_SOURCE_DIR}/src) find_library( hsart NAMES libhsa-runtime64.so libhsa-runtime64.so.1 HINTS ${HSA_PATH}/lib) target_link_libraries(ca PUBLIC ${hsart}) -target_compile_options(ca PUBLIC -Wall) +target_compile_options(ca PUBLIC -DDISABLE_REDUCED_GPU_BLOB_COPY -Wall) install(TARGETS ca RUNTIME DESTINATION bin) -#-------------------------------------CA---------------------------------------# \ No newline at end of file +#-------------------------------------CA---------------------------------------# diff --git a/hipamd/packaging/hip-hcc.txt b/hipamd/packaging/hip-hcc.txt index 63f3d73e67..21e138e1ed 100644 --- a/hipamd/packaging/hip-hcc.txt +++ b/hipamd/packaging/hip-hcc.txt @@ -36,11 +36,7 @@ set(CPACK_PACKAGE_FILE_NAME ${CPACK_PACKAGE_NAME}-${CPACK_PACKAGE_VERSION_MAJOR} set(CPACK_GENERATOR "TGZ;DEB;RPM") set(CPACK_BINARY_DEB "ON") set(CPACK_DEBIAN_PACKAGE_CONTROL_EXTRA "${PROJECT_BINARY_DIR}/postinst;${PROJECT_BINARY_DIR}/prerm") -if(@COMPILE_HIP_ATP_MARKER@) - set(CPACK_DEBIAN_PACKAGE_DEPENDS "hip-base (= ${CPACK_PACKAGE_VERSION}), ${HCC_PACKAGE_NAME} (= @HCC_PACKAGE_VERSION@), rocm-profiler, comgr (>= 1.1)") -else() - set(CPACK_DEBIAN_PACKAGE_DEPENDS "hip-base (= ${CPACK_PACKAGE_VERSION}), ${HCC_PACKAGE_NAME} (= @HCC_PACKAGE_VERSION@), comgr (>= 1.1)") -endif() +set(CPACK_DEBIAN_PACKAGE_DEPENDS "hip-base (= ${CPACK_PACKAGE_VERSION}), ${HCC_PACKAGE_NAME} (= @HCC_PACKAGE_VERSION@), comgr (>= 1.1)") set(CPACK_DEBIAN_PACKAGE_PROVIDES "hip_hcc") set(CPACK_DEBIAN_PACKAGE_REPLACES "hip_hcc") set(CPACK_DEBIAN_PACKAGE_CONFLICTS "hip_hcc") @@ -50,11 +46,7 @@ set(CPACK_RPM_POST_INSTALL_SCRIPT_FILE "${PROJECT_BINARY_DIR}/postinst") set(CPACK_RPM_PRE_UNINSTALL_SCRIPT_FILE "${PROJECT_BINARY_DIR}/prerm") set(CPACK_RPM_PACKAGE_AUTOREQPROV " no") string(REPLACE "-" "_" HIP_BASE_VERSION ${CPACK_PACKAGE_VERSION}) -if(@COMPILE_HIP_ATP_MARKER@) - set(CPACK_RPM_PACKAGE_REQUIRES "hip-base = ${HIP_BASE_VERSION}, ${HCC_PACKAGE_NAME} = @HCC_PACKAGE_VERSION@, rocm-profiler, comgr >= 1.1") -else() - set(CPACK_RPM_PACKAGE_REQUIRES "hip-base = ${HIP_BASE_VERSION}, ${HCC_PACKAGE_NAME} = @HCC_PACKAGE_VERSION@, comgr >= 1.1") -endif() +set(CPACK_RPM_PACKAGE_REQUIRES "hip-base = ${HIP_BASE_VERSION}, ${HCC_PACKAGE_NAME} = @HCC_PACKAGE_VERSION@, comgr >= 1.1") set(CPACK_RPM_PACKAGE_OBSOLETES "hip_hcc") set(CPACK_RPM_PACKAGE_CONFLICTS "hip_hcc") set(CPACK_RPM_EXCLUDE_FROM_AUTO_FILELIST_ADDITION "/opt") diff --git a/hipamd/samples/1_Utils/hipDispatchLatency/Makefile b/hipamd/samples/1_Utils/hipDispatchLatency/Makefile index 0616f01f0d..74945dc515 100644 --- a/hipamd/samples/1_Utils/hipDispatchLatency/Makefile +++ b/hipamd/samples/1_Utils/hipDispatchLatency/Makefile @@ -4,16 +4,17 @@ ifeq (,$(HIP_PATH)) endif HIPCC=$(HIP_PATH)/bin/hipcc -std=c++11 -EXE=hipDispatchLatency - CXXFLAGS = -O3 -all: test_kernel.code ${EXE} +all: test_kernel.code hipDispatchLatency.out hipDispatchEnqueueRateMT.out -$(EXE): hipDispatchLatency.cpp +hipDispatchLatency.out: hipDispatchLatency.cpp $(HIPCC) $(CXXFLAGS) hipDispatchLatency.cpp -o $@ +hipDispatchEnqueueRateMT.out: hipDispatchEnqueueRateMT.cpp + $(HIPCC) $(CXXFLAGS) hipDispatchEnqueueRateMT.cpp -o $@ + test_kernel.code: test_kernel.cpp $(HIP_PATH)/bin/hipcc --genco $(GENCO_FLAGS) $^ -o $@ clean: - rm -f *.o $(EXE) + rm -f *.o *.out diff --git a/hipamd/samples/1_Utils/hipDispatchLatency/hipDispatchEnqueueRateMT.cpp b/hipamd/samples/1_Utils/hipDispatchLatency/hipDispatchEnqueueRateMT.cpp new file mode 100644 index 0000000000..d1b5c2f3b5 --- /dev/null +++ b/hipamd/samples/1_Utils/hipDispatchLatency/hipDispatchEnqueueRateMT.cpp @@ -0,0 +1,167 @@ +/* +Copyright (c) 2020-present 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 "hip/hip_runtime.h" +#ifdef __HIP_PLATFORM_HCC__ +#include "hip/hip_ext.h" +#endif +#include +#include +#include +#include +#include +#include +#include + +#define NUM_GROUPS 1 +#define GROUP_SIZE 1 +#define WARMUP_RUN_COUNT 10 +#define TIMING_RUN_COUNT 100 +#define TOTAL_RUN_COUNT WARMUP_RUN_COUNT + TIMING_RUN_COUNT + +__global__ void EmptyKernel() {} + +// Helper to print various timing metrics +void print_timing(std::string test, std::array &results, int batch = 1) +{ + + float total_us = 0.0f, mean_us = 0.0f, stddev_us = 0.0f; + + // remove top outliers due to nature of variability across large number of multi-threaded runs + std::sort(results.begin(), results.end(), std::greater()); + auto start_iter = std::next(results.begin(), WARMUP_RUN_COUNT); + auto end_iter = results.end(); + + // mean + std::for_each(start_iter, end_iter, [&](const float &run_ms) { + total_us += (run_ms * 1000) / batch; + }); + mean_us = total_us / TIMING_RUN_COUNT; + + // stddev + total_us = 0; + std::for_each(start_iter, end_iter, [&](const float &run_ms) { + float dev_us = ((run_ms * 1000) / batch) - mean_us; + total_us += dev_us * dev_us; + }); + stddev_us = sqrt(total_us / TIMING_RUN_COUNT); + + printf("\n %s: %.1f us, std: %.1f us\n", test.c_str(), mean_us, stddev_us); +} + +// Measure time taken to enqueue a kernel on the GPU using hipModuleLaunchKernel +void hipModuleLaunchKernel_enqueue_rate(std::atomic_int* shared, int max_threads) +{ + //resources necessary for this thread + hipStream_t stream; + hipStreamCreate(&stream); + hipModule_t module; + hipFunction_t function; + hipModuleLoad(&module, "test_kernel.code"); + hipModuleGetFunction(&function, module, "test"); + void* kernel_params = nullptr; + std::array results; + + //synchronize all threads, before running + int tid = shared->fetch_add(1, std::memory_order_release); + while (max_threads != shared->load(std::memory_order_acquire)) {} + + for (auto i = 0; i < TOTAL_RUN_COUNT; ++i) { + auto start = std::chrono::high_resolution_clock::now(); + hipModuleLaunchKernel(function, 1, 1, 1, 1, 1, 1, 0, stream, &kernel_params, nullptr); + auto stop = std::chrono::high_resolution_clock::now(); + results[i] = std::chrono::duration(stop - start).count(); + } + print_timing("Thread ID : " + std::to_string(tid) + " , " + "hipModuleLaunchKernel enqueue rate", results); +} + +// Measure time taken to enqueue a kernel on the GPU using hipLaunchKernelGGL +void hipLaunchKernelGGL_enqueue_rate(std::atomic_int* shared, int max_threads) +{ + //resources necessary for this thread + hipStream_t stream; + hipStreamCreate(&stream); + std::array results; + + //synchronize all threads, before running + int tid = shared->fetch_add(1, std::memory_order_release); + while (max_threads != shared->load(std::memory_order_acquire)) {} + + for (auto i = 0; i < TOTAL_RUN_COUNT; ++i) { + auto start = std::chrono::high_resolution_clock::now(); + hipLaunchKernelGGL((EmptyKernel), dim3(NUM_GROUPS), dim3(GROUP_SIZE), 0, stream); + auto stop = std::chrono::high_resolution_clock::now(); + results[i] = std::chrono::duration(stop - start).count(); + } + print_timing("Thread ID : " + std::to_string(tid) + " , " + "hipLaunchKernelGGL enqueue rate", results); +} + +// Simple thread pool +struct thread_pool { + thread_pool(int total_threads) : max_threads(total_threads) {} + void start(std::function f) { + for (int i = 0; i < max_threads; ++i) { + threads.push_back(std::async(std::launch::async, f, &shared, max_threads)); + } + } + void finish() { + for (auto&&thread : threads) { + thread.get(); + } + threads.clear(); + shared = {0}; + } + ~thread_pool() { + finish(); + } +private: + std::atomic_int shared {0}; + std::vector> threads; + int max_threads = 1; +}; + + +int main(int argc, char* argv[]) +{ + if (argc != 3) { + std::cerr << "Run test as 'hipDispatchEnqueueRateMT <0-hipModuleLaunchKernel /1-hipLaunchKernelGGL>'\n"; + return -1; + } + + int max_threads = atoi(argv[1]); + int run_module_test = atoi(argv[2]); + if(max_threads < 1 || run_module_test < 0 || run_module_test > 1) { + std::cerr << "Invalid Input.\n"; + std::cerr << "Run test as 'hipDispatchEnqueueRateMT <0-hipModuleLaunchKernel /1-hipLaunchKernelGGL>'\n"; + return -1; + } + thread_pool task(max_threads); + + if(run_module_test == 0) { + task.start(hipModuleLaunchKernel_enqueue_rate); + task.finish(); + } else { + task.start(hipLaunchKernelGGL_enqueue_rate); + task.finish(); + } + + return 0; +} + diff --git a/hipamd/src/hip_hcc.cpp b/hipamd/src/hip_hcc.cpp index 175d301ee1..c9688408c8 100644 --- a/hipamd/src/hip_hcc.cpp +++ b/hipamd/src/hip_hcc.cpp @@ -39,6 +39,7 @@ THE SOFTWARE. #include #include #include +#include #include #include @@ -71,7 +72,6 @@ int HIP_API_BLOCKING = 0; int HIP_PRINT_ENV = 0; int HIP_TRACE_API = 0; std::string HIP_TRACE_API_COLOR("green"); -int HIP_PROFILE_API = 0; // TODO - DB_START/STOP need more testing. std::string HIP_DB_START_API; @@ -149,12 +149,10 @@ uint64_t recordApiTrace(TlsData *tls, std::string* fullStr, const std::string& a if ((tid < g_dbStartTriggers.size()) && (apiSeqNum >= g_dbStartTriggers[tid].nextTrigger())) { printf("info: resume profiling at %lu\n", apiSeqNum); - RESUME_PROFILING; g_dbStartTriggers.pop_back(); }; if ((tid < g_dbStopTriggers.size()) && (apiSeqNum >= g_dbStopTriggers[tid].nextTrigger())) { printf("info: stop profiling at %lu\n", apiSeqNum); - STOP_PROFILING; g_dbStopTriggers.pop_back(); }; @@ -1294,9 +1292,6 @@ void HipReadEnv() { "executes."); READ_ENV_S(release, HIP_TRACE_API_COLOR, 0, "Color to use for HIP_API. None/Red/Green/Yellow/Blue/Magenta/Cyan/White"); - READ_ENV_I(release, HIP_PROFILE_API, 0, - "Add HIP API markers to ATP file generated with CodeXL. 0x1=short API name, " - "0x2=full API name including args."); READ_ENV_S(release, HIP_DB_START_API, 0, "Comma-separated list of tid.api_seq_num for when to start debug and profiling."); READ_ENV_S(release, HIP_DB_STOP_API, 0, @@ -1372,14 +1367,6 @@ void HipReadEnv() { HIP_DB |= 0x1; } - if (HIP_PROFILE_API && !COMPILE_HIP_ATP_MARKER) { - fprintf(stderr, - "warning: env var HIP_PROFILE_API=0x%x but COMPILE_HIP_ATP_MARKER=0. (perhaps " - "enable COMPILE_HIP_ATP_MARKER in src code before compiling?)\n", - HIP_PROFILE_API); - HIP_PROFILE_API = 0; - } - if (HIP_DB) { fprintf(stderr, "HIP_DB=0x%x [%s]\n", HIP_DB, HIP_DB_string(HIP_DB).c_str()); } @@ -1423,11 +1410,6 @@ void HipReadEnv() { // This function creates a vector with only the GPU accelerators. // It is called with C++11 call_once, which provided thread-safety. void ihipInit() { -#if COMPILE_HIP_ATP_MARKER - amdtInitializeActivityLogger(); - amdtScopedMarker("ihipInit", "HIP", NULL); -#endif - HipReadEnv(); @@ -1617,7 +1599,7 @@ hipStream_t ihipSyncAndResolveStream(hipStream_t stream, bool lockAcquired) { void ihipPrintKernelLaunch(const char* kernelName, const grid_launch_parm* lp, const hipStream_t stream) { - if ((HIP_TRACE_API & (1 << TRACE_KCMD)) || HIP_PROFILE_API || + if ((HIP_TRACE_API & (1 << TRACE_KCMD)) || (COMPILE_HIP_DB & HIP_TRACE_API)) { GET_TLS(); std::stringstream os; @@ -1630,14 +1612,6 @@ void ihipPrintKernelLaunch(const char* kernelName, const grid_launch_parm* lp, std::string fullStr; recordApiTrace(tls, &fullStr, os.str()); } - - if (HIP_PROFILE_API == 0x1) { - std::string shortAtpString("hipLaunchKernel:"); - shortAtpString += kernelName; - MARKER_BEGIN(shortAtpString.c_str(), "HIP"); - } else if (HIP_PROFILE_API == 0x2) { - MARKER_BEGIN(os.str().c_str(), "HIP"); - } } } @@ -1696,9 +1670,6 @@ void ihipPostLaunchKernel(const char* kernelName, hipStream_t stream, grid_launc tprintf(DB_SYNC, "ihipPostLaunchKernel, unlocking stream\n"); stream->lockclose_postKernelCommand(kernelName, lp.av, unlockPostponed); - if (HIP_PROFILE_API) { - MARKER_END(); - } } //================================================================================================= @@ -1796,6 +1767,8 @@ const char* ihipErrorString(hipError_t hip_error) { return "hipErrorMissingConfiguration"; case hipErrorLaunchFailure: return "hipErrorLaunchFailure"; + case hipErrorCooperativeLaunchTooLarge: + return "hipErrorCooperativeLaunchTooLarge"; case hipErrorPriorLaunchFailure: return "hipErrorPriorLaunchFailure"; case hipErrorLaunchTimeOut: @@ -2478,29 +2451,17 @@ bool ihipStream_t::locked_copy2DAsync(void* dst, const void* src, size_t width, return retStatus; } -//------------------------------------------------------------------------------------------------- -//------------------------------------------------------------------------------------------------- -// Profiler, really these should live elsewhere: hipError_t hipProfilerStart() { HIP_INIT_API(hipProfilerStart); -#if COMPILE_HIP_ATP_MARKER - amdtResumeProfiling(AMDT_ALL_PROFILING); -#endif - return ihipLogStatus(hipSuccess); }; hipError_t hipProfilerStop() { HIP_INIT_API(hipProfilerStop); -#if COMPILE_HIP_ATP_MARKER - amdtStopProfiling(AMDT_ALL_PROFILING); -#endif - return ihipLogStatus(hipSuccess); }; - //------------------------------------------------------------------------------------------------- //------------------------------------------------------------------------------------------------- // HCC-specific accessor functions: @@ -2540,6 +2501,16 @@ hipError_t hipHccGetAcceleratorView(hipStream_t stream, hc::accelerator_view** a // TODO - add a contect sequence number for debug. Print operator<< ctx:0.1 (device.ctx) namespace hip_impl { + std::unordered_set& get_all_gpuarch() { + static std::unordered_set r{}; + static std::once_flag init; + std::call_once(init, []() { + for (int i=0; i < g_deviceCnt; i++){ + r.insert("hcc-amdgcn-amd-amdhsa--gfx"+std::to_string(g_deviceArray[i]->_props.gcnArch)); + }}); + return r; + } + std::vector all_hsa_agents() { std::vector r{}; std::vector visible_accelerators; diff --git a/hipamd/src/hip_hcc_internal.h b/hipamd/src/hip_hcc_internal.h index 0510015c42..ac63f49dba 100644 --- a/hipamd/src/hip_hcc_internal.h +++ b/hipamd/src/hip_hcc_internal.h @@ -63,7 +63,6 @@ extern int HIP_LAUNCH_BLOCKING; extern int HIP_API_BLOCKING; extern int HIP_PRINT_ENV; -extern int HIP_PROFILE_API; // extern int HIP_TRACE_API; extern int HIP_ATP; extern int HIP_DB; @@ -250,34 +249,6 @@ extern const char* API_COLOR_END; // Must be enabled at runtime with HIP_TRACE_API #define COMPILE_HIP_TRACE_API 0x3 - -// Compile code that generates trace markers for CodeXL ATP at HIP function begin/end. -// ATP is standard CodeXL format that includes timestamps for kernels, HSA RT APIs, and HIP APIs. -#ifndef COMPILE_HIP_ATP_MARKER -#define COMPILE_HIP_ATP_MARKER 0 -#endif - - -// Compile support for trace markers that are displayed on CodeXL GUI at start/stop of each function -// boundary. -// TODO - currently we print the trace message at the beginning. if we waited, we could also -// tls->tidInfo return codes, and any values returned through ptr-to-args (ie the pointers allocated -// by hipMalloc). -#if COMPILE_HIP_ATP_MARKER -#include "CXLActivityLogger.h" -#define MARKER_BEGIN(markerName, group) amdtBeginMarker(markerName, group, nullptr); -#define MARKER_END() amdtEndMarker(); -#define RESUME_PROFILING amdtResumeProfiling(AMDT_ALL_PROFILING); -#define STOP_PROFILING amdtStopProfiling(AMDT_ALL_PROFILING); -#else -// Swallow scoped markers: -#define MARKER_BEGIN(markerName, group) -#define MARKER_END() -#define RESUME_PROFILING -#define STOP_PROFILING -#endif - - //--- // HIP Trace modes - use with HIP_TRACE_API=... #define TRACE_ALL 0 // 0x01 @@ -336,22 +307,17 @@ static inline uint64_t getTicks() { return hc::get_system_ticks(); } //--- extern uint64_t recordApiTrace(TlsData *tls, std::string* fullStr, const std::string& apiStr); -#if COMPILE_HIP_ATP_MARKER || (COMPILE_HIP_TRACE_API & 0x1) +#if (COMPILE_HIP_TRACE_API & 0x1) #define API_TRACE(forceTrace, ...) \ GET_TLS(); \ uint64_t hipApiStartTick = 0; \ { \ tls->tidInfo.incApiSeqNum(); \ if (forceTrace || \ - (HIP_PROFILE_API || (COMPILE_HIP_DB && (HIP_TRACE_API & (1 << TRACE_ALL))))) { \ + (COMPILE_HIP_DB && (HIP_TRACE_API & (1 << TRACE_ALL)))) { \ std::string apiStr = std::string(__func__) + " (" + ToString(__VA_ARGS__) + ')'; \ std::string fullStr; \ hipApiStartTick = recordApiTrace(tls, &fullStr, apiStr); \ - if (HIP_PROFILE_API == 0x1) { \ - MARKER_BEGIN(__func__, "HIP") \ - } else if (HIP_PROFILE_API == 0x2) { \ - MARKER_BEGIN(fullStr.c_str(), "HIP"); \ - } \ } \ } @@ -398,9 +364,6 @@ extern uint64_t recordApiTrace(TlsData *tls, std::string* fullStr, const std::st tls->tidInfo.apiSeqNum(), __func__, localHipStatus, \ ihipErrorString(localHipStatus), ticks, API_COLOR_END); \ } \ - if (HIP_PROFILE_API) { \ - MARKER_END(); \ - } \ localHipStatus; \ }) diff --git a/hipamd/src/hip_memory.cpp b/hipamd/src/hip_memory.cpp index 444e41107a..8159f22a97 100644 --- a/hipamd/src/hip_memory.cpp +++ b/hipamd/src/hip_memory.cpp @@ -37,7 +37,6 @@ __device__ uint32_t __hip_device_page_flag[__HIP_NUM_PAGES]; namespace hip_internal { namespace { - inline const char* hsa_to_string(hsa_status_t err) noexcept { @@ -140,6 +139,8 @@ namespace { return r; }()}; + constexpr std::uint32_t is_cpu_owned{UINT32_MAX}; + inline hsa_amd_pointer_info_t info(const void* p) { @@ -149,13 +150,14 @@ namespace { const_cast(p), &r, nullptr, nullptr, nullptr), __FILE__, __func__, __LINE__); - r.size = is_large_BAR || (type(r.agentOwner) == HSA_DEVICE_TYPE_CPU) ? - UINT32_MAX : sizeof(hsa_amd_pointer_info_t); + if (type(r.agentOwner) == HSA_DEVICE_TYPE_CPU) r.size = is_cpu_owned; return r; } - constexpr size_t staging_sz{4 * 1024 * 1024}; // 2 Pages. + constexpr size_t staging_sz{4 * 1024 * 1024}; // 2 Pages. + constexpr size_t max_h2d_std_memcpy_sz{8 * 1024}; // 8 KiB. + constexpr size_t max_d2h_std_memcpy_sz{64}; // 1 cacheline. thread_local const std::unique_ptr staging_buffer{ []() { @@ -202,7 +204,7 @@ namespace { } // Unnamed namespace. inline -void do_copy(void* __restrict dst, const void* __restrict src, std::size_t n, +void do_copy(void* __restrict dst, const void* __restrict src, size_t n, hsa_agent_t da, hsa_agent_t sa) { hsa_signal_silent_store_relaxed(copy_signal, 1); throwing_result_check( @@ -224,14 +226,20 @@ void do_std_memcpy( inline void d2h_copy(void* __restrict dst, const void* __restrict src, size_t n, hsa_amd_pointer_info_t si) { - // TODO: characterise direct largeBAR reads from agent-allocated memory. - // if (si.size == UINT32_MAX) { - // return do_std_memcpy(dst, src, n); - // } - const auto di{info(dst)}; + const auto is_locked{di.type == HSA_EXT_POINTER_TYPE_LOCKED}; - if (di.type == HSA_EXT_POINTER_TYPE_LOCKED) { + if (!is_locked && si.size == is_cpu_owned) { + return do_std_memcpy(dst, src, n); + } + if (!is_locked && is_large_BAR && n <= max_d2h_std_memcpy_sz) { + return do_std_memcpy(dst, src, n); + } + if (di.type == HSA_EXT_POINTER_TYPE_HSA) { + return do_copy(dst, src, n, si.agentOwner, si.agentOwner); + } + + if (is_locked) { dst = static_cast(di.agentBaseAddress) + (static_cast(dst) - static_cast(di.hostBaseAddress)); @@ -247,7 +255,7 @@ void d2h_copy(void* __restrict dst, const void* __restrict src, size_t n, throwing_result_check(hsa_amd_memory_lock(dst, n, &si.agentOwner, 1, const_cast(&dst)), - __FILE__, __func__, __LINE__); + __FILE__, __func__, __LINE__); do_copy(dst, src, n, si.agentOwner, si.agentOwner); } @@ -256,16 +264,23 @@ void d2h_copy(void* __restrict dst, const void* __restrict src, size_t n, inline void h2d_copy(void* __restrict dst, const void* __restrict src, size_t n, hsa_amd_pointer_info_t di) { - if (di.size == UINT32_MAX) { + const auto si{info(const_cast(src))}; + const auto is_locked{si.type == HSA_EXT_POINTER_TYPE_LOCKED}; + + if (!is_locked && di.size == is_cpu_owned) { return do_std_memcpy(dst, src, n); } + if (!is_locked && is_large_BAR && n <= max_h2d_std_memcpy_sz) { + return do_std_memcpy(dst, src, n); + } + if (si.type == HSA_EXT_POINTER_TYPE_HSA) { + return do_copy(dst, src, n, di.agentOwner, di.agentOwner); + } - const auto si{info(const_cast(src))}; - - if (si.type == HSA_EXT_POINTER_TYPE_LOCKED) { + if (is_locked) { src = static_cast(si.agentBaseAddress) + - (static_cast(src) - - static_cast(si.hostBaseAddress)); + (static_cast(src) - + static_cast(si.hostBaseAddress)); do_copy(dst, src, n, di.agentOwner, di.agentOwner); } else if (n <= staging_sz) { @@ -279,7 +294,7 @@ void h2d_copy(void* __restrict dst, const void* __restrict src, size_t n, throwing_result_check(hsa_amd_memory_lock(const_cast(src), n, &di.agentOwner, 1, const_cast(&src)), - __FILE__, __func__, __LINE__); + __FILE__, __func__, __LINE__); do_copy(dst, src, n, di.agentOwner, di.agentOwner); } @@ -288,69 +303,36 @@ void h2d_copy(void* __restrict dst, const void* __restrict src, size_t n, inline void generic_copy(void* __restrict dst, const void* __restrict src, size_t n, hsa_amd_pointer_info_t di, hsa_amd_pointer_info_t si) { - if (di.size == UINT32_MAX && si.size == UINT32_MAX) { + if (di.size == is_cpu_owned && si.size == is_cpu_owned) { return do_std_memcpy(dst, src, n); } + if (di.size == is_cpu_owned) return d2h_copy(dst, src, n, si); + if (si.size == is_cpu_owned) return h2d_copy(dst, src, n, di); - std::unique_ptr lck0{ - nullptr, [](void* p) { hsa_amd_memory_unlock(p); }}; - std::unique_ptr lck1{nullptr, lck0.get_deleter()}; + throwing_result_check(hsa_amd_agents_allow_access(1u, &si.agentOwner, + nullptr, + di.agentBaseAddress), + __FILE__, __func__, __LINE__); - switch (si.type) { - case HSA_EXT_POINTER_TYPE_HSA: - if (di.type == HSA_EXT_POINTER_TYPE_HSA) { - hsa_memory_copy(dst, src, n); - return; // TODO: do_copy(dst, src, n, di.agentOwner, si.agentOwner); - } - - if (di.type == HSA_EXT_POINTER_TYPE_UNKNOWN || - di.type == HSA_EXT_POINTER_TYPE_LOCKED) { - return d2h_copy(dst, src, n, si); - } - break; - case HSA_EXT_POINTER_TYPE_LOCKED: - if (di.type == HSA_EXT_POINTER_TYPE_UNKNOWN) { - std::memcpy(dst, si.hostBaseAddress, n); - - return; - } - if (di.type == HSA_EXT_POINTER_TYPE_LOCKED) { - std::memcpy(di.hostBaseAddress, si.hostBaseAddress, n); - - return; - } - src = si.agentBaseAddress; - si.agentOwner = di.agentOwner; - break; - case HSA_EXT_POINTER_TYPE_UNKNOWN: - if (di.type == HSA_EXT_POINTER_TYPE_UNKNOWN) { - std::memcpy(dst, src, n); - - return; - } - if (di.type == HSA_EXT_POINTER_TYPE_LOCKED) { - std::memcpy(di.hostBaseAddress, src, n); - - return; - } - return h2d_copy(dst, src, n, di); - default: do_copy(dst, src, n, di.agentOwner, si.agentOwner); break; - } + return do_copy(dst, src, n, di.agentOwner, si.agentOwner); } inline void memcpy_impl(void* __restrict dst, const void* __restrict src, size_t n, - hipMemcpyKind k) noexcept { + hipMemcpyKind k) { switch (k) { case hipMemcpyHostToHost: std::memcpy(dst, src, n); break; - case hipMemcpyHostToDevice: - return is_large_BAR ? do_std_memcpy(dst, src, n) - : h2d_copy(dst, src, n, info(dst)); - case hipMemcpyDeviceToHost: - // TODO: characterise direct largeBAR reads from agent-allocated memory. - return /*is_large_BAR ? do_std_memcpy(dst, src, n) - : */d2h_copy(dst, src, n, info(src)); - case hipMemcpyDeviceToDevice: hsa_memory_copy(dst, src, n); break; + case hipMemcpyHostToDevice: return h2d_copy(dst, src, n, info(dst)); + case hipMemcpyDeviceToHost: return d2h_copy(dst, src, n, info(src)); + case hipMemcpyDeviceToDevice: { + const auto di{info(dst)}; + const auto si{info(src)}; + throwing_result_check(hsa_amd_agents_allow_access(1u, &si.agentOwner, + nullptr, + di.agentBaseAddress), + __FILE__, __func__, __LINE__); + return do_copy(dst, src, n, di.agentOwner, si.agentOwner); + } default: return generic_copy(dst, src, n, info(dst), info(src)); } } @@ -1290,7 +1272,7 @@ hipError_t hipMemcpyToSymbolAsync(void* dst, const void* src, size_t count, if (dst == nullptr) { return ihipLogStatus(hipErrorInvalidSymbol); } - + if (kind == hipMemcpyDeviceToHost || kind == hipMemcpyHostToHost) { return ihipLogStatus(hipErrorInvalidMemcpyDirection); } else if (kind == hipMemcpyDeviceToDevice) { @@ -1322,7 +1304,7 @@ hipError_t hipMemcpyFromSymbolAsync(void* dst, const void* src, size_t count, if (src == nullptr || dst == nullptr) { return ihipLogStatus(hipErrorInvalidSymbol); } - + if (kind == hipMemcpyHostToDevice || kind == hipMemcpyHostToHost) { return ihipLogStatus(hipErrorInvalidMemcpyDirection); } else if (kind == hipMemcpyDeviceToDevice) { @@ -1540,111 +1522,144 @@ hipError_t hipMemcpyAtoH(void* dst, hipArray* srcArray, size_t srcOffset, size_t return ihipLogStatus(e); } +int getByteSizeFromFormat(const hipChannelFormatDesc& desc){ + int byteSize =0; + switch (desc.f) { + case hipChannelFormatKindUnsigned: + switch (desc.x) { + case 32: + byteSize = sizeof(uint32_t); + break; + case 16: + byteSize = sizeof(uint16_t); + break; + case 8: + byteSize = sizeof(uint8_t); + break; + default: + byteSize = sizeof(uint32_t); + } + break; + case hipChannelFormatKindSigned: + switch (desc.x) { + case 32: + byteSize = sizeof(int32_t); + break; + case 16: + byteSize = sizeof(int16_t); + break; + case 8: + byteSize = sizeof(int8_t); + break; + default: + byteSize = sizeof(int32_t); + } + break; + case hipChannelFormatKindFloat: + switch (desc.x) { + case 32: + byteSize = sizeof(float); + break; + case 16: + byteSize = sizeof(_Float16); + break; + default: + byteSize = sizeof(float); + } + break; + case hipChannelFormatKindNone: + default: + break; + } + return byteSize; +} + hipError_t ihipMemcpy3D(const struct hipMemcpy3DParms* p, hipStream_t stream, bool isAsync) { hipError_t e = hipSuccess; if(p) { - size_t byteSize, width, height, depth, widthInBytes, srcPitch, dstPitch, ySize; - hipChannelFormatDesc desc; - void* srcPtr;void* dstPtr; + size_t dstByteSize, srcByteSize, copyWidth, copyHeight, copyDepth, widthInBytes, srcPitch, dstPitch, srcYsize, dstYsize; + size_t srcXoffset, srcYoffset, srcZoffset, dstXoffset, dstYoffset, dstZoffset; + size_t srcWidth, srcHeight, srcDepth, dstWidth, dstHeight, dstDepth; + + void* srcPtr, *dstPtr; + bool copyWidthUpdate= false; + copyDepth = p->extent.depth; + copyHeight = p->extent.height; + copyWidth = p->extent.width; // in bytes ? + dstXoffset = p->dstPos.x; + dstYoffset = p->dstPos.y; + dstZoffset = p->dstPos.z; + srcXoffset = p->srcPos.x; + srcYoffset = p->srcPos.y; + srcZoffset = p->srcPos.z; if (p->dstArray != nullptr) { - if (p->dstArray->isDrv == false) { - switch (p->dstArray->desc.f) { - case hipChannelFormatKindSigned: - byteSize = sizeof(int); - break; - case hipChannelFormatKindUnsigned: - byteSize = sizeof(unsigned int); - break; - case hipChannelFormatKindFloat: - byteSize = sizeof(float); - break; - case hipChannelFormatKindNone: - byteSize = sizeof(size_t); - break; - default: - byteSize = 0; - break; - } - depth = p->extent.depth; - height = p->extent.height; - width = p->extent.width; - widthInBytes = p->extent.width * byteSize; - srcPitch = p->srcPtr.pitch; - srcPtr = p->srcPtr.ptr; - ySize = p->srcPtr.ysize; - desc = p->dstArray->desc; - dstPtr = p->dstArray->data; - hsa_ext_image_data_info_t imageInfo; - if(hipTextureType2DLayered == p->dstArray->textureType) - GetImageInfo(HSA_EXT_IMAGE_GEOMETRY_2DA, width, height, 0, desc, imageInfo, depth); - else - GetImageInfo(HSA_EXT_IMAGE_GEOMETRY_3D, width, height, depth, desc, imageInfo); - dstPitch = imageInfo.size/(height == 0 ? 1 : height)/(depth == 0 ? 1 : depth); - } else { - depth = p->Depth; - height = p->Height; - widthInBytes = p->WidthInBytes; - width = p->dstArray->width; - hsa_ext_image_channel_order_t channelOrder; - switch(p->dstArray->NumChannels) { - case 2: - channelOrder = HSA_EXT_IMAGE_CHANNEL_ORDER_RG; - break; - case 3: - channelOrder = HSA_EXT_IMAGE_CHANNEL_ORDER_RGB; - break; - case 4: - channelOrder = HSA_EXT_IMAGE_CHANNEL_ORDER_RGBA; - break; - case 1: - default: - channelOrder = HSA_EXT_IMAGE_CHANNEL_ORDER_R; - break; - } - hsa_ext_image_channel_type_t channelType; - e = ihipArrayToImageFormat(p->dstArray->Format,channelType); - srcPitch = p->srcPitch; - srcPtr = (void*)p->srcHost; - ySize = p->srcHeight; - dstPtr = p->dstArray->data; - hsa_ext_image_data_info_t imageInfo; - if(hipTextureType2DLayered == p->dstArray->textureType) - GetImageInfo(HSA_EXT_IMAGE_GEOMETRY_2DA, width, height, 0, channelOrder, channelType, imageInfo, depth); - else - GetImageInfo(HSA_EXT_IMAGE_GEOMETRY_3D, width, height, depth, channelOrder, channelType, imageInfo); - dstPitch = imageInfo.size/(height == 0 ? 1 : height)/(depth == 0 ? 1 : depth); + if ((p->dstArray->isDrv == true) ||( p->dstPtr.ptr!= nullptr)){ + return hipErrorInvalidValue; + } + // Array destination + dstByteSize = getByteSizeFromFormat(p->dstArray->desc); + hipChannelFormatDesc desc; + desc = p->dstArray->desc; + dstPtr = p->dstArray->data; + dstWidth = p->dstArray->width; + dstHeight = p->dstArray->height; + dstDepth = p->dstArray->depth; + dstPitch = dstByteSize * alignUp(dstWidth, IMAGE_PITCH_ALIGNMENT); + if(!copyWidthUpdate) { + copyWidth = copyWidth * dstByteSize; + copyWidthUpdate = true; } } else { - // Non array destination - depth = p->extent.depth; - height = p->extent.height; - widthInBytes = p->extent.width; - srcPitch = p->srcPtr.pitch; - srcPtr = p->srcPtr.ptr; + //Non Array destination dstPtr = p->dstPtr.ptr; - ySize = p->srcPtr.ysize; + dstWidth = p->dstPtr.xsize; + dstHeight = p->dstPtr.ysize; dstPitch = p->dstPtr.pitch; } + if (p->srcArray != nullptr) { + if ((p->srcArray->isDrv == true) ||( p->srcPtr.ptr!= nullptr)){ + return hipErrorInvalidValue; + } + // Array source + srcByteSize = getByteSizeFromFormat(p->srcArray->desc); + hipChannelFormatDesc desc; + desc = p->srcArray->desc; + srcPtr = p->srcArray->data; + srcWidth = p->srcArray->width; + srcHeight = p->srcArray->height; + srcDepth = p->srcArray->depth; + srcPitch = srcByteSize * alignUp(srcWidth, IMAGE_PITCH_ALIGNMENT); + if(!copyWidthUpdate) { + copyWidth = copyWidth * srcByteSize; + copyWidthUpdate = true; + } + } else { + //Non Array source + srcPtr = p->srcPtr.ptr; + srcWidth = p->srcPtr.xsize; + srcHeight = p->srcPtr.ysize; + srcPitch = p->srcPtr.pitch; + } + stream = ihipSyncAndResolveStream(stream); try { - if((widthInBytes == dstPitch) && (widthInBytes == srcPitch)) { + if((copyWidth == dstPitch) && (copyWidth == srcPitch)&& (copyHeight == dstHeight) &&(copyHeight == srcHeight)) { if(isAsync) - stream->locked_copyAsync((void*)dstPtr, (void*)srcPtr, widthInBytes*height*depth, p->kind); + stream->locked_copyAsync((void*)dstPtr, (void*)srcPtr, copyWidth*copyHeight*copyDepth, p->kind); else - stream->locked_copySync((void*)dstPtr, (void*)srcPtr, widthInBytes*height*depth, p->kind, false); + stream->locked_copySync((void*)dstPtr, (void*)srcPtr, copyWidth*copyHeight*copyDepth, p->kind, false); } else { - for (int i = 0; i < depth; i++) { - for (int j = 0; j < height; j++) { - // TODO: p->srcPos or p->dstPos are not 0. + for (int i = 0; i < copyDepth; i++) { + for (int j = 0; j < copyHeight; j++) { unsigned char* src = - (unsigned char*)srcPtr + i * ySize * srcPitch + j * srcPitch; + (unsigned char*)srcPtr + (i + srcZoffset) * srcHeight * srcPitch + (j + srcYoffset) * srcPitch + srcXoffset; unsigned char* dst = - (unsigned char*)dstPtr + i * height * dstPitch + j * dstPitch; + (unsigned char*)dstPtr + (i + dstZoffset) * dstHeight * dstPitch + (j + dstYoffset) * dstPitch + dstXoffset; if(isAsync) - stream->locked_copyAsync(dst, src, widthInBytes, p->kind); + stream->locked_copyAsync(dst, src, copyWidth, p->kind); else - stream->locked_copySync(dst, src, widthInBytes, p->kind); + stream->locked_copySync(dst, src, copyWidth, p->kind); } } } diff --git a/hipamd/src/hip_module.cpp b/hipamd/src/hip_module.cpp index 8901f2c4aa..b527d6ccb0 100644 --- a/hipamd/src/hip_module.cpp +++ b/hipamd/src/hip_module.cpp @@ -134,10 +134,10 @@ extern hipError_t ihipGetDeviceProperties(hipDeviceProp_t* props, int device); return ihipLogStatus(hipStatus); \ } -hipError_t ihipModuleLaunchKernel(TlsData *tls, hipFunction_t f, uint32_t gridSizeX, - uint32_t gridSizeY, uint32_t gridSizeZ, - uint32_t blockSizeX, uint32_t blockSizeY, - uint32_t blockSizeZ, size_t sharedMemBytes, +hipError_t ihipModuleLaunchKernel(TlsData *tls, hipFunction_t f, uint32_t globalWorkSizeX, + uint32_t globalWorkSizeY, uint32_t globalWorkSizeZ, + uint32_t localWorkSizeX, uint32_t localWorkSizeY, + uint32_t localWorkSizeZ, size_t sharedMemBytes, hipStream_t hStream, void** kernelParams, void** extra, hipEvent_t startEvent, hipEvent_t stopEvent, uint32_t flags, bool isStreamLocked = 0, void** impCoopParams = 0) { @@ -146,14 +146,6 @@ hipError_t ihipModuleLaunchKernel(TlsData *tls, hipFunction_t f, uint32_t gridSi auto ctx = ihipGetTlsDefaultCtx(); hipError_t ret = hipSuccess; - size_t globalWorkSizeX = (size_t)gridSizeX * (size_t)blockSizeX; - size_t globalWorkSizeY = (size_t)gridSizeY * (size_t)blockSizeY; - size_t globalWorkSizeZ = (size_t)gridSizeZ * (size_t)blockSizeZ; - if(globalWorkSizeX > UINT32_MAX || globalWorkSizeY > UINT32_MAX || globalWorkSizeZ > UINT32_MAX) - { - return hipErrorInvalidConfiguration; - } - if (ctx == nullptr) { ret = hipErrorInvalidDevice; @@ -211,8 +203,8 @@ hipError_t ihipModuleLaunchKernel(TlsData *tls, hipFunction_t f, uint32_t gridSi lp.dynamic_group_mem_bytes = sharedMemBytes; // TODO - this should be part of preLaunchKernel. hStream = ihipPreLaunchKernel( - hStream, dim3(globalWorkSizeX/blockSizeX, globalWorkSizeY/blockSizeY, globalWorkSizeZ/blockSizeZ), - dim3(blockSizeX, blockSizeY, blockSizeZ), &lp, f->_name.c_str(), isStreamLocked); + hStream, dim3(globalWorkSizeX/localWorkSizeX, globalWorkSizeY/localWorkSizeY, globalWorkSizeZ/localWorkSizeZ), + dim3(localWorkSizeX, localWorkSizeY, localWorkSizeZ), &lp, f->_name.c_str(), isStreamLocked); hsa_kernel_dispatch_packet_t aql; @@ -221,9 +213,9 @@ hipError_t ihipModuleLaunchKernel(TlsData *tls, hipFunction_t f, uint32_t gridSi // aql.completion_signal._handle = 0; // aql.kernarg_address = 0; - aql.workgroup_size_x = blockSizeX; - aql.workgroup_size_y = blockSizeY; - aql.workgroup_size_z = blockSizeZ; + aql.workgroup_size_x = localWorkSizeX; + aql.workgroup_size_y = localWorkSizeY; + aql.workgroup_size_z = localWorkSizeZ; aql.grid_size_x = globalWorkSizeX; aql.grid_size_y = globalWorkSizeY; aql.grid_size_z = globalWorkSizeZ; @@ -283,8 +275,17 @@ hipError_t hipModuleLaunchKernel(hipFunction_t f, uint32_t gridDimX, uint32_t gr void** kernelParams, void** extra) { HIP_INIT_API(hipModuleLaunchKernel, f, gridDimX, gridDimY, gridDimZ, blockDimX, blockDimY, blockDimZ, sharedMemBytes, hStream, kernelParams, extra); + + size_t globalWorkSizeX = (size_t)gridDimX * (size_t)blockDimX; + size_t globalWorkSizeY = (size_t)gridDimY * (size_t)blockDimY; + size_t globalWorkSizeZ = (size_t)gridDimZ * (size_t)blockDimZ; + if(globalWorkSizeX > UINT32_MAX || globalWorkSizeY > UINT32_MAX || globalWorkSizeZ > UINT32_MAX) + { + return hipErrorInvalidConfiguration; + } + return ihipLogStatus(ihipModuleLaunchKernel(tls, - f, gridDimX, gridDimY, gridDimZ, blockDimX, blockDimY, + f, globalWorkSizeX, globalWorkSizeY, globalWorkSizeZ, blockDimX, blockDimY, blockDimZ, sharedMemBytes, hStream, kernelParams, extra, nullptr, nullptr, 0)); } @@ -297,11 +298,8 @@ hipError_t hipExtModuleLaunchKernel(hipFunction_t f, uint32_t globalWorkSizeX, HIP_INIT_API(hipExtModuleLaunchKernel, f, globalWorkSizeX, globalWorkSizeY, globalWorkSizeZ, localWorkSizeX, localWorkSizeY, localWorkSizeZ, sharedMemBytes, hStream, kernelParams, extra); - if(localWorkSizeX == 0 || localWorkSizeY == 0 || localWorkSizeZ == 0) - return hipErrorInvalidValue; - return ihipLogStatus(ihipModuleLaunchKernel(tls, - f, globalWorkSizeX/localWorkSizeX, globalWorkSizeY/localWorkSizeY, globalWorkSizeZ/localWorkSizeZ, localWorkSizeX, localWorkSizeY, + f, globalWorkSizeX, globalWorkSizeY, globalWorkSizeZ, localWorkSizeX, localWorkSizeY, localWorkSizeZ, sharedMemBytes, hStream, kernelParams, extra, startEvent, stopEvent, flags)); } @@ -314,11 +312,8 @@ hipError_t hipHccModuleLaunchKernel(hipFunction_t f, uint32_t globalWorkSizeX, HIP_INIT_API(hipHccModuleLaunchKernel, f, globalWorkSizeX, globalWorkSizeY, globalWorkSizeZ, localWorkSizeX, localWorkSizeY, localWorkSizeZ, sharedMemBytes, hStream, kernelParams, extra); - if(localWorkSizeX == 0 || localWorkSizeY == 0 || localWorkSizeZ == 0) - return hipErrorInvalidValue; - return ihipLogStatus(ihipModuleLaunchKernel(tls, - f, globalWorkSizeX/localWorkSizeX, globalWorkSizeY/localWorkSizeY, globalWorkSizeZ/localWorkSizeZ, localWorkSizeX, localWorkSizeY, + f, globalWorkSizeX, globalWorkSizeY, globalWorkSizeZ, localWorkSizeX, localWorkSizeY, localWorkSizeZ, sharedMemBytes, hStream, kernelParams, extra, startEvent, stopEvent, 0)); } @@ -364,14 +359,26 @@ hipError_t ihipExtLaunchMultiKernelMultiDevice(hipLaunchParams* launchParamsList } GET_TLS(); + + size_t globalWorkSizeX = 0, globalWorkSizeY = 0, globalWorkSizeZ = 0; + // launch kernels for each device for (int i = 0; i < numDevices; ++i) { const hipLaunchParams& lp = launchParamsList[i]; + globalWorkSizeX = (size_t)lp.gridDim.x * (size_t)lp.blockDim.x; + globalWorkSizeY = (size_t)lp.gridDim.y * (size_t)lp.blockDim.y; + globalWorkSizeZ = (size_t)lp.gridDim.z * (size_t)lp.blockDim.z; + + if(globalWorkSizeX > UINT32_MAX || globalWorkSizeY > UINT32_MAX || globalWorkSizeZ > UINT32_MAX) + { + return hipErrorInvalidConfiguration; + } + result = ihipModuleLaunchKernel(tls, kds[i], - lp.gridDim.x, - lp.gridDim.y, - lp.gridDim.z, + lp.gridDim.x * lp.blockDim.x, + lp.gridDim.y * lp.blockDim.y, + lp.gridDim.z * lp.blockDim.z, lp.blockDim.x, lp.blockDim.y, lp.blockDim.z, lp.sharedMem, lp.stream, lp.args, nullptr, nullptr, nullptr, 0, @@ -424,6 +431,14 @@ hipError_t ihipLaunchCooperativeKernel(const void* f, dim3 gridDim, return hipErrorInvalidConfiguration; } + size_t globalWorkSizeX = (size_t)gridDim.x * (size_t)blockDimX.x; + size_t globalWorkSizeY = (size_t)gridDim.y * (size_t)blockDimX.y; + size_t globalWorkSizeZ = (size_t)gridDim.z * (size_t)blockDimX.z; + if(globalWorkSizeX > UINT32_MAX || globalWorkSizeY > UINT32_MAX || globalWorkSizeZ > UINT32_MAX) + { + return hipErrorInvalidConfiguration; + } + // Prepare the kernel descriptor for initializing the GWS hipFunction_t gwsKD = ps.kernel_descriptor( reinterpret_cast(&init_gws), @@ -483,9 +498,9 @@ hipError_t ihipLaunchCooperativeKernel(const void* f, dim3 gridDim, // launch the main kernel result = ihipModuleLaunchKernel(tls, kd, - gridDim.x, - gridDim.y, - gridDim.z, + gridDim.x * blockDimX.x, + gridDim.y * blockDimX.y, + gridDim.z * blockDimX.z, blockDimX.x, blockDimX.y, blockDimX.z, sharedMemBytes, stream, kernelParams, nullptr, nullptr, nullptr, 0, true, impCoopParams); @@ -620,6 +635,8 @@ hipError_t ihipLaunchCooperativeKernelMultiDevice(hipLaunchParams* launchParamsL void* impCoopParams[1]; ulong prev_sum = 0; + + size_t globalWorkSizeX = 0, globalWorkSizeY = 0, globalWorkSizeZ = 0; // launch the main kernels for each device for (int i = 0; i < numDevices; ++i) { const hipLaunchParams& lp = launchParamsList[i]; @@ -636,10 +653,18 @@ hipError_t ihipLaunchCooperativeKernelMultiDevice(hipLaunchParams* launchParamsL impCoopParams[0] = &mg_info_ptr[i]; + globalWorkSizeX = (size_t)lp.gridDim.x * (size_t)lp.blockDim.x; + globalWorkSizeY = (size_t)lp.gridDim.y * (size_t)lp.blockDim.y; + globalWorkSizeZ = (size_t)lp.gridDim.z * (size_t)lp.blockDim.z; + if(globalWorkSizeX > UINT32_MAX || globalWorkSizeY > UINT32_MAX || globalWorkSizeZ > UINT32_MAX) + { + return hipErrorInvalidConfiguration; + } + result = ihipModuleLaunchKernel(tls, kds[i], - lp.gridDim.x, - lp.gridDim.y, - lp.gridDim.z, + lp.gridDim.x * lp.blockDim.x, + lp.gridDim.y * lp.blockDim.y, + lp.gridDim.z * lp.blockDim.z, lp.blockDim.x, lp.blockDim.y, lp.blockDim.z, lp.sharedMem, lp.stream, lp.args, nullptr, nullptr, nullptr, 0, @@ -1469,6 +1494,15 @@ hipError_t hipOccupancyMaxActiveBlocksPerMultiprocessor( tls, numBlocks, f, blockSize, dynSharedMemPerBlk)); } +hipError_t hipDrvOccupancyMaxActiveBlocksPerMultiprocessor( + int* numBlocks, hipFunction_t f, int blockSize, size_t dynSharedMemPerBlk) +{ + HIP_INIT_API(hipDrvOccupancyMaxActiveBlocksPerMultiprocessor, numBlocks, f, blockSize, dynSharedMemPerBlk); + + return ihipLogStatus(ihipOccupancyMaxActiveBlocksPerMultiprocessor( + tls, (uint32_t*) numBlocks, f, blockSize, dynSharedMemPerBlk)); +} + hipError_t hipOccupancyMaxActiveBlocksPerMultiprocessorWithFlags( uint32_t* numBlocks, hipFunction_t f, uint32_t blockSize, size_t dynSharedMemPerBlk, unsigned int flags) @@ -1479,6 +1513,15 @@ hipError_t hipOccupancyMaxActiveBlocksPerMultiprocessorWithFlags( tls, numBlocks, f, blockSize, dynSharedMemPerBlk)); } +hipError_t hipDrvOccupancyMaxActiveBlocksPerMultiprocessorWithFlags( + int* numBlocks, hipFunction_t f, int blockSize, size_t dynSharedMemPerBlk, + unsigned int flags) +{ + HIP_INIT_API(hipDrvOccupancyMaxActiveBlocksPerMultiprocessorWithFlags, numBlocks, f, blockSize, dynSharedMemPerBlk, flags); + return ihipLogStatus(ihipOccupancyMaxActiveBlocksPerMultiprocessor( + tls, (uint32_t*) numBlocks, f, blockSize, dynSharedMemPerBlk)); +} + hipError_t hipLaunchKernel( const void* func_addr, dim3 numBlocks, dim3 dimBlocks, void** args, size_t sharedMemBytes, hipStream_t stream) diff --git a/hipamd/src/hip_peer.cpp b/hipamd/src/hip_peer.cpp index 7781af1dbe..8fd66a52bb 100644 --- a/hipamd/src/hip_peer.cpp +++ b/hipamd/src/hip_peer.cpp @@ -128,17 +128,24 @@ hipError_t ihipEnablePeerAccess(TlsData* tls, hipCtx_t peerCtx, unsigned int fla if (thisCtx == peerCtx) { err = hipErrorInvalidDevice; // Can't enable peer access to self. } else if ((thisCtx != NULL) && (peerCtx != NULL)) { - LockedAccessor_CtxCrit_t peerCrit(peerCtx->criticalData()); - // Add thisCtx to peerCtx's access list so that new allocations on peer will be made - // visible to this device: - bool isNewPeer = peerCrit->addPeerWatcher(peerCtx, thisCtx); - if (isNewPeer) { - tprintf(DB_MEM, "device=%s can now see all memory allocated on peer=%s\n", - thisCtx->toString().c_str(), peerCtx->toString().c_str()); - am_memtracker_update_peers(peerCtx->getDevice()->_acc, peerCrit->peerCnt(), - peerCrit->peerAgents()); + + int canAccess = 0; + if ((hipSuccess != ihipDeviceCanAccessPeer(&canAccess,thisCtx,peerCtx)) || (canAccess == 0)){ + tprintf(DB_MEM, "device=%s can't access peer=%s\n",thisCtx->toString().c_str(), peerCtx->toString().c_str()); + err = hipErrorInvalidDevice; } else { - err = hipErrorPeerAccessAlreadyEnabled; + LockedAccessor_CtxCrit_t peerCrit(peerCtx->criticalData()); + // Add thisCtx to peerCtx's access list so that new allocations on peer will be made + // visible to this device: + bool isNewPeer = peerCrit->addPeerWatcher(peerCtx, thisCtx); + if (isNewPeer) { + tprintf(DB_MEM, "device=%s can now see all memory allocated on peer=%s\n", + thisCtx->toString().c_str(), peerCtx->toString().c_str()); + am_memtracker_update_peers(peerCtx->getDevice()->_acc, peerCrit->peerCnt(), + peerCrit->peerAgents()); + } else { + err = hipErrorPeerAccessAlreadyEnabled; + } } } else { err = hipErrorInvalidDevice; diff --git a/hipamd/src/hiprtc.cpp b/hipamd/src/hiprtc.cpp index 8dcb944f72..e9a516c339 100644 --- a/hipamd/src/hiprtc.cpp +++ b/hipamd/src/hiprtc.cpp @@ -50,7 +50,7 @@ THE SOFTWARE. #include #include -const char* hiprtcGetErrorString(hiprtcResult x) +extern "C" const char* hiprtcGetErrorString(hiprtcResult x) { switch (x) { case HIPRTC_SUCCESS: @@ -95,6 +95,21 @@ inline bool fileExists (const std::string& name) { } } // namespace hip_impl +namespace +{ + char* demangle(const char* x) + { + if (!x) return nullptr; + + int s{}; + char* tmp = abi::__cxa_demangle(x, nullptr, nullptr, &s); + + if (s != 0) return nullptr; + + return tmp; + } +} // Unnamed namespace. + namespace { struct Symbol { @@ -158,7 +173,7 @@ struct _hiprtcProgram { { using namespace std; - char* demangled = hip_impl::demangle(name.c_str()); + char* demangled = demangle(name.c_str()); name.assign(demangled == nullptr ? "" : demangled); free(demangled); @@ -352,7 +367,7 @@ namespace } } // Unnamed namespace. -hiprtcResult hiprtcAddNameExpression(hiprtcProgram p, const char* n) +extern "C" hiprtcResult hiprtcAddNameExpression(hiprtcProgram p, const char* n) { if (!n) return HIPRTC_ERROR_INVALID_INPUT; if (!isValidProgram(p)) return HIPRTC_ERROR_INVALID_PROGRAM; @@ -413,21 +428,6 @@ namespace }; } // Unnamed namespace. -namespace hip_impl -{ - char* demangle(const char* x) - { - if (!x) return nullptr; - - int s{}; - char* tmp = abi::__cxa_demangle(x, nullptr, nullptr, &s); - - if (s != 0) return nullptr; - - return tmp; - } -} // Namespace hip_impl. - namespace { const std::string& defaultTarget() @@ -492,7 +492,7 @@ namespace } } // Unnamed namespace. -hiprtcResult hiprtcCompileProgram(hiprtcProgram p, int n, const char** o) +extern "C" hiprtcResult hiprtcCompileProgram(hiprtcProgram p, int n, const char** o) { using namespace std; @@ -530,7 +530,7 @@ hiprtcResult hiprtcCompileProgram(hiprtcProgram p, int n, const char** o) return HIPRTC_SUCCESS; } -hiprtcResult hiprtcCreateProgram(hiprtcProgram* p, const char* src, +extern "C" hiprtcResult hiprtcCreateProgram(hiprtcProgram* p, const char* src, const char* name, int n, const char** hdrs, const char** incs) { @@ -548,14 +548,14 @@ hiprtcResult hiprtcCreateProgram(hiprtcProgram* p, const char* src, return HIPRTC_SUCCESS; } -hiprtcResult hiprtcDestroyProgram(hiprtcProgram* p) +extern "C" hiprtcResult hiprtcDestroyProgram(hiprtcProgram* p) { if (!p) return HIPRTC_SUCCESS; return _hiprtcProgram::destroy(*p); } -hiprtcResult hiprtcGetLoweredName(hiprtcProgram p, const char* n, +extern "C" hiprtcResult hiprtcGetLoweredName(hiprtcProgram p, const char* n, const char** ln) { using namespace std; @@ -576,7 +576,7 @@ hiprtcResult hiprtcGetLoweredName(hiprtcProgram p, const char* n, return HIPRTC_SUCCESS; } -hiprtcResult hiprtcGetProgramLog(hiprtcProgram p, char* l) +extern "C" hiprtcResult hiprtcGetProgramLog(hiprtcProgram p, char* l) { if (!l) return HIPRTC_ERROR_INVALID_INPUT; if (!isValidProgram(p)) return HIPRTC_ERROR_INVALID_PROGRAM; @@ -588,7 +588,7 @@ hiprtcResult hiprtcGetProgramLog(hiprtcProgram p, char* l) return HIPRTC_SUCCESS; } -hiprtcResult hiprtcGetProgramLogSize(hiprtcProgram p, std::size_t* sz) +extern "C" hiprtcResult hiprtcGetProgramLogSize(hiprtcProgram p, std::size_t* sz) { if (!sz) return HIPRTC_ERROR_INVALID_INPUT; if (!isValidProgram(p)) return HIPRTC_ERROR_INVALID_PROGRAM; @@ -599,7 +599,7 @@ hiprtcResult hiprtcGetProgramLogSize(hiprtcProgram p, std::size_t* sz) return HIPRTC_SUCCESS; } -hiprtcResult hiprtcGetCode(hiprtcProgram p, char* c) +extern "C" hiprtcResult hiprtcGetCode(hiprtcProgram p, char* c) { if (!c) return HIPRTC_ERROR_INVALID_INPUT; if (!isValidProgram(p)) return HIPRTC_ERROR_INVALID_PROGRAM; @@ -610,7 +610,7 @@ hiprtcResult hiprtcGetCode(hiprtcProgram p, char* c) return HIPRTC_SUCCESS; } -hiprtcResult hiprtcGetCodeSize(hiprtcProgram p, std::size_t* sz) +extern "C" hiprtcResult hiprtcGetCodeSize(hiprtcProgram p, std::size_t* sz) { if (!sz) return HIPRTC_ERROR_INVALID_INPUT; if (!isValidProgram(p)) return HIPRTC_ERROR_INVALID_PROGRAM; @@ -620,3 +620,15 @@ hiprtcResult hiprtcGetCodeSize(hiprtcProgram p, std::size_t* sz) return HIPRTC_SUCCESS; } + +extern "C" hiprtcResult hiprtcVersion(int* major, int* minor) +{ + if (major == nullptr || minor == nullptr) { + return HIPRTC_ERROR_INVALID_INPUT; + } + + *major = 9; + *minor = 0; + + return HIPRTC_SUCCESS; +} diff --git a/hipamd/src/program_state.inl b/hipamd/src/program_state.inl index 2e959a9928..d77abd9674 100644 --- a/hipamd/src/program_state.inl +++ b/hipamd/src/program_state.inl @@ -26,6 +26,7 @@ #include #include #include +#include #include #include #include @@ -202,7 +203,7 @@ public: std::function>; std::pair< std::mutex, - std::vector>> code_readers; + std::deque>> code_readers; program_state_impl() { // Create placeholder for each agent for the per-agent members. @@ -244,7 +245,8 @@ public: if (!valid(tmp)) break; for (auto&& bundle : bundles(tmp)) { - impl.code_object_blobs.second[elf][triple_to_hsa_isa(bundle.triple)].push_back(bundle.blob); + if(bundle.blob.size()) + impl.code_object_blobs.second[elf][triple_to_hsa_isa(bundle.triple)].push_back(bundle.blob); } blob_it += tmp.bundled_code_size; @@ -418,13 +420,17 @@ public: decltype(code_readers.second)::iterator it; { std::lock_guard lck{code_readers.first}; - it = code_readers.second.emplace(code_readers.second.end(), - move(file), move(tmp)); + code_readers.second.emplace_back(move(file), move(tmp)); + it = std::prev(code_readers.second.end()); } auto check_hsa_error = [](hsa_status_t s) { if (s != HSA_STATUS_SUCCESS) { - hip_throw(std::runtime_error{"error when loading code object"}); + const char* hsa_err_msg; + hsa_status_string(s, &hsa_err_msg); + hip_throw(std::runtime_error{ + std::string("error when loading code object: ") + + hsa_err_msg}); } }; diff --git a/hipamd/tests/src/hipEnvVarDriver.cpp b/hipamd/tests/src/hipEnvVarDriver.cpp index 07379f0878..c970cb7674 100644 --- a/hipamd/tests/src/hipEnvVarDriver.cpp +++ b/hipamd/tests/src/hipEnvVarDriver.cpp @@ -34,42 +34,52 @@ OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWA using namespace std; -const string directed_dir = "directed_tests" + string(PATH_SEPERATOR_STR) + "hipEnvVar"; -const string dir = "." + string(PATH_SEPERATOR_STR) + "hipEnvVar"; +const string directed_dir = string(".") + PATH_SEPERATOR_STR + "directed_tests" + PATH_SEPERATOR_STR + "hipEnvVar"; +const string dir = string(".") + PATH_SEPERATOR_STR + "hipEnvVar"; -int getDeviceNumber() { - char buff[512]; - std::this_thread::sleep_for(std::chrono::milliseconds(10)); - FILE* in = popen((directed_dir + " -c").c_str(), "r"); - if(fgets(buff, 512, in) == NULL){ - pclose(in); - //Check at same level - in = popen((dir + " -c").c_str(), "r"); +int readHipEnvVar(string flags, char* buff){ + + std::cout << "\nFinding hipEnvVar in " << directed_dir << "...\n"; + FILE* directed_in = popen((directed_dir + flags).c_str(), "r"); + + if(fgets(buff, 512, directed_in) == NULL){ + std::cout << "Finding hipEnvVar in " << dir << "...\n"; + FILE* in = popen((dir + flags).c_str(), "r"); if(fgets(buff, 512, in) == NULL){ + pclose(directed_in); pclose(in); return 1; } + pclose(in); + } + std::cout << "hipEnvVar Found!\n"; + pclose(directed_in); + return 0; +} + +int getDeviceNumber(bool print_err=true) { + char buff[512]; + std::this_thread::sleep_for(std::chrono::milliseconds(10)); + + if (readHipEnvVar(string(" -c"), buff)){ + strncpy(buff, "1", 512); + if (print_err){ + std::cerr << "The system cannot find hipEnvVar, using 1 as number of devices\n"; + } + } + if (print_err) { + std::cout << buff; } - cout << buff; - pclose(in); return atoi(buff); } // Query the current device ID remotely to hipEnvVar void getDevicePCIBusNumRemote(int deviceID, char* pciBusID) { std::this_thread::sleep_for(std::chrono::milliseconds(10)); - FILE* in = popen((directed_dir + " -d " + std::to_string(deviceID)).c_str(), "r"); - if(fgets(pciBusID, 100, in) == NULL){ - pclose(in); - //Check at same level - in = popen((dir + " -d").c_str(), "r"); - if(fgets(pciBusID, 100, in) == NULL){ - pclose(in); - return; - } + if (readHipEnvVar((" -d " + std::to_string(deviceID)), pciBusID)){ + std::cerr << "The system cannot find hipEnvVar\n"; } cout << pciBusID; - pclose(in); return; } @@ -78,15 +88,15 @@ void getDevicePCIBusNum(int deviceID, char* pciBusID) { hipDevice_t deviceT; hipDeviceGet(&deviceT, deviceID); - memset(pciBusID, 0, 100); - hipDeviceGetPCIBusId(pciBusID, 100, deviceT); + memset(pciBusID, 0, 512); + hipDeviceGetPCIBusId(pciBusID, 512, deviceT); } int main() { unsetenv(HIP_VISIBLE_DEVICES_STR); unsetenv(CUDA_VISIBLE_DEVICES_STR); std::vector devPCINum; - char pciBusID[100]; + char pciBusID[512]; // collect the device pci bus ID for all devices int totalDeviceNum = getDeviceNumber(); std::cout << "The total number of available devices is " << totalDeviceNum << std::endl @@ -116,27 +126,27 @@ int main() { // check when set an invalid device number setenv("HIP_VISIBLE_DEVICES", "1000,0,1", 1); setenv("CUDA_VISIBLE_DEVICES", "1000,0,1", 1); - assert(getDeviceNumber() == 0); + assert(getDeviceNumber(false) == 0); if (totalDeviceNum > 2) { setenv("HIP_VISIBLE_DEVICES", "0,1,1000,2", 1); setenv("CUDA_VISIBLE_DEVICES", "0,1,1000,2", 1); - assert(getDeviceNumber() == 2); + assert(getDeviceNumber(false) == 2); setenv("HIP_VISIBLE_DEVICES", "0,1,2", 1); setenv("CUDA_VISIBLE_DEVICES", "0,1,2", 1); - assert(getDeviceNumber() == 3); + assert(getDeviceNumber(false) == 3); // test if CUDA_VISIBLE_DEVICES will be accepted by the runtime unsetenv(HIP_VISIBLE_DEVICES_STR); unsetenv(CUDA_VISIBLE_DEVICES_STR); setenv("CUDA_VISIBLE_DEVICES", "0,1,2", 1); - assert(getDeviceNumber() == 3); + assert(getDeviceNumber(false) == 3); } setenv("HIP_VISIBLE_DEVICES", "-100,0,1", 1); setenv("CUDA_VISIBLE_DEVICES", "-100,0,1", 1); - assert(getDeviceNumber() == 0); + assert(getDeviceNumber(false) == 0); std::cout << "PASSED" << std::endl; return 0; -} +} \ No newline at end of file diff --git a/hipamd/tests/src/hiprtc/hiprtcGetTypeName.cpp b/hipamd/tests/src/hiprtc/hiprtcGetTypeName.cpp deleted file mode 100644 index 812229f81f..0000000000 --- a/hipamd/tests/src/hiprtc/hiprtcGetTypeName.cpp +++ /dev/null @@ -1,147 +0,0 @@ -/* -Copyright (c) 2015 - present 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. -*/ -/* HIT_START - * BUILD: %t %s ../test_common.cpp LINK_OPTIONS hiprtc EXCLUDE_HIP_PLATFORM nvcc - * TEST: %t - * HIT_END - */ - -#include - -#define HIPRTC_GET_TYPE_NAME -#include -#include - -#include -#include -#include - -static constexpr auto gpu_program{ -R"( -#include - -namespace N1 { struct S1_t { int i; double d; }; } -template -__global__ void f3(int *result) { *result = sizeof(T); } -)"}; - -// note: this structure is also defined in GPU code string. Should ideally -// be in a header file included by both GPU code string and by CPU code. -namespace N1 { struct S1_t { int i; double d; }; }; - -template -std::string getKernelNameForType(void) -{ - std::string type_name; - hiprtcGetTypeName(&type_name); - return std::string{"f3<"} + type_name + '>'; -} - -int main() -{ - using namespace std; - - hiprtcProgram prog; - hiprtcCreateProgram(&prog, gpu_program, "gpu_program.cu", 0, nullptr, - nullptr); - - vector name_vec; - vector expected_result; - - name_vec.push_back(getKernelNameForType()); - expected_result.push_back(sizeof(int)); - name_vec.push_back(getKernelNameForType()); - expected_result.push_back(sizeof(double)); - name_vec.push_back(getKernelNameForType()); - expected_result.push_back(sizeof(N1::S1_t)); - - for (auto&& x : name_vec) hiprtcAddNameExpression(prog, x.c_str()); - - hipDeviceProp_t props; - int device = 0; - hipGetDeviceProperties(&props, device); - std::string gfxName = "gfx" + std::to_string(props.gcnArch); - std::string sarg = "--gpu-architecture=" + gfxName; - const char* options[] = { - sarg.c_str() - }; - - hiprtcResult compileResult = hiprtcCompileProgram(prog, 1, options); - - size_t logSize; - hiprtcGetProgramLogSize(prog, &logSize); - - if (logSize) { - string log(logSize, '\0'); - hiprtcGetProgramLog(prog, &log[0]); - - cout << log << '\n'; - } - - if (compileResult != HIPRTC_SUCCESS) { failed("Compilation failed."); } - - size_t codeSize; - hiprtcGetCodeSize(prog, &codeSize); - - vector code(codeSize); - hiprtcGetCode(prog, code.data()); - - hipModule_t module; - hipModuleLoadDataEx(&module, code.data(), 0, nullptr, nullptr); - - hipDeviceptr_t dResult; - int hResult = 0; - hipMalloc(&dResult, sizeof(hResult)); - hipMemcpyHtoD(dResult, &hResult, sizeof(hResult)); - - for (size_t i = 0; i < name_vec.size(); ++i) { - const char *name; - hiprtcGetLoweredName(prog, name_vec[i].c_str(), &name); - - hipFunction_t kernel; - hipModuleGetFunction(&kernel, module, name); - - struct { hipDeviceptr_t a_; } args{dResult}; - - auto size = sizeof(args); - void* config[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, &args, - HIP_LAUNCH_PARAM_BUFFER_SIZE, &size, - HIP_LAUNCH_PARAM_END}; - - hipModuleLaunchKernel(kernel, - 1, 1, 1, - 1, 1, 1, - 0, nullptr, - nullptr, config); - - hipMemcpyDtoH(&hResult, dResult, sizeof(hResult)); - - if (expected_result[i] != hResult) { failed("Validation failed."); } - } - - hipFree(dResult); - hipModuleUnload(module); - - hiprtcDestroyProgram(&prog); - - passed(); -} diff --git a/hipamd/tests/src/kernel/hipShflTests.cpp b/hipamd/tests/src/kernel/hipShflTests.cpp new file mode 100644 index 0000000000..9b1cc73248 --- /dev/null +++ b/hipamd/tests/src/kernel/hipShflTests.cpp @@ -0,0 +1,130 @@ +/* +Copyright (c) 2015-present 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. +*/ +/* HIT_START + * BUILD: %t %s ../test_common.cpp + * TEST: %t + * HIT_END + */ + +#include +#include +#include "test_common.h" + +#define WIDTH 4 + +#define NUM (WIDTH * WIDTH) + +#define THREADS_PER_BLOCK_X 4 +#define THREADS_PER_BLOCK_Y 4 +#define THREADS_PER_BLOCK_Z 1 + +// Device (Kernel) function, it must be void +template +__global__ void matrixTranspose(T* out, T* in, const int width) { + int x = hipBlockDim_x * hipBlockIdx_x + hipThreadIdx_x; + T val = in[x]; + for (int i = 0; i < width; i++) { + for (int j = 0; j < width; j++) out[i * width + j] = __shfl(val, j * width + i); + } +} + +// CPU implementation of matrix transpose +template +void matrixTransposeCPUReference(T* output, T* input, const unsigned int width) { + for (unsigned int j = 0; j < width; j++) { + for (unsigned int i = 0; i < width; i++) { + output[i * width + j] = input[j * width + i]; + } + } +} + +template +void runTest() { + T* Matrix; + T* TransposeMatrix; + T* cpuTransposeMatrix; + + T* gpuMatrix; + T* gpuTransposeMatrix; + + hipDeviceProp_t devProp; + hipGetDeviceProperties(&devProp, 0); + + int i; + int errors; + + Matrix = (T*)malloc(NUM * sizeof(T)); + TransposeMatrix = (T*)malloc(NUM * sizeof(T)); + cpuTransposeMatrix = (T*)malloc(NUM * sizeof(T)); + + // initialize the input data + for (i = 0; i < NUM; i++) { + Matrix[i] = (T)i * 10l; + } + + // allocate the memory on the device side + hipMalloc((void**)&gpuMatrix, NUM * sizeof(T)); + hipMalloc((void**)&gpuTransposeMatrix, NUM * sizeof(T)); + + // Memory transfer from host to device + hipMemcpy(gpuMatrix, Matrix, NUM * sizeof(T), hipMemcpyHostToDevice); + + // Lauching kernel from host + hipLaunchKernelGGL(matrixTranspose, dim3(1), dim3(THREADS_PER_BLOCK_X * THREADS_PER_BLOCK_Y), 0, 0, + gpuTransposeMatrix, gpuMatrix, WIDTH); + + // Memory transfer from device to host + hipMemcpy(TransposeMatrix, gpuTransposeMatrix, NUM * sizeof(T), hipMemcpyDeviceToHost); + + // CPU MatrixTranspose computation + matrixTransposeCPUReference(cpuTransposeMatrix, Matrix, WIDTH); + + // verify the results + errors = 0; + double eps = 1.0E-6; + for (i = 0; i < NUM; i++) { + if (TransposeMatrix[i] != cpuTransposeMatrix[i]) { + errors++; + } + } + + // free the resources on device side + hipFree(gpuMatrix); + hipFree(gpuTransposeMatrix); + + // free the resources on host side + free(Matrix); + free(TransposeMatrix); + free(cpuTransposeMatrix); + + if (errors != 0) { + failed("Mismatch present"); + } +} + +int main() { + runTest(); + runTest(); + runTest(); + runTest(); + passed(); +} diff --git a/hipamd/tests/src/kernel/hipShflUpDownTest.cpp b/hipamd/tests/src/kernel/hipShflUpDownTest.cpp new file mode 100644 index 0000000000..553087ce45 --- /dev/null +++ b/hipamd/tests/src/kernel/hipShflUpDownTest.cpp @@ -0,0 +1,102 @@ +/* +Copyright (c) 2015-present 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. +*/ +/* HIT_START + * BUILD: %t %s ../test_common.cpp + * TEST: %t + * HIT_END + */ + +#include +#include +#include "test_common.h" + +template +__global__ void shflDownSum(T* a, int size) { + T val = a[threadIdx.x]; + for (int i = size / 2; i > 0; i /= 2) { + val += __shfl_down(val, i, size); + } + a[threadIdx.x] = val; +} + +template +__global__ void shflUpSum(T* a, int size) { + T val = a[threadIdx.x]; + for (int i = size / 2; i > 0; i /= 2) { + val += __shfl_up(val, i, size); + } + a[threadIdx.x] = val; +} + +template +void runTestShflUp() { + const int size = 32; + T a[size]; + T cpuSum = 0; + for (int i = 0; i < size; i++) { + a[i] = i; + cpuSum += a[i]; + } + T* d_a; + hipMalloc(&d_a, sizeof(T) * size); + hipMemcpy(d_a, &a, sizeof(T) * size, hipMemcpyDefault); + hipLaunchKernelGGL(shflUpSum, 1, size, 0, 0, d_a, size); + hipMemcpy(&a, d_a, sizeof(T) * size, hipMemcpyDefault); + if (a[size - 1] != cpuSum) { + hipFree(d_a); + failed("Shfl Up Sum did not match."); + } + hipFree(d_a); +} + +template +void runTestShflDown() { + const int size = 32; + T a[size]; + T cpuSum = 0; + for (int i = 0; i < size; i++) { + a[i] = i; + cpuSum += a[i]; + } + T* d_a; + hipMalloc(&d_a, sizeof(T) * size); + hipMemcpy(d_a, &a, sizeof(T) * size, hipMemcpyDefault); + hipLaunchKernelGGL(shflDownSum, 1, size, 0, 0, d_a, size); + hipMemcpy(&a, d_a, sizeof(T) * size, hipMemcpyDefault); + if (a[0] != cpuSum) { + hipFree(d_a); + failed("Shfl Up Sum did not match."); + } + hipFree(d_a); +} +int main() { + runTestShflUp(); + runTestShflUp(); + runTestShflUp(); + runTestShflUp(); + + runTestShflDown(); + runTestShflDown(); + runTestShflDown(); + runTestShflDown(); + passed(); +} diff --git a/hipamd/tests/src/runtimeApi/memory/hipMemcpy3D.cpp b/hipamd/tests/src/runtimeApi/memory/hipMemcpy3D.cpp new file mode 100644 index 0000000000..255a3490b6 --- /dev/null +++ b/hipamd/tests/src/runtimeApi/memory/hipMemcpy3D.cpp @@ -0,0 +1,110 @@ +/* +Copyright (c) 2015 - present 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. +*/ + +/* HIT_START + * BUILD: %t %s ../../test_common.cpp + * TEST: %t + * HIT_END + */ + +#include "test_common.h" + +template +void runTest(int width,int height,int depth, hipChannelFormatKind formatKind) +{ + unsigned int size = width * height * depth * sizeof(T); + T* hData = (T*) malloc(size); + 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; + } + } + } + printf("test- sizeof(T) =%d\n", sizeof(T)); + hipChannelFormatDesc channelDesc = hipCreateChannelDesc(sizeof(T)*8, 0, 0, 0, formatKind); + hipArray *arr,*arr1; + + HIPCHECK(hipMalloc3DArray(&arr, &channelDesc, make_hipExtent(width, height, depth), hipArrayDefault)); + HIPCHECK(hipMalloc3DArray(&arr1, &channelDesc, make_hipExtent(width, height, depth), hipArrayDefault)); + hipMemcpy3DParms myparms = {0}; + myparms.srcPos = make_hipPos(0,0,0); + myparms.dstPos = make_hipPos(0,0,0); + myparms.srcPtr = make_hipPitchedPtr(hData, width * sizeof(T), width, height); + myparms.dstArray = arr; + myparms.extent = make_hipExtent(width , height, depth); +#ifdef __HIP_PLATFORM_NVCC__ + myparms.kind = cudaMemcpyHostToDevice; +#else + myparms.kind = hipMemcpyHostToDevice; +#endif + HIPCHECK(hipMemcpy3D(&myparms)); + HIPCHECK(hipDeviceSynchronize()); + //Array to Array + memset(&myparms,0x0, sizeof(hipMemcpy3DParms)); + myparms.srcPos = make_hipPos(0,0,0); + myparms.dstPos = make_hipPos(0,0,0); + myparms.srcArray = arr; + myparms.dstArray = arr1; + myparms.extent = make_hipExtent(width, height, depth); +#ifdef __HIP_PLATFORM_NVCC__ + myparms.kind = cudaMemcpyDeviceToDevice; +#else + myparms.kind = hipMemcpyDeviceToDevice; +#endif + HIPCHECK(hipMemcpy3D(&myparms)); + HIPCHECK(hipDeviceSynchronize()); + + T *hOutputData = (T*) malloc(size); + memset(hOutputData, 0, size); + //Device to host + memset(&myparms,0x0, sizeof(hipMemcpy3DParms)); + myparms.srcPos = make_hipPos(0,0,0); + myparms.dstPos = make_hipPos(0,0,0); + myparms.dstPtr = make_hipPitchedPtr(hOutputData, width * sizeof(T), width, height); + myparms.srcArray = arr1; + myparms.extent = make_hipExtent(width, height, depth); +#ifdef __HIP_PLATFORM_NVCC__ + myparms.kind = cudaMemcpyDeviceToHost; +#else + myparms.kind = hipMemcpyDeviceToHost; +#endif + HIPCHECK(hipMemcpy3D(&myparms)); + HIPCHECK(hipDeviceSynchronize()); + + // Check result + HipTest::checkArray(hData,hOutputData,width,height,depth); + hipFreeArray(arr); + hipFreeArray(arr1); + free(hData); + free(hOutputData); +} + +int main(int argc, char **argv) +{ + for(int i=1;i<25;i++) + { + runTest(i,i,i, hipChannelFormatKindFloat); + runTest(i+1,i,i, hipChannelFormatKindSigned); + runTest(i,i+1,i, hipChannelFormatKindSigned); + } + passed(); +} diff --git a/hipamd/tests/src/runtimeApi/module/hipModuleLoadDataMultThreaded.cpp b/hipamd/tests/src/runtimeApi/module/hipModuleLoadDataMultThreaded.cpp new file mode 100644 index 0000000000..e73bbedba5 --- /dev/null +++ b/hipamd/tests/src/runtimeApi/module/hipModuleLoadDataMultThreaded.cpp @@ -0,0 +1,152 @@ +/* +Copyright (c) 2015-Present 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. +*/ + +/* HIT_START + * BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS -std=c++11 + * TEST: %t + * HIT_END + */ + +#include "hip/hip_runtime.h" +#include "hip/hip_runtime_api.h" +#include +#include +#include +#include +#include + +#include "test_common.h" + +#define LEN 64 +#define SIZE LEN << 2 +#define THREADS 2 +#define MAX_THREADS 16 + +#define FILENAME "vcpy_kernel.code" +#define kernel_name "hello_world" + +std::vector load_file() +{ + std::ifstream file(FILENAME, std::ios::binary | std::ios::ate); + std::streamsize fsize = file.tellg(); + file.seekg(0, std::ios::beg); + + std::vector buffer(fsize); + if (!file.read(buffer.data(), fsize)) { + failed("could not open code object '%s'\n", FILENAME); + } + return buffer; +} + +void run(const std::vector& buffer) { + hipDevice_t device; + HIPCHECK(hipDeviceGet(&device, 0)); + hipCtx_t context; + HIPCHECK(hipCtxCreate(&context, 0, device)); + + hipModule_t Module; + hipFunction_t Function; + HIPCHECK(hipModuleLoadData(&Module, &buffer[0])); + HIPCHECK(hipModuleGetFunction(&Function, Module, kernel_name)); + + float *A, *B, *Ad, *Bd; + A = new float[LEN]; + B = new float[LEN]; + + for (uint32_t i = 0; i < LEN; i++) { + A[i] = i * 1.0f; + B[i] = 0.0f; + } + + HIPCHECK(hipMalloc((void**)&Ad, SIZE)); + HIPCHECK(hipMalloc((void**)&Bd, SIZE)); + + HIPCHECK(hipMemcpy(Ad, A, SIZE, hipMemcpyHostToDevice)); + HIPCHECK(hipMemcpy(Bd, B, SIZE, hipMemcpyHostToDevice)); + + hipStream_t stream; + HIPCHECK(hipStreamCreate(&stream)); + + struct { + void* _Ad; + void* _Bd; + } args; + args._Ad = (void*) Ad; + args._Bd = (void*) Bd; + size_t size = sizeof(args); + + void* config[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, &args, HIP_LAUNCH_PARAM_BUFFER_SIZE, &size, + HIP_LAUNCH_PARAM_END}; + HIPCHECK(hipModuleLaunchKernel(Function, 1, 1, 1, LEN, 1, 1, 0, stream, NULL, (void**)&config)); + + HIPCHECK(hipStreamDestroy(stream)); + + HIPCHECK(hipModuleUnload(Module)); + + HIPCHECK(hipMemcpy(B, Bd, SIZE, hipMemcpyDeviceToHost)); + + for (uint32_t i = 0; i < LEN; i++) { + assert(A[i] == B[i]); + } + + hipFree(Ad); + hipFree(Bd); + delete A; + delete B; + hipCtxDestroy(context); + +} + +struct joinable_thread : std::thread +{ + template + joinable_thread(Xs&&... xs) : std::thread(std::forward(xs)...) // NOLINT + { + } + + joinable_thread& operator=(joinable_thread&& other) = default; + joinable_thread(joinable_thread&& other) = default; + + ~joinable_thread() + { + if(this->joinable()) + this->join(); + } +}; + +void run_multi_threads(uint32_t n, const std::vector& buffer) { + + std::vector threads; + + for (uint32_t i = 0; i < n; i++) { + threads.emplace_back(std::thread{[&, buffer] { + run(buffer); + }}); + } + +} + +int main() { + + HIPCHECK(hipInit(0)); + auto buffer = load_file(); + run_multi_threads(min(THREADS * std::thread::hardware_concurrency(), MAX_THREADS), buffer); + + passed(); +} diff --git a/hipamd/tests/src/test_common.cpp b/hipamd/tests/src/test_common.cpp index e7a2622662..1c0dcc8c34 100644 --- a/hipamd/tests/src/test_common.cpp +++ b/hipamd/tests/src/test_common.cpp @@ -37,10 +37,12 @@ int p_tests = -1; /*which tests to run. Interpretation is left to each test. de const char* HIP_VISIBLE_DEVICES_STR = "HIP_VISIBLE_DEVICES="; const char* CUDA_VISIBLE_DEVICES_STR = "CUDA_VISIBLE_DEVICES="; const char* PATH_SEPERATOR_STR = "\\"; +const char* NULL_DEVICE = "NUL:"; #else const char* HIP_VISIBLE_DEVICES_STR = "HIP_VISIBLE_DEVICES"; const char* CUDA_VISIBLE_DEVICES_STR = "CUDA_VISIBLE_DEVICES"; const char* PATH_SEPERATOR_STR = "/"; +const char* NULL_DEVICE = "/dev/null"; #endif namespace HipTest { diff --git a/hipamd/tests/src/test_common.h b/hipamd/tests/src/test_common.h index 426ea846b1..7d8c39e74c 100644 --- a/hipamd/tests/src/test_common.h +++ b/hipamd/tests/src/test_common.h @@ -105,6 +105,10 @@ THE SOFTWARE. #define pclose(x) _pclose(x) #define setenv(x,y,z) _putenv_s(x,y) #define unsetenv _putenv +#define fileno(x) _fileno(x) +#define dup(x) _dup(x) +#define dup2(x,y) _dup2(x,y) +#define close(x) _close(x) #else #define aligned_free(x) free(x) #endif @@ -124,6 +128,7 @@ extern int p_tests; extern const char* HIP_VISIBLE_DEVICES_STR; extern const char* CUDA_VISIBLE_DEVICES_STR; extern const char* PATH_SEPERATOR_STR; +extern const char* NULL_DEVICE; // ********************* CPP section ********************* #ifdef __cplusplus diff --git a/hipamd/tests/src/texture/simpleTexture2DLayered.cpp b/hipamd/tests/src/texture/simpleTexture2DLayered.cpp index e72ea3483a..f4d3aac1e5 100644 --- a/hipamd/tests/src/texture/simpleTexture2DLayered.cpp +++ b/hipamd/tests/src/texture/simpleTexture2DLayered.cpp @@ -21,7 +21,7 @@ THE SOFTWARE. */ /* HIT_START - * BUILD: %t %s ../test_common.cpp + * BUILD: %t %s ../test_common.cpp EXCLUDE_HIP_PLATFORM vdi * TEST: %t * HIT_END */ @@ -66,7 +66,7 @@ void runTest(int width,int height,int num_layers,texturechannelDesc; + hipChannelFormatDesc channelDesc = hipCreateChannelDesc(sizeof(T)*8, 0, 0, 0, hipChannelFormatKindSigned); hipArray *arr; - HIPCHECK(hipMalloc3DArray(&arr, &channelDesc, make_hipExtent(width, height, depth), hipArrayDefault)); + HIPCHECK(hipMalloc3DArray(&arr, &channelDesc, make_hipExtent(width, height, depth), hipArrayCubemap)); hipMemcpy3DParms myparms = {0}; myparms.srcPos = make_hipPos(0,0,0); myparms.dstPos = make_hipPos(0,0,0); @@ -100,7 +100,6 @@ void runTest(int width,int height,int depth,textureaddressMode[0] = hipAddressModeWrap; tex->addressMode[1] = hipAddressModeWrap; - tex->addressMode[2] = hipAddressModeWrap; tex->filterMode = hipFilterModePoint; tex->normalized = false;