diff --git a/projects/clr/hipamd/CMakeLists.txt b/projects/clr/hipamd/CMakeLists.txt index d99907d695..71901da4a4 100644 --- a/projects/clr/hipamd/CMakeLists.txt +++ b/projects/clr/hipamd/CMakeLists.txt @@ -98,6 +98,15 @@ if(NOT DEFINED BUILD_CLANG_HIPIFY) endif() endif() +# Check if we need to enable ATP marker +if(NOT DEFINED COMPILE_HIP_ATP_MARKER) + if(NOT DEFINED ENV{COMPILE_HIP_ATP_MARKER}) + set(COMPILE_HIP_ATP_MARKER 0) + else() + set(COMPILE_HIP_ATP_MARKER $ENV{COMPILE_HIP_ATP_MARKER}) + endif() +endif() + ############################# # Build steps ############################# @@ -109,13 +118,18 @@ endif() # Build hip_hcc if platform is hcc if(HIP_PLATFORM STREQUAL "hcc") include_directories(${PROJECT_SOURCE_DIR}/include) + if(COMPILE_HIP_ATP_MARKER) + include_directories(/opt/rocm/profiler/CXLActivityLogger/include) + set(CMAKE_CXX_FLAGS " -DCOMPILE_HIP_ATP_MARKER=1 ${CMAKE_CXX_FLAGS}") + set(CMAKE_C_FLAGS " -DCOMPILE_HIP_ATP_MARKER=1 ${CMAKE_C_FLAGS}") + endif() set(CMAKE_CXX_COMPILER "${HCC_HOME}/bin/hcc") set(CMAKE_C_COMPILER "${HCC_HOME}/bin/hcc") # Set HIP_HCC so we know this is HIP compile, some files are shared with HCC (staging_buffer). - set(CMAKE_CXX_FLAGS " -hc -I${HCC_HOME}/include -I${HSA_PATH}/include -I/opt/rocm/libhsakmt/include/libhsakmt -stdlib=libc++ -DHIP_HCC") - set(CMAKE_C_FLAGS " -hc -I${HCC_HOME}/include -I${HSA_PATH}/include -I/opt/rocm/libhsakmt/include/libhsakmt -stdlib=libc++ -DHIP_HCC") + set(CMAKE_CXX_FLAGS " -hc -I${HCC_HOME}/include -I${HSA_PATH}/include -I/opt/rocm/libhsakmt/include/libhsakmt -stdlib=libc++ -DHIP_HCC ${CMAKE_CXX_FLAGS}") + set(CMAKE_C_FLAGS " -hc -I${HCC_HOME}/include -I${HSA_PATH}/include -I/opt/rocm/libhsakmt/include/libhsakmt -stdlib=libc++ -DHIP_HCC ${CMAKE_C_FLAGS}") set(SOURCE_FILES src/device_util.cpp src/hip_hcc.cpp @@ -227,5 +241,12 @@ add_custom_target(pkg_hip_samples COMMAND ${CMAKE_COMMAND} . WORKING_DIRECTORY ${BUILD_DIR}) # Package: all +if(POLICY CMP0037) + cmake_policy(PUSH) + cmake_policy(SET CMP0037 OLD) +endif() add_custom_target(package DEPENDS pkg_hip_base pkg_hip_hcc pkg_hip_nvcc pkg_hip_doc pkg_hip_samples) +if(POLICY CMP0037) + cmake_policy(POP) +endif() diff --git a/projects/clr/hipamd/README.md b/projects/clr/hipamd/README.md index 5a2146e74d..4b69485ba3 100644 --- a/projects/clr/hipamd/README.md +++ b/projects/clr/hipamd/README.md @@ -12,16 +12,15 @@ Key features include: New projects can be developed directly in the portable HIP C++ language and can run on either NVIDIA or AMD platforms. Additionally, HIP provides porting tools which make it easy to port existing CUDA codes to the HIP layer, with no loss of performance as compared to the original CUDA application. HIP is not intended to be a drop-in replacement for CUDA, and developers should expect to do some manual coding and performance tuning work to complete the port. ## More Info: -- [Installation](INSTALL.md) and [clang-hipify](INSTALL.md#use-clang-hipify.md) +- [Installation](INSTALL.md) - [HIP FAQ](docs/markdown/hip_faq.md) - [HIP Kernel Language](docs/markdown/hip_kernel_language.md) - [HIP Runtime API (Doxygen)](http://gpuopen-professionalcompute-tools.github.io/HIP) - [HIP Porting Guide](docs/markdown/hip_porting_guide.md) - [HIP Terminology](docs/markdown/hip_terms.md) (including Rosetta Stone of GPU computing terms across CUDA/HIP/HC/AMP/OpenL) +- [clang-hipify](clang-hipify/README.md) - [Developer/CONTRIBUTING Info](CONTRIBUTING.md) - [Release Notes](RELEASE.md) - - ## How do I get set up? See the [Installation](INSTALL.md) notes. diff --git a/projects/clr/hipamd/RELEASE.md b/projects/clr/hipamd/RELEASE.md index dfec3285ae..f60dafd4b5 100644 --- a/projects/clr/hipamd/RELEASE.md +++ b/projects/clr/hipamd/RELEASE.md @@ -1,11 +1,10 @@ # Release notes -We have attempted to document known bugs and limitations - in particular the [HIP Kernel Language](docs/markdown/hip_kernel_language.md) document uses the phrase "Under Development", and the [HIP Runtime API bug list](http://gpuopen-professionalcompute-tools.github.io/HIP/bug.html) lists known bugs. Some of the key items we are working on: -- Tuning built-in functions, including shfl. -- Performance optimization. +We have attempted to document known bugs and limitations - in particular the [HIP Kernel Language](docs/markdown/hip_kernel_language.md) document uses the phrase "Under Development", and the [HIP Runtime API bug list](http://gpuopen-professionalcompute-tools.github.io/HIP/bug.html) lists known bugs. - -Stay tuned - the work for many of these features is already in-flight. +=================================================================================================== +Upcoming: +- hipLaunchKernel supports one-dimensional grid and/or block dims, without explicit cast to dim3 type. =================================================================================================== Release:0.90.00 diff --git a/projects/clr/hipamd/bin/hipcc b/projects/clr/hipamd/bin/hipcc index 7c894af53e..ce888fea8f 100755 --- a/projects/clr/hipamd/bin/hipcc +++ b/projects/clr/hipamd/bin/hipcc @@ -31,10 +31,6 @@ $verbose = 0 unless defined $verbose; $HIP_PATH=$ENV{'HIP_PATH'}; $HIP_PATH=dirname (dirname $0) unless defined $HIP_PATH; # use parent directory of hipcc -$CODEXL_PATH = $ENV{'CODEXL_PATH'}; -$CODEXL_PATH = "/opt/AMD/CodeXL" unless defined $CODEXL_PATH; -$marker_path = "$CODEXL_PATH/SDK/AMDTActivityLogger"; - #--- #HIP_PLATFORM controls whether to use NVCC or HCC for compilation: $HIP_PLATFORM= `$HIP_PATH/bin/hipconfig --platform`; @@ -60,6 +56,11 @@ if ($HIP_PLATFORM eq "hcc") { $ROCM_PATH=$ENV{'ROCM_PATH'}; $ROCM_PATH="/opt/rocm" unless defined $ROCM_PATH; + $HIP_ATP_MARKER=$ENV{'HIP_ATP_MARKER'}; + if ($HIP_ATP_MARKER) { + $marker_path = "$ROCM_PATH/profiler/CXLActivityLogger"; + } + $ROCM_TARGET=$ENV{'ROCM_TARGET'}; $ROCM_TARGET="fiji" unless defined $ROCM_TARGET; @@ -97,14 +98,16 @@ if ($HIP_PLATFORM eq "hcc") { # Add trace marker library: # TODO - once we cleanly separate the HIP API headers from HIP library headers this logic should move to CMakebuild option - apps do not need to see the marker library. - $marker_inc_path = "$marker_path/include"; - if (-e $marker_inc_path) { - $HIPCXXFLAGS .= " -I$marker_inc_path"; - } + if ($HIP_ATP_MARKER) { + $marker_inc_path = "$marker_path/include"; + if (-e $marker_inc_path) { + $HIPCXXFLAGS .= " -I$marker_inc_path"; + } - $marker_lib_path = "$marker_path/bin/x86_64"; - if (-e $marker_lib_path) { - $HIPLDFLAGS .= " -L$marker_lib_path -lAMDTActivityLogger"; + $marker_lib_path = "$marker_path/bin/x86_64"; + if (-e $marker_lib_path) { + $HIPLDFLAGS .= " -L$marker_lib_path -lCXLActivityLogger -Wl,--rpath=$marker_lib_path"; + } } # Add C++ libs for GCC. diff --git a/projects/clr/hipamd/bin/hipconfig b/projects/clr/hipamd/bin/hipconfig index 4f4559cb8e..4f65d3230c 100755 --- a/projects/clr/hipamd/bin/hipconfig +++ b/projects/clr/hipamd/bin/hipconfig @@ -1,7 +1,7 @@ #!/usr/bin/perl -w $HIP_VERSION_MAJOR = "0"; -$HIP_VERSION_MINOR = "90"; +$HIP_VERSION_MINOR = "92"; $HIP_VERSION_PATCH = "0"; use Getopt::Long; diff --git a/projects/clr/hipamd/bin/hipify b/projects/clr/hipamd/bin/hipify index ce934dff15..1d1d071a78 100755 --- a/projects/clr/hipamd/bin/hipify +++ b/projects/clr/hipamd/bin/hipify @@ -279,6 +279,8 @@ while (@ARGV) { $ft{'mem'} += s/\bcudaPointerAttributes\b/hipPointerAttribute_t/g; + $ft{'mem'} += s/\bcudaMemcpy2D\b/hipMemcpy2D/g; + $ft{'mem'} += s/\bcudaMemcpy2DToArray\b/hipMemcpy2DToArray/g; #-------- # Memory management: @@ -293,6 +295,9 @@ while (@ARGV) { $ft{'mem'} += s/\bcudaHostAllocMapped\b/hipHostMallocMapped/g; $ft{'mem'} += s/\bcudaHostAllocWriteCombined\b/hipHostMallocWriteCombined/g; + $ft{'mem'} += s/\bcudaMallocArray\b/hipMallocArray/g; + $ft{'mem'} += s/\bcudaMallocPitch\b/hipMallocPitch/g; + #-------- # Coordinate Indexing and Dimensions: @@ -493,8 +498,10 @@ while (@ARGV) { $ft{'tex'} += s/\bcudaFilterModePoint\b/hipFilterModePoint/g; $ft{'tex'} += s/\bcudaReadModeElementType\b/hipReadModeElementType/g; + $ft{'tex'} += s/\bcudaArray\b/hipArrary/g; $ft{'tex'} += s/\bcudaCreateChannelDesc\b/hipCreateChannelDesc/g; $ft{'tex'} += s/\bcudaBindTexture\b/hipBindTexture/g; + $ft{'tex'} += s/\bcudaBindTextureToArray\b/hipBindTextureToArray/g; $ft{'tex'} += s/\bcudaUnbindTexture\b/hipUnbindTexture/g; } diff --git a/projects/clr/hipamd/docs/doxygen-input/mainpage.txt b/projects/clr/hipamd/docs/doxygen-input/mainpage.txt index 5424464f5c..48177198cc 100644 --- a/projects/clr/hipamd/docs/doxygen-input/mainpage.txt +++ b/projects/clr/hipamd/docs/doxygen-input/mainpage.txt @@ -7,11 +7,11 @@ * The HIP interface makes it very easy to port existing CUDA apps to run on AMD GPUs, * or to develop new apps that can run on either CUDA or AMD GPUs from a common source base. * - * - HIP is very thin and has little or no performance impact over coding directly in CUDA NVCC or Kalmar HC mode. + * - HIP is very thin and has little or no performance impact over coding directly in CUDA NVCC or HCC HC mode. * - HIP allows developers to use the "best" development environment and tools on each target platform. * - HIP allows coding in a single-source C++ programming language including features such as templates, C++11 lambdas,and more. - * - "hipify" script automatically converts source from CUDA to HIP. - * - Developers can specialize for CUDA or Kalmar to tune for performance or handle tricky cases with #ifdef. + * - "hipify" tool automatically converts source from CUDA to HIP. + * - Developers can specialize for CUDA or HCC to tune for performance or handle tricky cases with #ifdef. * - See the @ref API. diff --git a/projects/clr/hipamd/docs/markdown/hip_faq.md b/projects/clr/hipamd/docs/markdown/hip_faq.md index 5e722383df..e71dcdfb36 100644 --- a/projects/clr/hipamd/docs/markdown/hip_faq.md +++ b/projects/clr/hipamd/docs/markdown/hip_faq.md @@ -1,13 +1,12 @@ # FAQ - - -**Table of Contents** *generated with [DocToc](https://github.com/thlorenz/doctoc)* - -- [What APIs does HIP support ?](#what-apis-does-hip-support-) -- [What is not supported?](#what-is-not-supported) +- [What APIs and features does HIP support ?](#Q1) +- [What is not supported?](#Q2) - [Run-time features](#run-time-features) - [Kernel language features](#kernel-language-features) +- [Is HIP a drop-in replacement for CUDA?](#Q3) +- [What version of CUDA is supported?](#Q4) +- [What libraries does HIP support?](#Q5) - [How does HIP compare with OpenCL?](#how-does-hip-compare-with-opencl) - [What hardware does HIP support?](#what-hardware-does-hip-support) - [Does Hipify automatically convert all source code?](#does-hipify-automatically-convert-all-source-code) @@ -20,38 +19,79 @@ - [What's the difference between HIP and hc?](#whats-the-difference-between-hip-and-hc) - [HIP detected my platform (hcc vs nvcc) incorrectly - what should I do?](#hip-detected-my-platform-hcc-vs-nvcc-incorrectly---what-should-i-do) - -================= -### What APIs does HIP support ? +### What APIs and features does HIP support? HIP provides the following: -- Devices (hipSetDevice(), hipGetDeviceProperties(), etc) -- Memory management (hipMalloc(), hipMemcpy(), hipFree()) -- Streams (hipStreamCreate(), etc.)---under development +- Devices (hipSetDevice(), hipGetDeviceProperties(), etc.) +- Memory management (hipMalloc(), hipMemcpy(), hipFree(), etc.) +- Streams (hipStreamCreate(), etc.) - Events (hipEventRecord(), hipEventElapsedTime(), etc.) - Kernel launching (hipLaunchKernel is a standard C/C++ function that replaces <<< >>>) -- CUDA-style kernel indexing -- Device-side math built-ins +- CUDA-style kernel coordinate functions (threadIdx, blockIdx, blockDim, gridDim) +- Most device-side math built-ins - Error reporting (hipGetLastError(), hipGetErrorString()) -The HIP documentation describes each API and its limitations, if any, compared with the equivalent CUDA API. +The HIP API documentation describes each API and its limitations, if any, compared with the equivalent CUDA API. -### What is not supported? +### What is not supported? #### Run-time features - Textures -- Dynamic parallelism -- Managed memory +- MemcpyToSymbol functions +- Dynamic parallelism (CUDA 5.0) +- Managed memory (CUDA 6.5) - Graphics interoperation with OpenGL or Direct3D - CUDA array, mipmappedArray and pitched memory - CUDA Driver API - + #### Kernel language features -- Device-side dynamic memory allocations (malloc, free, new, delete) -- Virtual functions, indirect functions and try/catch +- Device-side dynamic memory allocations (malloc, free, new, delete) (CUDA 4.0) +- Virtual functions, indirect functions and try/catch (CUDA 4.0) - `__prof_trigger` -- PTX assembly -- See the [HIP Kernel Language](hip_kernel_language.md) for more information. +- PTX assembly (CUDA 4.0) +- Several kernel features are under development. See the [HIP Kernel Language](hip_kernel_language.md) for more information. + +### Is HIP a drop-in replacement for CUDA? +No. HIP provides porting tools which do most of the work do convert CUDA code into portable C++ code that uses the HIP APIs. +Most developers will port their code from CUDA to HIP and then maintain the HIP version. +HIP code provides the same performance as coding in native CUDA, plus the benefit that the code can also run on AMD platforms. + +### What version of CUDA is supported? +HIP APIs and features do not map to a specific CUDA version. HIP provides a strong subset of functionality provided in CUDA, and the hipify tools can +scan code to identify any unsupported CUDA functions - this is very useful for identifying the specific features required by a given application. + +However, we can provide a rough summary of the features included in each CUDA SDK and the support level in HIP: + +- CUDA 4.0 and earlier : + - HIP supports CUDA 4.0 except for the limitations described [above](#Q2). +- CUDA 5.0 : + - Dynamic Parallelism (not supported) + - cuIpc functions (under development). +- CUDA 5.5 : + - CUPTI (not directly supported, [AMD GPUPerfAPI](http://developer.amd.com/tools-and-sdks/graphics-development/gpuperfapi/) can be used as an alternative in some cases) +- CUDA 6.0 + - Managed memory (under development) +- CUDA 6.5 + - __shfl instriniscs (supported) +- CUDA 7.0 + - Per-thread-streams (under development) + - C++11 (HCC supports all of C++11, all of C++14 and some C++17 features) +- CUDA 7.5 + - float16 (under development) +- CUDA 8.0 + - No new language features. + +### What libraries does HIP support? +HIP includes growing support for the 4 key math libraries using hcBlas, hcFft, hcrng, and hcsparse). +These offer pointer-based memory interfaces (as opposed to opaque buffers) and can be easily interfaces with other HCC code. Developers should use conditional compliation if portability to nvcc systems is desired - using calls to cu* routines on one path and hc* routines on the other. + +- [hcblas](https://bitbucket.org/multicoreware/hcblas) +- [hcfft](https://bitbucket.org/multicoreware/hcfft) +- [hcsparse](https://bitbucket.org/multicoreware/hcsparse) +- [hcrng](https://bitbucket.org/multicoreware/hcrng) + +Additionally, some of the cublas routines are automatically converted to hipblas equivalents by the clang-hipify tool. These APIs use cublas or hcblas depending on the platform, and replace the need +to use conditional compilation. ### How does HIP compare with OpenCL? Both AMD and Nvidia support OpenCL 1.2 on their devices, so developers can write portable code. @@ -65,7 +105,7 @@ HIP offers several benefits over OpenCL: - HIP offers an offline compilation model. ### What hardware does HIP support? -- For AMD platforms, HIP runs on the same hardware that the HCC "hc" mode supports - specifically AMD Kaveri, Carrizo and Fiji. +- For AMD platforms, HIP runs on the same hardware that the HCC "hc" mode supports. See the ROCM documentation for the list of supported platforms. - For Nvidia platforms, HIP requires Unified Memory and should run on a device which runs the CUDA SDK 6.0 or newer. We have tested the Nvidia Titan and K40. ### Does Hipify automatically convert all source code? @@ -79,7 +119,7 @@ Developers should always expect to perform some platform-specific tuning and opt NVCC is Nvidia's compiler driver for compiling "CUDA C++" code into PTX or device code for Nvidia GPUs. It's a closed-source binary product that comes with CUDA SDKs. ### What is HCC? -HCC is AMD's compiler driver which compiles "heterogenous C++" code into HSAIL or GCN device code for AMD GPUs. HCC is an open-source compiler based on recent versions of CLANG/LLVM. +HCC is AMD's compiler driver which compiles "heterogenous C++" code into HSAIL or GCN device code for AMD GPUs. It's an open-source compiler based on recent versions of CLANG/LLVM. ### Why use HIP rather than supporting CUDA directly? While HIP is a strong subset of the CUDA, it is a subset. The HIP layer allows that subset to be clearly defined and documented. diff --git a/projects/clr/hipamd/docs/markdown/hip_kernel_language.md b/projects/clr/hipamd/docs/markdown/hip_kernel_language.md index 9ac1a48a25..ec2530e46e 100644 --- a/projects/clr/hipamd/docs/markdown/hip_kernel_language.md +++ b/projects/clr/hipamd/docs/markdown/hip_kernel_language.md @@ -1,6 +1,4 @@ - - -**Table of Contents** *generated with [DocToc](https://github.com/thlorenz/doctoc)* +**Table of Contents** - [Introduction](#introduction) - [Function-Type Qualifiers](#function-type-qualifiers) @@ -45,8 +43,6 @@ - [In-Line Assembly](#in-line-assembly) - [C++ Support](#c-support) - - ## Introduction @@ -98,9 +94,9 @@ HIP parses the `__noinline__` and `__forceinline__` keywords and converts them t `__global__` functions are often referred to as *kernels,* and calling one is termed *launching the kernel.* These functions require the caller to specify an "execution configuration" that includes the grid and block dimensions. The execution configuration can also include other information for the launch, such as the amount of additional shared memory to allocate and the stream where the kernel should execute. HIP introduces a standard C++ calling convention to pass the execution configuration to the kernel (this convention replaces the Cuda <<< >>> syntax). In HIP, - Kernels launch with the "hipLaunchKernel" function - The first five parameters to hipLaunchKernel are the following: - - **symbol kernelName**: the name of the kernel to launch - - **dim3 gridDim**: 3D-grid dimensions - - **dim3 blockDim**: 3D-block dimensions + - **symbol kernelName**: the name of the kernel to launch. To support template kernels which contains "," use the HIP_KERNEL_NAME macro. The hipify tools insert this automatically. + - **dim3 gridDim**: 3D-grid dimensions specifying the number of blocks to launch. + - **dim3 blockDim**: 3D-block dimensions specifying the number of threads in each block. - **size_t dynamicShared**: amount of additional shared memory to allocate when launching the kernel (see [__shared__](#__shared__)) - **hipStream_t**: stream where the kernel should execute. A value of 0 corresponds to the NULL stream (see [Synchronization Functions](#synchronization-functions)). - Kernel arguments follow these first five parameters @@ -234,6 +230,9 @@ typedef struct dim3 { ## Memory-Fence Instructions HIP support for __threadfence(), __threadfence_block() and __threadfence_system() is under development. +The stubs for the threadfence routines are defined in hcc_details/hip_runtime.h. +Applications that use these threadfence features should disable both of the L1 and L2 caches by: +"export HSA_DISABLE_CACHE=1" ## Synchronization Functions The __syncthreads() built-in function is supported in HIP. The __syncthreads_count(int), __syncthreads_and(int) and __syncthreads_or(int) functions are under development. @@ -629,6 +628,9 @@ The compiler ensures that the kernel uses fewer registers than both allowed maxi HIP/hcc will parse the `launch_bounds` attribute but silently ignores the performance hint. Full support is under development. +The hcc compiler does not support the "--maxregcount" option like nvcc. Instead, users are encouraged to use the hip_launch_bounds directive since the parameters are more intuitive and portable than +micro-architecture details like registers, and also the directive allows per-kernel control rather than an entire file. hip_launch_bounds works on both hcc and nvcc targets. + ## Register Keyword The register keyword affects code generation in neither nvcc nor hcc. It’s deprecated in standard C++, so hcc will generate a warning. (nvcc silently ignores use of this keyword.) To disable the warning, you can pass the option `-Wno-deprecated-register` to hcc. @@ -636,23 +638,25 @@ The register keyword affects code generation in neither nvcc nor hcc. It’s de ## Pragma Unroll -hcc support for the unroll pragma is under development and is slated to arrive with the Lightning Compiler. - +Unroll with a bounds that is known at compile-time is supported. For example: ``` #pragma unroll 16 /* hint to compiler to unroll next loop by 16 */ for (int i=0; i<16; i++) ... ``` +``` +#pragma unroll 1 /* tell compiler to never unroll the loop */ +for (int i=0; i<16; i++) ... +``` + + +Unbounded loop unroll is under development on HCC compiler. ``` #pragma unroll /* hint to compiler to completely unroll next loop. */ for (int i=0; i<16; i++) ... ``` -``` -#pragma unroll 1 /* tell compiler to never unroll the loop */ -for (int i=0; i<16; i++) ... -``` ## In-Line Assembly diff --git a/projects/clr/hipamd/docs/markdown/hip_porting_guide.md b/projects/clr/hipamd/docs/markdown/hip_porting_guide.md index 50282c3eff..76cf86aa49 100644 --- a/projects/clr/hipamd/docs/markdown/hip_porting_guide.md +++ b/projects/clr/hipamd/docs/markdown/hip_porting_guide.md @@ -1,95 +1,50 @@ - - -**Table of Contents** *generated with [DocToc](https://github.com/thlorenz/doctoc)* - -- [HIP Porting Guide](#hip-porting-guide) -- [###Table of Contents](#table-of-contents) - - [Porting a New Cuda Project](#porting-a-new-cuda-project) - - [General Tips](#general-tips) - - [Scanning existing CUDA code to scope the porting effort](#scanning-existing-cuda-code-to-scope-the-porting-effort) - - [Converting a project "in-place"](#converting-a-project-in-place) - - [Distinguishing Compiler Modes](#distinguishing-compiler-modes) - - [Identifying HIP Target Platform](#identifying-hip-target-platform) - - [Identifying the Compiler: hcc or nvcc](#identifying-the-compiler-hcc-or-nvcc) - - [Identifying Current Compilation Pass: Host or Device](#identifying-current-compilation-pass-host-or-device) - - [Compiler Defines: Summary](#compiler-defines-summary) - - [Identifying Architecture Features](#identifying-architecture-features) - - [HIP_ARCH Defines](#hip_arch-defines) - - [Device-Architecture Properties](#device-architecture-properties) - - [Table of Architecture Properties](#table-of-architecture-properties) - - [Finding HIP](#finding-hip) - - [Compiler Options](#compiler-options) - - [Linking Issues](#linking-issues) - - [Linking With hipcc](#linking-with-hipcc) - - [-lm Option](#-lm-option) - - [Linking Code With Other Compilers](#linking-code-with-other-compilers) - - [libc++ and libstdc++](#libc-and-libstdc) - - [HIP Headers (hip_runtime.h, hip_runtime_api.h)](#hip-headers-hip_runtimeh-hip_runtime_apih) - - [Using a Standard C++ Compiler](#using-a-standard-c-compiler) - - [cuda.h](#cudah) - - [Choosing HIP File Extensions](#choosing-hip-file-extensions) - - [Workarounds](#workarounds) - - [warpSize](#warpsize) - - [Textures and Cache Control](#textures-and-cache-control) - - [More Tips](#more-tips) - - [hcc CPU Mode](#hcc-cpu-mode) - - [HIPTRACE Mode](#hiptrace-mode) - - [Environment Variables](#environment-variables) - - [Debugging hipcc](#debugging-hipcc) - - [What Does This Error Mean?](#what-does-this-error-mean) - - [/usr/include/c++/v1/memory:5172:15: error: call to implicitly deleted default constructor of 'std::__1::bad_weak_ptr' throw bad_weak_ptr();](#usrincludecv1memory517215-error-call-to-implicitly-deleted-default-constructor-of-std__1bad_weak_ptr-throw-bad_weak_ptr) - - [grid_launch kernel dispatch - fallback](#grid_launch-kernel-dispatch---fallback) - - [HIP Environment Variables](#hip-environment-variables) - - [Editor Highlighting](#editor-highlighting) - - - # HIP Porting Guide In addition to providing a portable C++ programmming environement for GPUs, HIP is designed to ease the porting of existing CUDA code into the HIP environment. This section describes the available tools and provides practical suggestions on how to port CUDA code and work through common issues. -###Table of Contents +## Table of Contents ================= - * [HIP Porting Guide](#hip-porting-guide) - * [Table of Contents](#table-of-contents) - * [Porting a New Cuda Project To HIP](#porting-a-new-cuda-project) - * [General Tips](#general-tips" aria-hidden="true"> ## Porting a New Cuda Project @@ -463,17 +418,13 @@ void myFunc () ``` +Additionally, many of the Rodinia benchmarks demonstrate how to modify hipified programs so that textures are not required - search for USE_TEXTURES define in the rodinia source directory. +For example, [here + Cuda programs that employ sampler hardware must either wait for hcc texture support or use more-sophisticated workarounds. ## More Tips -### hcc CPU Mode -Recent hcc versions support CPU accelerator targets. This feature enables some interesting possibilities for HIP porting: - -- hcc can run on any machine, including perhaps a cross-compiling environment on a machine also running nvcc -- Standard CPU debuggers can debug CPU code -- A single code path can run on an AMD or Nvidia GPU or CPU, but the CPU accelerator is a low-performance target---it’s just a single core and lacks SIMD acceleration - ### HIPTRACE Mode On an hcc/AMD platform, set the HIP_TRACE_API environment variable to see a textural API trace. Use the following bit mask: @@ -490,7 +441,7 @@ On hcc/AMD platforms, set the HIP_PRINT_ENV environment variable to 1 and run an - HIP_TRACE_API = 1: trace each HIP API call. Print the function name and return code to stderr as the program executes. - HIP_LAUNCH_BLOCKING = 0: make HIP APIs “host-synchronous” so they are blocked until any kernel launches or data-copy commands are complete (an alias is CUDA_LAUNCH_BLOCKING) -- KMDUMPISA = 1 : Will dump the GCN ISA for all kernels into the local directory. +- KMDUMPISA = 1 : Will dump the GCN ISA for all kernels into the local directory. (This flag is provided by HCC). ### Debugging hipcc diff --git a/projects/clr/hipamd/include/hcc_detail/hip_hcc.h b/projects/clr/hipamd/include/hcc_detail/hip_hcc.h index 7c47d1b2a9..635b65f384 100644 --- a/projects/clr/hipamd/include/hcc_detail/hip_hcc.h +++ b/projects/clr/hipamd/include/hcc_detail/hip_hcc.h @@ -130,7 +130,7 @@ class ihipDevice_t; // TODO - currently we print the trace message at the beginning. if we waited, we could also include return codes, and any values returned // through ptr-to-args (ie the pointers allocated by hipMalloc). #if COMPILE_HIP_ATP_MARKER -#include "AMDTActivityLogger.h" +#include "CXLActivityLogger.h" #define SCOPED_MARKER(markerName,group,userString) amdtScopedMarker(markerName, group, userString) #else // Swallow scoped markers: @@ -445,6 +445,7 @@ private: private: void enqueueBarrier(hsa_queue_t* queue, ihipSignal_t *depSignal); + void enqueueBarrier(hsa_queue_t* queue, hsa_signal_t *depSignal); void waitCopy(LockedAccessor_StreamCrit_t &crit, ihipSignal_t *signal); // The unsigned return is hipMemcpyKind @@ -466,6 +467,17 @@ inline std::ostream& operator<<(std::ostream& os, const ihipStream_t& s) return os; } +inline std::ostream & operator<<(std::ostream& os, const dim3& s) +{ + os << '{'; + os << s.x; + os << ','; + os << s.y; + os << ','; + os << s.z; + os << '}'; + return os; +} //---- // Internal event structure: @@ -594,7 +606,7 @@ public: // Data, set at initialization: unsigned _compute_units; StagingBuffer *_staging_buffer[2]; // one buffer for each direction. - + int isLargeBar; unsigned _device_flags; diff --git a/projects/clr/hipamd/include/hcc_detail/hip_runtime.h b/projects/clr/hipamd/include/hcc_detail/hip_runtime.h index 22095b342d..c29995ba2a 100644 --- a/projects/clr/hipamd/include/hcc_detail/hip_runtime.h +++ b/projects/clr/hipamd/include/hcc_detail/hip_runtime.h @@ -48,7 +48,7 @@ THE SOFTWARE. #ifdef __HCC__ #include -#if defined (GRID_LAUNCH_VERSION) and (GRID_LAUNCH_VERSION >= 20) +#if defined (GRID_LAUNCH_VERSION) and (GRID_LAUNCH_VERSION >= 20) // Use field names for grid_launch 2.0 structure, if HCC supports GL 2.0. #define USE_GRID_LAUNCH_20 1 #else @@ -496,6 +496,59 @@ __device__ float __dsqrt_rz(double x); * Kernel launching */ +/** + *------------------------------------------------------------------------------------------------- + *------------------------------------------------------------------------------------------------- + * @defgroup Memory Fence Functions + * @{ + * + * + * @warning The HIP memory fence functions are currently not supported yet. + * If any of those threadfence stubs are reached by the application, you should set "export HSA_DISABLE_CACHE=1" to disable L1 and L2 caches. + * + * + * On AMD platforms, the threadfence* routines are currently empty stubs. + */ + + /** + * @brief threadfence_block makes writes visible to threads running in same block. + * + * @Returns void + * + * @param void + * + * @warning __threadfence_block is a stub and map to no-op. + */ +__device__ void __threadfence_block(void); + + /** + * @brief threadfence makes wirtes visible to other threads running on same GPU. + * + * @Returns void + * + * @param void + * + * @warning __threadfence is a stub and map to no-op, application should set "export HSA_DISABLE_CACHE=1" to disable both L1 and L2 caches. + */ +__device__ void __threadfence(void); + +/** + * @brief threadfence_system makes writes to pinned system memory visible on host CPU. + * + * @Returns void + * + * @param void + * + * @warning __threadfence_system is a stub and map to no-op, application should set "export HSA_DISABLE_CACHE=1" to disable both L1 and L2 caches. + */ +__device__ void __threadfence_system(void); + + +// doxygen end Memory Fence +/** + * @} + */ + #define hipThreadIdx_x (hc_get_workitem_id(0)) #define hipThreadIdx_y (hc_get_workitem_id(1)) diff --git a/projects/clr/hipamd/include/hcc_detail/hip_runtime_api.h b/projects/clr/hipamd/include/hcc_detail/hip_runtime_api.h index 3fd06aae38..1ee7d70ef5 100644 --- a/projects/clr/hipamd/include/hcc_detail/hip_runtime_api.h +++ b/projects/clr/hipamd/include/hcc_detail/hip_runtime_api.h @@ -121,7 +121,6 @@ typedef struct dim3 { } dim3; - /** * Memory copy types * @@ -760,6 +759,20 @@ hipError_t hipHostRegister(void* hostPtr, size_t sizeBytes, unsigned int flags) */ hipError_t hipHostUnregister(void* hostPtr) ; +/** + * Allocates at least width (in bytes) * height bytes of linear memory + * Padding may occur to ensure alighnment requirements are met for the given row + * The change in width size due to padding will be returned in *pitch. + * Currently the alignment is set to 128 bytes + * + * @param[out] ptr Pointer to the allocated device memory + * @param[out] pitch Pitch for allocation (in bytes) + * @param[in] width Requested pitched allocation width (in bytes) + * @param[in] height Requested pitched allocation height + * @return Error code + */ + +hipError_t hipMallocPitch(void** ptr, size_t* pitch, size_t width, size_t height); /** * @brief Free memory allocated by the hcc hip memory allocation API. diff --git a/projects/clr/hipamd/include/hcc_detail/hip_texture.h b/projects/clr/hipamd/include/hcc_detail/hip_texture.h index d4c3403ccf..5712e5c333 100644 --- a/projects/clr/hipamd/include/hcc_detail/hip_texture.h +++ b/projects/clr/hipamd/include/hcc_detail/hip_texture.h @@ -38,9 +38,20 @@ THE SOFTWARE. //Texture - TODO - likely need to move this to a separate file only included with kernel compilation. #define hipTextureType1D 1 +typedef enum { + hipChannelFormatKindSigned = 0, + hipChannelFormatKindUnsigned, + hipChannelFormatKindFloat, + hipChannelFormatKindNone + +} hipChannelFormatKind; + typedef struct hipChannelFormatDesc { - // TODO - this has 4-5 well-defined fields, we could just copy... - int _dummy; + int x; + int y; + int z; + int w; + hipChannelFormatKind f; } hipChannelFormatDesc; typedef enum hipTextureReadMode @@ -67,14 +78,39 @@ struct texture : public textureReference { const T * _dataPtr; // pointer to underlying data. //texture() : filterMode(hipFilterModePoint), normalized(false), _dataPtr(NULL) {}; + unsigned int width; + unsigned int height; + }; #endif +typedef struct hipArray { + unsigned int width; + unsigned int height; + hipChannelFormatKind f; + void* data; //FIXME: generalize this +} hipArray; #define tex1Dfetch(_tex, _addr) (_tex._dataPtr[_addr]) +#define tex2D(_tex, _dx, _dy) \ + _tex._dataPtr[(unsigned int)_dx + (unsigned int)_dy*(_tex.width)] +hipError_t hipMallocArray(hipArray** array, const hipChannelFormatDesc* desc, + size_t width, size_t height = 0, unsigned int flags = 0); + +hipError_t hipFreeArray(hipArray* array); + // +// dpitch, spitch, and width in bytes +hipError_t hipMemcpy2D(void* dst, size_t dpitch, const void* src, size_t spitch, size_t width, size_t height, hipMemcpyKind kind); + +// wOffset, width, and spitch in bytes +hipError_t hipMemcpy2DToArray(hipArray* dst, size_t wOffset, size_t hOffset, const void* src, + size_t spitch, size_t width, size_t height, hipMemcpyKind kind); + +hipError_t hipMemcpyToArray(hipArray* dst, size_t wOffset, size_t hOffset, + const void* src, size_t count, hipMemcpyKind kind); /** @@ -125,11 +161,31 @@ hipChannelFormatDesc hipBindTexture(size_t *offset, struct textureReference *te * * **/ -template -hipChannelFormatDesc hipCreateChannelDesc() -{ - hipChannelFormatDesc desc; - return desc; +hipChannelFormatDesc hipCreateChannelDesc(int x, int y, int z, int w, hipChannelFormatKind f); + +// descriptors +template inline hipChannelFormatDesc hipCreateChannelDesc() { + return hipCreateChannelDesc(0, 0, 0, 0, hipChannelFormatKindNone); +} +template <> inline hipChannelFormatDesc hipCreateChannelDesc() { + int e = (int)sizeof(int) * 8; + return hipCreateChannelDesc(e, 0, 0, 0, hipChannelFormatKindSigned); +} +template <> inline hipChannelFormatDesc hipCreateChannelDesc() { + int e = (int)sizeof(unsigned int) * 8; + return hipCreateChannelDesc(e, 0, 0, 0, hipChannelFormatKindUnsigned); +} +template <> inline hipChannelFormatDesc hipCreateChannelDesc() { + int e = (int)sizeof(long) * 8; + return hipCreateChannelDesc(e, 0, 0, 0, hipChannelFormatKindSigned); +} +template <> inline hipChannelFormatDesc hipCreateChannelDesc() { + int e = (int)sizeof(unsigned long) * 8; + return hipCreateChannelDesc(e, 0, 0, 0, hipChannelFormatKindUnsigned); +} +template <> inline hipChannelFormatDesc hipCreateChannelDesc() { + int e = (int)sizeof(float) * 8; + return hipCreateChannelDesc(e, 0, 0, 0, hipChannelFormatKindFloat); } /* @@ -178,6 +234,13 @@ hipError_t hipBindTexture(size_t *offset, return hipBindTexture(offset, tex, devPtr, &tex.channelDesc, size); } +template +hipError_t hipBindTextureToArray(struct texture &tex, hipArray* array) { + tex.width = array->width; + tex.height = array->height; + tex._dataPtr = static_cast(array->data); + return hipSuccess; +} /* * @brief Unbinds the textuer bound to @p tex @@ -187,9 +250,9 @@ hipError_t hipBindTexture(size_t *offset, * @return #hipSuccess **/ template -hipError_t hipUnbindTexture(struct texture *tex) +hipError_t hipUnbindTexture(struct texture &tex) { - tex->_dataPtr = NULL; + tex._dataPtr = NULL; return hipSuccess; } diff --git a/projects/clr/hipamd/packaging/convert_md_to_html.sh b/projects/clr/hipamd/packaging/convert_md_to_html.sh new file mode 100755 index 0000000000..0f6237deb6 --- /dev/null +++ b/projects/clr/hipamd/packaging/convert_md_to_html.sh @@ -0,0 +1,61 @@ +#!/bin/bash +function die { + echo "${1-Died}." >&2 + exit 1 +} + +function cleanup { + rm -rf "$workdir" +} + +# parse arguments +hip_srcdir=$1 +html_destdir=$2 +[ "$hip_srcdir" != "" ] || [ "$html_destdir" != "" ] || die "Invalid arguments!" + +# create temporary directory for grip settings +workdir=`mktemp -d` +trap cleanup EXIT + +# setup grip +export GRIPURL=$hip_srcdir +export GRIPHOME=$workdir +echo "CACHE_DIRECTORY = '$html_destdir/asset'" > $workdir/settings.py +mkdir -p $html_destdir $html_destdir/clang-hipify $html_destdir/docs/markdown + +# convert all md files to html +pushd $hip_srcdir +for f in *.md clang-hipify/*.md docs/markdown/*.md; do grip --export --no-inline $f $html_destdir/${f%.*}.html; done +popd + +# convert absolute links to relative links +pushd $html_destdir +for f in *.html; do sed -i "s?$GRIPURL/??g" $f; done +for f in clang-hipify/*.html; do sed -i "s?$GRIPURL/?../?g" $f; done +for f in docs/markdown/*.html; do sed -i "s?$GRIPURL/?../../?g" $f; done +popd + +# update document titles +pushd $html_destdir +for f in *.html; do sed -i "s?.md - Grip??g" $f; done +for f in clang-hipify/*.html; do sed -i "s?.md - Grip??g" $f; done +for f in docs/markdown/*.html; do sed -i "s?.md - Grip??g" $f; done +popd + +# replace .md with .html in links +pushd $html_destdir +for f in *.html; do sed -i "s?.md\"?.html\"?g" $f; done +for f in *.html; do sed -i "s?.md#?.html#?g" $f; done +for f in clang-hipify/*.html; do sed -i "s?.md\"?.html\"?g" $f; done +for f in clang-hipify/*.html; do sed -i "s?.md#?.html#?g" $f; done +for f in docs/markdown/*.html; do sed -i "s?.md\"?.html\"?g" $f; done +for f in docs/markdown/*.html; do sed -i "s?.md#?.html#?g" $f; done +popd + +# replace github.io links +pushd $html_destdir +sed -i "s?http://gpuopen-professionalcompute-tools.github.io/HIP?docs/RuntimeAPI/html/index.html?g" README.html +sed -i "s?http://gpuopen-professionalcompute-tools.github.io/HIP?docs/RuntimeAPI/html/?g" RELEASE.html +popd + +exit 0 diff --git a/projects/clr/hipamd/packaging/hip_doc.txt b/projects/clr/hipamd/packaging/hip_doc.txt index 287bb8e1a8..bbcaf54ec8 100644 --- a/projects/clr/hipamd/packaging/hip_doc.txt +++ b/projects/clr/hipamd/packaging/hip_doc.txt @@ -3,7 +3,10 @@ project(hip_doc) add_custom_target(build_doxygen ALL COMMAND HIP_PATH=@hip_SOURCE_DIR@ doxygen @hip_SOURCE_DIR@/docs/doxygen-input/doxy.cfg) -install(DIRECTORY RuntimeAPI/html DESTINATION docs) +add_custom_target(convert_md_to_html ALL + COMMAND @hip_SOURCE_DIR@/packaging/convert_md_to_html.sh @hip_SOURCE_DIR@ ${PROJECT_BINARY_DIR}/md2html) +install(DIRECTORY RuntimeAPI/html DESTINATION docs/docs/RuntimeAPI) +install(DIRECTORY md2html/ DESTINATION docs) ############################# # Packaging steps diff --git a/projects/clr/hipamd/packaging/hip_hcc.txt b/projects/clr/hipamd/packaging/hip_hcc.txt index 556bade072..661d77e530 100644 --- a/projects/clr/hipamd/packaging/hip_hcc.txt +++ b/projects/clr/hipamd/packaging/hip_hcc.txt @@ -20,12 +20,12 @@ 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") -set(CPACK_DEBIAN_PACKAGE_DEPENDS "hip_base (= ${CPACK_PACKAGE_VERSION}), hcc_lc (= @HCC_VERSION@)") +set(CPACK_DEBIAN_PACKAGE_DEPENDS "hip_base (= ${CPACK_PACKAGE_VERSION}), hcc_lc (= @HCC_VERSION@), rocm-profiler") set(CPACK_BINARY_RPM "ON") set(CPACK_RPM_PACKAGE_ARCHITECTURE "x86_64") 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") -set(CPACK_RPM_PACKAGE_REQUIRES "hip_base = ${CPACK_PACKAGE_VERSION}, hcc_lc = @HCC_VERSION@") +set(CPACK_RPM_PACKAGE_REQUIRES "hip_base = ${CPACK_PACKAGE_VERSION}, hcc_lc = @HCC_VERSION@, rocm-profiler") set(CPACK_SOURCE_GENERATOR "TGZ") include(CPack) diff --git a/projects/clr/hipamd/src/hip_device.cpp b/projects/clr/hipamd/src/hip_device.cpp index 86e8563272..f16ae715c1 100644 --- a/projects/clr/hipamd/src/hip_device.cpp +++ b/projects/clr/hipamd/src/hip_device.cpp @@ -61,7 +61,7 @@ hipError_t hipGetDeviceCount(int *count) */ hipError_t hipDeviceSetCacheConfig ( hipFuncCache cacheConfig ) { - std::call_once(hip_initialized, ihipInit); + HIP_INIT_API(cacheConfig); // Nop, AMD does not support variable cache configs. @@ -75,7 +75,7 @@ hipError_t hipDeviceSetCacheConfig ( hipFuncCache cacheConfig ) */ hipError_t hipDeviceGetCacheConfig ( hipFuncCache *cacheConfig ) { - std::call_once(hip_initialized, ihipInit); + HIP_INIT_API(cacheConfig); *cacheConfig = hipFuncCachePreferNone; @@ -89,7 +89,7 @@ hipError_t hipDeviceGetCacheConfig ( hipFuncCache *cacheConfig ) */ hipError_t hipFuncSetCacheConfig ( hipFuncCache cacheConfig ) { - std::call_once(hip_initialized, ihipInit); + HIP_INIT_API(cacheConfig); // Nop, AMD does not support variable cache configs. @@ -104,7 +104,7 @@ hipError_t hipFuncSetCacheConfig ( hipFuncCache cacheConfig ) */ hipError_t hipDeviceSetSharedMemConfig ( hipSharedMemConfig config ) { - std::call_once(hip_initialized, ihipInit); + HIP_INIT_API(config); // Nop, AMD does not support variable shared mem configs. @@ -119,7 +119,7 @@ hipError_t hipDeviceSetSharedMemConfig ( hipSharedMemConfig config ) */ hipError_t hipDeviceGetSharedMemConfig ( hipSharedMemConfig * pConfig ) { - std::call_once(hip_initialized, ihipInit); + HIP_INIT_API(pConfig); *pConfig = hipSharedMemBankSizeFourByte; @@ -184,7 +184,7 @@ hipError_t hipDeviceReset(void) */ hipError_t hipDeviceGetAttribute(int* pi, hipDeviceAttribute_t attr, int device) { - std::call_once(hip_initialized, ihipInit); + HIP_INIT_API(attr, device); hipError_t e = hipSuccess; diff --git a/projects/clr/hipamd/src/hip_event.cpp b/projects/clr/hipamd/src/hip_event.cpp index 983b190644..1514fc5868 100644 --- a/projects/clr/hipamd/src/hip_event.cpp +++ b/projects/clr/hipamd/src/hip_event.cpp @@ -69,7 +69,7 @@ hipError_t hipEventCreate(hipEvent_t* event) //--- hipError_t hipEventRecord(hipEvent_t event, hipStream_t stream) { - std::call_once(hip_initialized, ihipInit); + HIP_INIT_API(event, stream); ihipEvent_t *eh = event._handle; if (eh && eh->_state != hipEventStatusUnitialized) { @@ -104,7 +104,7 @@ hipError_t hipEventRecord(hipEvent_t event, hipStream_t stream) //--- hipError_t hipEventDestroy(hipEvent_t event) { - std::call_once(hip_initialized, ihipInit); + HIP_INIT_API(event); event._handle->_state = hipEventStatusUnitialized; @@ -119,7 +119,7 @@ hipError_t hipEventDestroy(hipEvent_t event) //--- hipError_t hipEventSynchronize(hipEvent_t event) { - std::call_once(hip_initialized, ihipInit); + HIP_INIT_API(event); ihipEvent_t *eh = event._handle; @@ -148,7 +148,7 @@ hipError_t hipEventSynchronize(hipEvent_t event) //--- hipError_t hipEventElapsedTime(float *ms, hipEvent_t start, hipEvent_t stop) { - std::call_once(hip_initialized, ihipInit); + HIP_INIT_API(ms, start, stop); ihipEvent_t *start_eh = start._handle; ihipEvent_t *stop_eh = stop._handle; @@ -193,7 +193,7 @@ hipError_t hipEventElapsedTime(float *ms, hipEvent_t start, hipEvent_t stop) //--- hipError_t hipEventQuery(hipEvent_t event) { - std::call_once(hip_initialized, ihipInit); + HIP_INIT_API(event); ihipEvent_t *eh = event._handle; diff --git a/projects/clr/hipamd/src/hip_hcc.cpp b/projects/clr/hipamd/src/hip_hcc.cpp index a80db8fd2e..70890d9374 100644 --- a/projects/clr/hipamd/src/hip_hcc.cpp +++ b/projects/clr/hipamd/src/hip_hcc.cpp @@ -35,7 +35,7 @@ THE SOFTWARE. #include #include #include - +#include #include #include @@ -50,6 +50,9 @@ extern const char *ihipErrorString(hipError_t hip_error); const int release = 1; +#define MEMCPY_D2H_STAGING_VS_PININPLACE_COPY_THRESHOLD 4194304 +#define MEMCPY_H2D_DIRECT_VS_STAGING_COPY_THRESHOLD 65336 +#define MEMCPY_H2D_STAGING_VS_PININPLACE_COPY_THRESHOLD 1048576 int HIP_LAUNCH_BLOCKING = 0; @@ -60,6 +63,10 @@ int HIP_DB= 0; int HIP_STAGING_SIZE = 64; /* size of staging buffers, in KB */ int HIP_STAGING_BUFFERS = 2; // TODO - remove, two buffers should be enough. int HIP_PININPLACE = 0; +int HIP_OPTIMAL_MEM_TRANSFER = 0; //ENV Variable to test different memory transfer logics +int HIP_H2D_MEM_TRANSFER_THRESHOLD_DIRECT_OR_STAGING = 0; +int HIP_H2D_MEM_TRANSFER_THRESHOLD_STAGING_OR_PININPLACE = 0; +int HIP_D2H_MEM_TRANSFER_THRESHOLD = 0; int HIP_STREAM_SIGNALS = 2; /* number of signals to allocate at stream creation */ int HIP_VISIBLE_DEVICES = 0; /* Contains a comma-separated sequence of GPU identifiers */ @@ -359,6 +366,42 @@ void ihipStream_t::enqueueBarrier(hsa_queue_t* queue, ihipSignal_t *depSignal) hsa_signal_store_relaxed(queue->doorbell_signal, index); } +void ihipStream_t::enqueueBarrier(hsa_queue_t* queue, hsa_signal_t *depSignal) +{ + + // Obtain the write index for the command queue + uint64_t index = hsa_queue_load_write_index_relaxed(queue); + const uint32_t queueMask = queue->size - 1; + + // Define the barrier packet to be at the calculated queue index address + hsa_barrier_and_packet_t* barrier = &(((hsa_barrier_and_packet_t*)(queue->base_address))[index&queueMask]); + memset(barrier, 0, sizeof(hsa_barrier_and_packet_t)); + + // setup header + uint16_t header = HSA_PACKET_TYPE_BARRIER_AND << HSA_PACKET_HEADER_TYPE; + header |= 1 << HSA_PACKET_HEADER_BARRIER; + //header |= HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_ACQUIRE_FENCE_SCOPE; + //header |= HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_RELEASE_FENCE_SCOPE; + barrier->header = header; + + barrier->dep_signal[0].handle = 0; + barrier->dep_signal[1].handle = 0; + barrier->dep_signal[2].handle = 0; + barrier->dep_signal[3].handle = 0; + barrier->dep_signal[4].handle = 0; + + hsa_signal_t signal; + hsa_signal_create(1, 0, NULL, &signal); + *depSignal = signal; + barrier->completion_signal = signal; + + // TODO - check queue overflow, return error: + // Increment write index and ring doorbell to dispatch the kernel + hsa_queue_store_write_index_relaxed(queue, index+1); + hsa_signal_store_relaxed(queue->doorbell_signal, index); +} + + //-- //When the commands in a stream change types (ie kernel command follows a data command, @@ -429,15 +472,16 @@ int ihipStream_t::preCopyCommand(LockedAccessor_StreamCrit_t &crit, ihipSignal_t tprintf (DB_SYNC, "stream %p switch %s to %s (async copy dep on prev kernel)\n", this, ihipCommandName[crit->_last_command_type], ihipCommandName[copyType]); needSync = 1; - hsa_signal_t *hsaSignal = (static_cast (crit->_last_kernel_future.get_native_handle())); - if (hsaSignal) { + this->enqueueBarrier(static_cast(_av.get_hsa_queue()), waitSignal); +// hsa_signal_t *hsaSignal = (static_cast (crit->_last_kernel_future.get_native_handle())); +/* if (hsaSignal) { // Keep reference to the kernel future in order to keep the // dependent signal alive. _depFutures.push_back(crit->_last_kernel_future); *waitSignal = * hsaSignal; } else { assert(0); // if NULL signal, and we return 1, hsa_amd_memory_copy_async will fail. Confirm this never happens. - } + }*/ } else if (crit->_last_copy_signal) { needSync = 1; tprintf (DB_SYNC, "stream %p switch %s to %s (async copy dep on other copy #%lu)\n", @@ -579,11 +623,124 @@ ihipDevice_t::~ihipDevice_t() #define ErrorCheck(x) error_check(x, __LINE__, __FILE__) void error_check(hsa_status_t hsa_error_code, int line_num, std::string str) { - if (hsa_error_code != HSA_STATUS_SUCCESS) { + if ((hsa_error_code != HSA_STATUS_SUCCESS)&& (hsa_error_code != HSA_STATUS_INFO_BREAK)) { printf("HSA reported error!\n In file: %s\nAt line: %d\n", str.c_str(),line_num); } } +// CPU agent used for verification +hsa_agent_t cpu_agent_; +hsa_agent_t gpu_agent_; +int gpu_region_count; +// System region +hsa_amd_memory_pool_t sys_region_; +hsa_amd_memory_pool_t gpu_region_; + +hsa_status_t FindGpuDevice(hsa_agent_t agent, void* data) { + if (data == NULL) { + return HSA_STATUS_ERROR_INVALID_ARGUMENT; + } + + hsa_device_type_t hsa_device_type; + hsa_status_t hsa_error_code = + hsa_agent_get_info(agent, HSA_AGENT_INFO_DEVICE, &hsa_device_type); + if (hsa_error_code != HSA_STATUS_SUCCESS) { + return hsa_error_code; + } + + if (hsa_device_type == HSA_DEVICE_TYPE_GPU) { + *((hsa_agent_t*)data) = agent; + return HSA_STATUS_INFO_BREAK; + } + + return HSA_STATUS_SUCCESS; +} + +hsa_status_t FindCpuDevice(hsa_agent_t agent, void* data) { + if (data == NULL) { + return HSA_STATUS_ERROR_INVALID_ARGUMENT; + } + + hsa_device_type_t hsa_device_type; + hsa_status_t hsa_error_code = + hsa_agent_get_info(agent, HSA_AGENT_INFO_DEVICE, &hsa_device_type); + if (hsa_error_code != HSA_STATUS_SUCCESS) { + return hsa_error_code; + } + + if (hsa_device_type == HSA_DEVICE_TYPE_CPU) { + *((hsa_agent_t*)data) = agent; + return HSA_STATUS_INFO_BREAK; + } + + return HSA_STATUS_SUCCESS; +} + +hsa_status_t GetDeviceRegion(hsa_amd_memory_pool_t region, void* data) { + if (NULL == data) { + return HSA_STATUS_ERROR_INVALID_ARGUMENT; + } + + hsa_status_t err; + hsa_amd_segment_t segment; + uint32_t flag; + + err = hsa_amd_memory_pool_get_info(region, HSA_AMD_MEMORY_POOL_INFO_SEGMENT, &segment); + ErrorCheck(err); + if (HSA_AMD_SEGMENT_GLOBAL != segment) return HSA_STATUS_SUCCESS; + err = hsa_amd_memory_pool_get_info(region, HSA_AMD_MEMORY_POOL_INFO_GLOBAL_FLAGS, &flag); + ErrorCheck(err); + *((hsa_amd_memory_pool_t*)data) = region; + return HSA_STATUS_SUCCESS; +} + +hsa_status_t FindGlobalRegion(hsa_amd_memory_pool_t region, void* data) { + if (NULL == data) { + return HSA_STATUS_ERROR_INVALID_ARGUMENT; + } + + hsa_status_t err; + hsa_amd_segment_t segment; + uint32_t flag; + err = hsa_amd_memory_pool_get_info(region, HSA_AMD_MEMORY_POOL_INFO_SEGMENT, &segment); + ErrorCheck(err); + + err = hsa_amd_memory_pool_get_info(region, HSA_AMD_MEMORY_POOL_INFO_GLOBAL_FLAGS, &flag); + ErrorCheck(err); + if ((HSA_AMD_SEGMENT_GLOBAL == segment) && + (flag & HSA_AMD_MEMORY_POOL_GLOBAL_FLAG_FINE_GRAINED)) { + *((hsa_amd_memory_pool_t*)data) = region; + } + return HSA_STATUS_SUCCESS; +} + +void FindDeviceRegion() +{ + hsa_status_t err = hsa_iterate_agents(FindGpuDevice, &gpu_agent_); + ErrorCheck(err); + + err = hsa_amd_agent_iterate_memory_pools(gpu_agent_, GetDeviceRegion, &gpu_region_); + ErrorCheck(err); +} + +void FindSystemRegion() +{ + hsa_status_t err = hsa_iterate_agents(FindCpuDevice, &cpu_agent_); + ErrorCheck(err); + + err = hsa_amd_agent_iterate_memory_pools(cpu_agent_, FindGlobalRegion, &sys_region_); + ErrorCheck(err); +} + +int checkAccess(hsa_agent_t agent, hsa_amd_memory_pool_t pool) +{ + hsa_status_t err; + hsa_amd_memory_pool_access_t access; + err = hsa_amd_agent_memory_pool_get_info(agent, pool, HSA_AMD_AGENT_MEMORY_POOL_INFO_ACCESS, &access); + ErrorCheck(err); + return access; +} + hsa_status_t get_region_info(hsa_region_t region, void* data) { hsa_status_t err; @@ -720,6 +877,17 @@ hipError_t ihipDevice_t::getProperties(hipDeviceProp_t* prop) /* Computemode for HSA Devices is always : cudaComputeModeDefault */ prop->computeMode = 0; + FindSystemRegion(); + FindDeviceRegion(); + int access=checkAccess(cpu_agent_, gpu_region_); + if(0!= access){ + isLargeBar= 1; + } + else{ + isLargeBar=0; + } + + // Get Max Threads Per Multiprocessor HsaSystemProperties props; @@ -934,7 +1102,7 @@ static hsa_status_t findCpuAgent(hsa_agent_t agent, void *data) void ihipInit() { -#if COMPILE_TRACE_MARKER +#if COMPILE_HIP_ATP_MARKER amdtInitializeActivityLogger(); amdtScopedMarker("ihipInit", "HIP", NULL); #endif @@ -957,13 +1125,30 @@ void ihipInit() READ_ENV_I(release, HIP_STAGING_SIZE, 0, "Size of each staging buffer (in KB)" ); READ_ENV_I(release, HIP_STAGING_BUFFERS, 0, "Number of staging buffers to use in each direction. 0=use hsa_memory_copy."); READ_ENV_I(release, HIP_PININPLACE, 0, "For unpinned transfers, pin the memory in-place in chunks before doing the copy. Under development."); + READ_ENV_I(release, HIP_OPTIMAL_MEM_TRANSFER, 0, "For optimal memory transfers for unpinned memory.Under testing."); + READ_ENV_I(release, HIP_H2D_MEM_TRANSFER_THRESHOLD_DIRECT_OR_STAGING, 0, "Threshold value for H2D unpinned memory transfer decision between direct copy or staging buffer usage,Under testing."); + READ_ENV_I(release, HIP_H2D_MEM_TRANSFER_THRESHOLD_STAGING_OR_PININPLACE, 0, "Threshold value for H2D unpinned memory transfer decision between staging buffer usage or pininplace usage .Under testing."); + READ_ENV_I(release, HIP_D2H_MEM_TRANSFER_THRESHOLD, 0, "Threshold value for D2H unpinned memory transfer decision between staging buffer usage or pininplace usage .Under testing."); READ_ENV_I(release, HIP_STREAM_SIGNALS, 0, "Number of signals to allocate when new stream is created (signal pool will grow on demand)"); READ_ENV_I(release, HIP_VISIBLE_DEVICES, CUDA_VISIBLE_DEVICES, "Only devices whose index is present in the secquence are visible to HIP applications and they are enumerated in the order of secquence" ); READ_ENV_I(release, HIP_DISABLE_HW_KERNEL_DEP, 0, "Disable HW dependencies before kernel commands - instead wait for dependency on host. -1 means ignore these dependencies. (debug mode)"); READ_ENV_I(release, HIP_DISABLE_HW_COPY_DEP, 0, "Disable HW dependencies before copy commands - instead wait for dependency on host. -1 means ifnore these dependencies (debug mode)"); + if (HIP_OPTIMAL_MEM_TRANSFER && !HIP_H2D_MEM_TRANSFER_THRESHOLD_DIRECT_OR_STAGING) { + HIP_H2D_MEM_TRANSFER_THRESHOLD_DIRECT_OR_STAGING= MEMCPY_H2D_DIRECT_VS_STAGING_COPY_THRESHOLD; + fprintf (stderr, "warning: env var HIP_OPTIMAL_MEM_TRANSFER=0x%x but HIP_H2D_MEM_TRANSFER_THRESHOLD_DIRECT_OR_STAGING=0.Using default value for this.\n", HIP_OPTIMAL_MEM_TRANSFER); + } + if (HIP_OPTIMAL_MEM_TRANSFER && !HIP_H2D_MEM_TRANSFER_THRESHOLD_STAGING_OR_PININPLACE) { + HIP_H2D_MEM_TRANSFER_THRESHOLD_STAGING_OR_PININPLACE= MEMCPY_H2D_STAGING_VS_PININPLACE_COPY_THRESHOLD; + fprintf (stderr, "warning: env var HIP_OPTIMAL_MEM_TRANSFER=0x%x but HIP_H2D_MEM_TRANSFER_THRESHOLD_STAGING_OR_PININPLACE=0.Using default value for this.\n", HIP_OPTIMAL_MEM_TRANSFER); + } + + if (HIP_OPTIMAL_MEM_TRANSFER && !HIP_D2H_MEM_TRANSFER_THRESHOLD) { + HIP_D2H_MEM_TRANSFER_THRESHOLD= MEMCPY_D2H_STAGING_VS_PININPLACE_COPY_THRESHOLD; + fprintf (stderr, "warning: env var HIP_OPTIMAL_MEM_TRANSFER=0x%x but HIP_D2H_MEM_TRANSFER_THRESHOLD=0.Using default value for this.\n", HIP_OPTIMAL_MEM_TRANSFER); + } // Some flags have both compile-time and runtime flags - generate a warning if user enables the runtime flag but the compile-time flag is disabled. if (HIP_DB && !COMPILE_HIP_DB) { fprintf (stderr, "warning: env var HIP_DB=0x%x but COMPILE_HIP_DB=0. (perhaps enable COMPILE_HIP_DB in src code before compiling?)", HIP_DB); @@ -1085,13 +1270,35 @@ hipStream_t ihipSyncAndResolveStream(hipStream_t stream) } } +// HIP uses only 64 kernels. If the performance decrease, add more +uint32_t kernelCount = 0; +std::vector vCF(64); + +void incKernelCnt(hc::completion_future *cf){ + vCF[kernelCount] = cf; + kernelCount++; +} + +void decKernelCnt(){ + if(kernelCount > 63){ + uint32_t len = kernelCount; + for(uint32_t i =0;iwait(); + } + delete vCF[i]; + vCF[i] = NULL; + kernelCount--; + } + } +} // TODO - data-up to data-down: // Called just before a kernel is launched from hipLaunchKernel. // Allows runtime to track some information about the stream. hipStream_t ihipPreLaunchKernel(hipStream_t stream, dim3 grid, dim3 block, grid_launch_parm *lp) { - std::call_once(hip_initialized, ihipInit); + HIP_INIT_API(stream, grid, block, lp); stream = ihipSyncAndResolveStream(stream); #if USE_GRID_LAUNCH_20 lp->grid_dim.x = grid.x; @@ -1114,13 +1321,14 @@ hipStream_t ihipPreLaunchKernel(hipStream_t stream, dim3 grid, dim3 block, grid_ // *av = &stream->_av; lp->av = &stream->_av; lp->cf = new hc::completion_future; + incKernelCnt(lp->cf); // lp->av = static_cast(av); // lp->cf = static_cast(malloc(sizeof(hc::completion_future))); return (stream); } hipStream_t ihipPreLaunchKernel(hipStream_t stream, size_t grid, dim3 block, grid_launch_parm *lp) { - std::call_once(hip_initialized, ihipInit); + HIP_INIT_API(stream, grid, block, lp); stream = ihipSyncAndResolveStream(stream); #if USE_GRID_LAUNCH_20 lp->grid_dim.x = grid; @@ -1143,6 +1351,7 @@ hipStream_t ihipPreLaunchKernel(hipStream_t stream, size_t grid, dim3 block, gri // *av = &stream->_av; lp->av = &stream->_av; lp->cf = new hc::completion_future; + incKernelCnt(lp->cf); // lp->av = static_cast(av); // lp->cf = static_cast(malloc(sizeof(hc::completion_future))); return (stream); @@ -1150,7 +1359,7 @@ hipStream_t ihipPreLaunchKernel(hipStream_t stream, size_t grid, dim3 block, gri hipStream_t ihipPreLaunchKernel(hipStream_t stream, dim3 grid, size_t block, grid_launch_parm *lp) { - std::call_once(hip_initialized, ihipInit); + HIP_INIT_API(stream, grid, block, lp); stream = ihipSyncAndResolveStream(stream); #if USE_GRID_LAUNCH_20 lp->grid_dim.x = grid.x; @@ -1173,6 +1382,7 @@ hipStream_t ihipPreLaunchKernel(hipStream_t stream, dim3 grid, size_t block, gri // *av = &stream->_av; lp->av = &stream->_av; lp->cf = new hc::completion_future; + incKernelCnt(lp->cf); // lp->av = static_cast(av); // lp->cf = static_cast(malloc(sizeof(hc::completion_future))); return (stream); @@ -1180,7 +1390,7 @@ hipStream_t ihipPreLaunchKernel(hipStream_t stream, dim3 grid, size_t block, gri hipStream_t ihipPreLaunchKernel(hipStream_t stream, size_t grid, size_t block, grid_launch_parm *lp) { - std::call_once(hip_initialized, ihipInit); + HIP_INIT_API(stream, grid, block, lp); stream = ihipSyncAndResolveStream(stream); #if USE_GRID_LAUNCH_20 lp->grid_dim.x = grid; @@ -1203,6 +1413,7 @@ hipStream_t ihipPreLaunchKernel(hipStream_t stream, size_t grid, size_t block, g // *av = &stream->_av; lp->av = &stream->_av; lp->cf = new hc::completion_future; + incKernelCnt(lp->cf); // lp->av = static_cast(av); // lp->cf = static_cast(malloc(sizeof(hc::completion_future))); return (stream); @@ -1215,6 +1426,7 @@ void ihipPostLaunchKernel(hipStream_t stream, grid_launch_parm &lp) { // stream->lockclose_postKernelCommand(cf); stream->lockclose_postKernelCommand(*lp.cf); + decKernelCnt(); if (HIP_LAUNCH_BLOCKING) { tprintf(DB_SYNC, " stream:%p LAUNCH_BLOCKING for kernel completion\n", stream); } @@ -1226,8 +1438,8 @@ void ihipPostLaunchKernel(hipStream_t stream, grid_launch_parm &lp) // HIP API Implementation // // Implementor notes: -// _ All functions should call ihipInit as first action: -// std::call_once(hip_initialized, ihipInit); +// _ All functions should call HIP_INIT_API as first action: +// HIP_INIT_API(); // // - ALl functions should use ihipLogStatus to return error code (not return error directly). //================================================================================================= @@ -1337,7 +1549,6 @@ void ihipStream_t::setAsyncCopyAgents(unsigned kind, ihipCommand_t *commandType, void ihipStream_t::copySync(LockedAccessor_StreamCrit_t &crit, void* dst, const void* src, size_t sizeBytes, unsigned kind) { ihipDevice_t *device = this->getDevice(); - if (device == NULL) { throw ihipException(hipErrorInvalidDevice); } @@ -1374,16 +1585,32 @@ void ihipStream_t::copySync(LockedAccessor_StreamCrit_t &crit, void* dst, const if(!srcTracked){ if (HIP_STAGING_BUFFERS) { tprintf(DB_COPY1, "D2H && !dstTracked: staged copy H2D dst=%p src=%p sz=%zu\n", dst, src, sizeBytes); - - if (HIP_PININPLACE) { - device->_staging_buffer[0]->CopyHostToDevicePinInPlace(dst, src, sizeBytes, depSignalCnt ? &depSignal : NULL); - } else { - device->_staging_buffer[0]->CopyHostToDevice(dst, src, sizeBytes, depSignalCnt ? &depSignal : NULL); + if(HIP_OPTIMAL_MEM_TRANSFER) + { + if((device->isLargeBar)&&(sizeBytes < HIP_H2D_MEM_TRANSFER_THRESHOLD_DIRECT_OR_STAGING)){ + memcpy(dst,src,sizeBytes); + std::atomic_thread_fence(std::memory_order_release); + } + else{ + if(sizeBytes > HIP_H2D_MEM_TRANSFER_THRESHOLD_STAGING_OR_PININPLACE){ + //if (HIP_PININPLACE) { + device->_staging_buffer[0]->CopyHostToDevicePinInPlace(dst, src, sizeBytes, depSignalCnt ? &depSignal : NULL); + } else { + device->_staging_buffer[0]->CopyHostToDevice(dst, src, sizeBytes, depSignalCnt ? &depSignal : NULL); + } + // The copy waits for inputs and then completes before returning so can reset queue to empty: + this->wait(crit, true); + } } - - // The copy waits for inputs and then completes before returning so can reset queue to empty: - this->wait(crit, true); - } else { + else { + if (HIP_PININPLACE) { + device->_staging_buffer[0]->CopyHostToDevicePinInPlace(dst, src, sizeBytes, depSignalCnt ? &depSignal : NULL); + } else { + device->_staging_buffer[0]->CopyHostToDevice(dst, src, sizeBytes, depSignalCnt ? &depSignal : NULL); + } + } + } + else { // TODO - remove, slow path. tprintf(DB_COPY1, "H2D && ! srcTracked: am_copy dst=%p src=%p sz=%zu\n", dst, src, sizeBytes); #if USE_AV_COPY @@ -1418,8 +1645,22 @@ void ihipStream_t::copySync(LockedAccessor_StreamCrit_t &crit, void* dst, const if (HIP_STAGING_BUFFERS) { tprintf(DB_COPY1, "D2H && !dstTracked: staged copy D2H dst=%p src=%p sz=%zu\n", dst, src, sizeBytes); //printf ("staged-copy- read dep signals\n"); - device->_staging_buffer[1]->CopyDeviceToHost(dst, src, sizeBytes, depSignalCnt ? &depSignal : NULL); - + if(HIP_OPTIMAL_MEM_TRANSFER) + { + if(sizeBytes> HIP_D2H_MEM_TRANSFER_THRESHOLD){ + device->_staging_buffer[1]->CopyDeviceToHostPinInPlace(dst, src, sizeBytes, depSignalCnt ? &depSignal : NULL); + }else { + //printf ("staged-copy- read dep signals\n"); + device->_staging_buffer[1]->CopyDeviceToHost(dst, src, sizeBytes, depSignalCnt ? &depSignal : NULL); + } + }else + { + device->_staging_buffer[1]->CopyDeviceToHost(dst, src, sizeBytes, depSignalCnt ? &depSignal : NULL); + } + if(crit->_last_command_type == ihipCommandKernel){ + std::cout<<"Destroying depSignal MemcpySync"<wait(crit, true); @@ -1502,6 +1743,10 @@ void ihipStream_t::copySync(LockedAccessor_StreamCrit_t &crit, void* dst, const throw ihipException(hipErrorInvalidValue); } } + if(crit->_last_command_type == ihipCommandKernel){ + hsa_signal_destroy(depSignal); + } + } @@ -1575,6 +1820,9 @@ void ihipStream_t::copyAsync(void* dst, const void* src, size_t sizeBytes, unsig hsa_status_t hsa_status = hsa_amd_memory_async_copy(dst, dstAgent, src, srcAgent, sizeBytes, depSignalCnt, depSignalCnt ? &depSignal:0x0, ihip_signal->_hsa_signal); + if (crit->_last_command_type == ihipCommandKernel) { + hsa_signal_destroy(depSignal); + } if (hsa_status == HSA_STATUS_SUCCESS) { if (HIP_LAUNCH_BLOCKING) { @@ -1602,7 +1850,7 @@ void ihipStream_t::copyAsync(void* dst, const void* src, size_t sizeBytes, unsig //--- hipError_t hipHccGetAccelerator(int deviceId, hc::accelerator *acc) { - std::call_once(hip_initialized, ihipInit); + HIP_INIT_API(deviceId, acc); ihipDevice_t *d = ihipGetDevice(deviceId); hipError_t err; @@ -1622,7 +1870,7 @@ hipError_t hipHccGetAccelerator(int deviceId, hc::accelerator *acc) //--- hipError_t hipHccGetAcceleratorView(hipStream_t stream, hc::accelerator_view **av) { - std::call_once(hip_initialized, ihipInit); + HIP_INIT_API(stream, av); if (stream == hipStreamNull ) { ihipDevice_t *device = ihipGetTlsDefaultDevice(); diff --git a/projects/clr/hipamd/src/hip_memory.cpp b/projects/clr/hipamd/src/hip_memory.cpp index cabea6f759..94442f4698 100644 --- a/projects/clr/hipamd/src/hip_memory.cpp +++ b/projects/clr/hipamd/src/hip_memory.cpp @@ -36,7 +36,7 @@ THE SOFTWARE. */ hipError_t hipPointerGetAttributes(hipPointerAttribute_t *attributes, void* ptr) { - std::call_once(hip_initialized, ihipInit); + HIP_INIT_API(attributes, ptr); hipError_t e = hipSuccess; @@ -85,7 +85,7 @@ hipError_t hipPointerGetAttributes(hipPointerAttribute_t *attributes, void* ptr) */ hipError_t hipHostGetDevicePointer(void **devicePointer, void *hostPointer, unsigned flags) { - std::call_once(hip_initialized, ihipInit); + HIP_INIT_API(devicePointer, hostPointer, flags); hipError_t e = hipSuccess; @@ -198,6 +198,116 @@ hipError_t hipMallocHost(void** ptr, size_t sizeBytes) return hipHostMalloc(ptr, sizeBytes, 0); } +// width in bytes +hipError_t hipMallocPitch(void** ptr, size_t* pitch, size_t width, size_t height) { + + HIP_INIT_API(ptr, pitch, width, height); + + hipError_t hip_status = hipSuccess; + + if(width == 0 || height == 0) + return ihipLogStatus(hipErrorUnknown); + + // hardcoded 128 bytes + *pitch = ((((int)width-1)/128) + 1)*128; + const size_t sizeBytes = (*pitch)*height; + + auto device = ihipGetTlsDefaultDevice(); + + //err = hipMalloc(ptr, (*pitch)*height); + if (device) { + const unsigned am_flags = 0; + *ptr = hc::am_alloc(sizeBytes, device->_acc, am_flags); + + if (sizeBytes && (*ptr == NULL)) { + hip_status = hipErrorMemoryAllocation; + } else { + hc::am_memtracker_update(*ptr, device->_device_index, 0); + { + LockedAccessor_DeviceCrit_t crit(device->criticalData()); + if (crit->peerCnt() > 1) { // peerCnt includes self so only call allow_access if other peers involved: + hsa_status_t hsa_status = hsa_amd_agents_allow_access(crit->peerCnt(), crit->peerAgents(), NULL, *ptr); + if (hsa_status != HSA_STATUS_SUCCESS) { + hip_status = hipErrorMemoryAllocation; + } + } + } + } + } else { + hip_status = hipErrorMemoryAllocation; + } + + return ihipLogStatus(hip_status); + +} + +hipChannelFormatDesc hipCreateChannelDesc(int x, int y, int z, int w, hipChannelFormatKind f) { + hipChannelFormatDesc cd; + cd.x = x; cd.y = y; cd.z = z; cd.w = w; + cd.f = f; + return cd; +} + +hipError_t hipMallocArray(hipArray** array, const hipChannelFormatDesc* desc, + size_t width, size_t height, unsigned int flags) { + + HIP_INIT_API(array, desc, width, height, flags); + + hipError_t hip_status = hipSuccess; + + auto device = ihipGetTlsDefaultDevice(); + + *array = (hipArray*)malloc(sizeof(hipArray)); + array[0]->width = width; + array[0]->height = height; + + array[0]->f = desc->f; + + void ** ptr = &array[0]->data; + + if (device) { + const unsigned am_flags = 0; + const size_t size = width*height; + + switch(desc->f) { + case hipChannelFormatKindSigned: + *ptr = hc::am_alloc(size*sizeof(int), device->_acc, am_flags); + break; + case hipChannelFormatKindUnsigned: + *ptr = hc::am_alloc(size*sizeof(unsigned int), device->_acc, am_flags); + break; + case hipChannelFormatKindFloat: + *ptr = hc::am_alloc(size*sizeof(float), device->_acc, am_flags); + break; + case hipChannelFormatKindNone: + *ptr = hc::am_alloc(size*sizeof(size_t), device->_acc, am_flags); + break; + default: + hip_status = hipErrorUnknown; + break; + } + if (size && (*ptr == NULL)) { + hip_status = hipErrorMemoryAllocation; + } else { + hc::am_memtracker_update(*ptr, device->_device_index, 0); + { + LockedAccessor_DeviceCrit_t crit(device->criticalData()); + if (crit->peerCnt() > 1) { // peerCnt includes self so only call allow_access if other peers involved: + hsa_status_t hsa_status = hsa_amd_agents_allow_access(crit->peerCnt(), crit->peerAgents(), NULL, *ptr); + if (hsa_status != HSA_STATUS_SUCCESS) { + hip_status = hipErrorMemoryAllocation; + } + } + } + } + + } else { + hip_status = hipErrorMemoryAllocation; + } + + return ihipLogStatus(hip_status); +} + //--- hipError_t hipHostGetFlags(unsigned int* flagsPtr, void* hostPtr) @@ -333,7 +443,7 @@ hipError_t hipMemcpy(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind /** - * @result #hipSuccess, #hipErrorInvalidDevice, #hipErrorInvalidMemcpyDirection, + * @result #hipSuccess, #hipErrorInvalidDevice, #hipErrorInvalidMemcpyDirection, * @result #hipErrorInvalidValue : If dst==NULL or src==NULL, or other bad argument. * @warning on HCC hipMemcpyAsync does not support overlapped H2D and D2H copies. * @warning on HCC hipMemcpyAsync requires that any host pointers are pinned (ie via the hipMallocHost call). @@ -364,6 +474,108 @@ hipError_t hipMemcpyAsync(void* dst, const void* src, size_t sizeBytes, hipMemcp return ihipLogStatus(e); } +// dpitch, spitch, and width in bytes +hipError_t hipMemcpy2D(void* dst, size_t dpitch, const void* src, size_t spitch, + size_t width, size_t height, hipMemcpyKind kind) { + + HIP_INIT_API(dst, dpitch, src, spitch, width, height, kind); + + if(width > dpitch || width > spitch) + return ihipLogStatus(hipErrorUnknown); + + hipStream_t stream = ihipSyncAndResolveStream(hipStreamNull); + + hc::completion_future marker; + + hipError_t e = hipSuccess; + + try { + for(int i = 0; i < height; ++i) { + stream->locked_copySync((unsigned char*)dst + i*dpitch, (unsigned char*)src + i*spitch, width, kind); + } + } + catch (ihipException ex) { + e = ex._code; + } + + return ihipLogStatus(e); +} + +// wOffset, width, and spitch in bytes +hipError_t hipMemcpy2DToArray(hipArray* dst, size_t wOffset, size_t hOffset, const void* src, + size_t spitch, size_t width, size_t height, hipMemcpyKind kind) { + + HIP_INIT_API(dst, wOffset, hOffset, src, spitch, width, height, kind); + + hipStream_t stream = ihipSyncAndResolveStream(hipStreamNull); + + hc::completion_future marker; + + hipError_t e = hipSuccess; + + size_t byteSize; + if(dst) { + switch(dst[0].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; + } + } else { + return ihipLogStatus(hipErrorUnknown); + } + + if((wOffset + width > (dst->width * byteSize)) || width > spitch) { + return ihipLogStatus(hipErrorUnknown); + } + + size_t src_w = spitch; + size_t dst_w = (dst->width)*byteSize; + + try { + for(int i = 0; i < height; ++i) { + stream->locked_copySync((unsigned char*)dst->data + i*dst_w, (unsigned char*)src + i*src_w, width, kind); + } + } + catch (ihipException ex) { + e = ex._code; + } + + return ihipLogStatus(e); +} + +hipError_t hipMemcpyToArray(hipArray* dst, size_t wOffset, size_t hOffset, + const void* src, size_t count, hipMemcpyKind kind) { + + HIP_INIT_API(dst, wOffset, hOffset, src, count, kind); + + hipStream_t stream = ihipSyncAndResolveStream(hipStreamNull); + + hc::completion_future marker; + + hipError_t e = hipSuccess; + + try { + stream->locked_copySync((char *)dst->data + wOffset, src, count, kind); + } + catch (ihipException ex) { + e = ex._code; + } + + return ihipLogStatus(e); +} + // TODO-sync: function is async unless target is pinned host memory - then these are fully sync. /** @return #hipErrorInvalidValue @@ -491,7 +703,7 @@ hipError_t hipMemGetInfo (size_t *free, size_t *total) size_t deviceMemSize, hostMemSize, userMemSize; hc::am_memtracker_sizeinfo(hipDevice->_acc, &deviceMemSize, &hostMemSize, &userMemSize); printf ("deviceMemSize=%zu\n", deviceMemSize); - + *free = hipDevice->_props.totalGlobalMem - deviceMemSize; } @@ -525,7 +737,7 @@ hipError_t hipFree(void* ptr) } } else { // free NULL pointer succeeds and is common technique to initialize runtime - hipStatus = hipSuccess; + hipStatus = hipSuccess; } return ihipLogStatus(hipStatus); @@ -553,7 +765,7 @@ hipError_t hipHostFree(void* ptr) } } else { // free NULL pointer succeeds and is common technique to initialize runtime - hipStatus = hipSuccess; + hipStatus = hipSuccess; } return ihipLogStatus(hipStatus); @@ -566,5 +778,40 @@ hipError_t hipFreeHost(void* ptr) return hipHostFree(ptr); } +hipError_t hipFreeArray(hipArray* array) +{ + HIP_INIT_API(array); + hipError_t hipStatus = hipErrorInvalidDevicePointer; + + // Synchronize to ensure all work has finished. + ihipGetTlsDefaultDevice()->locked_waitAllStreams(); // ignores non-blocking streams, this waits for all activity to finish. + + if(array->data) { + hc::accelerator acc; + hc::AmPointerInfo amPointerInfo(NULL, NULL, 0, acc, 0, 0); + am_status_t status = hc::am_memtracker_getinfo(&amPointerInfo, array->data); + if(status == AM_SUCCESS){ + if(amPointerInfo._hostPointer == NULL){ + hc::am_free(array->data); + hipStatus = hipSuccess; + } + } + } + + return ihipLogStatus(hipStatus); +} + +// Stubs of threadfence operations +__device__ void __threadfence_block(void){ + // no-op +} + +__device__ void __threadfence(void){ + // no-op +} + +__device__ void __threadfence_system(void){ + // no-op +} diff --git a/projects/clr/hipamd/src/staging_buffer.cpp b/projects/clr/hipamd/src/staging_buffer.cpp index be5058e47b..c6c23089bd 100644 --- a/projects/clr/hipamd/src/staging_buffer.cpp +++ b/projects/clr/hipamd/src/staging_buffer.cpp @@ -88,42 +88,48 @@ void StagingBuffer::CopyHostToDevicePinInPlace(void* dst, const void* src, size_ THROW_ERROR (hipErrorInvalidValue); } int bufferIndex = 0; +#if 0 for (int64_t bytesRemaining=sizeBytes; bytesRemaining>0 ; bytesRemaining -= _bufferSize) { size_t theseBytes = (bytesRemaining > _bufferSize) ? _bufferSize : bytesRemaining; +#endif + size_t theseBytes= sizeBytes; + //tprintf (DB_COPY2, "H2D: waiting... on completion signal handle=%lu\n", _completion_signal[bufferIndex].handle); + //hsa_signal_wait_acquire(_completion_signal[bufferIndex], HSA_SIGNAL_CONDITION_LT, 1, UINT64_MAX, HSA_WAIT_STATE_ACTIVE); - tprintf (DB_COPY2, "H2D: waiting... on completion signal handle=%lu\n", _completion_signal[bufferIndex].handle); - hsa_signal_wait_acquire(_completion_signal[bufferIndex], HSA_SIGNAL_CONDITION_LT, 1, UINT64_MAX, HSA_WAIT_STATE_ACTIVE); + //void * masked_srcp = (void*) ((uintptr_t)srcp & (uintptr_t)(~0x3f)) ; // TODO + void *locked_srcp; + //hsa_status_t hsa_status = hsa_amd_memory_lock(masked_srcp, theseBytes, &_hsa_agent, 1, &locked_srcp); + hsa_status_t hsa_status = hsa_amd_memory_lock(const_cast (srcp), theseBytes, &_hsa_agent, 1, &locked_srcp); + //tprintf (DB_COPY2, "H2D: bytesRemaining=%zu: pin-in-place:%p+%zu bufferIndex[%d]\n", bytesRemaining, srcp, theseBytes, bufferIndex); + //printf ("status=%x srcp=%p, masked_srcp=%p, locked_srcp=%p\n", hsa_status, srcp, masked_srcp, locked_srcp); - void * masked_srcp = (void*) ((uintptr_t)srcp & (uintptr_t)(~0x3f)) ; // TODO - void *locked_srcp; - hsa_status_t hsa_status = hsa_amd_memory_lock(masked_srcp, theseBytes, &_hsa_agent, 1, &locked_srcp); - //hsa_status_t hsa_status = hsa_amd_memory_lock(const_cast (srcp), theseBytes, &_hsa_agent, 1, &locked_srcp); - tprintf (DB_COPY2, "H2D: bytesRemaining=%zu: pin-in-place:%p+%zu bufferIndex[%d]\n", bytesRemaining, srcp, theseBytes, bufferIndex); - printf ("status=%x srcp=%p, masked_srcp=%p, locked_srcp=%p\n", hsa_status, srcp, masked_srcp, locked_srcp); + if (hsa_status != HSA_STATUS_SUCCESS) { + THROW_ERROR (hipErrorRuntimeMemory); + } - if (hsa_status != HSA_STATUS_SUCCESS) { - THROW_ERROR (hipErrorRuntimeMemory); - } + hsa_signal_store_relaxed(_completion_signal[bufferIndex], 1); - hsa_signal_store_relaxed(_completion_signal[bufferIndex], 1); - - hsa_status = hsa_amd_memory_async_copy(dstp, _hsa_agent, locked_srcp, g_cpu_agent, theseBytes, waitFor ? 1:0, waitFor, _completion_signal[bufferIndex]); - tprintf (DB_COPY2, "H2D: bytesRemaining=%zu: async_copy %zu bytes %p to %p status=%x\n", bytesRemaining, theseBytes, _pinnedStagingBuffer[bufferIndex], dstp, hsa_status); - - if (hsa_status != HSA_STATUS_SUCCESS) { - THROW_ERROR (hipErrorRuntimeMemory); - } + hsa_status = hsa_amd_memory_async_copy(dstp, _hsa_agent, locked_srcp, g_cpu_agent, theseBytes, waitFor ? 1:0, waitFor, _completion_signal[bufferIndex]); + //tprintf (DB_COPY2, "H2D: bytesRemaining=%zu: async_copy %zu bytes %p to %p status=%x\n", bytesRemaining, theseBytes, _pinnedStagingBuffer[bufferIndex], dstp, hsa_status); + if (hsa_status != HSA_STATUS_SUCCESS) { + THROW_ERROR (hipErrorRuntimeMemory); + } + tprintf (DB_COPY2, "H2D: waiting... on completion signal handle=%lu\n", _completion_signal[bufferIndex].handle); + hsa_signal_wait_acquire(_completion_signal[bufferIndex], HSA_SIGNAL_CONDITION_LT, 1, UINT64_MAX, HSA_WAIT_STATE_ACTIVE); + hsa_amd_memory_unlock(const_cast (srcp)); +#if 0 srcp += theseBytes; dstp += theseBytes; if (++bufferIndex >= _numBuffers) { bufferIndex = 0; } - - // Assume subsequent commands are dependent on previous and don't need dependency after first copy submitted, HIP_ONESHOT_COPY_DEP=1 - waitFor = NULL; - } +#endif + // Assume subsequent commands are dependent on previous and don't need dependency after first copy submitted, HIP_ONESHOT_COPY_DEP=1 + waitFor = NULL; +#if 0 +// } // TODO - printf ("unpin the memory\n"); @@ -132,6 +138,7 @@ void StagingBuffer::CopyHostToDevicePinInPlace(void* dst, const void* src, size_ for (int i=0; i<_numBuffers; i++) { hsa_signal_wait_acquire(_completion_signal[i], HSA_SIGNAL_CONDITION_LT, 1, UINT64_MAX, HSA_WAIT_STATE_ACTIVE); } +#endif } @@ -194,6 +201,47 @@ void StagingBuffer::CopyHostToDevice(void* dst, const void* src, size_t sizeByte } } + +void StagingBuffer::CopyDeviceToHostPinInPlace(void* dst, const void* src, size_t sizeBytes, hsa_signal_t *waitFor) +{ + std::lock_guard l (_copy_lock); + + const char *srcp = static_cast (src); + char *dstp = static_cast (dst); + + for (int i=0; i<_numBuffers; i++) { + hsa_signal_store_relaxed(_completion_signal[i], 0); + } + + if (sizeBytes >= UINT64_MAX/2) { + THROW_ERROR (hipErrorInvalidValue); + } + int bufferIndex = 0; + size_t theseBytes= sizeBytes; + void *locked_destp; + + hsa_status_t hsa_status = hsa_amd_memory_lock(const_cast (dstp), theseBytes, &_hsa_agent, 1, &locked_destp); + + + if (hsa_status != HSA_STATUS_SUCCESS) { + THROW_ERROR (hipErrorRuntimeMemory); + } + + hsa_signal_store_relaxed(_completion_signal[bufferIndex], 1); + + hsa_status = hsa_amd_memory_async_copy(locked_destp,g_cpu_agent , srcp, _hsa_agent, theseBytes, waitFor ? 1:0, waitFor, _completion_signal[bufferIndex]); + + if (hsa_status != HSA_STATUS_SUCCESS) { + THROW_ERROR (hipErrorRuntimeMemory); + } + tprintf (DB_COPY2, "D2H: waiting... on completion signal handle=%lu\n", _completion_signal[bufferIndex].handle); + hsa_signal_wait_acquire(_completion_signal[bufferIndex], HSA_SIGNAL_CONDITION_LT, 1, UINT64_MAX, HSA_WAIT_STATE_ACTIVE); + hsa_amd_memory_unlock(const_cast (dstp)); + + // Assume subsequent commands are dependent on previous and don't need dependency after first copy submitted, HIP_ONESHOT_COPY_DEP=1 + waitFor = NULL; +} + //--- //Copies sizeBytes from src to dst, using either a copy to a staging buffer or a staged pin-in-place strategy //IN: dst - dest pointer - must be accessible from agent this buffer is associated with (via _hsa_agent). diff --git a/projects/clr/hipamd/tests/src/CMakeLists.txt b/projects/clr/hipamd/tests/src/CMakeLists.txt index 723b99a243..127ebb2507 100644 --- a/projects/clr/hipamd/tests/src/CMakeLists.txt +++ b/projects/clr/hipamd/tests/src/CMakeLists.txt @@ -178,6 +178,10 @@ build_hip_executable (hipTestMemcpyPin hipTestMemcpyPin.cpp) #build_hip_executable (hipDynamicShared hipDynamicShared.cpp) build_hip_executable (hipLaunchParm hipLaunchParm.cpp) +if (${HIP_PLATFORM} STREQUAL "hcc") + build_hip_executable (hipArray hipArray.cpp) +endif() + make_test(hipEventRecord --iterations 10) make_test(hipEnvVarDriver " " ) make_test(hipLaunchParm " ") @@ -204,6 +208,7 @@ if (${HIP_MULTI_GPU}) endif() if (${HIP_PLATFORM} STREQUAL "hcc") + make_test(hipArray " ") make_test(hipFuncSetDevice " ") endif() diff --git a/projects/clr/hipamd/tests/src/hipArray.cpp b/projects/clr/hipamd/tests/src/hipArray.cpp new file mode 100644 index 0000000000..49add786d8 --- /dev/null +++ b/projects/clr/hipamd/tests/src/hipArray.cpp @@ -0,0 +1,239 @@ +#include "hip_runtime.h" +#include "test_common.h" + +void printSep() +{ + printf ("======================================================================================\n"); +} + + +//--- +// Test copies of a matrix numW by numH +// The subroutine allocates memory , copies to device, runs a vector add kernel, copies back, and checks the result. +// +// IN: numW: number of elements in the 1st dimension used for allocation +// IN: numH: number of elements in the 2nd dimension used for allocation +// IN: usePinnedHost : If true, allocate host with hipHostMalloc and is pinned ; else allocate host memory with malloc. +// +template +void memcpy2Dtest(size_t numW, size_t numH, bool usePinnedHost) +{ + + size_t width = numW * sizeof(T); + size_t sizeElements = width * numH; + + printf("memcpy2Dtest: %s<%s> size=%lu (%6.2fMB) W: %d, H:%d, usePinnedHost: %d\n", + __func__, + TYPENAME(T), + sizeElements, sizeElements/1024.0/1024.0, + (int)numW, (int)numH, usePinnedHost); + + T *A_d, *B_d, *C_d; + T *A_h, *B_h, *C_h; + + size_t pitch_A, pitch_B, pitch_C; + + hipChannelFormatDesc desc = hipCreateChannelDesc(); + HipTest::initArrays2DPitch(&A_d, &B_d, &C_d, &pitch_A, &pitch_B, &pitch_C, numW, numH); + HipTest::initArraysForHost(&A_h, &B_h, &C_h, numW*numH, usePinnedHost); + unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, numW*numH); + + HIPCHECK (hipMemcpy2D (A_d, pitch_A, A_h, width, width, numH, hipMemcpyHostToDevice) ); + HIPCHECK (hipMemcpy2D (B_d, pitch_B, B_h, width, width, numH, hipMemcpyHostToDevice) ); + + hipLaunchKernel(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock), 0, 0, A_d, B_d, C_d, (pitch_C/sizeof(T))*numH); + + HIPCHECK (hipMemcpy2D (C_h, width, C_d, pitch_C, width, numH, hipMemcpyDeviceToHost) ); + + HIPCHECK ( hipDeviceSynchronize() ); + HipTest::checkVectorADD(A_h, B_h, C_h, numW*numH); + + HipTest::freeArrays (A_d, B_d, C_d, A_h, B_h, C_h, usePinnedHost); + + printf (" %s success\n", __func__); +} + +//--- +// Test copies of a matrix numW by numH into a hipArray data structure +// The subroutine allocates memory , copies to device, runs a vector add kernel, copies back, and checks the result. +// +// IN: numW: number of elements in the 1st dimension used for allocation +// IN: numH: number of elements in the 2nd dimension used for allocation. If this is 1, then the 1-dimensional copy API +// would be used +// IN: usePinnedHost : If true, allocate host with hipHostMalloc and is pinned ; else allocate host memory with malloc. +// IN: usePitch: If true, pads additional memory. This is only valid in the 2-dimensional case +// +template +void memcpyArraytest(size_t numW, size_t numH, bool usePinnedHost, bool usePitch=false) +{ + + size_t width = numW * sizeof(T); + size_t sizeElements = width * numH; + + printf("memcpyArraytest: %s<%s> size=%lu (%6.2fMB) W: %d, H: %d, usePinnedHost: %d, usePitch: %d\n", + __func__, + TYPENAME(T), + sizeElements, sizeElements/1024.0/1024.0, + (int)numW, (int)numH, usePinnedHost, usePitch); + + hipArray *A_d, *B_d, *C_d; + T *A_h, *B_h, *C_h; + + // 1D + if ((numW >= 1) && (numH == 1)) { + hipChannelFormatDesc desc = hipCreateChannelDesc(); + HipTest::initHIPArrays(&A_d, &B_d, &C_d, &desc, numW, 1, 0); + HipTest::initArraysForHost(&A_h, &B_h, &C_h, numW*numH, usePinnedHost); + unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, numW*numH); + + HIPCHECK (hipMemcpyToArray (A_d, 0, 0, (void *)A_h, width, hipMemcpyHostToDevice) ); + HIPCHECK (hipMemcpyToArray (B_d, 0, 0, (void *)B_h, width, hipMemcpyHostToDevice) ); + + hipLaunchKernel(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock), 0, 0, (T*)A_d->data, (T*)B_d->data, (T*)C_d->data, numW); + + HIPCHECK (hipMemcpy (C_h, C_d->data, width, hipMemcpyDeviceToHost) ); + + HIPCHECK ( hipDeviceSynchronize() ); + HipTest::checkVectorADD(A_h, B_h, C_h, numW); + + } + // 2D + else if ((numW >= 1) && (numH >= 1)) { + + + hipChannelFormatDesc desc = hipCreateChannelDesc(); + HipTest::initHIPArrays(&A_d, &B_d, &C_d, &desc, numW, numH, 0); + HipTest::initArraysForHost(&A_h, &B_h, &C_h, numW*numH, usePinnedHost); + unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, numW*numH); + + if (usePitch) { + T *A_p, *B_p, *C_p; + size_t pitch_A, pitch_B, pitch_C; + + HipTest::initArrays2DPitch(&A_p, &B_p, &C_p, &pitch_A, &pitch_B, &pitch_C, numW, numH); + HIPCHECK (hipMemcpy2D (A_p, pitch_A, A_h, width, width, numH, hipMemcpyHostToDevice) ); + HIPCHECK (hipMemcpy2D (B_p, pitch_B, B_h, width, width, numH, hipMemcpyHostToDevice) ); + + HIPCHECK (hipMemcpy2DToArray (A_d, 0, 0, (void *)A_p, pitch_A, width, numH, hipMemcpyDeviceToDevice) ); + HIPCHECK (hipMemcpy2DToArray (B_d, 0, 0, (void *)B_p, pitch_B, width, numH, hipMemcpyDeviceToDevice) ); + + hipFree(A_p); + hipFree(B_p); + hipFree(C_p); + } + else { + HIPCHECK (hipMemcpy2DToArray (A_d, 0, 0, (void *)A_h, width, width, numH, hipMemcpyHostToDevice) ); + HIPCHECK (hipMemcpy2DToArray (B_d, 0, 0, (void *)B_h, width, width, numH, hipMemcpyHostToDevice) ); + } + + hipLaunchKernel(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock), 0, 0, (T*)A_d->data, (T*)B_d->data, (T*)C_d->data, numW*numH); + + HIPCHECK (hipMemcpy2D ((void*)C_h, width, (void*)C_d->data, width, width, numH, hipMemcpyDeviceToHost) ); + + HIPCHECK ( hipDeviceSynchronize() ); + HipTest::checkVectorADD(A_h, B_h, C_h, numW*numH); + } + // Unknown + else { + HIPASSERT("Incompatible dimensions" && 0); + } + + hipFreeArray(A_d); + hipFreeArray(B_d); + hipFreeArray(C_d); + HipTest::freeArraysForHost(A_h, B_h, C_h, usePinnedHost); + + printf (" %s success\n", __func__); + +} + +//--- +//Try many different sizes to memory copy. +template +void memcpyArraytest_size(size_t maxElem=0, size_t offset=0) +{ + printf ("test: %s<%s>\n", __func__, TYPENAME(T)); + + int deviceId; + HIPCHECK(hipGetDevice(&deviceId)); + + size_t free, total; + HIPCHECK(hipMemGetInfo(&free, &total)); + + if (maxElem == 0) { + maxElem = free/sizeof(T)/5; + } + + printf (" device#%d: hipMemGetInfo: free=%zu (%4.2fMB) total=%zu (%4.2fMB) maxSize=%6.1fMB offset=%lu\n", + deviceId, free, (float)(free/1024.0/1024.0), total, (float)(total/1024.0/1024.0), maxElem*sizeof(T)/1024.0/1024.0, offset); + + // Test 1D + for (size_t elem=64; elem+offset<=maxElem; elem*=2) { + HIPCHECK ( hipDeviceReset() ); + memcpyArraytest(elem+offset, 1, 0); // unpinned host + HIPCHECK ( hipDeviceReset() ); + memcpyArraytest(elem+offset, 1, 1); // pinned host + } + + // Test 2D + size_t maxElem2D = sqrt(maxElem); + + for (size_t elem=64; elem+offset<=maxElem2D; elem*=2) { + HIPCHECK ( hipDeviceReset() ); + memcpyArraytest(elem+offset, elem+offset, 0, 1); // use pitch + } +} + +int main(int argc, char *argv[]) +{ + HipTest::parseStandardArguments(argc, argv, true); + + printf ("info: set device to %d\n", p_gpuDevice); + HIPCHECK(hipSetDevice(p_gpuDevice)); + + if (p_tests & 0x1) { + printf ("\n\n=== tests&1 (types)\n"); + printSep(); + HIPCHECK ( hipDeviceReset() ); + size_t width = N/6; + size_t height = N/6; + memcpy2Dtest(321, 211, 0); + memcpy2Dtest(322, 211, 0); + memcpy2Dtest(320, 211, 0); + memcpy2Dtest(323, 211, 0); + printf ("===\n\n\n"); + + printf ("\n\n=== tests&1 (types)\n"); + printSep(); + // 2D + memcpyArraytest(320, 211, 0, 0); + memcpyArraytest(322, 211, 0, 0); + memcpyArraytest(320, 211, 0, 0); + memcpyArraytest(320, 211, 0, 1); + memcpyArraytest(322, 211, 0, 1); + memcpyArraytest(320, 211, 0, 1); + printSep(); + // 1D + memcpyArraytest(320, 1, 0); + memcpyArraytest(322, 1, 0); + memcpyArraytest(320, 1, 0); + printf ("===\n\n\n"); + } + + if (p_tests & 0x4) { + printf ("\n\n=== tests&4 (test sizes and offsets)\n"); + printSep(); + HIPCHECK ( hipDeviceReset() ); + printSep(); + memcpyArraytest_size(0,0); + printSep(); + memcpyArraytest_size(0,64); + printSep(); + memcpyArraytest_size(1024*1024,13); + printSep(); + memcpyArraytest_size(1024*1024,50); + } + + passed(); + +} diff --git a/projects/clr/hipamd/tests/src/stress/hipStressAsync.cpp b/projects/clr/hipamd/tests/src/stress/hipStressAsync.cpp new file mode 100644 index 0000000000..3b8acb40a6 --- /dev/null +++ b/projects/clr/hipamd/tests/src/stress/hipStressAsync.cpp @@ -0,0 +1,64 @@ +/* +Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved. +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANNTY OF ANY KIND, EXPRESS OR +IMPLIED, INNCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANNY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +/* + * Test for checking the functionality of + * hipError_t hipDeviceSynchronize(); +*/ + +#include"hip_runtime.h" +#include + +#define _SIZE sizeof(int)*1024*1024 +#define NUM_STREAMS 20 +#define ITER 1<<10 + +__global__ void Iter(hipLaunchParm lp, int *Ad, int num){ + int tx = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x; + if(tx == 0){ + for(int i = 0; i +#include + +#define NUM_SIZE 8 +#define NUM_ITER 1 << 30 +static size_t size[NUM_SIZE]; + +__global__ void Add(hipLaunchParm lp, int *Ad){ + int tx = hipThreadIdx_x; + Ad[tx] = Ad[tx] + tx; +} + +void setup(){ + for(int i=0;i +#include + +#define NUM_SIZE 8 +#define NUM_ITER 1 << 30 +static size_t size[NUM_SIZE]; + +__global__ void Add(hipLaunchParm lp, int *Ad){ + int tx = hipThreadIdx_x; + Ad[tx] = Ad[tx] + tx; +} + +void setup(){ + for(int i=0;i + +#define _SIZE sizeof(int)*1024*1024 +#define NUM_STREAMS 20 +#define ITER 1<<10 + +__global__ void Iter(hipLaunchParm lp, int *Ad, int num){ + int tx = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x; + if(tx == 0){ + for(int i = 0; i -void initArrays(T **A_d, T **B_d, T **C_d, - T **A_h, T **B_h, T **C_h, - size_t N, bool usePinnedHost=false) +void initArraysForHost(T **A_h, T **B_h, T **C_h, + size_t N, bool usePinnedHost=false) { size_t Nbytes = N*sizeof(T); - if (A_d) { - HIPCHECK ( hipMalloc(A_d, Nbytes) ); - } - if (B_d) { - HIPCHECK ( hipMalloc(B_d, Nbytes) ); - } - if (C_d) { - HIPCHECK ( hipMalloc(C_d, Nbytes) ); - } - if (usePinnedHost) { if (A_h) { HIPCHECK ( hipHostMalloc((void**)A_h, Nbytes) ); @@ -173,7 +162,7 @@ void initArrays(T **A_d, T **B_d, T **C_d, *A_h = (T*)malloc(Nbytes); HIPASSERT(*A_h != NULL); } - + if (B_h) { *B_h = (T*)malloc(Nbytes); HIPASSERT(*B_h != NULL); @@ -185,7 +174,6 @@ void initArrays(T **A_d, T **B_d, T **C_d, } } - // Initialize the host data: for (size_t i=0; i -void freeArrays(T *A_d, T *B_d, T *C_d, - T *A_h, T *B_h, T *C_h, bool usePinnedHost) +void initArrays(T **A_d, T **B_d, T **C_d, + T **A_h, T **B_h, T **C_h, + size_t N, bool usePinnedHost=false) { + size_t Nbytes = N*sizeof(T); + if (A_d) { - HIPCHECK ( hipFree(A_d) ); + HIPCHECK ( hipMalloc(A_d, Nbytes) ); } if (B_d) { - HIPCHECK ( hipFree(B_d) ); + HIPCHECK ( hipMalloc(B_d, Nbytes) ); } if (C_d) { - HIPCHECK ( hipFree(C_d) ); + HIPCHECK ( hipMalloc(C_d, Nbytes) ); } + initArraysForHost(A_h, B_h, C_h, N, usePinnedHost); + +} + + +template +void freeArraysForHost(T *A_h, T *B_h, T *C_h, bool usePinnedHost) +{ if (usePinnedHost) { if (A_h) { HIPCHECK (hipHostFree(A_h)); @@ -231,9 +229,63 @@ void freeArrays(T *A_d, T *B_d, T *C_d, free (C_h); } } - + } +template +void freeArrays(T *A_d, T *B_d, T *C_d, + T *A_h, T *B_h, T *C_h, bool usePinnedHost) +{ + if (A_d) { + HIPCHECK ( hipFree(A_d) ); + } + if (B_d) { + HIPCHECK ( hipFree(B_d) ); + } + if (C_d) { + HIPCHECK ( hipFree(C_d) ); + } + + freeArraysForHost(A_h, B_h, C_h, usePinnedHost); +} + +#if defined(__HIP_PLATFORM_HCC__) +template +void initArrays2DPitch(T **A_d, T **B_d, T **C_d, + size_t *pitch_A, size_t *pitch_B, size_t *pitch_C, + size_t numW, size_t numH) +{ + if (A_d) { + HIPCHECK ( hipMallocPitch((void**)A_d, pitch_A, numW*sizeof(T), numH) ); + } + if (B_d) { + HIPCHECK ( hipMallocPitch((void**)B_d, pitch_B, numW*sizeof(T), numH) ); + } + if (C_d) { + HIPCHECK ( hipMallocPitch((void**)C_d, pitch_C, numW*sizeof(T), numH) ); + } + + HIPASSERT(*pitch_A == *pitch_B); + HIPASSERT(*pitch_A == *pitch_C) + +} + +inline void initHIPArrays(hipArray **A_d, hipArray **B_d, hipArray **C_d, + const hipChannelFormatDesc *desc, const size_t numW, const size_t numH, const unsigned int flags) +{ + + if (A_d) { + HIPCHECK( hipMallocArray(A_d, desc, numW, numH, flags)); + } + if (B_d) { + HIPCHECK( hipMallocArray(B_d, desc, numW, numH, flags)); + } + if (C_d) { + HIPCHECK( hipMallocArray(C_d, desc, numW, numH, flags)); + } + +} +#endif // Assumes C_h contains vector add of A_h + B_h // Calls the test "failed" macro if a mismatch is detected.