AI/rocm-systems
2
0
Rozštěpit 0
Zdrojový kód Úkoly Pull requesty Akce Balíčky Projekty Vydání Wiki Aktivita

Merge branch 'amd-develop' into amd-master

Procházet zdrojové kódy 
Change-Id: I0e856db61fa4a50e190bd1d4c464ceb4a709b550


[ROCm/hip commit: 6e76723e6e]
Tento commit je obsažen v:
Maneesh Gupta
2017-02-23 11:19:06 +05:30
rodič 5546eda08c e5682bd659
revize a12f087d2c
204 změnil soubory, kde provedl 21109 přidání a 4660 odebrání
Stáhněte soubor záplaty Stáhněte rozdílový soubor Rozbalit všechny soubory Sbalit všechny soubory
projects/hip/CMakeLists.txt
+52 -8
Zobrazit soubor
@@ -33,7 +33,7 @@ execute_process(COMMAND git show -s --format=@%ct
OUTPUT_VARIABLE HIP_VERSION_PATCH
OUTPUT_STRIP_TRAILING_WHITESPACE)
set(HIP_VERSION $HIP_VERSION_MAJOR.$HIP_VERSION_MINOR.$HIP_VERSION_PATCH)
set(HIP_VERSION ${HIP_VERSION_MAJOR}.${HIP_VERSION_MINOR}.${HIP_VERSION_PATCH})
add_to_config(_versionInfo HIP_VERSION_MAJOR)
add_to_config(_versionInfo HIP_VERSION_MINOR)
add_to_config(_versionInfo HIP_VERSION_PATCH)
@@ -158,7 +158,8 @@ if(HIP_PLATFORM STREQUAL "hcc")
set(HIP_HCC_BUILD_FLAGS "${HIP_HCC_BUILD_FLAGS} -DHIP_VERSION_MAJOR=${HIP_VERSION_MAJOR} -DHIP_VERSION_MINOR=${HIP_VERSION_MINOR} -DHIP_VERSION_PATCH=${HIP_VERSION_PATCH}")
# Add remaining flags
set(HIP_HCC_BUILD_FLAGS "${HIP_HCC_BUILD_FLAGS} -fPIC -hc -I${HCC_HOME}/include -I${HSA_PATH}/include -I/opt/rocm/libhsakmt/include -I/usr/local/include/c++/v1 -stdlib=libc++")
execute_process(COMMAND ${HCC_HOME}/bin/hcc-config --cxxflags OUTPUT_VARIABLE HCC_CXX_FLAGS)
set(HIP_HCC_BUILD_FLAGS "${HIP_HCC_BUILD_FLAGS} -fPIC ${HCC_CXX_FLAGS} -I${HSA_PATH}/include -I/opt/rocm/libhsakmt/include")
# Set compiler and compiler flags
set(CMAKE_CXX_COMPILER "${HCC_HOME}/bin/hcc")
@@ -175,23 +176,33 @@ if(HIP_PLATFORM STREQUAL "hcc")
src/hip_memory.cpp
src/hip_peer.cpp
src/hip_stream.cpp
src/hip_module.cpp)
src/hip_module.cpp
src/env.cpp)
set(SOURCE_FILES_DEVICE
src/device_util.cpp
src/hip_ldg.cpp
src/hip_fp16.cpp
src/device_functions.cpp)
src/device_functions.cpp
src/math_functions.cpp)
set(CMAKE_SHARED_LINKER_FLAGS "${CMAKE_SHARED_LINKER_FLAGS} -L${HCC_HOME}/lib -lmcwamp -Wl,-Bsymbolic -Wl,-rpath ${HCC_HOME}/lib")
execute_process(COMMAND ${HCC_HOME}/bin/hcc-config --ldflags OUTPUT_VARIABLE HCC_LD_FLAGS)
set(CMAKE_SHARED_LINKER_FLAGS "${CMAKE_SHARED_LINKER_FLAGS} ${HCC_LD_FLAGS} -Wl,-Bsymbolic")
set(CMAKE_SHARED_LINKER_FLAGS "${CMAKE_SHARED_LINKER_FLAGS} --amdgpu-target=gfx701 --amdgpu-target=gfx801 --amdgpu-target=gfx802 --amdgpu-target=gfx803")
add_library(hip_hcc SHARED ${SOURCE_FILES_RUNTIME})
target_link_libraries(hip_hcc c++ c++abi hc_am)
target_link_libraries(hip_hcc PRIVATE hc_am)
add_library(hip_hcc_static STATIC ${SOURCE_FILES_RUNTIME})
target_link_libraries(hip_hcc_static c++ c++abi hc_am)
target_link_libraries(hip_hcc_static PRIVATE hc_am)
add_dependencies(hip_hcc_static hip_hcc)
add_library(hip_device STATIC ${SOURCE_FILES_DEVICE})
add_dependencies(hip_device hip_hcc)
string(REPLACE " " ";" HCC_CXX_FLAGS_LIST ${HCC_CXX_FLAGS})
foreach(TARGET hip_hcc hip_hcc_static hip_device)
target_include_directories(${TARGET} SYSTEM INTERFACE $<INSTALL_INTERFACE:$<INSTALL_PREFIX>/include>;${HSA_PATH}/include)
endforeach()
target_link_libraries(hip_hcc INTERFACE hcc::hccrt;hcc::hc_am)
# 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)
@@ -216,7 +227,7 @@ add_custom_target(doc COMMAND HIP_PATH=${CMAKE_CURRENT_SOURCE_DIR} doxygen ${CMA
# Install hip_hcc if platform is hcc
if(HIP_PLATFORM STREQUAL "hcc")
install(TARGETS hip_hcc_static hip_hcc hip_device DESTINATION lib)
install(FILES ${CMAKE_CURRENT_SOURCE_DIR}/src/hip_ir.ll DESTINATION lib)
install(FILES ${CMAKE_CURRENT_SOURCE_DIR}/src/hip_hc.ll ${CMAKE_CURRENT_SOURCE_DIR}/src/hip_hc_gfx803.ll DESTINATION lib)
# Install .hipInfo
install(FILES ${PROJECT_BINARY_DIR}/.hipInfo DESTINATION lib)
@@ -235,6 +246,39 @@ if(NOT ${INSTALL_SOURCE} EQUAL 0)
install(DIRECTORY cmake DESTINATION .)
endif()
#############################
# hip-config
#############################
set(LIB_INSTALL_DIR ${CMAKE_INSTALL_PREFIX}/lib)
set(INCLUDE_INSTALL_DIR ${CMAKE_INSTALL_PREFIX}/include)
set(BIN_INSTALL_DIR ${CMAKE_INSTALL_PREFIX}/bin)
set(CONFIG_PACKAGE_INSTALL_DIR ${LIB_INSTALL_DIR}/cmake/hip)
if(HIP_PLATFORM STREQUAL "hcc")
install(TARGETS hip_hcc_static hip_hcc hip_device EXPORT hip-targets DESTINATION ${LIB_INSTALL_DIR})
install(EXPORT hip-targets DESTINATION ${CONFIG_PACKAGE_INSTALL_DIR} NAMESPACE hip::)
include(CMakePackageConfigHelpers)
configure_package_config_file(
hip-config.cmake.in
${CMAKE_CURRENT_BINARY_DIR}/hip-config.cmake
INSTALL_DESTINATION ${CONFIG_PACKAGE_INSTALL_DIR}
PATH_VARS LIB_INSTALL_DIR INCLUDE_INSTALL_DIR BIN_INSTALL_DIR
)
write_basic_package_version_file(
${CMAKE_CURRENT_BINARY_DIR}/hip-config-version.cmake
VERSION "${HIP_VERSION}"
COMPATIBILITY SameMajorVersion
)
install(
FILES
${CMAKE_CURRENT_BINARY_DIR}/hip-config.cmake
${CMAKE_CURRENT_BINARY_DIR}/hip-config-version.cmake
DESTINATION
${CONFIG_PACKAGE_INSTALL_DIR}
)
endif()
#############################
# Packaging steps
#############################
projects/hip/CONTRIBUTING.md
+1 -1
Zobrazit soubor
@@ -126,7 +126,7 @@ Differences or limitations of HIP APIs as compared to CUDA APIs should be clearl
- All HIP environment variables should begin with the keyword HIP_
Environment variables should be long enough to describe their purpose but short enough so they can be remembered - perhaps 10-20 characters, with 3-4 parts separated by underscores.
To see the list of current environment variables, along with their values, set HIP_PRINT_ENV and run any hip applications on ROCM platform .
To see the list of current environment variables, along with their values, set HIP_PRINT_ENV and run any hip applications on ROCm platform .
HIPCC or other tools may support additional environment variables which should follow the above convention.
projects/hip/INSTALL.md
+3 -3
Zobrazit soubor
@@ -22,14 +22,14 @@ HIP code can be developed either on AMD ROCm platform using hcc compiler, or a C
## AMD-hcc
* Install the [rocm](http://gpuopen.com/getting-started-with-boltzmann-components-platforms-installation/) packages. Rocm will install all of the necessary components, including the kernel driver, runtime software, HCC compiler, and HIP.
* Install the [rocm](http://gpuopen.com/getting-started-with-boltzmann-components-platforms-installation/) packages. ROCm will install all of the necessary components, including the kernel driver, runtime software, HCC compiler, and HIP.
* Default paths and environment variables:
* By default HIP looks for hcc in /opt/rocm/hcc (can be overridden by setting HCC_HOME environment variable)
* By default HIP looks for HSA in /opt/rocm/hsa (can be overridden by setting HSA_PATH environment variable)
* By default HIP is installed into /opt/rocm/hip (can be overridden by setting HIP_PATH environment variable).
* Optionally, consider adding /opt/rocm/bin to your path to make it easier to use the tools.
* Optionally, consider adding /opt/rocm/bin to your PATH to make it easier to use the tools.
## NVIDIA-nvcc
@@ -96,7 +96,7 @@ The native GCN target is included with upstream LLVM, and has also been integrat
Binary packages for the direct-to-isa package are included with the [rocm](http://gpuopen.com/getting-started-with-boltzmann-components-platforms-installation/) package.
Alternatively, this sections describes how to build it from source:
1. Install the rocm packages as described above.
1. Install the ROCm packages as described above.
2. Follow the instructions [here](https://github.com/RadeonOpenCompute/HCC-Native-GCN-ISA/wiki)
* In the make step for HCC, we recommend setting -DCMAKE_INSTALL_PREFIX.
* Set HCC_HOME environment variable before compiling HIP program to point to the native compiler:
projects/hip/bin/hccgenco.sh
+4 -6
Zobrazit soubor
@@ -46,17 +46,15 @@ done
printf "\nint main(){}\n" >> $hipgenisa_main
$HIP_PATH/bin/hipcc $hipgenisa_files -o $hipgenisa_dir/a.out
mv dump.* $hipgenisa_dir
$ROCM_PATH/hcc-lc/bin/llvm-mc -arch=amdgcn -mcpu=$ROCM_TARGET -filetype=obj $hipgenisa_dir/dump.isa -o $hipgenisa_dir/dump.o
$ROCM_PATH/llvm/bin/clang -target amdgcn--amdhsa $hipgenisa_dir/dump.o -o $hipgenisa_dir/dump.co
mv dump* $hipgenisa_dir
map_sym=""
kernels=$(objdump -t $hipgenisa_dir/dump.co | grep grid_launch_parm | sed 's/ \+/ /g; s/\t/ /g' | cut -d" " -f6)
kernels=$(objdump -t $hipgenisa_dir/dump-fiji.hsaco | grep grid_launch_parm | sed 's/ \+/ /g; s/\t/ /g' | cut -d" " -f6)
for mangled_sym in $kernels; do
real_sym=$(c++filt $(c++filt _$mangled_sym | cut -d: -f3 | sed 's/_functor//g') | cut -d\( -f1)
real_sym=$(c++filt $(c++filt $mangled_sym | cut -d: -f3 | sed 's/_functor//g') | cut -d\( -f1)
map_sym="--redefine-sym $mangled_sym=$real_sym $map_sym"
done
objcopy -F elf64-little $map_sym $hipgenisa_dir/dump.co $OUTPUT_FILE
objcopy -F elf64-little $map_sym $hipgenisa_dir/dump-fiji.hsaco $OUTPUT_FILE
rm $hipgenisa_files
rm -r $hipgenisa_dir
projects/hip/bin/hipcc
+67 -41
Zobrazit soubor
@@ -70,6 +70,11 @@ if ($verbose & 0x2) {
# set if user explicitly requests -stdlib=libc++. (else we default to libstdc++ for better interop with g++):
$setStdLib = 0; # TODO - set to 0
$target_gfx701 = 0;
$target_gfx801 = 0;
$target_gfx802 = 0;
$target_gfx803 = 0;
if ($HIP_PLATFORM eq "hcc") {
$HSA_PATH=$ENV{'HSA_PATH'} // "/opt/rocm/hsa";
@@ -92,6 +97,8 @@ if ($HIP_PLATFORM eq "hcc") {
$HIPCC=$HCC;
$HIPCXXFLAGS = $HCCFLAGS;
$HIPLDFLAGS = `${HCC_HOME}/bin/hcc-config --ldflags`;
#### GCC system includes workaround ####
$HCC_WA_FLAGS = " ";
if ($HCC_VERSION_MAJOR eq 1) {
@@ -99,50 +106,25 @@ if ($HIP_PLATFORM eq "hcc") {
my $GPP_CUR_VER = `g++ -dumpversion`;
$GCC_CUR_VER =~ s/\R//g;
$GPP_CUR_VER =~ s/\R//g;
if (${GCC_CUR_VER} eq ${GPP_CUR_VER}) {
$HCC_WA_FLAGS .= " -I/usr/include/x86_64-linux-gnu -I/usr/include/x86_64-linux-gnu/c++/${GCC_CUR_VER} -I/usr/include/c++/${GCC_CUR_VER} ";
my @GPP_VER_FIELDS = split('\.', $GPP_CUR_VER);
# Only include the libstdc++ headers and libraries flags explicitly if the g++ is older than version 5.
# That's because HCC already uses libstdc++ by default if a newer g++/libstdc++ is available
if (${GCC_CUR_VER} eq ${GPP_CUR_VER} and $GPP_VER_FIELDS[0] < 5) {
$HCC_WA_FLAGS .= " -stdlib=libstdc++ -I/usr/include/x86_64-linux-gnu -I/usr/include/x86_64-linux-gnu/c++/${GCC_CUR_VER} -I/usr/include/c++/${GCC_CUR_VER} ";
# Add C++ libs for GCC.
$HIPLDFLAGS .= " -lstdc++";
}
}
$HIPCXXFLAGS .= " -I$HIP_PATH/include/hip/hcc_detail/cuda";
$HIPCXXFLAGS .= " -I$HSA_PATH/include";
$HIPCXXFLAGS .= " -Wno-deprecated-register";
$HIPLDFLAGS = "-hc -L$HCC_HOME/lib -Wl,--rpath=$HCC_HOME/lib -lc++ -ldl -lpthread -Wl,--whole-archive -lmcwamp -Wl,--no-whole-archive";
# Suppress linker warnings in case HCC distribution contains OpenCL/SPIR symbols
$HOST_OSNAME= `cat /etc/os-release | grep "^ID\=" | cut -d= -f2 | tr -d '\n'`;
$HOST_OSVER= `cat /etc/os-release | grep "^VERSION_ID\=" | cut -d= -f2 | tr -d '\n'`;
if ($HOST_OSNAME eq "ubuntu" and $HOST_OSVER eq "\"16.04\"") {
# No additional flags required
} else {
$HIPLDFLAGS .= " -Wl,--defsym=_binary_kernel_spir_end=1 -Wl,--defsym=_binary_kernel_spir_start=1 -Wl,--defsym=_binary_kernel_cl_start=1 -Wl,--defsym=_binary_kernel_cl_end=1";
}
if ($HOST_OSNAME eq "fedora") {
$HIPCXXFLAGS .= " -I/usr/local/include/c++/v1";
}
# Satisfy HCC dependencies
if ($HOST_OSNAME eq "fedora") {
$HIPLDFLAGS .= " -lc++abi";
} else {
$HIPLDFLAGS .= " -lc++abi -lsupc++";
}
$HIPLDFLAGS .= " -lsupc++";
$HIPLDFLAGS .= " -L$HSA_PATH/lib -L$ROCM_PATH/lib -lhsa-runtime64 -lhc_am -lhsakmt";
# Handle ROCm target platform
if ($ROCM_TARGET eq "fiji") {
$HIPLDFLAGS .= " -amdgpu-target=AMD:AMDGPU:8:0:3";
}
if ($ROCM_TARGET eq "carrizo") {
$HIPLDFLAGS .= " -amdgpu-target=AMD:AMDGPU:8:0:1";
}
if ($ROCM_TARGET eq "hawaii") {
$HIPLDFLAGS .= " -amdgpu-target=AMD:AMDGPU:7:0:1";
}
if ($ROCM_TARGET eq "polaris") {
$HIPLDFLAGS .= " -amdgpu-target=AMD:AMDGPU:8:0:3";
}
# Add trace marker library:
# TODO - once we cleanly separate the HIP API headers from HIP library headers this logic should move to CMakebuild option - apps do not need to see the marker library.
if ($HIP_ATP_MARKER) {
@@ -157,8 +139,6 @@ if ($HIP_PLATFORM eq "hcc") {
$HIPLDFLAGS .= " -L$marker_lib_path -lCXLActivityLogger -Wl,--rpath=$marker_lib_path";
}
# Add C++ libs for GCC.
$HIPLDFLAGS .= " -lstdc++";
$HIPLDFLAGS .= " -lm";
if ($verbose & 0x2) {
@@ -222,7 +202,8 @@ if($HIP_PLATFORM eq "hcc"){
if(($HIP_PLATFORM eq "hcc")){
$EXPORT_LL=" ";
$ENV{HCC_EXTRA_LIBRARIES}="$HIP_PATH/lib/hip_ir.ll\n";
$ENV{HCC_EXTRA_LIBRARIES}="$HIP_PATH/lib/hip_hc.ll\n";
$ENV{HIP_HC_IR_FILE}="";
}
if($HIP_PLATFORM eq "nvcc"){
@@ -261,9 +242,24 @@ foreach $arg (@ARGV)
$HIPCXXFLAGS .= " -stdlib=libc++";
$setStdLib = 1;
}
if($arg eq '--amdgpu-target=gfx701')
{
$target_gfx701 = 1;
}
if($arg eq '--amdgpu-target=gfx801')
{
$target_gfx801 = 1;
}
if($arg eq '--amdgpu-target=gfx802')
{
$target_gfx802 = 1;
}
if($arg eq '--amdgpu-target=gfx803')
{
$target_gfx803 = 1;
}
if(($trimarg eq '-stdlib=libstdc++') and ($setStdLib eq 0))
{
$HIPCXXFLAGS .= " -stdlib=libstdc++";
$HIPCXXFLAGS .= $HCC_WA_FLAGS;
$setStdLib = 1;
}
@@ -322,6 +318,36 @@ foreach $arg (@ARGV)
$toolArgs .= " $arg" unless $swallowArg;
}
if($HIP_PLATFORM eq "hcc"){
$ENV{HCC_EXTRA_LIBRARIES}="$HIP_PATH/lib/hip_hc.ll\n";
# Handle ROCm target platform
if ($target_gfx701 eq 1) {
$HIPLDFLAGS .= " --amdgpu-target=gfx701";
$HIPCXXFLAGS .= " -D__HIP_ARCH_GFX701__=1 ";
}
if ($target_gfx801 eq 1) {
$HIPLDFLAGS .= " --amdgpu-target=gfx801";
$HIPCXXFLAGS .= " -D__HIP_ARCH_GFX801__=1 ";
}
if ($target_gfx802 eq 1) {
$HIPLDFLAGS .= " --amdgpu-target=gfx802";
$HIPCXXFLAGS .= " -D__HIP_ARCH_GFX802__=1 ";
}
if ($target_gfx803 eq 1) {
$HIPLDFLAGS .= " --amdgpu-target=gfx803";
$HIPCXXFLAGS .= " -D__HIP_ARCH_GFX803__=1 ";
$ENV{HCC_EXTRA_LIBRARIES_GFX803}="$HIP_PATH/lib/hip_hc_gfx803.ll\n";
}
if ($target_gfx701 eq 0 and $target_gfx801 eq 0 and $target_gfx802 eq 0 and $target_gfx803 eq 0)
{
$HIPLDFLAGS .= " --amdgpu-target=gfx701 --amdgpu-target=gfx801 --amdgpu-target=gfx802 --amdgpu-target=gfx803";
$ENV{HCC_EXTRA_LIBRARIES_GFX803}="$HIP_PATH/lib/hip_hc_gfx803.ll\n";
}
}
if ($hasC and $HIP_PLATFORM eq 'nvcc') {
$HIPCXXFLAGS .= " -x cu";
}
@@ -334,7 +360,6 @@ if ($buildDeps and $HIP_PLATFORM eq 'nvcc') {
if ($setStdLib eq 0 and $HIP_PLATFORM eq 'hcc')
{
$HIPCXXFLAGS .= " -stdlib=libstdc++";
$HIPCXXFLAGS .= $HCC_WA_FLAGS;
}
@@ -374,7 +399,8 @@ if ($printHipVersion) {
}
if ($runCmd) {
if ($HIP_PLATFORM eq "hcc" and exists($hipConfig{'HCC_VERSION'}) and $HCC_VERSION ne $hipConfig{'HCC_VERSION'}) {
print ("HIP ($HIP_PATH) was built using hcc $hipConfig{'HCC_VERSION'}, but you are using $HCC_HOME/hcc with version $HCC_VERSION from hipcc. Please rebuild HIP including cmake or update HCC_HOME variable.\n") && die ();
print ("HIP ($HIP_PATH) was built using hcc $hipConfig{'HCC_VERSION'}, but you are using $HCC_HOME/hcc with version $HCC_VERSION from hipcc. Please rebuild HIP including cmake or update HCC_HOME variable.\n") ;
die unless $ENV{'HIP_IGNORE_HCC_VERSION'};
}
system ("$CMD") and die ();
}
projects/hip/docs/markdown/CUDA_Runtime_API_functions_supported_by_HIP.md
+160 -54
Zobrazit soubor
@@ -1,34 +1,36 @@
**1. Device Management**
# CUDA Runtime API functions supported by HIP
## **1. Device Management**
| **CUDA** | **HIP** | **CUDA description** |
|-----------------------------------------------------------|-------------------------------|--------------------------------------------------------------------------------------------------------------------------------|
| `cudaChooseDevice` | `hipChooseDevice` | Select compute-device which best matches criteria. |
| `cudaChooseDevice` | `hipChooseDevice` | Select compute-device which best matches criteria. |
| `cudaDeviceGetAttribute` | `hipDeviceGetAttribute` | Returns information about the device. |
| `cudaDeviceGetByPCIBusId` | | Returns a handle to a compute device. |
| `cudaDeviceGetByPCIBusId` | `hipDeviceGetByPCIBusId` | Returns a handle to a compute device. |
| `cudaDeviceGetCacheConfig` | `hipDeviceGetCacheConfig` | Returns the preferred cache configuration for the current device. |
| `cudaDeviceGetLimit` | `hipDeviceGetLimit` | Returns resource limits. |
| `cudaDeviceGetPCIBusId` | | Returns a PCI Bus Id string for the device. |
| `cudaDeviceGetLimit` | `hipDeviceGetLimit` | Returns resource limits. |
| `cudaDeviceGetPCIBusId` | `hipDeviceGetPCIBusId` | Returns a PCI Bus Id string for the device. |
| `cudaDeviceGetSharedMemConfig` | `hipDeviceGetSharedMemConfig` | Returns the shared memory configuration for the current device. |
| `cudaDeviceGetStreamPriorityRange` | | Returns numerical values that correspond to the least and greatest stream priorities. |
| `cudaDeviceReset` | `hipDeviceReset` | Destroy all allocations and reset all state on the current device in the current process. |
| `cudaDeviceSetCacheConfig` | `hipDeviceSetCacheConfig` | Sets the preferred cache configuration for the current device. |
| `cudaDeviceSetLimit` | `hipDeviceSetLimit` | Set resource limits. |
| `cudaDeviceSetLimit` | `hipDeviceSetLimit` | Set resource limits. |
| `cudaDeviceSetSharedMemConfig` | `hipDeviceSetSharedMemConfig` | Sets the shared memory configuration for the current device. |
| `cudaDeviceSynchronize` | `hipDeviceSynchronize` | Wait for compute device to finish. |
| `cudaGetDevice` | `hipGetDevice` | Returns which device is currently being used. |
| `cudaGetDeviceCount` | `hipGetDeviceCount` | Returns the number of compute-capable devices. |
| `cudaGetDeviceFlags` | | Gets the flags for the current device. |
| `cudaGetDeviceProperties` | `hipGetDeviceProperties` | Returns information about the compute-device. |
| `cudaIpcCloseMemHandle` | | Close memory mapped with cudaIpcOpenMemHandle. |
| `cudaIpcGetEventHandle` | | Gets an interprocess handle for a previously allocated event. |
| `cudaIpcGetMemHandle` | | Gets an interprocess memory handle for an existing device memory allocation. |
| `cudaIpcOpenEventHandle` | | Opens an interprocess event handle for use in the current process. |
| `cudaIpcOpenMemHandle` | | Opens an interprocess memory handle exported from another process and returns a device pointer usable in the local process. |
| `cudaIpcCloseMemHandle` | `hipIpcCloseMemHandle` | Close memory mapped with cudaIpcOpenMemHandle. |
| `cudaIpcGetEventHandle` | `hipIpcGetEventHandle` | Gets an interprocess handle for a previously allocated event. |
| `cudaIpcGetMemHandle` | `hipIpcGetMemHandle` | Gets an interprocess memory handle for an existing device memory allocation. |
| `cudaIpcOpenEventHandle` | `hipIpcOpenEventHandle` | Opens an interprocess event handle for use in the current process. |
| `cudaIpcOpenMemHandle` | `hipIpcOpenMemHandle` | Opens an interprocess memory handle exported from another process and returns a device pointer usable in the local process. |
| `cudaSetDevice` | `hipSetDevice` | Set device to be used for GPU executions. |
| `cudaSetDeviceFlags` | `hipSetDeviceFlags` | Sets flags to be used for device executions. |
| `cudaSetValidDevices` | | Set a list of devices that can be used for CUDA. |`
| `cudaSetValidDevices` | | Set a list of devices that can be used for CUDA. |
**2. Error Handling**
## **2. Error Handling**
| **CUDA** | **HIP** | **CUDA description** |
|-----------------------------------------------------------|-------------------------------|--------------------------------------------------------------------------------------------------------------------------------|
@@ -37,12 +39,12 @@
| `cudaGetLastError` | `hipGetLastError` | Returns the last error from a runtime call. |
| `cudaPeekAtLastError` | `hipPeekAtLastError` | Returns the last error from a runtime call. |
**3. Stream Management**
## **3. Stream Management**
| **CUDA** | **HIP** | **CUDA description** |
|-----------------------------------------------------------|-------------------------------|--------------------------------------------------------------------------------------------------------------------------------|
| `cudaStreamAddCallback` | | Add a callback to a compute stream. |
| `cudaStreamAttachMemAsync` | | Attach managed memory to a stream asynchronously. |
| `cudaStreamAddCallback` | `hipStreamAddCallback` | Add a callback to a compute stream. |
| `cudaStreamAttachMemAsync` | | Attach managed memory to a stream asynchronously. |
| `cudaStreamCreate` | `hipStreamCreate` | Create an asynchronous stream. |
| `cudaStreamCreateWithFlags` | `hipStreamCreateWithFlags` | Create an asynchronous stream. |
| `cudaStreamCreateWithPriority` | | Create an asynchronous stream with the specified priority. |
@@ -53,7 +55,7 @@
| `cudaStreamSynchronize` | `hipStreamSynchronize` | Waits for stream tasks to complete. |
| `cudaStreamWaitEvent` | `hipStreamWaitEvent` | Make a compute stream wait on an event. |
**4. Event Management**
## **4. Event Management**
| **CUDA** | **HIP** | **CUDA description** |
|-----------------------------------------------------------|-------------------------------|--------------------------------------------------------------------------------------------------------------------------------|
@@ -65,7 +67,7 @@
| `cudaEventRecord` | `hipEventRecord` | Records an event. |
| `cudaEventSynchronize` | `hipEventSynchronize` | Waits for an event to complete. |
**5. Execution Control**
## **5. Execution Control**
| **CUDA** | **HIP** | **CUDA description** |
|-----------------------------------------------------------|-------------------------------|--------------------------------------------------------------------------------------------------------------------------------|
@@ -78,46 +80,54 @@
| `cudaSetDoubleForDevice` | | Converts a double argument to be executed on a device. |
| `cudaSetDoubleForHost` | | Converts a double argument after execution on a device. |
**6. Occupancy**
## **6. Occupancy**
| **CUDA** | **HIP** | **CUDA description** |
|-----------------------------------------------------------|-------------------------------|--------------------------------------------------------------------------------------------------------------------------------|
| `cudaOccupancyMaxActiveBlocksPerMultiprocessor` | | Returns occupancy for a device function. |
| `cudaOccupancyMaxActiveBlocksPerMultiprocessor` | `hipOccupancyMaxActiveBlocksPerMultiprocessor`| Returns occupancy for a device function. |
| `cudaOccupancyMaxActiveBlocksPerMultiprocessorWithFlags` | | Returns occupancy for a device function with the specified flags. |
**7. Memory Management**
## **7. Execution Control [deprecated since 7.0]**
| **CUDA** | **HIP** | **CUDA description** |
|-----------------------------------------------------------|-------------------------------|--------------------------------------------------------------------------------------------------------------------------------|
| `cudaConfigureCall` | | Configure a device-launch. |
| `cudaLaunch` | | Launches a device function. |
| `cudaSetupArgument` | | Configure a device launch. |
## **8. Memory Management**
| **CUDA** | **HIP** | **CUDA description** |
|-----------------------------------------------------------|-------------------------------|--------------------------------------------------------------------------------------------------------------------------------|
| `cudaArrayGetInfo` | | Gets info about the specified cudaArray. |
| `cudaFree` | `hipFree` | Frees memory on the device. |
| `cudaFreeArray` | | Frees an array on the device. |
| `cudaFreeArray` | `hipFreeArray` | Frees an array on the device. |
| `cudaFreeHost` | `hipHostFree` | Frees page-locked memory. |
| `cudaFreeMipmappedArray` | | Frees a mipmapped array on the device. |
| `cudaGetMipmappedArrayLevel` | | Gets a mipmap level of a CUDA mipmapped array. |
| `cudaGetSymbolAddress` | | Finds the address associated with a CUDA symbol. |
| `cudaGetSymbolSize` | | Finds the size of the object associated with a CUDA symbol. |
| `cudaHostAlloc` | `hipHostMalloc` | Allocates page-locked memory on the host. |
| `cudaHostGetDevicePointer` | `hipHostGetDevicePointer` | Passes back device pointer of mapped host memory allocated by cudaHostAlloc or registered by cudaHostRegister. |
| `cudaHostGetFlags` | `hipHostGetFlags` | Passes back flags used to allocate pinned host memory allocated by cudaHostAlloc. |
| `cudaHostAlloc` | `hipHostMalloc` | Allocates page-locked memory on the host. |
| `cudaHostGetDevicePointer` | `hipHostGetDevicePointer` | Passes back device pointer of mapped host memory allocated by cudaHostAlloc or registered by cudaHostRegister. |
| `cudaHostGetFlags` | `hipHostGetFlags` | Passes back flags used to allocate pinned host memory allocated by cudaHostAlloc. |
| `cudaHostRegister` | `hipHostRegister` | Registers an existing host memory range for use by CUDA. |
| `cudaHostUnregister` | `hipHostUnregister` | Unregisters a memory range that was registered with cudaHostRegister. |
| `cudaMalloc` | `hipMalloc` | Allocate memory on the device. |
| `cudaMalloc3D` | | Allocates logical 1D, 2D, or 3D memory objects on the device. |
| `cudaMalloc3DArray` | | Allocate an array on the device. |
| `cudaMallocArray` | | Allocate an array on the device. |
| `cudaMallocHost` | `hipHostMalloc` | Allocates page-locked memory on the host. |
| `cudaMallocArray` | `hipMallocArray` | Allocate an array on the device. |
| `cudaMallocHost` | `hipHostMalloc` | Allocates page-locked memory on the host. |
| `cudaMallocManaged` | | Allocates memory that will be automatically managed by the Unified Memory system. |
| `cudaMallocMipmappedArray` | | Allocate a mipmapped array on the device. |
| `cudaMallocPitch` | | Allocates pitched memory on the device. |
| `cudaMemGetInfo` | | Gets free and total device memory. |
| `cudaMemGetInfo` | `hipMemGetInfo` | Gets free and total device memory. |
| `cudaMemcpy` | `hipMemcpy` | Copies data between host and device. |
| `cudaMemcpy2D` | | Copies data between host and device. |
| `cudaMemcpy2D` | `hipMemcpy2D` | Copies data between host and device. |
| `cudaMemcpy2DArrayToArray` | | Copies data between host and device. |
| `cudaMemcpy2DAsync` | | Copies data between host and device. |
| `cudaMemcpy2DFromArray` | | Copies data between host and device. |
| `cudaMemcpy2DFromArrayAsync` | | Copies data between host and device. |
| `cudaMemcpy2DToArray` | | Copies data between host and device. |
| `cudaMemcpy2DToArray` | `hipMemcpy2DToArray` | Copies data between host and device. |
| `cudaMemcpy2DToArrayAsync` | | Copies data between host and device. |
| `cudaMemcpy3D` | | Copies data between 3D objects. |
| `cudaMemcpy3DAsync` | | Copies data between 3D objects. |
@@ -131,27 +141,27 @@
| `cudaMemcpyFromSymbolAsync` | | Copies data from the given symbol on the device. |
| `cudaMemcpyPeer` | `hipMemcpyPeer` | Copies memory between two devices. |
| `cudaMemcpyPeerAsync` | `hipMemcpyPeerAsync` | Copies memory between two devices asynchronously. |
| `cudaMemcpyToArray` | | Copies data between host and device. |
| `cudaMemcpyToArray` | `hipMemcpyToArray` | Copies data between host and device. |
| `cudaMemcpyToArrayAsync` | | Copies data between host and device. |
| `cudaMemcpyToSymbol` | `hipMemcpyToSymbol` | Copies data to the given symbol on the device. |
| `cudaMemcpyToSymbolAsync` | | Copies data to the given symbol on the device. |
| `cudaMemcpyToSymbolAsync` | `hipMemcpyToSymbolAsync` | Copies data to the given symbol on the device. |
| `cudaMemset` | `hipMemset` | Initializes or sets device memory to a value. |
| `cudaMemset2D` | | Initializes or sets device memory to a value. |
| `cudaMemset2DAsync` | | Initializes or sets device memory to a value. |
| `cudaMemset3D` | | Initializes or sets device memory to a value. |
| `cudaMemset3DAsync` | | Initializes or sets device memory to a value. |
| `cudaMemsetAsync` | `hipMemsetAsync` | Initializes or sets device memory to a value. |
| `make\_cudaExtent` | | Returns a cudaExtent based on input parameters. |
| `make\_cudaPitchedPtr` | | Returns a cudaPitchedPtr based on input parameters. |
| `make\_cudaPos` | | Returns a cudaPos based on input parameters. |
| `make_cudaExtent` | | Returns a cudaExtent based on input parameters. |
| `make_cudaPitchedPtr` | | Returns a cudaPitchedPtr based on input parameters. |
| `make_cudaPos` | | Returns a cudaPos based on input parameters. |
**8. Unified Addressing**
## **9. Unified Addressing**
| **CUDA** | **HIP** | **CUDA description** |
|-----------------------------------------------------------|-------------------------------|--------------------------------------------------------------------------------------------------------------------------------|
| `cudaPointerGetAttributes` | `hipPointerGetAttributes` | Returns attributes about a specified pointer. |
**9. Peer Device Memory Access**
## **10. Peer Device Memory Access**
| **CUDA** | **HIP** | **CUDA description** |
|-----------------------------------------------------------|-------------------------------|--------------------------------------------------------------------------------------------------------------------------------|
@@ -159,7 +169,7 @@
| `cudaDeviceDisablePeerAccess` | `hipDeviceDisablePeerAccess` | Disables direct access to memory allocations on a peer device. |
| `cudaDeviceEnablePeerAccess` | `hipDeviceEnablePeerAccess` | Enables direct access to memory allocations on a peer device. |
**10. OpenGL Interoperability**
## **11. OpenGL Interoperability**
| **CUDA** | **HIP** | **CUDA description** |
|-----------------------------------------------------------|-------------------------------|--------------------------------------------------------------------------------------------------------------------------------|
@@ -168,7 +178,7 @@
| `cudaGraphicsGLRegisterImage` | | Register an OpenGL texture or renderbuffer object. |
| `cudaWGLGetDevice` | | Gets the CUDA device associated with hGpu. |
**11. Graphics Interoperability**
## **12. Graphics Interoperability**
| **CUDA** | **HIP** | **CUDA description** |
|-----------------------------------------------------------|-------------------------------|--------------------------------------------------------------------------------------------------------------------------------|
@@ -180,7 +190,7 @@
| `cudaGraphicsUnmapResources` | | Unmap graphics resources. |
| `cudaGraphicsUnregisterResource` | | Unregisters a graphics resource for access by CUDA. |
**12. Texture Reference Management**
## **13. Texture Reference Management**
| **CUDA** | **HIP** | **CUDA description** |
|-----------------------------------------------------------|-------------------------------|--------------------------------------------------------------------------------------------------------------------------------|
@@ -194,14 +204,14 @@
| `cudaGetTextureReference` | | Get the texture reference associated with a symbol. |
| `cudaUnbindTexture` | | Unbinds a texture. |
**13. Surface Reference Management**
## **14. Surface Reference Management**
| **CUDA** | **HIP** | **CUDA description** |
|-----------------------------------------------------------|-------------------------------|--------------------------------------------------------------------------------------------------------------------------------|
| `cudaBindSurfaceToArray` | | Binds an array to a surface. |
| `cudaGetSurfaceReference` | | Get the surface reference associated with a symbol. |
**14. Texture Object Management**
## **15. Texture Object Management**
| **CUDA** | **HIP** | **CUDA description** |
|-----------------------------------------------------------|-------------------------------|--------------------------------------------------------------------------------------------------------------------------------|
@@ -211,7 +221,7 @@
| `cudaGetTextureObjectResourceViewDesc` | | Returns a texture object's resource view descriptor. |
| `cudaGetTextureObjectTextureDesc` | | Returns a texture object's texture descriptor. |
**15. Surface Object Management**
## **16. Surface Object Management**
| **CUDA** | **HIP** | **CUDA description** |
|-----------------------------------------------------------|-------------------------------|--------------------------------------------------------------------------------------------------------------------------------|
@@ -219,25 +229,24 @@
| `cudaDestroySurfaceObject` | | Destroys a surface object. |
| `cudaGetSurfaceObjectResourceDesc` | | Returns a surface object's resource descriptor Returns the resource descriptor for the surface object specified by surfObject. |
**16. Version Management**
## **17. Version Management**
| **CUDA** | **HIP** | **CUDA description** |
|-----------------------------------------------------------|-------------------------------|--------------------------------------------------------------------------------------------------------------------------------|
| `cudaDriverGetVersion` | `hipDriverGetVersion` | Returns the CUDA driver version. |
| `cudaRuntimeGetVersion` | | Returns the CUDA Runtime version. |
| `cudaRuntimeGetVersion` | `hipRuntimeGetVersion` | Returns the CUDA Runtime version. |
**17. C++ API Routines (7.0 contains, 7.5 doesn’t)**
> Will not support for HIP (probably)
## **18. C++ API Routines**
*(7.0 contains, 7.5 doesn’t)*
| **CUDA** | **HIP** | **CUDA description** |
|-----------------------------------------------------------|-------------------------------|--------------------------------------------------------------------------------------------------------------------------------|
| `cudaBindSurfaceToArray` | | Binds an array to a surface. |
| `cudaBindTexture` | | Binds a memory area to a texture. |
| `cudaBindTexture` | `hipBindTexture` | Binds a memory area to a texture. |
| `cudaBindTexture2D` | | Binds a 2D memory area to a texture. |
| `cudaBindTextureToArray` | | Binds an array to a texture. |
| `cudaBindTextureToMipmappedArray` | | Binds a mipmapped array to a texture. |
| `cudaCreateChannelDesc` | | Returns a channel descriptor using the specified format. |
| `cudaEventCreate` | | Creates an event object with the specified flags. |
| `cudaCreateChannelDesc` | `hipCreateChannelDesc` | Returns a channel descriptor using the specified format. |
| `cudaFuncGetAttributes` | | Find out attributes for a given function. |
| `cudaFuncSetCacheConfig` | | Sets the preferred cache configuration for a device function. |
| `cudaGetSymbolAddress` | | Finds the address associated with a CUDA symbol |
@@ -251,20 +260,117 @@
| `cudaMemcpyFromSymbolAsync` | | Copies data from the given symbol on the device. |
| `cudaMemcpyToSymbol` | | Copies data to the given symbol on the device. |
| `cudaMemcpyToSymbolAsync` | | Async copies data to the given symbol on the device. |
| `cudaOccupancyMaxActiveBlocksPerMultiprocessor` | | Returns occupancy for a device function. |
| `cudaOccupancyMaxActiveBlocksPerMultiprocessor` | `hipOccupancyMaxActiveBlocksPerMultiprocessor` | Returns occupancy for a device function. |
| `cudaOccupancyMaxActiveBlocksPerMultiprocessorWithFlags` | | Returns occupancy for a device function with the specified flags. |
| `cudaOccupancyMaxPotentialBlockSize` | | Returns grid and block size that achieves maximum potential occupancy for a device function. |
| `cudaOccupancyMaxPotentialBlockSize` | `hipOccupancyMaxPotentialBlockSize` | Returns grid and block size that achieves maximum potential occupancy for a device function. |
| `cudaOccupancyMaxPotentialBlockSizeVariableSMem` | | Returns grid and block size that achieves maximum potential occupancy for a device function. |
| `cudaOccupancyMaxPotentialBlockSizeVariableSMemWithFlags` | | Returns grid and block size that achieves maximum potential occupancy for a device function. |
| `cudaOccupancyMaxPotentialBlockSizeWithFlags` | | Returns grid and block size that achived maximum potential occupancy for a device function with the specified flags. |
| `cudaSetupArgument` | | Configure a device launch. |
| `cudaStreamAttachMemAsync` | | Attach memory to a stream asynchronously. |
| `cudaUnbindTexture` | | Unbinds a texture. |
| `cudaUnbindTexture` | `hipUnbindTexture` | Unbinds a texture. |
**18. Profiler Control**
## **19. Profiler Control**
| **CUDA** | **HIP** | **CUDA description** |
|-----------------------------------------------------------|-------------------------------|--------------------------------------------------------------------------------------------------------------------------------|
| `cudaProfilerInitialize` | | Initialize the CUDA profiler. |
| `cudaProfilerStart` | `hipProfilerStart` | Enable profiling. |
| `cudaProfilerStop` | `hipProfilerStop` | Disable profiling. |
# Data types used by CUDA Runtime API and supported by HIP
## **20. Data types**
| **type** | **CUDA** | **HIP** | **CUDA description** |
|--------------|--------------------------------------------|-------------------------------|--------------------------------------------------------------------------------------------------------------------------------|
| struct | `cudaChannelFormatDesc` | `hipChannelFormatDesc` | CUDA Channel format descriptor. |
| struct | `cudaDeviceProp` | `hipDeviceProp_t` | CUDA device properties. |
| struct | `cudaExtent` | | CUDA extent (width, height, depth). |
| struct | `cudaFuncAttributes` | | CUDA function attributes. |
| struct | `cudaIpcEventHandle_t` | `hipIpcEventHandle_t` | CUDA IPC event handle. |
| struct | `cudaIpcMemHandle_t` | `hipIpcMemHandle_t` | CUDA IPC memory handle. |
| struct | `cudaMemcpy3DParms` | | CUDA 3D memory copying parameters. |
| struct | `cudaMemcpy3DPeerParms` | | CUDA 3D cross-device memory copying parameters. |
| struct | `cudaPitchedPtr` | | CUDA Pitched memory pointer. |
| struct | `cudaPointerAttributes` | `hipPointerAttribute_t` | CUDA pointer attributes. |
| struct | `cudaPos` | | CUDA 3D position. |
| struct | `cudaResourceDesc` | | CUDA resource descriptor. |
| struct | `cudaResourceViewDesc` | | CUDA resource view descriptor. |
| struct | `cudaTextureDesc` | | CUDA texture descriptor. |
| struct | `surfaceReference` | | CUDA Surface reference. |
| struct | `textureReference` | `textureReference` | CUDA texture reference. |
| enum | `cudaChannelFormatKind` | `hipChannelFormatKind` | Channel format kind. |
| enum | `cudaComputeMode` | | CUDA device compute modes. |
| enum | `cudaDeviceAttr` | `hipDeviceAttribute_t` | CUDA device attributes. |
| enum | `cudaError` | `hipError_t` | CUDA Error types. |
| enum | `cudaError_t` | `hipError_t` | CUDA Error types. |
| enum | `cudaFuncCache` | `hipFuncCache_t` | CUDA function cache configurations. |
| enum | `cudaGraphicsCubeFace` | | CUDA graphics interop array indices for cube maps. |
| enum | `cudaGraphicsMapFlags` | | CUDA graphics interop map flags. |
| enum | `cudaGraphicsRegisterFlags` | | CUDA graphics interop register flags. |
| enum | `cudaMemcpyKind` | `hipMemcpyKind` | CUDA memory copy types. |
| enum | `cudaMemoryType` | `hipMemoryType` | CUDA memory types. |
| enum | `cudaOutputMode` | | CUDA Profiler Output modes. |
| enum | `cudaResourceType` | | CUDA resource types. |
| enum | `cudaResourceViewFormat` | | CUDA texture resource view formats. |
| enum | `cudaSharedMemConfig` | `hipSharedMemConfig` | CUDA shared memory configuration. |
| enum | `cudaSurfaceBoundaryMode` | | CUDA Surface boundary modes. |
| enum | `cudaSurfaceFormatMode` | | CUDA Surface format modes. |
| enum | `cudaTextureAddressMode` | | CUDA texture address modes. |
| enum | `cudaTextureFilterMode` | `hipTextureFilterMode` | CUDA texture filter modes. |
| enum | `cudaTextureReadMode` | `hipTextureReadMode` | CUDA texture read modes. |
| struct | `cudaArray` | `hipArray` | CUDA array [opaque]. |
| typedef | `cudaArray_t` | `hipArray *` | CUDA array pointer. |
| typedef | `cudaArray_const_t` | `const hipArray *` | CUDA array (as source copy argument). |
| enum | `cudaError` | `hipError_t` | CUDA Error types. |
| typedef | `cudaError_t` | `hipError_t` | CUDA Error types. |
| typedef | `cudaEvent_t` | `hipEvent_t` | CUDA event types. |
| typedef | `cudaGraphicsResource_t` | | CUDA graphics resource types. |
| typedef | `cudaMipmappedArray_t` | | CUDA mipmapped array. |
| typedef | `cudaMipmappedArray_const_t` | | CUDA mipmapped array (as source argument). |
| enum | `cudaOutputMode` | | CUDA output file modes. |
| typedef | `cudaOutputMode_t` | | CUDA output file modes. |
| typedef | `cudaStream_t` | `hipStream_t` | CUDA stream. |
| typedef | `cudaSurfaceObject_t` | | An opaque value that represents a CUDA Surface object. |
| typedef | `cudaTextureObject_t` | | An opaque value that represents a CUDA texture object. |
| typedef | `CUuuid_stcudaUUID_t` | | CUDA UUID types. |
| define | `CUDA_IPC_HANDLE_SIZE` | | CUDA IPC Handle Size. |
| define | `cudaArrayCubemap` | | Must be set in cudaMalloc3DArray to create a cubemap CUDA array. |
| define | `cudaArrayDefault` | | Default CUDA array allocation flag. |
| define | `cudaArrayLayered` | | Must be set in cudaMalloc3DArray to create a layered CUDA array. |
| define | `cudaArraySurfaceLoadStore` | | Must be set in cudaMallocArray or cudaMalloc3DArray in order to bind surfaces to the CUDA array. |
| define | `cudaArrayTextureGather` | | Must be set in cudaMallocArray or cudaMalloc3DArray in order to perform texture gather operations on the CUDA array. |
| define | `cudaDeviceBlockingSync` | `hipDeviceScheduleBlockingSync` | Device flag - Use blocking synchronization. Deprecated as of CUDA 4.0 and replaced with cudaDeviceScheduleBlockingSync. |
| define | `cudaDeviceLmemResizeToMax` | | Device flag - Keep local memory allocation after launch. |
| define | `cudaDeviceMapHost` | | Device flag - Support mapped pinned allocations. |
| define | `cudaDeviceMask` | | Device flags mask. |
| define | `cudaDevicePropDontCare` | | Empty device properties. |
| define | `cudaDeviceScheduleAuto` | `hipDeviceScheduleAuto` | Device flag - Automatic scheduling. |
| define | `cudaDeviceScheduleBlockingSync` | `hipDeviceScheduleBlockingSync` | Device flag - Use blocking synchronization. |
| define | `cudaDeviceScheduleMask` | `hipDeviceScheduleMask` | Device schedule flags mask. |
| define | `cudaDeviceScheduleSpin` | `hipDeviceScheduleSpin` | Device flag - Spin default scheduling. |
| define | `cudaDeviceScheduleYield` | `hipDeviceScheduleYield` | Device flag - Yield default scheduling. |
| define | `cudaEventBlockingSync` | `hipEventBlockingSync` | Event uses blocking synchronization. |
| define | `cudaEventDefault` | `hipEventDefault` | Default event flag. |
| define | `cudaEventDisableTiming` | `hipEventDisableTiming` | Event will not record timing data. |
| define | `cudaEventInterprocess` | `hipEventInterprocess` | Event is suitable for interprocess use. cudaEventDisableTiming must be set. |
| define | `cudaHostAllocDefault` | `hipHostMallocDefault` | Default page-locked allocation flag. |
| define | `cudaHostAllocMapped` | `hipHostMallocMapped` | Map allocation into device space. |
| define | `cudaHostAllocPortable` | `hipHostMallocPortable` | Pinned memory accessible by all CUDA contexts. |
| define | `cudaHostAllocWriteCombined` | `hipHostMallocWriteCombined` | Write-combined memory. |
| define | `cudaHostRegisterDefault` | `hipHostRegisterDefault` | Default host memory registration flag. |
| define | `cudaHostRegisterIoMemory` | `hipHostRegisterIoMemory` | Memory-mapped I/O space. |
| define | `cudaHostRegisterMapped` | `hipHostRegisterMapped` | Map registered memory into device space. |
| define | `cudaHostRegisterPortable` | `hipHostRegisterPortable` | Pinned memory accessible by all CUDA contexts. |
| define | `cudaIpcMemLazyEnablePeerAccess` | `hipIpcMemLazyEnablePeerAccess` | Automatically enable peer access between remote devices as needed. |
| define | `cudaMemAttachGlobal` | | Memory can be accessed by any stream on any device. |
| define | `cudaMemAttachHost` | | Memory cannot be accessed by any stream on any device. |
| define | `cudaMemAttachSingle` | | Memory can only be accessed by a single stream on the associated device. |
| define | `cudaOccupancyDefault` | | Default behavior. |
| define | `cudaOccupancyDisableCachingOverride` | | Assume global caching is enabled and cannot be automatically turned off. |
| define | `cudaPeerAccessDefault` | | Default peer addressing enable flag. |
| define | `cudaStreamDefault` | `hipStreamDefault` | Default stream flag. |
| define | `cudaStreamLegacy` | | Default stream flag. |
| define | `cudaStreamNonBlocking` | `hipStreamNonBlocking` | Stream does not synchronize with stream 0 (the NULL stream). |
| define | `cudaStreamPerThread` | | Per-thread stream handle. |
projects/hip/docs/markdown/device_md_gen.py
+508
Zobrazit soubor
@@ -0,0 +1,508 @@
"""
Copyright (c) 2015-2017 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.
"""
"""
1. This files uses Python3 to run
List of device functions:
acosf
acoshf
asinf
asinhf
atan2f
atanf
atanhf
cbrtf
ceilf
copysignf
cosf
coshf
cospif
cyl_bessel_i0f
cyl_bessel_i1f
erfcf
erfcinvf
erfcxf
erff
erfinvf
exp10f
exp2f
expf
expm1f
fabsf
fdimf
fdividef
floorf
fmaf
fmaxf
fminf
fmodf
frexpf
hypotf
ilogbf
isfinite
isinf
isnan
j0f
j1f
jnf
ldexpf
lgammaf
llrintf
llroundf
log10f
log1pf
logbf
lrintf
lroundf
modff
nanf
nearbyintf
nextafterf
norm3df
norm4df
normcdff
normcdfinvf
normf
powf
rcbrtf
remainderf
remquof
rhypotf
rintf
rnorm3df
rnorm4df
rnormf
roundf
rsqrtf
scalblnf
scalbnf
signbit
sincosf
sincospif
sinf
sinhf
sinpif
sqrtf
tanf
tanhf
tgammaf
truncf
y0f
y1f
ynf
acos
acosh
asin
asinh
atan
atan2
atanh
cbrt
ceil
copysign
cos
cosh
cospi
cyl_bessel_i0
cyl_bessel_i1
erf
erfc
erfcinv
erfcx
erfinv
exp
exp10
exp2
expm1
fabs
fdim
floor
fma
fmax
fmin
fmod
frexp
hypot
ilogb
isfinite
isinf
isnan
j0
j1
jn
ldexp
lgamma
llrint
llround
log
log10
log1p
log2
logb
lrint
lround
modf
nan
nearbyint
nextafter
norm
norm3d
norm4d
normcdf
normcdfinv
pow
rcbrt
remainder
remquo
rhypot
rint
rnorm
rnorm3d
rnorm4d
round
rsqrt
scalbln
scalbn
signbit
sin
sincos
sincospi
sinh
sinpi
sqrt
tan
tanh
tgamma
trunc
y0
y1
yn
__cosf
__exp10f
__expf
__fadd_rd
__fadd_rn
__fadd_ru
__fadd_rz
__fdiv_rd
__fdiv_rn
__fdiv_ru
__fdiv_rz
__fdividef
__fmaf_rd
__fmaf_rn
__fmaf_ru
__fmaf_rz
__fmul_rd
__fmul_rn
__fmul_ru
__fmul_rz
__frcp_rd
__frcp_rn
__frcp_ru
__frcp_rz
__frsqrt_rn
__fsqrt_rd
__fsqrt_rn
__fsqrt_ru
__fsqrt_rz
__fsub_rd
__fsub_rn
__fsub_ru
__log10f
__log2f
__logf
__powf
__saturatef
__sincosf
__sinf
__tanf
__dadd_rd
__dadd_rn
__dadd_ru
__dadd_rz
__ddiv_rd
__ddiv_rn
__ddiv_ru
__ddiv_rz
__dmul_rd
__dmul_rn
__dmul_ru
__dmul_rz
__drcp_rd
__drcp_rn
__drcp_ru
__drcp_rz
__dsqrt_rd
__dsqrt_rn
__dsqrt_ru
__dsqrt_rz
__dsub_rd
__dsub_rn
__dsub_ru
__dsub_rz
__fma_rd
__fma_rn
__fma_ru
__fma_rz
__brev
__brevll
__byte_perm
__clz
__clzll
__ffs
__ffsll
__hadd
__mul24
__mul64hi
__mulhi
__popc
__popcll
__rhadd
__sad
__uhadd
__umul24
__umul64hi
__umulhi
__urhadd
__usad
__double2float_rd
__double2float_rn
__double2float_ru
__double2float_rz
__double2hiint
__double2int_rd
__double2int_rn
__double2int_ru
__double2int_rz
__double2ll_rd
__double2ll_rn
__double2ll_ru
__double2ll_rz
__double2loint
__double2uint_rd
__double2uint_rn
__double2uint_ru
__double2uint_rz
__double2ull_rd
__double2ull_rn
__double2ull_ru
__double2ull_rz
__double_as_longlong
__float2half_rn
__half2float
__float2half_rn
__half2float
__float2int_rd
__float2int_rn
__float2int_ru
__float2int_rz
__float2ll_rd
__float2ll_rn
__float2ll_ru
__float2ll_rz
__float2uint_rd
__float2uint_rn
__float2uint_ru
__float2uint_rz
__float2ull_rd
__float2ull_rn
__float2ull_ru
__float2ull_rz
__float_as_int
__float_as_uint
__hiloint2double
__int2double_rn
__int2float_rd
__int2float_rn
__int2float_ru
__int2float_rz
__int_as_float
__ll2double_rd
__ll2double_rn
__ll2double_ru
__ll2double_rz
__ll2float_rd
__ll2float_rn
__ll2float_ru
__ll2float_rz
__longlong_as_double
__uint2double_rn
__uint2float_rd
__uint2float_rn
__uint2float_ru
__uint2float_rz
__uint_as_float
__ull2double_rd
__ull2double_rn
__ull2double_ru
__ull2double_rz
__ull2float_rd
__ull2float_rn
__ull2float_ru
__ull2float_rz
__heq
__hge
__hgt
__hisinf
__hisnan
__hle
__hlt
__hne
__hbeq2
__hbge2
__hbgt2
__hble2
__hblt2
__hbne2
__heq2
__hge2
__hgt2
__hisnan2
__hle2
__hlt2
__hne2
__float22half2_rn
__float2half
__float2half2_rn
__float2half_rd
__float2half_rn
__float2half_ru
__float2half_rz
__floats2half2_rn
__half22float2
__half2float
half2half2
__half2int_rd
__half2int_rn
__half2int_ru
__half2int_rz
__half2ll_rd
__half2ll_rn
__half2ll_ru
__half2ll_rz
__half2short_rd
__half2short_rn
__half2short_ru
__half2short_rz
__half2uint_rd
__half2uint_rn
__half2uint_ru
__half2uint_rz
__half2ull_rd
__half2ull_rn
__half2ull_ru
__half2ull_rz
__half2ushort_rd
__half2ushort_rn
__half2ushort_ru
__half2ushort_rz
__half_as_short
__half_as_ushort
__halves2half2
__high2float
__high2half
__high2half2
__highs2half2
__int2half_rd
__int2half_rn
__int2half_ru
__int2half_rz
__ll2half_rd
__ll2half_rn
__ll2half_ru
__ll2half_rz
__low2float
__low2half
__low2half2
__low2half2
__lowhigh2highlow
__lows2half2
__short2half_rd
__short2half_rn
__short2half_ru
__short2half_rz
__uint2half_rd
__uint2half_rn
__uint2half_ru
__uint2half_rz
__ull2half_rd
__ull2half_rn
__ull2half_ru
__ull2half_rz
__ushort2half_rd
__ushort2half_rn
__ushort2half_ru
__ushort2half_rz
__ushort_as_half
"""
# The dictionary is to place description of each device function. Expand it to all the device functions
deviceFuncDesc = {'acosf': "This function returns floating point of arc cosine from a floating point input"}
fnames = ["../../include/hip/hcc_detail/math_functions.h","../../include/hip/hcc_detail/device_functions.h","../../include/hip/hcc_detail/hip_fp16.h"]
markdownFileName = "./hip-math-api.md"
preamble = "# HIP MATH APIs Documentation \n"+\
"HIP supports most of the device functions supported by CUDA. Way to find the unsupported one is to search for the function and check its description\n" + \
"Note: This document is not human generated. Any changes to this file will be discarded. Please make changes to Python3 script docs/markdown/device_md_gen.py\n\n" + \
"## For Developers \n" + \
"If you add or fixed a device function, make sure to add a signature of the function and definition later.\n" + \
"For example, if you want to add `__device__ float __dotf(float4, float4)`, which does a dot product on 4 float vector components \n" + \
"The way to add to the header is, \n" + \
"```cpp \n" + \
"__device__ static float __dotf(float4, float4); \n" + \
"/*Way down in the file....*/\n" + \
"__device__ static inline float __dotf(float4 x, float4 y) { \n" + \
" /*implementation*/\n}\n" + \
"```\n\n" + \
"This helps python script to add the device function newly declared into markdown documentation (as it looks at functions with `;` at the end and `__device__` at the beginning)\n\n" + \
"The next step would be to add Description to `deviceFuncDesc` dictionary in python script.\n" + \
"From the above example, it can be writtern as,\n`deviceFuncDesc['__dotf'] = 'This functions takes 2 4 component float vector and outputs dot product across them'`\n\n"
def generateSnippet(name, description, signature):
return "### " + name + "\n" + \
"```cpp \n" + signature + "\n```\n" + \
"**Description:** " + description + "\n\n\n"
def getName(line):
l1 = line.split('(')
l2 = l1[0].split(' ')
return l2[-1]
with open(markdownFileName, 'w') as mdfd:
mdfd.truncate()
mdfd.write(preamble)
for fname in fnames:
with open(fname) as fd:
lines = fd.readlines()
for line in lines:
if line.find('HIP_FAST_MATH') != -1:
break;
if line.find('__device__') != -1 and line.find(';') != -1 and line.find('hip') == -1:
name = getName(line)
if line.find('//') == -1:
if name in deviceFuncDesc:
mdfd.write(generateSnippet(name, deviceFuncDesc[name], line))
else:
mdfd.write(generateSnippet(name, "Supported", line))
else:
mdfd.write(generateSnippet(name, "**NOT Supported**", line))
fd.close()
mdfd.close()
projects/hip/docs/markdown/hip-math-api.md
+3838
Zobrazit soubor
Rozdílový obsah nebyl zobrazen, protože je příliš veliký Načíst rozdílové porovnání
projects/hip/docs/markdown/hip_bugs.md
+46
Zobrazit soubor
@@ -0,0 +1,46 @@
# HIP Bugs
<!-- toc -->
- [Errors related to undefined reference to `__hcLaunchKernel__***__grid_launch_parm**](#errors-related-to-undefined-reference-to-hclaunchkernel__grid_launch_parm)
- [Application hangs after a hipLaunchKernel call](#what-if-i-see-application-hangs-after-a-hiplaunchkernel-call)
<!-- tocstop -->
### Errors related to undefined reference to `__hcLaunchKernel__***__grid_launch_parm**
Some common code practices may lead to hipcc generating a error with the form :
undefined reference to `__hcLaunchKernel__ZN15vecAddNamespace6vecAddIidEEv16grid_launch_parmPT0_S3_S3_T_
To workaround, try:
- Avoid calling hcLaunchKernel from a function with the __host__ attribute
__host__ MyFunc(…) {
hipLaunchKernel(myKernel, …)
- Avoid use of static with kernel definition:
static __global__ MyKernel
- Avoid defining kernels in anonymous namespace
namespace {
__global__ MyKernel …
- Avoid calling member functions
### What if I see application hangs after a hipLaunchKernel call?
If hipLaunchKernel takes parameters that request explicitly memcpy, then it will cause application hang.
Reason is that the hipLaunchKernel macro locks the stream.
If kernel paramters are actually function calls which invoke other hip apis (i.e. memcpy) to the same stream, then deadlock occurs.
To workaround, try:
Move the function calls so they occur outside the hipLaunchKernel macro, store results in temps, then use the tems inside the kernel.
```
// Example pseudo code causing system hang:
// "bottom[0]->gpu_data()" calls hipMemcpy() implicitly and using the same stream, cause deadlock condition.
hipLaunchKernel(HIP_KERNEL_NAME(LRNComputeDiff),dim3(CAFFE_GET_BLOCKS(n_threads)), dim3(CAFFE_HIP_NUM_THREADS), 0, 0, n_threads,
bottom[0]->gpu_data());
// Move "gpu_data()" ouside of hipLaunchKernel to avoid hang.
auto bot_gpu_data = bottom[0]->gpu_data();
hipLaunchKernel( LRNComputeDiff, dim3(CAFFE_GET_BLOCKS(n_threads)), dim3(CAFFE_HIP_NUM_THREADS), 0, 0, n_threads,
bot_gpu_data);
```
projects/hip/docs/markdown/hip_faq.md
+25 -25
Zobrazit soubor
@@ -50,7 +50,7 @@ At a high-level, the following features are not supported:
- Textures
- Dynamic parallelism (CUDA 5.0)
- Managed memory (CUDA 6.5)
- Graphics interoperation with OpenGL or Direct3D
- Graphics interoperability with OpenGL or Direct3D
- CUDA Driver API (Under Development)
- CUDA IPC Functions (Under Development)
- CUDA array, mipmappedArray and pitched memory
@@ -75,13 +75,13 @@ See the [API Support Table](CUDA_Runtime_API_functions_supported_by_HIP.md) for
### Is HIP a drop-in replacement for CUDA?
No. HIP provides porting tools which do most of the work do convert CUDA code into portable C++ code that uses the HIP APIs.
No. HIP provides porting tools which do most of the work to convert CUDA code into portable C++ code that uses the HIP APIs.
Most developers will port their code from CUDA to HIP and then maintain the HIP version.
HIP code provides the same performance as native CUDA code, plus the benefits of running on AMD platforms.
### What specific version of CUDA does HIP support?
HIP APIs and features do not map to a specific CUDA version. HIP provides a strong subset of functionality provided in CUDA, and the hipify tools can
scan code to identify any unsupported CUDA functions - this is very useful for identifying the specific features required by a given application.
HIP APIs and features do not map to a specific CUDA version. HIP provides a strong subset of functionality provided in CUDA, and the hipify tools can
scan code to identify any unsupported CUDA functions - this is useful for identifying the specific features required by a given application.
However, we can provide a rough summary of the features included in each CUDA SDK and the support level in HIP:
@@ -105,8 +105,8 @@ However, we can provide a rough summary of the features included in each CUDA SD
- TBD.
### What libraries does HIP support?
HIP includes growing support for the 4 key math libraries using hcBlas, hcFft, hcrng, and hcsparse).
These offer pointer-based memory interfaces (as opposed to opaque buffers) and can be easily interfaces with other HCC code. Developers should use conditional compilation if portability to nvcc systems is desired - using calls to cu* routines on one path and hc* routines on the other.
HIP includes growing support for the 4 key math libraries using hcBlas, hcFft, hcrng and hcsparse.
These offer pointer-based memory interfaces (as opposed to opaque buffers) and can be easily interfaced with other HCC applications. Developers should use conditional compilation if portability to nvcc systems is desired - using calls to cu* routines on one path and hc* routines on the other.
- [hcblas](https://bitbucket.org/multicoreware/hcblas)
- [hcfft](https://bitbucket.org/multicoreware/hcfft)
@@ -130,47 +130,47 @@ HIP offers several benefits over OpenCL:
### 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.
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 philosophy 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.
As a result, the OpenCL syntax is different from 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.
- 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 any device supporting CUDA SDK 6.0 or newer. We have tested the Nvidia Titan and Tesla K40.
### Does Hipify automatically convert all source code?
Typically, Hipify can automatically convert almost all run-time code, and the coordinate indexing device code (i.e. threadIdx.x -> hipThreadIdx_x).
Typically, hipify can automatically convert almost all run-time code, and the coordinate indexing device code ( threadIdx.x -> hipThreadIdx_x ).
Most device code needs no additional conversion, since HIP and CUDA have similar names for math and built-in functions.
The hipify-clang 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.
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.
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 compiler that is provided by the CUDA SDK.
### What is HCC?
HCC is AMD's compiler driver which compiles "heterogenous C++" code into HSAIL or GCN device code for AMD GPUs. It's an open-source compiler based on recent versions of CLANG/LLVM.
HCC is AMD's compiler driver which compiles "heterogeneous C++" code into HSAIL or GCN device code for AMD GPUs. It's an open-source compiler based on recent versions of CLANG/LLVM.
### Why use HIP rather than supporting CUDA directly?
While HIP is a strong subset of the CUDA, it is a subset. The HIP layer allows that subset to be clearly defined and documented.
Developers who code to the HIP API can be assured there code will remain portable across Nvidia and AMD platforms.
Developers who code to the HIP API can be assured their code will remain portable across Nvidia and AMD platforms.
In addition, HIP defines portable mechanisms to query architectural features, and supports a larger 64-bit wavesize which expands the return type for cross-lane functions like ballot and shuffle from 32-bit ints to 64-bit ints.
### Can I develop HIP code on an Nvidia CUDA platform?
Yes! HIP's CUDA path only exposes the APIs and functionality that work on both NVCC and HCC back-ends.
"Extra" APIs, parameters, and features which exist in CUDA but not in HCC will typically result in compile-time or run-time errors.
Developers need to use the HIP API for most accelerator code, and bracket any CUDA-specific code with appropriate ifdefs.
Yes. HIP's CUDA path only exposes the APIs and functionality that work on both NVCC and HCC back-ends.
"Extra" APIs, parameters, and features which exist in CUDA but not in HCC will typically result in compile- or run-time errors.
Developers need to use the HIP API for most accelerator code, and bracket any CUDA-specific code with preprocessor conditionals.
Developers concerned about portability should of course run on both platforms, and should expect to tune for performance.
In some cases CUDA has a richer set of modes for some APIs, and some C++ capabilities such as virtual functions - see the HIP @API documentation for more details.
### Can I develop HIP code on an AMD HCC platform?
Yes! HIP's HCC path only exposes the APIs and functions that work on both NVCC and HCC back ends. "Extra" APIs, parameters and features that appear in HCC but not CUDA will typically cause compile- or run-time errors. Developers must use the HIP API for most accelerator code and bracket any HCC-specific code with appropriate ifdefs. Those concerned about portability should, of course, test their code on both platforms and should tune it for performance. Typically, HCC supports a more modern set of C++11/C++14/C++17 features, so HIP developers who want portability should be careful when using advanced C++ features on the hc path.
Yes. HIP's HCC path only exposes the APIs and functions that work on both NVCC and HCC back ends. "Extra" APIs, parameters and features that appear in HCC but not CUDA will typically cause compile- or run-time errors. Developers must use the HIP API for most accelerator code and bracket any HCC-specific code with preprocessor conditionals. Those concerned about portability should, of course, test their code on both platforms and should tune it for performance. Typically, HCC supports a more modern set of C++11/C++14/C++17 features, so HIP developers who want portability should be careful when using advanced C++ features on the hc path.
### Can a HIP binary run on both AMD and Nvidia platforms?
HIP is a source-portable language that can be compiled to run on either the HCC or NVCC platform. HIP tools don't create a "fat binary" that can run on either platform, however.
@@ -183,9 +183,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.S
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 the generated object code
with GPU code compiled with HIP. Larger projects often contain a mixture of accelerator code (initially written in CUDA with nvcc) and 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 their preferred compiler.
@@ -197,7 +197,7 @@ Sometimes this isn't what you want - you can force HIP to recognize the platform
export HIP_PLATFORM=hcc
```
One symptom of this problem is the message "error: 'unknown error'(11) at square.hipref.cpp:56". This can occur if you have a CUDA installation on an AMD platform, and HIP incorrectly detects the platform as nvcc. HIP may be able to compile the application using the nvcc tool-chain, but will generate this error at runtime since the platform does not have a CUDA device. The fix is to set HIP_PLATFORM=hcc and rebuild the issue.
One symptom of this problem is the message "error: 'unknown error'(11) at square.hipref.cpp:56". This can occur if you have a CUDA installation on an AMD platform, and HIP incorrectly detects the platform as nvcc. HIP may be able to compile the application using the nvcc tool-chain, but will generate this error at runtime since the platform does not have a CUDA device. The fix is to set HIP_PLATFORM=hcc and rebuild.
If you see issues related to incorrect platform detection, please file an issue with the GitHub issue tracker so we can improve HIP's platform detection logic.
@@ -206,7 +206,7 @@ Yes. You can use HIP_PLATFORM to choose which path hipcc targets. This configur
### 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 .
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
@@ -217,10 +217,10 @@ If platform portability is important, use #ifdef __HIP_PLATFORM_NVCC__ to guard
### 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 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:
Also these functions can be used to extract HCC accelerator 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);
projects/hip/docs/markdown/hip_performance.md
+3 -3
Zobrazit soubor
@@ -6,18 +6,18 @@ Please note that this document lists possible ways for experimenting with HIP st
#### On Small BAR Setup
There are two possible ways to transfer data from Host to Device (H2D) and Device to Host(D2H)
There are two possible ways to transfer data from host-to-device (H2D) and device-to-host(D2H)
* Using Staging Buffers
* Using PinInPlace
#### On Large BAR Setup
There are three possible ways to transfer data from Host to Device (H2D)
There are three possible ways to transfer data from host-to-device (H2D)
* Using Staging Buffers
* Using PinInPlace
* Direct Memcpy
And there are two possible ways to transfer data from Device to Host (D2H)
And there are two possible ways to transfer data from device-to-host (D2H)
* Using Staging Buffers
* Using PinInPlace
projects/hip/docs/markdown/hip_porting_driver_api.md
+3 -3
Zobrazit soubor
@@ -1,7 +1,7 @@
# Porting CUDA Driver API
## Introduction to the CUDA Driver and Runtime APIs
CUDA provides a separate CUDA Driver and Runtime APIs. The two APis have significant overlap in functionality:
CUDA provides a separate CUDA Driver and Runtime APIs. The two APIs have significant overlap in functionality:
- Both APIs support events, streams, memory management, memory copy, and error handling.
- Both APIs deliver similar performance.
- Driver APIs calls begin with the prefix `cu` while Runtime APIs begin with the prefix `cuda`. For example, the Driver API API contains `cuEventCreate` while the Runtime API contains `cudaEventCreate`, with similar functionality.
@@ -90,14 +90,14 @@ the context. The current context is implicitly used by other APIs such as `hipS
The hipify tool will convert CUDA Driver APIs for streams, events, memory management to
the equivalent HIP driver calls. For example, `cuEventCreate` will be translated to
`hipEventCreate`. Hipify also converts error code from the Driver namespace and coding
convention to the equivalent HIP error code. Thus, HIP unifies the APis for these common functions.
convention to the equivalent HIP error code. Thus, HIP unifies the APIs for these common functions.
[hipify support for translating driver API is Under Development]
The memory copy APIs require additional explanation. The CUDA driver includes the memory
direction in the name of the API (ie `cuMemcpyH2D`) while the CUDA driver API provides
a single memory copy API with a parameter that specifies the direction and additionally
supports a "default" direction where the runtime determines the direction automatically.
HIP provides APis with both styles: for example, `hipMemcpyH2D` as well as `hipMemcpy`.
HIP provides APIs with both styles: for example, `hipMemcpyH2D` as well as `hipMemcpy`.
The first flavor may be faster in some cases since they avoid host overhead to detect the
different memory directions.
projects/hip/docs/markdown/hip_porting_guide.md
+1 -1
Zobrazit soubor
@@ -553,7 +553,7 @@ If you pass a ".cu" file, hcc will attempt to compile it as a Cuda language file
### HIP Environment Variables
On the HCC path, HIP provides a number of environment variables that control the behavior of HIP. Some of these are useful for appliction development (for example HIP_VISIBLE_DEVICES, HIP_LAUNCH_BLOCKING),
On the HCC path, HIP provides a number of environment variables that control the behavior of HIP. Some of these are useful for application development (for example HIP_VISIBLE_DEVICES, HIP_LAUNCH_BLOCKING),
some are useful for performance tuning or experimentation (for example HIP_STAGING*), and some are useful for debugging (HIP_DB). You can see the environment variables supported by HIP as well as
their current values and usage with the environment var "HIP_PRINT_ENV" - set this and then run any HIP application. For example:
projects/hip/docs/markdown/hip_profiling.md
+72 -7
Zobrazit soubor
@@ -92,11 +92,11 @@ HIP_PROFILE_API supports two levels of information.
#### Adding markers to applications
Markers can be used to define application-specific events that will be recorded in the ATP file and displayed in the CodeXL gui.
Markers can be used to define application-specific events that will be recorded in the ATP file and displayed in the CodeXL GUI.
This can be particularly useful for visualizing how the higher-level phases of application behavior relate to the lower level HIP APIs, kernel launches, and data transfers.
For example, an instrumented machine learning framework could show the beginning and ending of each layer in the network.
Markers have a specific begin and end time, and can be nested. Nested calls are displayed hierarchically in the CodeXL gui, with each level of the hierarchy occupying a different row.
Markers have a specific begin and end time, and can be nested. Nested calls are displayed hierarchically in the CodeXL GUI, with each level of the hierarchy occupying a different row.
The HIP APis are defined in "hip_profile.h":
```
@@ -131,7 +131,7 @@ The HIP marker API is only supported on ROCm platform. The marker macros are de
This [HIP sample](samples/2_Cookbook/2_Profiler/) shows the profiler marker API used in a small application.
More information on the marker API can be found in the profiler header file and PDF in a ROCM installation:
More information on the marker API can be found in the profiler header file and PDF in a ROCm installation:
- /opt/rocm/profiler/CXLActivityLogger/include/CXLActivityLogger.h
- /opt/rocm/profiler/CXLActivityLogger/doc/CXLActivityLogger.pdf
@@ -185,7 +185,7 @@ $ nvprof --profile-from-start-off ...
This feature is under development.
#### Reducing timeline trace output file size
If the application is already recording the HIP APIs, the HSA APIs are somewhat redundant and the ATP file size can be substantially reduced by not recording these APIs. HIP includes a text file that lists all of the HSA APis and can assist in this filtering:
If the application is already recording the HIP APIs, the HSA APIs are somewhat redundant and the ATP file size can be substantially reduced by not recording these APIs. HIP includes a text file that lists all of the HSA APIs and can assist in this filtering:
```
$ rocm-profiler -F hip/bin/hsa-api-filter-cxl.txt
@@ -273,7 +273,7 @@ None will disable use of color control codes for both the opening and closing an
This flag is primarily targeted to assist HIP development team in the development of the HIP runtime, but in some situations may be useful to HIP application developers as well.
The HIP debug information is designed to print important information during the execution of a HIP API. HIP provides
different color-coded levels of debug informaton:
different color-coded levels of debug information:
- api : Print the beginning and end of each HIP API, including the arguments and return codes. This is equivalent to setting HIP_TRACE_API=1.
- sync : Print multi-thread and other synchronization debug information.
- copy : Print which engine is doing the copy, which copy flavor is selected, information on source and destination memory.
@@ -315,7 +315,7 @@ libhsa-runtime64.so.1->hsaKmtDeregisterMemory(0x7f7776d3e010, 0x7f7776d3e010, 0x
```
Some key information from the trace above.
- The trace snippet shows the execution of a hipMemcpy API, bracketed by the first and last message in the trace output. The messages show the thread id and API sequence number (`1.17`). ltrace output intermixes messages from all threads, so the HIP debug information can be useful to determine which threads are executing.
- Thy trace snippet shows the execution of a hipMemcpy API, bracketed by the first and last message in the trace output. The messages show the thread id and API sequence number (`1.17`). ltrace output intermixes messages from all threads, so the HIP debug information can be useful to determine which threads are executing.
- The code flows through HIP APIs into ROCr (HSA) APIs (hsa*) and into the thunk (hsaKmt*) calls.
- The HCC runtime is "libmcwamp_hsa.so" and the HSA/ROCr runtime is "libhsa-runtime64.so".
- In this particular case, the memory copy is for unpinned memory, and the selected copy algorithm is to pin the host memory "in-place" before performing the copy. The signaling APIs and calls to pin ("lock", "register") the memory are readily apparent in the trace output.
@@ -349,7 +349,8 @@ These options cause HCC to serialize. Useful if you have libraries or code whic
$32 = {_shortTid = 1, _apiSeqNum = 803}
```
- HCC tracks all of the application memory allocations, including those from HIP and HC's "am_alloc". These can be printed by calling the function 'hc::am_memtracker_print()'.
- HCC tracks all of the application memory allocations, including those from HIP and HC's "am_alloc".
If the HCC runtime is built with debug information (HCC_RUNTIME_DEBUG=ON when building HCC), then calling the function 'hc::am_memtracker_print()' will show all memory allocations.
An optional argument specifies a void * targetPointer - the print routine will mark the allocation which contains the specified pointer with "-->" in the printed output.
This example shows a sample GDB session where we print the memory allocated by this process and mark a specified address by using the gdb "call" function..
The gdb syntax also supports using the variable name (in this case 'dst'):
@@ -361,4 +362,68 @@ TargetAddress:0x5ec7e9000
0x504cfc000-0x504cfc00f:: allocSeqNum:1 hostPointer:0x504cfc000 devicePointer:0x504cfc000 sizeBytes:16 isInDeviceMem:0 isAmManaged:1 appId:0 appAllocFlags:0 appPtr:(nil)
...
-->0x5ec7e9000-0x5f7e28fff:: allocSeqNum:488 hostPointer:(nil) devicePointer:0x5ec7e9000 sizeBytes:191102976 isInDeviceMem:1 isAmManaged:1 appId:0 appAllocFlags:0 appPtr:(nil)
```
To debug an explicit address, cast the address to (void*) :
```
(gdb) call hc::am_memtracker_print((void*)0x508c7f000)
```
- Debugging GPUVM fault.
For example:
```
Memory access fault by GPU node-1 on address 0x5924000. Reason: Page not present or supervisor privilege.
Program received signal SIGABRT, Aborted.
[Switching to Thread 0x7fffdffb5700 (LWP 14893)]
0x00007ffff2057c37 in __GI_raise (sig=sig@entry=6) at ../nptl/sysdeps/unix/sysv/linux/raise.c:56
56 ../nptl/sysdeps/unix/sysv/linux/raise.c: No such file or directory.
(gdb) bt
#0 0x00007ffff2057c37 in __GI_raise (sig=sig@entry=6) at ../nptl/sysdeps/unix/sysv/linux/raise.c:56
#1 0x00007ffff205b028 in __GI_abort () at abort.c:89
#2 0x00007ffff6f960eb in ?? () from /opt/rocm/hsa/lib/libhsa-runtime64.so.1
#3 0x00007ffff6f99ea5 in ?? () from /opt/rocm/hsa/lib/libhsa-runtime64.so.1
#4 0x00007ffff6f78107 in ?? () from /opt/rocm/hsa/lib/libhsa-runtime64.so.1
#5 0x00007ffff744f184 in start_thread (arg=0x7fffdffb5700) at pthread_create.c:312
#6 0x00007ffff211b37d in clone () at ../sysdeps/unix/sysv/linux/x86_64/clone.S:111
(gdb) info threads
Id Target Id Frame
4 Thread 0x7fffdd521700 (LWP 14895) "caffe" pthread_cond_wait@@GLIBC_2.3.2 () at ../nptl/sysdeps/unix/sysv/linux/x86_64/pthread_cond_wait.S:185
3 Thread 0x7fffddd22700 (LWP 14894) "caffe" pthread_cond_wait@@GLIBC_2.3.2 () at ../nptl/sysdeps/unix/sysv/linux/x86_64/pthread_cond_wait.S:185
* 2 Thread 0x7fffdffb5700 (LWP 14893) "caffe" 0x00007ffff2057c37 in __GI_raise (sig=sig@entry=6) at ../nptl/sysdeps/unix/sysv/linux/raise.c:56
1 Thread 0x7ffff7fa6ac0 (LWP 14892) "caffe" 0x00007ffff6f934d5 in ?? () from /opt/rocm/hsa/lib/libhsa-runtime64.so.1
(gdb) thread 1
[Switching to thread 1 (Thread 0x7ffff7fa6ac0 (LWP 14892))]
#0 0x00007ffff6f934d5 in ?? () from /opt/rocm/hsa/lib/libhsa-runtime64.so.1
(gdb) bt
#0 0x00007ffff6f934d5 in ?? () from /opt/rocm/hsa/lib/libhsa-runtime64.so.1
#1 0x00007ffff6f929ba in ?? () from /opt/rocm/hsa/lib/libhsa-runtime64.so.1
#2 0x00007fffe080beca in HSADispatch::waitComplete() () from /opt/rocm/hcc/lib/libmcwamp_hsa.so
#3 0x00007fffe080415f in HSADispatch::dispatchKernelAsync(Kalmar::HSAQueue*, void const*, int, bool) () from /opt/rocm/hcc/lib/libmcwamp_hsa.so
#4 0x00007fffe080238e in Kalmar::HSAQueue::dispatch_hsa_kernel(hsa_kernel_dispatch_packet_s const*, void const*, unsigned long, hc::completion_future*) () from /opt/rocm/hcc/lib/libmcwamp_hsa.so
#5 0x00007ffff7bb7559 in hipModuleLaunchKernel () from /opt/rocm/hip/lib/libhip_hcc.so
#6 0x00007ffff2e6cd2c in mlopen::HIPOCKernel::run (this=0x7fffffffb5a8, args=0x7fffffffb2a8, size=80) at /root/MIOpen/src/hipoc/hipoc_kernel.cpp:15
...
```
### General Debugging Tips
- The fault will be caught by the runtime but was actually generated by an asynchronous command running on the GPU. So, the GDB backtrace will show a path in the runtime, ie inside "GI_Raise" as shown in the example above.
- To determine the true location of the fault, force the kernels to execute synchronously by seeing the environment variables HCC_SERIALIZE_KERNEL=3 HCC_SERIALIZE_COPY=3. This will force HCC to wait for the kernel to finish executing before retuning. If the fault occurs during the execution of a kernel, you can see the code which launched the kernel inside the backtrace. A bit of guesswork is required to determine which thread is actually causing the issue - typically it will the thread which is waiting inside the libhsa-runtime64.so.
- VM faults inside kernels can be caused byi:
- incorrect code (ie a for loop which extends past array boundaries), i
- memory issues - kernel arguments which are invalid (null pointers, unregistered host pointers, bad pointers).
- synchronization issues
- compiler issues (incorrect code generation from the compiler)
- runtime issues
-- General debug tips:
- 'gdb --args' can be used to conviently pass the executable and arguments to gdb.
- From inside GDB, you can set environment variables "set env". Note the command does not use an '=' sign:
```
(gdb) set env HIP_DB 1
```
Setting HIP_PRINT_ENV=1 and then running a HIP application will print the HIP environment variables, their current values, and usage info.
Setting HCC_PRINT_ENV=1 and then running a HCC application will print the HCC environment variables, their current values, and usage info.
projects/hip/hip-config.cmake.in
+65
Zobrazit soubor
@@ -0,0 +1,65 @@
@PACKAGE_INIT@
include(CMakeFindDependencyMacro OPTIONAL RESULT_VARIABLE _CMakeFindDependencyMacro_FOUND)
if (NOT _CMakeFindDependencyMacro_FOUND)
macro(find_dependency dep)
if (NOT ${dep}_FOUND)
set(cmake_fd_version)
if (${ARGC} GREATER 1)
set(cmake_fd_version ${ARGV1})
endif()
set(cmake_fd_exact_arg)
if(${CMAKE_FIND_PACKAGE_NAME}_FIND_VERSION_EXACT)
set(cmake_fd_exact_arg EXACT)
endif()
set(cmake_fd_quiet_arg)
if(${CMAKE_FIND_PACKAGE_NAME}_FIND_QUIETLY)
set(cmake_fd_quiet_arg QUIET)
endif()
set(cmake_fd_required_arg)
if(${CMAKE_FIND_PACKAGE_NAME}_FIND_REQUIRED)
set(cmake_fd_required_arg REQUIRED)
endif()
find_package(${dep} ${cmake_fd_version}
${cmake_fd_exact_arg}
${cmake_fd_quiet_arg}
${cmake_fd_required_arg}
)
string(TOUPPER ${dep} cmake_dep_upper)
if (NOT ${dep}_FOUND AND NOT ${cmake_dep_upper}_FOUND)
set(${CMAKE_FIND_PACKAGE_NAME}_NOT_FOUND_MESSAGE "${CMAKE_FIND_PACKAGE_NAME} could not be found because dependency ${dep} could not be found.")
set(${CMAKE_FIND_PACKAGE_NAME}_FOUND False)
return()
endif()
set(cmake_fd_version)
set(cmake_fd_required_arg)
set(cmake_fd_quiet_arg)
set(cmake_fd_exact_arg)
endif()
endmacro()
endif()
set_and_check( hip_INCLUDE_DIR "@PACKAGE_INCLUDE_INSTALL_DIR@" )
set_and_check( hip_INCLUDE_DIRS "${hip_INCLUDE_DIR}" )
set_and_check( hip_LIB_INSTALL_DIR "@PACKAGE_LIB_INSTALL_DIR@" )
set_and_check( hip_BIN_INSTALL_DIR "@PACKAGE_BIN_INSTALL_DIR@" )
set_and_check(hip_HIPCC_EXECUTABLE "${hip_BIN_INSTALL_DIR}/hipcc")
set_and_check(hip_HIPCONFIG_EXECUTABLE "${hip_BIN_INSTALL_DIR}/hipconfig")
find_dependency(hcc)
include( "${CMAKE_CURRENT_LIST_DIR}/hip-targets.cmake" )
set( hip_LIBRARIES hip::hip_hcc)
set( hip_LIBRARY ${hip_LIBRARIES})
set(HIP_INCLUDE_DIR ${hip_INCLUDE_DIR})
set(HIP_INCLUDE_DIRS ${hip_INCLUDE_DIRS})
set(HIP_LIB_INSTALL_DIR ${hip_LIB_INSTALL_DIR})
set(HIP_BIN_INSTALL_DIR ${hip_BIN_INSTALL_DIR})
set(HIP_LIBRARIES ${hip_LIBRARIES})
set(HIP_LIBRARY ${hip_LIBRARY})
set(HIP_HIPCC_EXECUTABLE ${hip_HIPCC_EXECUTABLE})
set(HIP_HIPCONFIG_EXECUTABLE ${hip_HIPCONFIG_EXECUTABLE})
projects/hip/hipify-clang/README.md
+10
Zobrazit soubor
@@ -48,6 +48,16 @@ sudo dpkg -i cuda-repo-ubuntu1404_7.5-18_amd64.deb
sudo apt-get update && sudo apt-get install cuda-minimal-build-7-5 cuda-curand-dev-7-5
```
To set additional options like Language Selection (only "-x cuda" is supported), Preprocessor Definition (-D), Include Path (-I), etc., options delimiter "--" should be used before them, for instance:
```shell
./hipify-clang -print-stats sort_kernel.cu -- -x cuda -I/srv/git/HIP/include -I/usr/local/cuda-7.5/include -DX=1
```
Delimiter "--" is used to separate hipify-clang options (before the delimiter) from clang options (after the delimiter). It is strongly recommended to always specify the delimiter, even if there are no clang specific options at all, in order to avoid possible errors regarding compilation database; in such case delimeter should be the last option in hipify-clang's command line.
Option "-x clang" is also worth specifying in order to convert source CUDA files with extensions other than standard extensions (*.cu, *.cuh).
#### Disclaimer
The information contained herein is for informational purposes only, and is subject to change without notice. While every precaution has been taken in the preparation of this document, it may contain technical inaccuracies, omissions and typographical errors, and AMD is under no obligation to update or otherwise correct this information. Advanced Micro Devices, Inc. makes no representations or warranties with respect to the accuracy or completeness of the contents of this document, and assumes no liability of any kind, including the implied warranties of noninfringement, merchantability or fitness for particular purposes, with respect to the operation or use of AMD hardware, software or other products described herein. No license, including implied or arising by estoppel, to any intellectual property rights is granted by this document. Terms and limitations applicable to the purchase or use of AMD's products are as set forth in a signed agreement between the parties or in AMD's Standard Terms and Conditions of Sale.
projects/hip/hipify-clang/src/Cuda2Hip.cpp
+608 -324
Zobrazit soubor
Rozdílový obsah nebyl zobrazen, protože je příliš veliký Načíst rozdílové porovnání
projects/hip/include/hip/channel_descriptor.h
+36
Zobrazit soubor
@@ -0,0 +1,36 @@
/*
Copyright (c) 2015-2017 Advanced Micro Devices, Inc. All rights reserved.
Permission is hereby granted, free of charge, to any person obtaining a copy
of this software and associated documentation files (the "Software"), to deal
in the Software without restriction, including without limitation the rights
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
copies of the Software, and to permit persons to whom the Software is
furnished to do so, subject to the following conditions:
The above copyright notice and this permission notice shall be included in
all copies or substantial portions of the Software.
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
#pragma once
// Some standard header files, these are included by hc.hpp and so want to make them avail on both
// paths to provide a consistent include env and avoid "missing symbol" errors that only appears
// on NVCC path:
#if defined(__HIP_PLATFORM_HCC__) && !defined (__HIP_PLATFORM_NVCC__)
#include <hip/hcc_detail/channel_descriptor.h>
#elif defined(__HIP_PLATFORM_NVCC__) && !defined (__HIP_PLATFORM_HCC__)
#include <hip/nvcc_detail/channel_descriptor.h>
#else
#error("Must define exactly one of __HIP_PLATFORM_HCC__ or __HIP_PLATFORM_NVCC__");
#endif
projects/hip/include/hip/driver_types.h
+29
Zobrazit soubor
@@ -0,0 +1,29 @@
/*
Copyright (c) 2015-2017 Advanced Micro Devices, Inc. All rights reserved.
Permission is hereby granted, free of charge, to any person obtaining a copy
of this software and associated documentation files (the "Software"), to deal
in the Software without restriction, including without limitation the rights
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
copies of the Software, and to permit persons to whom the Software is
furnished to do so, subject to the following conditions:
The above copyright notice and this permission notice shall be included in
all copies or substantial portions of the Software.
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
#pragma once
#include <hip/hip_common.h>
#if defined(__HIP_PLATFORM_HCC__) && !defined (__HIP_PLATFORM_NVCC__)
#include <hip/hcc_detail/driver_types.h>
#elif defined(__HIP_PLATFORM_NVCC__) && !defined (__HIP_PLATFORM_HCC__)
#include "driver_types.h"
#else
#error("Must define exactly one of __HIP_PLATFORM_HCC__ or __HIP_PLATFORM_NVCC__");
#endif
projects/hip/include/hip/hcc_detail/channel_descriptor.h
+382
Zobrazit soubor
@@ -0,0 +1,382 @@
/*
Copyright (c) 2015-2017 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.
*/
#ifndef HIP_HCC_DETAIL_CHANNEL_DESCRIPTOR_H
#define HIP_HCC_DETAIL_CHANNEL_DESCRIPTOR_H
#include<hip/hcc_detail/driver_types.h>
#include<hip/hcc_detail/hip_vector_types.h>
#ifdef __cplusplus
hipChannelFormatDesc hipCreateChannelDesc(int x, int y, int z, int w, hipChannelFormatKind f);
static inline hipChannelFormatDesc hipCreateChannelDescHalf() {
int e = (int)sizeof(unsigned short) * 8;
return hipCreateChannelDesc(e, 0, 0, 0, hipChannelFormatKindFloat);
}
static inline hipChannelFormatDesc hipCreateChannelDescHalf1() {
int e = (int)sizeof(unsigned short) * 8;
return hipCreateChannelDesc(e, 0, 0, 0, hipChannelFormatKindFloat);
}
static inline hipChannelFormatDesc hipCreateChannelDescHalf2()
{
int e = (int)sizeof(unsigned short) * 8;
return hipCreateChannelDesc(e, 0, 0, 0, hipChannelFormatKindFloat);
}
template<typename T>
static inline hipChannelFormatDesc hipCreateChannelDesc() {
return hipCreateChannelDesc(0, 0, 0, 0, hipChannelFormatKindNone);
}
template<>
inline hipChannelFormatDesc hipCreateChannelDesc<char>()
{
int e = (int)sizeof(char) * 8;
return hipCreateChannelDesc(e, 0, 0, 0, hipChannelFormatKindSigned);
}
template<>
inline hipChannelFormatDesc hipCreateChannelDesc<signed char>()
{
int e = (int)sizeof(signed char) * 8;
return hipCreateChannelDesc(e, 0, 0, 0, hipChannelFormatKindUnsigned);
}
template<>
inline hipChannelFormatDesc hipCreateChannelDesc<unsigned char>()
{
int e = (int)sizeof(unsigned char) * 8;
return hipCreateChannelDesc(e, 0, 0, 0, hipChannelFormatKindUnsigned);
}
template<>
inline hipChannelFormatDesc hipCreateChannelDesc<uchar1>()
{
int e = (int)sizeof(unsigned char) * 8;
return hipCreateChannelDesc(e, 0, 0, 0, hipChannelFormatKindSigned);
}
template<>
inline hipChannelFormatDesc hipCreateChannelDesc<char1>()
{
int e = (int)sizeof(signed char) * 8;
return hipCreateChannelDesc(e, 0, 0, 0, hipChannelFormatKindSigned);
}
template<>
inline hipChannelFormatDesc hipCreateChannelDesc<uchar2>()
{
int e = (int)sizeof(unsigned char) * 8;
return hipCreateChannelDesc(e, e, 0, 0, hipChannelFormatKindSigned);
}
template<>
inline hipChannelFormatDesc hipCreateChannelDesc<char2>()
{
int e = (int)sizeof(signed char) * 8;
return hipCreateChannelDesc(e, e, 0, 0, hipChannelFormatKindSigned);
}
template<>
inline hipChannelFormatDesc hipCreateChannelDesc<uchar3>()
{
int e = (int)sizeof(unsigned char) * 8;
return hipCreateChannelDesc(e, e, e, 0, hipChannelFormatKindSigned);
}
template<>
inline hipChannelFormatDesc hipCreateChannelDesc<char3>()
{
int e = (int)sizeof(signed char) * 8;
return hipCreateChannelDesc(e, e, e, 0, hipChannelFormatKindSigned);
}
template<>
inline hipChannelFormatDesc hipCreateChannelDesc<uchar4>()
{
int e = (int)sizeof(unsigned char) * 8;
return hipCreateChannelDesc(e, e, e, e, hipChannelFormatKindSigned);
}
template<>
inline hipChannelFormatDesc hipCreateChannelDesc<char4>()
{
int e = (int)sizeof(signed char) * 8;
return hipCreateChannelDesc(e, e, e, e, hipChannelFormatKindSigned);
}
template<>
inline hipChannelFormatDesc hipCreateChannelDesc<unsigned short>()
{
int e = (int)sizeof(unsigned short) * 8;
return hipCreateChannelDesc(e, 0, 0, 0, hipChannelFormatKindUnsigned);
}
template<>
inline hipChannelFormatDesc hipCreateChannelDesc<signed short>()
{
int e = (int)sizeof(signed short) * 8;
return hipCreateChannelDesc(e, 0, 0, 0, hipChannelFormatKindSigned);
}
template<>
inline hipChannelFormatDesc hipCreateChannelDesc<ushort1>()
{
int e = (int)sizeof(unsigned short) * 8;
return hipCreateChannelDesc(e, 0, 0, 0, hipChannelFormatKindUnsigned);
}
template<>
inline hipChannelFormatDesc hipCreateChannelDesc<short1>()
{
int e = (int)sizeof(signed short) * 8;
return hipCreateChannelDesc(e, 0, 0, 0, hipChannelFormatKindSigned);
}
template<>
inline hipChannelFormatDesc hipCreateChannelDesc<ushort2>()
{
int e = (int)sizeof(unsigned short) * 8;
return hipCreateChannelDesc(e, e, 0, 0, hipChannelFormatKindUnsigned);
}
template<>
inline hipChannelFormatDesc hipCreateChannelDesc<short2>()
{
int e = (int)sizeof(signed short) * 8;
return hipCreateChannelDesc(e, e, 0, 0, hipChannelFormatKindSigned);
}
template<>
inline hipChannelFormatDesc hipCreateChannelDesc<ushort3>()
{
int e = (int)sizeof(unsigned short) * 8;
return hipCreateChannelDesc(e, e, e, 0, hipChannelFormatKindUnsigned);
}
template<>
inline hipChannelFormatDesc hipCreateChannelDesc<short3>()
{
int e = (int)sizeof(signed short) * 8;
return hipCreateChannelDesc(e, e, e, 0, hipChannelFormatKindSigned);
}
template<>
inline hipChannelFormatDesc hipCreateChannelDesc<ushort4>()
{
int e = (int)sizeof(unsigned short) * 8;
return hipCreateChannelDesc(e, e, e, e, hipChannelFormatKindUnsigned);
}
template<>
inline hipChannelFormatDesc hipCreateChannelDesc<short4>()
{
int e = (int)sizeof(signed short) * 8;
return hipCreateChannelDesc(e, e, e, e, hipChannelFormatKindSigned);
}
template<>
inline hipChannelFormatDesc hipCreateChannelDesc<unsigned int>()
{
int e = (int)sizeof(unsigned int) * 8;
return hipCreateChannelDesc(e, 0, 0, 0, hipChannelFormatKindUnsigned);
}
template<>
inline hipChannelFormatDesc hipCreateChannelDesc<signed int>()
{
int e = (int)sizeof(signed int) * 8;
return hipCreateChannelDesc(e, 0, 0, 0, hipChannelFormatKindSigned);
}
template<>
inline hipChannelFormatDesc hipCreateChannelDesc<uint1>()
{
int e = (int)sizeof(unsigned int) * 8;
return hipCreateChannelDesc(e, 0, 0, 0, hipChannelFormatKindUnsigned);
}
template<>
inline hipChannelFormatDesc hipCreateChannelDesc<int1>()
{
int e = (int)sizeof(signed int) * 8;
return hipCreateChannelDesc(e, 0, 0, 0, hipChannelFormatKindSigned);
}
template<>
inline hipChannelFormatDesc hipCreateChannelDesc<uint2>()
{
int e = (int)sizeof(unsigned int) * 8;
return hipCreateChannelDesc(e, e, 0, 0, hipChannelFormatKindUnsigned);
}
template<>
inline hipChannelFormatDesc hipCreateChannelDesc<int2>()
{
int e = (int)sizeof(signed int) * 8;
return hipCreateChannelDesc(e, e, 0, 0, hipChannelFormatKindSigned);
}
template<>
inline hipChannelFormatDesc hipCreateChannelDesc<uint3>()
{
int e = (int)sizeof(unsigned int) * 8;
return hipCreateChannelDesc(e, e, e, 0, hipChannelFormatKindUnsigned);
}
template<>
inline hipChannelFormatDesc hipCreateChannelDesc<int3>()
{
int e = (int)sizeof(signed int) * 8;
return hipCreateChannelDesc(e, e, e, 0, hipChannelFormatKindSigned);
}
template<>
inline hipChannelFormatDesc hipCreateChannelDesc<uint4>()
{
int e = (int)sizeof(unsigned int) * 8;
return hipCreateChannelDesc(e, e, e, e, hipChannelFormatKindUnsigned);
}
template<>
inline hipChannelFormatDesc hipCreateChannelDesc<int4>()
{
int e = (int)sizeof(signed int) * 8;
return hipCreateChannelDesc(e, e, e, e, hipChannelFormatKindSigned);
}
template<>
inline hipChannelFormatDesc hipCreateChannelDesc<float>()
{
int e = (int)sizeof(float) * 8;
return hipCreateChannelDesc(e, 0, 0, 0, hipChannelFormatKindFloat);
}
template<>
inline hipChannelFormatDesc hipCreateChannelDesc<float1>()
{
int e = (int)sizeof(float) * 8;
return hipCreateChannelDesc(e, 0, 0, 0, hipChannelFormatKindFloat);
}
template<>
inline hipChannelFormatDesc hipCreateChannelDesc<float2>()
{
int e = (int)sizeof(float) * 8;
return hipCreateChannelDesc(e, e, 0, 0, hipChannelFormatKindFloat);
}
template<>
inline hipChannelFormatDesc hipCreateChannelDesc<float3>()
{
int e = (int)sizeof(float) * 8;
return hipCreateChannelDesc(e, e, e, 0, hipChannelFormatKindFloat);
}
template<>
inline hipChannelFormatDesc hipCreateChannelDesc<float4>()
{
int e = (int)sizeof(float) * 8;
return hipCreateChannelDesc(e, e, e, e, hipChannelFormatKindFloat);
}
template<>
inline hipChannelFormatDesc hipCreateChannelDesc<unsigned long>()
{
int e = (int)sizeof(unsigned long) * 8;
return hipCreateChannelDesc(e, 0, 0, 0, hipChannelFormatKindUnsigned);
}
template<>
inline hipChannelFormatDesc hipCreateChannelDesc<signed long>()
{
int e = (int)sizeof(signed long) * 8;
return hipCreateChannelDesc(e, 0, 0, 0, hipChannelFormatKindSigned);
}
template<>
inline hipChannelFormatDesc hipCreateChannelDesc<ulong1>()
{
int e = (int)sizeof(unsigned long) * 8;
return hipCreateChannelDesc(e, 0, 0, 0, hipChannelFormatKindUnsigned);
}
template<>
inline hipChannelFormatDesc hipCreateChannelDesc<long1>()
{
int e = (int)sizeof(signed long) * 8;
return hipCreateChannelDesc(e, 0, 0, 0, hipChannelFormatKindSigned);
}
template<>
inline hipChannelFormatDesc hipCreateChannelDesc<ulong2>()
{
int e = (int)sizeof(unsigned long) * 8;
return hipCreateChannelDesc(e, e, 0, 0, hipChannelFormatKindUnsigned);
}
template<>
inline hipChannelFormatDesc hipCreateChannelDesc<long2>()
{
int e = (int)sizeof(signed long) * 8;
return hipCreateChannelDesc(e, e, 0, 0, hipChannelFormatKindSigned);
}
template<>
inline hipChannelFormatDesc hipCreateChannelDesc<ulong3>()
{
int e = (int)sizeof(unsigned long) * 8;
return hipCreateChannelDesc(e, e, e, 0, hipChannelFormatKindUnsigned);
}
template<>
inline hipChannelFormatDesc hipCreateChannelDesc<long3>()
{
int e = (int)sizeof(signed long) * 8;
return hipCreateChannelDesc(e, e, e, 0, hipChannelFormatKindSigned);
}
template<>
inline hipChannelFormatDesc hipCreateChannelDesc<ulong4>()
{
int e = (int)sizeof(unsigned long) * 8;
return hipCreateChannelDesc(e, e, e, e, hipChannelFormatKindUnsigned);
}
template<>
inline hipChannelFormatDesc hipCreateChannelDesc<long4>()
{
int e = (int)sizeof(signed long) * 8;
return hipCreateChannelDesc(e, e, e, e, hipChannelFormatKindSigned);
}
#else
struct hipChannelFormatDesc hipCreateChannelDesc(int x, int y, int z, int w, enum hipChannelFormatKind f);
#endif
#endif
projects/hip/include/hip/hcc_detail/device_functions.h
+569 -5
Zobrazit soubor
@@ -1,5 +1,5 @@
/*
Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved.
Copyright (c) 2015-2017 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
@@ -23,15 +23,579 @@ THE SOFTWARE.
#include <hip/hip_runtime.h>
#include <hip/hip_vector_types.h>
__device__ float __int_as_float (int x);
__device__ double __hiloint2double (int hi, int lo);
// Single Precision Fast Math
__device__ float __cosf(float x);
__device__ float __exp10f(float x);
__device__ float __expf(float x);
__device__ static float __fadd_rd(float x, float y);
__device__ static float __fadd_rn(float x, float y);
__device__ static float __fadd_ru(float x, float y);
__device__ static float __fadd_rz(float x, float y);
__device__ static float __fdiv_rd(float x, float y);
__device__ static float __fdiv_rn(float x, float y);
__device__ static float __fdiv_ru(float x, float y);
__device__ static float __fdiv_rz(float x, float y);
__device__ static float __fdividef(float x, float y);
__device__ float __fmaf_rd(float x, float y, float z);
__device__ float __fmaf_rn(float x, float y, float z);
__device__ float __fmaf_ru(float x, float y, float z);
__device__ float __fmaf_rz(float x, float y, float z);
__device__ static float __fmul_rd(float x, float y);
__device__ static float __fmul_rn(float x, float y);
__device__ static float __fmul_ru(float x, float y);
__device__ static float __fmul_rz(float x, float y);
__device__ float __frcp_rd(float x);
__device__ float __frcp_rn(float x);
__device__ float __frcp_ru(float x);
__device__ float __frcp_rz(float x);
__device__ float __frsqrt_rn(float x);
__device__ float __fsqrt_rd(float x);
__device__ float __fsqrt_rn(float x);
__device__ float __fsqrt_ru(float x);
__device__ float __fsqrt_rz(float x);
__device__ static float __fsub_rd(float x, float y);
__device__ static float __fsub_rn(float x, float y);
__device__ static float __fsub_ru(float x, float y);
__device__ float __log10f(float x);
__device__ float __log2f(float x);
__device__ float __logf(float x);
__device__ float __powf(float base, float exponent);
__device__ static float __saturatef(float x);
__device__ void __sincosf(float x, float *s, float *c);
__device__ float __sinf(float x);
__device__ float __tanf(float x);
/*
Double Precision Intrinsics
*/
__device__ static double __dadd_rd(double x, double y);
__device__ static double __dadd_rn(double x, double y);
__device__ static double __dadd_ru(double x, double y);
__device__ static double __dadd_rz(double x, double y);
__device__ static double __ddiv_rd(double x, double y);
__device__ static double __ddiv_rn(double x, double y);
__device__ static double __ddiv_ru(double x, double y);
__device__ static double __ddiv_rz(double x, double y);
__device__ static double __dmul_rd(double x, double y);
__device__ static double __dmul_rn(double x, double y);
__device__ static double __dmul_ru(double x, double y);
__device__ static double __dmul_rz(double x, double y);
__device__ double __drcp_rd(double x);
__device__ double __drcp_rn(double x);
__device__ double __drcp_ru(double x);
__device__ double __drcp_rz(double x);
__device__ double __dsqrt_rd(double x);
__device__ double __dsqrt_rn(double x);
__device__ double __dsqrt_ru(double x);
__device__ double __dsqrt_rz(double x);
__device__ static double __dsub_rd(double x, double y);
__device__ static double __dsub_rn(double x, double y);
__device__ static double __dsub_ru(double x, double y);
__device__ static double __dsub_rz(double x, double y);
__device__ double __fma_rd(double x, double y, double z);
__device__ double __fma_rn(double x, double y, double z);
__device__ double __fma_ru(double x, double y, double z);
__device__ double __fma_rz(double x, double y, double z);
// Single Precision Fast Math
extern __attribute__((const)) float __hip_fast_cosf(float) __asm("llvm.cos.f32");
extern __attribute__((const)) float __hip_fast_exp2f(float) __asm("llvm.exp2.f32");
__device__ float __hip_fast_exp10f(float);
__device__ float __hip_fast_expf(float);
__device__ float __hip_fast_frsqrt_rn(float);
extern __attribute__((const)) float __hip_fast_fsqrt_rd(float) __asm("llvm.sqrt.f32");
__device__ float __hip_fast_fsqrt_rn(float);
__device__ float __hip_fast_fsqrt_ru(float);
__device__ float __hip_fast_fsqrt_rz(float);
__device__ float __hip_fast_log10f(float);
extern __attribute__((const)) float __hip_fast_log2f(float) __asm("llvm.log2.f32");
__device__ float __hip_fast_logf(float);
__device__ float __hip_fast_powf(float, float);
__device__ void __hip_fast_sincosf(float,float*,float*);
extern __attribute__((const)) float __hip_fast_sinf(float) __asm("llvm.sin.f32");
__device__ float __hip_fast_tanf(float);
extern __attribute__((const)) float __hip_fast_fmaf(float,float,float) __asm("llvm.fma.f32");
extern __attribute__((const)) float __hip_fast_frcp(float) __asm("llvm.amdgcn.rcp.f32");
extern __attribute__((const)) double __hip_fast_dsqrt(double) __asm("llvm.sqrt.f64");
extern __attribute__((const)) double __hip_fast_fma(double,double,double) __asm("llvm.fma.f64");
extern __attribute__((const)) double __hip_fast_drcp(double) __asm("llvm.amdgcn.rcp.f64");
// Single Precision Fast Math
__device__ inline float __cosf(float x) {
return __hip_fast_cosf(x);
}
__device__ inline float __exp10f(float x) {
return __hip_fast_exp10f(x);
}
__device__ inline float __expf(float x) {
return __hip_fast_expf(x);
}
__device__ static inline float __fadd_rd(float x, float y) {
return x + y;
}
__device__ static inline float __fadd_rn(float x, float y) {
return x + y;
}
__device__ static inline float __fadd_ru(float x, float y) {
return x + y;
}
__device__ static inline float __fadd_rz(float x, float y) {
return x + y;
}
__device__ static inline float __fdiv_rd(float x, float y) {
return x / y;
}
__device__ static inline float __fdiv_rn(float x, float y) {
return x / y;
}
__device__ static inline float __fdiv_ru(float x, float y) {
return x / y;
}
__device__ static inline float __fdiv_rz(float x, float y) {
return x / y;
}
__device__ static inline float __fdividef(float x, float y) {
return x / y;
}
__device__ inline float __fmaf_rd(float x, float y, float z) {
return __hip_fast_fmaf(x, y, z);
}
__device__ inline float __fmaf_rn(float x, float y, float z) {
return __hip_fast_fmaf(x, y, z);
}
__device__ inline float __fmaf_ru(float x, float y, float z) {
return __hip_fast_fmaf(x, y, z);
}
__device__ inline float __fmaf_rz(float x, float y, float z) {
return __hip_fast_fmaf(x, y, z);
}
__device__ static inline float __fmul_rd(float x, float y) {
return x * y;
}
__device__ static inline float __fmul_rn(float x, float y) {
return x * y;
}
__device__ static inline float __fmul_ru(float x, float y) {
return x * y;
}
__device__ static inline float __fmul_rz(float x, float y) {
return x * y;
}
__device__ inline float __frcp_rd(float x) {
return __hip_fast_frcp(x);
}
__device__ inline float __frcp_rn(float x) {
return __hip_fast_frcp(x);
}
__device__ inline float __frcp_ru(float x) {
return __hip_fast_frcp(x);
}
__device__ inline float __frcp_rz(float x) {
return __hip_fast_frcp(x);
}
__device__ inline float __frsqrt_rn(float x) {
return __hip_fast_frsqrt_rn(x);
}
__device__ inline float __fsqrt_rd(float x) {
return __hip_fast_fsqrt_rd(x);
}
__device__ inline float __fsqrt_rn(float x) {
return __hip_fast_fsqrt_rn(x);
}
__device__ inline float __fsqrt_ru(float x) {
return __hip_fast_fsqrt_ru(x);
}
__device__ inline float __fsqrt_rz(float x) {
return __hip_fast_fsqrt_rz(x);
}
__device__ static inline float __fsub_rd(float x, float y) {
return x - y;
}
__device__ static inline float __fsub_rn(float x, float y) {
return x - y;
}
__device__ static inline float __fsub_ru(float x, float y) {
return x - y;
}
__device__ static inline float __fsub_rz(float x, float y) {
return x - y;
}
__device__ inline float __log10f(float x) {
return __hip_fast_log10f(x);
}
__device__ inline float __log2f(float x) {
return __hip_fast_log2f(x);
}
__device__ inline float __logf(float x) {
return __hip_fast_logf(x);
}
__device__ inline float __powf(float base, float exponent) {
return __hip_fast_powf(base, exponent);
}
__device__ static inline float __saturatef(float x) {
x = x > 1.0f ? 1.0f : x;
x = x < 0.0f ? 0.0f : x;
return x;
}
__device__ inline void __sincosf(float x, float *s, float *c) {
return __hip_fast_sincosf(x, s, c);
}
__device__ inline float __sinf(float x) {
return __hip_fast_sinf(x);
}
__device__ inline float __tanf(float x) {
return __hip_fast_tanf(x);
}
/*
Double Precision Intrinsics
*/
__device__ static inline double __dadd_rd(double x, double y) {
return x + y;
}
__device__ static inline double __dadd_rn(double x, double y) {
return x + y;
}
__device__ static inline double __dadd_ru(double x, double y) {
return x + y;
}
__device__ static inline double __dadd_rz(double x, double y) {
return x + y;
}
__device__ static inline double __ddiv_rd(double x, double y) {
return x / y;
}
__device__ static inline double __ddiv_rn(double x, double y) {
return x / y;
}
__device__ static inline double __ddiv_ru(double x, double y) {
return x / y;
}
__device__ static inline double __ddiv_rz(double x, double y) {
return x / y;
}
__device__ static inline double __dmul_rd(double x, double y) {
return x * y;
}
__device__ static inline double __dmul_rn(double x, double y) {
return x * y;
}
__device__ static inline double __dmul_ru(double x, double y) {
return x * y;
}
__device__ static inline double __dmul_rz(double x, double y) {
return x * y;
}
__device__ inline double __drcp_rd(double x) {
return __hip_fast_drcp(x);
}
__device__ inline double __drcp_rn(double x) {
return __hip_fast_drcp(x);
}
__device__ inline double __drcp_ru(double x) {
return __hip_fast_drcp(x);
}
__device__ inline double __drcp_rz(double x) {
return __hip_fast_drcp(x);
}
__device__ inline double __dsqrt_rd(double x) {
return __hip_fast_dsqrt(x);
}
__device__ inline double __dsqrt_rn(double x) {
return __hip_fast_dsqrt(x);
}
__device__ inline double __dsqrt_ru(double x) {
return __hip_fast_dsqrt(x);
}
__device__ inline double __dsqrt_rz(double x) {
return __hip_fast_dsqrt(x);
}
__device__ static inline double __dsub_rd(double x, double y) {
return x - y;
}
__device__ static inline double __dsub_rn(double x, double y) {
return x - y;
}
__device__ static inline double __dsub_ru(double x, double y) {
return x - y;
}
__device__ static inline double __dsub_rz(double x, double y) {
return x - y;
}
__device__ inline double __fma_rd(double x, double y, double z) {
return __hip_fast_fma(x, y, z);
}
__device__ inline double __fma_rn(double x, double y, double z) {
return __hip_fast_fma(x, y, z);
}
__device__ inline double __fma_ru(double x, double y, double z) {
return __hip_fast_fma(x, y, z);
}
__device__ inline double __fma_rz(double x, double y, double z) {
return __hip_fast_fma(x, y, z);
}
extern "C" unsigned int __hip_hc_ir_umul24_int(unsigned int, unsigned int);
extern "C" signed int __hip_hc_ir_mul24_int(signed int, signed int);
extern "C" signed int __hip_hc_ir_mulhi_int(signed int, signed int);
extern "C" unsigned int __hip_hc_ir_umulhi_int(unsigned int, unsigned int);
extern "C" unsigned int __hip_hc_ir_usad_int(unsigned int, unsigned int, unsigned int);
// integer intrinsic function __poc __clz __ffs __brev
__device__ unsigned int __brev( unsigned int x);
__device__ unsigned long long int __brevll( unsigned long long int x);
__device__ unsigned int __byte_perm(unsigned int x, unsigned int y, unsigned int s);
__device__ unsigned int __clz(int x);
__device__ unsigned int __clzll(long long int x);
__device__ unsigned int __ffs(int x);
__device__ unsigned int __ffsll(long long int x);
__device__ static unsigned int __hadd(int x, int y);
__device__ static int __mul24(int x, int y);
__device__ long long int __mul64hi(long long int x, long long int y);
__device__ static int __mulhi(int x, int y);
__device__ unsigned int __popc(unsigned int x);
__device__ unsigned int __popcll(unsigned long long int x);
__device__ static int __rhadd(int x, int y);
__device__ static unsigned int __sad(int x, int y, int z);
__device__ static unsigned int __uhadd(unsigned int x, unsigned int y);
__device__ static int __umul24(unsigned int x, unsigned int y);
__device__ unsigned long long int __umul64hi(unsigned long long int x, unsigned long long int y);
__device__ static unsigned int __umulhi(unsigned int x, unsigned int y);
__device__ static unsigned int __urhadd(unsigned int x, unsigned int y);
__device__ static unsigned int __usad(unsigned int x, unsigned int y, unsigned int z);
__device__ static inline unsigned int __hadd(int x, int y) {
int z = x + y;
int sign = z & 0x8000000;
int value = z & 0x7FFFFFFF;
return ((value) >> 1 || sign);
}
__device__ static inline int __mul24(int x, int y) {
return __hip_hc_ir_mul24_int(x, y);
}
__device__ static inline int __mulhi(int x, int y) {
return __hip_hc_ir_mulhi_int(x, y);
}
__device__ static inline int __rhadd(int x, int y) {
int z = x + y + 1;
int sign = z & 0x8000000;
int value = z & 0x7FFFFFFF;
return ((value) >> 1 || sign);
}
__device__ static inline unsigned int __sad(int x, int y, int z) {
return x > y ? x - y + z : y - x + z;
}
__device__ static inline unsigned int __uhadd(unsigned int x, unsigned int y) {
return (x + y) >> 1;
}
__device__ static inline int __umul24(unsigned int x, unsigned int y) {
return __hip_hc_ir_umul24_int(x, y);
}
__device__ static inline unsigned int __umulhi(unsigned int x, unsigned int y) {
return __hip_hc_ir_umulhi_int(x, y);
}
__device__ static inline unsigned int __urhadd(unsigned int x, unsigned int y) {
return (x + y + 1) >> 1;
}
__device__ static inline unsigned int __usad(unsigned int x, unsigned int y, unsigned int z)
{
return __hip_hc_ir_usad_int(x, y, z);
}
/*
Rounding modes are not yet supported in HIP
*/
__device__ float __double2float_rd(double x);
__device__ float __double2float_rn(double x);
__device__ float __double2float_ru(double x);
__device__ float __double2float_rz(double x);
__device__ int __double2hiint(double x);
__device__ int __double2int_rd(double x);
__device__ int __double2int_rn(double x);
__device__ int __double2int_ru(double x);
__device__ int __double2int_rz(double x);
__device__ long long int __double2ll_rd(double x);
__device__ long long int __double2ll_rn(double x);
__device__ long long int __double2ll_ru(double x);
__device__ long long int __double2ll_rz(double x);
__device__ int __double2loint(double x);
__device__ unsigned int __double2uint_rd(double x);
__device__ unsigned int __double2uint_rn(double x);
__device__ unsigned int __double2uint_ru(double x);
__device__ unsigned int __double2uint_rz(double x);
__device__ unsigned long long int __double2ull_rd(double x);
__device__ unsigned long long int __double2ull_rn(double x);
__device__ unsigned long long int __double2ull_ru(double x);
__device__ unsigned long long int __double2ull_rz(double x);
__device__ long long int __double_as_longlong(double x);
/*
__device__ unsigned short __float2half_rn(float x);
__device__ float __half2float(unsigned short);
The above device function are not a valid .
Use
__device__ __half __float2half_rn(float x);
__device__ float __half2float(__half);
from hip_fp16.h
CUDA implements half as unsigned short whereas, HIP doesn't.
*/
__device__ int __float2int_rd(float x);
__device__ int __float2int_rn(float x);
__device__ int __float2int_ru(float x);
__device__ int __float2int_rz(float x);
__device__ long long int __float2ll_rd(float x);
__device__ long long int __float2ll_rn(float x);
__device__ long long int __float2ll_ru(float x);
__device__ long long int __float2ll_rz(float x);
__device__ unsigned int __float2uint_rd(float x);
__device__ unsigned int __float2uint_rn(float x);
__device__ unsigned int __float2uint_ru(float x);
__device__ unsigned int __float2uint_rz(float x);
__device__ unsigned long long int __float2ull_rd(float x);
__device__ unsigned long long int __float2ull_rn(float x);
__device__ unsigned long long int __float2ull_ru(float x);
__device__ unsigned long long int __float2ull_rz(float x);
__device__ int __float_as_int(float x);
__device__ unsigned int __float_as_uint(float x);
__device__ double __hiloint2double(int hi, int lo);
__device__ double __int2double_rn(int x);
__device__ float __int2float_rd(int x);
__device__ float __int2float_rn(int x);
__device__ float __int2float_ru(int x);
__device__ float __int2float_rz(int x);
__device__ float __int_as_float(int x);
__device__ double __ll2double_rd(long long int x);
__device__ double __ll2double_rn(long long int x);
__device__ double __ll2double_ru(long long int x);
__device__ double __ll2double_rz(long long int x);
__device__ float __ll2float_rd(long long int x);
__device__ float __ll2float_rn(long long int x);
__device__ float __ll2float_ru(long long int x);
__device__ float __ll2float_rz(long long int x);
__device__ double __longlong_as_double(long long int x);
__device__ double __uint2double_rn(int x);
__device__ float __uint2float_rd(unsigned int x);
__device__ float __uint2float_rn(unsigned int x);
__device__ float __uint2float_ru(unsigned int x);
__device__ float __uint2float_rz(unsigned int x);
__device__ float __uint_as_float(unsigned int x);
__device__ double __ull2double_rd(unsigned long long int x);
__device__ double __ull2double_rn(unsigned long long int x);
__device__ double __ull2double_ru(unsigned long long int x);
__device__ double __ull2double_rz(unsigned long long int x);
__device__ float __ull2float_rd(unsigned long long int x);
__device__ float __ull2float_rn(unsigned long long int x);
__device__ float __ull2float_ru(unsigned long long int x);
__device__ float __ull2float_rz(unsigned long long int x);
__device__ char4 __hip_hc_add8pk(char4, char4);
__device__ char4 __hip_hc_sub8pk(char4, char4);
__device__ char4 __hip_hc_mul8pk(char4, char4);
extern __device__ double __longlong_as_double(long long int x);
extern __device__ long long int __double_as_longlong(double x);
#endif
projects/hip/include/hip/hcc_detail/driver_types.h
+41
Zobrazit soubor
@@ -0,0 +1,41 @@
/*
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.
*/
#ifndef HIP_HCC_DETAIL_DRIVER_TYPES_H
#define HIP_HCC_DETAIL_DRIVER_TYPES_H
enum hipChannelFormatKind
{
hipChannelFormatKindSigned = 0,
hipChannelFormatKindUnsigned = 1,
hipChannelFormatKindFloat = 2,
hipChannelFormatKindNone = 3
};
struct hipChannelFormatDesc
{
int x;
int y;
int z;
int w;
enum hipChannelFormatKind f;
};
#endif
projects/hip/include/hip/hcc_detail/hip_complex.h
+156 -5
Zobrazit soubor
@@ -1,5 +1,5 @@
/*
Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved.
Copyright (c) 2015-2017 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
@@ -20,11 +20,162 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
#ifndef HIPCOMPLEX_H
#define HIPCOMPLEX_H
#ifndef INCLUDE_HIP_HCC_DETAIL_HIP_COMPLEX_H
#define INCLUDE_HIP_HCC_DETAIL_HIP_COMPLEX_H
typedef float2 hipFloatComplex;
typedef double2 hipDoubleComplex;
#include "./hip_fp16.h"
#include "./hip_vector_types.h"
#if __cplusplus
#define COMPLEX_ADD_OP_OVERLOAD(type) \
__device__ __host__ static type operator + (const type& lhs, const type& rhs) { \
type ret; \
ret.x = lhs.x + rhs.x ; \
ret.y = lhs.y + rhs.y ; \
return ret; \
}
#define COMPLEX_SUB_OP_OVERLOAD(type) \
__device__ __host__ static type operator - (const type& lhs, const type& rhs) { \
type ret; \
ret.x = lhs.x - rhs.x; \
ret.y = lhs.y - rhs.y; \
return ret; \
}
#define COMPLEX_MUL_OP_OVERLOAD(type) \
__device__ __host__ static type operator * (const type& lhs, const type& rhs) { \
type ret; \
ret.x = lhs.x * rhs.x - lhs.y * rhs.y; \
ret.y = lhs.x * rhs.y + lhs.y * rhs.x; \
return ret; \
}
#define COMPLEX_DIV_OP_OVERLOAD(type) \
__device__ __host__ static type operator / (const type& lhs, const type& rhs) { \
type ret; \
ret.x = (lhs.x * rhs.x + lhs.y * rhs.y); \
ret.y = (rhs.x * lhs.y - lhs.x * rhs.y); \
ret.x = ret.x / (rhs.x * rhs.x + rhs.y * rhs.y); \
ret.y = ret.y / (rhs.x * rhs.x + rhs.y * rhs.y); \
return ret; \
}
#define COMPLEX_ADD_PREOP_OVERLOAD(type) \
__device__ __host__ static inline type& operator += (type& lhs, const type& rhs) { \
lhs.x += rhs.x; \
lhs.y += rhs.y; \
return lhs; \
}
#define COMPLEX_SUB_PREOP_OVERLOAD(type) \
__device__ __host__ static inline type& operator -= (type& lhs, const type& rhs) { \
lhs.x -= rhs.x; \
lhs.y -= rhs.y; \
return lhs; \
}
#define COMPLEX_MUL_PREOP_OVERLOAD(type) \
__device__ __host__ static inline type& operator *= (type& lhs, const type& rhs) { \
lhs = lhs * rhs; \
return lhs; \
}
#define COMPLEX_DIV_PREOP_OVERLOAD(type) \
__device__ __host__ static inline type& operator /= (type& lhs, const type& rhs) { \
lhs = lhs / rhs; \
return lhs; \
}
#define COMPLEX_SCALAR_PRODUCT(type, type1) \
__device__ __host__ static type operator * (const type& lhs, type1 rhs) { \
type ret; \
ret.x = lhs.x * rhs; \
ret.y = lhs.y * rhs; \
return ret; \
}
#endif
struct hipFloatComplex {
#ifdef __cplusplus
public:
__device__ __host__ hipFloatComplex() : x(0.0f), y(0.0f) {}
__device__ __host__ hipFloatComplex(float x) : x(x), y(0.0f) {}
__device__ __host__ hipFloatComplex(float x, float y) : x(x), y(y) {}
MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(hipFloatComplex, unsigned short)
MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(hipFloatComplex, signed short)
MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(hipFloatComplex, unsigned int)
MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(hipFloatComplex, signed int)
MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(hipFloatComplex, double)
MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(hipFloatComplex, unsigned long)
MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(hipFloatComplex, signed long)
MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(hipFloatComplex, unsigned long long)
MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(hipFloatComplex, signed long long)
#endif
float x, y;
} __attribute__((aligned(8)));
struct hipDoubleComplex {
#ifdef __cplusplus
public:
__device__ __host__ hipDoubleComplex() : x(0.0f), y(0.0f) {}
__device__ __host__ hipDoubleComplex(double x) : x(x), y(0.0f) {}
__device__ __host__ hipDoubleComplex(double x, double y) : x(x), y(y) {}
MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(hipDoubleComplex, unsigned short)
MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(hipDoubleComplex, signed short)
MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(hipDoubleComplex, unsigned int)
MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(hipDoubleComplex, signed int)
MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(hipDoubleComplex, float)
MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(hipDoubleComplex, unsigned long)
MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(hipDoubleComplex, signed long)
MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(hipDoubleComplex, unsigned long long)
MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(hipDoubleComplex, signed long long)
#endif
double x, y;
} __attribute__((aligned(16)));
#if __cplusplus
COMPLEX_ADD_OP_OVERLOAD(hipFloatComplex)
COMPLEX_SUB_OP_OVERLOAD(hipFloatComplex)
COMPLEX_MUL_OP_OVERLOAD(hipFloatComplex)
COMPLEX_DIV_OP_OVERLOAD(hipFloatComplex)
COMPLEX_ADD_PREOP_OVERLOAD(hipFloatComplex)
COMPLEX_SUB_PREOP_OVERLOAD(hipFloatComplex)
COMPLEX_MUL_PREOP_OVERLOAD(hipFloatComplex)
COMPLEX_DIV_PREOP_OVERLOAD(hipFloatComplex)
COMPLEX_SCALAR_PRODUCT(hipFloatComplex, unsigned short)
COMPLEX_SCALAR_PRODUCT(hipFloatComplex, signed short)
COMPLEX_SCALAR_PRODUCT(hipFloatComplex, unsigned int)
COMPLEX_SCALAR_PRODUCT(hipFloatComplex, signed int)
COMPLEX_SCALAR_PRODUCT(hipFloatComplex, float)
COMPLEX_SCALAR_PRODUCT(hipFloatComplex, unsigned long)
COMPLEX_SCALAR_PRODUCT(hipFloatComplex, signed long)
COMPLEX_SCALAR_PRODUCT(hipFloatComplex, double)
COMPLEX_SCALAR_PRODUCT(hipFloatComplex, signed long long)
COMPLEX_SCALAR_PRODUCT(hipFloatComplex, unsigned long long)
COMPLEX_ADD_OP_OVERLOAD(hipDoubleComplex)
COMPLEX_SUB_OP_OVERLOAD(hipDoubleComplex)
COMPLEX_MUL_OP_OVERLOAD(hipDoubleComplex)
COMPLEX_DIV_OP_OVERLOAD(hipDoubleComplex)
COMPLEX_ADD_PREOP_OVERLOAD(hipDoubleComplex)
COMPLEX_SUB_PREOP_OVERLOAD(hipDoubleComplex)
COMPLEX_MUL_PREOP_OVERLOAD(hipDoubleComplex)
COMPLEX_DIV_PREOP_OVERLOAD(hipDoubleComplex)
COMPLEX_SCALAR_PRODUCT(hipDoubleComplex, unsigned short)
COMPLEX_SCALAR_PRODUCT(hipDoubleComplex, signed short)
COMPLEX_SCALAR_PRODUCT(hipDoubleComplex, unsigned int)
COMPLEX_SCALAR_PRODUCT(hipDoubleComplex, signed int)
COMPLEX_SCALAR_PRODUCT(hipDoubleComplex, float)
COMPLEX_SCALAR_PRODUCT(hipDoubleComplex, unsigned long)
COMPLEX_SCALAR_PRODUCT(hipDoubleComplex, signed long)
COMPLEX_SCALAR_PRODUCT(hipDoubleComplex, double)
COMPLEX_SCALAR_PRODUCT(hipDoubleComplex, signed long long)
COMPLEX_SCALAR_PRODUCT(hipDoubleComplex, unsigned long long)
#endif
__device__ static inline float hipCrealf(hipFloatComplex z){
return z.x;
projects/hip/include/hip/hcc_detail/hip_fp16.h
+330 -92
Zobrazit soubor
@@ -1,5 +1,5 @@
/*
Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved.
Copyright (c) 2015-2017 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
@@ -20,8 +20,8 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
#ifndef HIP_FP16_H
#define HIP_FP16_H
#ifndef HIP_HCC_DETAIL_FP16_H
#define HIP_HCC_DETAIL_FP16_H
#include "hip/hip_runtime.h"
@@ -36,66 +36,211 @@ typedef struct __attribute__((aligned(4))){
};
} __half2;
struct holder{
union {
__half h;
unsigned short s;
};
};
typedef __half half;
typedef __half2 half2;
#define HINF 65504
/*
Half Arithmetic Functions
*/
__device__ __half __hadd(const __half a, const __half b);
__device__ __half __hadd_sat(__half a, __half b);
__device__ __half __hfma(__half a, __half b, __half c);
__device__ __half __hfma_sat(__half a, __half b, __half c);
__device__ __half __hmul(__half a, __half b);
__device__ __half __hmul_sat(__half a, __half b);
__device__ __half __hneg(__half a);
__device__ __half __hsub(__half a, __half b);
__device__ __half __hsub_sat(__half a, __half b);
__device__ __half hdiv(__half a, __half b);
static struct holder hInf = {HINF};
/*
Half2 Arithmetic Functions
*/
extern "C" __half __hip_hc_ir_hadd_half(__half, __half);
extern "C" __half __hip_hc_ir_hfma_half(__half, __half, __half);
extern "C" __half __hip_hc_ir_hmul_half(__half, __half);
extern "C" __half __hip_hc_ir_hsub_half(__half, __half);
__device__ static __half2 __hadd2(__half2 a, __half2 b);
__device__ static __half2 __hadd2_sat(__half2 a, __half2 b);
__device__ static __half2 __hfma2(__half2 a, __half2 b, __half2 c);
__device__ static __half2 __hfma2_sat(__half2 a, __half2 b, __half2 c);
__device__ static __half2 __hmul2(__half2 a, __half2 b);
__device__ static __half2 __hmul2_sat(__half2 a, __half2 b);
__device__ static __half2 __hsub2(__half2 a, __half2 b);
__device__ static __half2 __hneg2(__half2 a);
__device__ static __half2 __hsub2_sat(__half2 a, __half2 b);
__device__ static __half2 h2div(__half2 a, __half2 b);
/*
Half Comparision Functions
*/
__device__ bool __heq(__half a, __half b);
__device__ bool __hge(__half a, __half b);
__device__ bool __hgt(__half a, __half b);
__device__ bool __hisinf(__half a);
__device__ bool __hisnan(__half a);
__device__ bool __hle(__half a, __half b);
__device__ bool __hlt(__half a, __half b);
__device__ bool __hne(__half a, __half b);
/*
Half2 Comparision Functions
*/
__device__ bool __hbeq2(__half2 a, __half2 b);
__device__ bool __hbge2(__half2 a, __half2 b);
__device__ bool __hbgt2(__half2 a, __half2 b);
__device__ bool __hble2(__half2 a, __half2 b);
__device__ bool __hblt2(__half2 a, __half2 b);
__device__ bool __hbne2(__half2 a, __half2 b);
__device__ __half2 __heq2(__half2 a, __half2 b);
__device__ __half2 __hge2(__half2 a, __half2 b);
__device__ __half2 __hgt2(__half2 a, __half2 b);
__device__ __half2 __hisnan2(__half2 a);
__device__ __half2 __hle2(__half2 a, __half2 b);
__device__ __half2 __hlt2(__half2 a, __half2 b);
__device__ __half2 __hne2(__half2 a, __half2 b);
/*
Half Math Functions
*/
__device__ static __half hceil(const __half h);
__device__ static __half hcos(const __half h);
__device__ static __half hexp(const __half h);
__device__ static __half hexp10(const __half h);
__device__ static __half hexp2(const __half h);
__device__ static __half hfloor(const __half h);
__device__ static __half hlog(const __half h);
__device__ static __half hlog10(const __half h);
__device__ static __half hlog2(const __half h);
//__device__ static __half hrcp(const __half h);
__device__ static __half hrint(const __half h);
__device__ static __half hsin(const __half h);
__device__ static __half hsqrt(const __half a);
__device__ static __half htrunc(const __half a);
/*
Half2 Math Functions
*/
__device__ static __half2 h2ceil(const __half2 h);
__device__ static __half2 h2exp(const __half2 h);
__device__ static __half2 h2exp10(const __half2 h);
__device__ static __half2 h2exp2(const __half2 h);
__device__ static __half2 h2floor(const __half2 h);
__device__ static __half2 h2log(const __half2 h);
__device__ static __half2 h2log10(const __half2 h);
__device__ static __half2 h2log2(const __half2 h);
__device__ static __half2 h2rcp(const __half2 h);
__device__ static __half2 h2rsqrt(const __half2 h);
__device__ static __half2 h2sin(const __half2 h);
__device__ static __half2 h2sqrt(const __half2 h);
/*
Half Conversion And Data Movement
*/
__device__ __half2 __float22half2_rn(const float2 a);
__device__ __half __float2half(const float a);
__device__ __half2 __float2half2_rn(const float a);
__device__ __half __float2half_rd(const float a);
__device__ __half __float2half_rn(const float a);
__device__ __half __float2half_ru(const float a);
__device__ __half __float2half_rz(const float a);
__device__ __half2 __floats2half2_rn(const float a, const float b);
__device__ float2 __half22float2(const __half2 a);
__device__ float __half2float(const __half a);
__device__ __half2 half2half2(const __half a);
__device__ int __half2int_rd(__half h);
__device__ int __half2int_rn(__half h);
__device__ int __half2int_ru(__half h);
__device__ int __half2int_rz(__half h);
__device__ long long int __half2ll_rd(__half h);
__device__ long long int __half2ll_rn(__half h);
__device__ long long int __half2ll_ru(__half h);
__device__ long long int __half2ll_rz(__half h);
__device__ short __half2short_rd(__half h);
__device__ short __half2short_rn(__half h);
__device__ short __half2short_ru(__half h);
__device__ short __half2short_rz(__half h);
__device__ unsigned int __half2uint_rd(__half h);
__device__ unsigned int __half2uint_rn(__half h);
__device__ unsigned int __half2uint_ru(__half h);
__device__ unsigned int __half2uint_rz(__half h);
__device__ unsigned long long int __half2ull_rd(__half h);
__device__ unsigned long long int __half2ull_rn(__half h);
__device__ unsigned long long int __half2ull_ru(__half h);
__device__ unsigned long long int __half2ull_rz(__half h);
__device__ unsigned short int __half2ushort_rd(__half h);
__device__ unsigned short int __half2ushort_rn(__half h);
__device__ unsigned short int __half2ushort_ru(__half h);
__device__ unsigned short int __half2ushort_rz(__half h);
__device__ short int __half_as_short(const __half h);
__device__ unsigned short int __half_as_ushort(const __half h);
__device__ __half2 __halves2half2(const __half a, const __half b);
__device__ float __high2float(const __half2 a);
__device__ __half __high2half(const __half2 a);
__device__ __half2 __high2half2(const __half2 a);
__device__ __half2 __highs2half2(const __half2 a, const __half2 b);
__device__ __half __int2half_rd(int i);
__device__ __half __int2half_rn(int i);
__device__ __half __int2half_ru(int i);
__device__ __half __int2half_rz(int i);
__device__ __half __ll2half_rd(long long int i);
__device__ __half __ll2half_rn(long long int i);
__device__ __half __ll2half_ru(long long int i);
__device__ __half __ll2half_rz(long long int i);
__device__ float __low2float(const __half2 a);
__device__ __half __low2half(const __half2 a);
__device__ __half2 __low2half2(const __half2 a, const __half2 b);
__device__ __half2 __low2half2(const __half2 a);
__device__ __half2 __lowhigh2highlow(const __half2 a);
__device__ __half2 __lows2half2(const __half2 a, const __half2 b);
__device__ __half __short2half_rd(short int i);
__device__ __half __short2half_rn(short int i);
__device__ __half __short2half_ru(short int i);
__device__ __half __short2half_rz(short int i);
__device__ __half __uint2half_rd(unsigned int i);
__device__ __half __uint2half_rn(unsigned int i);
__device__ __half __uint2half_ru(unsigned int i);
__device__ __half __uint2half_rz(unsigned int i);
__device__ __half __ull2half_rd(unsigned long long int i);
__device__ __half __ull2half_rn(unsigned long long int i);
__device__ __half __ull2half_ru(unsigned long long int i);
__device__ __half __ull2half_rz(unsigned long long int i);
__device__ __half __ushort2half_rd(unsigned short int i);
__device__ __half __ushort2half_rn(unsigned short int i);
__device__ __half __ushort2half_ru(unsigned short int i);
__device__ __half __ushort2half_rz(unsigned short int i);
__device__ __half __ushort_as_half(const unsigned short int i);
extern "C" int __hip_hc_ir_hadd2_int(int, int);
extern "C" int __hip_hc_ir_hfma2_int(int, int, int);
extern "C" int __hip_hc_ir_hmul2_int(int, int);
extern "C" int __hip_hc_ir_hsub2_int(int, int);
__device__ static inline __half __hadd(const __half a, const __half b) {
return __hip_hc_ir_hadd_half(a, b);
}
extern "C" __half __hip_hc_ir_hceil_half(__half) __asm("llvm.ceil.f16");
extern "C" __half __hip_hc_ir_hcos_half(__half) __asm("llvm.cos.f16");
extern "C" __half __hip_hc_ir_hexp2_half(__half) __asm("llvm.exp2.f16");
extern "C" __half __hip_hc_ir_hfloor_half(__half) __asm("llvm.floor.f16");
extern "C" __half __hip_hc_ir_hlog2_half(__half) __asm("llvm.log2.f16");
extern "C" __half __hip_hc_ir_hrcp_half(__half) __asm("llvm.amdgcn.rcp.f16");
extern "C" __half __hip_hc_ir_hrint_half(__half) __asm("llvm.rint.f16");
extern "C" __half __hip_hc_ir_hrsqrt_half(__half) __asm("llvm.sqrt.f16");
extern "C" __half __hip_hc_ir_hsin_half(__half) __asm("llvm.sin.f16");
extern "C" __half __hip_hc_ir_hsqrt_half(__half) __asm("llvm.sqrt.f16");
extern "C" __half __hip_hc_ir_htrunc_half(__half) __asm("llvm.trunc.f16");
__device__ static inline __half __hadd_sat(__half a, __half b) {
return __hip_hc_ir_hadd_half(a, b);
}
__device__ static inline __half __hfma(__half a, __half b, __half c) {
return __hip_hc_ir_hfma_half(a, b, c);
}
__device__ static inline __half __hfma_sat(__half a, __half b, __half c) {
return __hip_hc_ir_hfma_half(a, b, c);
}
__device__ static inline __half __hmul(__half a, __half b) {
return __hip_hc_ir_hmul_half(a, b);
}
__device__ static inline __half __hmul_sat(__half a, __half b) {
return __hip_hc_ir_hmul_half(a, b);
}
__device__ static inline __half __hneg(__half a) {
return -a;
}
__device__ static inline __half __hsub(__half a, __half b) {
return __hip_hc_ir_hsub_half(a, b);
}
__device__ static inline __half __hsub_sat(__half a, __half b) {
return __hip_hc_ir_hsub_half(a, b);
}
__device__ static inline __half hdiv(__half a, __half b) {
return a/b;
}
extern "C" int __hip_hc_ir_h2ceil_int(int);
extern "C" int __hip_hc_ir_h2cos_int(int);
extern "C" int __hip_hc_ir_h2exp2_int(int);
extern "C" int __hip_hc_ir_h2floor_int(int);
extern "C" int __hip_hc_ir_h2log2_int(int);
extern "C" int __hip_hc_ir_h2rcp_int(int);
extern "C" int __hip_hc_ir_h2rsqrt_int(int);
extern "C" int __hip_hc_ir_h2sin_int(int);
extern "C" int __hip_hc_ir_h2sqrt_int(int);
extern "C" int __hip_hc_ir_h2trunc_int(int);
/*
Half2 Arithmetic Functions
@@ -163,70 +308,165 @@ __device__ static inline __half2 h2div(__half2 a, __half2 b) {
return c;
}
__device__ static inline __half hceil(const __half h) {
return __hip_hc_ir_hceil_half(h);
}
__device__ static inline __half hcos(const __half h) {
return __hip_hc_ir_hcos_half(h);
}
__device__ static inline __half hexp(const __half h) {
return __hip_hc_ir_hexp2_half(__hmul(h, 1.442694));
}
__device__ static inline __half hexp10(const __half h) {
return __hip_hc_ir_hexp2_half(__hmul(h, 3.3219281));
}
__device__ static inline __half hexp2(const __half h) {
return __hip_hc_ir_hexp2_half(h);
}
__device__ static inline __half hfloor(const __half h) {
return __hip_hc_ir_hfloor_half(h);
}
__device__ static inline __half hlog(const __half h) {
return __hmul(__hip_hc_ir_hlog2_half(h), 0.693147);
}
__device__ static inline __half hlog10(const __half h) {
return __hmul(__hip_hc_ir_hlog2_half(h), 0.301029);
}
__device__ static inline __half hlog2(const __half h) {
return __hip_hc_ir_hlog2_half(h);
}
/*
Half comparision Functions
__device__ static inline __half hrcp(const __half h) {
return __hip_hc_ir_hrcp_half(h);
}
*/
__device__ static inline bool __heq(__half a, __half b) {
return a == b ? true : false;
__device__ static inline __half hrint(const __half h) {
return __hip_hc_ir_hrint_half(h);
}
__device__ static inline bool __hge(__half a, __half b) {
return a >= b ? true : false;
__device__ static inline __half hrsqrt(const __half h) {
return __hip_hc_ir_hrsqrt_half(h);
}
__device__ static inline bool __hgt(__half a, __half b) {
return a > b ? true : false;
__device__ static inline __half hsin(const __half h) {
return __hip_hc_ir_hsin_half(h);
}
__device__ static inline bool __hisinf(__half a) {
return a == hInf.s ? true : false;
__device__ static inline __half hsqrt(const __half a) {
return __hip_hc_ir_hsqrt_half(a);
}
__device__ static inline bool __hisnan(__half a) {
return a > hInf.s ? true : false;
}
__device__ static inline bool __hle(__half a, __half b) {
return a <= b ? true : false;
}
__device__ static inline bool __hlt(__half a, __half b) {
return a < b ? true : false;
}
__device__ static inline bool __hne(__half a, __half b) {
return a != b ? true : false;
__device__ static inline __half htrunc(const __half a) {
return __hip_hc_ir_htrunc_half(a);
}
/*
Half2 Comparision Functions
Half2 Math Operations
*/
__device__ static inline bool __hbeq2(__half2 a, __half2 b) {
return (a.p[0] == b.p[0] ? true : false) && (a.p[1] == b.p[1] ? true : false);
__device__ static inline __half2 h2ceil(const __half2 h) {
__half2 a;
a.q = __hip_hc_ir_h2ceil_int(h.q);
return a;
}
__device__ static inline bool __hbge2(__half2 a, __half2 b) {
return (a.p[0] >= b.p[0] ? true : false) && (a.p[1] >= b.p[1] ? true : false);
__device__ static inline __half2 h2cos(const __half2 h) {
__half2 a;
a.q = __hip_hc_ir_h2cos_int(h.q);
return a;
}
__device__ static inline bool __hbgt2(__half2 a, __half2 b) {
return (a.p[0] > b.p[0] ? true : false) && (a.p[1] > b.p[1] ? true : false);
__device__ static inline __half2 h2exp(const __half2 h) {
__half2 factor;
factor.p[0] = 1.442694;
factor.p[1] = 1.442694;
factor.q = __hip_hc_ir_h2exp2_int(__hip_hc_ir_hmul2_int(h.q, factor.q));
return factor;
}
__device__ static inline bool __hble2(__half2 a, __half2 b) {
return (a.p[0] <= b.p[0] ? true : false) && (a.p[1] <= b.p[1] ? true : false);
__device__ static inline __half2 h2exp10(const __half2 h) {
__half2 factor;
factor.p[0] = 3.3219281;
factor.p[1] = 3.3219281;
factor.q = __hip_hc_ir_h2exp2_int(__hip_hc_ir_hmul2_int(h.q, factor.q));
return factor;
}
__device__ static inline bool __hblt2(__half2 a, __half2 b) {
return (a.p[0] < b.p[0] ? true : false) && (a.p[1] < b.p[1] ? true : false);
__device__ static inline __half2 h2exp2(const __half2 h) {
__half2 a;
a.q = __hip_hc_ir_h2exp2_int(h.q);
return a;
}
__device__ static inline bool __hbne2(__half2 a, __half2 b) {
return (a.p[0] != b.p[0] ? true : false) && (a.p[1] != b.p[1] ? true : false);
__device__ static inline __half2 h2floor(const __half2 h) {
__half2 a;
a.q = __hip_hc_ir_h2floor_int(h.q);
return a;
}
__device__ static inline __half2 h2log(const __half2 h) {
__half2 factor;
factor.p[0] = 0.693147;
factor.p[1] = 0.693147;
factor. q = __hip_hc_ir_hmul2_int(__hip_hc_ir_h2log2_int(h.q), factor.q);
return factor;
}
__device__ static inline __half2 h2log10(const __half2 h) {
__half2 factor;
factor.p[0] = 0.301029;
factor.p[1] = 0.301029;
factor.q = __hip_hc_ir_hmul2_int(__hip_hc_ir_h2log2_int(h.q), factor.q);
return factor;
}
__device__ static inline __half2 h2log2(const __half2 h) {
__half2 a;
a.q = __hip_hc_ir_h2log2_int(h.q);
return a;
}
__device__ static inline __half2 h2rcp(const __half2 h) {
__half2 a;
a.q = __hip_hc_ir_h2rcp_int(h.q);
return a;
}
__device__ static inline __half2 h2rsqrt(const __half2 h) {
__half2 a;
a.q = __hip_hc_ir_h2rsqrt_int(h.q);
return a;
}
__device__ static inline __half2 h2sin(const __half2 h) {
__half2 a;
a.q = __hip_hc_ir_h2sin_int(h.q);
return a;
}
__device__ static inline __half2 h2sqrt(const __half2 h) {
__half2 a;
a.q = __hip_hc_ir_h2sqrt_int(h.q);
return a;
}
__device__ static inline __half2 h2trunc(const __half2 h) {
__half2 a;
a.q = __hip_hc_ir_h2trunc_int(h.q);
return a;
}
#endif
#if __clang_major__ == 3
@@ -243,8 +483,6 @@ typedef struct __attribute__((aligned(4))){
} __half2;
#endif
projects/hip/include/hip/hcc_detail/hip_runtime.h
+38 -475
Zobrazit soubor
@@ -26,8 +26,8 @@ THE SOFTWARE.
*/
//#pragma once
#ifndef HIP_RUNTIME_H
#define HIP_RUNTIME_H
#ifndef HIP_HCC_DETAIL_RUNTIME_H
#define HIP_HCC_DETAIL_RUNTIME_H
//---
// Top part of file can be compiled with any compiler
@@ -69,6 +69,9 @@ extern int HIP_TRACE_API;
#include <hip/hcc_detail/hip_ldg.h>
#endif
#include <hip/hcc_detail/host_defines.h>
#include <hip/hcc_detail/math_functions.h>
#include <hip/hcc_detail/device_functions.h>
// TODO-HCC remove old definitions ; ~1602 hcc supports __HCC_ACCELERATOR__ define.
#if defined (__KALMAR_ACCELERATOR__) && !defined (__HCC_ACCELERATOR__)
#define __HCC_ACCELERATOR__ __KALMAR_ACCELERATOR__
@@ -121,206 +124,6 @@ extern int HIP_TRACE_API;
#define __HCC_C__
#endif
__device__ float acosf(float x);
__device__ float acoshf(float x);
__device__ float asinf(float x);
__device__ float asinhf(float x);
__device__ float atan2f(float y, float x);
__device__ float atanf(float x);
__device__ float atanhf(float x);
__device__ float cbrtf(float x);
__device__ float ceilf(float x);
__device__ float copysignf(float x, float y);
__device__ float coshf(float x);
__device__ float cyl_bessel_i0f(float x);
__device__ float cyl_bessel_i1f(float x);
__device__ float erfcf(float x);
__device__ float erfcinvf(float y);
__host__ float erfcinvf(float y);
__device__ float erfcxf(float x);
__host__ float erfcxf(float x);
__device__ float erff(float x);
__device__ float erfinvf(float y);
__host__ float erfinvf(float y);
__device__ float exp2f(float x);
__device__ float expm1f(float x);
__device__ float fabsf(float x);
__device__ float fdimf(float x, float y);
__device__ __host__ float fdividef(float x, float y);
__device__ float floorf(float x);
__device__ float fmaf(float x, float y, float z);
__device__ float fmaxf(float x, float y);
__device__ float fminf(float x, float y);
__device__ float fmodf(float x, float y);
__device__ float frexpf(float x, float y);
__device__ float hypotf(float x, float y);
__device__ float ilogbf(float x);
__host__ __device__ unsigned isfinite(float a);
__device__ unsigned isinf(float a);
__device__ unsigned isnan(float a);
__device__ float j0f(float x);
__device__ float j1f(float x);
__device__ float jnf(int n, float x);
__device__ float ldexpf(float x, int exp);
__device__ float lgammaf(float x);
__device__ long long int llrintf(float x);
__device__ long long int llroundf(float x);
__device__ float log1pf(float x);
__device__ float logbf(float x);
__device__ long int lrintf(float x);
__device__ long int lroundf(float x);
__device__ float modff(float x, float *iptr);
__device__ float nanf(const char* tagp);
__device__ float nearbyintf(float x);
__device__ float nextafterf(float x, float y);
__device__ float norm3df(float a, float b, float c);
__host__ float norm3df(float a, float b, float c);
__device__ float norm4df(float a, float b, float c, float d);
__host__ float norm4df(float a, float b, float c, float d);
__device__ float normcdff(float y);
__host__ float normcdff(float y);
__device__ float normcdfinvf(float y);
__host__ float normcdfinvf(float y);
__device__ float normf(int dim, const float *a);
__device__ float rcbrtf(float x);
__host__ float rcbrtf(float x);
__device__ float remainderf(float x, float y);
__device__ float remquof(float x, float y, int *quo);
__device__ float rhypotf(float x, float y);
__host__ float rhypotf(float x, float y);
__device__ float rintf(float x);
__device__ float rnorm3df(float a, float b, float c);
__host__ float rnorm3df(float a, float b, float c);
__device__ float rnorm4df(float a, float b, float c, float d);
__host__ float rnorm4df(float a, float b, float c, float d);
__device__ float rnormf(int dim, const float* a);
__host__ float rnormf(int dim, const float* a);
__device__ float roundf(float x);
__device__ float rsqrtf(float x);
__device__ float scalblnf(float x, long int n);
__device__ float scalbnf(float x, int n);
__host__ __device__ unsigned signbit(float a);
__device__ void sincospif(float x, float *sptr, float *cptr);
__host__ void sincospif(float x, float *sptr, float *cptr);
__device__ float sinhf(float x);
__device__ float sinpif(float x);
__device__ float sqrtf(float x);
__device__ float tanhf(float x);
__device__ float tgammaf(float x);
__device__ float truncf(float x);
__device__ float y0f(float x);
__device__ float y1f(float x);
__device__ float ynf(int n, float x);
__host__ __device__ float cospif(float x);
__host__ __device__ float sinpif(float x);
__device__ float sqrtf(float x);
__host__ __device__ float rsqrtf(float x);
__device__ double acos(double x);
__device__ double acosh(double x);
__device__ double asin(double x);
__device__ double asinh(double x);
__device__ double atan(double x);
__device__ double atan2(double y, double x);
__device__ double atanh(double x);
__device__ double cbrt(double x);
__device__ double ceil(double x);
__device__ double copysign(double x, double y);
__device__ double cos(double x);
__device__ double cosh(double x);
__host__ __device__ double cospi(double x);
__device__ double cyl_bessel_i0(double x);
__device__ double cyl_bessel_i1(double x);
__device__ double erf(double x);
__device__ double erfc(double x);
__device__ double erfcinv(double y);
__device__ double erfcx(double x);
__device__ double erfinv(double x);
__device__ double exp(double x);
__device__ double exp10(double x);
__device__ double exp2(double x);
__device__ double expm1(double x);
__device__ double fabs(double x);
__device__ double fdim(double x, double y);
__device__ double fdivide(double x, double y);
__device__ double floor(double x);
__device__ double fma(double x, double y, double z);
__device__ double fmax(double x, double y);
__device__ double fmin(double x, double y);
__device__ double fmod(double x, double y);
__device__ double frexp(double x, int *nptr);
__device__ double hypot(double x, double y);
__device__ double ilogb(double x);
__host__ __device__ unsigned isfinite(double x);
__device__ unsigned isinf(double x);
__device__ unsigned isnan(double x);
__device__ double j0(double x);
__device__ double j1(double x);
__device__ double jn(int n, double x);
__device__ double ldexp(double x, int exp);
__device__ double lgamma(double x);
__device__ long long llrint(double x);
__device__ long long llround(double x);
__device__ double log(double x);
__device__ double log10(double x);
__device__ double log1p(double x);
__device__ double log2(double x);
__device__ double logb(double x);
__device__ long int lrint(double x);
__device__ long int lround(double x);
__device__ double modf(double x, double *iptr);
__device__ double nan(const char* tagp);
__device__ double nearbyint(double x);
__device__ double nextafter(double x, double y);
__device__ double norm(int dim, const double* t);
__device__ double norm3d(double a, double b, double c);
__host__ double norm3d(double a, double b, double c);
__device__ double norm4d(double a, double b, double c, double d);
__host__ double norm4d(double a, double b, double c, double d);
__device__ double normcdf(double y);
__host__ double normcdf(double y);
__device__ double normcdfinv(double y);
__host__ double normcdfinv(double y);
__device__ double pow(double x, double y);
__device__ double rcbrt(double x);
__host__ double rcbrt(double x);
__device__ double remainder(double x, double y);
__device__ double remquo(double x, double y, int *quo);
__device__ double rhypot(double x, double y);
__host__ double rhypot(double x, double y);
__device__ double rint(double x);
__device__ double rnorm(int dim, const double* t);
__host__ double rnorm(int dim, const double* t);
__device__ double rnorm3d(double a, double b, double c);
__host__ double rnorm3d(double a, double b, double c);
__device__ double rnorm4d(double a, double b, double c, double d);
__host__ double rnorm4d(double a, double b, double c, double d);
__device__ double round(double x);
__host__ __device__ double rsqrt(double x);
__device__ double scalbln(double x, long int n);
__device__ double scalbn(double x, int n);
__host__ __device__ unsigned signbit(double a);
__device__ double sin(double a);
__device__ void sincos(double x, double *sptr, double *cptr);
__device__ void sincospi(double x, double *sptr, double *cptr);
__host__ void sincospi(double x, double *sptr, double *cptr);
__device__ double sinh(double x);
__host__ __device__ double sinpi(double x);
__device__ double sqrt(double x);
__device__ double tan(double x);
__device__ double tanh(double x);
__device__ double tgamma(double x);
__device__ double trunc(double x);
__device__ double y0(double x);
__device__ double y1(double y);
__device__ double yn(int n, double x);
__host__ double erfcinv(double y);
__host__ double erfcx(double x);
__host__ double erfinv(double y);
__host__ double fdivide(double x, double y);
// TODO - hipify-clang - change to use the function call.
//#define warpSize hc::__wavesize()
extern const int warpSize;
@@ -418,26 +221,7 @@ __device__ unsigned int atomicInc(unsigned int* address,
__device__ unsigned int atomicDec(unsigned int* address,
unsigned int val);
//__mul24 __umul24
__device__ int __mul24(int arg1, int arg2);
__device__ unsigned int __umul24(unsigned int arg1, unsigned int arg2);
// integer intrinsic function __poc __clz __ffs __brev
__device__ unsigned int __popc( unsigned int input);
__device__ unsigned int __popcll( unsigned long long int input);
__device__ unsigned int __clz(unsigned int input);
__device__ unsigned int __clzll(unsigned long long int input);
__device__ unsigned int __clz(int input);
__device__ unsigned int __clzll(long long int input);
__device__ unsigned int __ffs(unsigned int input);
__device__ unsigned int __ffsll(unsigned long long int input);
__device__ unsigned int __ffs(int input);
__device__ unsigned int __ffsll(long long int input);
__device__ unsigned int __brev( unsigned int input);
__device__ unsigned long long int __brevll( unsigned long long int input);
// warp vote function __all __any __ballot
// warp vote function __all __any __ballot
__device__ int __all( int input);
__device__ int __any( int input);
__device__ unsigned long long int __ballot( int input);
@@ -468,252 +252,6 @@ __host__ __device__ int max(int arg1, int arg2);
__device__ __attribute__((address_space(3))) void* __get_dynamicgroupbaseptr();
//TODO - add a couple fast math operations here, the set here will grow :
// Single Precision Precise Math
__device__ float __hip_precise_cosf(float);
__device__ float __hip_precise_exp10f(float);
__device__ float __hip_precise_expf(float);
__device__ float __hip_precise_frsqrt_rn(float);
__device__ float __hip_precise_fsqrt_rd(float);
__device__ float __hip_precise_fsqrt_rn(float);
__device__ float __hip_precise_fsqrt_ru(float);
__device__ float __hip_precise_fsqrt_rz(float);
__device__ float __hip_precise_log10f(float);
__device__ float __hip_precise_log2f(float);
__device__ float __hip_precise_logf(float);
__device__ float __hip_precise_powf(float, float);
__device__ void __hip_precise_sincosf(float,float*,float*);
__device__ float __hip_precise_sinf(float);
__device__ float __hip_precise_tanf(float);
// Double Precision Precise Math
__device__ double __hip_precise_dsqrt_rd(double);
__device__ double __hip_precise_dsqrt_rn(double);
__device__ double __hip_precise_dsqrt_ru(double);
__device__ double __hip_precise_dsqrt_rz(double);
// Single Precision Fast Math
extern __attribute__((const)) float __hip_fast_cosf(float) __asm("llvm.cos.f32");
extern __attribute__((const)) float __hip_fast_exp2f(float) __asm("llvm.exp2.f32");
__device__ float __hip_fast_exp10f(float);
__device__ float __hip_fast_expf(float);
__device__ float __hip_fast_frsqrt_rn(float);
extern __attribute__((const)) float __hip_fast_fsqrt_rd(float) __asm("llvm.sqrt.f32");
__device__ float __hip_fast_fsqrt_rn(float);
__device__ float __hip_fast_fsqrt_ru(float);
__device__ float __hip_fast_fsqrt_rz(float);
__device__ float __hip_fast_log10f(float);
extern __attribute__((const)) float __hip_fast_log2f(float) __asm("llvm.log2.f32");
__device__ float __hip_fast_logf(float);
__device__ float __hip_fast_powf(float, float);
__device__ void __hip_fast_sincosf(float,float*,float*);
extern __attribute__((const)) float __hip_fast_sinf(float) __asm("llvm.sin.f32");
__device__ float __hip_fast_tanf(float);
extern __attribute__((const)) float __hip_fast_fmaf(float,float,float) __asm("llvm.fma.f32");
extern __attribute__((const)) float __hip_fast_frcp(float) __asm("llvm.amdgcn.rcp.f32");
extern __attribute__((const)) double __hip_fast_dsqrt(double) __asm("llvm.sqrt.f64");
extern __attribute__((const)) double __hip_fast_fma(double,double,double) __asm("llvm.fma.f64");
extern __attribute__((const)) double __hip_fast_drcp(double) __asm("llvm.amdgcn.rcp.f64");
#ifdef HIP_FAST_MATH
// Single Precision Precise Math when enabled
__device__ inline float cosf(float x) {
return __hip_fast_cosf(x);
}
__device__ inline float exp10f(float x) {
return __hip_fast_exp10f(x);
}
__device__ inline float expf(float x) {
return __hip_fast_expf(x);
}
__device__ inline float log10f(float x) {
return __hip_fast_log10f(x);
}
__device__ inline float log2f(float x) {
return __hip_fast_log2f(x);
}
__device__ inline float logf(float x) {
return __hip_fast_logf(x);
}
__device__ inline float powf(float base, float exponent) {
return __hip_fast_powf(base, exponent);
}
__device__ inline void sincosf(float x, float *s, float *c) {
return __hip_fast_sincosf(x, s, c);
}
__device__ inline float sinf(float x) {
return __hip_fast_sinf(x);
}
__device__ inline float tanf(float x) {
return __hip_fast_tanf(x);
}
#else
__device__ float sinf(float);
__device__ float cosf(float);
__device__ float tanf(float);
__device__ void sincosf(float, float*, float*);
__device__ float logf(float);
__device__ float log2f(float);
__device__ float log10f(float);
__device__ float expf(float);
__device__ float exp10f(float);
__device__ float powf(float, float);
#endif
// Single Precision Fast Math
__device__ inline float __cosf(float x) {
return __hip_fast_cosf(x);
}
__device__ inline float __exp10f(float x) {
return __hip_fast_exp10f(x);
}
__device__ inline float __expf(float x) {
return __hip_fast_expf(x);
}
__device__ inline float __frsqrt_rn(float x) {
return __hip_fast_frsqrt_rn(x);
}
__device__ inline float __fsqrt_rd(float x) {
return __hip_fast_fsqrt_rd(x);
}
__device__ inline float __fsqrt_rn(float x) {
return __hip_fast_fsqrt_rn(x);
}
__device__ inline float __fsqrt_ru(float x) {
return __hip_fast_fsqrt_ru(x);
}
__device__ inline float __fsqrt_rz(float x) {
return __hip_fast_fsqrt_rz(x);
}
__device__ inline float __log10f(float x) {
return __hip_fast_log10f(x);
}
__device__ inline float __log2f(float x) {
return __hip_fast_log2f(x);
}
__device__ inline float __logf(float x) {
return __hip_fast_logf(x);
}
__device__ inline float __powf(float base, float exponent) {
return __hip_fast_powf(base, exponent);
}
__device__ inline void __sincosf(float x, float *s, float *c) {
return __hip_fast_sincosf(x, s, c);
}
__device__ inline float __sinf(float x) {
return __hip_fast_sinf(x);
}
__device__ inline float __tanf(float x) {
return __hip_fast_tanf(x);
}
__device__ inline float __fmaf_rd(float x, float y, float z) {
return __hip_fast_fmaf(x, y, z);
}
__device__ inline float __fmaf_rn(float x, float y, float z) {
return __hip_fast_fmaf(x, y, z);
}
__device__ inline float __fmaf_ru(float x, float y, float z) {
return __hip_fast_fmaf(x, y, z);
}
__device__ inline float __fmaf_rz(float x, float y, float z) {
return __hip_fast_fmaf(x, y, z);
}
__device__ inline float __frcp_rd(float x) {
return __hip_fast_frcp(x);
}
__device__ inline float __frcp_rn(float x) {
return __hip_fast_frcp(x);
}
__device__ inline float __frcp_ru(float x) {
return __hip_fast_frcp(x);
}
__device__ inline float __frcp_rz(float x) {
return __hip_fast_frcp(x);
}
__device__ inline double __dsqrt_rd(double x) {
return __hip_fast_dsqrt(x);
}
__device__ inline double __dsqrt_rn(double x) {
return __hip_fast_dsqrt(x);
}
__device__ inline double __dsqrt_ru(double x) {
return __hip_fast_dsqrt(x);
}
__device__ inline double __dsqrt_rz(double x) {
return __hip_fast_dsqrt(x);
}
__device__ inline double __fma_rd(double x, double y, double z) {
return __hip_fast_fma(x, y, z);
}
__device__ inline double __fma_rn(double x, double y, double z) {
return __hip_fast_fma(x, y, z);
}
__device__ inline double __fma_ru(double x, double y, double z) {
return __hip_fast_fma(x, y, z);
}
__device__ inline double __fma_rz(double x, double y, double z) {
return __hip_fast_fma(x, y, z);
}
__device__ inline double __drcp_rd(double x) {
return __hip_fast_drcp(x);
}
__device__ inline double __drcp_rn(double x) {
return __hip_fast_drcp(x);
}
__device__ inline double __drcp_ru(double x) {
return __hip_fast_drcp(x);
}
__device__ inline double __drcp_rz(double x) {
return __hip_fast_drcp(x);
}
/**
* CUDA 8 device function features
@@ -739,6 +277,10 @@ __device__ inline double __drcp_rz(double x) {
* On AMD platforms, the threadfence* routines are currently empty stubs.
*/
extern __attribute__((const)) __device__ void __hip_hc_threadfence() __asm("__llvm_fence_sc_dev");
extern __attribute__((const)) __device__ void __hip_hc_threadfence_block() __asm("__llvm_fence_sc_wg");
/**
* @brief threadfence_block makes writes visible to threads running in same block.
*
@@ -749,7 +291,9 @@ __device__ inline double __drcp_rz(double x) {
* @warning __threadfence_block is a stub and map to no-op.
*/
// __device__ void __threadfence_block(void);
extern "C" __device__ void __threadfence_block(void);
__device__ static inline void __threadfence_block(void) {
return __hip_hc_threadfence_block();
}
/**
* @brief threadfence makes wirtes visible to other threads running on same GPU.
@@ -761,7 +305,9 @@ extern "C" __device__ void __threadfence_block(void);
* @warning __threadfence is a stub and map to no-op, application should set "export HSA_DISABLE_CACHE=1" to disable both L1 and L2 caches.
*/
// __device__ void __threadfence(void) __attribute__((deprecated("Provided for compile-time compatibility, not yet functional")));
extern "C" __device__ void __threadfence(void);
__device__ static inline void __threadfence(void) {
return __hip_hc_threadfence();
}
/**
* @brief threadfence_system makes writes to pinned system memory visible on host CPU.
@@ -807,14 +353,31 @@ __device__ int __hip_move_dpp(int src, int dpp_ctrl, int row_mask, int bank_mask
#define hipGridDim_y (hc_get_num_groups(1))
#define hipGridDim_z (hc_get_num_groups(2))
//extern "C" __device__ void* memcpy(void* dst, void* src, size_t size);
//extern "C" __device__ void* memset(void* ptr, uint8_t val, size_t size);
extern "C" __device__ void* __hip_hc_memcpy(void* dst, void* src, size_t size);
extern "C" __device__ void* __hip_hc_memset(void* ptr, uint8_t val, size_t size);
extern "C" __device__ void* __hip_hc_malloc(size_t);
extern "C" __device__ void* __hip_hc_free(void *ptr);
//extern "C" __device__ void* malloc(size_t size);
//extern "C" __device__ void* free(void *ptr);
static inline __device__ void* malloc(size_t size)
{
return __hip_hc_malloc(size);
}
static inline __device__ void* free(void *ptr)
{
return __hip_hc_free(ptr);
}
static inline __device__ void* memcpy(void* dst, void* src, size_t size)
{
return __hip_hc_memcpy(dst, src, size);
}
static inline __device__ void* memset(void* ptr, uint8_t val, size_t size)
{
return __hip_hc_memset(ptr, val, size);
}
#define __syncthreads() hc_barrier(CLK_LOCAL_MEM_FENCE)
projects/hip/include/hip/hcc_detail/hip_runtime_api.h
+178 -16
Zobrazit soubor
@@ -1,5 +1,5 @@
/*
Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved.
Copyright (c) 2015-2017 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
@@ -21,8 +21,8 @@ THE SOFTWARE.
*/
//#pragma once
#ifndef HIP_RUNTIME_API_H
#define HIP_RUNTIME_API_H
#ifndef HIP_HCC_DETAIL_HIP_RUNTIME_API_H
#define HIP_HCC_DETAIL_HIP_RUNTIME_API_H
/**
* @file hcc_detail/hip_runtime_api.h
* @brief Contains C function APIs for HIP runtime. This file does not use any HCC builtin or special language extensions (-hc mode) ; those functions in hip_runtime.h.
@@ -30,9 +30,11 @@ THE SOFTWARE.
#include <stdint.h>
#include <stddef.h>
#include <iostream>
#include <hip/hcc_detail/host_defines.h>
#include <hip/hip_runtime_api.h>
#include <hip/hip_texture.h>
#if defined (__HCC__) && (__hcc_workweek__ < 16155)
#error("This version of HIP requires a newer version of HCC.");
@@ -71,14 +73,7 @@ typedef struct ihipIpcEventHandle_t *hipIpcEventHandle_t;
typedef struct ihipModule_t *hipModule_t;
struct ihipModuleSymbol_t{
uint64_t _object; // The kernel object.
uint32_t _groupSegmentSize;
uint32_t _privateSegmentSize;
char _name[64]; // TODO - review for performance cost. Name is just used for debug.
};
typedef struct ihipModuleSymbol_t hipFunction_t;
typedef struct ihipModuleSymbol_t *hipFunction_t;
typedef void* hipDeviceptr_t;
@@ -176,6 +171,12 @@ typedef enum hipMemcpyKind {
,hipMemcpyDefault = 4, ///< Runtime will automatically determine copy-kind based on virtual addresses.
} hipMemcpyKind;
typedef struct {
unsigned int width;
unsigned int height;
enum hipChannelFormatKind f;
void* data; //FIXME: generalize this
} hipArray;
@@ -364,7 +365,7 @@ hipError_t hipDeviceGetCacheConfig ( hipFuncCache_t *cacheConfig );
* Note: Currently, only hipLimitMallocHeapSize is available
*
*/
hipError_t hipDeviceGetLimit(size_t *pValue, hipLimit_t limit);
hipError_t hipDeviceGetLimit(size_t *pValue, enum hipLimit_t limit);
/**
@@ -510,12 +511,10 @@ const char *hipGetErrorString(hipError_t hipError);
* @{
*
* The following Stream APIs are not (yet) supported in HIP:
* - cudaStreamAddCallback
* - cudaStreamAttachMemAsync
* - cudaStreamCreateWithPriority
* - cudaStreamGetPriority
* - cudaStreamWaitEvent
*/
*/
/**
@@ -1160,7 +1159,7 @@ hipError_t hipMemcpyToSymbol(const char* symbolName, const void *src, size_t siz
*/
hipError_t hipMemcpyToSymbolAsync(const char* symbolName, const void *src, size_t sizeBytes, size_t offset, hipMemcpyKind kind, hipStream_t stream);
hipError_t hipMemcpyFromSymbolAsync(void *dst, const char* symbolName, size_t sizeBytes, size_t offset, hipMemcpyKind kind, hipStream_t stream);
/**
* @brief Copy data from src to dst asynchronously.
@@ -1234,6 +1233,88 @@ hipError_t hipMemsetAsync(void* dst, int value, size_t sizeBytes, hipStream_t st
**/
hipError_t hipMemGetInfo (size_t * free, size_t * total) ;
/**
* @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
*/
#if __cplusplus
hipError_t hipMallocArray(hipArray** array, const hipChannelFormatDesc* desc,
size_t width, size_t height = 0, unsigned int flags = 0);
#else
hipError_t hipMallocArray(hipArray** array, const struct hipChannelFormatDesc* desc,
size_t width, size_t height, unsigned int flags);
#endif
/**
* @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);
/**
* @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);
/**
* @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);
// doxygen end Memory
/**
* @}
@@ -1910,6 +1991,87 @@ hipError_t hipIpcCloseMemHandle(void *devPtr);
} /* extern "c" */
#endif
#ifdef __cplusplus
/*
* @brief hipBindTexture Binds size bytes of the memory area pointed to by @p devPtr to the texture reference tex.
*
* @p desc describes how the memory is interpreted when fetching values from the texture. The @p offset parameter is an optional byte offset as with the low-level
* hipBindTexture() function. Any memory previously bound to tex is unbound.
*
* @param[in] offset - Offset in bytes
* @param[out] tex - texture to bind
* @param[in] devPtr - Memory area on device
* @param[in] desc - Channel format
* @param[in] size - Size of the memory area pointed to by devPtr
* @return #hipSuccess, #hipErrorInvalidValue, #hipErrorMemoryFree, #hipErrorUnknown
**/
template <class T, int dim, enum hipTextureReadMode readMode>
hipError_t hipBindTexture(size_t *offset,
struct texture<T, dim, readMode> &tex,
const void *devPtr,
const struct hipChannelFormatDesc *desc,
size_t size=UINT_MAX)
{
tex._dataPtr = static_cast<const T*>(devPtr);
return hipSuccess;
}
/*
* @brief hipBindTexture Binds size bytes of the memory area pointed to by @p devPtr to the texture reference tex.
*
* @p desc describes how the memory is interpreted when fetching values from the texture. The @p offset parameter is an optional byte offset as with the low-level
* hipBindTexture() function. Any memory previously bound to tex is unbound.
*
* @param[in] offset - Offset in bytes
* @param[in] tex - texture to bind
* @param[in] devPtr - Memory area on device
* @param[in] size - Size of the memory area pointed to by devPtr
* @return #hipSuccess, #hipErrorInvalidValue, #hipErrorMemoryFree, #hipErrorUnknown
**/
template <class T, int dim, enum hipTextureReadMode readMode>
hipError_t hipBindTexture(size_t *offset,
struct texture<T, dim, readMode> &tex,
const void *devPtr,
size_t size=UINT_MAX)
{
return hipBindTexture(offset, tex, devPtr, &tex.channelDesc, size);
}
template <class T, int dim, enum hipTextureReadMode readMode>
hipError_t hipBindTextureToArray(struct texture<T, dim, readMode> &tex, hipArray* array) {
tex.width = array->width;
tex.height = array->height;
tex._dataPtr = static_cast<const T*>(array->data);
return hipSuccess;
}
/*
* @brief Unbinds the textuer bound to @p tex
*
* @param[in] tex - texture to unbind
*
* @return #hipSuccess
**/
template <class T, int dim, enum hipTextureReadMode readMode>
hipError_t hipUnbindTexture(struct texture<T, dim, readMode> &tex)
{
tex._dataPtr = NULL;
return hipSuccess;
}
// doxygen end Texture
/**
* @}
*/
#endif
/**
*-------------------------------------------------------------------------------------------------
*-------------------------------------------------------------------------------------------------
projects/hip/include/hip/hcc_detail/hip_texture.h
+5 -234
Zobrazit soubor
@@ -1,5 +1,5 @@
/*
Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved.
Copyright (c) 2015-2017 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
@@ -31,48 +31,17 @@ THE SOFTWARE.
*/
#include <limits.h>
#include <hip/hcc_detail/driver_types.h>
#include <hip/hcc_detail/channel_descriptor.h>
#include <hip/hcc_detail/texture_types.h>
//#include <hip/hcc_detail/hip_runtime.h>
//----
//Texture - TODO - likely need to move this to a separate file only included with kernel compilation.
#define hipTextureType1D 1
typedef enum {
hipChannelFormatKindSigned = 0,
hipChannelFormatKindUnsigned,
hipChannelFormatKindFloat,
hipChannelFormatKindNone
} hipChannelFormatKind;
typedef struct hipChannelFormatDesc {
int x;
int y;
int z;
int w;
hipChannelFormatKind f;
} hipChannelFormatDesc;
typedef enum hipTextureReadMode
{
hipReadModeElementType, ///< Read texture as specified element type
//! @warning cudaReadModeNormalizedFloat is not supported.
} hipTextureReadMode;
typedef enum hipTextureFilterMode
{
hipFilterModePoint, ///< Point filter mode.
//! @warning cudaFilterModeLinear is not supported.
} hipTextureFilterMode;
struct textureReference {
hipTextureFilterMode filterMode;
bool normalized;
hipChannelFormatDesc channelDesc;
};
#if __cplusplus
template <class T, int texType=hipTextureType1D, enum hipTextureReadMode=hipReadModeElementType>
template <class T, int texType=hipTextureType1D, hipTextureReadMode readMode=hipReadModeElementType>
struct texture : public textureReference {
const T * _dataPtr; // pointer to underlying data.
@@ -84,95 +53,12 @@ struct texture : public textureReference {
};
#endif
typedef struct {
unsigned int width;
unsigned int height;
hipChannelFormatKind f;
void* data; //FIXME: generalize this
} hipArray;
#define tex1Dfetch(_tex, _addr) (_tex._dataPtr[_addr])
#define tex2D(_tex, _dx, _dy) \
_tex._dataPtr[(unsigned int)_dx + (unsigned int)_dy*(_tex.width)]
/**
* @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);
/**
* @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);
/**
* @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);
/**
* @addtogroup API HIP API
* @{
@@ -212,120 +98,6 @@ hipChannelFormatDesc hipBindTexture(size_t *offset, struct textureReference *te
}
#endif
/**
* @brief Returns a channel descriptor using the specified format.
*
* @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
template <typename T> inline hipChannelFormatDesc hipCreateChannelDesc() {
return hipCreateChannelDesc(0, 0, 0, 0, hipChannelFormatKindNone);
}
template <> inline hipChannelFormatDesc hipCreateChannelDesc<int>() {
int e = (int)sizeof(int) * 8;
return hipCreateChannelDesc(e, 0, 0, 0, hipChannelFormatKindSigned);
}
template <> inline hipChannelFormatDesc hipCreateChannelDesc<unsigned int>() {
int e = (int)sizeof(unsigned int) * 8;
return hipCreateChannelDesc(e, 0, 0, 0, hipChannelFormatKindUnsigned);
}
template <> inline hipChannelFormatDesc hipCreateChannelDesc<long>() {
int e = (int)sizeof(long) * 8;
return hipCreateChannelDesc(e, 0, 0, 0, hipChannelFormatKindSigned);
}
template <> inline hipChannelFormatDesc hipCreateChannelDesc<unsigned long>() {
int e = (int)sizeof(unsigned long) * 8;
return hipCreateChannelDesc(e, 0, 0, 0, hipChannelFormatKindUnsigned);
}
template <> inline hipChannelFormatDesc hipCreateChannelDesc<float>() {
int e = (int)sizeof(float) * 8;
return hipCreateChannelDesc(e, 0, 0, 0, hipChannelFormatKindFloat);
}
/*
* @brief hipBindTexture Binds size bytes of the memory area pointed to by @p devPtr to the texture reference tex.
*
* @p desc describes how the memory is interpreted when fetching values from the texture. The @p offset parameter is an optional byte offset as with the low-level
* hipBindTexture() function. Any memory previously bound to tex is unbound.
*
* @param[in] offset - Offset in bytes
* @param[out] tex - texture to bind
* @param[in] devPtr - Memory area on device
* @param[in] desc - Channel format
* @param[in] size - Size of the memory area pointed to by devPtr
* @return #hipSuccess, #hipErrorInvalidValue, #hipErrorMemoryFree, #hipErrorUnknown
**/
template <class T, int dim, enum hipTextureReadMode readMode>
hipError_t hipBindTexture(size_t *offset,
struct texture<T, dim, readMode> &tex,
const void *devPtr,
const struct hipChannelFormatDesc *desc,
size_t size=UINT_MAX)
{
tex._dataPtr = static_cast<const T*>(devPtr);
return hipSuccess;
}
/*
* @brief hipBindTexture Binds size bytes of the memory area pointed to by @p devPtr to the texture reference tex.
*
* @p desc describes how the memory is interpreted when fetching values from the texture. The @p offset parameter is an optional byte offset as with the low-level
* hipBindTexture() function. Any memory previously bound to tex is unbound.
*
* @param[in] offset - Offset in bytes
* @param[in] tex - texture to bind
* @param[in] devPtr - Memory area on device
* @param[in] size - Size of the memory area pointed to by devPtr
* @return #hipSuccess, #hipErrorInvalidValue, #hipErrorMemoryFree, #hipErrorUnknown
**/
template <class T, int dim, enum hipTextureReadMode readMode>
hipError_t hipBindTexture(size_t *offset,
struct texture<T, dim, readMode> &tex,
const void *devPtr,
size_t size=UINT_MAX)
{
return hipBindTexture(offset, tex, devPtr, &tex.channelDesc, size);
}
template <class T, int dim, enum hipTextureReadMode readMode>
hipError_t hipBindTextureToArray(struct texture<T, dim, readMode> &tex, hipArray* array) {
tex.width = array->width;
tex.height = array->height;
tex._dataPtr = static_cast<const T*>(array->data);
return hipSuccess;
}
/*
* @brief Unbinds the textuer bound to @p tex
*
* @param[in] tex - texture to unbind
*
* @return #hipSuccess
**/
template <class T, int dim, enum hipTextureReadMode readMode>
hipError_t hipUnbindTexture(struct texture<T, dim, readMode> &tex)
{
tex._dataPtr = NULL;
return hipSuccess;
}
// doxygen end Texture
/**
* @}
*/
// End doxygen API:
/**
@@ -333,4 +105,3 @@ hipError_t hipUnbindTexture(struct texture<T, dim, readMode> &tex)
*/
#endif
projects/hip/include/hip/hcc_detail/hip_vector_types.h
+28 -28
Zobrazit soubor
@@ -1131,23 +1131,23 @@ struct longlong4 {
} __attribute__((aligned(32)));
#define DECLOP_MAKE_ONE_COMPONENT(comp, type) \
__device__ __host__ static inline type make_##type(comp x) { \
type ret; \
__device__ __host__ static inline struct type make_##type(comp x) { \
struct type ret; \
ret.x = x; \
return ret; \
}
#define DECLOP_MAKE_TWO_COMPONENT(comp, type) \
__device__ __host__ static inline type make_##type(comp x, comp y) { \
type ret; \
__device__ __host__ static inline struct type make_##type(comp x, comp y) { \
struct type ret; \
ret.x = x; \
ret.y = y; \
return ret; \
}
#define DECLOP_MAKE_THREE_COMPONENT(comp, type) \
__device__ __host__ static inline type make_##type(comp x, comp y, comp z) { \
type ret; \
__device__ __host__ static inline struct type make_##type(comp x, comp y, comp z) { \
struct type ret; \
ret.x = x; \
ret.y = y; \
ret.z = z; \
@@ -1155,8 +1155,8 @@ __device__ __host__ static inline type make_##type(comp x, comp y, comp z) { \
}
#define DECLOP_MAKE_FOUR_COMPONENT(comp, type) \
__device__ __host__ static inline type make_##type(comp x, comp y, comp z, comp w) { \
type ret; \
__device__ __host__ static inline struct type make_##type(comp x, comp y, comp z, comp w) { \
struct type ret; \
ret.x = x; \
ret.y = y; \
ret.z = z; \
@@ -1228,20 +1228,20 @@ DECLOP_MAKE_FOUR_COMPONENT(signed long, longlong4);
#if __cplusplus
#define DECLOP_1VAR_2IN_1OUT(type, op) \
__device__ __host__ static type operator op (const type& lhs, const type& rhs) { \
__device__ __host__ static inline type operator op (const type& lhs, const type& rhs) { \
type ret; \
ret.x = lhs.x op rhs.x; \
return ret; \
}
#define DECLOP_1VAR_SCALE_PRODUCT(type, type1) \
__device__ __host__ static type operator * (const type& lhs, type1 rhs) { \
__device__ __host__ static inline type operator * (const type& lhs, type1 rhs) { \
type ret; \
ret.x = lhs.x * rhs; \
return ret; \
} \
\
__device__ __host__ static type operator * (type1 lhs, const type& rhs) { \
__device__ __host__ static inline type operator * (type1 lhs, const type& rhs) { \
type ret; \
ret.x = lhs * rhs.x; \
return ret; \
@@ -1260,7 +1260,7 @@ __device__ __host__ static inline type& operator op (type& val) { \
}
#define DECLOP_1VAR_POSTOP(type, op) \
__device__ __host__ static type operator op (type& val, int i) { \
__device__ __host__ static inline type operator op (type& val, int i) { \
type ret; \
ret.x = val.x; \
val.x op; \
@@ -1273,7 +1273,7 @@ __device__ __host__ static inline bool operator op (type& lhs, type& rhs) { \
}
#define DECLOP_1VAR_1IN_1OUT(type, op) \
__device__ __host__ static type operator op(type& rhs) { \
__device__ __host__ static inline type operator op(type& rhs) { \
type ret; \
ret.x = op rhs.x; \
return ret; \
@@ -1289,7 +1289,7 @@ __device__ __host__ static inline bool operator op (type& rhs) { \
*/
#define DECLOP_2VAR_2IN_1OUT(type, op) \
__device__ __host__ static type operator op (const type& lhs, const type& rhs) { \
__device__ __host__ static inline type operator op (const type& lhs, const type& rhs) { \
type ret; \
ret.x = lhs.x op rhs.x; \
ret.y = lhs.y op rhs.y; \
@@ -1297,14 +1297,14 @@ __device__ __host__ static type operator op (const type& lhs, const type& rhs) {
}
#define DECLOP_2VAR_SCALE_PRODUCT(type, type1) \
__device__ __host__ static type operator * (const type& lhs, type1 rhs) { \
__device__ __host__ static inline type operator * (const type& lhs, type1 rhs) { \
type ret; \
ret.x = lhs.x * rhs; \
ret.y = lhs.y * rhs; \
return ret; \
} \
\
__device__ __host__ static type operator * (type1 lhs, const type& rhs) { \
__device__ __host__ static inline type operator * (type1 lhs, const type& rhs) { \
type ret; \
ret.x = lhs * rhs.x; \
ret.y = lhs * rhs.y; \
@@ -1326,7 +1326,7 @@ __device__ __host__ static inline type& operator op (type& val) { \
}
#define DECLOP_2VAR_POSTOP(type, op) \
__device__ __host__ static type operator op (type& val, int i) { \
__device__ __host__ static inline type operator op (type& val, int i) { \
type ret; \
ret.x = val.x; \
ret.y = val.y; \
@@ -1341,7 +1341,7 @@ __device__ __host__ static inline bool operator op (type& lhs, type& rhs) { \
}
#define DECLOP_2VAR_1IN_1OUT(type, op) \
__device__ __host__ static type operator op(type &rhs) { \
__device__ __host__ static inline type operator op(type &rhs) { \
type ret; \
ret.x = op rhs.x; \
ret.y = op rhs.y; \
@@ -1359,7 +1359,7 @@ __device__ __host__ static inline bool operator op (type &rhs) { \
*/
#define DECLOP_3VAR_2IN_1OUT(type, op) \
__device__ __host__ static type operator op (const type& lhs, const type& rhs) { \
__device__ __host__ static inline type operator op (const type& lhs, const type& rhs) { \
type ret; \
ret.x = lhs.x op rhs.x; \
ret.y = lhs.y op rhs.y; \
@@ -1368,7 +1368,7 @@ __device__ __host__ static type operator op (const type& lhs, const type& rhs) {
}
#define DECLOP_3VAR_SCALE_PRODUCT(type, type1) \
__device__ __host__ static type operator * (const type& lhs, type1 rhs) { \
__device__ __host__ static inline type operator * (const type& lhs, type1 rhs) { \
type ret; \
ret.x = lhs.x * rhs; \
ret.y = lhs.y * rhs; \
@@ -1376,7 +1376,7 @@ __device__ __host__ static type operator * (const type& lhs, type1 rhs) { \
return ret; \
} \
\
__device__ __host__ static type operator * (type1 lhs, const type& rhs) { \
__device__ __host__ static inline type operator * (type1 lhs, const type& rhs) { \
type ret; \
ret.x = lhs * rhs.x; \
ret.y = lhs * rhs.y; \
@@ -1401,7 +1401,7 @@ __device__ __host__ static inline type& operator op (type& val) { \
}
#define DECLOP_3VAR_POSTOP(type, op) \
__device__ __host__ static type operator op (type& val, int i) { \
__device__ __host__ static inline type operator op (type& val, int i) { \
type ret; \
ret.x = val.x; \
ret.y = val.y; \
@@ -1418,7 +1418,7 @@ __device__ __host__ static inline bool operator op (type& lhs, type& rhs) { \
}
#define DECLOP_3VAR_1IN_1OUT(type, op) \
__device__ __host__ static type operator op(type &rhs) { \
__device__ __host__ static inline type operator op(type &rhs) { \
type ret; \
ret.x = op rhs.x; \
ret.y = op rhs.y; \
@@ -1437,7 +1437,7 @@ __device__ __host__ static inline bool operator op (type &rhs) { \
*/
#define DECLOP_4VAR_2IN_1OUT(type, op) \
__device__ __host__ static type operator op ( const type& lhs, const type& rhs) { \
__device__ __host__ static inline type operator op ( const type& lhs, const type& rhs) { \
type ret; \
ret.x = lhs.x op rhs.x; \
ret.y = lhs.y op rhs.y; \
@@ -1447,7 +1447,7 @@ __device__ __host__ static type operator op ( const type& lhs, const type& rhs)
}
#define DECLOP_4VAR_SCALE_PRODUCT(type, type1) \
__device__ __host__ static type operator * (const type& lhs, type1 rhs) { \
__device__ __host__ static inline type operator * (const type& lhs, type1 rhs) { \
type ret; \
ret.x = lhs.x * rhs; \
ret.y = lhs.y * rhs; \
@@ -1456,7 +1456,7 @@ __device__ __host__ static type operator * (const type& lhs, type1 rhs) { \
return ret; \
} \
\
__device__ __host__ static type operator * (type1 lhs, const type& rhs) { \
__device__ __host__ static inline type operator * (type1 lhs, const type& rhs) { \
type ret; \
ret.x = lhs * rhs.x; \
ret.y = lhs * rhs.y; \
@@ -1484,7 +1484,7 @@ __device__ __host__ static inline type& operator op (type& val) { \
}
#define DECLOP_4VAR_POSTOP(type, op) \
__device__ __host__ static type operator op (type& val, int i) { \
__device__ __host__ static inline type operator op (type& val, int i) { \
type ret; \
ret.x = val.x; \
ret.y = val.y; \
@@ -1503,7 +1503,7 @@ __device__ __host__ static inline bool operator op (type& lhs, type& rhs) { \
}
#define DECLOP_4VAR_1IN_1OUT(type, op) \
__device__ __host__ static type operator op(type &rhs) { \
__device__ __host__ static inline type operator op(type &rhs) { \
type ret; \
ret.x = op rhs.x; \
ret.y = op rhs.y; \
projects/hip/include/hip/hcc_detail/host_defines.h
+1 -1
Zobrazit soubor
@@ -35,7 +35,7 @@ THE SOFTWARE.
#define __host__ __attribute__((cpu))
#define __device__ __attribute__((hc))
#define __global__ __attribute__((hc_grid_launch))
#define __global__ __attribute__((hc_grid_launch)) __attribute__((used))
#define __noinline__ __attribute__((noinline))
#define __forceinline__ __attribute__((always_inline))
projects/hip/include/hip/hcc_detail/math_functions.h
+263
Zobrazit soubor
@@ -0,0 +1,263 @@
/*
Copyright (c) 2015-2017 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.
*/
#ifndef HIP_HCC_DETAIL_MATH_FUNCTIONS_H
#define HIP_HCC_DETAIL_MATH_FUNCTIONS_H
#include <hip/hip_runtime.h>
#include <hip/hip_vector_types.h>
#include <hip/hcc_detail/device_functions.h>
__device__ float acosf(float x);
__device__ float acoshf(float x);
__device__ float asinf(float x);
__device__ float asinhf(float x);
__device__ float atan2f(float y, float x);
__device__ float atanf(float x);
__device__ float atanhf(float x);
__device__ float cbrtf(float x);
__device__ float ceilf(float x);
__device__ float copysignf(float x, float y);
__device__ float cosf(float x);
__device__ float coshf(float x);
__device__ float cospif(float x);
//__device__ float cyl_bessel_i0f(float x);
//__device__ float cyl_bessel_i1f(float x);
__device__ float erfcf(float x);
__device__ float erfcinvf(float y);
__device__ float erfcxf(float x);
__device__ float erff(float x);
__device__ float erfinvf(float y);
__device__ float exp10f(float x);
__device__ float exp2f(float x);
__device__ float expf(float x);
__device__ float expm1f(float x);
__device__ float fabsf(float x);
__device__ float fdimf(float x, float y);
__device__ float fdividef(float x, float y);
__device__ float floorf(float x);
__device__ float fmaf(float x, float y, float z);
__device__ float fmaxf(float x, float y);
__device__ float fminf(float x, float y);
__device__ float fmodf(float x, float y);
//__device__ float frexpf(float x, int* nptr);
__device__ float hypotf(float x, float y);
__device__ float ilogbf(float x);
__device__ int isfinite(float a);
__device__ unsigned isinf(float a);
__device__ unsigned isnan(float a);
__device__ float j0f(float x);
__device__ float j1f(float x);
__device__ float jnf(int n, float x);
__device__ float ldexpf(float x, int exp);
//__device__ float lgammaf(float x);
__device__ long long int llrintf(float x);
__device__ long long int llroundf(float x);
__device__ float log10f(float x);
__device__ float log1pf(float x);
__device__ float logbf(float x);
__device__ long int lrintf(float x);
__device__ long int lroundf(float x);
//__device__ float modff(float x, float *iptr);
__device__ float nanf(const char* tagp);
__device__ float nearbyintf(float x);
//__device__ float nextafterf(float x, float y);
__device__ float norm3df(float a, float b, float c);
__device__ float norm4df(float a, float b, float c, float d);
__device__ float normcdff(float y);
__device__ float normcdfinvf(float y);
__device__ float normf(int dim, const float *a);
__device__ float powf(float x, float y);
__device__ float rcbrtf(float x);
__device__ float remainderf(float x, float y);
__device__ float remquof(float x, float y, int *quo);
__device__ float rhypotf(float x, float y);
__device__ float rintf(float x);
__device__ float rnorm3df(float a, float b, float c);
__device__ float rnorm4df(float a, float b, float c, float d);
__device__ float rnormf(int dim, const float* a);
__device__ float roundf(float x);
__device__ float rsqrtf(float x);
__device__ float scalblnf(float x, long int n);
__device__ float scalbnf(float x, int n);
__device__ int signbit(float a);
__device__ void sincosf(float x, float *sptr, float *cptr);
__device__ void sincospif(float x, float *sptr, float *cptr);
__device__ float sinf(float x);
__device__ float sinhf(float x);
__device__ float sinpif(float x);
__device__ float sqrtf(float x);
__device__ float tanf(float x);
__device__ float tanhf(float x);
__device__ float tgammaf(float x);
__device__ float truncf(float x);
__device__ float y0f(float x);
__device__ float y1f(float x);
__device__ float ynf(int n, float x);
__device__ double acos(double x);
__device__ double acosh(double x);
__device__ double asin(double x);
__device__ double asinh(double x);
__device__ double atan(double x);
__device__ double atan2(double y, double x);
__device__ double atanh(double x);
__device__ double cbrt(double x);
__device__ double ceil(double x);
__device__ double copysign(double x, double y);
__device__ double cos(double x);
__device__ double cosh(double x);
__device__ double cospi(double x);
//__device__ double cyl_bessel_i0(double x);
//__device__ double cyl_bessel_i1(double x);
__device__ double erf(double x);
__device__ double erfc(double x);
__device__ double erfcinv(double y);
__device__ double erfcx(double x);
__device__ double erfinv(double x);
__device__ double exp(double x);
__device__ double exp10(double x);
__device__ double exp2(double x);
__device__ double expm1(double x);
__device__ double fabs(double x);
__device__ double fdim(double x, double y);
__device__ double floor(double x);
__device__ double fma(double x, double y, double z);
__device__ double fmax(double x, double y);
__device__ double fmin(double x, double y);
__device__ double fmod(double x, double y);
//__device__ double frexp(double x, int *nptr);
__device__ double hypot(double x, double y);
__device__ double ilogb(double x);
__device__ int isfinite(double x);
__device__ unsigned isinf(double x);
__device__ unsigned isnan(double x);
__device__ double j0(double x);
__device__ double j1(double x);
__device__ double jn(int n, double x);
__device__ double ldexp(double x, int exp);
__device__ double lgamma(double x);
__device__ long long llrint(double x);
__device__ long long llround(double x);
__device__ double log(double x);
__device__ double log10(double x);
__device__ double log1p(double x);
__device__ double log2(double x);
__device__ double logb(double x);
__device__ long int lrint(double x);
__device__ long int lround(double x);
//__device__ double modf(double x, double *iptr);
__device__ double nan(const char* tagp);
__device__ double nearbyint(double x);
__device__ double nextafter(double x, double y);
__device__ double norm(int dim, const double* t);
__device__ double norm3d(double a, double b, double c);
__device__ double norm4d(double a, double b, double c, double d);
__device__ double normcdf(double y);
__device__ double normcdfinv(double y);
__device__ double pow(double x, double y);
__device__ double rcbrt(double x);
__device__ double remainder(double x, double y);
//__device__ double remquo(double x, double y, int *quo);
__device__ double rhypot(double x, double y);
__device__ double rint(double x);
__device__ double rnorm(int dim, const double* t);
__device__ double rnorm3d(double a, double b, double c);
__device__ double rnorm4d(double a, double b, double c, double d);
__device__ double round(double x);
__device__ double rsqrt(double x);
__device__ double scalbln(double x, long int n);
__device__ double scalbn(double x, int n);
__device__ int signbit(double a);
__device__ double sin(double a);
__device__ void sincos(double x, double *sptr, double *cptr);
__device__ void sincospi(double x, double *sptr, double *cptr);
__device__ double sinh(double x);
__device__ double sinpi(double x);
__device__ double sqrt(double x);
__device__ double tan(double x);
__device__ double tanh(double x);
__device__ double tgamma(double x);
__device__ double trunc(double x);
__device__ double y0(double x);
__device__ double y1(double y);
__device__ double yn(int n, double x);
// ENDPARSER
#ifdef HIP_FAST_MATH
// Single Precision Precise Math when enabled
__device__ inline float cosf(float x) {
return __hip_fast_cosf(x);
}
__device__ inline float exp10f(float x) {
return __hip_fast_exp10f(x);
}
__device__ inline float expf(float x) {
return __hip_fast_expf(x);
}
__device__ inline float log10f(float x) {
return __hip_fast_log10f(x);
}
__device__ inline float log2f(float x) {
return __hip_fast_log2f(x);
}
__device__ inline float logf(float x) {
return __hip_fast_logf(x);
}
__device__ inline float powf(float base, float exponent) {
return __hip_fast_powf(base, exponent);
}
__device__ inline void sincosf(float x, float *s, float *c) {
return __hip_fast_sincosf(x, s, c);
}
__device__ inline float sinf(float x) {
return __hip_fast_sinf(x);
}
__device__ inline float tanf(float x) {
return __hip_fast_tanf(x);
}
#else
__device__ float sinf(float);
__device__ float cosf(float);
__device__ float tanf(float);
__device__ void sincosf(float, float*, float*);
__device__ float logf(float);
__device__ float log2f(float);
__device__ float log10f(float);
__device__ float expf(float);
__device__ float exp10f(float);
__device__ float powf(float, float);
#endif
#endif
projects/hip/include/hip/hcc_detail/texture_types.h
+42
Zobrazit soubor
@@ -0,0 +1,42 @@
/*
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.
*/
#ifndef HIP_HCC_DETAIL_TEXTURE_TYPES_H
#define HIP_HCC_DETAIL_TEXTURE_TYPES_H
#include<hip/hcc_detail/driver_types.h>
enum hipTextureReadMode
{
hipReadModeElementType = 0
};
enum hipTextureFilterMode
{
hipFilterModePoint = 0
};
struct textureReference {
enum hipTextureFilterMode filterMode;
unsigned normalized;
struct hipChannelFormatDesc channelDesc;
};
#endif
projects/hip/include/hip/math_functions.h
+36
Zobrazit soubor
@@ -0,0 +1,36 @@
/*
Copyright (c) 2015-2017 Advanced Micro Devices, Inc. All rights reserved.
Permission is hereby granted, free of charge, to any person obtaining a copy
of this software and associated documentation files (the "Software"), to deal
in the Software without restriction, including without limitation the rights
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
copies of the Software, and to permit persons to whom the Software is
furnished to do so, subject to the following conditions:
The above copyright notice and this permission notice shall be included in
all copies or substantial portions of the Software.
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
#pragma once
// Some standard header files, these are included by hc.hpp and so want to make them avail on both
// paths to provide a consistent include env and avoid "missing symbol" errors that only appears
// on NVCC path:
#if defined(__HIP_PLATFORM_HCC__) && !defined (__HIP_PLATFORM_NVCC__)
#include <hip/hcc_detail/math_functions.h>
#elif defined(__HIP_PLATFORM_NVCC__) && !defined (__HIP_PLATFORM_HCC__)
#include <hip/nvcc_detail/math_functions.h>
#else
#error("Must define exactly one of __HIP_PLATFORM_HCC__ or __HIP_PLATFORM_NVCC__");
#endif
projects/hip/include/hip/nvcc_detail/channel_descriptor.h
+25
Zobrazit soubor
@@ -0,0 +1,25 @@
/*
Copyright (c) 2015-2017 Advanced Micro Devices, Inc. All rights reserved.
Permission is hereby granted, free of charge, to any person obtaining a copy
of this software and associated documentation files (the "Software"), to deal
in the Software without restriction, including without limitation the rights
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
copies of the Software, and to permit persons to whom the Software is
furnished to do so, subject to the following conditions:
The above copyright notice and this permission notice shall be included in
all copies or substantial portions of the Software.
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
#pragma once
#include"channel_descriptor.h"
projects/hip/include/hip/nvcc_detail/hip_runtime_api.h
+25 -1
Zobrazit soubor
@@ -1,5 +1,5 @@
/*
Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved.
Copyright (c) 2015-2017 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
@@ -85,6 +85,10 @@ typedef CUdevice hipDevice_t;
typedef CUmodule hipModule_t;
typedef CUfunction hipFunction_t;
typedef CUdeviceptr hipDeviceptr_t;
typedef cudaChannelFormatKind hipChannelFormatKind;
typedef cudaChannelFormatDesc hipChannelFormatDesc;
typedef cudaTextureReadMode hipTextureReadMode;
typedef cudaArray hipArray;
// Flags that can be used with hipStreamCreateWithFlags
#define hipStreamDefault cudaStreamDefault
@@ -215,6 +219,14 @@ inline static hipError_t hipHostMalloc(void** ptr, size_t size, unsigned int fla
return hipCUDAErrorTohipError(cudaHostAlloc(ptr, size, flags));
}
inline static hipError_t hipMallocArray(hipArray** array, const hipChannelFormatDesc* desc, size_t width, size_t height, unsigned int flags) {
return hipCUDAErrorTohipError(cudaMallocArray(array, desc, width, height, flags));
}
inline static hipError_t hipFreeArray(hipArray* array) {
return hipCUDAErrorTohipError(cudaFreeArray(array));
}
inline static hipError_t hipHostGetDevicePointer(void** devPtr, void* hostPtr, unsigned int flags){
return hipCUDAErrorTohipError(cudaHostGetDevicePointer(devPtr, hostPtr, flags));
}
@@ -321,6 +333,18 @@ inline static hipError_t hipMemcpyToSymbolAsync(const void* symbol, const void*
return hipCUDAErrorTohipError(cudaMemcpyToSymbolAsync(symbol, src, sizeBytes, offset, hipMemcpyKindToCudaMemcpyKind(copyType)));
}
inline static hipError_t hipMemcpy2D(void* dst, size_t dpitch, const void* src, size_t spitch, size_t width, size_t height, hipMemcpyKind kind){
return hipCUDAErrorTohipError(cudaMemcpy2D(dst, dpitch, src, spitch, width, height, hipMemcpyKindToCudaMemcpyKind(kind)));
}
inline static 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){
return hipCUDAErrorTohipError(cudaMemcpy2DToArray(dst, wOffset, hOffset, src, spitch, width, height, hipMemcpyKindToCudaMemcpyKind(kind)));
}
inline static hipError_t hipMemcpyToArray(hipArray* dst, size_t wOffset, size_t hOffset, const void* src, size_t count, hipMemcpyKind kind) {
return hipCUDAErrorTohipError(cudaMemcpyToArray(dst, wOffset, hOffset, src, count, hipMemcpyKindToCudaMemcpyKind(kind)));
}
inline static hipError_t hipDeviceSynchronize() {
return hipCUDAErrorTohipError(cudaDeviceSynchronize());
}
projects/hip/include/hip/texture_types.h
+30
Zobrazit soubor
@@ -0,0 +1,30 @@
/*
Copyright (c) 2015-2017 Advanced Micro Devices, Inc. All rights reserved.
Permission is hereby granted, free of charge, to any person obtaining a copy
of this software and associated documentation files (the "Software"), to deal
in the Software without restriction, including without limitation the rights
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
copies of the Software, and to permit persons to whom the Software is
furnished to do so, subject to the following conditions:
The above copyright notice and this permission notice shall be included in
all copies or substantial portions of the Software.
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
#pragma once
#include <hip/hip_common.h>
#if defined(__HIP_PLATFORM_HCC__) && !defined (__HIP_PLATFORM_NVCC__)
#include <hip/hcc_detail/texture_types.h>
#elif defined(__HIP_PLATFORM_NVCC__) && !defined (__HIP_PLATFORM_HCC__)
#include "texture_types.h"
#else
#error("Must define exactly one of __HIP_PLATFORM_HCC__ or __HIP_PLATFORM_NVCC__");
#endif
projects/hip/packaging/hip_hcc.txt
+1 -1
Zobrazit soubor
@@ -5,7 +5,7 @@ install(FILES @PROJECT_BINARY_DIR@/libhip_hcc.so DESTINATION lib)
install(FILES @PROJECT_BINARY_DIR@/libhip_hcc_static.a DESTINATION lib)
install(FILES @PROJECT_BINARY_DIR@/libhip_device.a DESTINATION lib)
install(FILES @PROJECT_BINARY_DIR@/.hipInfo DESTINATION lib)
install(FILES @hip_SOURCE_DIR@/src/hip_ir.ll DESTINATION lib)
install(FILES @hip_SOURCE_DIR@/src/hip_hc.ll @hip_SOURCE_DIR@/src/hip_hc_gfx803.ll DESTINATION lib)
#############################
# Packaging steps
projects/hip/samples/0_Intro/square/square.hipref.cpp
+1 -1
Zobrazit soubor
@@ -81,7 +81,7 @@ int main(int argc, char *argv[])
const unsigned threadsPerBlock = 256;
printf ("info: launch 'vector_square' kernel\n");
hipLaunchKernel(HIP_KERNEL_NAME(vector_square), dim3(blocks), dim3(threadsPerBlock), 0, 0, C_d, A_d, N);
hipLaunchKernel(vector_square, dim3(blocks), dim3(threadsPerBlock), 0, 0, C_d, A_d, N);
printf ("info: copy Device2Host\n");
CHECK ( hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost));
projects/hip/samples/1_Utils/hipCommander/hipCommander.cpp
+1 -1
Zobrazit soubor
@@ -23,7 +23,7 @@ bool g_printedTiming = false;
// Cmdline parms:
int p_device = 0;
const char* p_command = "H2D; NullKernel; D2H";
const char* p_command = "setstream(1); H2D; NullKernel; D2H;";
const char* p_file = nullptr;
unsigned p_verbose = 0x0;
unsigned p_db = 0x0;
projects/hip/samples/1_Utils/hipCommander/perf/scripts/latency_2_d2h_h2d.hcm
+9
Zobrazit soubor
@@ -0,0 +1,9 @@
setstream(1);
loop(10); D2H; H2D; streamsync;D2H; H2D; streamsync; endloop(1);
loop(10); D2H; H2D; streamsync;D2H; H2D; streamsync; endloop(1);
loop(100); D2H; H2D; streamsync;D2H; H2D; streamsync; endloop(1);
loop(100); D2H; H2D; streamsync;D2H; H2D; streamsync; endloop(1);
loop(1000); D2H; H2D; streamsync;D2H; H2D; streamsync; endloop(1);
loop(1000); D2H; H2D; streamsync;D2H; H2D; streamsync; endloop(1);
loop(10000); D2H; H2D; streamsync;D2H; H2D; streamsync; endloop(1);
loop(10000); D2H; H2D; streamsync;D2H; H2D; streamsync; endloop(1);
projects/hip/samples/1_Utils/hipCommander/perf/scripts/latency_2_d2h_kernel.hcm
+9
Zobrazit soubor
@@ -0,0 +1,9 @@
setstream(1);
loop(10); D2H; NullKernel; streamsync;D2H; NullKernel; streamsync; endloop(1);
loop(10); D2H; NullKernel; streamsync;D2H; NullKernel; streamsync; endloop(1);
loop(100); D2H; NullKernel; streamsync;D2H; NullKernel; streamsync; endloop(1);
loop(100); D2H; NullKernel; streamsync;D2H; NullKernel; streamsync; endloop(1);
loop(1000); D2H; NullKernel; streamsync;D2H; NullKernel; streamsync; endloop(1);
loop(1000); D2H; NullKernel; streamsync;D2H; NullKernel; streamsync; endloop(1);
loop(10000); D2H; NullKernel; streamsync;D2H; NullKernel; streamsync; endloop(1);
loop(10000); D2H; NullKernel; streamsync;D2H; NullKernel; streamsync; endloop(1);
projects/hip/samples/1_Utils/hipCommander/perf/scripts/latency_2_d2h_sync_h2d.hcm
+9
Zobrazit soubor
@@ -0,0 +1,9 @@
setstream(1);
loop(10); D2H; streamsync; H2D; streamsync;D2H; streamsync; H2D; streamsync; endloop(1);
loop(10); D2H; streamsync; H2D; streamsync;D2H; streamsync; H2D; streamsync; endloop(1);
loop(100); D2H; streamsync; H2D; streamsync;D2H; streamsync; H2D; streamsync; endloop(1);
loop(100); D2H; streamsync; H2D; streamsync;D2H; streamsync; H2D; streamsync; endloop(1);
loop(1000); D2H; streamsync; H2D; streamsync;D2H; streamsync; H2D; streamsync; endloop(1);
loop(1000); D2H; streamsync; H2D; streamsync;D2H; streamsync; H2D; streamsync; endloop(1);
loop(10000); D2H; streamsync; H2D; streamsync;D2H; streamsync; H2D; streamsync; endloop(1);
loop(10000); D2H; streamsync; H2D; streamsync;D2H; streamsync; H2D; streamsync; endloop(1);
projects/hip/samples/1_Utils/hipCommander/perf/scripts/latency_2_d2h_sync_kernel.hcm
+9
Zobrazit soubor
@@ -0,0 +1,9 @@
setstream(1);
loop(10); D2H; streamsync; NullKernel; streamsync;D2H; streamsync; NullKernel; streamsync; endloop(1);
loop(10); D2H; streamsync; NullKernel; streamsync;D2H; streamsync; NullKernel; streamsync; endloop(1);
loop(100); D2H; streamsync; NullKernel; streamsync;D2H; streamsync; NullKernel; streamsync; endloop(1);
loop(100); D2H; streamsync; NullKernel; streamsync;D2H; streamsync; NullKernel; streamsync; endloop(1);
loop(1000); D2H; streamsync; NullKernel; streamsync;D2H; streamsync; NullKernel; streamsync; endloop(1);
loop(1000); D2H; streamsync; NullKernel; streamsync;D2H; streamsync; NullKernel; streamsync; endloop(1);
loop(10000); D2H; streamsync; NullKernel; streamsync;D2H; streamsync; NullKernel; streamsync; endloop(1);
loop(10000); D2H; streamsync; NullKernel; streamsync;D2H; streamsync; NullKernel; streamsync; endloop(1);
projects/hip/samples/1_Utils/hipCommander/perf/scripts/latency_2_h2d_d2h.hcm
+9
Zobrazit soubor
@@ -0,0 +1,9 @@
setstream(1);
loop(10); H2D; D2H; streamsync;H2D; D2H; streamsync; endloop(1);
loop(10); H2D; D2H; streamsync;H2D; D2H; streamsync; endloop(1);
loop(100); H2D; D2H; streamsync;H2D; D2H; streamsync; endloop(1);
loop(100); H2D; D2H; streamsync;H2D; D2H; streamsync; endloop(1);
loop(1000); H2D; D2H; streamsync;H2D; D2H; streamsync; endloop(1);
loop(1000); H2D; D2H; streamsync;H2D; D2H; streamsync; endloop(1);
loop(10000); H2D; D2H; streamsync;H2D; D2H; streamsync; endloop(1);
loop(10000); H2D; D2H; streamsync;H2D; D2H; streamsync; endloop(1);
projects/hip/samples/1_Utils/hipCommander/perf/scripts/latency_2_h2d_kernel.hcm
+9
Zobrazit soubor
@@ -0,0 +1,9 @@
setstream(1);
loop(10); H2D; NullKernel; streamsync;H2D; NullKernel; streamsync; endloop(1);
loop(10); H2D; NullKernel; streamsync; H2D; NullKernel; streamsync;endloop(1);
loop(100); H2D; NullKernel; streamsync; H2D; NullKernel; streamsync;endloop(1);
loop(100); H2D; NullKernel; streamsync; H2D; NullKernel; streamsync; endloop(1);
loop(1000); H2D; NullKernel; streamsync; H2D; NullKernel; streamsync; endloop(1);
loop(1000); H2D; NullKernel; streamsync; H2D; NullKernel; streamsync; endloop(1);
loop(10000); H2D; NullKernel; streamsync;H2D; NullKernel; streamsync; endloop(1);
loop(10000); H2D; NullKernel; streamsync; H2D; NullKernel; streamsync; endloop(1);
projects/hip/samples/1_Utils/hipCommander/perf/scripts/latency_2_h2d_sync_d2h.hcm
+9
Zobrazit soubor
@@ -0,0 +1,9 @@
setstream(1);
loop(10); H2D; streamsync; D2H; streamsync;H2D; streamsync; D2H; streamsync; endloop(1);
loop(10); H2D; streamsync; D2H; streamsync;H2D; streamsync; D2H; streamsync; endloop(1);
loop(100); H2D; streamsync; D2H; streamsync;H2D; streamsync; D2H; streamsync; endloop(1);
loop(100); H2D; streamsync; D2H; streamsync;H2D; streamsync; D2H; streamsync; endloop(1);
loop(1000); H2D; streamsync; D2H; streamsync;H2D; streamsync; D2H; streamsync; endloop(1);
loop(1000); H2D; streamsync; D2H; streamsync;H2D; streamsync; D2H; streamsync; endloop(1);
loop(10000); H2D; streamsync; D2H; streamsync;H2D; streamsync; D2H; streamsync; endloop(1);
loop(10000); H2D; streamsync; D2H; streamsync;H2D; streamsync; D2H; streamsync; endloop(1);
projects/hip/samples/1_Utils/hipCommander/perf/scripts/latency_2_h2d_sync_kernel.hcm
+9
Zobrazit soubor
@@ -0,0 +1,9 @@
setstream(1);
loop(10); H2D; streamsync; NullKernel; streamsync;H2D; streamsync; NullKernel; streamsync; endloop(1);
loop(10); H2D; streamsync; NullKernel; streamsync;H2D; streamsync; NullKernel; streamsync; endloop(1);
loop(100); H2D; streamsync; NullKernel; streamsync;H2D; streamsync; NullKernel; streamsync; endloop(1);
loop(100); H2D; streamsync; NullKernel; streamsync;H2D; streamsync; NullKernel; streamsync; endloop(1);
loop(1000); H2D; streamsync; NullKernel; streamsync;H2D; streamsync; NullKernel; streamsync; endloop(1);
loop(1000); H2D; streamsync; NullKernel; streamsync;H2D; streamsync; NullKernel; streamsync; endloop(1);
loop(10000); H2D; streamsync; NullKernel; streamsync;H2D; streamsync; NullKernel; streamsync; endloop(1);
loop(10000); H2D; streamsync; NullKernel; streamsync;H2D; streamsync; NullKernel; streamsync; endloop(1);
projects/hip/samples/1_Utils/hipCommander/perf/scripts/latency_2_kernel_d2h.hcm
+9
Zobrazit soubor
@@ -0,0 +1,9 @@
setstream(1);
loop(10); NullKernel; D2H; streamsync;NullKernel; D2H; streamsync; endloop(1);
loop(10); NullKernel; D2H; streamsync;NullKernel; D2H; streamsync; endloop(1);
loop(100); NullKernel; D2H; streamsync;NullKernel; D2H; streamsync; endloop(1);
loop(100); NullKernel; D2H; streamsync;NullKernel; D2H; streamsync; endloop(1);
loop(1000); NullKernel; D2H; streamsync;NullKernel; D2H; streamsync; endloop(1);
loop(1000); NullKernel; D2H; streamsync;NullKernel; D2H; streamsync; endloop(1);
loop(10000); NullKernel; D2H; streamsync;NullKernel; D2H; streamsync; endloop(1);
loop(10000); NullKernel; D2H; streamsync;NullKernel; D2H; streamsync; endloop(1);
projects/hip/samples/1_Utils/hipCommander/perf/scripts/latency_2_kernel_h2d.hcm
+9
Zobrazit soubor
@@ -0,0 +1,9 @@
setstream(1);
loop(10); NullKernel; H2D; streamsync;NullKernel; H2D; streamsync; endloop(1);
loop(10); NullKernel; H2D; streamsync;NullKernel; H2D; streamsync; endloop(1);
loop(100); NullKernel; H2D; streamsync;NullKernel; H2D; streamsync; endloop(1);
loop(100); NullKernel; H2D; streamsync;NullKernel; H2D; streamsync; endloop(1);
loop(1000); NullKernel; H2D; streamsync;NullKernel; H2D; streamsync; endloop(1);
loop(1000); NullKernel; H2D; streamsync;NullKernel; H2D; streamsync; endloop(1);
loop(10000); NullKernel; H2D; streamsync;NullKernel; H2D; streamsync; endloop(1);
loop(10000); NullKernel; H2D; streamsync;NullKernel; H2D; streamsync; endloop(1);
projects/hip/samples/1_Utils/hipCommander/perf/scripts/latency_2_kernel_sync_d2h.hcm
+9
Zobrazit soubor
@@ -0,0 +1,9 @@
setstream(1);
loop(10); NullKernel; streamsync; D2H; streamsync;NullKernel; streamsync; D2H; streamsync; endloop(1);
loop(10); NullKernel; streamsync; D2H; streamsync;NullKernel; streamsync; D2H; streamsync; endloop(1);
loop(100); NullKernel; streamsync; D2H; streamsync;NullKernel; streamsync; D2H; streamsync; endloop(1);
loop(100); NullKernel; streamsync; D2H; streamsync;NullKernel; streamsync; D2H; streamsync; endloop(1);
loop(1000); NullKernel; streamsync; D2H; streamsync;NullKernel; streamsync; D2H; streamsync; endloop(1);
loop(1000); NullKernel; streamsync; D2H; streamsync;NullKernel; streamsync; D2H; streamsync; endloop(1);
loop(10000); NullKernel; streamsync; D2H; streamsync;NullKernel; streamsync; D2H; streamsync; endloop(1);
loop(10000); NullKernel; streamsync; D2H; streamsync;NullKernel; streamsync; D2H; streamsync; endloop(1);
projects/hip/samples/1_Utils/hipCommander/perf/scripts/latency_2_kernel_sync_h2d.hcm
+9
Zobrazit soubor
@@ -0,0 +1,9 @@
setstream(1);
loop(10); NullKernel; streamsync; H2D; streamsync;NullKernel; streamsync; H2D; streamsync; endloop(1);
loop(10); NullKernel; streamsync; H2D; streamsync;NullKernel; streamsync; H2D; streamsync; endloop(1);
loop(100); NullKernel; streamsync; H2D; streamsync;NullKernel; streamsync; H2D; streamsync; endloop(1);
loop(100); NullKernel; streamsync; H2D; streamsync;NullKernel; streamsync; H2D; streamsync; endloop(1);
loop(1000); NullKernel; streamsync; H2D; streamsync;NullKernel; streamsync; H2D; streamsync; endloop(1);
loop(1000); NullKernel; streamsync; H2D; streamsync;NullKernel; streamsync; H2D; streamsync; endloop(1);
loop(10000); NullKernel; streamsync; H2D; streamsync;NullKernel; streamsync; H2D; streamsync; endloop(1);
loop(10000); NullKernel; streamsync; H2D; streamsync;NullKernel; streamsync; H2D; streamsync; endloop(1);
projects/hip/samples/1_Utils/hipCommander/perf/scripts/latency_2_sync.hcm
+9
Zobrazit soubor
@@ -0,0 +1,9 @@
setstream(1);
loop(10); streamsync; streamsync; endloop(1);
loop(10); streamsync; streamsync; endloop(1);
loop(100); streamsync; streamsync; endloop(1);
loop(100); streamsync; streamsync; endloop(1);
loop(1000); streamsync; streamsync; endloop(1);
loop(1000); streamsync; streamsync; endloop(1);
loop(10000); streamsync; streamsync; endloop(1);
loop(10000); streamsync; streamsync; endloop(1);
projects/hip/samples/1_Utils/hipCommander/perf/scripts/latency_2d2h.hcm
+10
Zobrazit soubor
@@ -0,0 +1,10 @@
setstream(1);
loop(10); D2H; streamsync; D2H; streamsync; endloop(1);
loop(10); D2H; streamsync; D2H; streamsync; endloop(1);
loop(100); D2H; streamsync; D2H; streamsync; endloop(1);
loop(100); D2H; streamsync; D2H; streamsync; endloop(1);
loop(1000); D2H;streamsync; D2H; streamsync; endloop(1);
loop(1000); D2H; streamsync; D2H; streamsync; endloop(1);
loop(1000); D2H; streamsync; D2H; streamsync; endloop(1);
loop(10000); D2H; streamsync; D2H; streamsync; endloop(1);
loop(10000); D2H; streamsync; D2H; streamsync; endloop(1);
projects/hip/samples/1_Utils/hipCommander/perf/scripts/latency_2d2h_wosync.hcm
+10
Zobrazit soubor
@@ -0,0 +1,10 @@
setstream(1);
loop(10); D2H; D2H; streamsync; endloop(1);
loop(10); D2H; D2H; streamsync; endloop(1);
loop(100); D2H; D2H; streamsync; endloop(1);
loop(100); D2H; D2H; streamsync; endloop(1);
loop(1000); D2H; D2H; streamsync; endloop(1);
loop(1000); D2H; D2H; streamsync; endloop(1);
loop(1000); D2H; D2H; streamsync; endloop(1);
loop(10000); D2H; D2H; streamsync; endloop(1);
loop(10000); D2H; D2H; streamsync; endloop(1);
projects/hip/samples/1_Utils/hipCommander/perf/scripts/latency_2h2d.hcm
+10
Zobrazit soubor
@@ -0,0 +1,10 @@
setstream(1);
loop(10); H2D; streamsync; H2D; streamsync; endloop(1);
loop(10); H2D; streamsync; H2D; streamsync; endloop(1);
loop(100); H2D; streamsync; H2D; streamsync; endloop(1);
loop(100); H2D; streamsync; H2D; streamsync; endloop(1);
loop(1000); H2D;streamsync; H2D; streamsync; endloop(1);
loop(1000); H2D; streamsync; H2D; streamsync; endloop(1);
loop(1000); H2D; streamsync; H2D; streamsync; endloop(1);
loop(10000); H2D; streamsync; H2D; streamsync; endloop(1);
loop(10000); H2D; streamsync; H2D; streamsync; endloop(1);
projects/hip/samples/1_Utils/hipCommander/perf/scripts/latency_2h2d_kernel.hcm
+9
Zobrazit soubor
@@ -0,0 +1,9 @@
setstream(1);
loop(10); H2D; streamsync; NullKernel; streamsync; H2D; streamsync; NullKernel; streamsync;endloop(1);
loop(10); H2D; streamsync; NullKernel; streamsync;H2D; streamsync; NullKernel; streamsync; endloop(1);
loop(100); H2D; streamsync; NullKernel; streamsync; H2D; streamsync; NullKernel; streamsync;endloop(1);
loop(100); H2D; streamsync; NullKernel; streamsync; H2D; streamsync; NullKernel; streamsync;endloop(1);
loop(1000); H2D; streamsync; NullKernel; streamsync; H2D; streamsync; NullKernel; streamsync;endloop(1);
loop(1000); H2D; streamsync; NullKernel; streamsync; H2D; streamsync; NullKernel; streamsync;endloop(1);
loop(10000); H2D; streamsync; NullKernel; streamsync;H2D; streamsync; NullKernel; streamsync; endloop(1);
loop(10000); H2D; streamsync; NullKernel; streamsync; H2D; streamsync; NullKernel; streamsync;endloop(1);
projects/hip/samples/1_Utils/hipCommander/perf/scripts/latency_2h2d_kernel_wosync.hcm
+9
Zobrazit soubor
@@ -0,0 +1,9 @@
setstream(1);
loop(10); H2D; NullKernel; streamsync;H2D; NullKernel; streamsync; endloop(1);
loop(10); H2D; NullKernel; streamsync; H2D; NullKernel; streamsync;endloop(1);
loop(100); H2D; NullKernel; streamsync; H2D; NullKernel; streamsync;endloop(1);
loop(100); H2D; NullKernel; streamsync; H2D; NullKernel; streamsync;endloop(1);
loop(1000); H2D; NullKernel; streamsync;H2D; NullKernel; streamsync; endloop(1);
loop(1000); H2D; NullKernel; streamsync;H2D; NullKernel; streamsync; endloop(1);
loop(10000); H2D; NullKernel; streamsync;H2D; NullKernel; streamsync; endloop(1);
loop(10000); H2D ; NullKernel; streamsync;H2D; NullKernel; streamsync; endloop(1);
projects/hip/samples/1_Utils/hipCommander/perf/scripts/latency_2h2d_wosync.hcm
+10
Zobrazit soubor
@@ -0,0 +1,10 @@
setstream(1);
loop(10); H2D; H2D; streamsync; endloop(1);
loop(10); H2D; H2D; streamsync; endloop(1);
loop(100); H2D; H2D; streamsync; endloop(1);
loop(100); H2D; H2D; streamsync; endloop(1);
loop(1000); H2D; H2D; streamsync; endloop(1);
loop(1000); H2D; H2D; streamsync; endloop(1);
loop(1000); H2D; H2D; streamsync; endloop(1);
loop(10000); H2D; H2D; streamsync; endloop(1);
loop(10000); H2D; H2D; streamsync; endloop(1);
projects/hip/samples/1_Utils/hipCommander/perf/scripts/latency_2kernels.hcm
+10
Zobrazit soubor
@@ -0,0 +1,10 @@
setstream(1);
loop(10); NullKernel; streamsync; NullKernel; streamsync; endloop(1);
loop(10); NullKernel; streamsync; NullKernel; streamsync; endloop(1);
loop(100); NullKernel; streamsync; NullKernel; streamsync; endloop(1);
loop(100); NullKernel; streamsync; NullKernel; streamsync; endloop(1);
loop(1000); NullKernel; streamsync; NullKernel; streamsync; endloop(1);
loop(1000); NullKernel; streamsync; NullKernel; streamsync; endloop(1);
loop(1000); NullKernel; streamsync; NullKernel; streamsync; endloop(1);
loop(10000); NullKernel; streamsync; NullKernel; streamsync; endloop(1);
loop(10000); NullKernel; streamsync; NullKernel; streamsync; endloop(1);
projects/hip/samples/1_Utils/hipCommander/perf/scripts/latency_2kernels_wosync.hcm
+10
Zobrazit soubor
@@ -0,0 +1,10 @@
setstream(1);
loop(10); NullKernel; NullKernel; streamsync; endloop(1);
loop(10); NullKernel; NullKernel; streamsync; endloop(1);
loop(100); NullKernel; NullKernel; streamsync; endloop(1);
loop(100); NullKernel; NullKernel; streamsync; endloop(1);
loop(1000); NullKernel; NullKernel; streamsync; endloop(1);
loop(1000); NullKernel; NullKernel; streamsync; endloop(1);
loop(1000); NullKernel; NullKernel; streamsync; endloop(1);
loop(10000); NullKernel; NullKernel; streamsync; endloop(1);
loop(10000); NullKernel; NullKernel; streamsync; endloop(1);
projects/hip/samples/1_Utils/hipCommander/perf/scripts/latency_3_d2h_h2d.hcm
+9
Zobrazit soubor
@@ -0,0 +1,9 @@
setstream(1);
loop(10); D2H; H2D; streamsync;D2H; H2D; streamsync; D2H; H2D; streamsync; endloop(1);
loop(10); D2H; H2D; streamsync;D2H; H2D; streamsync; D2H; H2D; streamsync; endloop(1);
loop(100); D2H; H2D; streamsync;D2H; H2D; streamsync; D2H; H2D; streamsync; endloop(1);
loop(100); D2H; H2D; streamsync;D2H; H2D; streamsync; D2H; H2D; streamsync; endloop(1);
loop(1000); D2H; H2D; streamsync;D2H; H2D; streamsync; D2H; H2D; streamsync; endloop(1);
loop(1000); D2H; H2D; streamsync;D2H; H2D; streamsync; D2H; H2D; streamsync; endloop(1);
loop(10000); D2H; H2D; streamsync;D2H; H2D; streamsync; D2H; H2D; streamsync; endloop(1);
loop(10000); D2H; H2D; streamsync;D2H; H2D; streamsync; D2H; H2D; streamsync; endloop(1);
projects/hip/samples/1_Utils/hipCommander/perf/scripts/latency_3_d2h_kernel.hcm
+9
Zobrazit soubor
@@ -0,0 +1,9 @@
setstream(1);
loop(10); D2H; NullKernel; streamsync;D2H; NullKernel; streamsync;streamsync; D2H; NullKernel; streamsync; endloop(1);
loop(10); D2H; NullKernel; streamsync;D2H; NullKernel; streamsync;streamsync; D2H; NullKernel; streamsync; endloop(1);
loop(100); D2H; NullKernel; streamsync;D2H; NullKernel; streamsync;streamsync; D2H; NullKernel; streamsync; endloop(1);
loop(100); D2H; NullKernel; streamsync;D2H; NullKernel; streamsync;streamsync; D2H; NullKernel; streamsync; endloop(1);
loop(1000); D2H; NullKernel; streamsync;D2H; NullKernel; streamsync;streamsync; D2H; NullKernel; streamsync; endloop(1);
loop(1000); D2H; NullKernel; streamsync;D2H; NullKernel; streamsync;streamsync; D2H; NullKernel; streamsync; endloop(1);
loop(10000); D2H; NullKernel; streamsync;D2H; NullKernel; streamsync;streamsync; D2H; NullKernel; streamsync; endloop(1);
loop(10000); D2H; NullKernel; streamsync;D2H; NullKernel; streamsync;streamsync; D2H; NullKernel; streamsync; endloop(1);
projects/hip/samples/1_Utils/hipCommander/perf/scripts/latency_3_d2h_sync_h2d.hcm
+9
Zobrazit soubor
@@ -0,0 +1,9 @@
setstream(1);
loop(10); D2H; streamsync; H2D; streamsync;D2H; streamsync; H2D; streamsync; D2H; streamsync; H2D;streamsync; endloop(1);
loop(10); D2H; streamsync; H2D; streamsync;D2H; streamsync; H2D; streamsync; D2H; streamsync; H2D;streamsync; endloop(1);
loop(100); D2H; streamsync; H2D; streamsync;D2H; streamsync; H2D; streamsync; D2H; streamsync; H2D;streamsync; endloop(1);
loop(100); D2H; streamsync; H2D; streamsync;D2H; streamsync; H2D; streamsync; D2H; streamsync; H2D;streamsync; endloop(1);
loop(1000); D2H; streamsync; H2D; streamsync;D2H; streamsync; H2D; streamsync; D2H; streamsync; H2D;streamsync; endloop(1);
loop(1000); D2H; streamsync; H2D; streamsync;D2H; streamsync; H2D; streamsync; D2H; streamsync; H2D;streamsync; endloop(1);
loop(10000); D2H; streamsync; H2D; streamsync;D2H; streamsync; H2D; streamsync; D2H; streamsync; H2D;streamsync; endloop(1);
loop(10000); D2H; streamsync; H2D; streamsync;D2H; streamsync; H2D; streamsync; D2H; streamsync; H2D;streamsync; endloop(1);
projects/hip/samples/1_Utils/hipCommander/perf/scripts/latency_3_d2h_sync_kernel.hcm
+9
Zobrazit soubor
@@ -0,0 +1,9 @@
setstream(1);
loop(10); D2H; streamsync; NullKernel; streamsync;D2H; streamsync; NullKernel; streamsync; D2H; streamsync; NullKernel;streamsync; endloop(1);
loop(10); D2H; streamsync; NullKernel; streamsync;D2H; streamsync; NullKernel; streamsync; D2H; streamsync; NullKernel;streamsync; endloop(1);
loop(100); D2H; streamsync; NullKernel; streamsync;D2H; streamsync; NullKernel; streamsync; D2H; streamsync; NullKernel;streamsync; endloop(1);
loop(100); D2H; streamsync; NullKernel; streamsync;D2H; streamsync; NullKernel; streamsync; D2H; streamsync; NullKernel;streamsync; endloop(1);
loop(1000); D2H; streamsync; NullKernel; streamsync;D2H; streamsync; NullKernel; streamsync; D2H; streamsync; NullKernel;streamsync; endloop(1);
loop(1000); D2H; streamsync; NullKernel; streamsync;D2H; streamsync; NullKernel; streamsync; D2H; streamsync; NullKernel;streamsync; endloop(1);
loop(10000); D2H; streamsync; NullKernel; streamsync;D2H; streamsync; NullKernel; streamsync; D2H; streamsync; NullKernel;streamsync; endloop(1);
loop(10000); D2H; streamsync; NullKernel; streamsync;D2H; streamsync; NullKernel; streamsync; D2H; streamsync; NullKernel;streamsync; endloop(1);
projects/hip/samples/1_Utils/hipCommander/perf/scripts/latency_3_h2d_d2h.hcm
+10
Zobrazit soubor
@@ -0,0 +1,10 @@
setstream(1);
loop(10); H2D; D2H; streamsync;H2D; D2H; streamsync; H2D; D2H; streamsync; endloop(1);
loop(10); H2D; D2H; streamsync;H2D; D2H; streamsync; H2D; D2H; streamsync; endloop(1);
loop(100); H2D; D2H; streamsync;H2D; D2H; streamsync; H2D; D2H; streamsync; endloop(1);
loop(100); H2D; D2H; streamsync;H2D; D2H; streamsync; H2D; D2H; streamsync; endloop(1);
loop(1000); H2D; D2H; streamsync;H2D; D2H; streamsync; H2D; D2H; streamsync; endloop(1);
loop(1000); H2D; D2H; streamsync;H2D; D2H; streamsync; H2D; D2H; streamsync; endloop(1);
loop(10000); H2D; D2H; streamsync;H2D; D2H; streamsync; H2D; D2H; streamsync; endloop(1);
loop(10000); H2D; D2H; streamsync;H2D; D2H; streamsync; H2D; D2H; streamsync; endloop(1);
projects/hip/samples/1_Utils/hipCommander/perf/scripts/latency_3_h2d_kernel.hcm
+9
Zobrazit soubor
@@ -0,0 +1,9 @@
setstream(1);
loop(10); H2D; NullKernel; streamsync;H2D; NullKernel; streamsync;streamsync; H2D; NullKernel; streamsync; endloop(1);
loop(10); H2D; NullKernel; streamsync;H2D; NullKernel; streamsync;streamsync; H2D; NullKernel; streamsync; endloop(1);
loop(100); H2D; NullKernel; streamsync;H2D; NullKernel; streamsync;streamsync; H2D; NullKernel; streamsync; endloop(1);
loop(100); H2D; NullKernel; streamsync;H2D; NullKernel; streamsync;streamsync; H2D; NullKernel; streamsync; endloop(1);
loop(1000); H2D; NullKernel; streamsync;H2D; NullKernel; streamsync;streamsync; H2D; NullKernel; streamsync; endloop(1);
loop(1000); H2D; NullKernel; streamsync;H2D; NullKernel; streamsync;streamsync; H2D; NullKernel; streamsync; endloop(1);
loop(10000); H2D; NullKernel; streamsync;H2D; NullKernel; streamsync;streamsync; H2D; NullKernel; streamsync; endloop(1);
loop(10000); H2D; NullKernel; streamsync;H2D; NullKernel; streamsync;streamsync; H2D; NullKernel; streamsync; endloop(1);
projects/hip/samples/1_Utils/hipCommander/perf/scripts/latency_3_h2d_sync_d2h.hcm
+9
Zobrazit soubor
@@ -0,0 +1,9 @@
setstream(1);
loop(10); H2D; streamsync; D2H; streamsync;H2D; streamsync; D2H; streamsync; H2D; streamsync; D2H;streamsync; endloop(1);
loop(10); H2D; streamsync; D2H; streamsync;H2D; streamsync; D2H; streamsync; H2D; streamsync; D2H;streamsync; endloop(1);
loop(100); H2D; streamsync; D2H; streamsync;H2D; streamsync; D2H; streamsync; H2D; streamsync; D2H;streamsync; endloop(1);
loop(100); H2D; streamsync; D2H; streamsync;H2D; streamsync; D2H; streamsync; H2D; streamsync; D2H;streamsync; endloop(1);
loop(1000); H2D; streamsync; D2H; streamsync;H2D; streamsync; D2H; streamsync; H2D; streamsync; D2H;streamsync; endloop(1);
loop(1000); H2D; streamsync; D2H; streamsync;H2D; streamsync; D2H; streamsync; H2D; streamsync; D2H;streamsync; endloop(1);
loop(10000); H2D; streamsync; D2H; streamsync;H2D; streamsync; D2H; streamsync; H2D; streamsync; D2H;streamsync; endloop(1);
loop(10000); H2D; streamsync; D2H; streamsync;H2D; streamsync; D2H; streamsync; H2D; streamsync; D2H;streamsync; endloop(1);
projects/hip/samples/1_Utils/hipCommander/perf/scripts/latency_3_h2d_sync_kernel.hcm
+9
Zobrazit soubor
@@ -0,0 +1,9 @@
setstream(1);
loop(10); H2D; streamsync; NullKernel; streamsync;H2D; streamsync; NullKernel; streamsync; H2D; streamsync; NullKernel;streamsync; endloop(1);
loop(10); H2D; streamsync; NullKernel; streamsync;H2D; streamsync; NullKernel; streamsync; H2D; streamsync; NullKernel;streamsync; endloop(1);
loop(100); H2D; streamsync; NullKernel; streamsync;H2D; streamsync; NullKernel; streamsync; H2D; streamsync; NullKernel;streamsync; endloop(1);
loop(100); H2D; streamsync; NullKernel; streamsync;H2D; streamsync; NullKernel; streamsync; H2D; streamsync; NullKernel;streamsync; endloop(1);
loop(1000); H2D; streamsync; NullKernel; streamsync;H2D; streamsync; NullKernel; streamsync; H2D; streamsync; NullKernel;streamsync; endloop(1);
loop(1000); H2D; streamsync; NullKernel; streamsync;H2D; streamsync; NullKernel; streamsync; H2D; streamsync; NullKernel;streamsync; endloop(1);
loop(10000); H2D; streamsync; NullKernel; streamsync;H2D; streamsync; NullKernel; streamsync; H2D; streamsync; NullKernel;streamsync; endloop(1);
loop(10000); H2D; streamsync; NullKernel; streamsync;H2D; streamsync; NullKernel; streamsync; H2D; streamsync; NullKernel;streamsync; endloop(1);
projects/hip/samples/1_Utils/hipCommander/perf/scripts/latency_3_kernel_d2h.hcm
+9
Zobrazit soubor
@@ -0,0 +1,9 @@
setstream(1);
loop(10); NullKernel; D2H; streamsync;NullKernel; D2H; streamsync; NullKernel; D2H; streamsync;endloop(1);
loop(10); NullKernel; D2H; streamsync;NullKernel; D2H; streamsync; NullKernel; D2H; streamsync;endloop(1);
loop(100); NullKernel; D2H; streamsync;NullKernel; D2H; streamsync; NullKernel; D2H; streamsync;endloop(1);
loop(100); NullKernel; D2H; streamsync;NullKernel; D2H; streamsync; NullKernel; D2H; streamsync;endloop(1);
loop(1000); NullKernel; D2H; streamsync;NullKernel; D2H; streamsync; NullKernel; D2H; streamsync;endloop(1);
loop(1000); NullKernel; D2H; streamsync;NullKernel; D2H; streamsync; NullKernel; D2H; streamsync;endloop(1);
loop(10000); NullKernel; D2H; streamsync;NullKernel; D2H; streamsync; NullKernel; D2H; streamsync;endloop(1);
loop(10000); NullKernel; D2H; streamsync;NullKernel; D2H; streamsync; NullKernel; D2H; streamsync;endloop(1);
projects/hip/samples/1_Utils/hipCommander/perf/scripts/latency_3_kernel_h2d.hcm
+9
Zobrazit soubor
@@ -0,0 +1,9 @@
setstream(1);
loop(10); NullKernel; H2D; streamsync;NullKernel; H2D; streamsync; NullKernel; H2D; streamsync;endloop(1);
loop(10); NullKernel; H2D; streamsync;NullKernel; H2D; streamsync; NullKernel; H2D; streamsync;endloop(1);
loop(100); NullKernel; H2D; streamsync;NullKernel; H2D; streamsync; NullKernel; H2D; streamsync;endloop(1);
loop(100); NullKernel; H2D; streamsync;NullKernel; H2D; streamsync; NullKernel; H2D; streamsync;endloop(1);
loop(1000); NullKernel; H2D; streamsync;NullKernel; H2D; streamsync; NullKernel; H2D; streamsync;endloop(1);
loop(1000); NullKernel; H2D; streamsync;NullKernel; H2D; streamsync; NullKernel; H2D; streamsync;endloop(1);
loop(10000); NullKernel; H2D; streamsync;NullKernel; H2D; streamsync; NullKernel; H2D; streamsync;endloop(1);
loop(10000); NullKernel; H2D; streamsync;NullKernel; H2D; streamsync; NullKernel; H2D; streamsync;endloop(1);
projects/hip/samples/1_Utils/hipCommander/perf/scripts/latency_3_kernel_sync_d2h.hcm
+9
Zobrazit soubor
@@ -0,0 +1,9 @@
setstream(1);
loop(10); NullKernel; streamsync; D2H; streamsync;NullKernel; streamsync; D2H; streamsync;NullKernel; streamsync; D2H; streamsync; endloop(1);
loop(10); NullKernel; streamsync; D2H; streamsync;NullKernel; streamsync; D2H; streamsync;NullKernel; streamsync; D2H; streamsync; endloop(1);
loop(100); NullKernel; streamsync; D2H; streamsync;NullKernel; streamsync; D2H; streamsync;NullKernel; streamsync; D2H; streamsync; endloop(1);
loop(100); NullKernel; streamsync; D2H; streamsync;NullKernel; streamsync; D2H; streamsync;NullKernel; streamsync; D2H; streamsync; endloop(1);
loop(1000); NullKernel; streamsync; D2H; streamsync;NullKernel; streamsync; D2H; streamsync;NullKernel; streamsync; D2H; streamsync; endloop(1);
loop(1000); NullKernel; streamsync; D2H; streamsync;NullKernel; streamsync; D2H; streamsync;NullKernel; streamsync; D2H; streamsync; endloop(1);
loop(10000); NullKernel; streamsync; D2H; streamsync;NullKernel; streamsync; D2H; streamsync;NullKernel; streamsync; D2H; streamsync; endloop(1);
loop(10000); NullKernel; streamsync; D2H; streamsync;NullKernel; streamsync; D2H; streamsync;NullKernel; streamsync; D2H; streamsync; endloop(1);
projects/hip/samples/1_Utils/hipCommander/perf/scripts/latency_3_kernel_sync_h2d.hcm
+9
Zobrazit soubor
@@ -0,0 +1,9 @@
setstream(1);
loop(10); NullKernel; streamsync; H2D; streamsync;NullKernel; streamsync; H2D; streamsync;NullKernel; streamsync; H2D; streamsync; endloop(1);
loop(10); NullKernel; streamsync; H2D; streamsync;NullKernel; streamsync; H2D; streamsync;NullKernel; streamsync; H2D; streamsync; endloop(1);
loop(100); NullKernel; streamsync; H2D; streamsync;NullKernel; streamsync; H2D; streamsync;NullKernel; streamsync; H2D; streamsync; endloop(1);
loop(100); NullKernel; streamsync; H2D; streamsync;NullKernel; streamsync; H2D; streamsync;NullKernel; streamsync; H2D; streamsync; endloop(1);
loop(1000); NullKernel; streamsync; H2D; streamsync;NullKernel; streamsync; H2D; streamsync;NullKernel; streamsync; H2D; streamsync; endloop(1);
loop(1000); NullKernel; streamsync; H2D; streamsync;NullKernel; streamsync; H2D; streamsync;NullKernel; streamsync; H2D; streamsync; endloop(1);
loop(10000); NullKernel; streamsync; H2D; streamsync;NullKernel; streamsync; H2D; streamsync;NullKernel; streamsync; H2D; streamsync; endloop(1);
loop(10000); NullKernel; streamsync; H2D; streamsync;NullKernel; streamsync; H2D; streamsync;NullKernel; streamsync; H2D; streamsync; endloop(1);
projects/hip/samples/1_Utils/hipCommander/perf/scripts/latency_3_sync.hcm
+9
Zobrazit soubor
@@ -0,0 +1,9 @@
setstream(1);
loop(10);streamsync; streamsync; streamsync; endloop(1);
loop(10);streamsync; streamsync; streamsync; endloop(1);
loop(100);streamsync; streamsync; streamsync; endloop(1);
loop(100);streamsync; streamsync; streamsync; endloop(1);
loop(1000);streamsync; streamsync; streamsync; endloop(1);
loop(1000);streamsync; streamsync; streamsync; endloop(1);
loop(10000);streamsync; streamsync; streamsync; endloop(1);
loop(10000);streamsync; streamsync; streamsync; endloop(1);
projects/hip/samples/1_Utils/hipCommander/perf/scripts/latency_3d2h.hcm
+9
Zobrazit soubor
@@ -0,0 +1,9 @@
setstream(1);
loop(10); D2H; streamsync; D2H; streamsync; D2H; streamsync; endloop(1);
loop(10); D2H; streamsync; D2H; streamsync; D2H; streamsync; endloop(1);
loop(100); D2H; streamsync; D2H; streamsync; D2H; streamsync; endloop(1);
loop(100); D2H; streamsync; D2H; streamsync; D2H; streamsync; endloop(1);
loop(1000); D2H;streamsync; D2H; streamsync; D2H; streamsync; endloop(1);
loop(1000); D2H; streamsync; D2H; streamsync; D2H; streamsync; endloop(1);
loop(10000); D2H; streamsync; D2H; streamsync; D2H; streamsync; endloop(1);
loop(10000); D2H; streamsync; D2H; streamsync; D2H; streamsync; endloop(1);
projects/hip/samples/1_Utils/hipCommander/perf/scripts/latency_3d2h_wosync.hcm
+9
Zobrazit soubor
@@ -0,0 +1,9 @@
setstream(1);
loop(10); D2H; D2H; D2H; streamsync; endloop(1);
loop(10); D2H; D2H; D2H; streamsync; endloop(1);
loop(100); D2H; D2H; D2H; streamsync; endloop(1);
loop(100); D2H; D2H; D2H; streamsync; endloop(1);
loop(1000); D2H; D2H; D2H; streamsync; endloop(1);
loop(1000); D2H; D2H; D2H; streamsync; endloop(1);
loop(10000); D2H; D2H; D2H; streamsync; endloop(1);
loop(10000); D2H; D2H; D2H;streamsync; endloop(1);
projects/hip/samples/1_Utils/hipCommander/perf/scripts/latency_3h2d.hcm
+10
Zobrazit soubor
@@ -0,0 +1,10 @@
setstream(1);
loop(10); H2D; streamsync;H2D;streamsync; H2D; streamsync; endloop(1);
loop(10); H2D; streamsync;H2D;streamsync; H2D; streamsync; endloop(1);
loop(100); H2D; streamsync;H2D; streamsync;H2D; streamsync; endloop(1);
loop(100); H2D;streamsync; H2D; streamsync;H2D; streamsync; endloop(1);
loop(1000); H2D;streamsync; H2D;streamsync; H2D; streamsync; endloop(1);
loop(1000); H2D;streamsync; H2D; streamsync;H2D; streamsync; endloop(1);
loop(1000); H2D;streamsync; H2D; streamsync;H2D; streamsync; endloop(1);
loop(10000); H2D;streamsync; H2D; streamsync;H2D; streamsync; endloop(1);
loop(10000); H2D;streamsync; H2D;streamsync; H2D; streamsync; endloop(1);
projects/hip/samples/1_Utils/hipCommander/perf/scripts/latency_3h2d_wosync.hcm
+10
Zobrazit soubor
@@ -0,0 +1,10 @@
setstream(1);
loop(10); H2D; H2D; H2D; streamsync; endloop(1);
loop(10); H2D; H2D; H2D; streamsync; endloop(1);
loop(100); H2D; H2D; H2D; streamsync; endloop(1);
loop(100); H2D; H2D; H2D; streamsync; endloop(1);
loop(1000); H2D; H2D; H2D; streamsync; endloop(1);
loop(1000); H2D; H2D; H2D; streamsync; endloop(1);
loop(1000); H2D; H2D; H2D; streamsync; endloop(1);
loop(10000); H2D; H2D; H2D; streamsync; endloop(1);
loop(10000); H2D; H2D; H2D; streamsync; endloop(1);
projects/hip/samples/1_Utils/hipCommander/perf/scripts/latency_3kernels.hcm
+10
Zobrazit soubor
@@ -0,0 +1,10 @@
setstream(1);
loop(10); NullKernel; streamsync; NullKernel; streamsync; NullKernel; streamsync; endloop(1);
loop(10); NullKernel; streamsync; NullKernel; streamsync; NullKernel; streamsync; endloop(1);
loop(100); NullKernel; streamsync; NullKernel; streamsync; NullKernel; streamsync; endloop(1);
loop(100); NullKernel; streamsync; NullKernel; streamsync; NullKernel; streamsync; endloop(1);
loop(1000); NullKernel; streamsync; NullKernel; streamsync; NullKernel; streamsync; endloop(1);
loop(1000); NullKernel; streamsync; NullKernel; streamsync; NullKernel; streamsync; endloop(1);
loop(1000); NullKernel; streamsync; NullKernel; streamsync; NullKernel; streamsync; endloop(1);
loop(10000); NullKernel; streamsync; NullKernel; streamsync; NullKernel; streamsync; endloop(1);
loop(10000); NullKernel; streamsync; NullKernel; streamsync; NullKernel; streamsync; endloop(1);
projects/hip/samples/1_Utils/hipCommander/perf/scripts/latency_3kernels_wosync.hcm
+10
Zobrazit soubor
@@ -0,0 +1,10 @@
setstream(1);
loop(10); NullKernel; NullKernel; NullKernel; streamsync; endloop(1);
loop(10); NullKernel; NullKernel; NullKernel; streamsync; endloop(1);
loop(100); NullKernel; NullKernel; NullKernel; streamsync; endloop(1);
loop(100); NullKernel; NullKernel; NullKernel; streamsync; endloop(1);
loop(1000); NullKernel; NullKernel; NullKernel; streamsync; endloop(1);
loop(1000); NullKernel; NullKernel; NullKernel; streamsync; endloop(1);
loop(1000); NullKernel; NullKernel; NullKernel; streamsync; endloop(1);
loop(10000); NullKernel; NullKernel; NullKernel; streamsync; endloop(1);
loop(10000); NullKernel; NullKernel; NullKernel; streamsync; endloop(1);
projects/hip/samples/1_Utils/hipCommander/perf/scripts/latency_4kernels.hcm
+10
Zobrazit soubor
@@ -0,0 +1,10 @@
setstream(1);
loop(10); NullKernel; NullKernel; NullKernel; NullKernel; streamsync; endloop(1);
loop(10); NullKernel; NullKernel; NullKernel; NullKernel; streamsync; endloop(1);
loop(100); NullKernel; NullKernel; NullKernel; NullKernel;streamsync; endloop(1);
loop(100); NullKernel; NullKernel; NullKernel; NullKernel; streamsync; endloop(1);
loop(1000); NullKernel; NullKernel; NullKernel; NullKernel; streamsync; endloop(1);
loop(1000); NullKernel; NullKernel; NullKernel; NullKernel; streamsync; endloop(1);
loop(1000); NullKernel; NullKernel; NullKernel; NullKernel; streamsync; endloop(1);
loop(10000); NullKernel; NullKernel; NullKernel; NullKernel; streamsync; endloop(1);
loop(10000); NullKernel; NullKernel; NullKernel; NullKernel; streamsync; endloop(1);
projects/hip/samples/1_Utils/hipCommander/perf/scripts/latency_5kernels.hcm
+10
Zobrazit soubor
@@ -0,0 +1,10 @@
setstream(1);
loop(10); NullKernel; NullKernel; NullKernel; NullKernel;NullKernel; streamsync; endloop(1);
loop(10); NullKernel; NullKernel; NullKernel; NullKernel; NullKernel; streamsync; endloop(1);
loop(100); NullKernel; NullKernel; NullKernel; NullKernel; NullKernel;streamsync; endloop(1);
loop(100); NullKernel; NullKernel; NullKernel; NullKernel; NullKernel; streamsync; endloop(1);
loop(1000); NullKernel; NullKernel; NullKernel; NullKernel; NullKernel; streamsync; endloop(1);
loop(1000); NullKernel; NullKernel; NullKernel; NullKernel; NullKernel; streamsync; endloop(1);
loop(1000); NullKernel; NullKernel; NullKernel; NullKernel; NullKernel; streamsync; endloop(1);
loop(10000); NullKernel; NullKernel; NullKernel; NullKernel; NullKernel; streamsync; endloop(1);
loop(10000); NullKernel; NullKernel; NullKernel; NullKernel; NullKernel; streamsync; endloop(1);
projects/hip/samples/1_Utils/hipCommander/perf/scripts/latency_6kernels.hcm
+10
Zobrazit soubor
@@ -0,0 +1,10 @@
setstream(1);
loop(10); NullKernel; NullKernel; NullKernel; NullKernel;NullKernel;NullKernel; streamsync; endloop(1);
loop(10); NullKernel; NullKernel; NullKernel; NullKernel; NullKernel; NullKernel;streamsync; endloop(1);
loop(100); NullKernel; NullKernel; NullKernel; NullKernel; NullKernel;NullKernel;streamsync; endloop(1);
loop(100); NullKernel; NullKernel; NullKernel; NullKernel; NullKernel; NullKernel;streamsync; endloop(1);
loop(1000); NullKernel; NullKernel; NullKernel; NullKernel; NullKernel; NullKernel;streamsync; endloop(1);
loop(1000); NullKernel; NullKernel; NullKernel; NullKernel; NullKernel; NullKernel;streamsync; endloop(1);
loop(1000); NullKernel; NullKernel; NullKernel; NullKernel; NullKernel; NullKernel;streamsync; endloop(1);
loop(10000); NullKernel; NullKernel; NullKernel; NullKernel; NullKernel; NullKernel;streamsync; endloop(1);
loop(10000); NullKernel; NullKernel; NullKernel; NullKernel; NullKernel; NullKernel;streamsync; endloop(1);
projects/hip/samples/1_Utils/hipCommander/perf/scripts/latency_d2h.hcm
+9
Zobrazit soubor
@@ -0,0 +1,9 @@
setstream(1);
loop(10); D2H; streamsync; endloop(1);
loop(10); D2H; streamsync; endloop(1);
loop(100); D2H; streamsync; endloop(1);
loop(100); D2H; streamsync; endloop(1);
loop(1000); D2H; streamsync; endloop(1);
loop(1000); D2H; streamsync; endloop(1);
loop(10000); D2H; streamsync; endloop(1);
loop(10000); D2H; streamsync; endloop(1);
projects/hip/samples/1_Utils/hipCommander/perf/scripts/latency_d2h_h2d.hcm
+9
Zobrazit soubor
@@ -0,0 +1,9 @@
setstream(1);
loop(10); D2H; H2D; streamsync; endloop(1);
loop(10); D2H; H2D; streamsync; endloop(1);
loop(100); D2H; H2D; streamsync; endloop(1);
loop(100); D2H; H2D; streamsync; endloop(1);
loop(1000); D2H; H2D; streamsync; endloop(1);
loop(1000); D2H; H2D; streamsync; endloop(1);
loop(10000); D2H; H2D; streamsync; endloop(1);
loop(10000); D2H; H2D; streamsync; endloop(1);
projects/hip/samples/1_Utils/hipCommander/perf/scripts/latency_d2h_kernel.hcm
+9
Zobrazit soubor
@@ -0,0 +1,9 @@
setstream(1);
loop(10); D2H; NullKernel; streamsync; endloop(1);
loop(10); D2H; NullKernel; streamsync; endloop(1);
loop(100); D2H; NullKernel; streamsync; endloop(1);
loop(100); D2H; NullKernel; streamsync; endloop(1);
loop(1000); D2H; NullKernel; streamsync; endloop(1);
loop(1000); D2H; NullKernel; streamsync; endloop(1);
loop(10000); D2H; NullKernel; streamsync; endloop(1);
loop(10000); D2H; NullKernel; streamsync; endloop(1);
projects/hip/samples/1_Utils/hipCommander/perf/scripts/latency_d2h_sync_h2d.hcm
+9
Zobrazit soubor
@@ -0,0 +1,9 @@
setstream(1);
loop(10); D2H; streamsync; H2D; streamsync; endloop(1);
loop(10); D2H; streamsync; H2D; streamsync; endloop(1);
loop(100); D2H; streamsync; H2D; streamsync; endloop(1);
loop(100); D2H; streamsync; H2D; streamsync; endloop(1);
loop(1000); D2H; streamsync; H2D; streamsync; endloop(1);
loop(1000); D2H; streamsync; H2D; streamsync; endloop(1);
loop(10000); D2H; streamsync; H2D; streamsync; endloop(1);
loop(10000); D2H; streamsync; H2D; streamsync; endloop(1);
projects/hip/samples/1_Utils/hipCommander/perf/scripts/latency_d2h_sync_kernel.hcm
+9
Zobrazit soubor
@@ -0,0 +1,9 @@
setstream(1);
loop(10); D2H; streamsync;NullKernel; streamsync; endloop(1);
loop(10); D2H; streamsync;NullKernel; streamsync; endloop(1);
loop(100); D2H; streamsync;NullKernel; streamsync; endloop(1);
loop(100); D2H; streamsync;NullKernel; streamsync; endloop(1);
loop(1000); D2H; streamsync;NullKernel; streamsync; endloop(1);
loop(1000); D2H; streamsync;NullKernel; streamsync; endloop(1);
loop(10000); D2H; streamsync;NullKernel; streamsync; endloop(1);
loop(10000); D2H; streamsync;NullKernel; streamsync; endloop(1);
projects/hip/samples/1_Utils/hipCommander/perf/scripts/latency_h2d.hcm
+10
Zobrazit soubor
@@ -0,0 +1,10 @@
setstream(1);
loop(10); H2D; streamsync; endloop(1);
loop(10); H2D; streamsync; endloop(1);
loop(100); H2D; streamsync; endloop(1);
loop(100); H2D; streamsync; endloop(1);
loop(1000); H2D; streamsync; endloop(1);
loop(1000); H2D; streamsync; endloop(1);
loop(1000); H2D; streamsync; endloop(1);
loop(10000); H2D; streamsync; endloop(1);
loop(10000); H2D; streamsync; endloop(1);
projects/hip/samples/1_Utils/hipCommander/perf/scripts/latency_h2d_10.hcm
+2
Zobrazit soubor
@@ -0,0 +1,2 @@
setstream(1);
loop(10); H2D; streamsync; endloop(1);
projects/hip/samples/1_Utils/hipCommander/perf/scripts/latency_h2d_d2h.hcm
+9
Zobrazit soubor
@@ -0,0 +1,9 @@
setstream(1);
loop(10); H2D; D2H; streamsync; endloop(1);
loop(10); H2D; D2H; streamsync; endloop(1);
loop(100); H2D; D2H; streamsync; endloop(1);
loop(100); H2D; D2H; streamsync; endloop(1);
loop(1000); H2D; D2H; streamsync; endloop(1);
loop(1000); H2D; D2H; streamsync; endloop(1);
loop(10000); H2D; D2H; streamsync; endloop(1);
loop(10000); H2D; D2H; streamsync; endloop(1);
projects/hip/samples/1_Utils/hipCommander/perf/scripts/latency_h2d_kernel.hcm
+9
Zobrazit soubor
@@ -0,0 +1,9 @@
setstream(1);
loop(10); H2D; NullKernel; streamsync; endloop(1);
loop(10); H2D; NullKernel; streamsync; endloop(1);
loop(100); H2D; NullKernel; streamsync; endloop(1);
loop(100); H2D; NullKernel; streamsync; endloop(1);
loop(1000); H2D; NullKernel; streamsync; endloop(1);
loop(1000); H2D; NullKernel; streamsync; endloop(1);
loop(10000); H2D; NullKernel; streamsync; endloop(1);
loop(10000); H2D; NullKernel; streamsync; endloop(1);
projects/hip/samples/1_Utils/hipCommander/perf/scripts/latency_h2d_kernel_d2h.hcm
+9
Zobrazit soubor
@@ -0,0 +1,9 @@
setstream(1);
loop(10);H2D; streamsync; NullKernel; D2H; streamsync;endloop(1);
loop(10); H2D; streamsync; NullKernel; D2H; streamsync; endloop(1);
loop(100); H2D; streamsync; NullKernel; D2H; streamsync; endloop(1);
loop(100); H2D; streamsync; NullKernel; D2H; streamsync; endloop(1);
loop(1000); H2D; streamsync; NullKernel; D2H; streamsync; endloop(1);
loop(1000); H2D; streamsync; NullKernel; D2H; streamsync; endloop(1);
loop(10000); H2D; streamsync; NullKernel; D2H; streamsync; endloop(1);
loop(10000); H2D; streamsync; NullKernel; D2H; streamsync; endloop(1);
projects/hip/samples/1_Utils/hipCommander/perf/scripts/latency_h2d_kernel_d2h_wosync.hcm
+9
Zobrazit soubor
@@ -0,0 +1,9 @@
setstream(1);
loop(10);H2D; NullKernel; D2H; streamsync;endloop(1);
loop(10); H2D; NullKernel; D2H; streamsync; endloop(1);
loop(100); H2D; NullKernel; D2H; streamsync; endloop(1);
loop(100); H2D; NullKernel; D2H; streamsync; endloop(1);
loop(1000); H2D; NullKernel; D2H; streamsync; endloop(1);
loop(1000); H2D; NullKernel; D2H; streamsync; endloop(1);
loop(10000); H2D; NullKernel; D2H; streamsync; endloop(1);
loop(10000); H2D; NullKernel; D2H; streamsync; endloop(1);
projects/hip/samples/1_Utils/hipCommander/perf/scripts/latency_h2d_kernel_wosync.hcm
+9
Zobrazit soubor
@@ -0,0 +1,9 @@
setstream(1);
loop(10); H2D; NullKernel; streamsync; endloop(1);
loop(10); H2D; NullKernel; streamsync; endloop(1);
loop(100); H2D; NullKernel; streamsync; endloop(1);
loop(100); H2D; NullKernel; streamsync; endloop(1);
loop(1000); H2D; NullKernel; streamsync; endloop(1);
loop(1000); H2D; NullKernel; streamsync; endloop(1);
loop(10000); H2D; NullKernel; streamsync; endloop(1);
loop(10000); H2D ; NullKernel; streamsync; endloop(1);
projects/hip/samples/1_Utils/hipCommander/perf/scripts/latency_h2d_sync_d2h.hcm
+9
Zobrazit soubor
@@ -0,0 +1,9 @@
setstream(1);
loop(10); H2D; streamsync; D2H; streamsync; endloop(1);
loop(10); H2D; streamsync; D2H; streamsync; endloop(1);
loop(100); H2D; streamsync; D2H; streamsync; endloop(1);
loop(100); H2D; streamsync; D2H; streamsync; endloop(1);
loop(1000); H2D; streamsync; D2H; streamsync; endloop(1);
loop(1000); H2D; streamsync; D2H; streamsync; endloop(1);
loop(10000); H2D; streamsync; D2H; streamsync; endloop(1);
loop(10000); H2D; streamsync; D2H; streamsync; endloop(1);
projects/hip/samples/1_Utils/hipCommander/perf/scripts/latency_h2d_sync_kernel_sync.hcm
+9
Zobrazit soubor
@@ -0,0 +1,9 @@
setstream(1);
loop(10); H2D; streamsync; NullKernel; streamsync; endloop(1);
loop(10); H2D; streamsync; NullKernel; streamsync; endloop(1);
loop(100); H2D; streamsync; NullKernel; streamsync; endloop(1);
loop(100); H2D; streamsync; NullKernel; streamsync; endloop(1);
loop(1000); H2D; streamsync; NullKernel; streamsync; endloop(1);
loop(1000); H2D; streamsync; NullKernel; streamsync; endloop(1);
loop(10000); H2D; streamsync; NullKernel; streamsync; endloop(1);
loop(10000); H2D; streamsync; NullKernel; streamsync; endloop(1);
projects/hip/samples/1_Utils/hipCommander/perf/scripts/latency_h2d_sync_kernel_sync_d2h.hcm
+9
Zobrazit soubor
@@ -0,0 +1,9 @@
setstream(1);
loop(10);H2D; streamsync; NullKernel;streamsync; D2H; streamsync;endloop(1);
loop(10); H2D; streamsync; NullKernel; streamsync; D2H; streamsync; endloop(1);
loop(100); H2D; streamsync; NullKernel; streamsync; D2H; streamsync; endloop(1);
loop(100); H2D; streamsync; NullKernel; streamsync; D2H; streamsync; endloop(1);
loop(1000); H2D; streamsync; NullKernel; streamsync; D2H; streamsync; endloop(1);
loop(1000); H2D; streamsync; NullKernel; streamsync; D2H; streamsync; endloop(1);
loop(10000); H2D; streamsync; NullKernel; streamsync; D2H; streamsync; endloop(1);
loop(10000); H2D; streamsync; NullKernel; streamsync; D2H; streamsync; endloop(1);
projects/hip/samples/1_Utils/hipCommander/perf/scripts/latency_kernel.hcm
+10
Zobrazit soubor
@@ -0,0 +1,10 @@
setstream(1);
loop(10); NullKernel; streamsync; endloop(1);
loop(10); NullKernel; streamsync; endloop(1);
loop(100); NullKernel; streamsync; endloop(1);
loop(100); NullKernel; streamsync; endloop(1);
loop(1000); NullKernel; streamsync; endloop(1);
loop(1000); NullKernel; streamsync; endloop(1);
loop(1000); NullKernel; streamsync; endloop(1);
loop(10000); NullKernel; streamsync; endloop(1);
loop(10000); NullKernel; streamsync; endloop(1);
projects/hip/samples/1_Utils/hipCommander/perf/scripts/latency_kernel_barrier.hcm
+9
Zobrazit soubor
@@ -0,0 +1,9 @@
setstream(1);
loop(10); NullKernel; streamsync; streamsync; endloop(1);
loop(10); NullKernel; streamsync; streamsync; endloop(1);
loop(100); NullKernel; streamsync; streamsync; endloop(1);
loop(100); NullKernel; streamsync; streamsync; endloop(1);
loop(1000); NullKernel; streamsync; streamsync; endloop(1);
loop(1000); NullKernel; streamsync; streamsync; endloop(1);
loop(10000); NullKernel; streamsync; streamsync; endloop(1);
loop(10000); NullKernel; streamsync; streamsync; endloop(1);
projects/hip/samples/1_Utils/hipCommander/perf/scripts/latency_kernel_d2h.hcm
+9
Zobrazit soubor
@@ -0,0 +1,9 @@
setstream(1);
loop(10); NullKernel; D2H; streamsync; endloop(1);
loop(10); NullKernel; D2H; streamsync; endloop(1);
loop(100); NullKernel; D2H; streamsync; endloop(1);
loop(100); NullKernel; D2H; streamsync; endloop(1);
loop(1000); NullKernel; D2H; streamsync; endloop(1);
loop(1000); NullKernel; D2H; streamsync; endloop(1);
loop(10000); NullKernel; D2H; streamsync; endloop(1);
loop(10000); NullKernel; D2H; streamsync; endloop(1);
projects/hip/samples/1_Utils/hipCommander/perf/scripts/latency_kernel_h2d.hcm
+9
Zobrazit soubor
@@ -0,0 +1,9 @@
setstream(1);
loop(10); NullKernel; H2D; streamsync; endloop(1);
loop(10); NullKernel; H2D; streamsync; endloop(1);
loop(100); NullKernel; H2D; streamsync; endloop(1);
loop(100); NullKernel; H2D; streamsync; endloop(1);
loop(1000); NullKernel; H2D; streamsync; endloop(1);
loop(1000); NullKernel; H2D; streamsync; endloop(1);
loop(10000); NullKernel; H2D; streamsync; endloop(1);
loop(10000); NullKernel; H2D; streamsync; endloop(1);

Některé soubory nejsou zobrazny, neboť je v této revizi změněno mnoho souborů Zobrazit více