From f58a6b4dd5c12871c2435d77741b15bc57ad30ba Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Fri, 15 Jul 2016 14:18:59 +0530 Subject: [PATCH 01/14] README.md: Fix link to clang-hipify documentation Change-Id: I4c668c0e3c5ee2b16dfc213b48ccd1dbaac7b2a2 [ROCm/clr commit: a181262ae548829507b123d0d799c0c05b3ace4f] --- projects/clr/hipamd/README.md | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/projects/clr/hipamd/README.md b/projects/clr/hipamd/README.md index 5a2146e74d..38d4d147ed 100644 --- a/projects/clr/hipamd/README.md +++ b/projects/clr/hipamd/README.md @@ -12,7 +12,7 @@ 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) and [clang-hipify](clang-hipify/README.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) From 5586cc02b471174a77cbc1365973a442bd9e4e3e Mon Sep 17 00:00:00 2001 From: Ben Sander Date: Mon, 18 Jul 2016 13:56:13 -0500 Subject: [PATCH 02/14] Update documentation. - Add more detailed feature comparison HIP vs specific CUDA SDK revs. - Remove reference to old codenames. - Fix TOC in some md files. Change-Id: If0f6a05dc61237c4710c4d4bf2d933d556dcaeae [ROCm/clr commit: 4b3102b57820b94ef4dc9f8021f0d5fddb63c768] --- projects/clr/hipamd/README.md | 5 +- projects/clr/hipamd/RELEASE.md | 9 +- .../hipamd/docs/doxygen-input/mainpage.txt | 6 +- projects/clr/hipamd/docs/markdown/hip_faq.md | 90 ++++++++---- .../docs/markdown/hip_kernel_language.md | 29 ++-- .../hipamd/docs/markdown/hip_porting_guide.md | 135 ++++++------------ 6 files changed, 132 insertions(+), 142 deletions(-) diff --git a/projects/clr/hipamd/README.md b/projects/clr/hipamd/README.md index 38d4d147ed..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](clang-hipify/README.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/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..478de7cda6 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 @@ -629,6 +625,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 +635,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 From bb2296d48b537cd89d087450b782277024d4a302 Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Tue, 19 Jul 2016 06:52:38 +0530 Subject: [PATCH 03/14] Set cmake policy CMP0037 to old behavior for newer cmake versions Change-Id: Ib2a7da53a238a489e73d6c006c50f12f07f866a1 [ROCm/clr commit: 89cbfadf8699640c6895077a72e2c2b5673d6e9b] --- projects/clr/hipamd/CMakeLists.txt | 7 +++++++ 1 file changed, 7 insertions(+) diff --git a/projects/clr/hipamd/CMakeLists.txt b/projects/clr/hipamd/CMakeLists.txt index d99907d695..65a920f4ad 100644 --- a/projects/clr/hipamd/CMakeLists.txt +++ b/projects/clr/hipamd/CMakeLists.txt @@ -227,5 +227,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() From 1d3164ac80e9b7f2d874514ddcce4599f4c9e3a7 Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Tue, 19 Jul 2016 14:30:03 +0530 Subject: [PATCH 04/14] Add markdown documentation to hip_docs package Change-Id: I4075e6baaf287356ac8b485cccc231bfb729f078 [ROCm/clr commit: 538eeb7cc06fb4327e5e273a44b6a72c5561982c] --- .../hipamd/packaging/convert_md_to_html.sh | 61 +++++++++++++++++++ projects/clr/hipamd/packaging/hip_doc.txt | 5 +- 2 files changed, 65 insertions(+), 1 deletion(-) create mode 100755 projects/clr/hipamd/packaging/convert_md_to_html.sh 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 From 2098afc04d151620c459e02f5f5227403e3729b5 Mon Sep 17 00:00:00 2001 From: Aditya Atluri Date: Tue, 19 Jul 2016 13:51:44 -0500 Subject: [PATCH 05/14] added fix for signal overflow in kernels Change-Id: Ie0b1f97f69b7d7b34e445f6f120472819be03a0e [ROCm/clr commit: 77d71346191d1d55b70cd7bc18cc9fdff1d6fe55] --- projects/clr/hipamd/src/hip_hcc.cpp | 27 +++++++++++++++++++++++++++ 1 file changed, 27 insertions(+) diff --git a/projects/clr/hipamd/src/hip_hcc.cpp b/projects/clr/hipamd/src/hip_hcc.cpp index a80db8fd2e..c70c35f23b 100644 --- a/projects/clr/hipamd/src/hip_hcc.cpp +++ b/projects/clr/hipamd/src/hip_hcc.cpp @@ -1085,6 +1085,28 @@ 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. @@ -1114,6 +1136,7 @@ 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); @@ -1143,6 +1166,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); @@ -1173,6 +1197,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); @@ -1203,6 +1228,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 +1241,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); } From 794d6bfda829fbbbeb16e5b2be84d29873341e59 Mon Sep 17 00:00:00 2001 From: Aditya Atluri Date: Tue, 19 Jul 2016 13:57:56 -0500 Subject: [PATCH 06/14] added kernel launch stress test Change-Id: Ib50d47e55079839ba31aabb524b70b704d7a46fe [ROCm/clr commit: 300ec2af1e6b3e6dbd253b88b2fdfb71707efae4] --- .../tests/src/stress/hipStressKernel.cpp | 61 +++++++++++++++++++ 1 file changed, 61 insertions(+) create mode 100644 projects/clr/hipamd/tests/src/stress/hipStressKernel.cpp diff --git a/projects/clr/hipamd/tests/src/stress/hipStressKernel.cpp b/projects/clr/hipamd/tests/src/stress/hipStressKernel.cpp new file mode 100644 index 0000000000..a6888eb51f --- /dev/null +++ b/projects/clr/hipamd/tests/src/stress/hipStressKernel.cpp @@ -0,0 +1,61 @@ +/* +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 WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#include"hip_runtime.h" +#include +#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 Date: Wed, 20 Jul 2016 10:56:25 +0530 Subject: [PATCH 07/14] Bump HIP version to 0.92.00 Change-Id: I4d653213dcf3ddf7d09d36433afe31e21d17a7ee [ROCm/clr commit: 8c4cd599030c88417eb12cbb4c4a9acc9eefe206] --- projects/clr/hipamd/bin/hipconfig | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) 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; From 997cb27be8f44df9c32d457f659e53744a264c57 Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Thu, 21 Jul 2016 12:29:56 +0530 Subject: [PATCH 08/14] Merge branch 'hiparray' into amd-develop Change-Id: I63ca7b1db7b593ac5cfb3fd7cd5d08d6e4075a4c [ROCm/clr commit: 7d5cffdc177d52915e15975caaf5f764ad8820b6] --- projects/clr/hipamd/bin/hipify | 7 + .../include/hcc_detail/hip_runtime_api.h | 14 + .../hipamd/include/hcc_detail/hip_texture.h | 81 +++++- projects/clr/hipamd/src/hip_memory.cpp | 235 +++++++++++++++++ projects/clr/hipamd/tests/src/CMakeLists.txt | 5 + projects/clr/hipamd/tests/src/hipArray.cpp | 239 ++++++++++++++++++ projects/clr/hipamd/tests/src/test_common.h | 96 +++++-- 7 files changed, 646 insertions(+), 31 deletions(-) create mode 100644 projects/clr/hipamd/tests/src/hipArray.cpp 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/include/hcc_detail/hip_runtime_api.h b/projects/clr/hipamd/include/hcc_detail/hip_runtime_api.h index 3fd06aae38..daf737563e 100644 --- a/projects/clr/hipamd/include/hcc_detail/hip_runtime_api.h +++ b/projects/clr/hipamd/include/hcc_detail/hip_runtime_api.h @@ -760,6 +760,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/src/hip_memory.cpp b/projects/clr/hipamd/src/hip_memory.cpp index cabea6f759..95429b26e6 100644 --- a/projects/clr/hipamd/src/hip_memory.cpp +++ b/projects/clr/hipamd/src/hip_memory.cpp @@ -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) @@ -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 @@ -566,5 +778,28 @@ 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); +} 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/test_common.h b/projects/clr/hipamd/tests/src/test_common.h index 47e5f63a5b..cfde590f78 100644 --- a/projects/clr/hipamd/tests/src/test_common.h +++ b/projects/clr/hipamd/tests/src/test_common.h @@ -142,22 +142,11 @@ vectorADD(hipLaunchParm lp, template -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. From 68b1a199c4c187d6cb15faeb1b7bd0b7094e7a52 Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Thu, 21 Jul 2016 16:02:51 +0530 Subject: [PATCH 09/14] Fix using ATP markers Change-Id: If2d04f80b580237426c569737551e2001a8cd35a [ROCm/clr commit: dffed956fbb69e287f3f259f3d885a5e0e951ed1] --- projects/clr/hipamd/CMakeLists.txt | 18 +++++++++++-- projects/clr/hipamd/bin/hipcc | 25 +++++++++++-------- .../clr/hipamd/include/hcc_detail/hip_hcc.h | 2 +- projects/clr/hipamd/src/hip_hcc.cpp | 2 +- 4 files changed, 32 insertions(+), 15 deletions(-) diff --git a/projects/clr/hipamd/CMakeLists.txt b/projects/clr/hipamd/CMakeLists.txt index 65a920f4ad..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 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/include/hcc_detail/hip_hcc.h b/projects/clr/hipamd/include/hcc_detail/hip_hcc.h index 7c47d1b2a9..e79ba9e1dc 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: diff --git a/projects/clr/hipamd/src/hip_hcc.cpp b/projects/clr/hipamd/src/hip_hcc.cpp index c70c35f23b..6334b6004d 100644 --- a/projects/clr/hipamd/src/hip_hcc.cpp +++ b/projects/clr/hipamd/src/hip_hcc.cpp @@ -934,7 +934,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 From 21ae09859d9e57bf6dee33105bc1515772b676af Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Fri, 22 Jul 2016 15:45:32 +0530 Subject: [PATCH 10/14] hip_hcc package now depends on rocm-profiler Change-Id: I80d1c6048cc18c47c2024efb90368b17139a09ad [ROCm/clr commit: 5c1818774130bd6b705d76b27041a8b7a03a9080] --- projects/clr/hipamd/packaging/hip_hcc.txt | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) 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) From 64d0ffb40402a8b982c00880fb84f3d3a2df1603 Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Fri, 22 Jul 2016 15:46:55 +0530 Subject: [PATCH 11/14] Replace calls to ihipInit with use of HIP_INIT_API macro Change-Id: Iabf7df79f0238a8ddffea4607fe945df36642850 [ROCm/clr commit: b4854708197bad5d583f86dde58fc67af246031c] --- projects/clr/hipamd/include/hcc_detail/hip_hcc.h | 11 +++++++++++ .../hipamd/include/hcc_detail/hip_runtime_api.h | 1 - projects/clr/hipamd/src/hip_device.cpp | 12 ++++++------ projects/clr/hipamd/src/hip_event.cpp | 10 +++++----- projects/clr/hipamd/src/hip_hcc.cpp | 16 ++++++++-------- projects/clr/hipamd/src/hip_memory.cpp | 4 ++-- 6 files changed, 32 insertions(+), 22 deletions(-) diff --git a/projects/clr/hipamd/include/hcc_detail/hip_hcc.h b/projects/clr/hipamd/include/hcc_detail/hip_hcc.h index e79ba9e1dc..a3f9c73a13 100644 --- a/projects/clr/hipamd/include/hcc_detail/hip_hcc.h +++ b/projects/clr/hipamd/include/hcc_detail/hip_hcc.h @@ -466,6 +466,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: 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 daf737563e..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 * 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 6334b6004d..5bdbefcdfe 100644 --- a/projects/clr/hipamd/src/hip_hcc.cpp +++ b/projects/clr/hipamd/src/hip_hcc.cpp @@ -1113,7 +1113,7 @@ void decKernelCnt(){ // 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; @@ -1143,7 +1143,7 @@ hipStream_t ihipPreLaunchKernel(hipStream_t stream, dim3 grid, dim3 block, grid_ } 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; @@ -1174,7 +1174,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; @@ -1205,7 +1205,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; @@ -1253,8 +1253,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). //================================================================================================= @@ -1629,7 +1629,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; @@ -1649,7 +1649,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 95429b26e6..d458256fdf 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; From 2c945dba23c29d080ce544637db2f01e46dfc9d1 Mon Sep 17 00:00:00 2001 From: pensun Date: Fri, 22 Jul 2016 10:40:58 -0500 Subject: [PATCH 12/14] Add empty stubs for threadfence family routines, changes include: - stubs and documentation in include/hcc_details/hip_runtime.h - stubs with "no-op" in src/hip_memory.cpp - document update in hip_kernel_language.md, add suggestions to disable L1 and L2 caches when using the threadfence routines. Change-Id: Ic0753170f802003055bca9d7476d7f48817b98b7 [ROCm/clr commit: 6db08e51357f8ff8abe3a77950a2df0432613ee1] --- .../docs/markdown/hip_kernel_language.md | 3 + .../hipamd/include/hcc_detail/hip_runtime.h | 55 ++++++++++++++++++- projects/clr/hipamd/src/hip_memory.cpp | 22 ++++++-- 3 files changed, 74 insertions(+), 6 deletions(-) diff --git a/projects/clr/hipamd/docs/markdown/hip_kernel_language.md b/projects/clr/hipamd/docs/markdown/hip_kernel_language.md index 478de7cda6..ec2530e46e 100644 --- a/projects/clr/hipamd/docs/markdown/hip_kernel_language.md +++ b/projects/clr/hipamd/docs/markdown/hip_kernel_language.md @@ -230,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. 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/src/hip_memory.cpp b/projects/clr/hipamd/src/hip_memory.cpp index d458256fdf..94442f4698 100644 --- a/projects/clr/hipamd/src/hip_memory.cpp +++ b/projects/clr/hipamd/src/hip_memory.cpp @@ -295,7 +295,7 @@ hipError_t hipMallocArray(hipArray** array, const hipChannelFormatDesc* desc, 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; + hip_status = hipErrorMemoryAllocation; } } } @@ -443,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). @@ -703,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; } @@ -737,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); @@ -765,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); @@ -802,4 +802,16 @@ hipError_t hipFreeArray(hipArray* array) 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 +} From 776ea7a1f0b0c587b79822fbb89083a9f30e6096 Mon Sep 17 00:00:00 2001 From: Aditya Atluri Date: Sat, 23 Jul 2016 14:54:20 -0500 Subject: [PATCH 13/14] Partial fix async after kernel launch signal issue Change-Id: Ib48d6564379160035bded9493b93663fba361710 [ROCm/clr commit: 1704006bed0524899f02e9bbbb7334b629882182] --- .../clr/hipamd/include/hcc_detail/hip_hcc.h | 1 + projects/clr/hipamd/src/hip_hcc.cpp | 56 ++++++++++++++-- .../tests/src/stress/hipStressAsync.cpp | 64 +++++++++++++++++++ .../tests/src/stress/hipStressChain.cpp | 64 +++++++++++++++++++ .../hipamd/tests/src/stress/hipStressSync.cpp | 63 ++++++++++++++++++ 5 files changed, 243 insertions(+), 5 deletions(-) create mode 100644 projects/clr/hipamd/tests/src/stress/hipStressAsync.cpp create mode 100644 projects/clr/hipamd/tests/src/stress/hipStressChain.cpp create mode 100644 projects/clr/hipamd/tests/src/stress/hipStressSync.cpp diff --git a/projects/clr/hipamd/include/hcc_detail/hip_hcc.h b/projects/clr/hipamd/include/hcc_detail/hip_hcc.h index a3f9c73a13..10c6cf5ff0 100644 --- a/projects/clr/hipamd/include/hcc_detail/hip_hcc.h +++ b/projects/clr/hipamd/include/hcc_detail/hip_hcc.h @@ -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 diff --git a/projects/clr/hipamd/src/hip_hcc.cpp b/projects/clr/hipamd/src/hip_hcc.cpp index 5bdbefcdfe..097cd702e4 100644 --- a/projects/clr/hipamd/src/hip_hcc.cpp +++ b/projects/clr/hipamd/src/hip_hcc.cpp @@ -359,6 +359,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 +465,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", @@ -1364,7 +1401,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); } @@ -1446,7 +1482,10 @@ void ihipStream_t::copySync(LockedAccessor_StreamCrit_t &crit, void* dst, const 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(crit->_last_command_type == ihipCommandKernel){ + std::cout<<"Destroying depSignal MemcpySync"<wait(crit, true); @@ -1529,6 +1568,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); + } + } @@ -1602,6 +1645,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) { 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 + +#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 Date: Sat, 9 Jul 2016 19:29:55 +0530 Subject: [PATCH 14/14] D2H and H2D unpinned memory transfer support Change-Id: If6d6c970f435e5d917d5cc6cddc2ee2918cd1c37 Conflicts: src/hip_hcc.cpp [ROCm/clr commit: d11d65d401c4f68affea02783e92fc6309c39025] --- .../clr/hipamd/include/hcc_detail/hip_hcc.h | 2 +- projects/clr/hipamd/src/hip_hcc.cpp | 201 ++++++++++++++++-- projects/clr/hipamd/src/staging_buffer.cpp | 94 ++++++-- 3 files changed, 260 insertions(+), 37 deletions(-) diff --git a/projects/clr/hipamd/include/hcc_detail/hip_hcc.h b/projects/clr/hipamd/include/hcc_detail/hip_hcc.h index 10c6cf5ff0..635b65f384 100644 --- a/projects/clr/hipamd/include/hcc_detail/hip_hcc.h +++ b/projects/clr/hipamd/include/hcc_detail/hip_hcc.h @@ -606,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/src/hip_hcc.cpp b/projects/clr/hipamd/src/hip_hcc.cpp index 097cd702e4..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 */ @@ -616,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; @@ -757,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; @@ -994,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); @@ -1437,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 @@ -1481,10 +1645,21 @@ 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); 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).