AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

Merge branch 'amd-develop' into amd-master

Browse Source 
Change-Id: I5b104cbe9aa97f4152fcd7d5d86b606f69492f70
This commit is contained in:
Maneesh Gupta
2016-09-30 12:57:47 +05:30
parent 9df8d9b399 67c7f0ab5a
commit 91c34c976a
100 changed files with 3025 additions and 1558 deletions
Download Patch File Download Diff File Expand all files Collapse all files
hipamd/CMakeLists.txt
+43 -7
View File
@@ -145,6 +145,9 @@ add_to_config(_buildInfo COMPILE_HIP_ATP_MARKER)
#############################
# Build steps
#############################
# Rebuild cmake cache updates .buildInfo and .version
add_custom_target(update_build_and_version_info ALL COMMAND make rebuild_cache)
# Build clang hipify if enabled
if(BUILD_CLANG_HIPIFY)
add_subdirectory(clang-hipify)
@@ -153,18 +156,23 @@ endif()
# Build hip_hcc if platform is hcc
if(HIP_PLATFORM STREQUAL "hcc")
include_directories(${PROJECT_SOURCE_DIR}/include)
set(HIP_HCC_BUILD_FLAGS)
if(COMPILE_HIP_ATP_MARKER)
include_directories(/opt/rocm/profiler/CXLActivityLogger/include)
set(CMAKE_CXX_FLAGS " -DCOMPILE_HIP_ATP_MARKER=1 ${CMAKE_CXX_FLAGS}")
set(CMAKE_C_FLAGS " -DCOMPILE_HIP_ATP_MARKER=1 ${CMAKE_C_FLAGS}")
set(HIP_HCC_BUILD_FLAGS "${HIP_HCC_BUILD_FLAGS} -DCOMPILE_HIP_ATP_MARKER=1")
endif()
# Add HIP_VERSION to CMAKE_<LANG>_FLAGS
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/libhsakmt -stdlib=libc++")
# Set compiler and compiler flags
set(CMAKE_CXX_COMPILER "${HCC_HOME}/bin/hcc")
set(CMAKE_C_COMPILER "${HCC_HOME}/bin/hcc")
# Set HIP_HCC so we know this is HIP compile, some files are shared with HCC (staging_buffer).
set(CMAKE_CXX_FLAGS " -fPIC -hc -I${HCC_HOME}/include -I${HSA_PATH}/include -I/opt/rocm/libhsakmt/include/libhsakmt -stdlib=libc++ -DHIP_HCC ${CMAKE_CXX_FLAGS}")
set(CMAKE_C_FLAGS " -fPIC -hc -I${HCC_HOME}/include -I${HSA_PATH}/include -I/opt/rocm/libhsakmt/include/libhsakmt -stdlib=libc++ -DHIP_HCC ${CMAKE_C_FLAGS}")
set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} ${HIP_HCC_BUILD_FLAGS}")
set(CMAKE_C_FLAGS "${CMAKE_C_FLAGS} ${HIP_HCC_BUILD_FLAGS}")
set(SOURCE_FILES src/device_util.cpp
src/hip_hcc.cpp
@@ -177,7 +185,6 @@ if(HIP_PLATFORM STREQUAL "hcc")
src/hip_peer.cpp
src/hip_stream.cpp
src/hip_fp16.cpp
src/unpinned_copy_engine.cpp
src/hip_module.cpp)
if(${HIP_LIB_TYPE} EQUAL 0)
@@ -188,6 +195,13 @@ if(HIP_PLATFORM STREQUAL "hcc")
add_library(hip_hcc SHARED ${SOURCE_FILES})
endif()
# Generate hcc_version.txt
add_custom_target(query_hcc_version COMMAND ${HCC_HOME}/bin/hcc --version > ${PROJECT_BINARY_DIR}/hcc_version.tmp)
add_custom_target(check_hcc_version COMMAND ${CMAKE_COMMAND} -E copy_if_different ${PROJECT_BINARY_DIR}/hcc_version.tmp ${PROJECT_BINARY_DIR}/hcc_version.txt DEPENDS query_hcc_version)
set_source_files_properties(${PROJECT_BINARY_DIR}/hcc_version.txt PROPERTIES GENERATED TRUE)
set_source_files_properties(${SOURCE_FILES} PROPERTIES OBJECT_DEPENDS ${PROJECT_BINARY_DIR}/hcc_version.txt)
add_dependencies(hip_hcc check_hcc_version update_build_and_version_info)
# Generate .buildInfo
file(WRITE "${PROJECT_BINARY_DIR}/.buildInfo" ${_buildInfo})
endif()
@@ -300,3 +314,25 @@ if(POLICY CMP0037)
cmake_policy(POP)
endif()
#############################
# Testing steps
#############################
# Target: test
set(BUILD_DIR ${CMAKE_CURRENT_BINARY_DIR}/hip_tests)
configure_file(tests/hip_tests.txt ${BUILD_DIR}/CMakeLists.txt @ONLY)
if(POLICY CMP0037)
cmake_policy(PUSH)
cmake_policy(SET CMP0037 OLD)
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} .
COMMAND make -j ${DASH_JAY}
COMMAND make test
WORKING_DIRECTORY ${BUILD_DIR}
DEPENDS install_for_test)
if(POLICY CMP0037)
cmake_policy(POP)
endif()
hipamd/RELEASE.md
+1
View File
@@ -7,6 +7,7 @@ Upcoming:
number of commands (>1K) without synchronizing.
- Register keyword now silently ignored on HCC (previously would emit warning).
- Doc updates: Add some more frequently asked questions to FAQ, fix TOC in some files, review.
- Cookbook.
===================================================================================================
hipamd/bin/hipcc
+7 -6
View File
@@ -60,6 +60,7 @@ parse_config_file("$HIP_PATH/lib/.buildInfo", \%hipConfig);
#HIP_PLATFORM controls whether to use NVCC or HCC for compilation:
$HIP_PLATFORM= `$HIP_PATH/bin/hipconfig --platform` // "hcc";
$HIP_VERSION= `$HIP_PATH/bin/hipconfig --version`;
($HIP_VERSION_MAJOR, $HIP_VERSION_MINOR, $HIP_VERSION_PATCH) = split(/\./, $HIP_VERSION);
if ($verbose & 0x2) {
print ("HIP_PATH=$HIP_PATH\n");
@@ -151,14 +152,14 @@ if ($HIP_PLATFORM eq "hcc") {
$HIPCC="$CUDA_PATH/bin/nvcc";
$HIPCXXFLAGS .= " -I$CUDA_PATH/include";
$HIPLDFLAGS = "-lcuda -lcudart";
$HIPLDFLAGS = "-lcuda -lcudart -L$CUDA_PATH/lib64";
} else {
printf ("error: unknown HIP_PLATFORM = '$HIP_PLATFORM'");
exit (-1);
}
# Add paths to common HIP includes:
$HIPCXXFLAGS .= " -I$HIP_PATH/include" ;
$HIPCXXFLAGS .= " -I$HIP_PATH/include -DHIP_VERSION_MAJOR=$HIP_VERSION_MAJOR -DHIP_VERSION_MINOR=$HIP_VERSION_MINOR -DHIP_VERSION_PATCH=$HIP_VERSION_PATCH" ;
my $compileOnly = 0;
my $needCXXFLAGS = 0; # need to add CXX flags to compile step
@@ -224,12 +225,12 @@ foreach $arg (@ARGV)
$needLDFLAGS = 1;
}
if($arg eq '-stdlib=libc++' and $setStdLib eq 0)
if(($arg eq '-stdlib=libc++') and ($setStdLib eq 0))
{
$HIPCXXFLAGS .= " -stdlib=libc++";
$setStdLib = 1;
}
if($arg eq '-stdlib=libstdc++' and $setStdLib eq 0)
if(($arg eq '-stdlib=libstdc++') and ($setStdLib eq 0))
{
$HIPCXXFLAGS .= " -stdlib=libstdc++";
$setStdLib = 1;
@@ -296,7 +297,7 @@ if ($needHipHcc) {
$HIP_LIB_TYPE = $hipConfig{'HIP_LIB_TYPE'} // 0;
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/unpinned_copy_engine.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";
$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";
} elsif ($HIP_LIB_TYPE eq 1) {
$HIPLDFLAGS .= " -L$HIP_PATH/lib -lhip_hcc" ;
} else {
@@ -332,7 +333,7 @@ if ($printHipVersion) {
}
if ($runCmd) {
if ($HIP_PLATFORM eq "hcc" and exists($hipConfig{'HCC_VERSION'}) and $HCC_VERSION ne $hipConfig{'HCC_VERSION'}) {
print ("HIP was built using $hipConfig{'HCC_VERSION'}, but you are using $HCC_VERSION. Please rebuild HIP.\n") && die ();
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 ();
}
system ("$CMD") and die ();
}
hipamd/bin/hipify
+1 -1
View File
@@ -288,7 +288,7 @@ while (@ARGV) {
$ft{'mem'} += s/\bcudaMallocHost\b/hipHostMalloc/g; # note conversion to standard hipHost* naming convention
$ft{'mem'} += s/\bcudaFree\b/hipFree/g;
$ft{'mem'} += s/\bcudaFreeHost\b/hipHostFree/g; # note conversion to standard hipHost* naming convention
$ft{'mem'} += s/\bcudaHostAlloc\b/hipHostAlloc/g;
$ft{'mem'} += s/\bcudaHostAlloc\b/hipHostMalloc/g;
$ft{'mem'} += s/\bcudaHostGetDevicePointer\b/hipHostGetDevicePointer/g;
$ft{'mem'} += s/\bcudaHostAllocDefault\b/hipHostMallocDefault/g;
$ft{'mem'} += s/\bcudaHostAllocPortable\b/hipHostMallocPortable/g;
hipamd/cmake/FindHIP.cmake
+73 -62
View File
@@ -193,9 +193,15 @@ macro(HIP_GET_SOURCES_AND_OPTIONS _sources _cmake_options _hipcc_options _hcc_op
foreach(arg ${ARGN})
if("x${arg}" STREQUAL "xHIPCC_OPTIONS")
set(_hipcc_found_options TRUE)
set(_hcc_found_options FALSE)
set(_nvcc_found_options FALSE)
elseif("x${arg}" STREQUAL "xHCC_OPTIONS")
set(_hipcc_found_options FALSE)
set(_hcc_found_options TRUE)
set(_nvcc_found_options FALSE)
elseif("x${arg}" STREQUAL "xNVCC_OPTIONS")
set(_hipcc_found_options FALSE)
set(_hcc_found_options FALSE)
set(_nvcc_found_options TRUE)
elseif(
"x${arg}" STREQUAL "xEXCLUDE_FROM_ALL" OR
@@ -235,9 +241,9 @@ endmacro()
function(HIP_COMPUTE_BUILD_PATH path build_path)
# Convert to cmake style paths
file(TO_CMAKE_PATH "${path}" bpath)
if (IS_ABSOLUTE "${bpath}")
if(IS_ABSOLUTE "${bpath}")
string(FIND "${bpath}" "${CMAKE_CURRENT_BINARY_DIR}" _binary_dir_pos)
if (_binary_dir_pos EQUAL 0)
if(_binary_dir_pos EQUAL 0)
file(RELATIVE_PATH bpath "${CMAKE_CURRENT_BINARY_DIR}" "${bpath}")
else()
file(RELATIVE_PATH bpath "${CMAKE_CURRENT_SOURCE_DIR}" "${bpath}")
@@ -316,7 +322,7 @@ endmacro()
###############################################################################
# MACRO: Prepare cmake commands for the target
###############################################################################
macro(HIP_PREPARE_TARGET_COMMANDS _target _format _generated_files)
macro(HIP_PREPARE_TARGET_COMMANDS _target _format _generated_files _source_files)
set(_hip_flags "")
set(_hip_build_configuration "${CMAKE_BUILD_TYPE}")
if(HIP_HOST_COMPILATION_CPP)
@@ -378,6 +384,7 @@ macro(HIP_PREPARE_TARGET_COMMANDS _target _format _generated_files)
# Reset the output variable
set(_hip_generated_files "")
set(_hip_source_files "")
# Iterate over all arguments and create custom commands for all source files
foreach(file ${ARGN})
@@ -391,77 +398,79 @@ macro(HIP_PREPARE_TARGET_COMMANDS _target _format _generated_files)
set(host_flag TRUE)
endif()
if (NOT host_flag)
# Determine output directory
HIP_COMPUTE_BUILD_PATH("${file}" hip_build_path)
set(hip_compile_output_dir "${CMAKE_CURRENT_BINARY_DIR}/CMakeFiles/${_target}.dir/${hip_build_path}")
if(NOT host_flag)
# Determine output directory
HIP_COMPUTE_BUILD_PATH("${file}" hip_build_path)
set(hip_compile_output_dir "${CMAKE_CURRENT_BINARY_DIR}/CMakeFiles/${_target}.dir/${hip_build_path}")
get_filename_component(basename ${file} NAME)
set(generated_file_path "${hip_compile_output_dir}/${CMAKE_CFG_INTDIR}")
set(generated_file_basename "${_target}_generated_${basename}${generated_extension}")
get_filename_component(basename ${file} NAME)
set(generated_file_path "${hip_compile_output_dir}/${CMAKE_CFG_INTDIR}")
set(generated_file_basename "${_target}_generated_${basename}${generated_extension}")
# Set file names
set(generated_file "${generated_file_path}/${generated_file_basename}")
set(cmake_dependency_file "${hip_compile_output_dir}/${generated_file_basename}.depend")
set(custom_target_script_pregen "${hip_compile_output_dir}/${generated_file_basename}.cmake.pre-gen")
set(custom_target_script "${hip_compile_output_dir}/${generated_file_basename}.cmake")
# Set file names
set(generated_file "${generated_file_path}/${generated_file_basename}")
set(cmake_dependency_file "${hip_compile_output_dir}/${generated_file_basename}.depend")
set(custom_target_script_pregen "${hip_compile_output_dir}/${generated_file_basename}.cmake.pre-gen")
set(custom_target_script "${hip_compile_output_dir}/${generated_file_basename}.cmake")
# Set properties for object files
set_source_files_properties("${generated_file}"
PROPERTIES
EXTERNAL_OBJECT true # This is an object file not to be compiled, but only be linked
)
# Set properties for object files
set_source_files_properties("${generated_file}"
PROPERTIES
EXTERNAL_OBJECT true # This is an object file not to be compiled, but only be linked
)
# Don't add CMAKE_CURRENT_SOURCE_DIR if the path is already an absolute path
get_filename_component(file_path "${file}" PATH)
if(IS_ABSOLUTE "${file_path}")
set(source_file "${file}")
else()
set(source_file "${CMAKE_CURRENT_SOURCE_DIR}/${file}")
endif()
# Don't add CMAKE_CURRENT_SOURCE_DIR if the path is already an absolute path
get_filename_component(file_path "${file}" PATH)
if(IS_ABSOLUTE "${file_path}")
set(source_file "${file}")
else()
set(source_file "${CMAKE_CURRENT_SOURCE_DIR}/${file}")
endif()
# Bring in the dependencies
HIP_INCLUDE_HIPCC_DEPENDENCIES(${cmake_dependency_file})
# Bring in the dependencies
HIP_INCLUDE_HIPCC_DEPENDENCIES(${cmake_dependency_file})
# Configure the build script
configure_file("${HIP_run_hipcc}" "${custom_target_script_pregen}" @ONLY)
file(GENERATE
OUTPUT "${custom_target_script}"
INPUT "${custom_target_script_pregen}"
)
set(main_dep DEPENDS ${source_file})
set(verbose_output "$(VERBOSE)")
# Configure the build script
configure_file("${HIP_run_hipcc}" "${custom_target_script_pregen}" @ONLY)
file(GENERATE
OUTPUT "${custom_target_script}"
INPUT "${custom_target_script_pregen}"
)
set(main_dep DEPENDS ${source_file})
set(verbose_output "$(VERBOSE)")
# Create up the comment string
file(RELATIVE_PATH generated_file_relative_path "${CMAKE_BINARY_DIR}" "${generated_file}")
set(hip_build_comment_string "Building HIPCC object ${generated_file_relative_path}")
# Create up the comment string
file(RELATIVE_PATH generated_file_relative_path "${CMAKE_BINARY_DIR}" "${generated_file}")
set(hip_build_comment_string "Building HIPCC object ${generated_file_relative_path}")
# Build the generated file and dependency file
add_custom_command(
OUTPUT ${generated_file}
# These output files depend on the source_file and the contents of cmake_dependency_file
${main_dep}
DEPENDS ${HIP_HIPCC_DEPEND}
DEPENDS ${custom_target_script}
# Make sure the output directory exists before trying to write to it.
COMMAND ${CMAKE_COMMAND} -E make_directory "${generated_file_path}"
COMMAND ${CMAKE_COMMAND} ARGS
-D verbose:BOOL=${verbose_output}
-D build_configuration:STRING=${_hip_build_configuration}
-D "generated_file:STRING=${generated_file}"
-P "${custom_target_script}"
WORKING_DIRECTORY "${hip_compile_output_dir}"
COMMENT "${hip_build_comment_string}"
)
# Build the generated file and dependency file
add_custom_command(
OUTPUT ${generated_file}
# These output files depend on the source_file and the contents of cmake_dependency_file
${main_dep}
DEPENDS ${HIP_HIPCC_DEPEND}
DEPENDS ${custom_target_script}
# Make sure the output directory exists before trying to write to it.
COMMAND ${CMAKE_COMMAND} -E make_directory "${generated_file_path}"
COMMAND ${CMAKE_COMMAND} ARGS
-D verbose:BOOL=${verbose_output}
-D build_configuration:STRING=${_hip_build_configuration}
-D "generated_file:STRING=${generated_file}"
-P "${custom_target_script}"
WORKING_DIRECTORY "${hip_compile_output_dir}"
COMMENT "${hip_build_comment_string}"
)
# Make sure the build system knows the file is generated
set_source_files_properties(${generated_file} PROPERTIES GENERATED TRUE)
list(APPEND _hip_generated_files ${generated_file})
# Make sure the build system knows the file is generated
set_source_files_properties(${generated_file} PROPERTIES GENERATED TRUE)
list(APPEND _hip_generated_files ${generated_file})
list(APPEND _hip_source_files ${file})
endif()
endforeach()
# Set the return parameter
set(${_generated_files} ${_hip_generated_files})
set(${_source_files} ${_hip_source_files})
endmacro()
###############################################################################
@@ -470,7 +479,8 @@ endmacro()
macro(HIP_ADD_EXECUTABLE hip_target)
# Separate the sources from the options
HIP_GET_SOURCES_AND_OPTIONS(_sources _cmake_options _hipcc_options _hcc_options _nvcc_options ${ARGN})
HIP_PREPARE_TARGET_COMMANDS(${hip_target} OBJ _generated_files ${_sources} HIPCC_OPTIONS ${_hipcc_options} HCC_OPTIONS ${_hcc_options} NVCC_OPTIONS ${_nvcc_options})
HIP_PREPARE_TARGET_COMMANDS(${hip_target} OBJ _generated_files _source_files ${_sources} HIPCC_OPTIONS ${_hipcc_options} HCC_OPTIONS ${_hcc_options} NVCC_OPTIONS ${_nvcc_options})
list(REMOVE_ITEM _sources ${_source_files})
set(CMAKE_HIP_LINK_EXECUTABLE "${HIP_HIPCC_EXECUTABLE} <FLAGS> <CMAKE_CXX_LINK_FLAGS> <LINK_FLAGS> <OBJECTS> -o <TARGET>")
add_executable(${hip_target} ${_cmake_options} ${_generated_files} ${_sources})
set_target_properties(${hip_target} PROPERTIES LINKER_LANGUAGE HIP)
@@ -482,7 +492,8 @@ endmacro()
macro(HIP_ADD_LIBRARY hip_target)
# Separate the sources from the options
HIP_GET_SOURCES_AND_OPTIONS(_sources _cmake_options _hipcc_options _hcc_options _nvcc_options ${ARGN})
HIP_PREPARE_TARGET_COMMANDS(${hip_target} OBJ _generated_files ${_sources} ${_cmake_options} HIPCC_OPTIONS ${_hipcc_options} HCC_OPTIONS ${_hcc_options} NVCC_OPTIONS ${_nvcc_options})
HIP_PREPARE_TARGET_COMMANDS(${hip_target} OBJ _generated_files _source_files ${_sources} ${_cmake_options} HIPCC_OPTIONS ${_hipcc_options} HCC_OPTIONS ${_hcc_options} NVCC_OPTIONS ${_nvcc_options})
list(REMOVE_ITEM _sources ${_source_files})
add_library(${hip_target} ${_cmake_options} ${_generated_files} ${_sources})
set_target_properties(${hip_target} PROPERTIES LINKER_LANGUAGE ${HIP_C_OR_CXX})
endmacro()
hipamd/docs/markdown/CUDA_Runtime_API_functions_supported_by_HIP.md
+1 -1
View File
@@ -49,7 +49,7 @@
| `cudaStreamDestroy` | `hipStreamDestroy` | Destroys and cleans up an asynchronous stream. |
| `cudaStreamGetFlags` | `hipStreamGetFlags` | Query the flags of a stream. |
| `cudaStreamGetPriority` | | Query the priority of a stream. |
| `cudaStreamQuery` | | Queries an asynchronous stream for completion status. |
| `cudaStreamQuery` | `hipStreamQuery` | Queries an asynchronous stream for completion status. |
| `cudaStreamSynchronize` | `hipStreamSynchronize` | Waits for stream tasks to complete. |
| `cudaStreamWaitEvent` | `hipStreamWaitEvent` | Make a compute stream wait on an event. |
hipamd/docs/markdown/hip_faq.md
+44 -5
View File
@@ -10,6 +10,7 @@
- [What specific version of CUDA does HIP support?](#what-specific-version-of-cuda-does-hip-support)
- [What libraries does HIP support?](#what-libraries-does-hip-support)
- [How does HIP compare with OpenCL?](#how-does-hip-compare-with-opencl)
- [How does porting CUDA to HIP compare to porting CUDA to OpenCL?]
- [What hardware does HIP support?](#what-hardware-does-hip-support)
- [Does Hipify automatically convert all source code?](#does-hipify-automatically-convert-all-source-code)
- [What is NVCC?](#what-is-nvcc)
@@ -102,7 +103,7 @@ However, we can provide a rough summary of the features included in each CUDA SD
- CUDA 7.5
- float16 (under development)
- CUDA 8.0
- No new language features.
- TBD.
### What libraries does HIP support?
HIP includes growing support for the 4 key math libraries using hcBlas, hcFft, hcrng, and hcsparse).
@@ -127,16 +128,29 @@ HIP offers several benefits over OpenCL:
- HIP provides device-level control over memory allocation and placement.
- HIP offers an offline compilation model.
### How does porting CUDA to HIP compare to porting CUDA to OpenCL?
Both HIP and CUDA are dialects of C++, and thus porting between them is relatively straightforward.
Both dialects support templates, classes, lambdas, and other C++ constructs.
As one example, the hipify tool was originally a perl script that used simple text conversions from CUDA to HIP.
HIP and CUDA provide similar math library calls as well. In summary, the HIP philospohy was to make the HIP language close enough to CUDA that the porting effort is relatively simple.
This reduces the potential for error, and also makes it easy to automate the translation. HIP's goal is to quickly get the ported program running on both platforms with little manual intervention,
so that the programmer can focus on performance optimizations.
There have been several tools that have attempted to convert CUDA into OpenCL, such as CU2CL. OpenCL is a C99-based kernel language (rather than C++) and also does not support single-source compilation.
As a result, the OpenCL syntax is quite different than CUDA, and the porting tools have to perform some heroic transformations to bridge this gap.
The tools also struggle with more complex CUDA applications, in particular those that use templates, classes, or other C++ features inside the kernel.
### What hardware does HIP support?
- For AMD platforms, HIP runs on the same hardware that the HCC "hc" mode supports. See the ROCM documentation for the list of supported platforms.
- For Nvidia platforms, HIP requires Unified Memory and should run on a device which runs the CUDA SDK 6.0 or newer. We have tested the Nvidia Titan and K40.
### Does Hipify automatically convert all source code?
Typically, Hipify can automatically convert almost all run-time code, and the coordinate indexing device code.
Typically, Hipify can automatically convert almost all run-time code, and the coordinate indexing device code (i.e. threadIdx.x -> hipThreadIdx_x).
Most device code needs no additional conversion, since HIP and CUDA have similar names for math and built-in functions.
HIP currently requires manual addition of one more arguments to the kernel so that the host can communicate the execution configuration to the device.
The clang-hipify tool will automatically modify the kernel signature as needed (automating a step that used to be done manually)
Additional porting may be required to deal with architecture feature queries or with CUDA capabilities that HIP doesn't support.
Developers should always expect to perform some platform-specific tuning and optimization.
In general, developers should always expect to perform some platform-specific tuning and optimization.
### What is NVCC?
NVCC is Nvidia's compiler driver for compiling "CUDA C++" code into PTX or device code for Nvidia GPUs. It's a closed-source binary product that comes with CUDA SDKs.
@@ -172,7 +186,9 @@ A C++ dialect, hc is supported by the AMD HCC compiler. It provides C++ run time
### On HCC, can I link HIP code with host code compiled with another compiler such as gcc, icc, or clang ?
Yes! HIP/HCC generates the object code which conforms to the GCC ABI, and also links with libstdc++. This means you can compile host code with the compiler of your choice and link this
with GPU code compiler with HIP. Larger projects often contain a mixture of accelerator code (initially written in CUDA with nvcc) plus host code (compiled with gcc, icc, or clang). These projects
can convert the accelerator code to HIP, compile that code with hipcc, and link with object code from the preferred compiler.
can convert the accelerator code to HIP, compile that code with hipcc, and link with object code from the preferred compiler.S
### HIP detected my platform (hcc vs nvcc) incorrectly - what should I do?
@@ -189,6 +205,29 @@ If you see issues related to incorrect platform detection, please file an issue
### Can I install both CUDA SDK and HCC on same machine?
Yes. You can use HIP_PLATFORM to choose which path hipcc targets. This configuration can be useful when using HIP to develop an application which is portable to both AMD and NVIDIA.
### On CUDA, can I mix CUDA code with HIP code?
Yes. Most HIP data structures (hipStream_t, hipEvent_t) are typedefs to CUDA equivalents and can be intermixed. Both CUDA and HIP use integer device ids .
One notable exception is that hipError_t is a new type, and cannot be used where a cudaError_t is expected. In these cases, refactor the code to remove the expectation. Alternatively, hip_runtime_api.h defines functions which convert between the error code spaces:
hipErrorToCudaError
hipCUDAErrorTohipError
hipCUResultTohipError
If platform portability is important, use #ifdef __HIP_PLATFORM_NVCC__ to guard the CUDA-specific code.
### On HCC, can I use HC functionality with HIP?
Yes.
The code can include hc.hpp and use HC functions inside the kernel. A typical use case is to use AMD-specific hardware features such as the permute, swizzle, or DPP operations.
The "-stdlib=libc++" must be passed to hipcc in order to compile hc.hpp. See the 'bit_extract' sample for an example.
Also these functions can be used to extract HCC acclerator and accelerator_view structures from the HIP deviceId and hipStream_t:
hipHccGetAccelerator(int deviceId, hc::accelerator *acc);
hipError_t hipHccGetAcceleratorView(hipStream_t stream, hc::accelerator_view **av);
If platform portability is important, use #ifdef __HIP_PLATFORM_HIPCC__ to guard the HCC-specific code.
### 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.
hipamd/include/hcc_detail/hcc_acc.h
+4
View File
@@ -7,11 +7,15 @@
#include <hc.hpp>
/**
* @brief Return hc::accelerator associated with the specified deviceId
* @return #hipSuccess, #hipErrorInvalidDevice
*/
hipError_t hipHccGetAccelerator(int deviceId, hc::accelerator *acc);
/**
* @brief Return hc::accelerator_view associated with the specified stream
*
* If stream is 0, the accelerator_view for the default stream is returned.
* @return #hipSuccess
*/
hipError_t hipHccGetAcceleratorView(hipStream_t stream, hc::accelerator_view **av);
#endif
hipamd/include/hcc_detail/hip_hcc.h
+50 -109
View File
@@ -1,5 +1,5 @@
/*
Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved.
Link errors represented as this: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
@@ -21,8 +21,8 @@ THE SOFTWARE.
#define HIP_HCC_H
#include <hc.hpp>
#include <hsa/hsa.h>
#include "hip/hcc_detail/hip_util.h"
#include "hip/hcc_detail/unpinned_copy_engine.h"
#if defined(__HCC__) && (__hcc_workweek__ < 16354)
@@ -31,13 +31,6 @@ THE SOFTWARE.
// #define USE_MEMCPYTOSYMBOL
//
//Use the new HCC accelerator_view::copy instead of am_copy
#define USE_AV_COPY (__hcc_workweek__ >= 16351)
// Compile peer-to-peer support.
// >= 2 : use HCC hc:accelerator::get_is_peer
// >= 3 : use hc::am_memtracker_update_peers(...)
#define USE_PEER_TO_PEER 3
//---
@@ -55,8 +48,6 @@ extern int HIP_ATP_MARKER;
extern int HIP_ATP;
extern int HIP_DB;
extern int HIP_STAGING_SIZE; /* size of staging buffers, in KB */
extern int HIP_STAGING_BUFFERS; // TODO - remove, two buffers should be enough.
extern int HIP_PININPLACE;
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 */
@@ -64,7 +55,6 @@ extern int HIP_VISIBLE_DEVICES; /* Contains a comma-separated sequence of GPU id
//---
// Chicken bits for disabling functionality to work around potential issues:
extern int HIP_DISABLE_HW_KERNEL_DEP;
extern int HIP_DISABLE_HW_COPY_DEP;
//---
//Extern tls
@@ -98,15 +88,6 @@ extern const char *API_COLOR_END;
#define CTX_THREAD_SAFE 1
// If FORCE_COPY_DEP=1 , HIP runtime will add
// synchronization for copy commands in the same stream, regardless of command type.
// If FORCE_COPY_DEP=0 data copies of the same kind (H2H, H2D, D2H, D2D) are assumed to be implicitly ordered.
// ROCR runtime implementation currently provides this guarantee when using SDMA queues but not
// when using shader queues.
// TODO - measure if this matters for performance, in particular for back-to-back small copies.
// If not, we can simplify the copy dependency tracking by collapsing to a single Copy type, and always forcing dependencies for copy commands.
#define FORCE_SAMEDIR_COPY_DEP 1
// Compile debug trace mode - this prints debug messages to stderr when env var HIP_DB is set.
// May be set to 0 to remove debug if checks - possible code size and performance difference?
@@ -128,8 +109,17 @@ extern const char *API_COLOR_END;
#endif
// #include CPP files to produce one object file
#define ONE_OBJECT_FILE 0
#define DB_SHOW_TID 1
#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;
#endif
// Compile support for trace markers that are displayed on CodeXL GUI at start/stop of each function boundary.
@@ -150,7 +140,8 @@ extern const char *API_COLOR_END;
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) {\
fprintf (stderr, "%s<<hip-api: %s\n%s" , API_COLOR, s.c_str(), API_COLOR_END);\
COMPUTE_TID_STR\
fprintf (stderr, "%s<<hip-api:%s %s\n%s" , API_COLOR, tid_ss.str().c_str(), s.c_str(), API_COLOR_END);\
}\
SCOPED_MARKER(s.c_str(), "HIP", NULL);\
}\
@@ -208,12 +199,15 @@ static const char *dbName [] =
KNRM "hip-copy2",
};
#if COMPILE_HIP_DB
#define tprintf(trace_level, ...) {\
if (HIP_DB & (1<<(trace_level))) {\
fprintf (stderr, " %s:", dbName[trace_level]); \
fprintf (stderr, __VA_ARGS__);\
fprintf (stderr, "%s", KNRM); \
char msgStr[1000];\
snprintf(msgStr, 2000, __VA_ARGS__);\
COMPUTE_TID_STR\
fprintf (stderr, " %s%s:%s%s", dbName[trace_level], tid_ss.str().c_str(), msgStr, KNRM); \
}\
}
#else
@@ -246,40 +240,6 @@ extern "C" {
const hipStream_t hipStreamNull = 0x0;
enum ihipCommand_t {
ihipCommandCopyH2H,
ihipCommandCopyH2D,
ihipCommandCopyD2H,
ihipCommandCopyD2D,
ihipCommandCopyP2P,
ihipCommandKernel,
};
static const char* ihipCommandName[] = {
"CopyH2H", "CopyH2D", "CopyD2H", "CopyD2D", "CopyP2P", "Kernel"
};
typedef uint64_t SIGSEQNUM;
//---
// Small wrapper around signals.
// Designed to be used from stream.
// TODO-someday refactor this class so it can be stored in a vector<>
// we already store the index here so we can use for garbage collection.
struct ihipSignal_t {
hsa_signal_t _hsaSignal; // hsa signal handle
int _index; // Index in pool, used for garbage collection.
SIGSEQNUM _sigId; // unique sequentially increasing ID.
ihipSignal_t();
~ihipSignal_t();
void release();
};
// Used to remove lock, for performance or stimulating bugs.
class FakeMutex
{
@@ -318,18 +278,21 @@ public:
_autoUnlock(autoUnlock)
{
tprintf(DB_SYNC, "lock critical data %s.%p\n", typeid(T).name(), _criticalData);
_criticalData->_mutex.lock();
};
~LockedAccessor()
{
if (_autoUnlock) {
tprintf(DB_SYNC, "auto-unlock critical data %s.%p\n",typeid(T).name(), _criticalData);
_criticalData->_mutex.unlock();
}
}
void unlock()
{
tprintf(DB_SYNC, "unlock critical data %s.%p\n", typeid(T).name(), _criticalData);
_criticalData->_mutex.unlock();
}
@@ -391,56 +354,34 @@ class ihipStreamCriticalBase_t : public LockedBase<MUTEX_TYPE>
{
public:
ihipStreamCriticalBase_t(hc::accelerator_view av) :
_last_command_type(ihipCommandCopyH2H),
_last_copy_signal(NULL),
_signalCursor(0),
_oldest_live_sig_id(1),
_streamSigId(0),
_kernelCnt(0),
_signalCnt(0),
_av(av)
{
_signalPool.resize(HIP_STREAM_SIGNALS > 0 ? HIP_STREAM_SIGNALS : 1);
};
~ihipStreamCriticalBase_t() {
_signalPool.clear();
}
ihipStreamCriticalBase_t<StreamMutex> * mlock() { LockedBase<MUTEX_TYPE>::lock(); return this;};
public:
// Critical Data:
ihipCommand_t _last_command_type; // type of the last command
// signal of last copy command sent to the stream.
// May be NULL, indicating the previous command has completley finished and future commands don't need to create a dependency.
// Copy can be either H2D or D2H.
ihipSignal_t *_last_copy_signal;
hc::completion_future _last_kernel_future; // Completion future of last kernel command sent to GPU.
// Signal pool:
int _signalCursor;
SIGSEQNUM _oldest_live_sig_id; // oldest live seq_id, anything < this can be allocated.
std::deque<ihipSignal_t> _signalPool; // Pool of signals for use by this stream.
uint32_t _signalCnt; // Count of inflight commands using signals from the signal pool.
// Each copy may use 1-2 signals depending on command transitions:
// 2 are required if a barrier packet is inserted.
// TODO - remove _kernelCnt mechanism:
uint32_t _kernelCnt; // Count of inflight kernels in this stream. Reset at ::wait().
SIGSEQNUM _streamSigId; // Monotonically increasing unique signal id.
hc::accelerator_view _av;
std::vector<hc::completion_future*> _cfs;
};
// if HIP code needs to acquire locks for both ihipCtx_t and ihipStream_t, it should first acquire the lock
// for the ihipCtx_t and then for the individual streams. The locks should not be acquired in reverse order
// or deadlock may occur. In some cases, it may be possible to reduce the range where the locks must be held.
// HIP routines should avoid acquiring and releasing the same lock during the execution of a single HIP API.
typedef ihipStreamCriticalBase_t<StreamMutex> ihipStreamCritical_t;
typedef LockedAccessor<ihipStreamCritical_t> LockedAccessor_StreamCrit_t;
//---
// Internal stream structure.
class ihipStream_t {
public:
@@ -449,21 +390,18 @@ typedef uint64_t SeqNum_t ;
~ihipStream_t();
// kind is hipMemcpyKind
void copySync (LockedAccessor_StreamCrit_t &crit, void* dst, const void* src, size_t sizeBytes, unsigned kind, bool resolveOn = true);
void locked_copySync (void* dst, const void* src, size_t sizeBytes, unsigned kind, bool resolveOn = true);
void copyAsync(void* dst, const void* src, size_t sizeBytes, unsigned kind);
void locked_copyAsync(void* dst, const void* src, size_t sizeBytes, unsigned kind);
int preCopyCommand(LockedAccessor_StreamCrit_t &crit, ihipSignal_t *lastCopy, hsa_signal_t *waitSignal, ihipCommand_t copyType);
//---
// Member functions that begin with locked_ are thread-safe accessors - these acquire / release the critical mutex.
LockedAccessor_StreamCrit_t lockopen_preKernelCommand();
void lockclose_postKernelCommand(hc::completion_future &kernel_future);
void lockclose_postKernelCommand(hc::accelerator_view *av);
void locked_reclaimSignals(SIGSEQNUM sigNum);
void locked_wait(bool assertQueueEmpty=false);
hc::accelerator_view* locked_getAv() { LockedAccessor_StreamCrit_t crit(_criticalData); return &(crit->_av); };
@@ -471,8 +409,6 @@ typedef uint64_t SeqNum_t ;
void locked_waitEvent(hipEvent_t event);
void locked_recordEvent(hipEvent_t event);
void addCFtoStream(LockedAccessor_StreamCrit_t &crit, hc::completion_future* cf);
void waitOnAllCFs(LockedAccessor_StreamCrit_t &crit);
//---
@@ -485,9 +421,6 @@ typedef uint64_t SeqNum_t ;
uint32_t groupSegmentSize, uint32_t sharedMemBytes,
void *kernarg, size_t kernSize, uint64_t kernel);
// Non-threadsafe accessors - must be protected by high-level stream lock with accessor passed to function.
SIGSEQNUM lastCopySeqId (LockedAccessor_StreamCrit_t &crit) const { return crit->_last_copy_signal ? crit->_last_copy_signal->_sigId : 0; };
ihipSignal_t * allocSignal (LockedAccessor_StreamCrit_t &crit);
//-- Non-racy accessors:
@@ -504,17 +437,16 @@ public:
private:
void enqueueBarrier(hsa_queue_t* queue, ihipSignal_t *depSignal, ihipSignal_t *completionSignal);
void waitCopy(LockedAccessor_StreamCrit_t &crit, ihipSignal_t *signal);
// The unsigned return is hipMemcpyKind
unsigned resolveMemcpyDirection(bool srcTracked, bool dstTracked, bool srcInDeviceMem, bool dstInDeviceMem);
void setAsyncCopyAgents(unsigned kind, ihipCommand_t *commandType, hsa_agent_t *srcAgent, hsa_agent_t *dstAgent);
bool canSeePeerMemory(const ihipCtx_t *thisCtx, ihipCtx_t *dstCtx, ihipCtx_t *srcCtx);
private: // Data
// Critical Data. THis MUST be accessed through LockedAccessor_StreamCrit_t
// Critical Data - MUST be accessed through LockedAccessor_StreamCrit_t
ihipStreamCritical_t _criticalData;
ihipCtx_t *_ctx; // parent context that owns this stream.
@@ -545,8 +477,6 @@ struct ihipEvent_t {
hc::completion_future _marker;
uint64_t _timestamp; // store timestamp, may be set on host or by marker.
SIGSEQNUM _copySeqId;
} ;
@@ -575,7 +505,7 @@ public:
unsigned _computeUnits;
hipDeviceProp_t _props; // saved device properties.
UnpinnedCopyEngine *_stagingBuffer[2]; // one buffer for each direction.
// TODO - report this through device properties, base on HCC API call.
int _isLargeBar;
ihipCtx_t *_primaryCtx;
@@ -614,10 +544,11 @@ public:
// Peer Accessor classes:
bool isPeer(const ihipCtx_t *peer); // returns Trus if peer has access to memory physically located on this device.
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;
uint32_t peerCnt() const { return _peerCnt; };
hsa_agent_t *peerAgents() const { return _peerAgents; };
@@ -633,6 +564,7 @@ private:
//--- Peer Tracker:
// 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<ihipCtx_t*> _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.)
@@ -676,6 +608,8 @@ public: // Functions:
// TODO - review uses of getWriteableDevice(), can these be converted to getDevice()
ihipDevice_t *getWriteableDevice() const { return _device; };
std::string toString() const;
public: // Data
// The NULL stream is used if no other stream is specified.
// Default stream has special synchronization properties with other streams.
@@ -761,5 +695,12 @@ inline std::ostream& operator<<(std::ostream& os, const hipEvent_t& e)
return os;
}
inline std::ostream& operator<<(std::ostream& os, const ihipCtx_t* c)
{
os << "ctx:" << static_cast<const void*> (c)
<< " dev:" << c->getDevice()->_deviceId;
return os;
}
#endif
hipamd/include/hcc_detail/hip_runtime.h
+11 -1
View File
@@ -33,10 +33,13 @@ THE SOFTWARE.
//#include <cstring>
#if __cplusplus
#include <cmath>
#else
#include <math.h>
#include <string.h>
#include <stddef.h>
#endif
// Define NVCC_COMPAT for CUDA compatibility
#define NVCC_COMPAT
#define CUDA_SUCCESS hipSuccess
@@ -493,6 +496,13 @@ __device__ float __dsqrt_rd(double x);
__device__ float __dsqrt_rn(double x);
__device__ float __dsqrt_ru(double x);
__device__ float __dsqrt_rz(double x);
/**
* CUDA 8 device function features
*/
/**
* Kernel launching
*/
hipamd/include/hcc_detail/hip_runtime_api.h
+184 -133
View File
@@ -267,7 +267,7 @@ hipError_t hipGetDevice(int *deviceId);
*
* @returns #hipSuccess, #hipErrorNoDevice
*
*
*
* Returns in @p *count the number of devices that have ability to run compute commands. If there are no such devices, then @ref hipGetDeviceCount will return #hipErrorNoDevice.
* If 1 or more devices can be found, then hipGetDeviceCount returns #hipSuccess.
*/
@@ -279,7 +279,7 @@ hipError_t hipGetDeviceCount(int *count);
* @param [out] pi pointer to value to return
* @param [in] attr attribute to query
* @param [in] deviceId which device to query for information
*
*
* @returns #hipSuccess, #hipErrorInvalidDevice, #hipErrorInvalidValue
*/
hipError_t hipDeviceGetAttribute(int* pi, hipDeviceAttribute_t attr, int deviceId);
@@ -328,7 +328,7 @@ hipError_t hipDeviceGetCacheConfig ( hipFuncCache *cacheConfig );
* @brief Set Cache configuration for a specific function
*
* @param [in] config;
*
*
* @returns #hipSuccess, #hipErrorInitializationError
* Note: AMD devices and recent Nvidia GPUS do not support reconfigurable cache. This hint is ignored on those architectures.
*
@@ -458,27 +458,13 @@ const char *hipGetErrorString(hipError_t hipError);
* @{
*
* The following Stream APIs are not (yet) supported in HIP:
* - cudaStreamAddCallback
* - cudaStreamAttachMemAsync
* - cudaStreamCreateWithPriority
* - cudaStreamGetPriority
* - cudaStreamWaitEvent
*/
/**
* @brief Create an asynchronous stream.
*
* @param[in, out] stream Pointer to new stream
* @param[in ] flags to control stream creation.
* @return #hipSuccess, #hipErrorInvalidValue
*
* Create a new asynchronous stream. @p stream returns an opaque handle that can be used to reference the newly
* created stream in subsequent hipStream* commands. The stream is allocated on the heap and will remain allocated
*
* even if the handle goes out-of-scope. To release the memory used by the stream, applicaiton must call hipStreamDestroy.
* Flags controls behavior of the stream. See #hipStreamDefault, #hipStreamNonBlocking.
*/
hipError_t hipStreamCreateWithFlags(hipStream_t *stream, unsigned int flags);
/**
* @brief Create an asynchronous stream.
@@ -490,15 +476,82 @@ hipError_t hipStreamCreateWithFlags(hipStream_t *stream, unsigned int flags);
* created stream in subsequent hipStream* commands. The stream is allocated on the heap and will remain allocated
* even if the handle goes out-of-scope. To release the memory used by the stream, applicaiton must call hipStreamDestroy.
*
* @return #hipSuccess, #hipErrorInvalidValue
*
* @see hipStreamDestroy
*
* @return
*
* @see hipStreamCreateWithFlags, hipStreamSynchronize, hipStreamWaitEvent, hipStreamDestroy
*/
hipError_t hipStreamCreate(hipStream_t *stream);
/**
* @brief Create an asynchronous stream.
*
* @param[in, out] stream Pointer to new stream
* @param[in ] flags to control stream creation.
* @return #hipSuccess, #hipErrorInvalidValue
*
* Create a new asynchronous stream. @p stream returns an opaque handle that can be used to reference the newly
* created stream in subsequent hipStream* commands. The stream is allocated on the heap and will remain allocated
* even if the handle goes out-of-scope. To release the memory used by the stream, applicaiton must call hipStreamDestroy.
* Flags controls behavior of the stream. See #hipStreamDefault, #hipStreamNonBlocking.
*
*
* @see hipStreamCreate, hipStreamSynchronize, hipStreamWaitEvent, hipStreamDestroy
*/
hipError_t hipStreamCreateWithFlags(hipStream_t *stream, unsigned int flags);
/**
* @brief Destroys the specified stream.
*
* @param[in, out] stream Valid pointer to hipStream_t. This function writes the memory with the newly created stream.
* @return #hipSuccess #hipErrorInvalidResourceHandle
*
* Destroys the specified stream.
*
* If commands are still executing on the specified stream, some may complete execution before the queue is deleted.
*
* The queue may be destroyed while some commands are still inflight, or may wait for all commands queued to the stream
* before destroying it.
*
* @see hipStreamCreate, hipStreamCreateWithFlags, hipStreamQuery, hipStreamWaitEvent, hipStreamSynchronize
*/
hipError_t hipStreamDestroy(hipStream_t stream);
/**
* @brief Return #hipSuccess if all of the operations in the specified @p stream have completed, or #hipErrorNotReady if not.
*
* @param[in] stream stream to query
*
* @return #hipSuccess, #hipErrorNotReady, #hipErrorInvalidResourceHandle
*
* This is thread-safe and returns a snapshot of the current state of the queue. However, if other host threads are sending work to the stream,
* the status may change immediately after the function is called. It is typically used for debug.
*
* @see hipStreamCreate, hipStreamCreateWithFlags, hipStreamWaitEvent, hipStreamSynchronize, hipStreamDestroy
*/
hipError_t hipStreamQuery(hipStream_t stream);
/**
* @brief Wait for all commands in stream to complete.
*
* @param[in] stream stream identifier.
*
* @return #hipSuccess, #hipErrorInvalidResourceHandle
*
* If the null stream is specified, this command blocks until all
* This command honors the hipDeviceLaunchBlocking flag, which controls whether the wait is active or blocking.
* This command is host-synchronous : the host will block until the stream is empty.
*
* @see hipStreamCreate, hipStreamCreateWithFlags, hipStreamWaitEvent, hipStreamDestroy
*
*/
hipError_t hipStreamSynchronize(hipStream_t stream);
/**
* @brief Make the specified compute stream wait for an event
*
@@ -512,67 +565,25 @@ hipError_t hipStreamCreate(hipStream_t *stream);
* All future work submitted to @p stream will wait until @p event reports completion before beginning execution.
* This function is host-asynchronous and the function may return before the wait has completed.
*
* @see hipStreamCreate, hipStreamCreateWithFlags, hipStreamSynchronize, hipStreamDestroy
*
*/
hipError_t hipStreamWaitEvent(hipStream_t stream, hipEvent_t event, unsigned int flags);
/**
* @brief Return #hipSuccess if all of the operations in the specified @p stream have completed, or #hipErrorNotReady if not.
*
* @param[in] stream stream to query
*
* @return #hipSuccess, #hipErrorNotReady
*
* This is thread-safe and returns a snapshot of the current state of the queue. However, if other host threads are sending work to the stream,
* the status may change immediately after the function is called. It is typically used for debug.
*/
hipError_t hipStreamQuery(hipStream_t stream);
/**
* @brief Wait for all commands in stream to complete.
*
* If the null stream is specified, this command blocks until all
*
* This command honors the hipDeviceLaunchBlocking flag, which controls whether the wait is active or blocking.
*
* This command is host-synchronous : the host will block until the stream is empty.
*
* TODO
*/
hipError_t hipStreamSynchronize(hipStream_t stream);
/**
* @brief Destroys the specified stream.
*
* @param[in, out] stream Valid pointer to hipStream_t. This function writes the memory with the newly created stream.
* @return #hipSuccess
*
* Destroys the specified stream.
*
* If commands are still executing on the specified stream, some may complete execution before the queue is deleted.
*
* The queue may be destroyed while some commands are still inflight, or may wait for all commands queued to the stream
* before destroying it.
*/
hipError_t hipStreamDestroy(hipStream_t stream);
/**
* @brief Return flags associated with this stream.
*
* @param[in] stream
* @param[in,out] flags
* @param[in] stream stream to be queried
* @param[in,out] flags Pointer to an unsigned integer in which the stream's flags are returned
* @return #hipSuccess, #hipErrorInvalidValue, #hipErrorInvalidResourceHandle
*
* @returns #hipSuccess #hipErrorInvalidValue #hipErrorInvalidResourceHandle
*
* Return flags associated with this stream in *@p flags.
*
* @see hipStreamCreateWithFlags
*
* @returns #hipSuccess
* @see hipStreamCreateWithFlags
*/
hipError_t hipStreamGetFlags(hipStream_t stream, unsigned int *flags);
@@ -766,23 +777,10 @@ hipError_t hipPointerGetAttributes(hipPointerAttribute_t *attributes, void* ptr)
*
* @return #hipSuccess
*
* @see hipMallocPitch, hipFree, hipMallocArray, hipFreeArray, hipMalloc3D, hipMalloc3DArray, hipMallocHost, hipFreeHost, hipHostAlloc
* @see hipMallocPitch, hipFree, hipMallocArray, hipFreeArray, hipMalloc3D, hipMalloc3DArray, hipHostFree, hipHostMalloc
*/
hipError_t hipMalloc(void** ptr, size_t size) ;
/**
* @brief Allocate pinned host memory
*
* @param[out] ptr Pointer to the allocated host pinned memory
* @param[in] size Requested memory size
*
* @return #hipSuccess, #hipErrorMemoryAllocation
*
* @see hipMalloc, hipMallocPitch, hipMallocArray, hipMalloc3D, hipMalloc3DArray, hipHostAlloc, hipFree, hipFreeArray, hipMallocHost, hipFreeHost, hipHostAlloc
*/
hipError_t hipMallocHost(void** ptr, size_t size) __attribute__((deprecated("use hipHostMalloc instead"))) ;
/**
* @brief Allocate device accessible page locked host memory
*
@@ -792,21 +790,20 @@ hipError_t hipMallocHost(void** ptr, size_t size) __attribute__((deprecated("use
*
* @return #hipSuccess, #hipErrorMemoryAllocation
*
* @see hipSetDeviceFlags, hipMallocHost, hipFreeHost
* @see hipSetDeviceFlags, hipHostFree
*/
hipError_t hipHostMalloc(void** ptr, size_t size, unsigned int flags) ;
hipError_t hipHostAlloc(void** ptr, size_t size, unsigned int flags) __attribute__((deprecated("use hipHostMalloc instead"))) ;;
/**
* @brief Get Device pointer from Host Pointer allocated through hipHostAlloc
* @brief Get Device pointer from Host Pointer allocated through hipHostMalloc
*
* @param[out] dstPtr Device Pointer mapped to passed host pointer
* @param[in] hstPtr Host Pointer allocated through hipHostAlloc
* @param[in] hstPtr Host Pointer allocated through hipHostMalloc
* @param[in] flags Flags to be passed for extension
*
*
* @return #hipSuccess, #hipErrorInvalidValue, #hipErrorMemoryAllocation
*
* @see hipSetDeviceFlags, hipHostAlloc
* @see hipSetDeviceFlags, hipHostMalloc
*/
hipError_t hipHostGetDevicePointer(void** devPtr, void* hstPtr, unsigned int flags) ;
@@ -817,7 +814,7 @@ hipError_t hipHostGetDevicePointer(void** devPtr, void* hstPtr, unsigned int fla
* @param[in] hostPtr Host Pointer allocated through hipHostMalloc
* @return #hipSuccess, #hipErrorInvalidValue
*
* @see hipHostAlloc
* @see hipHostMalloc
*/
hipError_t hipHostGetFlags(unsigned int* flagsPtr, void* hostPtr) ;
@@ -850,7 +847,7 @@ hipError_t hipHostGetFlags(unsigned int* flagsPtr, void* hostPtr) ;
* from the other registered memory region.
*
* @return #hipSuccess, #hipErrorMemoryAllocation
*
*
* @see hipHostUnregister, hipHostGetFlags, hipHostGetDevicePointer
*/
hipError_t hipHostRegister(void* hostPtr, size_t sizeBytes, unsigned int flags) ;
@@ -876,8 +873,8 @@ hipError_t hipHostUnregister(void* hostPtr) ;
* @param[in] width Requested pitched allocation width (in bytes)
* @param[in] height Requested pitched allocation height
* @return Error code
*
* @see hipMalloc, hipFree, hipMallocArray, hipFreeArray, hipMallocHost, hipFreeHost, hipMalloc3D, hipMalloc3DArray, hipHostAlloc
*
* @see hipMalloc, hipFree, hipMallocArray, hipFreeArray, hipHostFree, hipMalloc3D, hipMalloc3DArray, hipHostMalloc
*/
hipError_t hipMallocPitch(void** ptr, size_t* pitch, size_t width, size_t height);
@@ -891,24 +888,10 @@ hipError_t hipMallocPitch(void** ptr, size_t* pitch, size_t width, size_t height
* @return #hipSuccess
* @return #hipErrorInvalidDevicePointer (if pointer is invalid, including host pointers allocated with hipHostMalloc)
*
* @see hipMalloc, hipMallocPitch, hipMallocArray, hipFreeArray, hipMallocHost, hipFreeHost, hipMalloc3D, hipMalloc3DArray, hipHostAlloc
* @see hipMalloc, hipMallocPitch, hipMallocArray, hipFreeArray, hipHostFree, hipMalloc3D, hipMalloc3DArray, hipHostMalloc
*/
hipError_t hipFree(void* ptr);
/**
* @brief Free memory allocated by the hcc hip host memory allocation API. [Deprecated.]
*
* @param[in] ptr Pointer to memory to be freed
* @return #hipSuccess,
* #hipErrorInvalidValue (if pointer is invalid, including device pointers allocated with hipMalloc)
* @see hipHostFree
*/
hipError_t hipFreeHost(void* ptr) __attribute__((deprecated("use hipHostFree instead"))) ;
/**
* @brief Free memory allocated by the hcc hip host memory allocation API
* This API performs an implicit hipDeviceSynchronize() call.
@@ -918,12 +901,10 @@ hipError_t hipFreeHost(void* ptr) __attribute__((deprecated("use hipHostFree ins
* @return #hipSuccess,
* #hipErrorInvalidValue (if pointer is invalid, including device pointers allocated with hipMalloc)
*
* @see hipMalloc, hipMallocPitch, hipFree, hipMallocArray, hipFreeArray, hipMallocHost, hipMalloc3D, hipMalloc3DArray, hipHostAlloc
* @see hipMalloc, hipMallocPitch, hipFree, hipMallocArray, hipFreeArray, hipMalloc3D, hipMalloc3DArray, hipHostMalloc
*/
hipError_t hipHostFree(void* ptr);
/**
* @brief Copy data from src to dst.
*
@@ -1102,7 +1083,7 @@ hipError_t hipMemset(void* dst, int value, size_t sizeBytes );
* @return #hipSuccess, #hipErrorInvalidValue, #hipErrorMemoryFree
*/
#if __cplusplus
hipError_t hipMemsetAsync(void* dst, int value, size_t sizeBytes, hipStream_t = 0 );
hipError_t hipMemsetAsync(void* dst, int value, size_t sizeBytes, hipStream_t stream = 0 );
#else
hipError_t hipMemsetAsync(void* dst, int value, size_t sizeBytes, hipStream_t stream);
#endif
@@ -1112,7 +1093,7 @@ hipError_t hipMemsetAsync(void* dst, int value, size_t sizeBytes, hipStream_t st
* Return snapshot of free memory, and total allocatable memory on the device.
*
* Returns in *free a snapshot of the current free memory.
* @returns #hipSuccess, #hipErrorInvalidDevice, #hipErrorInvalidValue (if free != NULL due to bugs)
* @returns #hipSuccess, #hipErrorInvalidDevice, #hipErrorInvalidValue
* @warning On HCC, the free memory only accounts for memory allocated by this process and may be optimistic.
**/
hipError_t hipMemGetInfo (size_t * free, size_t * total) ;
@@ -1256,10 +1237,21 @@ hipError_t hipInit(unsigned int flags) ;
* @param [in] flags
* @param [in] associated device handle
*
* @returns #hipSuccess, #hipErrorInvalidContext
* @return #hipSuccess
*
* @see hipCtxDestroy, hipCtxGetFlags, hipCtxPopCurrent, hipCtxGetCurrent, hipCtxPushCurrent, hipCtxSetCacheConfig, hipCtxSynchronize, hipCtxGetDevice
*/
hipError_t hipCtxCreate(hipCtx_t *ctx, unsigned int flags, hipDevice_t device);
/**
* @brief Destroy a HIP context.
*
* @param [in] ctx Context to destroy
*
* @returns #hipSuccess, #hipErrorInvalidValue
*
* @see hipCtxCreate, hipCtxGetFlags, hipCtxPopCurrent, hipCtxGetCurrent,hipCtxSetCurrent, hipCtxPushCurrent, hipCtxSetCacheConfig, hipCtxSynchronize , hipCtxGetDevice
*/
hipError_t hipCtxDestroy(hipCtx_t ctx);
/**
@@ -1267,9 +1259,10 @@ hipError_t hipCtxDestroy(hipCtx_t ctx);
*
* @param [out] ctx
*
* @returns #hipSuccess
* @returns #hipSuccess, #hipErrorInvalidContext
*
* @see hipCtxCreate, hipCtxDestroy, hipCtxGetFlags, hipCtxSetCurrent, hipCtxGetCurrent, hipCtxPushCurrent, hipCtxSetCacheConfig, hipCtxSynchronize, hipCtxGetDevice
*/
hipError_t hipCtxPopCurrent(hipCtx_t* ctx);
/**
@@ -1278,8 +1271,9 @@ hipError_t hipCtxPopCurrent(hipCtx_t* ctx);
* @param [in] ctx
*
* @returns #hipSuccess, #hipErrorInvalidContext
*
* @see hipCtxCreate, hipCtxDestroy, hipCtxGetFlags, hipCtxPopCurrent, hipCtxGetCurrent, hipCtxPushCurrent, hipCtxSetCacheConfig, hipCtxSynchronize , hipCtxGetDevice
*/
hipError_t hipCtxPushCurrent(hipCtx_t ctx);
/**
@@ -1287,9 +1281,10 @@ hipError_t hipCtxPushCurrent(hipCtx_t ctx);
*
* @param [in] ctx
*
* @returns #hipSuccess
* @returns #hipSuccess, #hipErrorInvalidContext
*
* @see hipCtxCreate, hipCtxDestroy, hipCtxGetFlags, hipCtxPopCurrent, hipCtxGetCurrent, hipCtxPushCurrent, hipCtxSetCacheConfig, hipCtxSynchronize , hipCtxGetDevice
*/
hipError_t hipCtxSetCurrent(hipCtx_t ctx);
/**
@@ -1297,9 +1292,10 @@ hipError_t hipCtxSetCurrent(hipCtx_t ctx);
*
* @param [out] ctx
*
* @returns #hipSuccess
* @returns #hipSuccess, #hipErrorInvalidContext
*
* @see hipCtxCreate, hipCtxDestroy, hipCtxGetDevice, hipCtxGetFlags, hipCtxPopCurrent, hipCtxPushCurrent, hipCtxSetCacheConfig, hipCtxSynchronize, hipCtxGetDevice
*/
hipError_t hipCtxGetCurrent(hipCtx_t* ctx);
/**
@@ -1308,6 +1304,8 @@ hipError_t hipCtxGetCurrent(hipCtx_t* ctx);
* @param [out] device
*
* @returns #hipSuccess, #hipErrorInvalidContext
*
* @see hipCtxCreate, hipCtxDestroy, hipCtxGetFlags, hipCtxPopCurrent, hipCtxGetCurrent, hipCtxPushCurrent, hipCtxSetCacheConfig, hipCtxSynchronize
*/
hipError_t hipCtxGetDevice(hipDevice_t *device);
@@ -1315,53 +1313,81 @@ hipError_t hipCtxGetDevice(hipDevice_t *device);
/**
* @brief Returns the approximate HIP api version.
*
* @param [in] ctx Context to check
* @param [out] apiVersion
*
* @return #hipSuccess
*
* @warning The HIP feature set does not correspond to an exact CUDA SDK api revision.
* This function always set *apiVersion to 4 as an approximation though HIP supports
* some features which were introduced in later CUDA SDK revisions.
* HIP apps code should not rely on the api revision number here and should
* use arch feature flags to test device capabilities or conditional compilation.
*
* @see hipCtxCreate, hipCtxDestroy, hipCtxGetDevice, hipCtxGetFlags, hipCtxPopCurrent, hipCtxPushCurrent, hipCtxSetCacheConfig, hipCtxSynchronize, hipCtxGetDevice
*/
hipError_t hipCtxGetApiVersion (hipCtx_t ctx,int *apiVersion);
/**
* @brief Set Cache configuration for a specific function
*
* Note: AMD devices and recent Nvidia GPUS do not support reconfigurable cache. This hint is ignored on those architectures.
* @param [out] cacheConfiguration
*
* @return #hipSuccess
*
* @warning AMD devices and recent Nvidia GPUS do not support reconfigurable cache. This hint is ignored on those architectures.
*
* @see hipCtxCreate, hipCtxDestroy, hipCtxGetFlags, hipCtxPopCurrent, hipCtxGetCurrent, hipCtxSetCurrent, hipCtxPushCurrent, hipCtxSetCacheConfig, hipCtxSynchronize, hipCtxGetDevice
*/
hipError_t hipCtxGetCacheConfig ( hipFuncCache *cacheConfig );
/**
* @brief Set L1/Shared cache partition.
*
* @param [in] cacheConfiguration
*
* Note: AMD devices and recent Nvidia GPUS do not support reconfigurable cache. This hint is ignored on those architectures.
* @return #hipSuccess
*
* @warning AMD devices and recent Nvidia GPUS do not support reconfigurable cache. This hint is ignored on those architectures.
*
* @see hipCtxCreate, hipCtxDestroy, hipCtxGetFlags, hipCtxPopCurrent, hipCtxGetCurrent, hipCtxSetCurrent, hipCtxPushCurrent, hipCtxSetCacheConfig, hipCtxSynchronize, hipCtxGetDevice
*/
hipError_t hipCtxSetCacheConfig ( hipFuncCache cacheConfig );
/**
* @brief Set Shared memory bank configuration.
*
* Note: AMD devices and recent Nvidia GPUS do not support shared cache banking, and the hint is ignored on those architectures.
* @param [in] sharedMemoryConfiguration
*
* @return #hipSuccess
*
* @warning AMD devices and recent Nvidia GPUS do not support shared cache banking, and the hint is ignored on those architectures.
*
* @see hipCtxCreate, hipCtxDestroy, hipCtxGetFlags, hipCtxPopCurrent, hipCtxGetCurrent, hipCtxSetCurrent, hipCtxPushCurrent, hipCtxSetCacheConfig, hipCtxSynchronize, hipCtxGetDevice
*/
hipError_t hipCtxSetSharedMemConfig ( hipSharedMemConfig config );
/**
* @brief Get Shared memory bank configuration.
*
* Note: AMD devices and recent Nvidia GPUS do not support shared cache banking, and the hint is ignored on those architectures.
* @param [out] sharedMemoryConfiguration
*
* @return #hipSuccess
*
* @warning AMD devices and recent Nvidia GPUS do not support shared cache banking, and the hint is ignored on those architectures.
*
* @see hipCtxCreate, hipCtxDestroy, hipCtxGetFlags, hipCtxPopCurrent, hipCtxGetCurrent, hipCtxSetCurrent, hipCtxPushCurrent, hipCtxSetCacheConfig, hipCtxSynchronize, hipCtxGetDevice
*/
hipError_t hipCtxGetSharedMemConfig ( hipSharedMemConfig * pConfig );
/**
* @brief Blocks until the default context has completed all preceding requested tasks.
*
* This function waits for all streams on the default context to complete execution, and then returns.
* @return #hipSuccess
*
* @returns #hipSuccess.
* @warning This function waits for all streams on the default context to complete execution, and then returns.
*
* @see hipCtxCreate, hipCtxDestroy, hipCtxGetFlags, hipCtxPopCurrent, hipCtxGetCurrent, hipCtxSetCurrent, hipCtxPushCurrent, hipCtxSetCacheConfig, hipCtxGetDevice
*/
hipError_t hipCtxSynchronize ( void );
@@ -1370,7 +1396,9 @@ hipError_t hipCtxSynchronize ( void );
*
* @param [out] flags
*
* @returns #hipSuccess.
* @returns #hipSuccess
*
* @see hipCtxCreate, hipCtxDestroy, hipCtxPopCurrent, hipCtxGetCurrent, hipCtxGetCurrent, hipCtxSetCurrent, hipCtxPushCurrent, hipCtxSetCacheConfig, hipCtxSynchronize, hipCtxGetDevice
*/
hipError_t hipCtxGetFlags ( unsigned int* flags );
@@ -1385,8 +1413,9 @@ hipError_t hipCtxGetFlags ( unsigned int* flags );
* @param [in] peerCtx
* @param [in] flags
*
* Returns #hipSuccess, #hipErrorInvalidDevice, #hipErrorInvalidValue,
* @returns #hipErrorPeerAccessAlreadyEnabled if peer access is already enabled for this device.
* @returns #hipSuccess, #hipErrorInvalidDevice, #hipErrorInvalidValue, #hipErrorPeerAccessAlreadyEnabled
*
* @see hipCtxCreate, hipCtxDestroy, hipCtxGetFlags, hipCtxPopCurrent, hipCtxGetCurrent, hipCtxSetCurrent, hipCtxPushCurrent, hipCtxSetCacheConfig, hipCtxSynchronize, hipCtxGetDevice
* @warning PeerToPeer support is experimental.
*/
hipError_t hipCtxEnablePeerAccess (hipCtx_t peerCtx, unsigned int flags);
@@ -1399,9 +1428,12 @@ hipError_t hipCtxEnablePeerAccess (hipCtx_t peerCtx, unsigned int flags);
* @param [in] peerCtx
*
* @returns #hipSuccess, #hipErrorPeerAccessNotEnabled
*
* @see hipCtxCreate, hipCtxDestroy, hipCtxGetFlags, hipCtxPopCurrent, hipCtxGetCurrent, hipCtxSetCurrent, hipCtxPushCurrent, hipCtxSetCacheConfig, hipCtxSynchronize, hipCtxGetDevice
* @warning PeerToPeer support is experimental.
*/
hipError_t hipCtxDisablePeerAccess (hipCtx_t peerCtx);
// doxygen end Context Management
/**
* @}
@@ -1457,6 +1489,10 @@ hipError_t hipDeviceTotalMem (size_t *bytes,hipDevice_t device);
/**
* @brief Returns the approximate HIP driver version.
*
* @param [out] driverVersion
*
* @returns #hipSuccess, #hipErrorInavlidValue
*
* @warning The HIP feature set does not correspond to an exact CUDA SDK driver revision.
* This function always set *driverVersion to 4 as an approximation though HIP supports
@@ -1464,9 +1500,24 @@ hipError_t hipDeviceTotalMem (size_t *bytes,hipDevice_t device);
* HIP apps code should not rely on the driver revision number here and should
* use arch feature flags to test device capabilities or conditional compilation.
*
* @see hipRuntimeGetVersion
*/
hipError_t hipDriverGetVersion(int *driverVersion) ;
/**
* @brief Returns the approximate HIP Runtime version.
*
* @param [out] runtimeVersion
*
* @returns #hipSuccess, #hipErrorInavlidValue
*
* @warning On HIP/HCC path this function returns HIP runtime patch version however on
* HIP/NVCC path this function return CUDA runtime version.
*
* @see hipDriverGetVersion
*/
hipError_t hipRuntimeGetVersion(int *runtimeVersion) ;
/**
* @brief Loads code object from file into a hipModule_t
*
hipamd/include/hcc_detail/hip_texture.h
+72 -10
View File
@@ -97,18 +97,78 @@ typedef struct hipArray {
#define tex2D(_tex, _dx, _dy) \
_tex._dataPtr[(unsigned int)_dx + (unsigned int)_dy*(_tex.width)]
/**
* @brief Allocate an array on the device.
*
* @param[out] array Pointer to allocated array in device memory
* @param[in] desc Requested channel format
* @param[in] width Requested array allocation width
* @param[in] height Requested array allocation height
* @param[in] flags Requested properties of allocated array
* @return #hipSuccess, #hipErrorMemoryAllocation
*
* @see hipMalloc, hipMallocPitch, hipFree, hipFreeArray, hipHostMalloc, hipHostFree
*/
hipError_t hipMallocArray(hipArray** array, const hipChannelFormatDesc* desc,
size_t width, size_t height = 0, unsigned int flags = 0);
/**
* @brief Frees an array on the device.
*
* @param[in] array Pointer to array to free
* @return #hipSuccess, #hipErrorInvalidValue, #hipErrorInitializationError
*
* @see hipMalloc, hipMallocPitch, hipFree, hipMallocArray, hipHostMalloc, hipHostFree
*/
hipError_t hipFreeArray(hipArray* array);
//
// dpitch, spitch, and width in bytes
/**
* @brief Copies data between host and device.
*
* @param[in] dst Destination memory address
* @param[in] dpitch Pitch of destination memory
* @param[in] src Source memory address
* @param[in] spitch Pitch of source memory
* @param[in] width Width of matrix transfer (columns in bytes)
* @param[in] height Height of matrix transfer (rows)
* @param[in] kind Type of transfer
* @return #hipSuccess, #hipErrorInvalidValue, #hipErrorInvalidPitchValue, #hipErrorInvalidDevicePointer, #hipErrorInvalidMemcpyDirection
*
* @see hipMemcpy, hipMemcpyToArray, hipMemcpy2DToArray, hipMemcpyFromArray, hipMemcpyToSymbol, hipMemcpyAsync
*/
hipError_t hipMemcpy2D(void* dst, size_t dpitch, const void* src, size_t spitch, size_t width, size_t height, hipMemcpyKind kind);
// wOffset, width, and spitch in bytes
/**
* @brief Copies data between host and device.
*
* @param[in] dst Destination memory address
* @param[in] dpitch Pitch of destination memory
* @param[in] src Source memory address
* @param[in] spitch Pitch of source memory
* @param[in] width Width of matrix transfer (columns in bytes)
* @param[in] height Height of matrix transfer (rows)
* @param[in] kind Type of transfer
* @return #hipSuccess, #hipErrorInvalidValue, #hipErrorInvalidPitchValue, #hipErrorInvalidDevicePointer, #hipErrorInvalidMemcpyDirection
*
* @see hipMemcpy, hipMemcpyToArray, hipMemcpy2D, hipMemcpyFromArray, hipMemcpyToSymbol, hipMemcpyAsync
*/
hipError_t hipMemcpy2DToArray(hipArray* dst, size_t wOffset, size_t hOffset, const void* src,
size_t spitch, size_t width, size_t height, hipMemcpyKind kind);
/**
* @brief Copies data between host and device.
*
* @param[in] dst Destination memory address
* @param[in] dpitch Pitch of destination memory
* @param[in] src Source memory address
* @param[in] spitch Pitch of source memory
* @param[in] width Width of matrix transfer (columns in bytes)
* @param[in] height Height of matrix transfer (rows)
* @param[in] kind Type of transfer
* @return #hipSuccess, #hipErrorInvalidValue, #hipErrorInvalidPitchValue, #hipErrorInvalidDevicePointer, #hipErrorInvalidMemcpyDirection
*
* @see hipMemcpy, hipMemcpy2DToArray, hipMemcpy2D, hipMemcpyFromArray, hipMemcpyToSymbol, hipMemcpyAsync
*/
hipError_t hipMemcpyToArray(hipArray* dst, size_t wOffset, size_t hOffset,
const void* src, size_t count, hipMemcpyKind kind);
@@ -152,15 +212,17 @@ hipChannelFormatDesc hipBindTexture(size_t *offset, struct textureReference *te
}
#endif
/*
* @brief Returns a channel descriptor with format f and number of bits of each ocmponent x,y,z and w.
/**
* @brief Returns a channel descriptor using the specified format.
*
* @par Parameters
* None.
* @return Channel descriptor
* @param[in] x X component
* @param[in] y Y component
* @param[in] z Z component
* @param[in] w W component
* @param[in] f Channel format
* @return Channel descriptor with format f
*
*
**/
*/
hipChannelFormatDesc hipCreateChannelDesc(int x, int y, int z, int w, hipChannelFormatKind f);
// descriptors
hipamd/include/hcc_detail/hip_vector_types.h
+147 -2
View File
@@ -32,11 +32,11 @@ THE SOFTWARE.
#error("This version of HIP requires a newer version of HCC.");
#endif
#if __cplusplus
#if __HCC__
#include <hc_short_vector.hpp>
using namespace hc::short_vector;
#endif
//-- Signed
// Define char vector types
@@ -115,6 +115,151 @@ typedef hc::short_vector::double2 double2;
typedef hc::short_vector::double3 double3;
typedef hc::short_vector::double4 double4;
#else
#define __hip_align(name, val, data) \
__attribute__((aligned(val))) name \
{ data }
struct __hip_align(char1, 1, signed char x;);
struct __hip_align(uchar1, 1, unsigned char x;);
struct __hip_align(char2, 2, signed char x; signed char y;);
struct __hip_align(uchar2, 2, unsigned char x; unsigned char y;);
struct char3
{
signed char x, y, z;
};
struct uchar3
{
unsigned char x, y, z;
};
struct __hip_align(char4, 4, signed char x; signed char y; signed char z; signed char w;);
struct __hip_align(uchar4, 4, unsigned char x; unsigned char y; unsigned char z; unsigned char w;);
struct __hip_align(short1, 2, signed short x;);
struct __hip_align(ushort1, 2, unsigned short x;);
struct __hip_align(short2, 4, signed short x; signed short y;);
struct __hip_align(ushort2, 4, unsigned short x; unsigned short y;);
struct short3
{
signed short x, y, z;
};
struct ushort3
{
unsigned short x, y, z;
};
struct __hip_align(short4, 8, signed short x; signed short y; signed short z; signed short w;);
struct __hip_align(ushort4, 8, unsigned short x; unsigned short y; unsigned short z; unsigned short w;);
struct __hip_align(int1, 4, signed int x;);
struct __hip_align(uint1, 4, unsigned int x;);
struct __hip_align(int2, 8, signed int x; signed int y;);
struct __hip_align(uint2, 8, unsigned int x; unsigned int y;);
struct int3{
signed int x, y, z;
};
struct uint3{
unsigned int x, y, z;
};
struct __hip_align(int4, 16, signed int x; signed int y; signed int z; signed int w;);
struct __hip_align(uint4, 16, unsigned int x; unsigned int y; unsigned int z; unsigned int w;);
struct __hip_align(long1, 8, long int x;);
struct __hip_align(ulong1, 8, unsigned long x;);
struct __hip_align(long2, 16, long int x; long int y;);
struct __hip_align(ulong2, 16, unsigned long x; unsigned long y;);
struct long3{
long int x, y, z;
};
struct ulong3{
unsigned long x, y, z;
};
struct __hip_align(long4, 32, long int x; long int y; long int z; long int w;);
struct __hip_align(ulong4, 32, unsigned long x; unsigned long y; unsigned long z; unsigned long w;);
struct float1
{
float x;
};
struct __hip_align(float2, 8, float x; float y;);
struct float3
{
float x, y, z;
};
struct __hip_align(float4, 16, float x; float y; float z; float w;);
struct __hip_align(longlong1, 16, long long int x;);
struct __hip_align(ulonglong1, 16, unsigned long long int x;);
struct __attribute__((aligned(32))) longlong2
{
long long int x, y;
};
struct __attribute__((aligned(32))) ulonglong2
{
unsigned long long int x, y;
};
struct longlong3
{
long long int x, y, z;
};
struct ulonglong3
{
unsigned long long int x, y, z;
};
struct __attribute__((aligned(64))) longlong4
{
long long int x, y, z, w;
};
struct __attribute__((aligned(64))) ulonglong4
{
unsigned long long int x, y, z, w;
};
struct double1
{
double x;
};
struct __attribute__((aligned(16))) double2
{
double x, y;
};
struct double3
{
double x, y, z;
};
struct __attribute__((aligned(32))) double4
{
double x, y, z, w;
};
#endif
#if __HCC__
#include"hip/hcc_detail/host_defines.h"
#define __HIP_DEVICE__ __device__ __host__
hipamd/include/hcc_detail/unpinned_copy_engine.h
-71
View File
@@ -1,71 +0,0 @@
/*
Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved.
Permission is hereby granted, free of charge, to any person obtaining a copy
of this software and associated documentation files (the "Software"), to deal
in the Software without restriction, including without limitation the rights
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
copies of the Software, and to permit persons to whom the Software is
furnished to do so, subject to the following conditions:
The above copyright notice and this permission notice shall be included in
all copies or substantial portions of the Software.
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANNTY OF ANY KIND, EXPRESS OR
IMPLIED, INNCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANNY CLAIM, DAMAGES OR OTHER
LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
//#pragma once
#ifndef STAGING_BUFFER_H
#define STAGING_BUFFER_H
#include "hsa.h"
//-------------------------------------------------------------------------------------------------
// An optimized "staging buffer" used to implement Host-To-Device and Device-To-Host copies.
// Some GPUs may not be able to directly access host memory, and in these cases we need to
// stage the copy through a pinned staging buffer. For example, the CopyHostToDevice
// uses the CPU to copy to a pinned "staging buffer", and then use the GPU DMA engine to copy
// from the staging buffer to the final destination. The copy is broken into buffer-sized chunks
// to limit the size of the buffer and also to provide better performance by overlapping the CPU copies
// with the DMA copies.
//
// PinInPlace is another algorithm which pins the host memory "in-place", and copies it with the DMA
// engine. This routine is under development.
//
// Staging buffer provides thread-safe access via a mutex.
struct UnpinnedCopyEngine {
static const int _max_buffers = 4;
UnpinnedCopyEngine(hsa_agent_t hsaAgent,hsa_agent_t cpuAgent, size_t bufferSize, int numBuffers,int thresholdH2D_directStaging,int thresholdH2D_stagingPinInPlace,int thresholdD2H) ;
~UnpinnedCopyEngine();
void CopyHostToDevice(int tempIndex,int isLargeBar,void* dst, const void* src, size_t sizeBytes, hsa_signal_t *waitFor);
void CopyHostToDevicePinInPlace(void* dst, const void* src, size_t sizeBytes, hsa_signal_t *waitFor);
void CopyDeviceToHost (int tempIndex,void* dst, const void* src, size_t sizeBytes, hsa_signal_t *waitFor);
void CopyDeviceToHostPinInPlace(void* dst, const void* src, size_t sizeBytes, hsa_signal_t *waitFor);
void CopyPeerToPeer( void* dst, hsa_agent_t dstAgent, const void* src, hsa_agent_t srcAgent, size_t sizeBytes, hsa_signal_t *waitFor);
private:
hsa_agent_t _hsaAgent;
hsa_agent_t _cpuAgent;
size_t _bufferSize; // Size of the buffers.
int _numBuffers;
char *_pinnedStagingBuffer[_max_buffers];
hsa_signal_t _completion_signal[_max_buffers];
hsa_signal_t _completion_signal2[_max_buffers]; // P2P needs another set of signals.
std::mutex _copy_lock; // provide thread-safe access
int _hipH2DTransferThresholdDirectOrStaging;
int _hipH2DTransferThresholdStagingOrPininplace;
int _hipD2HTransferThreshold;
};
#endif
hipamd/include/nvcc_detail/hip_runtime_api.h
+49 -25
View File
@@ -50,6 +50,12 @@ hipMemcpyHostToHost
} hipTextureFilterMode;*/
#define hipFilterModePoint cudaFilterModePoint
//! Flags that can be used with hipEventCreateWithFlags:
#define hipEventDefault cudaEventDefault
#define hipEventBlockingSync cudaEventBlockingSync
#define hipEventDisableTiming cudaEventDisableTiming
#define hipEventInterprocess cudaEventInterprocess
#define hipHostMallocDefault cudaHostAllocDefault
#define hipHostMallocPortable cudaHostAllocPortable
#define hipHostMallocMapped cudaHostAllocMapped
@@ -64,6 +70,8 @@ hipMemcpyHostToHost
typedef cudaEvent_t hipEvent_t;
typedef cudaStream_t hipStream_t;
typedef cudaIpcEventHandle_t hipIpcEventHandle_t;
typedef cudaIpcMemHandle_t hipIpcMemHandle_t;
typedef CUcontext hipCtx_t;
typedef CUsharedconfig hipSharedMemConfig;
typedef CUfunc_cache hipFuncCache;
@@ -138,15 +146,17 @@ switch(hError) {
}
inline static cudaMemcpyKind hipMemcpyKindToCudaMemcpyKind(hipMemcpyKind kind) {
switch(kind) {
case hipMemcpyHostToHost:
return cudaMemcpyHostToHost;
case hipMemcpyHostToDevice:
return cudaMemcpyHostToDevice;
case hipMemcpyDeviceToHost:
return cudaMemcpyDeviceToHost;
default:
return cudaMemcpyDefault;
switch(kind) {
case hipMemcpyHostToHost:
return cudaMemcpyHostToHost;
case hipMemcpyHostToDevice:
return cudaMemcpyHostToDevice;
case hipMemcpyDeviceToHost:
return cudaMemcpyDeviceToHost;
case hipMemcpyDeviceToDevice:
return cudaMemcpyDeviceToDevice;
default:
return cudaMemcpyDefault;
}
}
@@ -175,12 +185,6 @@ inline static hipError_t hipFree(void* ptr) {
return hipCUDAErrorTohipError(cudaFree(ptr));
}
inline static hipError_t hipMallocHost(void** ptr, size_t size) __attribute__((deprecated("use hipHostMalloc instead")));
inline static hipError_t hipMallocHost(void** ptr, size_t size) {
return hipCUDAErrorTohipError(cudaMallocHost(ptr, size));
}
inline static hipError_t hipHostMalloc(void** ptr, size_t size, unsigned int flags){
return hipCUDAErrorTohipError(cudaHostAlloc(ptr, size, flags));
}
@@ -201,11 +205,6 @@ inline static hipError_t hipHostUnregister(void* ptr){
return hipCUDAErrorTohipError(cudaHostUnregister(ptr));
}
inline static hipError_t hipFreeHost(void* ptr) __attribute__((deprecated("use hipHostFree instead")));
inline static hipError_t hipFreeHost(void* ptr) {
return hipCUDAErrorTohipError(cudaFreeHost(ptr));
}
inline static hipError_t hipHostFree(void* ptr) {
return hipCUDAErrorTohipError(cudaFreeHost(ptr));
}
@@ -237,13 +236,13 @@ inline static hipError_t hipChooseDevice( int* device, const hipDeviceProp_t* pr
return hipCUDAErrorTohipError(cudaChooseDevice(device,&cdprop));
}
inline static hipError_t hipMemcpyHtoD(hipDeviceptr_t dst,
inline static hipError_t hipMemcpyHtoD(hipDeviceptr_t dst,
void* src, size_t size)
{
return hipCUResultTohipError(cuMemcpyHtoD(dst, src, size));
}
inline static hipError_t hipMemcpyDtoH(void* dst,
inline static hipError_t hipMemcpyDtoH(void* dst,
hipDeviceptr_t src, size_t size)
{
return hipCUResultTohipError(cuMemcpyDtoH(dst, src, size));
@@ -255,13 +254,13 @@ inline static hipError_t hipMemcpyDtoD(hipDeviceptr_t dst,
return hipCUResultTohipError(cuMemcpyDtoD(dst, src, size));
}
inline static hipError_t hipMemcpyHtoDAsync(hipDeviceptr_t dst,
inline static hipError_t hipMemcpyHtoDAsync(hipDeviceptr_t dst,
void* src, size_t size, hipStream_t stream)
{
return hipCUResultTohipError(cuMemcpyHtoDAsync(dst, src, size, stream));
}
inline static hipError_t hipMemcpyDtoHAsync(void* dst,
inline static hipError_t hipMemcpyDtoHAsync(void* dst,
hipDeviceptr_t src, size_t size, hipStream_t stream)
{
return hipCUResultTohipError(cuMemcpyDtoH(dst, src, size));
@@ -306,6 +305,26 @@ inline static hipError_t hipGetDevice(int * device){
return hipCUDAErrorTohipError(cudaGetDevice(device));
}
inline static hipError_t hipIpcCloseMemHandle(void *devPtr){
return hipCUDAErrorTohipError(cudaIpcCloseMemHandle(devPtr));
}
inline static hipError_t hipIpcGetEventHandle(hipIpcEventHandle_t* handle, hipEvent_t event){
return hipCUDAErrorTohipError(cudaIpcGetEventHandle(handle, event));
}
inline static hipError_t hipIpcGetMemHandle(hipIpcMemHandle_t* handle, void* devPtr){
return hipCUDAErrorTohipError(cudaIpcGetMemHandle(handle, devPtr));
}
inline static hipError_t hipIpcOpenEventHandle(hipEvent_t* event, hipIpcEventHandle_t handle){
return hipCUDAErrorTohipError(cudaIpcOpenEventHandle(event, handle));
}
inline static hipError_t hipIpcOpenMemHandle(void** devPtr, hipIpcMemHandle_t handle, unsigned int flags){
return hipCUDAErrorTohipError(cudaIpcOpenMemHandle(devPtr, handle, flags));
}
inline static hipError_t hipMemset(void* devPtr,int value, size_t count) {
return hipCUDAErrorTohipError(cudaMemset(devPtr, value, count));
}
@@ -546,6 +565,11 @@ inline static hipError_t hipDriverGetVersion(int *driverVersion)
return hipCUDAErrorTohipError(err);
}
inline static hipError_t hipRuntimeGetVersion(int *runtimeVersion)
{
return hipCUDAErrorTohipError(cudaRuntimeGetVersion(runtimeVersion));
}
inline static hipError_t hipDeviceCanAccessPeer ( int* canAccessPeer, int device, int peerDevice )
{
return hipCUDAErrorTohipError(cudaDeviceCanAccessPeer(canAccessPeer, device, peerDevice));
@@ -729,7 +753,7 @@ inline static hipError_t hipModuleLaunchKernel(hipFunction_t f,
unsigned int sharedMemBytes, hipStream_t stream,
void **kernelParams, void **extra)
{
return hipCUResultTohipError(cuLaunchKernel(f,
return hipCUResultTohipError(cuLaunchKernel(f,
gridDimX, gridDimY, gridDimZ,
blockDimX, blockDimY, blockDimZ,
sharedMemBytes, stream, kernelParams, extra));
hipamd/samples/1_Utils/hipBusBandwidth/ResultDatabase.cpp
+11 -3
View File
@@ -7,16 +7,22 @@
using namespace std;
#define SORT_RETAIN_ATTS_ORDER 1
bool ResultDatabase::Result::operator<(const Result &rhs) const
{
if (test < rhs.test)
return true;
if (test > rhs.test)
return false;
#if (SORT_RETAIN_ATTS_ORDER == 0)
// For ties, sort by the value of the attribute:
if (atts < rhs.atts)
return true;
if (atts > rhs.atts)
return false;
#endif
return false; // less-operator returns false on equal
}
@@ -189,7 +195,8 @@ void ResultDatabase::AddResult(const string &test_orig,
void ResultDatabase::DumpDetailed(ostream &out)
{
vector<Result> sorted(results);
sort(sorted.begin(), sorted.end());
stable_sort(sorted.begin(), sorted.end());
const int testNameW = 24 ;
const int attW = 12;
@@ -281,7 +288,8 @@ void ResultDatabase::DumpDetailed(ostream &out)
void ResultDatabase::DumpSummary(ostream &out)
{
vector<Result> sorted(results);
sort(sorted.begin(), sorted.end());
stable_sort(sorted.begin(), sorted.end());
const int testNameW = 24 ;
const int attW = 12;
@@ -377,7 +385,7 @@ void ResultDatabase::DumpCsv(string fileName)
bool emptyFile;
vector<Result> sorted(results);
sort(sorted.begin(), sorted.end());
stable_sort(sorted.begin(), sorted.end());
//Check to see if the file is empty - if so, add the headers
emptyFile = this->IsFileEmpty(fileName);
hipamd/samples/1_Utils/hipBusBandwidth/hipBusBandwidth.cpp
+2 -2
View File
@@ -49,8 +49,8 @@ std::string sizeToString(int size)
using namespace std;
stringstream ss;
if (size < 0) {
// char (09, horiz tab) lexically sorts before " " so will cause Byte values to be displayed before kB.
ss << char(0x09)/*tab*/ << setfill('0') << setw(3) << -size << "B";
// char (-) lexically sorts before " " so will cause Byte values to be displayed before kB.
ss << "+" << setfill('0') << setw(3) << -size << "By";
} else {
ss << size << "kB";
}
hipamd/samples/1_Utils/hipDispatchLatency/Makefile
+5 -3
View File
@@ -6,10 +6,12 @@ HIPCC=$(HIP_PATH)/bin/hipcc
EXE=hipDispatchLatency
all: install
CXXFLAGS = -O3
$(EXE): hipDispatchLatency.cpp
$(HIPCC) hipDispatchLatency.cpp ResultDatabase.cpp -o $@
all: ${EXE}
$(EXE): hipDispatchLatency.cpp ResultDatabase.cpp
$(HIPCC) $(CXXFLAGS) hipDispatchLatency.cpp ResultDatabase.cpp -o $@
install: $(EXE)
cp $(EXE) $(HIP_PATH)/bin
hipamd/samples/1_Utils/hipDispatchLatency/ResultDatabase.cpp
+9 -5
View File
@@ -253,10 +253,12 @@ void ResultDatabase::DumpDetailed(ostream &out)
out << endl;
}
out << endl
<< "Note: Any results marked with (*) had missing values." << endl
<< " This can occur on systems with a mixture of" << endl
<< " device types or architectural capabilities." << endl;
if (0) {
out << endl
<< "Note: Any results marked with (*) had missing values." << endl
<< " This can occur on systems with a mixture of" << endl
<< " device types or architectural capabilities." << endl;
}
}
@@ -330,9 +332,11 @@ void ResultDatabase::DumpSummary(ostream &out)
out << endl;
}
out << endl
if (0) {
out << endl
<< "Note: results marked with (*) had missing values such as" << endl
<< "might occur with a mixture of architectural capabilities." << endl;
}
}
// ****************************************************************************
hipamd/samples/1_Utils/hipDispatchLatency/hipDispatchLatency.cpp
+23 -16
View File
@@ -30,15 +30,22 @@ if(status != hipSuccess){ \
#define LEN 1024*1024
#define SIZE LEN * sizeof(float)
#define ITER 5120
#define ITER 10120
__global__ void One(hipLaunchParm lp, float* Ad){
// HCC optimizes away fully NULL kernel calls, so run one that is nearly null:
__global__ void NearlyNull(hipLaunchParm lp, float* Ad){
if (Ad) {
Ad[0] = 42;
}
}
int main(){
hipError_t err;
float *A, *Ad;
float *A;
float *Ad = NULL;
A = new float[LEN];
@@ -50,11 +57,10 @@ int main(){
err = hipStreamCreate(&stream);
check("Creating stream",err);
err = hipMalloc(&Ad, SIZE);
check("Allocating Ad memory on device", err);
err = hipMemcpy(Ad, A, SIZE, hipMemcpyHostToDevice);
check("Doing memory copy from A to Ad", err);
//err = hipMalloc(&Ad, SIZE);
//check("Allocating Ad memory on device", err);
//err = hipMemcpy(Ad, A, SIZE, hipMemcpyHostToDevice);
//check("Doing memory copy from A to Ad", err);
float mS = 0;
hipEvent_t start, stop;
@@ -63,15 +69,16 @@ int main(){
ResultDatabase resultDB[8];
hipEventRecord(start);
hipLaunchKernel(HIP_KERNEL_NAME(One), dim3(LEN/512), dim3(512), 0, 0, Ad);
hipLaunchKernel(NearlyNull, dim3(LEN/512), dim3(512), 0, 0, Ad);
hipEventRecord(stop);
hipEventElapsedTime(&mS, start, stop);
resultDB[0].AddResult(std::string("First Kernel Launch"), "", "uS", mS*1000);
// std::cout<<"First Kernel Launch: \t\t"<<mS*1000<<" uS"<<std::endl;
resultDB[0].DumpSummary(std::cout);
hipEventRecord(start);
hipLaunchKernel(HIP_KERNEL_NAME(One), dim3(LEN/512), dim3(512), 0, 0, Ad);
hipLaunchKernel(NearlyNull, dim3(LEN/512), dim3(512), 0, 0, Ad);
hipEventRecord(stop);
hipEventElapsedTime(&mS, start, stop);
resultDB[1].AddResult(std::string("Second Kernel Launch"), "", "uS", mS*1000);
@@ -79,7 +86,7 @@ int main(){
resultDB[1].DumpSummary(std::cout);
hipEventRecord(start);
for(int i=0;i<ITER;i++){
hipLaunchKernel(HIP_KERNEL_NAME(One), dim3(LEN/512), dim3(512), 0, 0, Ad);
hipLaunchKernel(NearlyNull, dim3(LEN/512), dim3(512), 0, 0, Ad);
}
hipDeviceSynchronize();
hipEventRecord(stop);
@@ -91,7 +98,7 @@ int main(){
hipEventRecord(start);
for(int i=0;i<ITER;i++){
hipLaunchKernel(HIP_KERNEL_NAME(One), dim3(LEN/512), dim3(512), 0, 0, Ad);
hipLaunchKernel(NearlyNull, dim3(LEN/512), dim3(512), 0, 0, Ad);
}
hipEventRecord(stop);
hipDeviceSynchronize();
@@ -103,7 +110,7 @@ int main(){
hipEventRecord(start);
for(int i=0;i<ITER;i++){
hipLaunchKernel(HIP_KERNEL_NAME(One), dim3(LEN/512), dim3(512), 0, stream, Ad);
hipLaunchKernel(NearlyNull, dim3(LEN/512), dim3(512), 0, stream, Ad);
hipDeviceSynchronize();
}
hipEventRecord(stop);
@@ -114,7 +121,7 @@ int main(){
hipDeviceSynchronize();
hipEventRecord(start);
for(int i=0;i<ITER;i++){
hipLaunchKernel(HIP_KERNEL_NAME(One), dim3(LEN/512), dim3(512), 0, stream, Ad);
hipLaunchKernel(NearlyNull, dim3(LEN/512), dim3(512), 0, stream, Ad);
}
hipDeviceSynchronize();
hipEventRecord(stop);
@@ -126,7 +133,7 @@ int main(){
hipEventRecord(start);
for(int i=0;i<ITER;i++){
hipLaunchKernel(HIP_KERNEL_NAME(One), dim3(LEN/512), dim3(512), 0, 0, Ad);
hipLaunchKernel(NearlyNull, dim3(LEN/512), dim3(512), 0, 0, Ad);
}
hipEventRecord(stop);
hipEventElapsedTime(&mS, start, stop);
@@ -137,7 +144,7 @@ int main(){
hipEventRecord(start);
for(int i=0;i<ITER;i++){
hipLaunchKernel(HIP_KERNEL_NAME(One), dim3(LEN/512), dim3(512), 0, stream, Ad);
hipLaunchKernel(NearlyNull, dim3(LEN/512), dim3(512), 0, stream, Ad);
}
hipEventRecord(stop);
hipEventElapsedTime(&mS, start, stop);
hipamd/samples/1_Utils/hipInfo/hipInfo.cpp
+1 -1
View File
@@ -65,7 +65,7 @@ double bytesToGB(size_t s)
void printDeviceProp (int deviceId)
{
using namespace std;
const int w1 = 30;
const int w1 = 34;
cout << left;
hipamd/samples/2_Cookbook/0_MatrixTranspose/Makefile
+36
View File
@@ -0,0 +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=./exe
.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
hipamd/samples/2_Cookbook/0_MatrixTranspose/MatrixTranspose.cpp
+137
View File
@@ -0,0 +1,137 @@
/*
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<iostream>
// hip header file
#include "hip_runtime.h"
#define WIDTH 1024
#define HEIGHT 1024
#define NUM (WIDTH*HEIGHT)
#define THREADS_PER_BLOCK_X 16
#define THREADS_PER_BLOCK_Y 16
#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,
const int height)
{
int x = hipBlockDim_x * hipBlockIdx_x + hipThreadIdx_x;
int y = hipBlockDim_y * hipBlockIdx_y + hipThreadIdx_y;
out[y * width + x] = in[x * height + y];
}
// CPU implementation of matrix transpose
void matrixTransposeCPUReference(
float * output,
float * input,
const unsigned int width,
const unsigned int height)
{
for(unsigned int j=0; j < height; j++)
{
for(unsigned int i=0; i < width; i++)
{
output[i*height + 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;
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));
// Memory transfer from host to device
hipMemcpy(gpuMatrix, Matrix, NUM*sizeof(float), hipMemcpyHostToDevice);
// Lauching kernel from host
hipLaunchKernel(matrixTranspose,
dim3(WIDTH/THREADS_PER_BLOCK_X, HEIGHT/THREADS_PER_BLOCK_Y),
dim3(THREADS_PER_BLOCK_X, THREADS_PER_BLOCK_Y),
0, 0,
gpuTransposeMatrix , gpuMatrix, WIDTH ,HEIGHT);
// Memory transfer from device to host
hipMemcpy(TransposeMatrix, gpuTransposeMatrix, NUM*sizeof(float), hipMemcpyDeviceToHost);
// CPU MatrixTranspose computation
matrixTransposeCPUReference(cpuTransposeMatrix, Matrix, WIDTH, HEIGHT);
// verify the results
errors = 0;
double eps = 1.0E-6;
for (i = 0; i < NUM; i++) {
if (std::abs(TransposeMatrix[i] - cpuTransposeMatrix[i]) > 0 ) {
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;
}
hipamd/samples/2_Cookbook/0_MatrixTranspose/Readme.md
+100
View File
@@ -0,0 +1,100 @@
## Writing first HIP program ###
This tutorial shows how to get write simple HIP application. We will write the simplest Matrix Transpose program.
## HIP Introduction:
HIP is a C++ runtime API and kernel language that allows developers to create portable applications that can run on AMD and other GPU’s. Our goal was to rise above the lowest-common-denominator paths and deliver a solution that allows you, the developer, to use essential hardware features and maximize your application’s performance on GPU hardware.
## Requirement:
For hardware requirement and software installation [Installation](https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP/INSTALL.md)
## prerequiste knowledge:
Programmers familiar with CUDA, OpenCL will be able to quickly learn and start coding with the HIP API. In case you are not, don't worry. You choose to start with the best one. We'll be explaining everything assuming you are completely new to gpgpu programming.
## Simple Matrix Transpose
Here is simple example showing how to write your first program in HIP.
In order to use the HIP framework, we need to add the "hip_runtime.h" header file. SInce its c++ api you can add any header file you have been using earlier while writing your c/c++ program. For gpgpu programming, we have host(microprocessor) and the device(gpu).
## Device-side code
We will work on device side code first, Here is simple example showing a snippet of HIP device side code:
`__global__ void matrixTranspose(hipLaunchParm lp, `
` float *out, `
` float *in, `
` const int width, `
` const int height) `
`{ `
` int x = hipBlockDim_x * hipBlockIdx_x + hipThreadIdx_x; `
` int y = hipBlockDim_y * hipBlockIdx_y + hipThreadIdx_y; `
` `
` out[y * width + x] = in[x * height + y]; `
`} `
`__global__` keyword is the Function-Type Qualifiers, it is used with functions that are executed on device and are called/launched from the hosts.
other function-type qualifiers are:
`__device__` functions are Executed on the device and Called from the device only
`__host__` functions are Executed on the host and Called from the host
`__host__` can combine with `__device__`, in which case the function compiles for both the host and device. These functions cannot use the HIP grid coordinate functions (for example, "hipThreadIdx_x", will talk about it latter). A possible workaround is to pass the necessary coordinate info as an argument to the function.
`__host__` cannot combine with `__global__`.
`__global__` functions are often referred to as *kernels, and calling one is termed *launching the kernel*.
Next keyword is `void`. HIP `__global__` functions must have a `void` return type, and the first parameter to a HIP `__global__` function must have the type `hipLaunchParm`, which is for execution configuration. Global functions require the caller to specify an "execution configuration" that includes the grid and block dimensions. The execution configuration can also include other information for the launch, such as the amount of additional shared memory to allocate and the stream where the kernel should execute.
After `hipLaunchParm`, Kernel arguments follows next(i.e., `float *out, float *in, const int width, const int height`).
The kernel function begins with
` int x = hipBlockDim_x * hipBlockIdx_x + hipThreadIdx_x;`
` int y = hipBlockDim_y * hipBlockIdx_y + hipThreadIdx_y;`
here the keyword hipBlockIdx_x, hipBlockIdx_y and hipBlockIdx_z(not used here) are the built-in functions to identify the threads in a block. The keyword hipBlockDim_x, hipBlockDim_y and hipBlockDim_z(not used here) are to identify the dimensions of the block.
We are familiar with rest of the code on device-side.
## Host-side code
Now, we'll see how to call the kernel from the host. Inside the main() function, we first defined the pointers(for both, the host-side as well as device). The declaration of device pointer is similar to that of the host. Next, we have `hipDeviceProp_t`, it is the pre-defined struct for hip device properties. This is followed by `hipGetDeviceProperties(&devProp, 0)` It is used to extract the device information. The first parameter is the struct, second parameter is the device number to get properties for. Next line print the name of the device.
We allocated memory to the Matrix on host side by using malloc and initiallized it. While in order to allocate memory on device side we will be using `hipMalloc`, it's quiet similar to that of malloc instruction. After this, we will copy the data to the allocated memory on device-side using `hipMemcpy`.
` hipMemcpy(gpuMatrix, Matrix, NUM*sizeof(float), hipMemcpyHostToDevice);`
here the first parameter is the destination pointer, second is the source pointer, third is the size of memory copy and the last specify the direction on memory copy(which is in this case froom host to device). While in order to transfer memory from device to host, use `hipMemcpyDeviceToHost` and for device to device memory copy use `hipMemcpyDeviceToDevice`.
Now, we'll see how to launch the kernel.
` hipLaunchKernel(matrixTranspose, `
` dim3(WIDTH/THREADS_PER_BLOCK_X, HEIGHT/THREADS_PER_BLOCK_Y), `
` dim3(THREADS_PER_BLOCK_X, THREADS_PER_BLOCK_Y), `
` 0, 0, `
` gpuTransposeMatrix , gpuMatrix, WIDTH ,HEIGHT); `
HIP introduces a standard C++ calling convention to pass the execution configuration to the kernel (this convention replaces the `Cuda <<< >>>` syntax). In HIP,
- Kernels launch with the `"hipLaunchKernel"` function
- The first five parameters to hipLaunchKernel are the following:
- **symbol kernelName**: the name of the kernel to launch. To support template kernels which contains "," use the HIP_KERNEL_NAME macro. In current application it's "matrixTranspose".
- **dim3 gridDim**: 3D-grid dimensions specifying the number of blocks to launch. In MatrixTranspose sample, it's "dim3(WIDTH/THREADS_PER_BLOCK_X, HEIGHT/THREADS_PER_BLOCK_Y)".
- **dim3 blockDim**: 3D-block dimensions specifying the number of threads in each block.In MatrixTranspose sample, it's "dim3(THREADS_PER_BLOCK_X, THREADS_PER_BLOCK_Y)".
- **size_t dynamicShared**: amount of additional shared memory to allocate when launching the kernel. In MatrixTranspose sample, it's '0'.
- **hipStream_t**: stream where the kernel should execute. A value of 0 corresponds to the NULL stream.In MatrixTranspose sample, it's '0'.
- Kernel arguments follow these first five parameters. Here, these are "gpuTransposeMatrix , gpuMatrix, WIDTH ,HEIGHT".
Next, we'll copy the computed values/data back to the device using the `hipMemcpy`. Here the last parameter will be `hipMemcpyDeviceToHost`
After, copying the data from device to memory, we will verify it with the one we computed with the cpu reference funtion.
Finally, we will free the memory allocated earlier by using free() for host while for devices we will use `hipFree`.
## How to build and run:
Use the make command and execute it using ./exe
Use hipcc to build the application, which is using hcc on AMD and nvcc on nvidia.
## More Info:
- [HIP FAQ](https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP/docs/markdown/hip_faq.md)
- [HIP Kernel Language](https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP/docs/markdown/hip_kernel_language.md)
- [HIP Runtime API (Doxygen)](http://gpuopen-professionalcompute-tools.github.io/HIP)
- [HIP Porting Guide](https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP/docs/markdown/hip_porting_guide.md)
- [HIP Terminology](https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP/docs/markdown/hip_terms.md) (including Rosetta Stone of GPU computing terms across CUDA/HIP/HC/AMP/OpenL)
- [clang-hipify](https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP/clang-hipify/README.md)
- [Developer/CONTRIBUTING Info](https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP/CONTRIBUTING.md)
- [Release Notes](https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP/RELEASE.md)
hipamd/samples/2_Cookbook/1_hipEvent/Makefile
+36
View File
@@ -0,0 +1,36 @@
HIP_PATH?= $(wildcard /opt/rocm/hip)
ifeq (,$(HIP_PATH))
HIP_PATH=../../..
endif
HIPCC=$(HIP_PATH)/bin/hipcc
TARGET=hcc
SOURCES = hipEvent.cpp
OBJECTS = $(SOURCES:.cpp=.o)
EXECUTABLE=./exe
.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
hipamd/samples/2_Cookbook/1_hipEvent/Readme.md
+74
View File
@@ -0,0 +1,74 @@
## Using hipEvents to measure performance ###
This tutorial is follow-up of the previous one where we learn how to write our first hip program, in which we compute Matrix Transpose. In this tutorial, we'll explain how to use the hipEvent to get the performance score for memory transfer and kernel execution time.
## Introduction:
Memory transfer and kernel execution are the most important parameter in parallel computing (specially HPC and machine learning). Memory bottlenecks is the main problem why we are not able to get the highest performance, therefore obtaining the memory transfer timing and kernel execution timing plays key role in application optimization.
## Requirement:
For hardware requirement and software installation [Installation](https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP/INSTALL.md)
## prerequiste knowledge:
Programmers familiar with CUDA, OpenCL will be able to quickly learn and start coding with the HIP API. In case you are not, don't worry. You choose to start with the best one. We'll be explaining everything assuming you are completely new to gpgpu programming.
## Simple Matrix Transpose
We will be using the Simple Matrix Transpose application from the previous tutorial and modify it to learn how to get the performance score for memory transfer and kernel execution time.
## hipEnvent_t
We'll learn how to use the event management functionality of HIP runtime api. In the same sourcecode, we used for MatrixTranspose we will declare the following events as follows:
` hipEvent_t start, stop;`
We'll create the event with the help of following code:
` hipEventCreate(&start);`
` hipEventCreate(&stop);`
We'll use the "eventMs" variable to store the time taken value:
` float eventMs = 1.0f;`
## Time taken measurement by using hipEvents:
We'll start the timer by calling:
` hipEventRecord(start, NULL);`
in this, the first parameter is the hipEvent_t, will will mark the start of the time from which the measurement has to be performed, while the second parameter has to be of the type hipStream_t. In current situation, we have passed NULL (the default stream). We will learn about the `hipStream_t` in more detail latter.
Now, we'll have the operation for which we need to compute the time taken. For the case of memory transfer, we'll place the `hipMemcpy`:
` hipMemcpy(gpuMatrix, Matrix, NUM*sizeof(float), hipMemcpyHostToDevice);`
and for kernel execution time we'll use `hipKernelLaunch`:
` hipLaunchKernel(matrixTranspose, `
` dim3(WIDTH/THREADS_PER_BLOCK_X, HEIGHT/THREADS_PER_BLOCK_Y), `
` dim3(THREADS_PER_BLOCK_X, THREADS_PER_BLOCK_Y), `
` 0, 0, `
` gpuTransposeMatrix , gpuMatrix, WIDTH ,HEIGHT); `
Now to mark the end of the eventRecord, we will again use the hipEventRecord by passing the stop event:
` hipEventRecord(stop, NULL);`
Will synchronize the event with the help of:
` hipEventSynchronize(stop);`
In order to calculate the time taken by measuring the difference of occurance marked by the start and stop event, we'll use:
` hipEventElapsedTime(&eventMs, start, stop);`
Here the first parameter will store the time taken value, second parameter is the starting marker for the event while the third one is marking the end.
We can print the value of time take comfortably since eventMs is float variable.
## How to build and run:
Use the make command and execute it using ./exe
Use hipcc to build the application, which is using hcc on AMD and nvcc on nvidia.
## More Info:
- [HIP FAQ](https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP/docs/markdown/hip_faq.md)
- [HIP Kernel Language](https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP/docs/markdown/hip_kernel_language.md)
- [HIP Runtime API (Doxygen)](http://gpuopen-professionalcompute-tools.github.io/HIP)
- [HIP Porting Guide](https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP/docs/markdown/hip_porting_guide.md)
- [HIP Terminology](https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP/docs/markdown/hip_terms.md) (including Rosetta Stone of GPU computing terms across CUDA/HIP/HC/AMP/OpenL)
- [clang-hipify](https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP/clang-hipify/README.md)
- [Developer/CONTRIBUTING Info](https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP/CONTRIBUTING.md)
- [Release Notes](https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP/RELEASE.md)
hipamd/samples/2_Cookbook/1_hipEvent/hipEvent.cpp
+174
View File
@@ -0,0 +1,174 @@
/*
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<iostream>
// hip header file
#include "hip_runtime.h"
#define WIDTH 1024
#define HEIGHT 1024
#define NUM (WIDTH*HEIGHT)
#define THREADS_PER_BLOCK_X 16
#define THREADS_PER_BLOCK_Y 16
#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,
const int height)
{
int x = hipBlockDim_x * hipBlockIdx_x + hipThreadIdx_x;
int y = hipBlockDim_y * hipBlockIdx_y + hipThreadIdx_y;
out[y * width + x] = in[x * height + y];
}
// CPU implementation of matrix transpose
void matrixTransposeCPUReference(
float * output,
float * input,
const unsigned int width,
const unsigned int height)
{
for(unsigned int j=0; j < height; j++)
{
for(unsigned int i=0; i < width; i++)
{
output[i*height + 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, HEIGHT/THREADS_PER_BLOCK_Y),
dim3(THREADS_PER_BLOCK_X, THREADS_PER_BLOCK_Y),
0, 0,
gpuTransposeMatrix , gpuMatrix, WIDTH ,HEIGHT);
// 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, HEIGHT);
// verify the results
errors = 0;
double eps = 1.0E-6;
for (i = 0; i < NUM; i++) {
if (std::abs(TransposeMatrix[i] - cpuTransposeMatrix[i]) > 0 ) {
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;
}
hipamd/samples/2_Cookbook/2_HIP_ATP_MARKER/Makefile
+36
View File
@@ -0,0 +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=./exe
.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
hipamd/samples/2_Cookbook/2_HIP_ATP_MARKER/MatrixTranspose.cpp
+174
View File
@@ -0,0 +1,174 @@
/*
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<iostream>
// hip header file
#include "hip_runtime.h"
#define WIDTH 1024
#define HEIGHT 1024
#define NUM (WIDTH*HEIGHT)
#define THREADS_PER_BLOCK_X 16
#define THREADS_PER_BLOCK_Y 16
#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,
const int height)
{
int x = hipBlockDim_x * hipBlockIdx_x + hipThreadIdx_x;
int y = hipBlockDim_y * hipBlockIdx_y + hipThreadIdx_y;
out[y * width + x] = in[x * height + y];
}
// CPU implementation of matrix transpose
void matrixTransposeCPUReference(
float * output,
float * input,
const unsigned int width,
const unsigned int height)
{
for(unsigned int j=0; j < height; j++)
{
for(unsigned int i=0; i < width; i++)
{
output[i*height + 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, HEIGHT/THREADS_PER_BLOCK_Y),
dim3(THREADS_PER_BLOCK_X, THREADS_PER_BLOCK_Y),
0, 0,
gpuTransposeMatrix , gpuMatrix, WIDTH ,HEIGHT);
// 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, HEIGHT);
// verify the results
errors = 0;
double eps = 1.0E-6;
for (i = 0; i < NUM; i++) {
if (std::abs(TransposeMatrix[i] - cpuTransposeMatrix[i]) > 0 ) {
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;
}
hipamd/samples/2_Cookbook/2_HIP_ATP_MARKER/Readme.md
+51
View File
@@ -0,0 +1,51 @@
## 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.
## Introduction:
CodeXL and rocm-profiler are the tool used for profiling the application, which is of prominent use in optimizing the application by means of finding the memory bottlenecks and etc.
## Requirement:
[CodeXL Installation](http://gpuopen.com/compute-product/codexl/)
## prerequiste knowledge:
Programmers familiar with CUDA, OpenCL will be able to quickly learn and start coding with the HIP API. In case you are not, don't worry. You choose to start with the best one. We'll be explaining everything assuming you are completely new to gpgpu programming.
## Simple Matrix Transpose
We will be using the Simple Matrix Transpose source code from the previous tutorial as it is.
## 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 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 <outputATPFileName> -A <applicationName> <applicationArguments>`
##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:
`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.
## More Info:
- [HIP FAQ](https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP/docs/markdown/hip_faq.md)
- [HIP Kernel Language](https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP/docs/markdown/hip_kernel_language.md)
- [HIP Runtime API (Doxygen)](http://gpuopen-professionalcompute-tools.github.io/HIP)
- [HIP Porting Guide](https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP/docs/markdown/hip_porting_guide.md)
- [HIP Terminology](https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP/docs/markdown/hip_terms.md) (including Rosetta Stone of GPU computing terms across CUDA/HIP/HC/AMP/OpenL)
- [clang-hipify](https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP/clang-hipify/README.md)
- [Developer/CONTRIBUTING Info](https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP/CONTRIBUTING.md)
- [Release Notes](https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP/RELEASE.md)
hipamd/samples/2_Cookbook/3_shared_memory/Makefile
+36
View File
@@ -0,0 +1,36 @@
HIP_PATH?= $(wildcard /opt/rocm/hip)
ifeq (,$(HIP_PATH))
HIP_PATH=../../..
endif
HIPCC=$(HIP_PATH)/bin/hipcc
TARGET=hcc
SOURCES = sharedMemory.cpp
OBJECTS = $(SOURCES:.cpp=.o)
EXECUTABLE=./exe
.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
hipamd/samples/2_Cookbook/3_shared_memory/sharedMemory.cpp
+144
View File
@@ -0,0 +1,144 @@
/*
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<iostream>
// hip header file
#include "hip_runtime.h"
#define WIDTH 1024
#define HEIGHT 1024
#define NUM (WIDTH*HEIGHT)
#define THREADS_PER_BLOCK_X 16
#define THREADS_PER_BLOCK_Y 16
#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,
const int height)
{
__shared__ float sharedMem[16*16];
int x = hipBlockDim_x * hipBlockIdx_x + hipThreadIdx_x;
int y = hipBlockDim_y * hipBlockIdx_y + hipThreadIdx_y;
sharedMem[y * width + x] = in[x * height + y];
__syncthreads();
out[y * width + x] = sharedMem[y * width + x];
}
// CPU implementation of matrix transpose
void matrixTransposeCPUReference(
float * output,
float * input,
const unsigned int width,
const unsigned int height)
{
for(unsigned int j=0; j < height; j++)
{
for(unsigned int i=0; i < width; i++)
{
output[i*height + 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;
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));
// Memory transfer from host to device
hipMemcpy(gpuMatrix, Matrix, NUM*sizeof(float), hipMemcpyHostToDevice);
// Lauching kernel from host
hipLaunchKernel(matrixTranspose,
dim3(WIDTH/THREADS_PER_BLOCK_X, HEIGHT/THREADS_PER_BLOCK_Y),
dim3(THREADS_PER_BLOCK_X, THREADS_PER_BLOCK_Y),
0, 0,
gpuTransposeMatrix , gpuMatrix, WIDTH ,HEIGHT);
// Memory transfer from device to host
hipMemcpy(TransposeMatrix, gpuTransposeMatrix, NUM*sizeof(float), hipMemcpyDeviceToHost);
// CPU MatrixTranspose computation
matrixTransposeCPUReference(cpuTransposeMatrix, Matrix, WIDTH, HEIGHT);
// verify the results
errors = 0;
double eps = 1.0E-6;
for (i = 0; i < NUM; i++) {
if (std::abs(TransposeMatrix[i] - cpuTransposeMatrix[i]) > 0 ) {
printf("%d cpu: %f gpu %f\n",i,cpuTransposeMatrix[i],TransposeMatrix[i]);
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;
}
hipamd/samples/2_Cookbook/4_shfl/Makefile
+36
View File
@@ -0,0 +1,36 @@
HIP_PATH?= $(wildcard /opt/rocm/hip)
ifeq (,$(HIP_PATH))
HIP_PATH=../../..
endif
HIPCC=$(HIP_PATH)/bin/hipcc
TARGET=hcc
SOURCES = shfl.cpp
OBJECTS = $(SOURCES:.cpp=.o)
EXECUTABLE=./exe
.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
hipamd/samples/2_Cookbook/4_shfl/shfl.cpp
+143
View File
@@ -0,0 +1,143 @@
/*
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<iostream>
// hip header file
#include "hip_runtime.h"
#define WIDTH 4
#define HEIGHT 4
#define NUM (WIDTH*HEIGHT)
#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,
const int height)
{
int x = hipBlockDim_x * hipBlockIdx_x + hipThreadIdx_x;
//int y = hipBlockDim_y * hipBlockIdx_y + hipThreadIdx_y;
float val = in[x];
for(int i=0;i<width;i++)
{
for(int j=0;j<width;j++)
out[i*height + j] = __shfl(val,j*width + i);
}
}
// CPU implementation of matrix transpose
void matrixTransposeCPUReference(
float * output,
float * input,
const unsigned int width,
const unsigned int height)
{
for(unsigned int j=0; j < height; j++)
{
for(unsigned int i=0; i < width; i++)
{
output[i*height + 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;
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));
// Memory transfer from host to device
hipMemcpy(gpuMatrix, Matrix, NUM*sizeof(float), hipMemcpyHostToDevice);
// Lauching kernel from host
hipLaunchKernel(matrixTranspose,
dim3(1),
dim3(THREADS_PER_BLOCK_X * THREADS_PER_BLOCK_Y),
0, 0,
gpuTransposeMatrix , gpuMatrix, WIDTH ,HEIGHT);
// Memory transfer from device to host
hipMemcpy(TransposeMatrix, gpuTransposeMatrix, NUM*sizeof(float), hipMemcpyDeviceToHost);
// CPU MatrixTranspose computation
matrixTransposeCPUReference(cpuTransposeMatrix, Matrix, WIDTH, HEIGHT);
// verify the results
errors = 0;
double eps = 1.0E-6;
for (i = 0; i < NUM; i++) {
if (std::abs(TransposeMatrix[i] - cpuTransposeMatrix[i]) > 0 ) {
printf("%d cpu: %f gpu %f\n",i,cpuTransposeMatrix[i],TransposeMatrix[i]);
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;
}
hipamd/samples/2_Cookbook/5_2dshfl/2dshfl.cpp
+139
View File
@@ -0,0 +1,139 @@
/*
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<iostream>
// hip header file
#include "hip_runtime.h"
#define WIDTH 4
#define HEIGHT 4
#define NUM (WIDTH*HEIGHT)
#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,
const int height)
{
int x = hipBlockDim_x * hipBlockIdx_x + hipThreadIdx_x;
int y = hipBlockDim_y * hipBlockIdx_y + hipThreadIdx_y;
float val = in[y*width + x];
out[x*height + y] = __shfl(val,y*width + x);
}
// CPU implementation of matrix transpose
void matrixTransposeCPUReference(
float * output,
float * input,
const unsigned int width,
const unsigned int height)
{
for(unsigned int j=0; j < height; j++)
{
for(unsigned int i=0; i < width; i++)
{
output[i*height + 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;
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));
// Memory transfer from host to device
hipMemcpy(gpuMatrix, Matrix, NUM*sizeof(float), hipMemcpyHostToDevice);
// Lauching kernel from host
hipLaunchKernel(matrixTranspose,
dim3(1),
dim3(THREADS_PER_BLOCK_X , THREADS_PER_BLOCK_Y),
0, 0,
gpuTransposeMatrix , gpuMatrix, WIDTH ,HEIGHT);
// Memory transfer from device to host
hipMemcpy(TransposeMatrix, gpuTransposeMatrix, NUM*sizeof(float), hipMemcpyDeviceToHost);
// CPU MatrixTranspose computation
matrixTransposeCPUReference(cpuTransposeMatrix, Matrix, WIDTH, HEIGHT);
// verify the results
errors = 0;
double eps = 1.0E-6;
for (i = 0; i < NUM; i++) {
if (std::abs(TransposeMatrix[i] - cpuTransposeMatrix[i]) > 0 ) {
printf("%d cpu: %f gpu %f\n",i,cpuTransposeMatrix[i],TransposeMatrix[i]);
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;
}
hipamd/samples/2_Cookbook/5_2dshfl/Makefile
+36
View File
@@ -0,0 +1,36 @@
HIP_PATH?= $(wildcard /opt/rocm/hip)
ifeq (,$(HIP_PATH))
HIP_PATH=../../..
endif
HIPCC=$(HIP_PATH)/bin/hipcc
TARGET=hcc
SOURCES = 2dshfl.cpp
OBJECTS = $(SOURCES:.cpp=.o)
EXECUTABLE=./exe
.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
hipamd/src/hip_context.cpp
+19 -2
View File
@@ -79,12 +79,29 @@ hipError_t hipDeviceGet(hipDevice_t *device, int deviceId)
hipError_t hipDriverGetVersion(int *driverVersion)
{
HIP_INIT_API(driverVersion);
hipError_t e = hipSuccess;
if (driverVersion) {
*driverVersion = 4;
}
else {
e = hipErrorInvalidValue;
}
return ihipLogStatus(hipSuccess);
return ihipLogStatus(e);
}
hipError_t hipRuntimeGetVersion(int *runtimeVersion)
{
HIP_INIT_API(runtimeVersion);
hipError_t e = hipSuccess;
if (runtimeVersion) {
*runtimeVersion = HIP_VERSION_PATCH;
}
else {
e = hipErrorInvalidValue;
}
return ihipLogStatus(e);
}
hipError_t hipCtxDestroy(hipCtx_t ctx)
hipamd/src/hip_device.cpp
+34 -14
View File
@@ -33,11 +33,15 @@ hipError_t hipGetDevice(int *deviceId)
auto ctx = ihipGetTlsDefaultCtx();
if (ctx == nullptr) {
e = hipErrorInvalidDevice; // TODO, check error code.
*deviceId = -1;
} else {
*deviceId = ctx->getDevice()->_deviceId;
if(deviceId != nullptr){
if (ctx == nullptr) {
e = hipErrorInvalidDevice; // TODO, check error code.
*deviceId = -1;
} else {
*deviceId = ctx->getDevice()->_deviceId;
}
}else{
e = hipErrorInvalidDevice;
}
return ihipLogStatus(e);
@@ -48,13 +52,20 @@ hipError_t hipGetDeviceCount(int *count)
{
HIP_INIT_API(count);
*count = g_deviceCnt;
hipError_t e = hipSuccess;
if (*count > 0) {
return ihipLogStatus(hipSuccess);
if(count != nullptr) {
*count = g_deviceCnt;
if (*count > 0) {
e = ihipLogStatus(hipSuccess);
} else {
e = ihipLogStatus(hipErrorNoDevice);
}
} else {
return ihipLogStatus(hipErrorNoDevice);
e = ihipLogStatus(hipErrorNoDevice);
}
return e;
}
hipError_t hipDeviceSetCacheConfig ( hipFuncCache cacheConfig )
@@ -144,6 +155,8 @@ hipError_t hipDeviceGetAttribute(int* pi, hipDeviceAttribute_t attr, int device)
hipError_t e = hipSuccess;
if(pi != nullptr) {
auto * hipDevice = ihipGetDevice(device);
hipDeviceProp_t *prop = &hipDevice->_props;
if (hipDevice) {
@@ -204,6 +217,9 @@ hipError_t hipDeviceGetAttribute(int* pi, hipDeviceAttribute_t attr, int device)
} else {
e = hipErrorInvalidDevice;
}
}else{
e = hipErrorInvalidDevice;
}
return ihipLogStatus(e);
}
@@ -213,12 +229,16 @@ hipError_t hipGetDeviceProperties(hipDeviceProp_t* props, int device)
hipError_t e;
auto * hipDevice = ihipGetDevice(device);
if (hipDevice) {
if(props != nullptr){
auto * hipDevice = ihipGetDevice(device);
if (hipDevice) {
// copy saved props
*props = hipDevice->_props;
e = hipSuccess;
} else {
*props = hipDevice->_props;
e = hipSuccess;
} else {
e = hipErrorInvalidDevice;
}
}else{
e = hipErrorInvalidDevice;
}
hipamd/src/hip_event.cpp
-2
View File
@@ -39,7 +39,6 @@ hipError_t ihipEventCreate(hipEvent_t* event, unsigned flags)
eh->_stream = NULL;
eh->_flags = flags;
eh->_timestamp = 0;
eh->_copySeqId = 0;
*event = eh; // TODO - allocat the event directly, no copy needed.
} else {
e = hipErrorInvalidValue;
@@ -123,7 +122,6 @@ hipError_t hipEventSynchronize(hipEvent_t event)
return ihipLogStatus(hipSuccess);
} else {
event->_marker.wait((event->_flags & hipEventBlockingSync) ? hc::hcWaitModeBlocked : hc::hcWaitModeActive);
event->_stream->locked_reclaimSignals(event->_copySeqId);
return ihipLogStatus(hipSuccess);
}
hipamd/src/hip_hcc.cpp
+114 -584
View File
@@ -67,24 +67,8 @@ int HIP_TRACE_API= 0;
std::string HIP_TRACE_API_COLOR("green");
int HIP_ATP_MARKER= 0;
int HIP_DB= 0;
int HIP_STAGING_SIZE = 64; /* size of staging buffers, in KB */
int HIP_STAGING_BUFFERS = 2; // TODO - remove, two buffers should be enough.
int HIP_PININPLACE = 0;
int HIP_OPTIMAL_MEM_TRANSFER = 0; //ENV Variable to test different memory transfer logics
int HIP_H2D_MEM_TRANSFER_THRESHOLD_DIRECT_OR_STAGING = 0;
int HIP_H2D_MEM_TRANSFER_THRESHOLD_STAGING_OR_PININPLACE = 0;
int HIP_D2H_MEM_TRANSFER_THRESHOLD = 0;
int HIP_STREAM_SIGNALS = 32; /* number of signals to allocate at stream creation */
int HIP_VISIBLE_DEVICES = 0; /* Contains a comma-separated sequence of GPU identifiers */
//---
// Chicken bits for disabling functionality to work around potential issues:
int HIP_DISABLE_HW_KERNEL_DEP = 0;
int HIP_DISABLE_HW_COPY_DEP = 0;
int HIP_NUM_KERNELS_INFLIGHT = 128;
std::once_flag hip_initialized;
@@ -99,9 +83,6 @@ std::vector<int> g_hip_visible_devices;
hsa_agent_t g_cpu_agent;
// TODO, remove these if possible:
hsa_agent_t gpu_agent_;
hsa_amd_memory_pool_t gpu_pool_;
//=================================================================================================
// Thread-local storage:
@@ -169,29 +150,6 @@ hipError_t ihipSynchronize(void)
return (hipSuccess);
}
//=================================================================================================
// ihipSignal_t:
//=================================================================================================
//
//---
ihipSignal_t::ihipSignal_t() : _sigId(0)
{
if (hsa_signal_create(0/*value*/, 0, NULL, &_hsaSignal) != HSA_STATUS_SUCCESS) {
throw ihipException(hipErrorRuntimeMemory);
}
//tprintf (DB_SIGNAL, " allocated hsa_signal=%lu\n", (_hsaSignal.handle));
}
//---
ihipSignal_t::~ihipSignal_t()
{
tprintf (DB_SIGNAL, " destroy hsa_signal #%lu (#%lu)\n", (_hsaSignal.handle), _sigId);
if (hsa_signal_destroy(_hsaSignal) != HSA_STATUS_SUCCESS) {
throw ihipException(hipErrorRuntimeOther);
}
};
//=================================================================================================
// ihipStream_t:
@@ -213,74 +171,16 @@ ihipStream_t::~ihipStream_t()
}
//---
//TODO - this function is dangerous since it does not propertly account
//for younger commands which may be depending on the signals we are reclaiming.
//Will fix when we move to HCC management of copy signals.
void ihipStream_t::locked_reclaimSignals(SIGSEQNUM sigNum)
{
LockedAccessor_StreamCrit_t crit(_criticalData);
tprintf(DB_SIGNAL, "reclaim signal #%lu\n", sigNum);
// Mark all signals older and including this one as available for re-allocation.
crit->_oldest_live_sig_id = sigNum+1;
}
//---
void ihipStream_t::waitCopy(LockedAccessor_StreamCrit_t &crit, ihipSignal_t *signal)
{
SIGSEQNUM sigNum = signal->_sigId;
tprintf(DB_SYNC, "waitCopy signal:#%lu\n", sigNum);
hsa_signal_wait_acquire(signal->_hsaSignal, HSA_SIGNAL_CONDITION_LT, 1, UINT64_MAX, HSA_WAIT_STATE_ACTIVE);
tprintf(DB_SIGNAL, "waitCopy reclaim signal #%lu\n", sigNum);
// Mark all signals older and including this one as available for reclaim
if (sigNum > crit->_oldest_live_sig_id) {
crit->_oldest_live_sig_id = sigNum+1; // TODO, +1 here seems dangerous.
}
}
//Wait for all kernel and data copy commands in this stream to complete.
//This signature should be used in routines that already have locked the stream mutex
void ihipStream_t::wait(LockedAccessor_StreamCrit_t &crit, bool assertQueueEmpty)
{
if (! assertQueueEmpty) {
tprintf (DB_SYNC, "stream %p wait for queue-empty..\n", this);
// crit->_av.wait();
waitOnAllCFs(crit);
}
if (crit->_last_copy_signal) {
tprintf (DB_SYNC, "stream %p wait for lastCopy:#%lu...\n", this, lastCopySeqId(crit) );
this->waitCopy(crit, crit->_last_copy_signal);
crit->_av.wait();
}
crit->_kernelCnt = 0;
// Reset the stream to "empty" - next command will not set up an inpute dependency on any older signal.
crit->_last_command_type = ihipCommandCopyH2D;
crit->_last_copy_signal = NULL;
// crit->_signalCnt = 0;
}
void ihipStream_t::addCFtoStream(LockedAccessor_StreamCrit_t &crit, hc::completion_future *cf)
{
crit->_cfs.push_back(cf);
}
void ihipStream_t::waitOnAllCFs(LockedAccessor_StreamCrit_t &crit)
{
for(uint32_t i=0;i<crit->_cfs.size();i++){
if(crit->_cfs[i] != NULL){
crit->_cfs[i]->wait();
delete crit->_cfs[i];
}
}
crit->_cfs.clear();
}
//---
@@ -293,7 +193,6 @@ void ihipStream_t::locked_wait(bool assertQueueEmpty)
};
#if USE_AV_COPY
// Causes current stream to wait for specified event to complete:
void ihipStream_t::locked_waitEvent(hipEvent_t event)
{
@@ -302,7 +201,6 @@ void ihipStream_t::locked_waitEvent(hipEvent_t event)
// TODO - check state of event here:
crit->_av.create_blocking_marker(event->_marker);
}
#endif
// Create a marker in this stream.
// Save state in the event so it can track the status of the event.
@@ -312,7 +210,6 @@ void ihipStream_t::locked_recordEvent(hipEvent_t event)
LockedAccessor_StreamCrit_t crit(_criticalData);
event->_marker = crit->_av.create_marker();
event->_copySeqId = lastCopySeqId(crit);
}
//=============================================================================
@@ -336,131 +233,19 @@ ihipCtx_t * ihipStream_t::getCtx() const
};
#define HIP_NUM_SIGNALS_PER_STREAM 32
//---
// Allocate a new signal from the signal pool.
// Returned signals have value of 0.
// Signals are intended for use in this stream and are always reclaimed "in-order".
ihipSignal_t *ihipStream_t::allocSignal(LockedAccessor_StreamCrit_t &crit)
{
int numToScan = crit->_signalPool.size();
crit->_signalCnt++;
if(crit->_signalCnt == HIP_STREAM_SIGNALS){
this->wait(crit);
crit->_signalCnt = 0;
}
return &crit->_signalPool[crit->_signalCnt];
do {
auto thisCursor = crit->_signalCursor;
if (++crit->_signalCursor == crit->_signalPool.size()) {
crit->_signalCursor = 0;
}
if (crit->_signalPool[thisCursor]._sigId < crit->_oldest_live_sig_id) {
SIGSEQNUM oldSigId = crit->_signalPool[thisCursor]._sigId;
crit->_signalPool[thisCursor]._index = thisCursor;
crit->_signalPool[thisCursor]._sigId = ++crit->_streamSigId; // allocate it.
tprintf(DB_SIGNAL, "allocatSignal #%lu at pos:%i (old sigId:%lu < oldest_live:%lu)\n",
crit->_signalPool[thisCursor]._sigId,
thisCursor, oldSigId, crit->_oldest_live_sig_id);
return &crit->_signalPool[thisCursor];
}
} while (--numToScan) ;
assert(numToScan == 0);
// Have to grow the pool:
crit->_signalCursor = crit->_signalPool.size(); // set to the beginning of the new entries:
if (crit->_signalCursor > 10000) {
fprintf (stderr, "warning: signal pool size=%d, may indicate runaway number of inflight commands\n", crit->_signalCursor);
}
crit->_signalPool.resize(crit->_signalPool.size() * 2);
tprintf (DB_SIGNAL, "grow signal pool to %zu entries, cursor=%d\n", crit->_signalPool.size(), crit->_signalCursor);
return allocSignal(crit); // try again,
// Should never reach here.
assert(0);
}
//---
void ihipStream_t::enqueueBarrier(hsa_queue_t* queue, ihipSignal_t *depSignal, ihipSignal_t *completionSignal)
{
// Obtain the write index for the command queue
uint64_t index = hsa_queue_load_write_index_relaxed(queue);
const uint32_t queueMask = queue->size - 1;
// Define the barrier packet to be at the calculated queue index address
hsa_barrier_and_packet_t* barrier = &(((hsa_barrier_and_packet_t*)(queue->base_address))[index&queueMask]);
memset(barrier, 0, sizeof(hsa_barrier_and_packet_t));
// setup header
uint16_t header = HSA_PACKET_TYPE_BARRIER_AND << HSA_PACKET_HEADER_TYPE;
header |= 1 << HSA_PACKET_HEADER_BARRIER;
//header |= HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_ACQUIRE_FENCE_SCOPE;
//header |= HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_RELEASE_FENCE_SCOPE;
barrier->header = header;
barrier->dep_signal[0].handle = depSignal ? depSignal->_hsaSignal.handle: 0;
barrier->completion_signal.handle = completionSignal ? completionSignal->_hsaSignal.handle : 0;
// TODO - check queue overflow, return error:
// Increment write index and ring doorbell to dispatch the kernel
hsa_queue_store_write_index_relaxed(queue, index+1);
hsa_signal_store_relaxed(queue->doorbell_signal, index);
}
int HIP_NUM_KERNELS_INFLIGHT = 128;
//--
//When the commands in a stream change types (ie kernel command follows a data command,
//or data command follows a kernel command), then we need to add a barrier packet
//into the stream to mimic CUDA stream semantics. (some hardware uses separate
//queues for data commands and kernel commands, and no implicit ordering is provided).
//
// Lock the stream to prevent other threads from intervening.
LockedAccessor_StreamCrit_t ihipStream_t::lockopen_preKernelCommand()
{
LockedAccessor_StreamCrit_t crit(_criticalData, false/*no unlock at destruction*/);
if(crit->_kernelCnt > HIP_NUM_KERNELS_INFLIGHT){
this->wait(crit);
this->wait(crit);
crit->_kernelCnt = 0;
}
crit->_kernelCnt++;
// If switching command types, we need to add a barrier packet to synchronize things.
if (crit->_last_command_type != ihipCommandKernel) {
if (crit->_last_copy_signal) {
hsa_queue_t * q = (hsa_queue_t*) (crit->_av.get_hsa_queue());
if (HIP_DISABLE_HW_KERNEL_DEP == 0) {
this->enqueueBarrier(q, crit->_last_copy_signal, NULL);
tprintf (DB_SYNC, "stream %p switch %s to %s (barrier pkt inserted with wait on #%lu)\n",
this, ihipCommandName[crit->_last_command_type], ihipCommandName[ihipCommandKernel], crit->_last_copy_signal->_sigId)
} else if (HIP_DISABLE_HW_KERNEL_DEP>0) {
tprintf (DB_SYNC, "stream %p switch %s to %s (HOST wait for previous...)\n",
this, ihipCommandName[crit->_last_command_type], ihipCommandName[ihipCommandKernel]);
this->waitCopy(crit, crit->_last_copy_signal);
} else if (HIP_DISABLE_HW_KERNEL_DEP==-1) {
tprintf (DB_SYNC, "stream %p switch %s to %s (IGNORE dependency)\n",
this, ihipCommandName[crit->_last_command_type], ihipCommandName[ihipCommandKernel]);
}
}
crit->_last_command_type = ihipCommandKernel;
}
return crit;
}
@@ -468,13 +253,12 @@ LockedAccessor_StreamCrit_t ihipStream_t::lockopen_preKernelCommand()
//---
// Must be called after kernel finishes, this releases the lock on the stream so other commands can submit.
void ihipStream_t::lockclose_postKernelCommand(hc::completion_future &kernelFuture)
void ihipStream_t::lockclose_postKernelCommand(hc::accelerator_view *av)
{
// We locked _criticalData in the lockopen_preKernelCommand() so OK to access here:
_criticalData._last_kernel_future = kernelFuture;
if (HIP_LAUNCH_BLOCKING) {
kernelFuture.wait();
// TODO - fix this so it goes through proper stream::wait() call.
av->wait(); // direct wait OK since we know the stream is locked.
tprintf(DB_SYNC, " %s LAUNCH_BLOCKING for kernel completion\n", ToString(this).c_str());
}
@@ -483,54 +267,6 @@ void ihipStream_t::lockclose_postKernelCommand(hc::completion_future &kernelFutu
//---
// Called whenever a copy command is set to the stream.
// Examines the last command sent to this stream and returns a signal to wait on, if required.
int ihipStream_t::preCopyCommand(LockedAccessor_StreamCrit_t &crit, ihipSignal_t *copyCompletionSignal, hsa_signal_t *waitSignal, ihipCommand_t copyType)
{
int needSync = 0;
waitSignal->handle = 0;
// If switching command types, we need to add a barrier packet to synchronize things.
if (FORCE_SAMEDIR_COPY_DEP || (crit->_last_command_type != copyType)) {
if (crit->_last_command_type == ihipCommandKernel) {
tprintf (DB_SYNC, "stream %p switch %s to %s (async copy dep on prev kernel)\n",
this, ihipCommandName[crit->_last_command_type], ihipCommandName[copyType]);
needSync = 1;
ihipSignal_t *depSignal = allocSignal(crit);
hsa_signal_store_relaxed(depSignal->_hsaSignal,1);
this->enqueueBarrier(static_cast<hsa_queue_t*>(crit->_av.get_hsa_queue()), NULL, depSignal);
*waitSignal = depSignal->_hsaSignal;
} else if (crit->_last_copy_signal) {
needSync = 1;
tprintf (DB_SYNC, "stream %p switch %s to %s (async copy dep on other copy #%lu)\n",
this, ihipCommandName[crit->_last_command_type], ihipCommandName[copyType], crit->_last_copy_signal->_sigId);
*waitSignal = crit->_last_copy_signal->_hsaSignal;
}
if (HIP_DISABLE_HW_COPY_DEP && needSync) {
if (HIP_DISABLE_HW_COPY_DEP == -1) {
tprintf (DB_SYNC, "IGNORE copy dependency\n")
} else {
tprintf (DB_SYNC, "HOST-wait for copy dependency\n")
// do the wait here on the host, and disable the device-side command resolution.
hsa_signal_wait_acquire(*waitSignal, HSA_SIGNAL_CONDITION_LT, 1, UINT64_MAX, HSA_WAIT_STATE_ACTIVE);
needSync = 0;
}
}
crit->_last_command_type = copyType;
}
crit->_last_copy_signal = copyCompletionSignal;
return needSync;
}
// Precursor: the stream is already locked,specifically so this routine can enqueue work into the specified av.
void ihipStream_t::launchModuleKernel(
@@ -543,7 +279,7 @@ void ihipStream_t::launchModuleKernel(
uint32_t gridDimY,
uint32_t gridDimZ,
uint32_t groupSegmentSize,
uint32_t privateSegmentSize,
uint32_t privateSegmentSize,
void *kernarg,
size_t kernSize,
uint64_t kernel){
@@ -646,6 +382,18 @@ void ihipCtxCriticalBase_t<CtxMutex>::resetPeers(ihipCtx_t *thisDevice)
}
template<>
void ihipCtxCriticalBase_t<CtxMutex>::printPeers(FILE *f) const
{
for (auto iter = _peers.begin(); iter!=_peers.end(); iter++) {
fprintf (f, "%s ", (*iter)->toString().c_str());
};
}
template<>
void ihipCtxCriticalBase_t<CtxMutex>::addStream(ihipStream_t *stream)
{
@@ -675,8 +423,6 @@ ihipDevice_t::ihipDevice_t(unsigned deviceId, unsigned deviceCnt, hc::accelerato
initProperties(&_props);
_stagingBuffer[0] = new UnpinnedCopyEngine(_hsaAgent,g_cpu_agent, HIP_STAGING_SIZE*1024, HIP_STAGING_BUFFERS,HIP_H2D_MEM_TRANSFER_THRESHOLD_DIRECT_OR_STAGING,HIP_H2D_MEM_TRANSFER_THRESHOLD_STAGING_OR_PININPLACE,HIP_D2H_MEM_TRANSFER_THRESHOLD);
_stagingBuffer[1] = new UnpinnedCopyEngine(_hsaAgent,g_cpu_agent, HIP_STAGING_SIZE*1024, HIP_STAGING_BUFFERS,HIP_H2D_MEM_TRANSFER_THRESHOLD_DIRECT_OR_STAGING,HIP_H2D_MEM_TRANSFER_THRESHOLD_STAGING_OR_PININPLACE,HIP_D2H_MEM_TRANSFER_THRESHOLD);
_primaryCtx = new ihipCtx_t(this, deviceCnt, hipDeviceMapHost);
}
@@ -684,12 +430,8 @@ ihipDevice_t::ihipDevice_t(unsigned deviceId, unsigned deviceCnt, hc::accelerato
ihipDevice_t::~ihipDevice_t()
{
for (int i=0; i<2; i++) {
if (_stagingBuffer[i]) {
delete _stagingBuffer[i];
_stagingBuffer[i] = NULL;
}
}
delete _primaryCtx;
_primaryCtx = NULL;
}
@@ -763,14 +505,6 @@ hsa_status_t GetDevicePool(hsa_amd_memory_pool_t pool, void* data) {
return HSA_STATUS_SUCCESS;
}
void FindDevicePool()
{
hsa_status_t err = hsa_iterate_agents(FindGpuDevice, &gpu_agent_);
ErrorCheck(err);
err = hsa_amd_agent_iterate_memory_pools(gpu_agent_, GetDevicePool, &gpu_pool_);
ErrorCheck(err);
}
int checkAccess(hsa_agent_t agent, hsa_amd_memory_pool_t pool)
{
@@ -923,14 +657,7 @@ hipError_t ihipDevice_t::initProperties(hipDeviceProp_t* prop)
/* Computemode for HSA Devices is always : cudaComputeModeDefault */
prop->computeMode = 0;
FindDevicePool();
int access=checkAccess(g_cpu_agent, gpu_pool_);
if (0!= access){
_isLargeBar= 1;
} else {
_isLargeBar=0;
}
_isLargeBar = _acc.has_cpu_accessible_am();
// Get Max Threads Per Multiprocessor
@@ -1070,6 +797,13 @@ void ihipCtx_t::locked_reset()
};
//---
std::string ihipCtx_t::toString() const
{
std::ostringstream ss;
ss << this;
return ss.str();
};
//----
@@ -1260,6 +994,7 @@ void ihipInit()
READ_ENV_I(release, HIP_PRINT_ENV, 0, "Print HIP environment variables.");
//-- READ HIP_PRINT_ENV env first, since it has impact on later env var reading
// TODO: In HIP/hcc, this variable blocks after both kernel commmands and data transfer. Maybe should be bit-mask for each command type?
READ_ENV_I(release, HIP_LAUNCH_BLOCKING, CUDA_LAUNCH_BLOCKING, "Make HIP APIs 'host-synchronous', so they block until any kernel launches or data copy commands complete. Alias: CUDA_LAUNCH_BLOCKING." );
READ_ENV_I(release, HIP_DB, 0, "Print various debug info. Bitmask, see hip_hcc.cpp for more information.");
if ((HIP_DB & (1<<DB_API)) && (HIP_TRACE_API == 0)) {
@@ -1270,35 +1005,11 @@ void ihipInit()
READ_ENV_I(release, HIP_TRACE_API, 0, "Trace each HIP API call. Print function name and return code to stderr as program executes.");
READ_ENV_S(release, HIP_TRACE_API_COLOR, 0, "Color to use for HIP_API. None/Red/Green/Yellow/Blue/Magenta/Cyan/White");
READ_ENV_I(release, HIP_ATP_MARKER, 0, "Add HIP function begin/end to ATP file generated with CodeXL");
READ_ENV_I(release, HIP_STAGING_SIZE, 0, "Size of each staging buffer (in KB)" );
READ_ENV_I(release, HIP_STAGING_BUFFERS, 0, "Number of staging buffers to use in each direction. 0=use hsa_memory_copy.");
READ_ENV_I(release, HIP_PININPLACE, 0, "For unpinned transfers, pin the memory in-place in chunks before doing the copy. Under development.");
READ_ENV_I(release, HIP_OPTIMAL_MEM_TRANSFER, 0, "For optimal memory transfers for unpinned memory.Under testing.");
READ_ENV_I(release, HIP_H2D_MEM_TRANSFER_THRESHOLD_DIRECT_OR_STAGING, 0, "Threshold value for H2D unpinned memory transfer decision between direct copy or staging buffer usage,Under testing.");
READ_ENV_I(release, HIP_H2D_MEM_TRANSFER_THRESHOLD_STAGING_OR_PININPLACE, 0, "Threshold value for H2D unpinned memory transfer decision between staging buffer usage or pininplace usage .Under testing.");
READ_ENV_I(release, HIP_D2H_MEM_TRANSFER_THRESHOLD, 0, "Threshold value for D2H unpinned memory transfer decision between staging buffer usage or pininplace usage .Under testing.");
READ_ENV_I(release, HIP_STREAM_SIGNALS, 0, "Number of signals to allocate when new stream is created (signal pool will grow on demand)");
READ_ENV_I(release, HIP_VISIBLE_DEVICES, CUDA_VISIBLE_DEVICES, "Only devices whose index is present in the secquence are visible to HIP applications and they are enumerated in the order of secquence" );
READ_ENV_I(release, HIP_DISABLE_HW_KERNEL_DEP, 0, "Disable HW dependencies before kernel commands - instead wait for dependency on host. -1 means ignore these dependencies. (debug mode)");
READ_ENV_I(release, HIP_DISABLE_HW_COPY_DEP, 0, "Disable HW dependencies before copy commands - instead wait for dependency on host. -1 means ifnore these dependencies (debug mode)");
READ_ENV_I(release, HIP_NUM_KERNELS_INFLIGHT, 128, "Number of kernels per stream ");
if (HIP_OPTIMAL_MEM_TRANSFER && !HIP_H2D_MEM_TRANSFER_THRESHOLD_DIRECT_OR_STAGING) {
HIP_H2D_MEM_TRANSFER_THRESHOLD_DIRECT_OR_STAGING= MEMCPY_H2D_DIRECT_VS_STAGING_COPY_THRESHOLD;
fprintf (stderr, "warning: env var HIP_OPTIMAL_MEM_TRANSFER=0x%x but HIP_H2D_MEM_TRANSFER_THRESHOLD_DIRECT_OR_STAGING=0.Using default value for this.\n", HIP_OPTIMAL_MEM_TRANSFER);
}
if (HIP_OPTIMAL_MEM_TRANSFER && !HIP_H2D_MEM_TRANSFER_THRESHOLD_STAGING_OR_PININPLACE) {
HIP_H2D_MEM_TRANSFER_THRESHOLD_STAGING_OR_PININPLACE= MEMCPY_H2D_STAGING_VS_PININPLACE_COPY_THRESHOLD;
fprintf (stderr, "warning: env var HIP_OPTIMAL_MEM_TRANSFER=0x%x but HIP_H2D_MEM_TRANSFER_THRESHOLD_STAGING_OR_PININPLACE=0.Using default value for this.\n", HIP_OPTIMAL_MEM_TRANSFER);
}
if (HIP_OPTIMAL_MEM_TRANSFER && !HIP_D2H_MEM_TRANSFER_THRESHOLD) {
HIP_D2H_MEM_TRANSFER_THRESHOLD= MEMCPY_D2H_STAGING_VS_PININPLACE_COPY_THRESHOLD;
fprintf (stderr, "warning: env var HIP_OPTIMAL_MEM_TRANSFER=0x%x but HIP_D2H_MEM_TRANSFER_THRESHOLD=0.Using default value for this.\n", HIP_OPTIMAL_MEM_TRANSFER);
}
// Some flags have both compile-time and runtime flags - generate a warning if user enables the runtime flag but the compile-time flag is disabled.
if (HIP_DB && !COMPILE_HIP_DB) {
fprintf (stderr, "warning: env var HIP_DB=0x%x but COMPILE_HIP_DB=0. (perhaps enable COMPILE_HIP_DB in src code before compiling?)", HIP_DB);
@@ -1436,7 +1147,6 @@ void ihipPrintKernelLaunch(const char *kernelName, const grid_launch_parm *lp, c
}
}
// TODO - data-up to data-down:
// Called just before a kernel is launched from hipLaunchKernel.
// Allows runtime to track some information about the stream.
hipStream_t ihipPreLaunchKernel(hipStream_t stream, dim3 grid, dim3 block, grid_launch_parm *lp, const char *kernelNameStr)
@@ -1454,8 +1164,7 @@ hipStream_t ihipPreLaunchKernel(hipStream_t stream, dim3 grid, dim3 block, grid_
auto crit = stream->lockopen_preKernelCommand();
lp->av = &(crit->_av);
lp->cf = new hc::completion_future;
stream->addCFtoStream(crit, lp->cf);
lp->cf = nullptr;
ihipPrintKernelLaunch(kernelNameStr, lp, stream);
return (stream);
@@ -1477,8 +1186,7 @@ hipStream_t ihipPreLaunchKernel(hipStream_t stream, size_t grid, dim3 block, gri
auto crit = stream->lockopen_preKernelCommand();
lp->av = &(crit->_av);
lp->cf = new hc::completion_future;
stream->addCFtoStream(crit, lp->cf);
lp->cf = nullptr;
ihipPrintKernelLaunch(kernelNameStr, lp, stream);
return (stream);
}
@@ -1499,8 +1207,7 @@ hipStream_t ihipPreLaunchKernel(hipStream_t stream, dim3 grid, size_t block, gri
auto crit = stream->lockopen_preKernelCommand();
lp->av = &(crit->_av);
lp->cf = new hc::completion_future;
stream->addCFtoStream(crit, lp->cf);
lp->cf = nullptr;
ihipPrintKernelLaunch(kernelNameStr, lp, stream);
return (stream);
}
@@ -1521,9 +1228,8 @@ hipStream_t ihipPreLaunchKernel(hipStream_t stream, size_t grid, size_t block, g
auto crit = stream->lockopen_preKernelCommand();
lp->av = &(crit->_av);
lp->cf = new hc::completion_future; // TODO, is this necessary?
lp->cf = nullptr;
stream->addCFtoStream(crit, lp->cf);
ihipPrintKernelLaunch(kernelNameStr, lp, stream);
return (stream);
@@ -1535,7 +1241,7 @@ 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)
{
stream->lockclose_postKernelCommand(*(lp.cf));
stream->lockclose_postKernelCommand(lp.av);
}
@@ -1642,8 +1348,38 @@ void ihipSetTs(hipEvent_t e)
}
// Returns true if thisCtx 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 booth include thisCtx.
bool ihipStream_t::canSeePeerMemory(const ihipCtx_t *thisCtx, ihipCtx_t *dstCtx, ihipCtx_t *srcCtx)
{
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;
};
}
{
LockedAccessor_CtxCrit_t ctxCrit(srcCtx->criticalData());
tprintf(DB_SYNC, "srcCrit lock succeeded\n");
if (!ctxCrit->isPeer(thisCtx)) {
return false;
};
}
return true;
};
// 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)
{
hipMemcpyKind kind = hipMemcpyDefault;
@@ -1672,30 +1408,8 @@ unsigned ihipStream_t::resolveMemcpyDirection(bool srcTracked, bool dstTracked,
}
// Setup the copyCommandType and the copy agents (for hsa_amd_memory_async_copy)
// srcPhysAcc is the physical location of the src data. For many copies this is the
void ihipStream_t::setAsyncCopyAgents(unsigned kind, ihipCommand_t *commandType, hsa_agent_t *srcAgent, hsa_agent_t *dstAgent)
{
// current* represents the device associated with the specified stream.
const ihipDevice_t *streamDevice = this->getDevice();
hsa_agent_t streamAgent = streamDevice->_hsaAgent;
// ROCR runtime logic is :
// - If both src and dst are cpu agent, launch thread and memcpy. We want to avoid this.
// - If either/both src or dst is a gpu agent, use the first gpu agent’s DMA engine to perform the copy.
switch (kind) {
//case hipMemcpyHostToHost : *commandType = ihipCommandCopyH2H; *srcAgent=streamAgent; *dstAgent=streamAgent; break; // TODO - enable me, for async copy use SDMA.
case hipMemcpyHostToHost : *commandType = ihipCommandCopyH2H; *srcAgent=g_cpu_agent; *dstAgent=g_cpu_agent; break;
case hipMemcpyHostToDevice : *commandType = ihipCommandCopyH2D; *srcAgent=g_cpu_agent; *dstAgent=streamAgent; break;
case hipMemcpyDeviceToHost : *commandType = ihipCommandCopyD2H; *srcAgent=streamAgent; *dstAgent=g_cpu_agent; break;
case hipMemcpyDeviceToDevice : *commandType = ihipCommandCopyD2D; *srcAgent=streamAgent; *dstAgent=streamAgent; break;
default: throw ihipException(hipErrorInvalidMemcpyDirection);
};
}
void ihipStream_t::copySync(LockedAccessor_StreamCrit_t &crit, void* dst, const void* src, size_t sizeBytes, unsigned kind, bool resolveOn)
// 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)
{
ihipCtx_t *ctx = this->getCtx();
const ihipDevice_t *device = ctx->getDevice();
@@ -1707,193 +1421,43 @@ void ihipStream_t::copySync(LockedAccessor_StreamCrit_t &crit, void* dst, const
hc::accelerator acc;
hc::AmPointerInfo dstPtrInfo(NULL, NULL, 0, acc, 0, 0);
hc::AmPointerInfo srcPtrInfo(NULL, NULL, 0, acc, 0, 0);
bool dstTracked = (hc::am_memtracker_getinfo(&dstPtrInfo, dst) == AM_SUCCESS);
bool srcTracked = (hc::am_memtracker_getinfo(&srcPtrInfo, src) == AM_SUCCESS);
bool srcInDeviceMem = srcPtrInfo._isInDeviceMem;
bool dstInDeviceMem = dstPtrInfo._isInDeviceMem;
// Resolve default to a specific Kind so we know which algorithm to use:
if (kind == hipMemcpyDefault && resolveOn) {
kind = resolveMemcpyDirection(srcTracked, dstTracked, srcInDeviceMem, dstInDeviceMem);
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;
};
hsa_signal_t depSignal;
bool copyEngineCanSeeSrcAndDest = false;
if (kind == hipMemcpyDeviceToDevice) {
#if USE_PEER_TO_PEER>=2
// Lock to prevent another thread from modifying peer list while we are trying to look at it.
LockedAccessor_CtxCrit_t dcrit(ctx->criticalData());
// FIXME - this assumes peer access only from primary context.
// Would need to change the tracker to store a void * parameter that we could map to the ctx where the pointer is allocated.
if (dcrit->isPeer(ihipGetPrimaryCtx(dstPtrInfo._appId)) && (dcrit->isPeer(ihipGetPrimaryCtx(srcPtrInfo._appId)))) {
copyEngineCanSeeSrcAndDest = true;
// 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());
}
#endif
};
{
LockedAccessor_StreamCrit_t crit (_criticalData);
crit->_av.copy_ext(src, dst, sizeBytes, hcCopyDir, srcPtrInfo, dstPtrInfo, forceHostCopyEngine);
}
if (kind == hipMemcpyHostToDevice) {
int depSignalCnt = preCopyCommand(crit, NULL, &depSignal, ihipCommandCopyH2D);
if(!srcTracked){
if (HIP_STAGING_BUFFERS) {
tprintf(DB_COPY1, "D2H && !dstTracked: staged copy H2D dst=%p src=%p sz=%zu\n", dst, src, sizeBytes);
if(HIP_OPTIMAL_MEM_TRANSFER)
{
device->_stagingBuffer[0]->CopyHostToDevice(1,device->_isLargeBar,dst, src, sizeBytes, depSignalCnt ? &depSignal : NULL);
}
else {
if (HIP_PININPLACE) {
device->_stagingBuffer[0]->CopyHostToDevicePinInPlace(dst, src, sizeBytes, depSignalCnt ? &depSignal : NULL);
} else {
device->_stagingBuffer[0]->CopyHostToDevice(0,0,dst, src, sizeBytes, depSignalCnt ? &depSignal : NULL);
}
}
// The copy waits for inputs and then completes before returning so can reset queue to empty:
this->wait(crit, true);
}
else {
// TODO - remove, slow path.
tprintf(DB_COPY1, "H2D && ! srcTracked: am_copy dst=%p src=%p sz=%zu\n", dst, src, sizeBytes);
#if USE_AV_COPY
crit->_av.copy(src,dst,sizeBytes);
#else
hc::am_copy(dst, src, sizeBytes);
#endif
}
} else {
// This is H2D copy, and source is pinned host memory : we can copy directly w/o using staging buffer.
hsa_agent_t dstAgent = *(static_cast<hsa_agent_t*>(dstPtrInfo._acc.get_hsa_agent()));
hsa_agent_t srcAgent = *(static_cast<hsa_agent_t*>(srcPtrInfo._acc.get_hsa_agent()));
ihipSignal_t *ihipSignal = allocSignal(crit);
hsa_signal_t copyCompleteSignal = ihipSignal->_hsaSignal;
hsa_signal_store_relaxed(copyCompleteSignal, 1);
void *devPtrSrc = srcPtrInfo._devicePointer;
tprintf(DB_COPY1, "HSA Async_copy dst=%p src=%p sz=%zu\n", dst, src, sizeBytes);
hsa_status_t hsa_status = hsa_amd_memory_async_copy(dst, dstAgent, devPtrSrc, g_cpu_agent, sizeBytes, depSignalCnt, depSignalCnt ? &depSignal:0x0, copyCompleteSignal);
// This is sync copy, so let's wait for copy right here:
if (hsa_status == HSA_STATUS_SUCCESS) {
waitCopy(crit, ihipSignal); // wait for copy, and return to pool.
} else {
throw ihipException(hipErrorInvalidValue);
}
}
} else if (kind == hipMemcpyDeviceToHost) {
int depSignalCnt = preCopyCommand(crit, NULL, &depSignal, ihipCommandCopyD2H);
if (!dstTracked){
if (HIP_STAGING_BUFFERS) {
tprintf(DB_COPY1, "D2H && !dstTracked: staged copy D2H dst=%p src=%p sz=%zu\n", dst, src, sizeBytes);
//printf ("staged-copy- read dep signals\n");
if(HIP_OPTIMAL_MEM_TRANSFER)
{
//printf ("staged-copy- read dep signals\n");
device->_stagingBuffer[1]->CopyDeviceToHost(1,dst, src, sizeBytes, depSignalCnt ? &depSignal : NULL);
}
else
{
device->_stagingBuffer[1]->CopyDeviceToHost(0,dst, src, sizeBytes, depSignalCnt ? &depSignal : NULL);
}
// The copy completes before returning so can reset queue to empty:
this->wait(crit, true);
} else {
// TODO - remove, slow path.
tprintf(DB_COPY1, "D2H && !dstTracked: am_copy dst=%p src=%p sz=%zu\n", dst, src, sizeBytes);
#if USE_AV_COPY
crit->_av.copy(src, dst, sizeBytes);
#else
hc::am_copy(dst, src, sizeBytes);
#endif
}
} else {
// This is D2H copy, and destination is pinned host memory : we can copy directly w/o using staging buffer.
hsa_agent_t dstAgent = *(static_cast<hsa_agent_t*>(dstPtrInfo._acc.get_hsa_agent()));
hsa_agent_t srcAgent = *(static_cast<hsa_agent_t*>(srcPtrInfo._acc.get_hsa_agent()));
ihipSignal_t *ihipSignal = allocSignal(crit);
hsa_signal_t copyCompleteSignal = ihipSignal->_hsaSignal;
hsa_signal_store_relaxed(copyCompleteSignal, 1);
void *devPtrDst = dstPtrInfo._devicePointer;
tprintf(DB_COPY1, "HSA Async_copy dst=%p src=%p sz=%zu\n", dst, src, sizeBytes);
hsa_status_t hsa_status = hsa_amd_memory_async_copy(devPtrDst, g_cpu_agent, src, srcAgent, sizeBytes, depSignalCnt, depSignalCnt ? &depSignal:0x0, copyCompleteSignal);
// This is sync copy, so let's wait for copy right here:
if (hsa_status == HSA_STATUS_SUCCESS) {
waitCopy(crit, ihipSignal); // wait for copy, and return to pool.
} else {
throw ihipException(hipErrorInvalidValue);
}
}
} else if (kind == hipMemcpyHostToHost) {
int depSignalCnt = preCopyCommand(crit, NULL, &depSignal, ihipCommandCopyH2H);
if (depSignalCnt) {
// host waits before doing host memory copy.
hsa_signal_wait_acquire(depSignal, HSA_SIGNAL_CONDITION_LT, 1, UINT64_MAX, HSA_WAIT_STATE_ACTIVE);
}
memcpy(dst, src, sizeBytes);
} else if ((kind == hipMemcpyDeviceToDevice) && !copyEngineCanSeeSrcAndDest) {
int depSignalCnt = preCopyCommand(crit, NULL, &depSignal, ihipCommandCopyP2P);
if (HIP_STAGING_BUFFERS) {
tprintf(DB_COPY1, "P2P but engine can't see both pointers: staged copy P2P dst=%p src=%p sz=%zu\n", dst, src, sizeBytes);
//printf ("staged-copy- read dep signals\n");
hsa_agent_t dstAgent = * (static_cast<hsa_agent_t*> (dstPtrInfo._acc.get_hsa_agent()));
hsa_agent_t srcAgent = * (static_cast<hsa_agent_t*> (srcPtrInfo._acc.get_hsa_agent()));
device->_stagingBuffer[1]->CopyPeerToPeer(dst, dstAgent, src, srcAgent, sizeBytes, depSignalCnt ? &depSignal : NULL);
// The copy completes before returning so can reset queue to empty:
this->wait(crit, true);
} else {
assert(0); // currently no fallback for this path.
}
} else {
// If not special case - these can all be handled by the hsa async copy:
ihipCommand_t commandType;
hsa_agent_t srcAgent, dstAgent;
setAsyncCopyAgents(kind, &commandType, &srcAgent, &dstAgent);
int depSignalCnt = preCopyCommand(crit, NULL, &depSignal, commandType);
// Get a completion signal:
ihipSignal_t *ihipSignal = allocSignal(crit);
hsa_signal_t copyCompleteSignal = ihipSignal->_hsaSignal;
hsa_signal_store_relaxed(copyCompleteSignal, 1);
tprintf(DB_COPY1, "HSA Async_copy dst=%p src=%p sz=%zu\n", dst, src, sizeBytes);
hsa_status_t hsa_status = hsa_amd_memory_async_copy(dst, dstAgent, src, srcAgent, sizeBytes, depSignalCnt, depSignalCnt ? &depSignal:0x0, copyCompleteSignal);
// This is sync copy, so let's wait for copy right here:
if (hsa_status == HSA_STATUS_SUCCESS) {
waitCopy(crit, ihipSignal); // wait for copy, and return to pool.
} else {
throw ihipException(hipErrorInvalidValue);
}
}
}
// Sync copy that acquires lock:
void ihipStream_t::locked_copySync(void* dst, const void* src, size_t sizeBytes, unsigned kind, bool resolveOn)
void ihipStream_t::locked_copyAsync(void* dst, const void* src, size_t sizeBytes, unsigned kind)
{
LockedAccessor_StreamCrit_t crit (_criticalData);
copySync(crit, dst, src, sizeBytes, kind, resolveOn);
}
void ihipStream_t::copyAsync(void* dst, const void* src, size_t sizeBytes, unsigned kind)
{
LockedAccessor_StreamCrit_t crit(_criticalData);
const ihipCtx_t *ctx = this->getCtx();
@@ -1908,12 +1472,12 @@ void ihipStream_t::copyAsync(void* dst, const void* src, size_t sizeBytes, unsig
/* As this is a CPU op, we need to wait until all
the commands in current stream are finished.
*/
LockedAccessor_StreamCrit_t crit(_criticalData);
this->wait(crit);
memcpy(dst, src, sizeBytes);
} else {
bool trueAsync = true;
hc::accelerator acc;
hc::AmPointerInfo dstPtrInfo(NULL, NULL, 0, acc, 0, 0);
@@ -1922,49 +1486,32 @@ void ihipStream_t::copyAsync(void* dst, const void* src, size_t sizeBytes, unsig
bool srcTracked = (hc::am_memtracker_getinfo(&srcPtrInfo, src) == AM_SUCCESS);
bool copyEngineCanSeeSrcAndDest = true;
if (kind == hipMemcpyDeviceToDevice) {
copyEngineCanSeeSrcAndDest = canSeePeerMemory(ctx, ihipGetPrimaryCtx(dstPtrInfo._appId), ihipGetPrimaryCtx(srcPtrInfo._appId));
}
// "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) {
trueAsync = false;
}
if (dstTracked && srcTracked && copyEngineCanSeeSrcAndDest) {
LockedAccessor_StreamCrit_t crit(_criticalData);
if (kind == hipMemcpyDefault) {
bool srcInDeviceMem = (srcTracked && srcPtrInfo._isInDeviceMem);
bool dstInDeviceMem = (dstTracked && dstPtrInfo._isInDeviceMem);
kind = resolveMemcpyDirection(srcTracked, dstTracked, srcPtrInfo._isInDeviceMem, dstPtrInfo._isInDeviceMem);
}
// Perform asynchronous copy:
try {
crit->_av.copy_async(src, dst, sizeBytes);
} catch (Kalmar::runtime_exception) {
throw ihipException(hipErrorRuntimeOther);
};
ihipSignal_t *ihip_signal = allocSignal(crit);
hsa_signal_store_relaxed(ihip_signal->_hsaSignal, 1);
if (HIP_LAUNCH_BLOCKING) {
tprintf(DB_SYNC, "LAUNCH_BLOCKING for completion of hipMemcpyAsync(%zu)\n", sizeBytes);
this->wait(crit);
}
if(trueAsync == true){
ihipCommand_t commandType;
hsa_agent_t srcAgent, dstAgent;
setAsyncCopyAgents(kind, &commandType, &srcAgent, &dstAgent);
hsa_signal_t depSignal;
int depSignalCnt = preCopyCommand(crit, ihip_signal, &depSignal, commandType);
tprintf (DB_SYNC, " copy-async, waitFor=%lu completion=#%lu(%lu)\n", depSignalCnt? depSignal.handle:0x0, ihip_signal->_sigId, ihip_signal->_hsaSignal.handle);
hsa_status_t hsa_status = hsa_amd_memory_async_copy(dst, dstAgent, src, srcAgent, sizeBytes, depSignalCnt, depSignalCnt ? &depSignal:0x0, ihip_signal->_hsaSignal);
if (hsa_status == HSA_STATUS_SUCCESS) {
if (HIP_LAUNCH_BLOCKING) {
tprintf(DB_SYNC, "LAUNCH_BLOCKING for completion of hipMemcpyAsync(%zu)\n", sizeBytes);
this->wait(crit);
}
} else {
// This path can be hit if src or dst point to unpinned host memory.
// TODO-stream - does async-copy fall back to sync if input pointers are not pinned?
throw ihipException(hipErrorInvalidValue);
}
} else {
copySync(crit, dst, src, sizeBytes, kind);
locked_copySync(dst, src, sizeBytes, kind);
}
}
}
@@ -1973,9 +1520,6 @@ void ihipStream_t::copyAsync(void* dst, const void* src, size_t sizeBytes, unsig
//-------------------------------------------------------------------------------------------------
// HCC-specific accessor functions:
/**
* @return #hipSuccess, #hipErrorInvalidDevice
*/
//---
hipError_t hipHccGetAccelerator(int deviceId, hc::accelerator *acc)
{
@@ -1993,9 +1537,6 @@ hipError_t hipHccGetAccelerator(int deviceId, hc::accelerator *acc)
}
/**
* @return #hipSuccess
*/
//---
hipError_t hipHccGetAcceleratorView(hipStream_t stream, hc::accelerator_view **av)
{
@@ -2012,16 +1553,5 @@ hipError_t hipHccGetAcceleratorView(hipStream_t stream, hc::accelerator_view **a
return ihipLogStatus(err);
}
// TODO - review signal / error reporting code.
// TODO - describe naming convention. ihip _. No accessors. No early returns from functions. Set status to success at top, only set error codes in implementation. No tabs.
// Caps convention _ or camelCase
// if { }
// Should use ihip* data structures inside code rather than app-facing hip. For example, use ihipCtx_t (rather than hipDevice_t), ihipStream_t (rather than hipStream_t).
// locked_
// TODO - describe MT strategy
//
//// TODO - add identifier numbers for streams and devices to help with debugging.
#if ONE_OBJECT_FILE
#include "unpinned_copy_engine.cpp"
#endif
//TODO - add a contect sequence number for debug. Print operator<< ctx:0.1 (device.ctx)
hipamd/src/hip_memory.cpp
+15 -27
View File
@@ -104,7 +104,7 @@ hipError_t hipMalloc(void** ptr, size_t sizeBytes)
if (sizeBytes == 0)
{
*ptr = NULL;
return ihipLogStatus(hip_status);
return ihipLogStatus(hipSuccess);
}
auto ctx = ihipGetTlsDefaultCtx();
@@ -172,20 +172,6 @@ hipError_t hipHostMalloc(void** ptr, size_t sizeBytes, unsigned int flags)
return ihipLogStatus(hip_status);
}
//---
// TODO - remove me, this is deprecated.
hipError_t hipHostAlloc(void** ptr, size_t sizeBytes, unsigned int flags)
{
return hipHostMalloc(ptr, sizeBytes, flags);
};
//---
// TODO - remove me, this is deprecated.
hipError_t hipMallocHost(void** ptr, size_t sizeBytes)
{
return hipHostMalloc(ptr, sizeBytes, 0);
}
// width in bytes
hipError_t hipMallocPitch(void** ptr, size_t* pitch, size_t width, size_t height)
{
@@ -522,7 +508,7 @@ hipError_t hipMemcpyAsync(void* dst, const void* src, size_t sizeBytes, hipMemcp
e= hipErrorInvalidValue;
} else if (stream) {
try {
stream->copyAsync(dst, src, sizeBytes, kind);
stream->locked_copyAsync(dst, src, sizeBytes, kind);
}
catch (ihipException ex) {
e = ex._code;
@@ -548,7 +534,7 @@ hipError_t hipMemcpyHtoDAsync(hipDeviceptr_t dst, void* src, size_t sizeBytes, h
e= hipErrorInvalidValue;
} else if (stream) {
try {
stream->copyAsync((void*)dst, src, sizeBytes, kind);
stream->locked_copyAsync((void*)dst, src, sizeBytes, kind);
}
catch (ihipException ex) {
e = ex._code;
@@ -575,7 +561,7 @@ hipError_t hipMemcpyDtoDAsync(hipDeviceptr_t dst, hipDeviceptr_t src, size_t siz
e= hipErrorInvalidValue;
} else if (stream) {
try {
stream->copyAsync((void*)dst, (void*)src, sizeBytes, kind);
stream->locked_copyAsync((void*)dst, (void*)src, sizeBytes, kind);
}
catch (ihipException ex) {
e = ex._code;
@@ -601,7 +587,7 @@ hipError_t hipMemcpyDtoHAsync(void* dst, hipDeviceptr_t src, size_t sizeBytes, h
e= hipErrorInvalidValue;
} else if (stream) {
try {
stream->copyAsync(dst, (void*)src, sizeBytes, kind);
stream->locked_copyAsync(dst, (void*)src, sizeBytes, kind);
}
catch (ihipException ex) {
e = ex._code;
@@ -613,6 +599,7 @@ hipError_t hipMemcpyDtoHAsync(void* dst, hipDeviceptr_t src, size_t sizeBytes, h
return ihipLogStatus(e);
}
// TODO - review and optimize
hipError_t hipMemcpy2D(void* dst, size_t dpitch, const void* src, size_t spitch,
size_t width, size_t height, hipMemcpyKind kind) {
@@ -785,7 +772,7 @@ hipError_t hipMemsetAsync(void* dst, int value, size_t sizeBytes, hipStream_t s
}
}
stream->lockclose_postKernelCommand(cf);
stream->lockclose_postKernelCommand(&crit->_av);
if (HIP_LAUNCH_BLOCKING) {
@@ -835,9 +822,10 @@ hipError_t hipMemset(void* dst, int value, size_t sizeBytes )
e = hipErrorInvalidValue;
}
}
// TODO - is hipMemset supposed to be async?
cf.wait();
stream->lockclose_postKernelCommand(cf);
stream->lockclose_postKernelCommand(&crit->_av);
if (HIP_LAUNCH_BLOCKING) {
@@ -864,6 +852,9 @@ hipError_t hipMemGetInfo (size_t *free, size_t *total)
if (total) {
*total = device->_props.totalGlobalMem;
}
else {
e = hipErrorInvalidValue;
}
if (free) {
// TODO - replace with kernel-level for reporting free memory:
@@ -872,6 +863,9 @@ hipError_t hipMemGetInfo (size_t *free, size_t *total)
*free = device->_props.totalGlobalMem - deviceMemSize;
}
else {
e = hipErrorInvalidValue;
}
} else {
e = hipErrorInvalidDevice;
@@ -934,12 +928,6 @@ hipError_t hipHostFree(void* ptr)
return ihipLogStatus(hipStatus);
};
// TODO - deprecated function.
hipError_t hipFreeHost(void* ptr)
{
return hipHostFree(ptr);
}
hipError_t hipFreeArray(hipArray* array)
{
HIP_INIT_API(array);
hipamd/src/hip_peer.cpp
+11 -23
View File
@@ -47,11 +47,7 @@ hipError_t hipDeviceCanAccessPeer (int* canAccessPeer, hipCtx_t thisCtx, hipCtx_
if (thisCtx == peerCtx) {
*canAccessPeer = 0;
} else {
#if USE_PEER_TO_PEER>=2
*canAccessPeer = peerCtx->getDevice()->_acc.get_is_peer(thisCtx->getDevice()->_acc);
#else
*canAccessPeer = 0;
#endif
}
} else {
@@ -69,18 +65,11 @@ hipError_t hipDeviceCanAccessPeer (int* canAccessPeer, hipCtx_t thisCtx, hipCtx_
// Remove this device from peer device peerlist.
hipError_t ihipDisablePeerAccess (hipCtx_t peerCtx)
{
HIP_INIT_API(peerCtx);
hipError_t err = hipSuccess;
auto thisCtx = ihipGetTlsDefaultCtx();
if ((thisCtx != NULL) && (peerCtx != NULL)) {
#if USE_PEER_TO_PEER>=2
// Return true if thisCtx can access peerCtx's memory:
bool canAccessPeer = peerCtx->getDevice()->_acc.get_is_peer(thisCtx->getDevice()->_acc);
#else
bool canAccessPeer = 0;
#endif
if (! canAccessPeer) {
err = hipErrorInvalidDevice; // P2P not allowed between these devices.
@@ -90,10 +79,8 @@ hipError_t ihipDisablePeerAccess (hipCtx_t peerCtx)
LockedAccessor_CtxCrit_t peerCrit(peerCtx->criticalData());
bool changed = peerCrit->removePeer(thisCtx);
if (changed) {
#if USE_PEER_TO_PEER>=3
// Update the peers for all memory already saved in the tracker:
am_memtracker_update_peers(peerCtx->getDevice()->_acc, peerCrit->peerCnt(), peerCrit->peerAgents());
#endif
} else {
err = hipErrorPeerAccessNotEnabled; // never enabled P2P access.
}
@@ -102,16 +89,15 @@ hipError_t ihipDisablePeerAccess (hipCtx_t peerCtx)
err = hipErrorInvalidDevice;
}
return ihipLogStatus(err);
return err;
};
//---
// Allow the current device to see all memory allocated on peerDevice.
// Allow the current device to see all memory allocated on peerCtx.
// This should add this device to the peer-device peer list.
hipError_t ihipEnablePeerAccess (hipCtx_t peerCtx, unsigned int flags)
{
HIP_INIT_API(peerCtx, flags);
hipError_t err = hipSuccess;
if (flags != 0) {
@@ -122,11 +108,10 @@ hipError_t ihipEnablePeerAccess (hipCtx_t peerCtx, unsigned int flags)
err = hipErrorInvalidDevice; // Can't enable peer access to self.
} else if ((thisCtx != NULL) && (peerCtx != NULL)) {
LockedAccessor_CtxCrit_t peerCrit(peerCtx->criticalData());
// Add thisCtx to peerCtx's access list so that new allocations on peer will be made visible to this device:
bool isNewPeer = peerCrit->addPeer(thisCtx);
if (isNewPeer) {
#if USE_PEER_TO_PEER>=3
am_memtracker_update_peers(peerCtx->getDevice()->_acc, peerCrit->peerCnt(), peerCrit->peerAgents());
#endif
} else {
err = hipErrorPeerAccessAlreadyEnabled;
}
@@ -135,7 +120,7 @@ hipError_t ihipEnablePeerAccess (hipCtx_t peerCtx, unsigned int flags)
}
}
return ihipLogStatus(err);
return err;
}
@@ -144,6 +129,7 @@ hipError_t hipMemcpyPeer (void* dst, hipCtx_t dstCtx, const void* src, hipCtx_t
{
HIP_INIT_API(dst, dstCtx, src, srcCtx, sizeBytes);
// 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);
};
@@ -153,6 +139,8 @@ hipError_t hipMemcpyPeer (void* dst, hipCtx_t dstCtx, const void* src, hipCtx_t
hipError_t hipMemcpyPeerAsync (void* dst, hipCtx_t dstDevice, const void* src, hipCtx_t srcDevice, size_t sizeBytes, hipStream_t stream)
{
HIP_INIT_API(dst, dstDevice, src, srcDevice, sizeBytes, stream);
// 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);
};
@@ -175,7 +163,7 @@ hipError_t hipDeviceDisablePeerAccess (int peerDeviceId)
{
HIP_INIT_API(peerDeviceId);
return ihipDisablePeerAccess(ihipGetPrimaryCtx(peerDeviceId));
return ihipLogStatus(ihipDisablePeerAccess(ihipGetPrimaryCtx(peerDeviceId)));
}
@@ -183,7 +171,7 @@ hipError_t hipDeviceEnablePeerAccess (int peerDeviceId, unsigned int flags)
{
HIP_INIT_API(peerDeviceId, flags);
return ihipEnablePeerAccess(ihipGetPrimaryCtx(peerDeviceId), flags);
return ihipLogStatus(ihipEnablePeerAccess(ihipGetPrimaryCtx(peerDeviceId), flags));
}
@@ -204,12 +192,12 @@ hipError_t hipCtxEnablePeerAccess (hipCtx_t peerCtx, unsigned int flags)
{
HIP_INIT_API(peerCtx, flags);
return ihipEnablePeerAccess(peerCtx, flags);
return ihipLogStatus(ihipEnablePeerAccess(peerCtx, flags));
}
hipError_t hipCtxDisablePeerAccess (hipCtx_t peerCtx)
{
HIP_INIT_API(peerCtx);
return ihipDisablePeerAccess(peerCtx);
return ihipLogStatus(ihipDisablePeerAccess(peerCtx));
}
hipamd/src/hip_stream.cpp
-9
View File
@@ -74,12 +74,6 @@ hipError_t hipStreamCreate(hipStream_t *stream)
}
#if USE_AV_COPY==0
//---
/**
* @bug This function conservatively waits for all work in the specified stream to complete.
*/
#endif
hipError_t hipStreamWaitEvent(hipStream_t stream, hipEvent_t event, unsigned int flags)
{
HIP_INIT_API(stream, event, flags);
@@ -93,14 +87,11 @@ hipError_t hipStreamWaitEvent(hipStream_t stream, hipEvent_t event, unsigned int
bool fastWait = false;
#if USE_AV_COPY
if (stream != hipStreamNull) {
stream->locked_waitEvent(event);
fastWait = true; // don't use the slow host-side synchronization.
}
// TODO - clean up if/else logic when USE_AV_COPY enabled.
#endif
if (!fastWait) {
// TODO-hcc Convert to use create_blocking_marker(...) functionality.
hipamd/src/unpinned_copy_engine.cpp
-413
View File
@@ -1,413 +0,0 @@
/*
Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved.
Permission is hereby granted, free of charge, to any person obtaining a copy
of this software and associated documentation files (the "Software"), to deal
in the Software without restriction, including without limitation the rights
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
copies of the Software, and to permit persons to whom the Software is
furnished to do so, subject to the following conditions:
The above copyright notice and this permission notice shall be included in
all copies or substantial portions of the Software.
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANNTY OF ANY KIND, EXPRESS OR
IMPLIED, INNCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANNY CLAIM, DAMAGES OR OTHER
LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
#include <hc_am.hpp>
#include "hsa_ext_amd.h"
#include "hcc_detail/unpinned_copy_engine.h"
#ifdef HIP_HCC
#include "hcc_detail/hip_runtime.h"
#include "hcc_detail/hip_hcc.h"
#define THROW_ERROR(e) throw ihipException(e)
#else
#define THROW_ERROR(e) throw
#define tprintf(trace_level, ...)
#endif
void errorCheck(hsa_status_t hsa_error_code, int line_num, std::string str) {
if ((hsa_error_code != HSA_STATUS_SUCCESS)&& (hsa_error_code != HSA_STATUS_INFO_BREAK)) {
printf("HSA reported error!\n In file: %s\nAt line: %d\n", str.c_str(),line_num);
}
}
#define ErrorCheck(x) errorCheck(x, __LINE__, __FILE__)
hsa_amd_memory_pool_t sys_pool_;
hsa_status_t findGlobalPool(hsa_amd_memory_pool_t pool, void* data) {
if (NULL == data) {
return HSA_STATUS_ERROR_INVALID_ARGUMENT;
}
hsa_status_t err;
hsa_amd_segment_t segment;
uint32_t flag;
err = hsa_amd_memory_pool_get_info(pool, HSA_AMD_MEMORY_POOL_INFO_SEGMENT, &segment);
ErrorCheck(err);
err = hsa_amd_memory_pool_get_info(pool, HSA_AMD_MEMORY_POOL_INFO_GLOBAL_FLAGS, &flag);
ErrorCheck(err);
if ((HSA_AMD_SEGMENT_GLOBAL == segment) &&
(flag & HSA_AMD_MEMORY_POOL_GLOBAL_FLAG_COARSE_GRAINED)) {
*((hsa_amd_memory_pool_t*)data) = pool;
}
return HSA_STATUS_SUCCESS;
}
//-------------------------------------------------------------------------------------------------
UnpinnedCopyEngine::UnpinnedCopyEngine(hsa_agent_t hsaAgent, hsa_agent_t cpuAgent, size_t bufferSize, int numBuffers, int thresholdH2DDirectStaging,int thresholdH2DStagingPinInPlace,int thresholdD2H) :
_hsaAgent(hsaAgent),
_cpuAgent(cpuAgent),
_bufferSize(bufferSize),
_numBuffers(numBuffers > _max_buffers ? _max_buffers : numBuffers),
_hipH2DTransferThresholdDirectOrStaging(thresholdH2DDirectStaging),
_hipH2DTransferThresholdStagingOrPininplace(thresholdH2DStagingPinInPlace),
_hipD2HTransferThreshold(thresholdD2H)
{
hsa_status_t err = hsa_amd_agent_iterate_memory_pools(_cpuAgent, findGlobalPool, &sys_pool_);
ErrorCheck(err);
for (int i=0; i<_numBuffers; i++) {
// TODO - experiment with alignment here.
err = hsa_amd_memory_pool_allocate(sys_pool_, _bufferSize, 0, (void**)(&_pinnedStagingBuffer[i]));
ErrorCheck(err);
if ((err != HSA_STATUS_SUCCESS) || (_pinnedStagingBuffer[i] == NULL)) {
THROW_ERROR(hipErrorMemoryAllocation);
}
err = hsa_amd_agents_allow_access(1, &hsaAgent, NULL, _pinnedStagingBuffer[i]);
ErrorCheck(err);
hsa_signal_create(0, 0, NULL, &_completion_signal[i]);
hsa_signal_create(0, 0, NULL, &_completion_signal2[i]);
}
};
//---
UnpinnedCopyEngine::~UnpinnedCopyEngine()
{
for (int i=0; i<_numBuffers; i++) {
if (_pinnedStagingBuffer[i]) {
hsa_amd_memory_pool_free(_pinnedStagingBuffer[i]);
_pinnedStagingBuffer[i] = NULL;
}
hsa_signal_destroy(_completion_signal[i]);
hsa_signal_destroy(_completion_signal2[i]);
}
}
//---
//Copies sizeBytes from src to dst, using either a copy to a staging buffer or a staged pin-in-place strategy
//IN: dst - dest pointer - must be accessible from host CPU.
//IN: src - src pointer for copy. Must be accessible from agent this buffer is associated with (via _hsaAgent)
//IN: waitFor - hsaSignal to wait for - the copy will begin only when the specified dependency is resolved. May be NULL indicating no dependency.
void UnpinnedCopyEngine::CopyHostToDevicePinInPlace(void* dst, const void* src, size_t sizeBytes, hsa_signal_t *waitFor)
{
std::lock_guard<std::mutex> l (_copy_lock);
const char *srcp = static_cast<const char*> (src);
char *dstp = static_cast<char*> (dst);
for (int i=0; i<_numBuffers; i++) {
hsa_signal_store_relaxed(_completion_signal[i], 0);
}
if (sizeBytes >= UINT64_MAX/2) {
THROW_ERROR (hipErrorInvalidValue);
}
int bufferIndex = 0;
size_t theseBytes= sizeBytes;
//tprintf (DB_COPY2, "H2D: waiting... on completion signal handle=%lu\n", _completion_signal[bufferIndex].handle);
//hsa_signal_wait_acquire(_completion_signal[bufferIndex], HSA_SIGNAL_CONDITION_LT, 1, UINT64_MAX, HSA_WAIT_STATE_ACTIVE);
//void * masked_srcp = (void*) ((uintptr_t)srcp & (uintptr_t)(~0x3f)) ; // TODO
void *locked_srcp;
//hsa_status_t hsa_status = hsa_amd_memory_lock(masked_srcp, theseBytes, &_hsaAgent, 1, &locked_srcp);
hsa_status_t hsa_status = hsa_amd_memory_lock(const_cast<char*> (srcp), theseBytes, &_hsaAgent, 1, &locked_srcp);
//tprintf (DB_COPY2, "H2D: bytesRemaining=%zu: pin-in-place:%p+%zu bufferIndex[%d]\n", bytesRemaining, srcp, theseBytes, bufferIndex);
//printf ("status=%x srcp=%p, masked_srcp=%p, locked_srcp=%p\n", hsa_status, srcp, masked_srcp, locked_srcp);
if (hsa_status != HSA_STATUS_SUCCESS) {
THROW_ERROR (hipErrorRuntimeMemory);
}
hsa_signal_store_relaxed(_completion_signal[bufferIndex], 1);
hsa_status = hsa_amd_memory_async_copy(dstp, _hsaAgent, locked_srcp, _cpuAgent, theseBytes, waitFor ? 1:0, waitFor, _completion_signal[bufferIndex]);
//tprintf (DB_COPY2, "H2D: bytesRemaining=%zu: async_copy %zu bytes %p to %p status=%x\n", bytesRemaining, theseBytes, _pinnedStagingBuffer[bufferIndex], dstp, hsa_status);
if (hsa_status != HSA_STATUS_SUCCESS) {
THROW_ERROR (hipErrorRuntimeMemory);
}
tprintf (DB_COPY2, "H2D: waiting... on completion signal handle=%lu\n", _completion_signal[bufferIndex].handle);
hsa_signal_wait_acquire(_completion_signal[bufferIndex], HSA_SIGNAL_CONDITION_LT, 1, UINT64_MAX, HSA_WAIT_STATE_ACTIVE);
hsa_amd_memory_unlock(const_cast<char*> (srcp));
// Assume subsequent commands are dependent on previous and don't need dependency after first copy submitted, HIP_ONESHOT_COPY_DEP=1
waitFor = NULL;
}
//---
//Copies sizeBytes from src to dst, using either a copy to a staging buffer or a staged pin-in-place strategy
//IN: dst - dest pointer - must be accessible from host CPU.
//IN: src - src pointer for copy. Must be accessible from agent this buffer is associated with (via _hsaAgent)
//IN: waitFor - hsaSignal to wait for - the copy will begin only when the specified dependency is resolved. May be NULL indicating no dependency.
void UnpinnedCopyEngine::CopyHostToDevice(int tempIndex,int isLargeBar,void* dst, const void* src, size_t sizeBytes, hsa_signal_t *waitFor)
{
if((tempIndex==1)&&(isLargeBar)&&(sizeBytes < _hipH2DTransferThresholdDirectOrStaging)){
memcpy(dst,src,sizeBytes);
std::atomic_thread_fence(std::memory_order_release);
}
else if((tempIndex==1) && (sizeBytes > _hipH2DTransferThresholdStagingOrPininplace)){
CopyHostToDevicePinInPlace(dst, src, sizeBytes, waitFor);
}
else
{
std::lock_guard<std::mutex> l (_copy_lock);
const char *srcp = static_cast<const char*> (src);
char *dstp = static_cast<char*> (dst);
for (int i=0; i<_numBuffers; i++) {
hsa_signal_store_relaxed(_completion_signal[i], 0);
}
if (sizeBytes >= UINT64_MAX/2) {
THROW_ERROR (hipErrorInvalidValue);
}
int bufferIndex = 0;
for (int64_t bytesRemaining=sizeBytes; bytesRemaining>0 ; bytesRemaining -= _bufferSize) {
size_t theseBytes = (bytesRemaining > _bufferSize) ? _bufferSize : bytesRemaining;
tprintf (DB_COPY2, "H2D: waiting... on completion signal handle=%lu\n", _completion_signal[bufferIndex].handle);
hsa_signal_wait_acquire(_completion_signal[bufferIndex], HSA_SIGNAL_CONDITION_LT, 1, UINT64_MAX, HSA_WAIT_STATE_ACTIVE);
tprintf (DB_COPY2, "H2D: bytesRemaining=%zu: copy %zu bytes %p to stagingBuf[%d]:%p\n", bytesRemaining, theseBytes, srcp, bufferIndex, _pinnedStagingBuffer[bufferIndex]);
// TODO - use uncached memcpy, someday.
memcpy(_pinnedStagingBuffer[bufferIndex], srcp, theseBytes);
hsa_signal_store_relaxed(_completion_signal[bufferIndex], 1);
hsa_status_t hsa_status = hsa_amd_memory_async_copy(dstp, _hsaAgent, _pinnedStagingBuffer[bufferIndex], _cpuAgent, theseBytes, waitFor ? 1:0, waitFor, _completion_signal[bufferIndex]);
tprintf (DB_COPY2, "H2D: bytesRemaining=%zu: async_copy %zu bytes %p to %p status=%x\n", bytesRemaining, theseBytes, _pinnedStagingBuffer[bufferIndex], dstp, hsa_status);
if (hsa_status != HSA_STATUS_SUCCESS) {
THROW_ERROR ((hipErrorRuntimeMemory));
}
srcp += theseBytes;
dstp += theseBytes;
if (++bufferIndex >= _numBuffers) {
bufferIndex = 0;
}
// Assume subsequent commands are dependent on previous and don't need dependency after first copy submitted, HIP_ONESHOT_COPY_DEP=1
waitFor = NULL;
}
for (int i=0; i<_numBuffers; i++) {
hsa_signal_wait_acquire(_completion_signal[i], HSA_SIGNAL_CONDITION_LT, 1, UINT64_MAX, HSA_WAIT_STATE_ACTIVE);
}
}
}
void UnpinnedCopyEngine::CopyDeviceToHostPinInPlace(void* dst, const void* src, size_t sizeBytes, hsa_signal_t *waitFor)
{
std::lock_guard<std::mutex> l (_copy_lock);
const char *srcp = static_cast<const char*> (src);
char *dstp = static_cast<char*> (dst);
for (int i=0; i<_numBuffers; i++) {
hsa_signal_store_relaxed(_completion_signal[i], 0);
}
if (sizeBytes >= UINT64_MAX/2) {
THROW_ERROR (hipErrorInvalidValue);
}
int bufferIndex = 0;
size_t theseBytes= sizeBytes;
void *locked_destp;
hsa_status_t hsa_status = hsa_amd_memory_lock(const_cast<char*> (dstp), theseBytes, &_hsaAgent, 1, &locked_destp);
if (hsa_status != HSA_STATUS_SUCCESS) {
THROW_ERROR (hipErrorRuntimeMemory);
}
hsa_signal_store_relaxed(_completion_signal[bufferIndex], 1);
hsa_status = hsa_amd_memory_async_copy(locked_destp,_cpuAgent , srcp, _hsaAgent, theseBytes, waitFor ? 1:0, waitFor, _completion_signal[bufferIndex]);
if (hsa_status != HSA_STATUS_SUCCESS) {
THROW_ERROR (hipErrorRuntimeMemory);
}
tprintf (DB_COPY2, "D2H: waiting... on completion signal handle=%lu\n", _completion_signal[bufferIndex].handle);
hsa_signal_wait_acquire(_completion_signal[bufferIndex], HSA_SIGNAL_CONDITION_LT, 1, UINT64_MAX, HSA_WAIT_STATE_ACTIVE);
hsa_amd_memory_unlock(const_cast<char*> (dstp));
// Assume subsequent commands are dependent on previous and don't need dependency after first copy submitted, HIP_ONESHOT_COPY_DEP=1
waitFor = NULL;
}
//---
//Copies sizeBytes from src to dst, using either a copy to a staging buffer or a staged pin-in-place strategy
//IN: dst - dest pointer - must be accessible from agent this buffer is associated with (via _hsaAgent).
//IN: src - src pointer for copy. Must be accessible from host CPU.
//IN: waitFor - hsaSignal to wait for - the copy will begin only when the specified dependency is resolved. May be NULL indicating no dependency.
void UnpinnedCopyEngine::CopyDeviceToHost(int tempIndex,void* dst, const void* src, size_t sizeBytes, hsa_signal_t *waitFor)
{
if((tempIndex==1) && (sizeBytes> _hipD2HTransferThreshold)){
CopyDeviceToHostPinInPlace(dst, src, sizeBytes, waitFor);
}
else
{
std::lock_guard<std::mutex> l (_copy_lock);
const char *srcp0 = static_cast<const char*> (src);
char *dstp1 = static_cast<char*> (dst);
for (int i=0; i<_numBuffers; i++) {
hsa_signal_store_relaxed(_completion_signal[i], 0);
}
if (sizeBytes >= UINT64_MAX/2) {
THROW_ERROR (hipErrorInvalidValue);
}
int64_t bytesRemaining0 = sizeBytes; // bytes to copy from dest into staging buffer.
int64_t bytesRemaining1 = sizeBytes; // bytes to copy from staging buffer into final dest
while (bytesRemaining1 > 0)
{
// First launch the async copies to copy from dest to host
for (int bufferIndex = 0; (bytesRemaining0>0) && (bufferIndex < _numBuffers); bytesRemaining0 -= _bufferSize, bufferIndex++) {
size_t theseBytes = (bytesRemaining0 > _bufferSize) ? _bufferSize : bytesRemaining0;
tprintf (DB_COPY2, "D2H: bytesRemaining0=%zu async_copy %zu bytes src:%p to staging:%p\n", bytesRemaining0, theseBytes, srcp0, _pinnedStagingBuffer[bufferIndex]);
hsa_signal_store_relaxed(_completion_signal[bufferIndex], 1);
hsa_status_t hsa_status = hsa_amd_memory_async_copy(_pinnedStagingBuffer[bufferIndex], _cpuAgent, srcp0, _hsaAgent, theseBytes, waitFor ? 1:0, waitFor, _completion_signal[bufferIndex]);
if (hsa_status != HSA_STATUS_SUCCESS) {
THROW_ERROR (hipErrorRuntimeMemory);
}
srcp0 += theseBytes;
// Assume subsequent commands are dependent on previous and don't need dependency after first copy submitted, HIP_ONESHOT_COPY_DEP=1
waitFor = NULL;
}
// Now unload the staging buffers:
for (int bufferIndex=0; (bytesRemaining1>0) && (bufferIndex < _numBuffers); bytesRemaining1 -= _bufferSize, bufferIndex++) {
size_t theseBytes = (bytesRemaining1 > _bufferSize) ? _bufferSize : bytesRemaining1;
tprintf (DB_COPY2, "D2H: wait_completion[%d] bytesRemaining=%zu\n", bufferIndex, bytesRemaining1);
hsa_signal_wait_acquire(_completion_signal[bufferIndex], HSA_SIGNAL_CONDITION_LT, 1, UINT64_MAX, HSA_WAIT_STATE_ACTIVE);
tprintf (DB_COPY2, "D2H: bytesRemaining1=%zu copy %zu bytes stagingBuf[%d]:%p to dst:%p\n", bytesRemaining1, theseBytes, bufferIndex, _pinnedStagingBuffer[bufferIndex], dstp1);
memcpy(dstp1, _pinnedStagingBuffer[bufferIndex], theseBytes);
dstp1 += theseBytes;
}
}
}
}
//---
//Copies sizeBytes from src to dst, using either a copy to a staging buffer or a staged pin-in-place strategy
//IN: dst - dest pointer - must be accessible from agent this buffer is associated with (via _hsaAgent).
//IN: src - src pointer for copy. Must be accessible from host CPU.
//IN: waitFor - hsaSignal to wait for - the copy will begin only when the specified dependency is resolved. May be NULL indicating no dependency.
void UnpinnedCopyEngine::CopyPeerToPeer(void* dst, hsa_agent_t dstAgent, const void* src, hsa_agent_t srcAgent, size_t sizeBytes, hsa_signal_t *waitFor)
{
std::lock_guard<std::mutex> l (_copy_lock);
const char *srcp0 = static_cast<const char*> (src);
char *dstp1 = static_cast<char*> (dst);
for (int i=0; i<_numBuffers; i++) {
hsa_signal_store_relaxed(_completion_signal[i], 0);
hsa_signal_store_relaxed(_completion_signal2[i], 0);
}
if (sizeBytes >= UINT64_MAX/2) {
THROW_ERROR (hipErrorInvalidValue);
}
int64_t bytesRemaining0 = sizeBytes; // bytes to copy from dest into staging buffer.
int64_t bytesRemaining1 = sizeBytes; // bytes to copy from staging buffer into final dest
while (bytesRemaining1 > 0) {
// First launch the async copies to copy from dest to host
for (int bufferIndex = 0; (bytesRemaining0>0) && (bufferIndex < _numBuffers); bytesRemaining0 -= _bufferSize, bufferIndex++) {
size_t theseBytes = (bytesRemaining0 > _bufferSize) ? _bufferSize : bytesRemaining0;
// Wait to make sure we are not overwriting a buffer before it has been drained:
hsa_signal_wait_acquire(_completion_signal2[bufferIndex], HSA_SIGNAL_CONDITION_LT, 1, UINT64_MAX, HSA_WAIT_STATE_ACTIVE);
tprintf (DB_COPY2, "P2P: bytesRemaining0=%zu async_copy %zu bytes src:%p to staging:%p\n", bytesRemaining0, theseBytes, srcp0, _pinnedStagingBuffer[bufferIndex]);
hsa_signal_store_relaxed(_completion_signal[bufferIndex], 1);
hsa_status_t hsa_status = hsa_amd_memory_async_copy(_pinnedStagingBuffer[bufferIndex], _cpuAgent, srcp0, srcAgent, theseBytes, waitFor ? 1:0, waitFor, _completion_signal[bufferIndex]);
if (hsa_status != HSA_STATUS_SUCCESS) {
THROW_ERROR (hipErrorRuntimeMemory);
}
srcp0 += theseBytes;
// Assume subsequent commands are dependent on previous and don't need dependency after first copy submitted, HIP_ONESHOT_COPY_DEP=1
waitFor = NULL;
}
// Now unload the staging buffers:
for (int bufferIndex=0; (bytesRemaining1>0) && (bufferIndex < _numBuffers); bytesRemaining1 -= _bufferSize, bufferIndex++) {
size_t theseBytes = (bytesRemaining1 > _bufferSize) ? _bufferSize : bytesRemaining1;
tprintf (DB_COPY2, "P2P: wait_completion[%d] bytesRemaining=%zu\n", bufferIndex, bytesRemaining1);
bool hostWait = 0; // TODO - remove me
if (hostWait) {
// Host-side wait, should not be necessary:
hsa_signal_wait_acquire(_completion_signal[bufferIndex], HSA_SIGNAL_CONDITION_LT, 1, UINT64_MAX, HSA_WAIT_STATE_ACTIVE);
}
tprintf (DB_COPY2, "P2P: bytesRemaining1=%zu copy %zu bytes stagingBuf[%d]:%p to device:%p\n", bytesRemaining1, theseBytes, bufferIndex, _pinnedStagingBuffer[bufferIndex], dstp1);
hsa_signal_store_relaxed(_completion_signal2[bufferIndex], 1);
hsa_status_t hsa_status = hsa_amd_memory_async_copy(dstp1, dstAgent, _pinnedStagingBuffer[bufferIndex], _cpuAgent /*not used*/, theseBytes,
hostWait ? 0:1, hostWait ? NULL : &_completion_signal[bufferIndex],
_completion_signal2[bufferIndex]);
dstp1 += theseBytes;
}
}
// Wait for the staging-buffer to dest copies to complete:
for (int i=0; i<_numBuffers; i++) {
hsa_signal_wait_acquire(_completion_signal2[i], HSA_SIGNAL_CONDITION_LT, 1, UINT64_MAX, HSA_WAIT_STATE_ACTIVE);
}
}
hipamd/tests/README.md
+3 -1
View File
@@ -51,9 +51,11 @@ ctest -R Memcpy
```
### If a test fails:
### If a test fails - how to debug a test
Extract the commandline from the testing log:
(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
hipamd/tests/hip_tests.txt
+13
View File
@@ -0,0 +1,13 @@
cmake_minimum_required(VERSION 2.8.3)
project(hip_tests)
# Setup
set(HIP_PATH @CMAKE_INSTALL_PREFIX@)
set(ENV{HIP_PATH} ${HIP_PATH})
set(HIP_SRC_PATH @hip_SOURCE_DIR@)
set(CMAKE_MODULE_PATH "${HIP_PATH}/cmake" ${CMAKE_MODULE_PATH})
include(${HIP_SRC_PATH}/tests/hit/HIT.cmake)
# Add tests
include_directories(${HIP_SRC_PATH}/tests/src)
hit_add_directory_recursive(${HIP_SRC_PATH}/tests/src "directed_tests")
hipamd/tests/hit/HIT.cmake
+219
View File
@@ -0,0 +1,219 @@
include(CTest)
find_package(HIP REQUIRED)
#-------------------------------------------------------------------------------
# Helper macro to parse BUILD instructions
macro(PARSE_BUILD_COMMAND _target _sources _hipcc_options _hcc_options _nvcc_options _exclude_platforms _dir)
set(${_target})
set(${_sources})
set(${_hipcc_options})
set(${_hcc_options})
set(${_nvcc_options})
set(${_exclude_platforms})
set(_target_found FALSE)
set(_hipcc_options_found FALSE)
set(_hcc_options_found FALSE)
set(_nvcc_options_found FALSE)
set(_exclude_platforms_found FALSE)
foreach(arg ${ARGN})
if(NOT _target_found)
set(_target_found TRUE)
set(${_target} ${arg})
elseif("x${arg}" STREQUAL "xHIPCC_OPTIONS")
set(_hipcc_options_found TRUE)
set(_hcc_options_found FALSE)
set(_nvcc_options_found FALSE)
set(_exclude_platforms_found FALSE)
elseif("x${arg}" STREQUAL "xHCC_OPTIONS")
set(_hipcc_options_found FALSE)
set(_hcc_options_found TRUE)
set(_nvcc_options_found FALSE)
set(_exclude_platforms_found FALSE)
elseif("x${arg}" STREQUAL "xNVCC_OPTIONS")
set(_hipcc_options_found FALSE)
set(_hcc_options_found FALSE)
set(_nvcc_options_found TRUE)
set(_exclude_platforms_found FALSE)
elseif("x${arg}" STREQUAL "xEXCLUDE_HIP_PLATFORM")
set(_hipcc_options_found FALSE)
set(_hcc_options_found FALSE)
set(_nvcc_options_found FALSE)
set(_exclude_platforms_found TRUE)
else()
if(_hipcc_options_found)
list(APPEND ${_hipcc_options} ${arg})
elseif(_hcc_options_found)
list(APPEND ${_hcc_options} ${arg})
elseif(_nvcc_options_found)
list(APPEND ${_nvcc_options} ${arg})
elseif(_exclude_platforms_found)
set(${_exclude_platforms} ${arg})
else()
list(APPEND ${_sources} "${_dir}/${arg}")
endif()
endif()
endforeach()
endmacro()
# Helper macro to parse RUN instructions
macro(PARSE_RUN_COMMAND _target _arguments _exclude_platforms)
set(${_target})
set(${_arguments} " ")
set(${_exclude_platforms})
set(_target_found FALSE)
set(_exclude_platforms_found FALSE)
foreach(arg ${ARGN})
if(NOT _target_found)
set(_target_found TRUE)
set(${_target} ${arg})
elseif("x${arg}" STREQUAL "xEXCLUDE_HIP_PLATFORM")
set(_exclude_platforms_found TRUE)
else()
if(_exclude_platforms_found)
set(${_exclude_platforms} ${arg})
else()
list(APPEND ${_arguments} ${arg})
endif()
endif()
endforeach()
endmacro()
# Helper macro to parse RUN_NAMED instructions
macro(PARSE_RUN_NAMED_COMMAND _target _testname _arguments _exclude_platforms)
set(${_target})
set(${_arguments} " ")
set(${_exclude_platforms})
set(_target_found FALSE)
set(_testname_found FALSE)
set(_exclude_platforms_found FALSE)
foreach(arg ${ARGN})
if(NOT _target_found)
set(_target_found TRUE)
set(${_target} ${arg})
elseif(NOT _testname_found)
set(_testname_found TRUE)
set(${_testname} ${arg})
elseif("x${arg}" STREQUAL "xEXCLUDE_HIP_PLATFORM")
set(_exclude_platforms_found TRUE)
else()
if(_exclude_platforms_found)
set(${_exclude_platforms} ${arg})
else()
list(APPEND ${_arguments} ${arg})
endif()
endif()
endforeach()
endmacro()
# Helper macro to insert key/value pair into "hashmap"
macro(INSERT_INTO_MAP _map _key _value)
set("${_map}_${_key}" "${_value}")
endmacro()
# Helper macro to read key/value pair from "hashmap"
macro(READ_FROM_MAP _map _key _value)
set(${_value} "${${_map}_${_key}}")
endmacro()
# Helper macro to create a test
macro(MAKE_TEST exe)
string(REPLACE " " "" smush_args ${ARGN})
set(testname ${exe}${smush_args}.tst)
add_test(NAME ${testname} COMMAND ${PROJECT_BINARY_DIR}/${exe} ${ARGN})
set_tests_properties(${testname} PROPERTIES PASS_REGULAR_EXPRESSION "PASSED")
endmacro()
macro(MAKE_NAMED_TEST exe testname)
add_test(NAME ${testname} COMMAND ${PROJECT_BINARY_DIR}/${exe} ${ARGN})
set_tests_properties(${testname} PROPERTIES PASS_REGULAR_EXPRESSION "PASSED")
endmacro()
#-------------------------------------------------------------------------------
# Macro: HIT_ADD_FILES used to scan+add multiple files for testing.
macro(HIT_ADD_FILES _dir _label)
foreach (file ${ARGN})
# Build tests
execute_process(COMMAND ${HIP_SRC_PATH}/tests/hit/parser --buildCMDs ${file}
OUTPUT_VARIABLE _contents
ERROR_QUIET
WORKING_DIRECTORY ${_dir}
OUTPUT_STRIP_TRAILING_WHITESPACE)
string(REGEX REPLACE "\n" ";" _contents "${_contents}")
foreach(_cmd ${_contents})
string(REGEX REPLACE " " ";" _cmd "${_cmd}")
parse_build_command(_target _sources _hipcc_options _hcc_options _nvcc_options _exclude_platforms ${_dir} ${_cmd})
string(REGEX REPLACE "/" "." target ${_label}/${_target})
insert_into_map("_exclude" "${target}" "${_exclude_platforms}")
if(_exclude_platforms STREQUAL "all" OR _exclude_platforms STREQUAL ${HIP_PLATFORM})
else()
set_source_files_properties(${_sources} PROPERTIES HIP_SOURCE_PROPERTY_FORMAT 1)
hip_add_executable(${target} ${_sources} HIPCC_OPTIONS ${_hipcc_options} HCC_OPTIONS ${_hcc_options} NVCC_OPTIONS ${_nvcc_options})
set_target_properties(${target} PROPERTIES OUTPUT_NAME ${_target} RUNTIME_OUTPUT_DIRECTORY ${_label})
endif()
endforeach()
# Add tests
execute_process(COMMAND ${HIP_SRC_PATH}/tests/hit/parser --runCMDs ${file}
OUTPUT_VARIABLE _contents
ERROR_QUIET
WORKING_DIRECTORY ${_dir}
OUTPUT_STRIP_TRAILING_WHITESPACE)
string(REGEX REPLACE "\n" ";" _contents "${_contents}")
foreach(_cmd ${_contents})
string(REGEX REPLACE " " ";" _cmd "${_cmd}")
parse_run_command(_target _arguments _exclude_platforms ${_cmd})
string(REGEX REPLACE "/" "." target ${_label}/${_target})
read_from_map("_exclude" "${target}" _exclude_platforms_from_build)
if(_exclude_platforms STREQUAL "all" OR _exclude_platforms STREQUAL ${HIP_PLATFORM} OR
_exclude_platforms_from_build STREQUAL "all" OR _exclude_platforms_from_build STREQUAL ${HIP_PLATFORM})
else()
make_test(${_label}/${_target} ${_arguments})
endif()
endforeach()
# Add named tests
execute_process(COMMAND ${HIP_SRC_PATH}/tests/hit/parser --runNamedCMDs ${file}
OUTPUT_VARIABLE _contents
ERROR_QUIET
WORKING_DIRECTORY ${_dir}
OUTPUT_STRIP_TRAILING_WHITESPACE)
string(REGEX REPLACE "\n" ";" _contents "${_contents}")
foreach(_cmd ${_contents})
string(REGEX REPLACE " " ";" _cmd "${_cmd}")
parse_run_named_command(_target _testname _arguments _exclude_platforms ${_cmd})
string(REGEX REPLACE "/" "." target ${_label}/${_target})
read_from_map("_exclude" "${target}" _exclude_platforms_from_build)
if(_exclude_platforms STREQUAL "all" OR _exclude_platforms STREQUAL ${HIP_PLATFORM} OR
_exclude_platforms_from_build STREQUAL "all" OR _exclude_platforms_from_build STREQUAL ${HIP_PLATFORM})
else()
make_named_test(${_label}/${_target} ${_label}/${_testname}.tst ${_arguments})
endif()
endforeach()
endforeach()
endmacro()
# Macro: HIT_ADD_DIRECTORY to scan+add all files in a directory for testing
macro(HIT_ADD_DIRECTORY _dir _label)
execute_process(COMMAND ${CMAKE_COMMAND} -E make_directory ${_label} WORKING_DIRECTORY ${PROJECT_BINARY_DIR})
file(GLOB files "${_dir}/*.c*")
hit_add_files(${_dir} ${_label} ${files})
endmacro()
# Macro: HIT_ADD_DIRECTORY_RECURSIVE to scan+add all files in a directory+subdirectories for testing
macro(HIT_ADD_DIRECTORY_RECURSIVE _dir _label)
execute_process(COMMAND ${CMAKE_COMMAND} -E make_directory ${_label} WORKING_DIRECTORY ${PROJECT_BINARY_DIR})
file(GLOB children RELATIVE ${_dir} ${_dir}/*)
set(dirlist "")
foreach(child ${children})
if(IS_DIRECTORY ${_dir}/${child})
list(APPEND dirlist ${child})
else()
hit_add_files(${_dir} ${_label} ${child})
endif()
endforeach()
foreach(child ${dirlist})
hit_add_directory_recursive(${_dir}/${child} ${_label}/${child})
endforeach()
endmacro()
# vim: ts=4:sw=4:expandtab:smartindent
hipamd/tests/hit/parser
Executable
+95
View File
@@ -0,0 +1,95 @@
#!/usr/bin/perl -w
use 5.006; use v5.10.1;
use File::Basename;
use File::Spec;
# Scan input file for HIT information
sub parse_file {
my $file = shift;
(my $exe = $file) =~ s/\.[^.]+$//g;
my (@buildCMDs, @runCMDs, @runNamedCMDs);
if (open (SOURCE, '<:encoding(UTF-8)', "$file")) {
while (<SOURCE>) {
my $line=$_;
# Look for BUILD instructions
if ($line =~ /^ \* BUILD:/) {
$line =~ s/^ \* BUILD: //g; # Remove " * BUILD: "
$line =~ s/%s/$file/g; # Substitute %s -> filename
$line =~ s/%t/$exe/g; # Substitute %t -> targetname
$line =~ s/\R//g; # Remove line endings
push @buildCMDs, $line;
}
# Look for RUN instructions
if ($line =~ /^ \* RUN:/) {
$line =~ s/^ \* RUN: //g; # Remove " * RUN: "
$line =~ s/%s/$file/g; # Substitute %s -> filename
$line =~ s/%t/$exe/g; # Subsitute %t -> targetname
$line =~ s/\R//g; # Remove line endings
push @runCMDs, $line;
}
# Look for RUN_NAMED instructions
if ($line =~ /^ \* RUN_NAMED:/) {
$line =~ s/^ \* RUN_NAMED: //g; # Remove " * RUN_NAMED: "
$line =~ s/%s/$file/g; # Substitute %s -> filename
$line =~ s/%t/$exe/g; # Subsitute %t -> targetname
$line =~ s/\R//g; # Remove line endings
push @runNamedCMDs, $line;
}
}
close(SOURCE);
}
return (\@buildCMDs, \@runCMDs, \@runNamedCMDs);
}
# Exit if no arguments specified
if(scalar @ARGV == 0){
print "No Arguments passed, exiting ...\n";
exit(-1);
}
# Parse command
my @options = ();
my $retBuildCMDs = 0;
my $retRunCMDs = 0;
my $retRunNamedCMDs = 0;
foreach $arg (@ARGV) {
if ($retBuildCMDs or $retRunCMDs or $retRunNamedCMDs) {
push (@options, $arg);
}
if ($arg eq '--buildCMDs') {
$retBuildCMDs = 1;
}
if ($arg eq '--runCMDs') {
$retRunCMDs = 1;
}
if ($arg eq '--runNamedCMDs') {
$retRunNamedCMDs = 1;
}
}
# Atleast one command needs to be specified
if (($retBuildCMDs eq 0) and ($retRunCMDs eq 0) and ($retRunNamedCMDs eq 0)) {
die "Usage: $0 <--buildCMDs|--runCMDs|--runNamedCMDs> FILENAMEs\n";
}
# Iterate over input files
foreach $file (@options) {
# Convert absolute path to path relative to working directory
my $relfile = File::Spec->abs2rel($file);
my ($buildCMDs, $runCMDs, $runNamedCMDs) = parse_file("$relfile");
if ($retBuildCMDs) {
# print "BuildCMDs:\n";
print "$_\n" for @$buildCMDs;
}
if ($retRunCMDs) {
# print "RunCMDs:\n";
print "$_\n" for @$runCMDs;
}
if ($retRunNamedCMDs) {
# print "RunNamedCMDs:\n";
print "$_\n" for @$runNamedCMDs;
}
}
# vim: ts=4:sw=4:expandtab:smartindent
hipamd/tests/src/CMakeLists.txt
+2 -2
View File
@@ -184,7 +184,7 @@ build_hip_executable_libcpp(hipHcc hipHcc.cpp)
# __workweek fix.
#build_hip_executable_libcpp(hipPointerAttrib hipPointerAttrib.cpp)
build_hip_executable(hipHostAlloc hipHostAlloc.cpp)
build_hip_executable(hipHostMalloc hipHostMalloc.cpp)
build_hip_executable(hipHostGetFlags hipHostGetFlags.cpp)
build_hip_executable(hipHostRegister hipHostRegister.cpp)
build_hip_executable(hipRandomMemcpyAsync hipRandomMemcpyAsync.cpp)
@@ -208,7 +208,7 @@ make_test(hipLaunchParm " ")
#make_test(hipPointerAttrib " ")
make_test(hipHostAlloc " ")
make_test(hipHostMalloc " ")
# BS- comment out since test appears broken - asks for device pointer but pointer was never allocated.
#make_test(hipHostGetFlags " ")
make_test(hipHcc " ")
hipamd/tests/src/Functional/Negative/Device/hipDeviceGetAttribute.cpp
+21
View File
@@ -0,0 +1,21 @@
#include<hip/hip_runtime.h>
#include<hip/hip_runtime_api.h>
#include"hipDeviceUtil.h"
int main()
{
int pi;
int attr = 0;
// hipDeviceAttribute_t attr = hipDeviceAttributeMaxThreadsPerBlock;
HIP_CHECK(hipDeviceGetAttribute(nullptr, hipDeviceAttribute_t(attr), 0), hipDeviceGetAttribute);
HIP_CHECK(hipDeviceGetAttribute(&pi, hipDeviceAttribute_t(attr), 0), hipDeviceGetAttribute);
attr = -1;
HIP_CHECK(hipDeviceGetAttribute(nullptr, hipDeviceAttribute_t(attr), 0), hipDeviceGetAttribute);
HIP_CHECK(hipDeviceGetAttribute(&pi, hipDeviceAttribute_t(attr), 0), hipDeviceGetAttribute);
attr = 0;
HIP_CHECK(hipDeviceGetAttribute(nullptr, hipDeviceAttribute_t(attr), -1), hipDeviceGetAttribute);
HIP_CHECK(hipDeviceGetAttribute(&pi, hipDeviceAttribute_t(attr), -1), hipDeviceGetAttribute);
attr = -1;
HIP_CHECK(hipDeviceGetAttribute(nullptr, hipDeviceAttribute_t(attr), -1), hipDeviceGetAttribute);
HIP_CHECK(hipDeviceGetAttribute(&pi, hipDeviceAttribute_t(attr), -1), hipDeviceGetAttribute);
}
hipamd/tests/src/Functional/Negative/Device/hipDeviceUtil.h
+10
View File
@@ -0,0 +1,10 @@
#ifndef HIPDEVICEUTIL_H
#define HIPDEVICEUTIL_H
#include<hip/hip_runtime_api.h>
#include<iostream>
#define HIP_CHECK(status, func) \
std::cout<<#func<<" returned "<<hipGetErrorString(status)<<" in "<<__func__<<" at "<<__LINE__<<" in file "<<__FILE__<<std::endl;
#endif
hipamd/tests/src/Functional/Negative/Device/hipGetDevice.cpp
+10
View File
@@ -0,0 +1,10 @@
#include<hip/hip_runtime.h>
#include<hip/hip_runtime_api.h>
#include"hipDeviceUtil.h"
int main()
{
int device;
HIP_CHECK(hipGetDevice(NULL), hipGetDevice);
HIP_CHECK(hipGetDevice(&device), hipGetDevice);
}
hipamd/tests/src/Functional/Negative/Device/hipGetDeviceCount.cpp
+10
View File
@@ -0,0 +1,10 @@
#include<hip/hip_runtime.h>
#include<hip/hip_runtime_api.h>
#include"hipDeviceUtil.h"
int main()
{
int deviceCnt;
HIP_CHECK(hipGetDeviceCount(&deviceCnt), hipGetDeviceCount);
HIP_CHECK(hipGetDeviceCount(0), hipGetDeviceCount);
}
hipamd/tests/src/Functional/Negative/Device/hipGetDeviceProperties.cpp
+14
View File
@@ -0,0 +1,14 @@
#include<hip/hip_runtime.h>
#include<hip/hip_runtime_api.h>
#include"hipDeviceUtil.h"
int main()
{
hipDeviceProp_t props;
HIP_CHECK(hipGetDeviceProperties(&props, 0), hipGetDeviceProperties);
HIP_CHECK(hipGetDeviceProperties(NULL, 0), hipGetDeviceProperties);
HIP_CHECK(hipGetDeviceProperties(NULL, -1), hipGetDeviceProperties);
HIP_CHECK(hipGetDeviceProperties(&props, -1), hipGetDeviceProperties);
HIP_CHECK(hipGetDeviceProperties(NULL, 1024), hipGetDeviceProperties);
HIP_CHECK(hipGetDeviceProperties(&props, 1024), hipGetDeviceProperties);
}
hipamd/tests/src/Functional/Negative/Device/hipSetDevice.cpp
+10
View File
@@ -0,0 +1,10 @@
#include<hip/hip_runtime.h>
#include<hip/hip_runtime_api.h>
#include"hipDeviceUtil.h"
int main()
{
HIP_CHECK(hipSetDevice(0), hipSetDevice);
HIP_CHECK(hipSetDevice(1026), hipSetDevice);
HIP_CHECK(hipSetDevice(-1), hipSetDevice);
}
hipamd/tests/src/context/hipCtx_simple.cpp
+6
View File
@@ -20,6 +20,12 @@ 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 HCC_OPTIONS --stdlib=libc++
* RUN: %t
* HIT_END
*/
#include "hip_runtime.h"
#include "test_common.h"
hipamd/tests/src/deviceLib/hipMathFunctions.cpp
+11
View File
@@ -19,6 +19,17 @@ 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: %tHost %s hipSinglePrecisionMathHost.cpp hipDoublePrecisionMathHost.cpp ../test_common.cpp
* BUILD: %tDevice %s hipSinglePrecisionMathDevice.cpp hipDoublePrecisionMathDevice.cpp ../test_common.cpp
* BUILD: hipIntrinsics %s hipSinglePrecisionIntrinsics.cpp hipDoublePrecisionIntrinsics.cpp hipIntegerIntrinsics.cpp ../test_common.cpp
* RUN: %tHost
* RUN: %tDevice
* RUN: hipIntrinsics
* HIT_END
*/
#include "hip_runtime.h"
#include "test_common.h"
hipamd/tests/src/deviceLib/hipSimpleAtomicsTest.cpp
+6
View File
@@ -17,6 +17,12 @@ 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
* RUN: %t
* HIT_END
*/
// includes, system
#include <stdlib.h>
#include <stdio.h>
hipamd/tests/src/deviceLib/hipTestDevice.cpp
+7
View File
@@ -16,6 +16,13 @@ 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
* RUN: %t
* HIT_END
*/
#include"test_common.h"
#include"hip_runtime.h"
#include"hip_runtime_api.h"
hipamd/tests/src/deviceLib/hipTestDeviceDouble.cpp
+7
View File
@@ -16,6 +16,13 @@ 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
* RUN: %t
* HIT_END
*/
#include"test_common.h"
#include"hip_runtime.h"
#include"hip_runtime_api.h"
hipamd/tests/src/deviceLib/hip_anyall.cpp
+5
View File
@@ -20,6 +20,11 @@ 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
* RUN: %t
* HIT_END
*/
#include <stdio.h>
#include <iostream>
hipamd/tests/src/deviceLib/hip_ballot.cpp
+6
View File
@@ -17,6 +17,12 @@ 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
* RUN: %t
* HIT_END
*/
#include <iostream>
#include <hip_runtime.h>
hipamd/tests/src/deviceLib/hip_brev.cpp
+6
View File
@@ -20,6 +20,12 @@ 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
* RUN: %t
* HIT_END
*/
#include <assert.h>
#include <stdio.h>
#include <algorithm>
hipamd/tests/src/deviceLib/hip_clz.cpp
+6
View File
@@ -20,6 +20,12 @@ 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
* RUN: %t
* HIT_END
*/
#include <assert.h>
#include <stdio.h>
#include <algorithm>
hipamd/tests/src/deviceLib/hip_ffs.cpp
+6
View File
@@ -20,6 +20,12 @@ 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
* RUN: %t
* HIT_END
*/
#include <assert.h>
#include <stdio.h>
#include <algorithm>
hipamd/tests/src/deviceLib/hip_popc.cpp
+6
View File
@@ -20,6 +20,12 @@ 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
* RUN: %t
* HIT_END
*/
#include <assert.h>
#include <stdio.h>
#include <algorithm>
hipamd/tests/src/deviceLib/hip_test_ldg.cpp
+7
View File
@@ -19,6 +19,13 @@ 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 --gpu-architecture=sm_35
* RUN: %t
* HIT_END
*/
#include <assert.h>
#include <stdio.h>
#include <algorithm>
hipamd/tests/src/hipArray.cpp
+6
View File
@@ -1,3 +1,9 @@
/* HIT_START
* BUILD: %t %s test_common.cpp EXCLUDE_HIP_PLATFORM nvcc
* RUN: %t EXCLUDE_HIP_PLATFORM
* HIT_END
*/
#include "hip_runtime.h"
#include "test_common.h"
hipamd/tests/src/hipDynamicShared.cpp
+6
View File
@@ -20,6 +20,12 @@ 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
* RUN: %t EXCLUDE_HIP_PLATFORM nvcc
* HIT_END
*/
#include <hip_runtime.h>
#include "test_common.h"
hipamd/tests/src/hipEnvVar.cpp
+4
View File
@@ -20,6 +20,10 @@ 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
* HIT_END
*/
#include <stdio.h>
#include <stdlib.h>
hipamd/tests/src/hipEnvVarDriver.cpp
+25 -11
View File
@@ -15,19 +15,21 @@ EVENT SHALL THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR
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 <iostream>
#include <vector>
#include <stdio.h>
#include <stdlib.h>
#include <assert.h>
#include <string>
#include <hip_runtime.h>
using namespace std;
//./hipEnvVar -c -d 0 -h
//putenv("SomeVariable=SomeValue");
//putenv("export HIP_VISIBLE_DEVICES=0,1,2,3");
int getDeviceNumber(){
FILE *in;
char buff[512];
@@ -40,7 +42,8 @@ int getDeviceNumber(){
return atoi(buff);
}
int getDevicePCIBusNum(int deviceID){
// Query the current device ID remotely to hipEnvVar
int getDevicePCIBusNumRemote(int deviceID){
FILE *in;
char buff[512];
string str = "./hipEnvVar -d ";
@@ -53,6 +56,19 @@ int getDevicePCIBusNum(int deviceID){
return atoi(buff);
}
// Query the current device ID locally
int getDevicePCIBusNum(int deviceID){
hipSetDevice(deviceID);
hipDeviceProp_t devProp;
hipGetDeviceProperties(&devProp, deviceID);
if (devProp.major < 1) {
printf("%d does not support HIP\n", deviceID);
return -1;
}
return devProp.pciBusID;
}
int main() {
unsetenv("HIP_VISIBLE_DEVICES");
unsetenv("CUDA_VISIBLE_DEVICES");
@@ -72,11 +88,9 @@ int main() {
for (int i = 0; i < totalDeviceNum ; i++) {
setenv("HIP_VISIBLE_DEVICES",(char*)std::to_string(i).c_str(),1);
setenv("CUDA_VISIBLE_DEVICES",(char*)std::to_string(i).c_str(),1);
//cout<<"HIP_VISIBLE_DEVICES is "<<i<<" data in vector is "<<devPCINum[i]<<endl;
//std::cout <<"Returned pci number is"<< getDevicePCIBusNum(0) << std::endl;
if (devPCINum[i] != getDevicePCIBusNum(0)) {
std::cout << "The returned PciBusID is not correct"
<< std::endl;
if (devPCINum[i] != getDevicePCIBusNumRemote(0)) {
std::cout << "The returned PciBusID is not correct"<< std::endl;
std::cout << "Expected "<< devPCINum[i] << ", but get " << getDevicePCIBusNum << endl;
exit(-1);
} else {
continue;
hipamd/tests/src/hipEventRecord.cpp
+5
View File
@@ -23,6 +23,11 @@ THE SOFTWARE.
// Through manual inspection of the reported timestamps, can determine if recording a NULL event forces synchronization :
// set
/* HIT_START
* BUILD: %t %s test_common.cpp
* RUN: %t --iterations 10
* HIT_END
*/
#include "hip_runtime.h"
#include "test_common.h"
hipamd/tests/src/hipFuncDeviceSynchronize.cpp
+6
View File
@@ -22,6 +22,12 @@ THE SOFTWARE.
* hipError_t hipDeviceSynchronize();
*/
/* HIT_START
* BUILD: %t %s test_common.cpp
* RUN: %t
* HIT_END
*/
#include"test_common.h"
#define _SIZE sizeof(int)*1024*1024
hipamd/tests/src/hipFuncGetDevice.cpp
+6
View File
@@ -22,6 +22,12 @@ THE SOFTWARE.
* hipError_t hipGetDevice(int *device);
*/
/* HIT_START
* BUILD: %t %s test_common.cpp
* RUN: %t
* HIT_END
*/
#include "test_common.h"
int main()
hipamd/tests/src/hipFuncSetDevice.cpp
+6
View File
@@ -17,6 +17,12 @@ OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
/* HIT_START
* BUILD: %t %s test_common.cpp
* RUN: %t EXCLUDE_HIP_PLATFORM
* HIT_END
*/
#include "test_common.h"
int main(){
hipamd/tests/src/hipFuncSetDeviceFlags.cpp
+6
View File
@@ -17,6 +17,12 @@ OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
/* HIT_START
* BUILD: %t %s test_common.cpp
* RUN: %t
* HIT_END
*/
#include "test_common.h"
int main()
hipamd/tests/src/hipGetDeviceAttribute.cpp
+9 -3
View File
@@ -21,6 +21,12 @@ THE SOFTWARE.
*/
// Test the device info API extensions for HIP:
/* HIT_START
* BUILD: %t %s test_common.cpp
* RUN: %t EXCLUDE_HIP_PLATFORM nvcc
* HIT_END
*/
#include <stdio.h>
#include <iostream>
#include <hip_runtime.h>
@@ -70,15 +76,15 @@ int main(int argc, char *argv[])
CHECK(test_hipDeviceGetAttribute(deviceId, hipDeviceAttributeMemoryClockRate, props.memoryClockRate));
CHECK(test_hipDeviceGetAttribute(deviceId, hipDeviceAttributeMemoryBusWidth, props.memoryBusWidth));
CHECK(test_hipDeviceGetAttribute(deviceId, hipDeviceAttributeMultiprocessorCount, props.multiProcessorCount));
CHECK(test_hipDeviceGetAttribute(deviceId, hipDeviceAttributeIsMultiGpuBoard, props.isMultiGpuBoard));
CHECK(test_hipDeviceGetAttribute(deviceId, hipDeviceAttributeIsMultiGpuBoard, props.isMultiGpuBoard));//
CHECK(test_hipDeviceGetAttribute(deviceId, hipDeviceAttributeComputeMode, props.computeMode));
CHECK(test_hipDeviceGetAttribute(deviceId, hipDeviceAttributeL2CacheSize, props.l2CacheSize));
CHECK(test_hipDeviceGetAttribute(deviceId, hipDeviceAttributeMaxThreadsPerMultiProcessor, props.maxThreadsPerMultiProcessor));
CHECK(test_hipDeviceGetAttribute(deviceId, hipDeviceAttributeComputeCapabilityMajor, props.major));
CHECK(test_hipDeviceGetAttribute(deviceId, hipDeviceAttributeComputeCapabilityMinor, props.minor));
CHECK(test_hipDeviceGetAttribute(deviceId, hipDeviceAttributeComputeCapabilityMinor, props.minor));//
CHECK(test_hipDeviceGetAttribute(deviceId, hipDeviceAttributeConcurrentKernels, props.concurrentKernels));
CHECK(test_hipDeviceGetAttribute(deviceId, hipDeviceAttributePciBusId, props.pciBusID));
CHECK(test_hipDeviceGetAttribute(deviceId, hipDeviceAttributePciDeviceId, props.pciDeviceID));
CHECK(test_hipDeviceGetAttribute(deviceId, hipDeviceAttributePciDeviceId, props.pciDeviceID));//
CHECK(test_hipDeviceGetAttribute(deviceId, hipDeviceAttributeMaxSharedMemoryPerMultiprocessor, props.maxSharedMemoryPerMultiProcessor));
passed();
hipamd/tests/src/hipHcc.cpp
+6
View File
@@ -21,6 +21,12 @@ THE SOFTWARE.
*/
// Test the HCC-specific API extensions for HIP:
/* HIT_START
* BUILD: %t %s HCC_OPTIONS --stdlib=libc++
* RUN: %t
* HIT_END
*/
#include <stdio.h>
#include <iostream>
#include <hip_runtime.h>
hipamd/tests/src/hipHostGetFlags.cpp
+6
View File
@@ -20,6 +20,12 @@ 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
* RUN: %t
* HIT_END
*/
#include"test_common.h"
#include<malloc.h>
hipamd/tests/src/hipHostAlloc.cpp → hipamd/tests/src/hipHostMalloc.cpp
+6
View File
@@ -20,6 +20,12 @@
THE SOFTWARE.
*/
/* HIT_START
* BUILD: %t %s test_common.cpp
* RUN: %t
* HIT_END
*/
#include"test_common.h"
#define LEN 1024*1024
hipamd/tests/src/hipHostRegister.cpp
+6
View File
@@ -17,6 +17,12 @@ 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
* RUN: %t
* HIT_END
*/
#include"test_common.h"
#include<malloc.h>
hipamd/tests/src/hipLaunchParm.cpp
+6
View File
@@ -17,6 +17,12 @@ 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
* RUN: %t
* HIT_END
*/
#include"hip_runtime.h"
#include"test_common.h"
#include"hip_runtime_api.h"
hipamd/tests/src/hipPeerToPeer_simple.cpp
+25 -1
View File
@@ -22,6 +22,14 @@ THE SOFTWARE.
// Simple test for memset.
// Also serves as a template for other tests.
/* HIT_START
* BUILD: %t %s test_common.cpp
* RUN: %t EXCLUDE_HIP_PLATFORM all
* RUN: %t --memcpyWithPeer EXCLUDE_HIP_PLATFORM all
* RUN: %t --mirrorPeers EXCLUDE_HIP_PLATFORM all
* HIT_END
*/
#include "hip_runtime.h"
#include "test_common.h"
@@ -33,6 +41,15 @@ int p_peerDevice = -1; // explicly specify which peer to use, else use p_gpuDe
int g_currentDevice;
int g_peerDevice;
void help(char *argv[])
{
printf ("usage: %s [OPTIONS]\n", argv[0]);
printf (" --memcpyWithPeer : Perform memcpy with peer.\n");
printf (" --mirrorPeersi : Mirror memory onto both default device and peerdevice. If 0, memory is mapped only on the default device.\n");
printf (" --peerDevice N : Set peer device.\n");
};
void parseMyArguments(int argc, char *argv[])
{
int more_argc = HipTest::parseStandardArguments(argc, argv, false);
@@ -40,7 +57,10 @@ void parseMyArguments(int argc, char *argv[])
for (int i = 1; i < more_argc; i++) {
const char *arg = argv[i];
if (!strcmp(arg, "--memcpyWithPeer")) {
if (!strcmp(arg, "--help")) {
help(argv);
exit(-1);
} else if (!strcmp(arg, "--memcpyWithPeer")) {
p_memcpyWithPeer = true;
} else if (!strcmp(arg, "--mirrorPeers")) {
p_mirrorPeers = true;
@@ -90,10 +110,12 @@ void enablePeerFirst()
setupPeerTests();
// Always enable g_currentDevice to see the allocations on peerDevice.
HIPCHECK(hipSetDevice(g_currentDevice));
HIPCHECK(hipDeviceEnablePeerAccess(g_peerDevice, 0));
if (p_mirrorPeers) {
// Mirror peers allows the peer device to see the allocations on currentDevice.
int canAccessPeer;
HIPCHECK(hipDeviceCanAccessPeer(&canAccessPeer, g_peerDevice, g_currentDevice));
assert(canAccessPeer);
@@ -122,6 +144,8 @@ void enablePeerFirst()
// Device0 push to device1, using P2P:
// NOTE : if p_mirrorPeers=0 and p_memcpyWithPeer=1, then peer device does not have mapping for A_d1 and we need to use a
// a host staging copy for the P2P access.
HIPCHECK (hipSetDevice(p_memcpyWithPeer ? g_peerDevice : g_currentDevice));
HIPCHECK (hipMemcpy(A_d1, A_d0, Nbytes, hipMemcpyDefault)); // This is P2P copy.
hipamd/tests/src/hipRandomMemcpyAsync.cpp
+6
View File
@@ -17,6 +17,12 @@ 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
* RUN: %t
* HIT_END
*/
#include <assert.h>
#include <stdio.h>
#include <stdlib.h>
hipamd/tests/src/hipTestMemcpyPin.cpp
+6
View File
@@ -17,6 +17,12 @@ 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
* RUN: %t
* HIT_END
*/
#include<hip_runtime.h>
#include"test_common.h"
hipamd/tests/src/kernel/hipGridLaunch.cpp
+6
View File
@@ -21,6 +21,12 @@ THE SOFTWARE.
*/
// Test the Grid_Launch syntax.
/* HIT_START
* BUILD: %t %s ../test_common.cpp
* RUN: %t
* HIT_END
*/
#include "hip_runtime.h"
#include "test_common.h"
hipamd/tests/src/kernel/hipLanguageExtensions.cpp
+6
View File
@@ -21,6 +21,12 @@ THE SOFTWARE.
*/
// Collection of code to make sure that various features in the hip kernel language compile.
/* HIT_START
* BUILD: %t %s ../test_common.cpp HCC_OPTIONS --stdlib=libc++
* RUN: %t
* HIT_END
*/
#include <hip_runtime.h>
#include <test_common.h>
hipamd/tests/src/kernel/launch_bounds.cpp
+5
View File
@@ -20,6 +20,11 @@ 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
* RUN: %t
* HIT_END
*/
// Test launch bounds and initialization conditions.
hipamd/tests/src/runtimeApi/memory/hipMemcpy.cpp
+10
View File
@@ -19,6 +19,16 @@ 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
* RUN_NAMED: %t hipMemcpy-modes --tests 0x1
* RUN_NAMED: %t hipMemcpy-size --tests 0x6
* RUN_NAMED: %t hipMemcpy-multithreaded --tests 0x8
* HIT_END
*/
#include "hip_runtime.h"
#include "hip_runtime.h"
#include "test_common.h"
hipamd/tests/src/runtimeApi/memory/hipMemcpyAll.cpp
+4
View File
@@ -17,6 +17,10 @@ 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
* HIT_END
*/
#include<hip_runtime.h>
#include<iostream>
hipamd/tests/src/runtimeApi/memory/hipMemcpy_simple.cpp
+8
View File
@@ -19,6 +19,14 @@ LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
/* HIT_START
* BUILD: %t %s ../../test_common.cpp
* RUN: %t
* RUN_NAMED: %t hipMemcpyAsync-simple --async
* HIT_END
*/
#include "hip_runtime.h"
#include "test_common.h"
hipamd/tests/src/runtimeApi/memory/hipMemoryAllocate.cpp
+7
View File
@@ -16,6 +16,13 @@ LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
/* HIT_START
* BUILD: %t %s ../../test_common.cpp
* RUN: %t
* HIT_END
*/
#include"test_common.h"
#define SIZE 1024*1024*256
hipamd/tests/src/runtimeApi/memory/hipMemset.cpp
+12
View File
@@ -22,6 +22,18 @@ THE SOFTWARE.
// Simple test for memset.
// Also serves as a template for other tests.
/* HIT_START
* BUILD: %t %s ../../test_common.cpp
* RUN: %t
* //Small copy
* RUN: %t -N 10 --memsetval 0x42
* // Oddball size
* RUN: %t -N 10013 --memsetval 0x5a
* // Big copy
* RUN: %t -N 256M --memsetval 0xa6
* HIT_END
*/
#include "hip_runtime.h"
#include "test_common.h"
hipamd/tests/src/runtimeApi/multiThread/hipMultiThreadDevice.cpp
+8
View File
@@ -1,3 +1,11 @@
/* HIT_START
* BUILD: %t %s ../../test_common.cpp
* RUN_NAMED: %t hipMultiThreadDevice-serial --tests 0x1
* RUN_NAMED: %t hipMultiThreadDevice-pyramid --tests 0x4
* RUN_NAMED: %t hipMultiThreadDevice-nearzero --tests 0x10
* HIT_END
*/
#include <hip_runtime_api.h>
#include "test_common.h"
hipamd/tests/src/runtimeApi/multiThread/hipMultiThreadStreams1.cpp
+7 -2
View File
@@ -19,11 +19,16 @@ 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
* RUN: %t
* HIT_END
*/
#include "hip_runtime.h"
#include "test_common.h"
void printSep()
{
printf ("======================================================================================\n");
hipamd/tests/src/runtimeApi/multiThread/hipMultiThreadStreams2.cpp
+15 -1
View File
@@ -20,11 +20,19 @@ 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
* RUN: %t
* HIT_END
*/
#include<iostream>
#include"test_common.h"
#include<thread>
#define N 1000
template<typename T>
__global__ void Inc(hipLaunchParm lp, T *Array){
int tx = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x;
@@ -84,7 +92,10 @@ void run(size_t size, hipStream_t stream1, hipStream_t stream2){
HIPASSERT(Ehh[10] = Ahh[10] + 1.0f);
}
int main(int argc, char **argv){
int main(int argc, char **argv)
{
iterations = 100;
HipTest::parseStandardArguments(argc, argv, true);
@@ -94,6 +105,8 @@ int main(int argc, char **argv){
}
const size_t size = N * sizeof(float);
for (int i=0; i< iterations; i++) {
std::thread t1(run1, size, stream[0]);
std::thread t2(run1, size, stream[0]);
@@ -103,6 +116,7 @@ int main(int argc, char **argv){
t2.join();
// std::cout<<"T2"<<std::endl;
t3.join();
}
passed();
}
hipamd/tests/src/runtimeApi/stream/hipAPIStreamDisable.cpp
+5
View File
@@ -17,6 +17,11 @@ OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
/* HIT_START
* BUILD: %t %s ../../test_common.cpp
* HIT_END
*/
#include<iostream>
#include"test_common.h"
hipamd/tests/src/runtimeApi/stream/hipAPIStreamEnable.cpp
+5
View File
@@ -17,6 +17,11 @@ OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
/* HIT_START
* BUILD: %t %s ../../test_common.cpp
* HIT_END
*/
//#define HIP_API_PER_THREAD_DEFAULT_STREAM
#include<iostream>
hipamd/tests/src/runtimeApi/stream/hipStreamL5.cpp
+6
View File
@@ -20,6 +20,12 @@ 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
* RUN: %t
* HIT_END
*/
#include "test_common.h"
#include "hipStream.h"
hipamd/tests/src/runtimeApi/stream/hipStreamWaitEvent.cpp
+6
View File
@@ -17,6 +17,12 @@ 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
* RUN: %t
* HIT_END
*/
// Test under-development. Calls async mem-copy API, experiment with functionality.
#include "hip_runtime.h"