diff --git a/projects/hip/.gitignore b/projects/hip/.gitignore index 67fa40f563..22cd23f2c6 100644 --- a/projects/hip/.gitignore +++ b/projects/hip/.gitignore @@ -1,4 +1,5 @@ .* +!.gitignore *.o *.exe *.swp diff --git a/projects/hip/.vimrc b/projects/hip/.vimrc index ed64acd347..019afa57e6 100644 --- a/projects/hip/.vimrc +++ b/projects/hip/.vimrc @@ -1,4 +1 @@ -:set tabstop=4 -:set shiftwidth=4 -:set expandtab -:set smartindent +:set makeprg=make\ -C\ build.hcc-LC.db diff --git a/projects/hip/CMakeLists.txt b/projects/hip/CMakeLists.txt index 9ea2e0f877..ccd390fbe5 100644 --- a/projects/hip/CMakeLists.txt +++ b/projects/hip/CMakeLists.txt @@ -142,6 +142,7 @@ if(NOT DEFINED COMPILE_HIP_ATP_MARKER) endif() add_to_config(_buildInfo COMPILE_HIP_ATP_MARKER) + ############################# # Build steps ############################# @@ -164,7 +165,7 @@ if(HIP_PLATFORM STREQUAL "hcc") set(HIP_HCC_BUILD_FLAGS "${HIP_HCC_BUILD_FLAGS} -DHIP_VERSION_MAJOR=${HIP_VERSION_MAJOR} -DHIP_VERSION_MINOR=${HIP_VERSION_MINOR} -DHIP_VERSION_PATCH=${HIP_VERSION_PATCH}") # Add remaining flags - set(HIP_HCC_BUILD_FLAGS "${HIP_HCC_BUILD_FLAGS} -fPIC -hc -I${HCC_HOME}/include -I${HSA_PATH}/include -I/opt/rocm/libhsakmt/include -stdlib=libc++") + set(HIP_HCC_BUILD_FLAGS "${HIP_HCC_BUILD_FLAGS} -fPIC -hc -I${HCC_HOME}/include -I${HSA_PATH}/include -I/opt/rocm/libhsakmt/include -I/usr/local/include/c++/v1 -stdlib=libc++") # Set compiler and compiler flags set(CMAKE_CXX_COMPILER "${HCC_HOME}/bin/hcc") @@ -337,7 +338,7 @@ endif() add_custom_target(install_for_test COMMAND "${CMAKE_COMMAND}" --build . --target install WORKING_DIRECTORY ${CMAKE_BINARY_DIR}) execute_process(COMMAND getconf _NPROCESSORS_ONLN OUTPUT_VARIABLE DASH_JAY OUTPUT_STRIP_TRAILING_WHITESPACE) -add_custom_target(test COMMAND ${CMAKE_COMMAND} . +add_custom_target(test COMMAND ${CMAKE_COMMAND} -DCMAKE_BUILD_TYPE=${CMAKE_BUILD_TYPE} . COMMAND make -j ${DASH_JAY} COMMAND make test WORKING_DIRECTORY ${BUILD_DIR} diff --git a/projects/hip/CONTRIBUTING.md b/projects/hip/CONTRIBUTING.md index d535ccac39..f6ed47acef 100644 --- a/projects/hip/CONTRIBUTING.md +++ b/projects/hip/CONTRIBUTING.md @@ -124,6 +124,11 @@ Differences or limitations of HIP APIs as compared to CUDA APIs should be clearl - ihipLogStatus should only be called from top-level HIP APIs,and should be called to log and return the error code. The error code is used by the GetLastError and PeekLastError functions - if a HIP API simply returns, then the error will not be logged correctly. +- All HIP environment variables should begin with the keyword HIP_ + Environment variables should be long enough to describe their purpose but short enough so they can be remembered - perhaps 10-20 characters, with 3-4 parts separated by underscores. + To see the list of current environment variables, along with their values, set HIP_PRINT_ENV and run any hip applications on ROCM platform . + HIPCC or other tools may support additional environment variables which should follow the above convention. + #### Presubmit Testing: diff --git a/projects/hip/RELEASE.md b/projects/hip/RELEASE.md index d1f79bc3c6..a7c770f611 100644 --- a/projects/hip/RELEASE.md +++ b/projects/hip/RELEASE.md @@ -13,8 +13,25 @@ Upcoming: ## Revision History: +=================================================================================================== +Release:1.0 +Date: 2016.11.8 +- Initial implementation for FindHIP.cmake +- HIP library now installs as a static library by default +- Added support for HIP context and HIP module APIs +- Major changes to HIP signal & memory management implementation +- Support for complex data type and math functions +- clang-hipify is now known as hipify-clang +- Added several new HIP samples +- Preliminary support for new APIs: hipMemcpyToSymbol, hipDeviceGetLimit, hipRuntimeGetVersion +- Added support for async memcpy driver API (for example hipMemcpyHtoDAsync) +- Support for memory management device functions: malloc, free, memcpy & memset +- Removed deprecated HIP runtime header locations. Please include "hip/hip_runtime.h" instead of "hip_runtime.h". You can use `find . -type f -exec sed -i 's:#include "hip_runtime.h":#include "hip/hip_runtime.h":g' {} +` to replace all such references + + =================================================================================================== Release:0.92.00 +Date: 2016.8.14 - hipLaunchKernel supports one-dimensional grid and/or block dims, without explicit cast to dim3 type (actually in 0.90.00) - fp16 software support - Support for Hawaii dGPUs using environment variable ROCM_TARGET=hawaii diff --git a/projects/hip/bin/hipcc b/projects/hip/bin/hipcc index 21453634d6..dee0894869 100755 --- a/projects/hip/bin/hipcc +++ b/projects/hip/bin/hipcc @@ -79,7 +79,7 @@ if ($HIP_PLATFORM eq "hcc") { $ROCM_PATH=$ENV{'ROCM_PATH'} // "/opt/rocm"; - $HIP_ATP_MARKER=$ENV{'HIP_ATP_MARKER'}; + $HIP_ATP_MARKER=$ENV{'HIP_ATP_MARKER'} // 1; $marker_path = "$ROCM_PATH/profiler/CXLActivityLogger"; $ROCM_TARGET=$ENV{'ROCM_TARGET'} // "fiji"; @@ -116,9 +116,16 @@ if ($HIP_PLATFORM eq "hcc") { } else { $HIPLDFLAGS .= " -Wl,--defsym=_binary_kernel_spir_end=1 -Wl,--defsym=_binary_kernel_spir_start=1 -Wl,--defsym=_binary_kernel_cl_start=1 -Wl,--defsym=_binary_kernel_cl_end=1"; } + if ($HOST_OSNAME eq "fedora") { + $HIPCXXFLAGS .= " -I/usr/local/include/c++/v1"; + } # Satisfy HCC dependencies - $HIPLDFLAGS .= " -lc++abi -lsupc++"; + if ($HOST_OSNAME eq "fedora") { + $HIPLDFLAGS .= " -lc++abi"; + } else { + $HIPLDFLAGS .= " -lc++abi -lsupc++"; + } $HIPLDFLAGS .= " -L$HSA_PATH/lib -L$ROCM_PATH/lib -lhsa-runtime64 -lhc_am -lhsakmt"; # Handle ROCm target platform @@ -273,10 +280,9 @@ foreach $arg (@ARGV) # Process HIPCC options here: if ($arg =~ m/^--hipcc/) { $swallowArg = 1; - if ($arg eq "--hipcc_explicit_lib") { - # Some environments (ie cmake tests) already link their own hip_hcc.o, so don't add here: - $needHipHcc = 0; - } + #if $arg eq "--hipcc_profile") { # Example argument here, hipcc + # + #} } else { push (@options, $arg); } @@ -314,14 +320,15 @@ if ($setStdLib eq 0 and $HIP_PLATFORM eq 'hcc') } if ($needHipHcc) { - $HIP_LIB_TYPE = $hipConfig{'HIP_LIB_TYPE'} // 0; + $HIP_LIB_TYPE = $hipConfig{'HIP_LIB_TYPE'} // 1; + # TODO - remove the old sea-of-objects solution: if ($HIP_LIB_TYPE eq 0) { - $HIPLDFLAGS .= " $HIP_PATH/lib/device_util.cpp.o $HIP_PATH/lib/hip_device.cpp.o $HIP_PATH/lib/hip_error.cpp.o $HIP_PATH/lib/hip_event.cpp.o $HIP_PATH/lib/hip_hcc.cpp.o $HIP_PATH/lib/hip_memory.cpp.o $HIP_PATH/lib/hip_peer.cpp.o $HIP_PATH/lib/hip_stream.cpp.o $HIP_PATH/lib/hip_ldg.cpp.o $HIP_PATH/lib/hip_fp16.cpp.o $HIP_PATH/lib/hip_context.cpp.o $HIP_PATH/lib/hip_module.cpp.o"; + substr($HIPLDFLAGS,0,0) = " $HIP_PATH/lib/device_util.cpp.o $HIP_PATH/lib/hip_device.cpp.o $HIP_PATH/lib/hip_error.cpp.o $HIP_PATH/lib/hip_event.cpp.o $HIP_PATH/lib/hip_hcc.cpp.o $HIP_PATH/lib/hip_memory.cpp.o $HIP_PATH/lib/hip_peer.cpp.o $HIP_PATH/lib/hip_stream.cpp.o $HIP_PATH/lib/hip_ldg.cpp.o $HIP_PATH/lib/hip_fp16.cpp.o $HIP_PATH/lib/hip_context.cpp.o $HIP_PATH/lib/hip_module.cpp.o "; } elsif ($HIP_LIB_TYPE eq 1) { - $HIPLDFLAGS .= " -L$HIP_PATH/lib -lhip_hcc" ; + substr($HIPLDFLAGS,0,0) = " -L$HIP_PATH/lib -lhip_hcc " ; } else { - $HIPLDFLAGS .= " -L$HIP_PATH/lib -Wl,--rpath=$HIP_PATH/lib -lhip_hcc"; + substr($HIPLDFLAGS,0,0) = " -L$HIP_PATH/lib -Wl,--rpath=$HIP_PATH/lib -lhip_hcc "; } } @@ -353,7 +360,7 @@ if ($printHipVersion) { } if ($runCmd) { if ($HIP_PLATFORM eq "hcc" and exists($hipConfig{'HCC_VERSION'}) and $HCC_VERSION ne $hipConfig{'HCC_VERSION'}) { - print ("HIP ($HIP_PATH) was built using hcc $hipConfig{'HCC_VERSION'}, but you are using hcc $HCC_VERSION. Please rebuild HIP including cmake.\n") && die (); + print ("HIP ($HIP_PATH) was built using hcc $hipConfig{'HCC_VERSION'}, but you are using $HCC_HOME/hcc with version $HCC_VERSION from hipcc. Please rebuild HIP including cmake or update HCC_HOME variable.\n") && die (); } system ("$CMD") and die (); } diff --git a/projects/hip/bin/hipdemangleatp b/projects/hip/bin/hipdemangleatp new file mode 100755 index 0000000000..456ea9ae15 --- /dev/null +++ b/projects/hip/bin/hipdemangleatp @@ -0,0 +1,19 @@ +#!/bin/bash + +# usage: hipdemangleatp.sh ATP_FILE + +# HIP kernels +kernels=$(grep grid_launch_parm $1 | cut -d" " -f1 | sort | uniq) +for mangled_sym in $kernels; do + real_sym=$(c++filt -p $(c++filt _$mangled_sym | cut -d: -f3 | sed 's/_functor//g' | sed 's/ /\\\ /g')) + #echo "$mangled_sym => $real_sym" >> $1.log + sed -i "s/$mangled_sym/$real_sym/g" $1 +done + +# HC kernels +kernels=$(grep cxxamp_trampoline $1 | cut -d" " -f1 | sort | uniq) +for mangled_sym in $kernels; do + real_sym=$(echo $mangled_sym | sed "s/^/_/g; s/_EC_/_$/g" | c++filt -p | cut -d\( -f1 | cut -d" " -f1 --complement | sed 's/ /\\\ /g') + #echo "$mangled_sym => $real_sym" >> $1.log + sed -i "s/$mangled_sym/$real_sym/g" $1 +done diff --git a/projects/hip/docs/markdown/hip_faq.md b/projects/hip/docs/markdown/hip_faq.md index f2fa3346cc..70ad94ba43 100644 --- a/projects/hip/docs/markdown/hip_faq.md +++ b/projects/hip/docs/markdown/hip_faq.md @@ -229,43 +229,7 @@ If platform portability is important, use #ifdef __HIP_PLATFORM_HIPCC__ to guard ### How do I trace HIP application flow? -#### Using CodeXL markers for HIP Functions -HIP can generate markers at function being/end which are displayed on the CodeXL timeline view. -To do this, you need to install ROCm-Profiler and enable HIP to generate the markers: - -1. Install ROCm-Profiler -Installing HIP from the [rocm](http://gpuopen.com/getting-started-with-boltzmann-components-platforms-installation/) pre-built packages, installs the ROCm-Profiler as well. -Alternatively, you can build ROCm-Profiler using the instructions [here](https://github.com/RadeonOpenCompute/ROCm-Profiler#building-the-rocm-profiler). - -2. Build HIP with ATP markers enabled -HIP pre-built packages are enabled with ATP marker support by default. -To enable ATP marker support when building HIP from source, use the option ```-DCOMPILE_HIP_ATP_MARKER=1``` during the cmake configure step. - -3. Set HIP_ATP_MARKER -```shell -export HIP_ATP_MARKER=1 -``` - -4. Recompile the target application - -5. Run with profiler enabled to generate ATP file. -```shell -# Use profile to generate timeline view: -/opt/rocm/bin/rocm-profiler -o -A - -Or -/opt/rocm/bin/rocm-profiler -e HIP_ATP_MARKER=1 -o -A -``` - -#### Using HIP_TRACE_API -You can also print the HIP function strings to stderr using HIP_TRACE_API environment variable. This can also be combined with the more detailed debug information provided -by the HIP_DB switch. For example: -```shell -# Trace to stderr showing being/end of each function (with arguments) + intermediate debug trace during the execution of each function. -HIP_TRACE_API=1 HIP_DB=0x2 ./myHipApp -``` - -Note this trace mode uses colors. "less -r" can handle raw control characters and will display the debug output in proper colors. +See the [HIP Profiling Guide](hip_porting_guide.md) for more information. ### What if HIP generates error of "symbol multiply defined!" only on AMD machine? Unlike CUDA, in HCC, for functions defined in the header files, the keyword of "__forceinline__" does not imply "static". diff --git a/projects/hip/docs/markdown/hip_kernel_language.md b/projects/hip/docs/markdown/hip_kernel_language.md index 3b13cde08a..f84868987c 100644 --- a/projects/hip/docs/markdown/hip_kernel_language.md +++ b/projects/hip/docs/markdown/hip_kernel_language.md @@ -233,8 +233,11 @@ typedef struct dim3 { ## Memory-Fence Instructions HIP supports __threadfence() and __threadfence_block(). -Applications that use threadfence_system can disable the L1 and L2 caches on the GPU by: -"export HSA_DISABLE_CACHE=1". See the hip_porting_guide.md#threadfence_system for more information. +HIP provides workaround for threadfence_system() under HCC path. +To enable the workaround, HIP should be built with environment variable HIP_COHERENT_HOST_ALLOC enabled. +In addition,the kernels that use __threadfence_system() should be modified as follows: +- The kernel should only operate on finegrained system memory; which should be allocated with hipHostMalloc(). +- Remove all memcpy for those allocated finegrained system memory regions. ## 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/hip/docs/markdown/hip_porting_guide.md b/projects/hip/docs/markdown/hip_porting_guide.md index 621726ee5f..c530df5098 100644 --- a/projects/hip/docs/markdown/hip_porting_guide.md +++ b/projects/hip/docs/markdown/hip_porting_guide.md @@ -564,7 +564,7 @@ HIP_LAUNCH_BLOCKING = 0 : Make HIP APIs 'host-synchronous', so they HIP_DB = 0 : Print various debug info. Bitmask, see hip_hcc.cpp for more information. HIP_TRACE_API = 0 : Trace each HIP API call. Print function name and return code to stderr as program executes. HIP_TRACE_API_COLOR = green : Color to use for HIP_API. None/Red/Green/Yellow/Blue/Magenta/Cyan/White -HIP_ATP_MARKER = 0 : Add HIP function begin/end to ATP file generated with CodeXL +HIP_PROFILE_API = 0 : Add HIP function begin/end to ATP file generated with CodeXL HIP_VISIBLE_DEVICES = 0 : Only devices whose index is present in the secquence are visible to HIP applications and they are enumerated in the order of secquence HIP_NUM_KERNELS_INFLIGHT = 128 : Number of kernels per stream diff --git a/projects/hip/docs/markdown/hip_profiling.md b/projects/hip/docs/markdown/hip_profiling.md new file mode 100644 index 0000000000..e4b88945e5 --- /dev/null +++ b/projects/hip/docs/markdown/hip_profiling.md @@ -0,0 +1,95 @@ +# Profiling HIP Code + +HIP provides several capabilities to support debugging and profiling. Profiling information can be displayed to stderr or viewed in the CodeXl visualization tool. + +### Usign CodeXL to profile a HIP Application +By defauly, CodeXL can trace all kernel commands, data transfer commands, and HSA Runtime (ROCr) API calls. +/opt/rocm/bin/rocm-profiler -o -A + +### Using CodeXL markers for HIP Functions +HIP can generate markers at function being/end which are displayed on the CodeXL timeline view. +HIP 1.0 compiles marker support by default, and you can enable it by setting the HIP_PROFILE_API environment variable and then running the rocm-profiler: + +```shell + +# Use profile to generate timeline view: +export HIP_PROFILE_API=1 +/opt/rocm/bin/rocm-profiler -o -A + +Or +/opt/rocm/bin/rocm-profiler -e HIP_PROFILE_API=1 -o -A +``` + +#### Developer Builds +For developer builds, you must enable marker support manually when compiling HIP. + +1. Build HIP with ATP markers enabled +HIP pre-built packages are enabled with ATP marker support by default. +To enable ATP marker support when building HIP from source, use the option ```-DCOMPILE_HIP_ATP_MARKER=1``` during the cmake configure step. + +2. Install ROCm-Profiler +Installing HIP from the [rocm](http://gpuopen.com/getting-started-with-boltzmann-components-platforms-installation/) pre-built packages, installs the ROCm-Profiler as well. +Alternatively, you can build ROCm-Profiler using the instructions [here](https://github.com/RadeonOpenCompute/ROCm-Profiler#building-the-rocm-profiler). + +3. Recompile the target application + +Then follow the steps above to collect a marker-enabled trace. + + +### Using HIP_TRACE_API +You can also print the HIP function strings to stderr using HIP_TRACE_API environment variable. This can also be combined with the more detailed debug information provided +by the HIP_DB switch. For example: +```shell +# Trace to stderr showing being/end of each function (with arguments) + intermediate debug trace during the execution of each function. +HIP_TRACE_API=1 HIP_DB=0x2 ./myHipApp +``` + +#### Color +Note this trace mode uses colors. "less -r" can handle raw control characters and will display the debug output in proper colors. +You can change the color used for the trace mode with the HIP_TRACE_API_COLOR environment variable. Possible values are None/Red/Green/Yellow/Blue/Magenta/Cyan/White. +None will disable use of color control codes and may be useful when saving the trace file or when a pure text trace is desired. + +#### + + +### Using HIP_DB + +This flag is primarily targeted to assist HIP development team in the development of the HIP runtime, but in some situations may be useful to HIP application developers as well. +The HIP debug information is designed to print important information during the execution of a HIP API. HIP provides +different color-coded levels of debug informaton: + - api : Print the beginning and end of each HIP API, including the arguments and return codes. + - sync : Print multi-thread and other synchronization debug information. + - copy : Print which engine is doing the copy, which copy flavor is selected, information on source and destination memory. + - mem : Print information about memory allocation - which pointers are allocated, where they are allocated, peer mappings, and more. + +DB_MEM format is flags separated by '+' sign, or a hex code for the bitmask. Generally the + format is preferred. +For example: +```shell +HIP_DB=api+copy+mem my-application +HIP_DB=0xF my-application +``` +HIP_DB=1 same as HIP_TRACE_API=1 + + + + +Trace provides quick look at API. +Explain output of +Reference the cookbook example. +Command-line profile. +/// disable profiling at the start of the application you can start CodeXLGpuProfiler with the --startdisabled flag. + +Can use strace interleaved with HSA Debug calls . + +HIP_PROFILE_API=1 +HIP_PROFILE_API=2 : Will show the full API in the trace. This can be useful for lower-level debugging when you want to see all the parameters that are passed to a specific API. + +demangle atp + +Write how to collect performance counters. +- include how to compute bandwidth for copy and kernel activity. + +- How to disable HSA APIs. +- Do I need to use profiler with HSA enabled? Do I need to enable HSA profiling on the command line? + +Offline compile, how to visualize. diff --git a/projects/hip/hipify-clang/src/Cuda2Hip.cpp b/projects/hip/hipify-clang/src/Cuda2Hip.cpp index 33e38b2c4b..87a69e8cb9 100644 --- a/projects/hip/hipify-clang/src/Cuda2Hip.cpp +++ b/projects/hip/hipify-clang/src/Cuda2Hip.cpp @@ -423,7 +423,11 @@ struct cuda2hipMap { cuda2hipRename["cuMemHostRegister_v2"] = {"hipHostRegister", CONV_MEM, API_DRIVER}; cuda2hipRename["cuMemHostUnregister"] = {"hipHostUnregister", CONV_MEM, API_DRIVER}; - + // Profiler + // unsupported yet by HIP + // cuda2hipRename["cuProfilerInitialize"] = {"hipProfilerInitialize", CONV_OTHER, API_DRIVER}; + cuda2hipRename["cuProfilerStart"] = {"hipProfilerStart", CONV_OTHER, API_DRIVER}; + cuda2hipRename["cuProfilerStop"] = {"hipProfilerStop", CONV_OTHER, API_DRIVER}; /////////////////////////////// CUDA RT API /////////////////////////////// // Error API @@ -1606,7 +1610,11 @@ private: } } XStr.clear(); - OS << "hipLaunchKernel(HIP_KERNEL_NAME(" << calleeName << "),"; + if (calleeName.find(',') != StringRef::npos) { + SmallString<128> tmpData; + calleeName = Twine("HIP_KERNEL_NAME(" + calleeName + ")").toStringRef(tmpData); + } + OS << "hipLaunchKernel(" << calleeName << ","; const CallExpr *config = launchKernel->getConfig(); DEBUG(dbgs() << "Kernel config arguments:" << "\n"); SourceManager *SM = Result.SourceManager; diff --git a/projects/hip/include/hip/hcc_detail/hip_runtime.h b/projects/hip/include/hip/hcc_detail/hip_runtime.h index 1bda07eb7d..b1edef18d7 100644 --- a/projects/hip/include/hip/hcc_detail/hip_runtime.h +++ b/projects/hip/include/hip/hcc_detail/hip_runtime.h @@ -556,10 +556,19 @@ extern "C" __device__ void __threadfence(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. + * @warning __threadfence_system is a stub and map to no-op. */ -__device__ void __threadfence_system(void) __attribute__((deprecated("Provided for compile-time compatibility, not yet functional"))); +__device__ void __threadfence_system(void) __attribute__((deprecated("Provided with workaround configuration, see hip_kernel_language.md for details"))); +__device__ unsigned __hip_ds_bpermute(int index, unsigned src); +__device__ float __hip_ds_bpermutef(int index, float src); +__device__ unsigned __hip_ds_permute(int index, unsigned src); +__device__ float __hip_ds_permutef(int index, float src); + +__device__ unsigned __hip_ds_swizzle(unsigned int src, int pattern); +__device__ float __hip_ds_swizzlef(float src, int pattern); + +__device__ int __hip_move_dpp(int src, int dpp_ctrl, int row_mask, int bank_mask, bool bound_ctrl); // doxygen end Fence Fence /** diff --git a/projects/hip/include/hip/hcc_detail/hip_runtime_api.h b/projects/hip/include/hip/hcc_detail/hip_runtime_api.h index 8c575eedc0..82eba49771 100644 --- a/projects/hip/include/hip/hcc_detail/hip_runtime_api.h +++ b/projects/hip/include/hip/hcc_detail/hip_runtime_api.h @@ -1278,7 +1278,11 @@ hipError_t hipDeviceEnablePeerAccess (int peerDeviceId, unsigned int flags); hipError_t hipDeviceDisablePeerAccess (int peerDeviceId); -#ifdef PEER_NON_UNIFIED +#ifndef USE_PEER_NON_UNIFIED +#define USE_PEER_NON_UNIFIED 1 +#endif + +#if USE_PEER_NON_UNIFIED==1 /** * @brief Copies memory from one device to memory on another device. * @@ -1591,6 +1595,16 @@ hipError_t hipDeviceGetName(char *name,int len,hipDevice_t device); */ hipError_t hipDeviceGetPCIBusId (int *pciBusId,int len,hipDevice_t device); +/** + * @brief Returns a handle to a compute device. + * @param [out] device handle + * @param [in] PCI Bus ID + * + * @returns #hipSuccess, #hipErrorInavlidDevice, #hipErrorInvalidValue + */ +hipError_t hipDeviceGetByPCIBusId ( int* device,const int* pciBusId ); + + /** * @brief Returns the total amount of memory on the device. * @param [out] bytes @@ -1737,11 +1751,24 @@ hipError_t hipModuleLaunchKernel(hipFunction_t f, * * @warning The cudaProfilerInitialize API format for "configFile" is not supported. * - * On AMD platforms, hipProfilerStart and hipProfilerStop require installation of AMD's GPU - * perf counter API and defining GPU_PERF */ +// TODO - expand descriptions: +/** + * @brief Start recording of profiling information + * @warning : hipProfilerStart API is under development. + */ +hipError_t hipProfilerStart(); + + +/** + * @brief Stop recording of profiling information. + * @warning : hipProfilerStop API is under development. + */ +hipError_t hipProfilerStop(); + + /** * @} */ diff --git a/projects/hip/include/hip/hip_profile.h b/projects/hip/include/hip/hip_profile.h new file mode 100644 index 0000000000..489143adfd --- /dev/null +++ b/projects/hip/include/hip/hip_profile.h @@ -0,0 +1,38 @@ +/* +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. +*/ + +#pragma once + +#if not defined (ENABLE_HIP_PROFILE) +#define ENABLE_HIP_PROFILE 1 +#endif + +#if defined(__HIP_PLATFORM_HCC__) and (ENABLE_HIP_PROFILE==1) +#include +#define HIP_SCOPED_MARKER(markerName, group) amdtScopedMarker __scopedMarker(markerName, group, nullptr); +#define HIP_BEGIN_MARKER(markerName, group) amdtBeginMarker(markerName, group, nullptr); +#define HIP_END_MARKER() amdtEndMarker(); +#else +#define HIP_SCOPED_MARKER(markerName, group) +#define HIP_BEGIN_MARKER(markerName, group) +#define HIP_END_MARKER() +#endif diff --git a/projects/hip/include/hip/nvcc_detail/hip_runtime_api.h b/projects/hip/include/hip/nvcc_detail/hip_runtime_api.h index a632e57f97..0d15dfcb01 100644 --- a/projects/hip/include/hip/nvcc_detail/hip_runtime_api.h +++ b/projects/hip/include/hip/nvcc_detail/hip_runtime_api.h @@ -24,6 +24,7 @@ THE SOFTWARE. #include #include +#include #ifdef __cplusplus extern "C" { @@ -83,6 +84,10 @@ typedef CUmodule hipModule_t; typedef CUfunction hipFunction_t; typedef CUdeviceptr hipDeviceptr_t; +// Flags that can be used with hipStreamCreateWithFlags +#define hipStreamDefault cudaStreamDefault +#define hipStreamNonBlocking cudaStreamNonBlocking + //typedef cudaChannelFormatDesc hipChannelFormatDesc; #define hipChannelFormatDesc cudaChannelFormatDesc @@ -106,7 +111,7 @@ switch(cuError) { case cudaErrorHostMemoryNotRegistered : return hipErrorHostMemoryNotRegistered ; case cudaErrorUnsupportedLimit : return hipErrorUnsupportedLimit ; default : return hipErrorUnknown; // Note - translated error. -}; +} } inline static hipError_t hipCUResultTohipError(CUresult cuError) { //TODO Populate further @@ -120,7 +125,7 @@ switch(cuError) { case CUDA_ERROR_INVALID_CONTEXT : return hipErrorInvalidContext ; case CUDA_ERROR_NOT_INITIALIZED : return hipErrorNotInitialized ; default : return hipErrorUnknown; // Note - translated error. -}; +} } // TODO match the error enum names of hip and cuda @@ -167,8 +172,8 @@ inline static cudaMemcpyKind hipMemcpyKindToCudaMemcpyKind(hipMemcpyKind kind) { /** * Stream CallBack struct */ -typedef void(* hipStreamCallback_t)(hipStream_t stream, hipError_t status, void* userData); - +#define HIPRT_CB CUDART_CB +typedef void(HIPRT_CB * hipStreamCallback_t)(hipStream_t stream, hipError_t status, void* userData); inline static hipError_t hipInit(unsigned int flags) { return hipCUResultTohipError(cuInit(flags)); @@ -319,11 +324,11 @@ inline static hipError_t hipDeviceSynchronize() { } inline static const char* hipGetErrorString(hipError_t error){ - return cudaGetErrorString( hipErrorToCudaError(error) ); + return cudaGetErrorString(hipErrorToCudaError(error)); } inline static const char* hipGetErrorName(hipError_t error){ - return cudaGetErrorName( hipErrorToCudaError(error) ); + return cudaGetErrorName(hipErrorToCudaError(error)); } inline static hipError_t hipGetDeviceCount(int * count){ @@ -585,8 +590,7 @@ inline static hipError_t hipStreamQuery(hipStream_t stream) inline static hipError_t hipStreamAddCallback(hipStream_t stream, hipStreamCallback_t callback, void *userData, unsigned int flags) { - return hipCUDAErrorTohipError(cudaStreamAddCallback(cudaStream_t stream, - cudaStreamCallback_t callback, void *userData, unsigned int flags)); + return hipCUDAErrorTohipError(cudaStreamAddCallback(stream, (cudaStreamCallback_t)callback, userData, flags)); } inline static hipError_t hipDriverGetVersion(int *driverVersion) @@ -611,12 +615,12 @@ inline static hipError_t hipDeviceCanAccessPeer ( int* canAccessPeer, int devic inline static hipError_t hipDeviceDisablePeerAccess ( int peerDevice ) { - return hipCUDAErrorTohipError(cudaDeviceDisablePeerAccess ( peerDevice )); -}; + return hipCUDAErrorTohipError(cudaDeviceDisablePeerAccess(peerDevice)); +} inline static hipError_t hipDeviceEnablePeerAccess ( int peerDevice, unsigned int flags ) { - return hipCUDAErrorTohipError(cudaDeviceEnablePeerAccess ( peerDevice, flags )); + return hipCUDAErrorTohipError(cudaDeviceEnablePeerAccess(peerDevice, flags)); } inline static hipError_t hipCtxDisablePeerAccess ( hipCtx_t peerCtx ) @@ -626,22 +630,33 @@ inline static hipError_t hipCtxDisablePeerAccess ( hipCtx_t peerCtx ) inline static hipError_t hipCtxEnablePeerAccess ( hipCtx_t peerCtx, unsigned int flags ) { - return hipCUResultTohipError(cuCtxEnablePeerAccess ( peerCtx, flags )); + return hipCUResultTohipError(cuCtxEnablePeerAccess(peerCtx, flags)); } inline static hipError_t hipMemcpyPeer ( void* dst, int dstDevice, const void* src, int srcDevice, size_t count ) { - return hipCUDAErrorTohipError(cudaMemcpyPeer ( dst, dstDevice, src, srcDevice, count )); -}; + return hipCUDAErrorTohipError(cudaMemcpyPeer(dst, dstDevice, src, srcDevice, count)); +} inline static hipError_t hipMemcpyPeerAsync ( void* dst, int dstDevice, const void* src, int srcDevice, size_t count, hipStream_t stream=0 ) { - return hipCUDAErrorTohipError(cudaMemcpyPeerAsync ( dst, dstDevice, src, srcDevice, count, stream )); -}; + return hipCUDAErrorTohipError(cudaMemcpyPeerAsync(dst, dstDevice, src, srcDevice, count, stream)); +} + +// Profile APIs: +inline hipError_t hipProfilerStart() +{ + return hipCUDAErrorTohipError(cudaProfilerStart()); +} + +inline hipError_t hipProfilerStop() +{ + return hipCUDAErrorTohipError(cudaProfilerStop()); +} inline static hipError_t hipSetDeviceFlags (unsigned int flags) { - return hipCUDAErrorTohipError(cudaSetDeviceFlags( flags )); + return hipCUDAErrorTohipError(cudaSetDeviceFlags(flags)); } inline static hipError_t hipEventCreateWithFlags(hipEvent_t* event, unsigned int flags) @@ -656,62 +671,62 @@ inline static hipError_t hipEventQuery(hipEvent_t event) inline static hipError_t hipCtxCreate(hipCtx_t *ctx, unsigned int flags, hipDevice_t device) { - return hipCUResultTohipError(cuCtxCreate ( ctx,flags,device )); + return hipCUResultTohipError(cuCtxCreate(ctx,flags,device)); } inline static hipError_t hipCtxDestroy(hipCtx_t ctx) { - return hipCUResultTohipError(cuCtxDestroy ( ctx )); + return hipCUResultTohipError(cuCtxDestroy(ctx)); } inline static hipError_t hipCtxPopCurrent(hipCtx_t* ctx) { - return hipCUResultTohipError(cuCtxPopCurrent ( ctx )); + return hipCUResultTohipError(cuCtxPopCurrent(ctx)); } inline static hipError_t hipCtxPushCurrent(hipCtx_t ctx) { - return hipCUResultTohipError(cuCtxPushCurrent ( ctx )); + return hipCUResultTohipError(cuCtxPushCurrent(ctx)); } inline static hipError_t hipCtxSetCurrent(hipCtx_t ctx) { - return hipCUResultTohipError(cuCtxSetCurrent ( ctx )); + return hipCUResultTohipError(cuCtxSetCurrent(ctx)); } inline static hipError_t hipCtxGetCurrent(hipCtx_t* ctx) { - return hipCUResultTohipError(cuCtxGetCurrent ( ctx )); + return hipCUResultTohipError(cuCtxGetCurrent(ctx)); } inline static hipError_t hipCtxGetDevice(hipDevice_t *device) { - return hipCUResultTohipError(cuCtxGetDevice ( device )); + return hipCUResultTohipError(cuCtxGetDevice(device)); } inline static hipError_t hipCtxGetApiVersion (hipCtx_t ctx,int *apiVersion) { - return hipCUResultTohipError(cuCtxGetApiVersion ( ctx,(unsigned int*)apiVersion )); + return hipCUResultTohipError(cuCtxGetApiVersion(ctx,(unsigned int*)apiVersion)); } -inline static hipError_t hipCtxGetCacheConfig ( hipFuncCache *cacheConfig ) +inline static hipError_t hipCtxGetCacheConfig (hipFuncCache *cacheConfig) { - return hipCUResultTohipError(cuCtxGetCacheConfig ( cacheConfig )); + return hipCUResultTohipError(cuCtxGetCacheConfig(cacheConfig)); } -inline static hipError_t hipCtxSetCacheConfig ( hipFuncCache cacheConfig ) +inline static hipError_t hipCtxSetCacheConfig (hipFuncCache cacheConfig) { - return hipCUResultTohipError(cuCtxSetCacheConfig ( cacheConfig )); + return hipCUResultTohipError(cuCtxSetCacheConfig(cacheConfig)); } -inline static hipError_t hipCtxSetSharedMemConfig ( hipSharedMemConfig config ) +inline static hipError_t hipCtxSetSharedMemConfig (hipSharedMemConfig config) { - return hipCUResultTohipError(cuCtxSetSharedMemConfig ( config )); + return hipCUResultTohipError(cuCtxSetSharedMemConfig(config)); } inline static hipError_t hipCtxGetSharedMemConfig ( hipSharedMemConfig * pConfig ) { - return hipCUResultTohipError(cuCtxGetSharedMemConfig ( pConfig )); + return hipCUResultTohipError(cuCtxGetSharedMemConfig(pConfig)); } inline static hipError_t hipCtxSynchronize ( void ) @@ -721,7 +736,7 @@ inline static hipError_t hipCtxSynchronize ( void ) inline static hipError_t hipCtxGetFlags ( unsigned int* flags ) { - return hipCUResultTohipError(cuCtxGetFlags ( flags )); + return hipCUResultTohipError(cuCtxGetFlags(flags)); } inline static hipError_t hipCtxDetach(hipCtx_t ctx) @@ -749,6 +764,11 @@ inline static hipError_t hipDeviceGetPCIBusId(int *pciBusId,int len,hipDevice_t return hipCUResultTohipError(cuDeviceGetPCIBusId((char*)pciBusId,len,device)); } +inline static hipError_t hipDeviceGetByPCIBusId(int* device, const int *pciBusId) +{ + return hipCUDAErrorTohipError(cudaDeviceGetByPCIBusId(device,(char*)pciBusId)); +} + inline static hipError_t hipDeviceGetLimit(size_t *pValue, hipLimit_t limit) { return hipCUDAErrorTohipError(cudaDeviceGetLimit(pValue, limit)); @@ -798,6 +818,8 @@ inline static hipError_t hipModuleLaunchKernel(hipFunction_t f, sharedMemBytes, stream, kernelParams, extra)); } + + #ifdef __cplusplus } #endif diff --git a/projects/hip/packaging/hip_hcc.txt b/projects/hip/packaging/hip_hcc.txt index 5801554f7c..d7b0877b62 100644 --- a/projects/hip/packaging/hip_hcc.txt +++ b/projects/hip/packaging/hip_hcc.txt @@ -9,6 +9,7 @@ else() install(FILES @PROJECT_BINARY_DIR@/libhip_hcc.so DESTINATION lib) endif() install(FILES @PROJECT_BINARY_DIR@/.hipInfo DESTINATION lib) +install(FILES @hip_SOURCE_DIR@/src/hip_ir.ll DESTINATION lib) ############################# # Packaging steps diff --git a/projects/hip/samples/0_Intro/hcc_dialects/.gitignore b/projects/hip/samples/0_Intro/hcc_dialects/.gitignore new file mode 100644 index 0000000000..bce1cdf193 --- /dev/null +++ b/projects/hip/samples/0_Intro/hcc_dialects/.gitignore @@ -0,0 +1,5 @@ +vadd_amp_arrayview +vadd_hc_am +vadd_hc_array +vadd_hc_arrayview +vadd_hip diff --git a/projects/hip/samples/0_Intro/hcc_dialects/Makefile b/projects/hip/samples/0_Intro/hcc_dialects/Makefile index 3b5ceca7f0..4a514b6691 100644 --- a/projects/hip/samples/0_Intro/hcc_dialects/Makefile +++ b/projects/hip/samples/0_Intro/hcc_dialects/Makefile @@ -5,8 +5,8 @@ OPT=-O2 HCC_CFLAGS= `$(HCC_HOME)/bin/hcc-config --cxxflags` ${OPT} HCC_LDFLAGS= `$(HCC_HOME)/bin/hcc-config --ldflags` ${OPT} -CPPAMP_CFLAGS= -std=c++amp -stdlib=libc++ -I$(HCC_HOME)/include -CPPAMP_LDFLAGS= -std=c++amp -L$(HCC_HOME)/lib -Wl,--rpath=$(HCC_HOME)/lib -lc++ -lc++abi -ldl -lpthread -Wl,--whole-archive -lmcwamp -Wl,--no-whole-archive +CPPAMP_CFLAGS= `$(HCC_HOME)/bin/clamp-config --cxxflags` +CPPAMP_LDFLAGS= `$(HCC_HOME)/bin/clamp-config --ldflags` HIP_PATH?= $(wildcard /opt/rocm/hip) ifeq (,$(HIP_PATH)) diff --git a/projects/hip/samples/0_Intro/module_api/.gitignore b/projects/hip/samples/0_Intro/module_api/.gitignore new file mode 100644 index 0000000000..c1d81e043f --- /dev/null +++ b/projects/hip/samples/0_Intro/module_api/.gitignore @@ -0,0 +1,5 @@ +runKernel.hip.out +vcpy_isa.code +vcpy_isa.hsaco +vcpy_kernel.co +vcpy_kernel.code diff --git a/projects/hip/samples/0_Intro/square/Makefile b/projects/hip/samples/0_Intro/square/Makefile index 89921c2072..1e8cdba080 100644 --- a/projects/hip/samples/0_Intro/square/Makefile +++ b/projects/hip/samples/0_Intro/square/Makefile @@ -1,7 +1,4 @@ HIP_PATH?= $(wildcard /opt/rocm/hip) -ifeq (,$(HIP_PATH)) - HIP_PATH=../../.. -endif HIPCC=$(HIP_PATH)/bin/hipcc all: square.hip.out @@ -11,9 +8,10 @@ square.cuda.out : square.cu #hipify square.cu > square.cpp # Then review & finish port in square.cpp +# square.hip.out: square.hipref.cpp - $(HIPCC) square.hipref.cpp -o $@ + $(HIPCC) $(CXXFLAGS) square.hipref.cpp -o $@ diff --git a/projects/hip/samples/1_Utils/hipInfo/hipInfo.cpp b/projects/hip/samples/1_Utils/hipInfo/hipInfo.cpp index 0403162bd1..42a879e732 100644 --- a/projects/hip/samples/1_Utils/hipInfo/hipInfo.cpp +++ b/projects/hip/samples/1_Utils/hipInfo/hipInfo.cpp @@ -133,6 +133,15 @@ void printDeviceProp (int deviceId) } } cout << endl; + cout << setw(w1) << "non-peers: "; + for (int i=0; i - -// hip header file -#include "hip/hip_runtime.h" - -#define WIDTH 1024 - -#define NUM (WIDTH*WIDTH) - -#define THREADS_PER_BLOCK_X 4 -#define THREADS_PER_BLOCK_Y 4 -#define THREADS_PER_BLOCK_Z 1 - -// Device (Kernel) function, it must be void -// hipLaunchParm provides the execution configuration -__global__ void matrixTranspose(hipLaunchParm lp, - float *out, - float *in, - const int width) -{ - int x = hipBlockDim_x * hipBlockIdx_x + hipThreadIdx_x; - int y = hipBlockDim_y * hipBlockIdx_y + hipThreadIdx_y; - - out[y * width + x] = in[x * width + y]; -} - -// CPU implementation of matrix transpose -void matrixTransposeCPUReference( - float * output, - float * input, - const unsigned int width) -{ - for(unsigned int j=0; j < width; j++) - { - for(unsigned int i=0; i < width; i++) - { - output[i*width + j] = input[j*width + i]; - } - } -} - -int main() { - - float* Matrix; - float* TransposeMatrix; - float* cpuTransposeMatrix; - - float* gpuMatrix; - float* gpuTransposeMatrix; - - hipDeviceProp_t devProp; - hipGetDeviceProperties(&devProp, 0); - - std::cout << "Device name " << devProp.name << std::endl; - - hipEvent_t start, stop; - hipEventCreate(&start); - hipEventCreate(&stop); - float eventMs = 1.0f; - - int i; - int errors; - - Matrix = (float*)malloc(NUM * sizeof(float)); - TransposeMatrix = (float*)malloc(NUM * sizeof(float)); - cpuTransposeMatrix = (float*)malloc(NUM * sizeof(float)); - - // initialize the input data - for (i = 0; i < NUM; i++) { - Matrix[i] = (float)i*10.0f; - } - - // allocate the memory on the device side - hipMalloc((void**)&gpuMatrix, NUM * sizeof(float)); - hipMalloc((void**)&gpuTransposeMatrix, NUM * sizeof(float)); - - // Record the start event - hipEventRecord(start, NULL); - - // Memory transfer from host to device - hipMemcpy(gpuMatrix, Matrix, NUM*sizeof(float), hipMemcpyHostToDevice); - - // Record the stop event - hipEventRecord(stop, NULL); - hipEventSynchronize(stop); - - hipEventElapsedTime(&eventMs, start, stop); - - printf ("hipMemcpyHostToDevice time taken = %6.3fms\n", eventMs); - - // Record the start event - hipEventRecord(start, NULL); - - // Lauching kernel from host - hipLaunchKernel(matrixTranspose, - dim3(WIDTH/THREADS_PER_BLOCK_X, WIDTH/THREADS_PER_BLOCK_Y), - dim3(THREADS_PER_BLOCK_X, THREADS_PER_BLOCK_Y), - 0, 0, - gpuTransposeMatrix , gpuMatrix, WIDTH); - - // Record the stop event - hipEventRecord(stop, NULL); - hipEventSynchronize(stop); - - hipEventElapsedTime(&eventMs, start, stop); - - printf ("kernel Execution time = %6.3fms\n", eventMs); - - // Record the start event - hipEventRecord(start, NULL); - - // Memory transfer from device to host - hipMemcpy(TransposeMatrix, gpuTransposeMatrix, NUM*sizeof(float), hipMemcpyDeviceToHost); - - // Record the stop event - hipEventRecord(stop, NULL); - hipEventSynchronize(stop); - - hipEventElapsedTime(&eventMs, start, stop); - - printf ("hipMemcpyDeviceToHost time taken = %6.3fms\n", eventMs); - - // CPU MatrixTranspose computation - matrixTransposeCPUReference(cpuTransposeMatrix, Matrix, WIDTH); - - // verify the results - errors = 0; - double eps = 1.0E-6; - for (i = 0; i < NUM; i++) { - if (std::abs(TransposeMatrix[i] - cpuTransposeMatrix[i]) > eps ) { - errors++; - } - } - if (errors!=0) { - printf("FAILED: %d errors\n",errors); - } else { - printf ("PASSED!\n"); - } - - //free the resources on device side - hipFree(gpuMatrix); - hipFree(gpuTransposeMatrix); - - //free the resources on host side - free(Matrix); - free(TransposeMatrix); - free(cpuTransposeMatrix); - - return errors; -} diff --git a/projects/hip/samples/2_Cookbook/2_Profiler/Makefile b/projects/hip/samples/2_Cookbook/2_Profiler/Makefile new file mode 100644 index 0000000000..db2d008182 --- /dev/null +++ b/projects/hip/samples/2_Cookbook/2_Profiler/Makefile @@ -0,0 +1,53 @@ +HIP_PATH?= $(wildcard /opt/rocm/hip) + +HIPCC=$(HIP_PATH)/bin/hipcc + + +HIPPROFILER=/opt/rocm/bin/rocm-profiler +PROFILER_OPT=-A -o MT.atp -e HIP_PROFILE_API=1 +HIPPROFILER_POST_CMD=$(HIP_PATH)/bin/hipdemangleatp MT.atp + +TARGET=hcc + +SOURCES = MatrixTranspose.cpp +OBJECTS = $(SOURCES:.cpp=.o) + +EXECUTABLE=./MatrixTranspose + +.PHONY: test + + +all: $(EXECUTABLE) profile + + + +OPT =-g +CXXFLAGS =$(OPT) +CXX=$(HIPCC) + + +$(EXECUTABLE): $(OBJECTS) + $(HIPCC) $(OBJECTS) -o $@ + + +profile: $(EXECUTABLE) + $(HIPPROFILER) $(PROFILER_OPT) $(EXECUTABLE) + $(HIPPROFILER_POST_CMD) + + +# Pass option to control start and stop iterations for profiling - see MatrixTranspose.cpp for implementation: +# Note we start profiler in --startdisabled mode - no timing collected until app enabled it via hipProfilerStart() +profile_trigger: $(EXECUTABLE) + $(HIPPROFILER) $(PROFILER_OPT) --startdisabled $(EXECUTABLE) 3 6 + $(HIPPROFILER_POST_CMD) + + +run: $(EXECUTABLE) + $(EXECUTABLE) + + +clean: + rm -f $(EXECUTABLE) + rm -f $(OBJECTS) + rm -f $(HIP_PATH)/src/*.o + diff --git a/projects/hip/samples/2_Cookbook/2_Profiler/MatrixTranspose.cpp b/projects/hip/samples/2_Cookbook/2_Profiler/MatrixTranspose.cpp new file mode 100644 index 0000000000..3747bb4ec5 --- /dev/null +++ b/projects/hip/samples/2_Cookbook/2_Profiler/MatrixTranspose.cpp @@ -0,0 +1,233 @@ +/* +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 header file +#include "hip/hip_runtime.h" +#include "hip/hip_profile.h" + +#define WIDTH 1024 + +#define NUM (WIDTH*WIDTH) + +#define THREADS_PER_BLOCK_X 4 +#define THREADS_PER_BLOCK_Y 4 +#define THREADS_PER_BLOCK_Z 1 + +#define ITERATIONS 10 + +// Cmdline parms to control start and stop triggers +int startTriggerIteration=-1; +int stopTriggerIteration=-1; + +// Device (Kernel) function, it must be void +// hipLaunchParm provides the execution configuration +__global__ void matrixTranspose(hipLaunchParm lp, + float *out, + float *in, + const int width) +{ + int x = hipBlockDim_x * hipBlockIdx_x + hipThreadIdx_x; + int y = hipBlockDim_y * hipBlockIdx_y + hipThreadIdx_y; + + out[y * width + x] = in[x * width + y]; +} + +// CPU implementation of matrix transpose +void matrixTransposeCPUReference( + float * output, + float * input, + const unsigned int width) +{ + for(unsigned int j=0; j < width; j++) + { + for(unsigned int i=0; i < width; i++) + { + output[i*width + j] = input[j*width + i]; + } + } +} + + +// Use a separate function to demonstrate how to use function name as part of scoped marker: +void runGPU(float *Matrix, float *TransposeMatrix, + float* gpuMatrix, float* gpuTransposeMatrix) { + + // __func__ is a standard C++ macro which expands to the name of the function, in this case "runGPU" + HIP_SCOPED_MARKER(__func__, "MyGroup"); + + for (int i=0; i= 2) { + startTriggerIteration = atoi(argv[1]); + printf ("info : will start tracing at iteration:%d\n", startTriggerIteration); + } + if (argc >= 3) { + stopTriggerIteration = atoi(argv[2]); + printf ("info : will stop tracing at iteration:%d\n", stopTriggerIteration); + } + + float* Matrix; + float* TransposeMatrix; + float* cpuTransposeMatrix; + + float* gpuMatrix; + float* gpuTransposeMatrix; + + hipDeviceProp_t devProp; + hipGetDeviceProperties(&devProp, 0); + + std::cout << "Device name " << devProp.name << std::endl; + + { + // Show example of how to create a "scoped marker". + // The scoped marker records the time spent inside the { scope } of the marker - the begin timestamp is at the + // beginning of the code scope, and the end is recorded when the SCOPE exits. This can be viewed in CodeXL + // timeline relative to other GPU and CPU events. + // This marker captures the time spent in setup including host allocation, initialization, and device memory allocation. + HIP_SCOPED_MARKER("Setup", "MyGroup"); + + + + Matrix = (float*)malloc(NUM * sizeof(float)); + TransposeMatrix = (float*)malloc(NUM * sizeof(float)); + cpuTransposeMatrix = (float*)malloc(NUM * sizeof(float)); + + // initialize the input data + for (int i = 0; i < NUM; i++) { + Matrix[i] = (float)i*10.0f; + } + + + // allocate the memory on the device side + hipMalloc((void**)&gpuMatrix, NUM * sizeof(float)); + hipMalloc((void**)&gpuTransposeMatrix, NUM * sizeof(float)); + + // FYI, the scoped-marker will be destroyed here when the scope exits, and will record its "end" timestamp. + } + + runGPU(Matrix, TransposeMatrix, gpuMatrix, gpuTransposeMatrix); + + + // show how to use explicit begin/end markers: + // We begin the timed region with HIP_BEGIN_MARKER, passing in the markerName and group: + // The region will stop when HIP_END_MARKER is called + // This is another way to mark begin/end - as an alternative to scoped markers. + HIP_BEGIN_MARKER("Check&TearDown", "MyGroup"); + + int errors = 0; + + // CPU MatrixTranspose computation + matrixTransposeCPUReference(cpuTransposeMatrix, Matrix, WIDTH); + + // verify the results + double eps = 1.0E-6; + for (int i = 0; i < NUM; i++) { + if (std::abs(TransposeMatrix[i] - cpuTransposeMatrix[i]) > eps ) { + errors++; + } + } + if (errors!=0) { + printf("FAILED: %d errors\n",errors); + } else { + printf ("PASSED!\n"); + } + + //free the resources on device side + hipFree(gpuMatrix); + hipFree(gpuTransposeMatrix); + + //free the resources on host side + free(Matrix); + free(TransposeMatrix); + free(cpuTransposeMatrix); + + // This ends the last marker started in this thread, in this case "Check&TearDown" + HIP_END_MARKER(); + + return errors; +} diff --git a/projects/hip/samples/2_Cookbook/2_HIP_ATP_MARKER/Readme.md b/projects/hip/samples/2_Cookbook/2_Profiler/Readme.md similarity index 83% rename from projects/hip/samples/2_Cookbook/2_HIP_ATP_MARKER/Readme.md rename to projects/hip/samples/2_Cookbook/2_Profiler/Readme.md index de1a800572..92a8be228e 100644 --- a/projects/hip/samples/2_Cookbook/2_HIP_ATP_MARKER/Readme.md +++ b/projects/hip/samples/2_Cookbook/2_Profiler/Readme.md @@ -1,6 +1,6 @@ ## Using hipEvents to measure performance ### -This tutorial is follow-up of the previous two tutorial where we learn how to write our first hip program, in which we compute Matrix Transpose and in second one, we added feature to measure time taken for memory transfer and kernel execution. In this tutorial, we won't make amy changes to the source code. We'll explain how to use the codexl/rocm-profiler for hip timeline tracing. +This tutorial is follow-up of the previous two tutorial where we learn how to write our first hip program, in which we compute Matrix Transpose and in second one, we added feature to measure time taken for memory transfer and kernel execution. In this tutorial, we'll explain how to use the codexl/rocm-profiler for hip timeline tracing. Also, we will augment the source code with additional markers so we can see the high-level application flow alongside the information that CodeXL automatically collects. ## Introduction: @@ -24,15 +24,11 @@ HIP can generate markers at function being/end which are displayed on the CodeXL 1. Install ROCm-Profiler Installing HIP from the rocm pre-built packages, installs the ROCm-Profiler as well. Alternatively, you can build ROCm-Profiler using the instructions given below. -2. Build HIP with ATP markers enabled HIP pre-built packages are enabled with ATP marker support by default. To enable ATP marker support when building HIP from source, use the option -DCOMPILE_HIP_ATP_MARKER=1 during the cmake configure step. -3. Set HIP_ATP_MARKER -`export HIP_ATP_MARKER=1` - -4. Recompile the target application - -5. Run with profiler enabled to generate ATP file. -`/opt/rocm/bin/rocm-profiler -o -A ` +2. Run with profiler enabled to generate ATP file. +(These steps are also captured in the Makefile) +The HIP_PROFILE_API enables display of the HIP APIs on the CodeXL trimeline view. +`/opt/rocm/bin/rocm-profiler -o -A -e HIP_PROFILE_API=1 ` ##Using HIP_TRACE_API diff --git a/projects/hip/samples/2_Cookbook/2_HIP_ATP_MARKER/Makefile b/projects/hip/samples/2_Cookbook/8_peer2peer/Makefile similarity index 82% rename from projects/hip/samples/2_Cookbook/2_HIP_ATP_MARKER/Makefile rename to projects/hip/samples/2_Cookbook/8_peer2peer/Makefile index d3630a1c19..a1dad7d1da 100644 --- a/projects/hip/samples/2_Cookbook/2_HIP_ATP_MARKER/Makefile +++ b/projects/hip/samples/2_Cookbook/8_peer2peer/Makefile @@ -1,36 +1,36 @@ -HIP_PATH?= $(wildcard /opt/rocm/hip) -ifeq (,$(HIP_PATH)) - HIP_PATH=../../.. -endif - -HIPCC=$(HIP_PATH)/bin/hipcc - -TARGET=hcc - -SOURCES = MatrixTranspose.cpp -OBJECTS = $(SOURCES:.cpp=.o) - -EXECUTABLE=./MatrixTranspose - -.PHONY: test - - -all: $(EXECUTABLE) test - -CXXFLAGS =-g -CXX=$(HIPCC) - - -$(EXECUTABLE): $(OBJECTS) - $(HIPCC) $(OBJECTS) -o $@ - - -test: $(EXECUTABLE) - $(EXECUTABLE) - - -clean: - rm -f $(EXECUTABLE) - rm -f $(OBJECTS) - rm -f $(HIP_PATH)/src/*.o - +HIP_PATH?= $(wildcard /opt/rocm/hip) +ifeq (,$(HIP_PATH)) + HIP_PATH=../../.. +endif + +HIPCC=$(HIP_PATH)/bin/hipcc + +TARGET=hcc + +SOURCES = peer2peer.cpp +OBJECTS = $(SOURCES:.cpp=.o) + +EXECUTABLE=./peer2peer + +.PHONY: test + + +all: $(EXECUTABLE) test + +CXXFLAGS =-g +CXX=$(HIPCC) + + +$(EXECUTABLE): $(OBJECTS) + $(HIPCC) $(OBJECTS) -o $@ + + +test: $(EXECUTABLE) + $(EXECUTABLE) + + +clean: + rm -f $(EXECUTABLE) + rm -f $(OBJECTS) + rm -f $(HIP_PATH)/src/*.o + diff --git a/projects/hip/samples/2_Cookbook/8_peer2peer/peer2peer.cpp b/projects/hip/samples/2_Cookbook/8_peer2peer/peer2peer.cpp new file mode 100644 index 0000000000..624de56cb0 --- /dev/null +++ b/projects/hip/samples/2_Cookbook/8_peer2peer/peer2peer.cpp @@ -0,0 +1,241 @@ +/* +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 WARRANUMTY OF ANY KIND, EXPRESS OR +IMPLIED, INUMCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNUMESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANUMY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER INUM AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR INUM CONUMECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#include +#include +#include +#define WIDTH 32 + +#define NUM (WIDTH*WIDTH) + +#define THREADS_PER_BLOCK_X 4 +#define THREADS_PER_BLOCK_Y 4 +#define THREADS_PER_BLOCK_Z 1 + +using namespace std; + +#define KNRM "\x1B[0m" +#define KRED "\x1B[31m" + +#define failed(...) \ + printf ("%serror: ", KRED);\ + printf (__VA_ARGS__);\ + printf ("\n");\ + printf ("error: TEST FAILED\n%s", KNRM );\ + abort(); + +#define HIPCHECK(error) \ +{\ + hipError_t localError = error; \ + if (localError != hipSuccess) { \ + printf("%serror: '%s'(%d) from %s at %s:%d%s\n", \ + KRED, hipGetErrorString(localError), localError,\ + #error,__FILE__, __LINE__, KNRM); \ + failed("API returned error code.");\ + }\ +} + +void checkPeer2PeerSupport() +{ + int gpuCount; + int canAccessPeer; + int p2pCapableDeviceCount=0; + + HIPCHECK(hipGetDeviceCount(&gpuCount)); + + if (gpuCount < 2) + printf("Peer2Peer application requires atleast 2 gpu devices"); + + for (int currentGpu=0; currentGpu eps ) { + printf("%d cpu: %f gpu peered data %f\n",i,randArray[i],TransposeMatrix[1][i]); + errors++; + } + } + if (errors!=0) { + printf("FAILED: %d errors\n",errors); + } else { + printf ("Peer2Peer PASSED!\n"); + } + + free(randArray); + for(int i=0;i<2;i++){ + hipFree(data[i]); + hipFree(gpuTransposeMatrix[i]); + free(TransposeMatrix[i]); + } + + HIPCHECK(hipSetDevice(peerGpu)); + HIPCHECK(hipDeviceReset()); + + HIPCHECK(hipSetDevice(currentGpu)); + HIPCHECK(hipDeviceReset()); + + return 0; +} diff --git a/projects/hip/src/hip_device.cpp b/projects/hip/src/hip_device.cpp index 9d577f5313..29ab0805b8 100644 --- a/projects/hip/src/hip_device.cpp +++ b/projects/hip/src/hip_device.cpp @@ -340,8 +340,28 @@ hipError_t hipDeviceTotalMem (size_t *bytes,hipDevice_t device) return ihipLogStatus(e); } +hipError_t hipDeviceGetByPCIBusId (int* device, const int* pciBusId ) +{ + HIP_INIT_API(device,pciBusId); + hipDeviceProp_t tempProp; + int deviceCount; + hipError_t e = hipErrorInvalidValue; + hipGetDeviceCount( &deviceCount ); + *device = 0; + for (int i=0; i< deviceCount; i++) { + hipGetDeviceProperties( &tempProp, i ); + if(tempProp.pciBusID == *pciBusId) { + *device =i; + e = hipSuccess; + break; + } + } + return ihipLogStatus(e); +} + hipError_t hipChooseDevice( int* device, const hipDeviceProp_t* prop ) { + HIP_INIT_API(device,prop); hipDeviceProp_t tempProp; int deviceCount; int inPropCount=0; diff --git a/projects/hip/src/hip_hcc.cpp b/projects/hip/src/hip_hcc.cpp index 7fa25334bc..e5619f723e 100644 --- a/projects/hip/src/hip_hcc.cpp +++ b/projects/hip/src/hip_hcc.cpp @@ -48,6 +48,9 @@ THE SOFTWARE. #include "trace_helper.h" +#ifndef USE_COPY_EXT_V2 +#define USE_COPY_EXT_V2 0 +#endif //================================================================================================= //Global variables: @@ -62,12 +65,27 @@ int HIP_LAUNCH_BLOCKING = 0; int HIP_PRINT_ENV = 0; int HIP_TRACE_API= 0; std::string HIP_TRACE_API_COLOR("green"); -int HIP_ATP_MARKER= 0; +int HIP_PROFILE_API= 0; + +// TODO - DB_START/STOP need more testing. +std::string HIP_DB_START_API; +std::string HIP_DB_STOP_API; int HIP_DB= 0; int HIP_VISIBLE_DEVICES = 0; /* Contains a comma-separated sequence of GPU identifiers */ int HIP_NUM_KERNELS_INFLIGHT = 128; int HIP_WAIT_MODE = 0; +int HIP_FORCE_P2P_HOST = 0; +int HIP_DENY_PEER_ACCESS = 0; + +// Force async copies to actually use the synchronous copy interface. +int HIP_FORCE_SYNC_COPY = 0; + + + + + + #define HIP_USE_PRODUCT_NAME 0 //#define DISABLE_COPY_EXT 1 @@ -84,6 +102,15 @@ std::vector g_hip_visible_devices; hsa_agent_t g_cpu_agent; unsigned g_numLogicalThreads; +std::atomic g_lastShortTid(1); + +// Indexed by short-tid: +// +std::vector g_dbStartTriggers; +std::vector g_dbStopTriggers; + + + /* Implementation of malloc and free device functions. @@ -163,22 +190,79 @@ __device__ void* __hip_hc_free(void *ptr) return nullptr; } +__device__ unsigned __hip_ds_bpermute(int index, unsigned src) { + return hc::__amdgcn_ds_bpermute(index, src); +} +__device__ float __hip_ds_bpermutef(int index, float src) { + return hc::__amdgcn_ds_bpermute(index, src); +} + +__device__ unsigned __hip_ds_permute(int index, unsigned src) { + return hc::__amdgcn_ds_permute(index, src); +} + +__device__ float __hip_ds_permutef(int index, float src) { + return hc::__amdgcn_ds_permute(index, src); +} + +__device__ unsigned __hip_ds_swizzle(unsigned int src, int pattern) { + return hc::__amdgcn_ds_swizzle(src, pattern); +} + +__device__ float __hip_ds_swizzlef(float src, int pattern) { + return hc::__amdgcn_ds_swizzle(src, pattern); +} + +__device__ int __hip_move_dpp(int src, int dpp_ctrl, int row_mask, int bank_mask, bool bound_ctrl) { + return hc::__amdgcn_move_dpp(src, dpp_ctrl, row_mask, bank_mask, bound_ctrl); +} //================================================================================================= // Thread-local storage: //================================================================================================= // This is the implicit context used by all HIP commands. // It can be set by hipSetDevice or by the CTX manipulation commands: - thread_local hipError_t tls_lastHipError = hipSuccess; +thread_local ShortTid tls_shortTid; + + //================================================================================================= // Top-level "free" functions: //================================================================================================= +void recordApiTrace(std::string *fullStr, const std::string &apiStr) +{ + auto apiSeqNum = tls_shortTid.incApiSeqNum(); + auto tid = tls_shortTid.tid(); + + if ((tid < g_dbStartTriggers.size()) && (apiSeqNum >= g_dbStartTriggers[tid].nextTrigger())) { + printf ("info: resume profiling at %lu\n", apiSeqNum); + RESUME_PROFILING; + g_dbStartTriggers.pop_back(); + }; + if ((tid < g_dbStopTriggers.size()) && (apiSeqNum >= g_dbStopTriggers[tid].nextTrigger())) { + printf ("info: stop profiling at %lu\n", apiSeqNum); + STOP_PROFILING; + g_dbStopTriggers.pop_back(); + }; + + fullStr->reserve(16 + apiStr.length()); + *fullStr = std::to_string(tid) + "."; + *fullStr += std::to_string(apiSeqNum); + *fullStr += " "; + *fullStr += apiStr; + + + if (COMPILE_HIP_DB && HIP_TRACE_API) { + fprintf (stderr, "%s<c_str(), API_COLOR_END); + } +} + + static inline bool ihipIsValidDevice(unsigned deviceIndex) { // deviceIndex is unsigned so always > 0 @@ -230,6 +314,23 @@ hipError_t ihipSynchronize(void) return (hipSuccess); } +//================================================================================================= +// ihipStream_t: +//================================================================================================= +ShortTid::ShortTid() : + _apiSeqNum(0) +{ + _shortTid = g_lastShortTid.fetch_add(1); + + if (HIP_DB & (1<::recomputePeerAgents() template<> -bool ihipCtxCriticalBase_t::isPeer(const ihipCtx_t *peer) +bool ihipCtxCriticalBase_t::isPeerWatcher(const ihipCtx_t *peer) { auto match = std::find(_peers.begin(), _peers.end(), peer); return (match != std::end(_peers)); @@ -459,12 +558,14 @@ bool ihipCtxCriticalBase_t::isPeer(const ihipCtx_t *peer) template<> -bool ihipCtxCriticalBase_t::addPeer(ihipCtx_t *peer) +bool ihipCtxCriticalBase_t::addPeerWatcher(const ihipCtx_t *thisCtx, ihipCtx_t *peerWatcher) { - auto match = std::find(_peers.begin(), _peers.end(), peer); + auto match = std::find(_peers.begin(), _peers.end(), peerWatcher); if (match == std::end(_peers)) { // Not already a peer, let's update the list: - _peers.push_back(peer); + tprintf(DB_COPY, "addPeerWatcher. Allocations on %s now visible to peerWatcher %s.\n", + thisCtx->toString().c_str(), peerWatcher->toString().c_str()); + _peers.push_back(peerWatcher); recomputePeerAgents(); return true; } @@ -475,12 +576,14 @@ bool ihipCtxCriticalBase_t::addPeer(ihipCtx_t *peer) template<> -bool ihipCtxCriticalBase_t::removePeer(ihipCtx_t *peer) +bool ihipCtxCriticalBase_t::removePeerWatcher(const ihipCtx_t *thisCtx, ihipCtx_t *peerWatcher) { - auto match = std::find(_peers.begin(), _peers.end(), peer); + auto match = std::find(_peers.begin(), _peers.end(), peerWatcher); if (match != std::end(_peers)) { // Found a valid peer, let's remove it. - _peers.remove(peer); + tprintf(DB_COPY, "removePeerWatcher. Allocations on %s no longer visible to former peerWatcher %s.\n", + thisCtx->toString().c_str(), peerWatcher->toString().c_str()); + _peers.remove(peerWatcher); recomputePeerAgents(); return true; } else { @@ -490,16 +593,17 @@ bool ihipCtxCriticalBase_t::removePeer(ihipCtx_t *peer) template<> -void ihipCtxCriticalBase_t::resetPeers(ihipCtx_t *thisDevice) +void ihipCtxCriticalBase_t::resetPeerWatchers(ihipCtx_t *thisCtx) { + tprintf(DB_COPY, "resetPeerWatchers for context=%s\n", thisCtx->toString().c_str()); _peers.clear(); _peerCnt = 0; - addPeer(thisDevice); // peer-list always contains self agent. + addPeerWatcher(thisCtx, thisCtx); // peer-list always contains self agent. } template<> -void ihipCtxCriticalBase_t::printPeers(FILE *f) const +void ihipCtxCriticalBase_t::printPeerWatchers(FILE *f) const { for (auto iter = _peers.begin(); iter!=_peers.end(); iter++) { fprintf (f, "%s ", (*iter)->toString().c_str()); @@ -904,7 +1008,7 @@ void ihipCtx_t::locked_reset() // Reset peer list to just me: - crit->resetPeers(this); + crit->resetPeerWatchers(this); // Reset and release all memory stored in the tracker: // Reset will remove peer mapping so don't need to do this explicitly. @@ -1006,43 +1110,14 @@ void ihipReadEnv_I(int *var_ptr, const char *var_name1, const char *var_name2, c env = getenv(var_name2); } - // TODO: Refactor this code so it is a separate call rather than being part of ihipReadEnv_I, which should only read integers. - // Check if the environment variable is either HIP_VISIBLE_DEVICES or CUDA_LAUNCH_BLOCKING, which - // contains a sequence of comma-separated device IDs - if (!(strcmp(var_name1,"HIP_VISIBLE_DEVICES") && strcmp(var_name2, "CUDA_VISIBLE_DEVICES")) && env){ - // Parse the string stream of env and store the device ids to g_hip_visible_devices global variable - std::string str = env; - std::istringstream ss(str); - std::string device_id; - // Clean up the defult value - g_hip_visible_devices.clear(); - g_visible_device = true; - // Read the visible device numbers - while (std::getline(ss, device_id, ',')) { - if (atoi(device_id.c_str()) >= 0) { - g_hip_visible_devices.push_back(atoi(device_id.c_str())); - } else { // Any device number after invalid number will not present - break; - } - } - // Print out the number of ids - if (HIP_PRINT_ENV) { - printf ("%-30s = ", var_name1); - for(int i=0;i(var_ptr) = env; } if (HIP_PRINT_ENV) { printf ("%-30s = %s : %s\n", var_name1, var_ptr->c_str(), description); @@ -1065,6 +1140,25 @@ void ihipReadEnv_S(std::string *var_ptr, const char *var_name1, const char *var_ } +void ihipReadEnv_Callback(void *var_ptr, const char *var_name1, const char *var_name2, const char *description, std::string (*setterCallback)(void * var_ptr, const char * env)) +{ + char * env = getenv(var_name1); + + // Check second name if first not defined, used to allow HIP_ or CUDA_ env vars. + if ((env == NULL) && strcmp(var_name2, "0")) { + env = getenv(var_name2); + } + + std::string var_string = "0"; + if (env) { + var_string = setterCallback(var_ptr, env); + } + if (HIP_PRINT_ENV) { + printf ("%-30s = %s : %s\n", var_name1, var_string.c_str(), description); + } +} + + #if defined (DEBUG) #define READ_ENV_I(_build, _ENV_VAR, _ENV_VAR2, _description) \ @@ -1075,6 +1169,10 @@ void ihipReadEnv_S(std::string *var_ptr, const char *var_name1, const char *var_ if ((_build == release) || (_build == debug) {\ ihipReadEnv_S(&_ENV_VAR, #_ENV_VAR, #_ENV_VAR2, _description);\ }; +#define READ_ENV_C(_build, _ENV_VAR, _ENV_VAR2, _description, _callback) \ + if ((_build == release) || (_build == debug) {\ + ihipReadEnv_Callback(&_ENV_VAR, #_ENV_VAR, #_ENV_VAR2, _description, _callback);\ + }; #else @@ -1087,10 +1185,155 @@ void ihipReadEnv_S(std::string *var_ptr, const char *var_name1, const char *var_ if (_build == release) {\ ihipReadEnv_S(&_ENV_VAR, #_ENV_VAR, #_ENV_VAR2, _description);\ }; +#define READ_ENV_C(_build, _ENV_VAR, _ENV_VAR2, _description, _callback) \ + if (_build == release) {\ + ihipReadEnv_Callback(&_ENV_VAR, #_ENV_VAR, #_ENV_VAR2, _description, _callback);\ + }; #endif +static void tokenize(const std::string &s, char delim, std::vector *tokens) +{ + std::stringstream ss; + ss.str(s); + std::string item; + while (getline(ss, item, delim)) { + item.erase (std::remove (item.begin(), item.end(), ' '), item.end()); // remove whitespace. + tokens->push_back(item); + } +} + +static void trim(std::string *s) +{ + // trim whitespace from beginning and end: + const char *t = "\t\n\r\f\v"; + s->erase(0, s->find_first_not_of(t)); + s->erase(s->find_last_not_of(t)+1); +} + +static void ltrim(std::string *s) +{ + // trim whitespace from beginning + const char *t = "\t\n\r\f\v"; + s->erase(0, s->find_first_not_of(t)); +} + + +// TODO - change last arg to pointer. +void parseTrigger(std::string triggerString, std::vector &profTriggers ) +{ + std::vector tidApiTokens; + tokenize(std::string(triggerString), ',', &tidApiTokens); + for (auto t=tidApiTokens.begin(); t != tidApiTokens.end(); t++) { + std::vector oneToken; + //std::cout << "token=" << *t << "\n"; + tokenize(std::string(*t), '.', &oneToken); + int tid = 1; + uint64_t apiTrigger = 0; + if (oneToken.size() == 1) { + // the case with just apiNum + apiTrigger = std::strtoull(oneToken[0].c_str(), nullptr, 0); + } else if (oneToken.size() == 2) { + // the case with tid.apiNum + tid = std::strtoul(oneToken[0].c_str(), nullptr, 0); + apiTrigger = std::strtoull(oneToken[1].c_str(), nullptr, 0); + } else { + throw ihipException(hipErrorRuntimeOther); // TODO -> bad env var? + } + + if (tid > 10000) { + throw ihipException(hipErrorRuntimeOther); // TODO -> bad env var? + } else { + profTriggers.resize(tid+1); + //std::cout << "tid:" << tid << " add: " << apiTrigger << "\n"; + profTriggers[tid].add(apiTrigger); + } + } + + + for (int tid=1; tid (var_ptr); + + std::string e(envVarString); + trim(&e); + if (!e.empty() && isdigit(e.c_str()[0])) { + long int v = strtol(envVarString, NULL, 0); + *var_ptr_int = (int) (v); + } else { + *var_ptr_int = 0; + std::vector tokens; + tokenize(e, '+', &tokens); + for (auto t=tokens.begin(); t!= tokens.end(); t++) { + for (int i=0; ic_str(), dbName[i]._shortName)) { + *var_ptr_int |= (1<= 0) { + g_hip_visible_devices.push_back(atoi(device_id.c_str())); + } else { // Any device number after invalid number will not present + break; + } + } + + std::string valueString; + // Print out the number of ids + for(int i=0;igrid_dim << " groupDim:" << lp->group_dim << " sharedMem:+" << lp->dynamic_group_mem_bytes << " " << *stream; + if (HIP_PROFILE_API == 0x1) { + std::string shortAtpString("hipLaunchKernel:"); + shortAtpString += kernelName; + MARKER_BEGIN(shortAtpString.c_str(), "HIP"); + } else if (HIP_PROFILE_API == 0x2) { + MARKER_BEGIN(os.str().c_str(), "HIP"); + } if (COMPILE_HIP_DB && HIP_TRACE_API) { std::cerr << API_COLOR << os.str() << API_COLOR_END << std::endl; } - SCOPED_MARKER(os.str().c_str(), "HIP", NULL); } } @@ -1379,7 +1627,10 @@ hipStream_t ihipPreLaunchKernel(hipStream_t stream, size_t grid, size_t block, g //This releases the lock on the stream. void ihipPostLaunchKernel(hipStream_t stream, grid_launch_parm &lp) { + tprintf(DB_SYNC, "ihipPostLaunchKernel, unlocking stream\n"); + stream->lockclose_postKernelCommand(lp.av); + MARKER_END(); } @@ -1486,29 +1737,44 @@ void ihipSetTs(hipEvent_t e) } -// Returns true if thisCtx can see the memory allocated on dstCtx and srcCtx. +// Returns true if copyEngineCtx can see the memory allocated on dstCtx and srcCtx. // The peer-list for a context controls which contexts have access to the memory allocated on that context. // So we check dstCtx's and srcCtx's peerList to see if the both include thisCtx. -bool ihipStream_t::canSeePeerMemory(const ihipCtx_t *thisCtx, ihipCtx_t *dstCtx, ihipCtx_t *srcCtx) +bool ihipStream_t::canSeeMemory(const ihipCtx_t *copyEngineCtx, const hc::AmPointerInfo *dstPtrInfo, const hc::AmPointerInfo *srcPtrInfo) { - tprintf (DB_COPY1, "Checking if direct copy can be used. thisCtx:%s; dstCtx:%s ; srcCtx:%s\n", - thisCtx->toString().c_str(), dstCtx->toString().c_str(), srcCtx->toString().c_str()); - // Use blocks to control scope of critical sections. - { - LockedAccessor_CtxCrit_t ctxCrit(dstCtx->criticalData()); - tprintf(DB_SYNC, "dstCrit lock succeeded\n"); - if (!ctxCrit->isPeer(thisCtx)) { - return false; - }; + // Make sure this is a device-to-device copy with all memory available to the requested copy engine + // + // TODO - pointer-info stores a deviceID not a context,may have some unusual side-effects here: + if (dstPtrInfo->_sizeBytes == 0) { + return false; + } else { + ihipCtx_t *dstCtx = ihipGetPrimaryCtx(dstPtrInfo->_appId); + if (copyEngineCtx != dstCtx) { + // Only checks peer list if contexts are different + LockedAccessor_CtxCrit_t ctxCrit(dstCtx->criticalData()); + //tprintf(DB_SYNC, "dstCrit lock succeeded\n"); + if (!ctxCrit->isPeerWatcher(copyEngineCtx)) { + return false; + }; + } } - { - LockedAccessor_CtxCrit_t ctxCrit(srcCtx->criticalData()); - tprintf(DB_SYNC, "srcCrit lock succeeded\n"); - if (!ctxCrit->isPeer(thisCtx)) { - return false; - }; + + + // TODO - pointer-info stores a deviceID not a context,may have some unusual side-effects here: + if (srcPtrInfo->_sizeBytes == 0) { + return false; + } else { + ihipCtx_t *srcCtx = ihipGetPrimaryCtx(srcPtrInfo->_appId); + if (copyEngineCtx != srcCtx) { + // Only checks peer list if contexts are different + LockedAccessor_CtxCrit_t ctxCrit(srcCtx->criticalData()); + //tprintf(DB_SYNC, "srcCrit lock succeeded\n"); + if (!ctxCrit->isPeerWatcher(copyEngineCtx)) { + return false; + }; + } } return true; @@ -1517,7 +1783,7 @@ bool ihipStream_t::canSeePeerMemory(const ihipCtx_t *thisCtx, ihipCtx_t *dstCtx, #define CASE_STRING(X) case X: return #X ;break; -const char* memcpyStr(unsigned memKind) +const char* hipMemcpyStr(unsigned memKind) { switch (memKind) { CASE_STRING(hipMemcpyHostToHost); @@ -1529,38 +1795,81 @@ const char* memcpyStr(unsigned memKind) }; } +const char* hcMemcpyStr(hc::hcCommandKind memKind) +{ + using namespace hc; + switch (memKind) { + CASE_STRING(hcMemcpyHostToHost); + CASE_STRING(hcMemcpyHostToDevice); + CASE_STRING(hcMemcpyDeviceToHost); + CASE_STRING(hcMemcpyDeviceToDevice); + //CASE_STRING(hcMemcpyDefault); + default : return ("unknown memcpyKind"); + }; +} + // Resolve hipMemcpyDefault to a known type. -// TODO - review why is this so complicated, does this need srcTracked and dstTracked? -unsigned ihipStream_t::resolveMemcpyDirection(bool srcTracked, bool dstTracked, bool srcInDeviceMem, bool dstInDeviceMem) +unsigned ihipStream_t::resolveMemcpyDirection(bool srcInDeviceMem, bool dstInDeviceMem) { hipMemcpyKind kind = hipMemcpyDefault; - if(!srcTracked && !dstTracked) - { - kind = hipMemcpyHostToHost; - } - if(!srcTracked && dstTracked) - { - if(dstInDeviceMem) { kind = hipMemcpyHostToDevice; } - else{ kind = hipMemcpyHostToHost; } - } - if (srcTracked && !dstTracked) { - if(srcInDeviceMem) { kind = hipMemcpyDeviceToHost; } - else { kind = hipMemcpyHostToHost; } - } - if (srcTracked && dstTracked) { - if(srcInDeviceMem && dstInDeviceMem) { kind = hipMemcpyDeviceToDevice; } - if(srcInDeviceMem && !dstInDeviceMem) { kind = hipMemcpyDeviceToHost; } - if(!srcInDeviceMem && !dstInDeviceMem) { kind = hipMemcpyHostToHost; } - if(!srcInDeviceMem && dstInDeviceMem) { kind = hipMemcpyHostToDevice; } - } + + if( srcInDeviceMem && dstInDeviceMem) { kind = hipMemcpyDeviceToDevice; } + if( srcInDeviceMem && !dstInDeviceMem) { kind = hipMemcpyDeviceToHost; } + if(!srcInDeviceMem && !dstInDeviceMem) { kind = hipMemcpyHostToHost; } + if(!srcInDeviceMem && dstInDeviceMem) { kind = hipMemcpyHostToDevice; } assert (kind != hipMemcpyDefault); return kind; } +// hipMemKind must be "resolved" to a specific direction - cannot be default. +void ihipStream_t::resolveHcMemcpyDirection(unsigned hipMemKind, + const hc::AmPointerInfo *dstPtrInfo, + const hc::AmPointerInfo *srcPtrInfo, + hc::hcCommandKind *hcCopyDir, + ihipCtx_t **copyDevice, + bool *forceUnpinnedCopy) +{ + // Ignore what the user tells us and always resolve the direction: + // Some apps apparently rely on this. + hipMemKind = resolveMemcpyDirection(srcPtrInfo->_isInDeviceMem, dstPtrInfo->_isInDeviceMem); + + + switch (hipMemKind) { + case hipMemcpyHostToHost: *hcCopyDir = hc::hcMemcpyHostToHost; break; + case hipMemcpyHostToDevice: *hcCopyDir = hc::hcMemcpyHostToDevice; break; + case hipMemcpyDeviceToHost: *hcCopyDir = hc::hcMemcpyDeviceToHost; break; + case hipMemcpyDeviceToDevice: *hcCopyDir = hc::hcMemcpyDeviceToDevice; break; + default: throw ihipException(hipErrorRuntimeOther); + }; + + if (srcPtrInfo->_isInDeviceMem) { + *copyDevice = ihipGetPrimaryCtx(srcPtrInfo->_appId); + } else if (dstPtrInfo->_isInDeviceMem) { + *copyDevice = ihipGetPrimaryCtx(dstPtrInfo->_appId); + } else { + *copyDevice = nullptr; + } + + *forceUnpinnedCopy = false; + if (canSeeMemory(*copyDevice, dstPtrInfo, srcPtrInfo)) { + + if (HIP_FORCE_P2P_HOST & 0x1) { + *forceUnpinnedCopy = true; + tprintf (DB_COPY, "P2P. Copy engine (dev:%d) can see src and dst but HIP_FORCE_P2P_HOST=0, forcing copy through staging buffers.\n", (*copyDevice)->getDeviceNum()); + } else { + tprintf (DB_COPY, "P2P. Copy engine (dev:%d) can see src and dst.\n", (*copyDevice)->getDeviceNum()); + } + } else { + *forceUnpinnedCopy = true; + tprintf (DB_COPY, "P2P: copy engine(dev:%d) cannot see both host and device pointers - forcing copy with unpinned engine.\n", (*copyDevice)->getDeviceNum()); + } +} + + // TODO - remove kind parm from here or use it below? void ihipStream_t::locked_copySync(void* dst, const void* src, size_t sizeBytes, unsigned kind, bool resolveOn) { @@ -1577,42 +1886,24 @@ void ihipStream_t::locked_copySync(void* dst, const void* src, size_t sizeBytes, bool dstTracked = (hc::am_memtracker_getinfo(&dstPtrInfo, dst) == AM_SUCCESS); bool srcTracked = (hc::am_memtracker_getinfo(&srcPtrInfo, src) == AM_SUCCESS); - if (kind == hipMemcpyDefault) { - kind = resolveMemcpyDirection(srcTracked, dstTracked, srcPtrInfo._isInDeviceMem, dstPtrInfo._isInDeviceMem); - } + hc::hcCommandKind hcCopyDir; - switch (kind) { - case hipMemcpyHostToHost: hcCopyDir = hc::hcMemcpyHostToHost; break; - case hipMemcpyHostToDevice: hcCopyDir = hc::hcMemcpyHostToDevice; break; - case hipMemcpyDeviceToHost: hcCopyDir = hc::hcMemcpyDeviceToHost; break; - case hipMemcpyDeviceToDevice: hcCopyDir = hc::hcMemcpyDeviceToDevice; break; - default: throw ihipException(hipErrorRuntimeOther); - }; - - - // If this is P2P access, we need to check to see if the copy agent (specified by the stream where the copy is enqueued) - // has peer access enabled to both the source and dest. If this is true, then the copy agent can see both pointers - // and we can perform the access with the copy engine from the current stream. If not true, then we will copy through the host. (forceHostCopyEngine=true). - bool forceHostCopyEngine = false; - if (hcCopyDir == hc::hcMemcpyDeviceToDevice) { - if (!canSeePeerMemory(ctx, ihipGetPrimaryCtx(dstPtrInfo._appId), ihipGetPrimaryCtx(srcPtrInfo._appId))) { - forceHostCopyEngine = true; - tprintf (DB_COPY1, "Forcing use of host copy engine.\n"); - } else { - tprintf (DB_COPY1, "Will use SDMA engine on streamDevice=%s.\n", ctx->toString().c_str()); - } - }; - - tprintf (DB_COPY1, "locked_copy dir=%s dst=%p src=%p sz=%zu\n", memcpyStr(kind), src, dst, sizeBytes); + ihipCtx_t *copyDevice; + bool forceUnpinnedCopy; + resolveHcMemcpyDirection(kind, &dstPtrInfo, &srcPtrInfo, &hcCopyDir, ©Device, &forceUnpinnedCopy); { LockedAccessor_StreamCrit_t crit (_criticalData); -#if DISABLE_COPY_EXT -#warning ("Disabled copy_ext path, P2P host staging copies will not work") - // Note - peer-to-peer copies which require host staging will not work in this path. - crit->_av.copy(src, dst, sizeBytes); + tprintf (DB_COPY, "copySync copyDev:%d dst=%p(home_dev:%d, tracked:%d, isDevMem:%d) src=%p(home_dev:%d, tracked:%d, isDevMem:%d) sz=%zu dir=%s forceUnpinnedCopy=%d\n", + copyDevice ? copyDevice->getDeviceNum():-1, + dst, dstPtrInfo._appId, dstTracked, dstPtrInfo._isInDeviceMem, + src, srcPtrInfo._appId, srcTracked, srcPtrInfo._isInDeviceMem, + sizeBytes, hcMemcpyStr(hcCopyDir), forceUnpinnedCopy); + +#if USE_COPY_EXT_V2 + crit->_av.copy_ext(src, dst, sizeBytes, hcCopyDir, srcPtrInfo, dstPtrInfo, copyDevice ? ©Device->getDevice()->_acc : nullptr, forceUnpinnedCopy); #else - crit->_av.copy_ext(src, dst, sizeBytes, hcCopyDir, srcPtrInfo, dstPtrInfo, forceHostCopyEngine); + crit->_av.copy_ext(src, dst, sizeBytes, hcCopyDir, srcPtrInfo, dstPtrInfo, forceUnpinnedCopy); #endif } } @@ -1624,12 +1915,12 @@ void ihipStream_t::locked_copyAsync(void* dst, const void* src, size_t sizeBytes const ihipCtx_t *ctx = this->getCtx(); if ((ctx == nullptr) || (ctx->getDevice() == nullptr)) { - tprintf (DB_COPY1, "locked_copyAsync bad ctx or device\n"); + tprintf (DB_COPY, "locked_copyAsync bad ctx or device\n"); throw ihipException(hipErrorInvalidDevice); } if (kind == hipMemcpyHostToHost) { - tprintf (DB_COPY1, "locked_copyAsync: H2H with memcpy"); + tprintf (DB_COPY, "locked_copyAsync: H2H with memcpy"); // TODO - consider if we want to perhaps use the GPU SDMA engines anyway, to avoid the host-side sync here and keep everything flowing on the GPU. /* As this is a CPU op, we need to wait until all @@ -1649,24 +1940,39 @@ void ihipStream_t::locked_copyAsync(void* dst, const void* src, size_t sizeBytes bool srcTracked = (hc::am_memtracker_getinfo(&srcPtrInfo, src) == AM_SUCCESS); - bool copyEngineCanSeeSrcAndDest = true; - if ((kind == hipMemcpyDeviceToDevice) || - ((kind == hipMemcpyDefault) && srcTracked && dstTracked)) { - copyEngineCanSeeSrcAndDest = canSeePeerMemory(ctx, ihipGetPrimaryCtx(dstPtrInfo._appId), ihipGetPrimaryCtx(srcPtrInfo._appId)); - } + hc::hcCommandKind hcCopyDir; + ihipCtx_t *copyDevice; + bool forceUnpinnedCopy; + resolveHcMemcpyDirection(kind, &dstPtrInfo, &srcPtrInfo, &hcCopyDir, ©Device, &forceUnpinnedCopy); - tprintf (DB_COPY1, "locked_copyAsync: async memcpy dstTracked=%d srcTracked=%d copyEngineCanSeeSrcAndDest=%d\n", - dstTracked, srcTracked, copyEngineCanSeeSrcAndDest); + tprintf (DB_COPY, "copyASync copyEngine_dev:%d dst=%p(home_dev:%d, tracked:%d, isDevMem:%d) src=%p(home_dev:%d, tracked:%d, isDevMem:%d) sz=%zu dir=%s. forceUnpinnedCopy=%d \n", + copyDevice->getDeviceNum(), + dst, dstPtrInfo._appId, dstTracked, dstPtrInfo._isInDeviceMem, + src, srcPtrInfo._appId, srcTracked, srcPtrInfo._isInDeviceMem, + sizeBytes, hcMemcpyStr(hcCopyDir), forceUnpinnedCopy); // "tracked" really indicates if the pointer's virtual address is available in the GPU address space. // If both pointers are not tracked, we need to fall back to a sync copy. - if (dstTracked && srcTracked && copyEngineCanSeeSrcAndDest) { + if (dstTracked && srcTracked && !forceUnpinnedCopy && copyDevice/*code below assumes this is !nullptr*/) { LockedAccessor_StreamCrit_t crit(_criticalData); - // Perform asynchronous copy: + // Perform fast asynchronous copy - we know copyDevice != NULL based on check above try { - crit->_av.copy_async(src, dst, sizeBytes); + if (HIP_FORCE_SYNC_COPY) { +#if USE_COPY_EXT_V2 + crit->_av.copy_ext (src, dst, sizeBytes, hcCopyDir, srcPtrInfo, dstPtrInfo, ©Device->getDevice()->_acc, forceUnpinnedCopy); +#else + crit->_av.copy_ext (src, dst, sizeBytes, hcCopyDir, srcPtrInfo, dstPtrInfo, forceUnpinnedCopy); +#endif + + } else { +#if USE_COPY_EXT_V2 + crit->_av.copy_async_ext(src, dst, sizeBytes, hcCopyDir, srcPtrInfo, dstPtrInfo, ©Device->getDevice()->_acc); +#else + crit->_av.copy_async(src, dst, sizeBytes); +#endif + } } catch (Kalmar::runtime_exception) { throw ihipException(hipErrorRuntimeOther); }; @@ -1678,12 +1984,42 @@ void ihipStream_t::locked_copyAsync(void* dst, const void* src, size_t sizeBytes } } else { - // TODO - call copy_ext directly here? - locked_copySync(dst, src, sizeBytes, kind); + LockedAccessor_StreamCrit_t crit(_criticalData); +#if USE_COPY_EXT_V2 + crit->_av.copy_ext(src, dst, sizeBytes, hcCopyDir, srcPtrInfo, dstPtrInfo, copyDevice ? ©Device->getDevice()->_acc : nullptr, forceUnpinnedCopy); +#else + crit->_av.copy_ext(src, dst, sizeBytes, hcCopyDir, srcPtrInfo, dstPtrInfo, forceUnpinnedCopy); +#endif } } } + +//------------------------------------------------------------------------------------------------- +//------------------------------------------------------------------------------------------------- +//Profiler, really these should live elsewhere: +hipError_t hipProfilerStart() +{ + HIP_INIT_API(); +#if COMPILE_HIP_ATP_MARKER + amdtResumeProfiling(AMDT_ALL_PROFILING); +#endif + + return ihipLogStatus(hipSuccess); +}; + + +hipError_t hipProfilerStop() +{ + HIP_INIT_API(); +#if COMPILE_HIP_ATP_MARKER + amdtStopProfiling(AMDT_ALL_PROFILING); +#endif + + return ihipLogStatus(hipSuccess); +}; + + //------------------------------------------------------------------------------------------------- //------------------------------------------------------------------------------------------------- // HCC-specific accessor functions: diff --git a/projects/hip/src/hip_hcc.h b/projects/hip/src/hip_hcc.h index bb06e23c67..0040263194 100644 --- a/projects/hip/src/hip_hcc.h +++ b/projects/hip/src/hip_hcc.h @@ -43,26 +43,70 @@ THE SOFTWARE. //static const int debug = 0; extern const int release; +// TODO - this blocks both kernels and memory ops. Perhaps should have separate env var for kernels? extern int HIP_LAUNCH_BLOCKING; extern int HIP_PRINT_ENV; -extern int HIP_ATP_MARKER; +extern int HIP_PROFILE_API; //extern int HIP_TRACE_API; extern int HIP_ATP; extern int HIP_DB; extern int HIP_STAGING_SIZE; /* size of staging buffers, in KB */ extern int HIP_STREAM_SIGNALS; /* number of signals to allocate at stream creation */ extern int HIP_VISIBLE_DEVICES; /* Contains a comma-separated sequence of GPU identifiers */ +extern int HIP_FORCE_P2P_HOST; //--- // Chicken bits for disabling functionality to work around potential issues: extern int HIP_DISABLE_HW_KERNEL_DEP; + +// Class to assign a short TID to each new thread, for HIP debugging purposes. +class ShortTid { +public: + + ShortTid() ; + + int tid() const { return _shortTid; }; + uint64_t incApiSeqNum() { return ++_apiSeqNum; }; + uint64_t apiSeqNum() const { return _apiSeqNum; }; + +private: + int _shortTid; + + // monotonically increasing API sequence number for this threa. + uint64_t _apiSeqNum; +}; + +struct ProfTrigger { + + static const uint64_t MAX_TRIGGER = std::numeric_limits::max(); + + void print (int tid) { + std::cout << "Enabling tracing for "; + for (auto iter=_profTrigger.begin(); iter != _profTrigger.end(); iter++) { + std::cout << "tid:" << tid << "." << *iter << ","; + } + std::cout << "\n"; + }; + + uint64_t nextTrigger() { return _profTrigger.empty() ? MAX_TRIGGER : _profTrigger.back(); }; + void add(uint64_t trigger) { _profTrigger.push_back(trigger); }; + void sort() { std::sort (_profTrigger.begin(), _profTrigger.end(), std::greater()); }; +private: + std::vector _profTrigger; +}; + + + //--- //Extern tls extern thread_local hipError_t tls_lastHipError; +extern thread_local ShortTid tls_shortTid; +extern std::vector g_dbStartTriggers; +extern std::vector g_dbStopTriggers; //--- //Forward defs: @@ -112,41 +156,42 @@ extern const char *API_COLOR_END; #endif -#define DB_SHOW_TID 0 - -#if DB_SHOW_TID -#define COMPUTE_TID_STR \ - std::stringstream tid_ss;\ - std::stringstream tid_ss_num;\ - tid_ss_num << std::this_thread::get_id();\ - tid_ss << " tid:" << std::hex << std::stoull(tid_ss_num.str()); -#else -#define COMPUTE_TID_STR std::stringstream tid_ss; +// Compile code that force hipHostMalloc only allocates finegrained system memory. +#ifndef HIP_COHERENT_HOST_ALLOC +#define HIP_COHERENT_HOST_ALLOC 0 #endif + // Compile support for trace markers that are displayed on CodeXL GUI at start/stop of each function boundary. // TODO - currently we print the trace message at the beginning. if we waited, we could also include return codes, and any values returned // through ptr-to-args (ie the pointers allocated by hipMalloc). #if COMPILE_HIP_ATP_MARKER #include "CXLActivityLogger.h" -#define SCOPED_MARKER(markerName,group,userString) amdtScopedMarker(markerName, group, userString) +#define MARKER_BEGIN(markerName,group) amdtBeginMarker(markerName, group, nullptr); +#define MARKER_END() amdtEndMarker(); +#define RESUME_PROFILING amdtResumeProfiling(AMDT_ALL_PROFILING); +#define STOP_PROFILING amdtStopProfiling(AMDT_ALL_PROFILING); #else // Swallow scoped markers: -#define SCOPED_MARKER(markerName,group,userString) +#define MARKER_BEGIN(markerName,group) +#define MARKER_END() +#define RESUME_PROFILING +#define STOP_PROFILING #endif +extern void recordApiTrace(std::string *fullStr, const std::string &apiStr); + #if COMPILE_HIP_ATP_MARKER || (COMPILE_HIP_TRACE_API & 0x1) #define API_TRACE(...)\ {\ - if (HIP_ATP_MARKER || (COMPILE_HIP_DB && HIP_TRACE_API)) {\ - std::string s = std::string(__func__) + " (" + ToString(__VA_ARGS__) + ')';\ - if (COMPILE_HIP_DB && HIP_TRACE_API) {\ - COMPUTE_TID_STR\ - fprintf (stderr, "%s<>%s\n", (localHipStatus == 0) ? API_COLOR:KRED, __func__, localHipStatus, ihipErrorString(localHipStatus), API_COLOR_END);\ + fprintf(stderr, " %ship-api tid:%d.%lu %-30s ret=%2d (%s)>>%s\n", (localHipStatus == 0) ? API_COLOR:KRED, tls_shortTid.tid(),tls_shortTid.apiSeqNum(), __func__, localHipStatus, ihipErrorString(localHipStatus), API_COLOR_END);\ }\ + if (HIP_PROFILE_API) { MARKER_END(); }\ localHipStatus;\ }) @@ -187,10 +233,9 @@ extern const char *API_COLOR_END; #define DB_API 0 /* 0x01 - shortcut to enable HIP_TRACE_API on single switch */ #define DB_SYNC 1 /* 0x02 - trace synchronization pieces */ #define DB_MEM 2 /* 0x04 - trace memory allocation / deallocation */ -#define DB_COPY1 3 /* 0x08 - trace memory copy commands. . */ +#define DB_COPY 3 /* 0x08 - trace memory copy and peer commands. . */ #define DB_SIGNAL 4 /* 0x10 - trace signal pool commands */ -#define DB_COPY2 5 /* 0x20 - trace memory copy commands. Detailed. */ -#define DB_MAX_BITPOS 5 +#define DB_MAX_FLAG 5 // When adding a new debug flag, also add to the char name table below. // @@ -204,20 +249,18 @@ static const DbName dbName [] = {KGRN, "api"}, // not used, {KYEL, "sync"}, {KCYN, "mem"}, - {KMAG, "copy1"}, + {KMAG, "copy"}, {KRED, "signal"}, - {KNRM, "copy2"}, }; - + #if COMPILE_HIP_DB #define tprintf(trace_level, ...) {\ if (HIP_DB & (1<<(trace_level))) {\ char msgStr[1000];\ snprintf(msgStr, 2000, __VA_ARGS__);\ - COMPUTE_TID_STR\ - fprintf (stderr, " %ship-%s%s:%s%s", dbName[trace_level]._color, dbName[trace_level]._shortName, tid_ss.str().c_str(), msgStr, KNRM); \ + fprintf (stderr, " %ship-%s tid:%d:%s%s", dbName[trace_level]._color, dbName[trace_level]._shortName, tls_shortTid.tid(), msgStr, KNRM); \ }\ } #else @@ -431,7 +474,7 @@ public: void launchModuleKernel(hc::accelerator_view av, hsa_signal_t signal, uint32_t blockDimX, uint32_t blockDimY, uint32_t blockDimZ, uint32_t gridDimX, uint32_t gridDimY, uint32_t gridDimZ, - uint32_t groupSegmentSize, uint32_t sharedMemBytes, + uint32_t groupSegmentSize, uint32_t sharedMemBytes, void *kernarg, size_t kernSize, uint64_t kernel); @@ -453,9 +496,14 @@ private: // The unsigned return is hipMemcpyKind - unsigned resolveMemcpyDirection(bool srcTracked, bool dstTracked, bool srcInDeviceMem, bool dstInDeviceMem); + unsigned resolveMemcpyDirection(bool srcInDeviceMem, bool dstInDeviceMem); + void resolveHcMemcpyDirection(unsigned hipMemKind, + const hc::AmPointerInfo *dstPtrInfo, const hc::AmPointerInfo *srcPtrInfo, + hc::hcCommandKind *hcCopyDir, + ihipCtx_t **copyDevice, + bool *forceUnpinnedCopy); - bool canSeePeerMemory(const ihipCtx_t *thisCtx, ihipCtx_t *dstCtx, ihipCtx_t *srcCtx); + bool canSeeMemory(const ihipCtx_t *thisCtx, const hc::AmPointerInfo *dstInfo, const hc::AmPointerInfo *srcInfo); private: // Data @@ -559,16 +607,18 @@ public: // Peer Accessor classes: - bool isPeer(const ihipCtx_t *peer); // returns True if peer has access to memory physically located on this device. - bool addPeer(ihipCtx_t *peer); - bool removePeer(ihipCtx_t *peer); - void resetPeers(ihipCtx_t *thisDevice); - void printPeers(FILE *f) const; + bool isPeerWatcher(const ihipCtx_t *peer); // returns True if peer has access to memory physically located on this device. + bool addPeerWatcher(const ihipCtx_t *thisCtx, ihipCtx_t *peer); + bool removePeerWatcher(const ihipCtx_t *thisCtx, ihipCtx_t *peer); + void resetPeerWatchers(ihipCtx_t *thisDevice); + void printPeerWatchers(FILE *f) const; uint32_t peerCnt() const { return _peerCnt; }; hsa_agent_t *peerAgents() const { return _peerAgents; }; + // TODO - move private + std::list _peers; // list of enabled peer devices. friend class LockedAccessor; private: @@ -580,7 +630,6 @@ private: // These reflect the currently Enabled set of peers for this GPU: // Enabled peers have permissions to access the memory physically allocated on this device. // Note the peers always contain the self agent for easy interfacing with HSA APIs. - std::list _peers; // list of enabled peer devices. uint32_t _peerCnt; // number of enabled peers hsa_agent_t *_peerAgents; // efficient packed array of enabled agents (to use for allocations.) private: @@ -619,11 +668,12 @@ public: // Functions: ihipCtxCritical_t &criticalData() { return _criticalData; }; // TODO, move private. Fix P2P. const ihipDevice_t *getDevice() const { return _device; }; + int getDeviceNum() const { return _device->_deviceId; }; // TODO - review uses of getWriteableDevice(), can these be converted to getDevice() ihipDevice_t *getWriteableDevice() const { return _device; }; - std::string toString() const; + std::string toString() const; public: // Data // The NULL stream is used if no other stream is specified. @@ -712,10 +762,16 @@ inline std::ostream& operator<<(std::ostream& os, const hipEvent_t& e) inline std::ostream& operator<<(std::ostream& os, const ihipCtx_t* c) { - os << "ctx:" << static_cast (c) - << " dev:" << c->getDevice()->_deviceId; + os << "ctx:" << static_cast (c) + << ".dev:" << c->getDevice()->_deviceId; return os; } +// Helper functions that are used across src files: +namespace hip_internal { + hipError_t memcpyAsync (void* dst, const void* src, size_t sizeBytes, hipMemcpyKind kind, hipStream_t stream); +}; + + #endif diff --git a/projects/hip/src/hip_memory.cpp b/projects/hip/src/hip_memory.cpp index 01f940a408..672b9f2ee2 100644 --- a/projects/hip/src/hip_memory.cpp +++ b/projects/hip/src/hip_memory.cpp @@ -119,14 +119,31 @@ hipError_t hipMalloc(void** ptr, size_t sizeBytes) 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->_deviceId, 0); + int peerCnt=0; { LockedAccessor_CtxCrit_t crit(ctx->criticalData()); - if (crit->peerCnt()) { - hsa_amd_agents_allow_access(crit->peerCnt(), crit->peerAgents(), NULL, *ptr); + // the peerCnt always stores self so make sure the trace actually + peerCnt = crit->peerCnt(); + tprintf(DB_MEM, " allocated device_mem ptr:%p size:%zu on dev:%d and allowed %d other peer(s) access\n", + *ptr, sizeBytes, device->_deviceId, peerCnt-1); + if (peerCnt > 1) { + + //printf ("peer self access\n"); + + // TODOD - remove me: + for (auto iter = crit->_peers.begin(); iter!=crit->_peers.end(); iter++) { + tprintf (DB_MEM, " allow access to peer: %s%s\n", (*iter)->toString().c_str(), (iter == crit->_peers.begin()) ? " (self)":""); + }; + + hsa_status_t e = hsa_amd_agents_allow_access(crit->peerCnt(), crit->peerAgents(), NULL, *ptr); + if (e != HSA_STATUS_SUCCESS) { + hip_status = hipErrorMemoryAllocation; + } } } } @@ -134,11 +151,25 @@ hipError_t hipMalloc(void** ptr, size_t sizeBytes) hip_status = hipErrorMemoryAllocation; } - //printf (" hipMalloc allocated %p\n", *ptr); return ihipLogStatus(hip_status); } +void ihipReadSingleEnv(int *var_ptr, const char *var_name1, const char *description) +{ + char * env = getenv(var_name1); + + // Default is set when variable is initialized (at top of this file), so only override if we find + // an environment variable. + if (env) { + long int v = strtol(env, NULL, 0); + *var_ptr = (int) (v); + } + if (HIP_PRINT_ENV) { + printf ("%-30s = %2d : %s\n", var_name1, *var_ptr, description); + } +} + hipError_t hipHostMalloc(void** ptr, size_t sizeBytes, unsigned int flags) { HIP_INIT_API(ptr, sizeBytes, flags); @@ -147,31 +178,58 @@ hipError_t hipHostMalloc(void** ptr, size_t sizeBytes, unsigned int flags) auto ctx = ihipGetTlsDefaultCtx(); - if(ctx){ - // am_alloc requires writeable __acc, perhaps could be refactored? - auto device = ctx->getWriteableDevice(); - if((flags == hipHostMallocDefault)|| (flags == hipHostMallocPortable)){ - *ptr = hc::am_alloc(sizeBytes, device->_acc, amHostPinned); - if(sizeBytes < 1 && (*ptr == NULL)){ - hip_status = hipErrorMemoryAllocation; - } else { - hc::am_memtracker_update(*ptr, device->_deviceId, amHostPinned); - } - tprintf(DB_MEM, " %s: pinned ptr=%p\n", __func__, *ptr); - } else if(flags & hipHostMallocMapped){ - *ptr = hc::am_alloc(sizeBytes, device->_acc, amHostPinned); - if(sizeBytes && (*ptr == NULL)){ - hip_status = hipErrorMemoryAllocation; - }else{ - hc::am_memtracker_update(*ptr, device->_deviceId, flags); - { - LockedAccessor_CtxCrit_t crit(ctx->criticalData()); - if (crit->peerCnt()) { - hsa_amd_agents_allow_access(crit->peerCnt(), crit->peerAgents(), NULL, *ptr); + if (sizeBytes == 0) { + hip_status = hipSuccess; + // TODO - should size of 0 return err or be siliently ignored? + } else if ((ctx==nullptr) || (ptr == nullptr)) { + hip_status = hipErrorInvalidValue; + } else { + unsigned trueFlags = flags; + if (flags == hipHostMallocDefault) { + trueFlags = hipHostMallocMapped | hipHostMallocWriteCombined; + } + + const unsigned supportedFlags = hipHostMallocPortable | hipHostMallocMapped | hipHostMallocWriteCombined; + + // Read from environment variable of HIP_COHERENT_HOST_ALLOC + int coherent_alloc=0; + ihipReadSingleEnv(&coherent_alloc, "HIP_COHERENT_HOST_ALLOC", "Flag to force allocate finegrained system memory"); + + if (flags & ~supportedFlags) { + hip_status = hipErrorInvalidValue; + } + else { + auto device = ctx->getWriteableDevice(); + if(coherent_alloc){ + // Force to allocate finedgrained system memory + *ptr = hc::am_alloc(sizeBytes, device->_acc, amHostPinned); + if(sizeBytes < 1 && (*ptr == NULL)){ + hip_status = hipErrorMemoryAllocation; + } else { + hc::am_memtracker_update(*ptr, device->_deviceId, amHostCoherent); + } + tprintf(DB_MEM, " %s: finegrained system memory ptr=%p\n", __func__, *ptr); + } + else{ + // TODO - am_alloc requires writeable __acc, perhaps could be refactored? + // TODO - hipHostMallocMapped is be ignored on ROCM - all memory is mapped to host address space as WC. + *ptr = hc::am_alloc(sizeBytes, device->_acc, amHostPinned); + if (*ptr == NULL) { + hip_status = hipErrorMemoryAllocation; + } else { + hc::am_memtracker_update(*ptr, device->_deviceId, flags); + // TODO-hipHostMallocPortable should map the host memory into all contexts, regardless of peer status. + int peerCnt=0; + { + LockedAccessor_CtxCrit_t crit(ctx->criticalData()); + peerCnt = crit->peerCnt(); + if (peerCnt > 1) { + hsa_amd_agents_allow_access(crit->peerCnt(), crit->peerAgents(), NULL, *ptr); + } } + tprintf(DB_MEM, "allocated pinned_host ptr:%p size:%zu on dev:%d and allow access to %d other peer(s)\n", *ptr, sizeBytes, device->_deviceId, peerCnt-1); } } - tprintf(DB_MEM, " %s: pinned ptr=%p\n", __func__, *ptr); } } return ihipLogStatus(hip_status); @@ -354,6 +412,8 @@ hipError_t hipHostRegister(void *hostPtr, size_t sizeBytes, unsigned int flags) vecAcc.push_back(ihipGetDevice(i)->_acc); } am_status = hc::am_memory_host_lock(device->_acc, hostPtr, sizeBytes, &vecAcc[0], vecAcc.size()); + + tprintf(DB_MEM, " %s registered ptr=%p\n", __func__, hostPtr); if(am_status == AM_SUCCESS){ hip_status = hipSuccess; } else { @@ -377,6 +437,7 @@ hipError_t hipHostUnregister(void *hostPtr) }else{ auto device = ctx->getWriteableDevice(); am_status_t am_status = hc::am_memory_host_unlock(device->_acc, hostPtr); + tprintf(DB_MEM, " %s unregistered ptr=%p\n", __func__, hostPtr); if(am_status != AM_SUCCESS){ hip_status = hipErrorHostMemoryNotRegistered; } @@ -398,6 +459,7 @@ hipError_t hipMemcpyToSymbol(const char* symbolName, const void *src, size_t cou hc::accelerator acc = ctx->getDevice()->_acc; void *ptr = acc.get_symbol_address(symbolName); + tprintf(DB_MEM, " symbol '%s' resolved to address:%p\n", symbolName, ptr); if(ptr == nullptr) { @@ -427,6 +489,7 @@ hipError_t hipMemcpyToSymbolAsync(const char* symbolName, const void *src, size_ hc::accelerator acc = ctx->getDevice()->_acc; void *ptr = acc.get_symbol_address(symbolName); + tprintf(DB_MEM, " symbol '%s' resolved to address:%p\n", symbolName, ptr); if(ptr == nullptr) { @@ -555,10 +618,13 @@ hipError_t hipMemcpyHtoH(void* dst, void* src, size_t sizeBytes) return ihipLogStatus(e); } -hipError_t hipMemcpyAsync(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind kind, hipStream_t stream) -{ - HIP_INIT_API(dst, src, sizeBytes, kind, stream); + +// Internal copy sync: +namespace hip_internal { + +hipError_t memcpyAsync (void* dst, const void* src, size_t sizeBytes, hipMemcpyKind kind, hipStream_t stream) +{ hipError_t e = hipSuccess; stream = ihipSyncAndResolveStream(stream); @@ -577,86 +643,39 @@ hipError_t hipMemcpyAsync(void* dst, const void* src, size_t sizeBytes, hipMemcp e = hipErrorInvalidValue; } - return ihipLogStatus(e); + return e; } +} // end namespace hip_internal + + +hipError_t hipMemcpyAsync(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind kind, hipStream_t stream) +{ + HIP_INIT_API(dst, src, sizeBytes, kind, stream); + + return ihipLogStatus(hip_internal::memcpyAsync(dst, src, sizeBytes, kind, stream)); + +} + hipError_t hipMemcpyHtoDAsync(hipDeviceptr_t dst, void* src, size_t sizeBytes, hipStream_t stream) { HIP_INIT_API(dst, src, sizeBytes, stream); - hipError_t e = hipSuccess; - - stream = ihipSyncAndResolveStream(stream); - - hipMemcpyKind kind = hipMemcpyHostToDevice; - - if ((dst == NULL) || (src == NULL)) { - e= hipErrorInvalidValue; - } else if (stream) { - try { - stream->locked_copyAsync((void*)dst, src, sizeBytes, kind); - } - catch (ihipException ex) { - e = ex._code; - } - } else { - e = hipErrorInvalidValue; - } - - return ihipLogStatus(e); + return ihipLogStatus(hip_internal::memcpyAsync(dst, src, sizeBytes, hipMemcpyHostToDevice, stream)); } hipError_t hipMemcpyDtoDAsync(hipDeviceptr_t dst, hipDeviceptr_t src, size_t sizeBytes, hipStream_t stream) { HIP_INIT_API(dst, src, sizeBytes, stream); - hipError_t e = hipSuccess; - - hipMemcpyKind kind = hipMemcpyDeviceToDevice; - - stream = ihipSyncAndResolveStream(stream); - - - if ((dst == NULL) || (src == NULL)) { - e= hipErrorInvalidValue; - } else if (stream) { - try { - stream->locked_copyAsync((void*)dst, (void*)src, sizeBytes, kind); - } - catch (ihipException ex) { - e = ex._code; - } - } else { - e = hipErrorInvalidValue; - } - - return ihipLogStatus(e); + return ihipLogStatus(hip_internal::memcpyAsync(dst, src, sizeBytes, hipMemcpyDeviceToDevice, stream)); } hipError_t hipMemcpyDtoHAsync(void* dst, hipDeviceptr_t src, size_t sizeBytes, hipStream_t stream) { HIP_INIT_API(dst, src, sizeBytes, stream); - hipError_t e = hipSuccess; - - stream = ihipSyncAndResolveStream(stream); - - hipMemcpyKind kind = hipMemcpyDeviceToHost; - - if ((dst == NULL) || (src == NULL)) { - e= hipErrorInvalidValue; - } else if (stream) { - try { - stream->locked_copyAsync(dst, (void*)src, sizeBytes, kind); - } - catch (ihipException ex) { - e = ex._code; - } - } else { - e = hipErrorInvalidValue; - } - - return ihipLogStatus(e); + return ihipLogStatus(hip_internal::memcpyAsync(dst, src, sizeBytes, hipMemcpyDeviceToHost, stream)); } // TODO - review and optimize diff --git a/projects/hip/src/hip_peer.cpp b/projects/hip/src/hip_peer.cpp index e66a0d2971..b7dca06e5f 100644 --- a/projects/hip/src/hip_peer.cpp +++ b/projects/hip/src/hip_peer.cpp @@ -35,22 +35,27 @@ THE SOFTWARE. // public APIs are thin wrappers which call into this internal implementations. // TODO - actually not yet - currently the integer deviceId flavors just call the context APIs. need to fix. -/** - * HCC returns 0 in *canAccessPeer ; Need to update this function when RT supports P2P - */ -//--- -hipError_t hipDeviceCanAccessPeer (int* canAccessPeer, hipCtx_t thisCtx, hipCtx_t peerCtx) -{ - HIP_INIT_API(canAccessPeer, thisCtx, peerCtx); + +hipError_t ihipDeviceCanAccessPeer (int* canAccessPeer, hipCtx_t thisCtx, hipCtx_t peerCtx) +{ hipError_t err = hipSuccess; if ((thisCtx != NULL) && (peerCtx != NULL)) { + if (thisCtx == peerCtx) { *canAccessPeer = 0; + tprintf(DB_MEM, "Can't be peer to self. (this=%s, peer=%s)\n", + thisCtx->toString().c_str(), peerCtx->toString().c_str()); + } else if (HIP_FORCE_P2P_HOST & 0x2) { + *canAccessPeer = false; + tprintf(DB_MEM, "HIP_FORCE_P2P_HOST denies peer access this=%s peer=%s canAccessPeer=%d\n", + thisCtx->toString().c_str(), peerCtx->toString().c_str(), *canAccessPeer); } else { - *canAccessPeer = peerCtx->getDevice()->_acc.get_is_peer(thisCtx->getDevice()->_acc); + *canAccessPeer = peerCtx->getDevice()->_acc.get_is_peer(thisCtx->getDevice()->_acc); + tprintf(DB_MEM, "deviceCanAccessPeer this=%s peer=%s canAccessPeer=%d\n", + thisCtx->toString().c_str(), peerCtx->toString().c_str(), *canAccessPeer); } } else { @@ -59,7 +64,19 @@ hipError_t hipDeviceCanAccessPeer (int* canAccessPeer, hipCtx_t thisCtx, hipCtx_ } - return ihipLogStatus(err); + return err; +} + + +/** + * HCC returns 0 in *canAccessPeer ; Need to update this function when RT supports P2P + */ +//--- +hipError_t hipDeviceCanAccessPeer (int* canAccessPeer, hipCtx_t thisCtx, hipCtx_t peerCtx) +{ + HIP_INIT_API(canAccessPeer, thisCtx, peerCtx); + + return ihipLogStatus(ihipDeviceCanAccessPeer(canAccessPeer, thisCtx, peerCtx)); } @@ -80,8 +97,10 @@ hipError_t ihipDisablePeerAccess (hipCtx_t peerCtx) err = hipErrorInvalidDevice; // Can't disable peer access to self. } else { LockedAccessor_CtxCrit_t peerCrit(peerCtx->criticalData()); - bool changed = peerCrit->removePeer(thisCtx); + bool changed = peerCrit->removePeerWatcher(peerCtx, thisCtx); if (changed) { + tprintf(DB_MEM, "device %s disable access to memory allocated on peer:%s\n", + thisCtx->toString().c_str(), peerCtx->toString().c_str()); // Update the peers for all memory already saved in the tracker: am_memtracker_update_peers(peerCtx->getDevice()->_acc, peerCrit->peerCnt(), peerCrit->peerAgents()); } else { @@ -112,8 +131,10 @@ hipError_t ihipEnablePeerAccess (hipCtx_t peerCtx, unsigned int flags) } else if ((thisCtx != NULL) && (peerCtx != NULL)) { LockedAccessor_CtxCrit_t peerCrit(peerCtx->criticalData()); // Add thisCtx to peerCtx's access list so that new allocations on peer will be made visible to this device: - bool isNewPeer = peerCrit->addPeer(thisCtx); + bool isNewPeer = peerCrit->addPeerWatcher(peerCtx, thisCtx); if (isNewPeer) { + tprintf(DB_MEM, "device=%s can now see all memory allocated on peer=%s\n", + thisCtx->toString().c_str(), peerCtx->toString().c_str()); am_memtracker_update_peers(peerCtx->getDevice()->_acc, peerCrit->peerCnt(), peerCrit->peerAgents()); } else { err = hipErrorPeerAccessAlreadyEnabled; @@ -134,7 +155,7 @@ hipError_t hipMemcpyPeer (void* dst, hipCtx_t dstCtx, const void* src, hipCtx_t // TODO - move to ihip memory copy implementaion. // HCC has a unified memory architecture so device specifiers are not required. - return hipMemcpy(dst, src, sizeBytes, hipMemcpyDefault); + return ihipLogStatus(hipMemcpy(dst, src, sizeBytes, hipMemcpyDefault)); }; @@ -145,7 +166,7 @@ hipError_t hipMemcpyPeerAsync (void* dst, hipCtx_t dstDevice, const void* src, h // TODO - move to ihip memory copy implementaion. // HCC has a unified memory architecture so device specifiers are not required. - return hipMemcpyAsync(dst, src, sizeBytes, hipMemcpyDefault, stream); + return ihipLogStatus(hip_internal::memcpyAsync(dst, src, sizeBytes, hipMemcpyDefault, stream)); }; @@ -158,7 +179,7 @@ hipError_t hipMemcpyPeerAsync (void* dst, hipCtx_t dstDevice, const void* src, h hipError_t hipDeviceCanAccessPeer (int* canAccessPeer, int deviceId, int peerDeviceId) { HIP_INIT_API(canAccessPeer, deviceId, peerDeviceId); - return hipDeviceCanAccessPeer(canAccessPeer, ihipGetPrimaryCtx(deviceId), ihipGetPrimaryCtx(peerDeviceId)); + return ihipLogStatus(ihipDeviceCanAccessPeer(canAccessPeer, ihipGetPrimaryCtx(deviceId), ihipGetPrimaryCtx(peerDeviceId))); } @@ -181,14 +202,14 @@ hipError_t hipDeviceEnablePeerAccess (int peerDeviceId, unsigned int flags) hipError_t hipMemcpyPeer (void* dst, int dstDevice, const void* src, int srcDevice, size_t sizeBytes) { HIP_INIT_API(dst, dstDevice, src, srcDevice, sizeBytes); - return hipMemcpyPeer(dst, ihipGetPrimaryCtx(dstDevice), src, ihipGetPrimaryCtx(srcDevice), sizeBytes); + return ihipLogStatus(hipMemcpyPeer(dst, ihipGetPrimaryCtx(dstDevice), src, ihipGetPrimaryCtx(srcDevice), sizeBytes)); } hipError_t hipMemcpyPeerAsync (void* dst, int dstDevice, const void* src, int srcDevice, size_t sizeBytes, hipStream_t stream) { HIP_INIT_API(dst, dstDevice, src, srcDevice, sizeBytes, stream); - return hipMemcpyPeerAsync(dst, ihipGetPrimaryCtx(dstDevice), src, ihipGetPrimaryCtx(srcDevice), sizeBytes, stream); + return ihipLogStatus(hip_internal::memcpyAsync(dst, src, sizeBytes, hipMemcpyDefault, stream)); } hipError_t hipCtxEnablePeerAccess (hipCtx_t peerCtx, unsigned int flags) diff --git a/projects/hip/tests/README.md b/projects/hip/tests/README.md index 56bb4e7edd..223bd149dc 100644 --- a/projects/hip/tests/README.md +++ b/projects/hip/tests/README.md @@ -53,11 +53,11 @@ ctest -R Memcpy ### If a test fails - how to debug a test -Extract the commandline from the testing log: +Find the test and commandline that fail: (From the test build directory, perhaps hip/tests/build) -$ grep -A3 -m2 hipMemcpy-size Testing/Temporary/LastTest.log -36/47 Testing: hipMemcpy-size -36/47 Test: hipMemcpy-size -Command: "/home/bensander/git/compute/external/hip/hip/tests/b6.hcc-LC.debug/runtimeApi/memory/hipMemcpy" "--tests" "0x6" -Directory: /home/bensander/git/compute/external/hip/hip/tests/b6.hcc-LC.debug/runtimeApi/memory +grep -IR hipMemcpy-modes -IR ../tests/ +../tests/src/runtimeApi/memory/hipMemcpy.cpp: * RUN_NAMED: %t hipMemcpy-modes --tests 0x1 + + + diff --git a/projects/hip/tests/src/deviceLib/hip_anyall.cpp b/projects/hip/tests/src/deviceLib/hip_anyall.cpp index aa025dad43..a562b7810e 100644 --- a/projects/hip/tests/src/deviceLib/hip_anyall.cpp +++ b/projects/hip/tests/src/deviceLib/hip_anyall.cpp @@ -45,11 +45,15 @@ int main(int argc, char *argv[]) { int warpSize, pshift; hipDeviceProp_t devProp; hipGetDeviceProperties(&devProp, 0); - if(strncmp(devProp.name,"Fiji",1)==0) -{ warpSize =64; - pshift =6; -} - else {warpSize =32; pshift=5;} + warpSize = devProp.warpSize; + + int w = warpSize; + pshift = 0; + while (w >>= 1) ++pshift; + + printf ("warpSize=%d pshift=%d\n", warpSize, pshift); + + int anycount =0; int allcount =0; int Num_Threads_per_Block = 1024; diff --git a/projects/hip/tests/src/deviceLib/hip_ballot.cpp b/projects/hip/tests/src/deviceLib/hip_ballot.cpp index d6df069351..236ceb57fe 100644 --- a/projects/hip/tests/src/deviceLib/hip_ballot.cpp +++ b/projects/hip/tests/src/deviceLib/hip_ballot.cpp @@ -48,9 +48,11 @@ int main(int argc, char *argv[]) hipDeviceProp_t devProp; hipGetDeviceProperties(&devProp, 0); - if(strncmp(devProp.name,"Fiji",1)==0) - {warpSize = 64; pshift =6;} - else {warpSize =32; pshift =5;} + warpSize = devProp.warpSize; + + int w = warpSize; + pshift = 0; + while (w >>= 1) ++pshift; unsigned int Num_Threads_per_Block = 512; unsigned int Num_Blocks_per_Grid = 1; diff --git a/projects/hip/tests/src/hipFuncDeviceSynchronize.cpp b/projects/hip/tests/src/hipFuncDeviceSynchronize.cpp index 6d05253536..930bc37b8b 100644 --- a/projects/hip/tests/src/hipFuncDeviceSynchronize.cpp +++ b/projects/hip/tests/src/hipFuncDeviceSynchronize.cpp @@ -35,6 +35,7 @@ THE SOFTWARE. __global__ void Iter(hipLaunchParm lp, int *Ad, int num){ int tx = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x; + // Kernel loop designed to execute very slowly... ... ... so we can test timing-related behavior below if(tx == 0){ for(int i = 0; i +#include "hip/hip_runtime.h" + +__global__ void Kernel(hipLaunchParm lp,volatile float* hostRes) +{ + int tid = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x; + hostRes[tid] = tid + 1; + __threadfence_system(); + // expecting that the data is getting flushed to host here! + // time waster for-loop (sleep) + for (int timeWater = 0; timeWater < 100000000; timeWater++); +} + +int main() +{ + size_t blocks = 2; + volatile float* hostRes; + hipHostMalloc((void**)&hostRes,blocks*sizeof(float),hipHostMallocMapped); + hostRes[0]=0; + hostRes[1]=0; + hipLaunchKernel(HIP_KERNEL_NAME(Kernel), dim3(1), dim3(blocks), 0, 0, hostRes); + int eleCounter = 0; + while (eleCounter < blocks) + { + // blocks until the value changes + while(hostRes[eleCounter] == 0); + printf("%f\n", hostRes[eleCounter]);; + eleCounter++; + } + hipHostFree((void *)hostRes); + return 0; +} + diff --git a/projects/hip/tests/src/runtimeApi/memory/hipMemoryAllocateCoherentDriver.cpp b/projects/hip/tests/src/runtimeApi/memory/hipMemoryAllocateCoherentDriver.cpp new file mode 100644 index 0000000000..dc512b41f8 --- /dev/null +++ b/projects/hip/tests/src/runtimeApi/memory/hipMemoryAllocateCoherentDriver.cpp @@ -0,0 +1,60 @@ +/* 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. */ + +/* HIT_START + * BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS -std=c++11 + * RUN: %t + * HIT_END + */ + +#include +#include +#include +#include +#include +#include +#include +#include "hip/hip_runtime.h" +using namespace std; + +string getRes(){ + FILE *in; + char buff[512], buff_2[512]; + string str = "./hipMemoryAllocateCoherent"; + if(!(in = popen(str.c_str(), "r"))){ + exit(1); + } + fgets(buff, sizeof(buff), in); + fgets(buff_2, sizeof(buff_2), in); + string str_buff = buff; + str_buff += buff_2; + pclose(in); + return str_buff; +} + +int main() { + setenv("HIP_COHERENT_HOST_ALLOC","1000,0,1",1); + string output = getRes(); + istringstream buffer(output); + double res1, res2; + buffer >> res1; + buffer >> res2; + if((res2-res1*2)>0.000001) + exit(1); + std::cout << "PASSED" << std::endl; + return 0; +}