Merge branch 'amd-develop' into amd-master
Change-Id: Ib81211b5b3dc734ae9051b4aac9dfd7c871d470c
[ROCm/clr commit: 83dd6b4bec]
This commit is contained in:
Maneesh Gupta
@@ -63,6 +63,7 @@ if(HIP_PLATFORM STREQUAL "hcc")
|
||||
set(HCC_HOME $ENV{HCC_HOME} CACHE PATH "Path to which HCC has been installed")
|
||||
endif()
|
||||
endif()
|
||||
|
||||
if(DEFINED ENV{HIP_DEVELOPER})
|
||||
add_to_config(_buildInfo HCC_HOME)
|
||||
endif()
|
||||
@@ -178,6 +179,7 @@ if(HIP_PLATFORM STREQUAL "hcc")
|
||||
src/hip_peer.cpp
|
||||
src/hip_stream.cpp
|
||||
src/hip_module.cpp
|
||||
src/grid_launch.cpp
|
||||
src/env.cpp)
|
||||
|
||||
set(SOURCE_FILES_DEVICE
|
||||
@@ -189,9 +191,11 @@ if(HIP_PLATFORM STREQUAL "hcc")
|
||||
|
||||
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")
|
||||
#find_package(LLVM HINTS ${HCC_HOME}/compiler/lib/cmake)
|
||||
set(CMAKE_SHARED_LINKER_FLAGS "${CMAKE_SHARED_LINKER_FLAGS} --amdgpu-target=gfx701 --amdgpu-target=gfx801 --amdgpu-target=gfx802 --amdgpu-target=gfx803 --amdgpu-target=gfx900")
|
||||
add_library(hip_hcc SHARED ${SOURCE_FILES_RUNTIME})
|
||||
target_link_libraries(hip_hcc PRIVATE hc_am)
|
||||
#target_link_libraries(hip_hcc PUBLIC LLVMAMDGPUUtils)
|
||||
add_library(hip_hcc_static STATIC ${SOURCE_FILES_RUNTIME})
|
||||
target_link_libraries(hip_hcc_static PRIVATE hc_am)
|
||||
add_dependencies(hip_hcc_static hip_hcc)
|
||||
|
||||
@@ -10,7 +10,7 @@ if [ $# = 0 ]; then
|
||||
fi
|
||||
|
||||
: ${ROCM_PATH:=/opt/rocm}
|
||||
: ${ROCM_TARGET:=fiji}
|
||||
: ${ROCM_TARGET:=gfx803}
|
||||
|
||||
INPUT_FILES=""
|
||||
OUTPUT_FILE=""
|
||||
@@ -45,16 +45,17 @@ for inputfile in $INPUT_FILES; do
|
||||
done
|
||||
printf "\nint main(){}\n" >> $hipgenisa_main
|
||||
|
||||
$HIP_PATH/bin/hipcc $hipgenisa_files -o $hipgenisa_dir/a.out
|
||||
$HIP_PATH/bin/hipcc -DGENERIC_GRID_LAUNCH=0 $hipgenisa_files -o $hipgenisa_dir/a.out
|
||||
mv dump* $hipgenisa_dir
|
||||
|
||||
hsaco_file="dump-$ROCM_TARGET.hsaco"
|
||||
map_sym=""
|
||||
kernels=$(objdump -t $hipgenisa_dir/dump-fiji.hsaco | grep grid_launch_parm | sed 's/ \+/ /g; s/\t/ /g' | cut -d" " -f6)
|
||||
kernels=$(objdump -t $hipgenisa_dir/$hsaco_file | 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)
|
||||
map_sym="--redefine-sym $mangled_sym=$real_sym $map_sym"
|
||||
done
|
||||
objcopy -F elf64-little $map_sym $hipgenisa_dir/dump-fiji.hsaco $OUTPUT_FILE
|
||||
objcopy -F elf64-little $map_sym $hipgenisa_dir/$hsaco_file $OUTPUT_FILE
|
||||
|
||||
rm $hipgenisa_files
|
||||
rm -r $hipgenisa_dir
|
||||
|
||||
@@ -92,7 +92,7 @@ if ($HIP_PLATFORM eq "hcc") {
|
||||
$HIP_ATP_MARKER=$ENV{'HIP_ATP_MARKER'} // 1;
|
||||
$marker_path = "$ROCM_PATH/profiler/CXLActivityLogger";
|
||||
|
||||
$ROCM_TARGET=$ENV{'ROCM_TARGET'} // "fiji";
|
||||
$ROCM_TARGET=$ENV{'ROCM_TARGET'} // "gfx803";
|
||||
|
||||
# HCC* may be used to compile src/hip_hcc.o (and also feed the HIPCXXFLAGS below)
|
||||
$HCC = "$HCC_HOME/bin/hcc";
|
||||
@@ -103,6 +103,7 @@ if ($HIP_PLATFORM eq "hcc") {
|
||||
|
||||
$HIPLDFLAGS = `${HCC_HOME}/bin/hcc-config --ldflags`;
|
||||
|
||||
|
||||
#### GCC system includes workaround ####
|
||||
$HCC_WA_FLAGS = " ";
|
||||
if ($HCC_VERSION_MAJOR eq 1) {
|
||||
@@ -122,12 +123,21 @@ if ($HIP_PLATFORM eq "hcc") {
|
||||
}
|
||||
}
|
||||
|
||||
# Force -stdlib=libc++ on UB14.04
|
||||
$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 "\"14.04\"") {
|
||||
$HIPCXXFLAGS .= " -stdlib=libc++";
|
||||
$setStdLib = 1;
|
||||
}
|
||||
|
||||
$HIPCXXFLAGS .= " -I$HIP_PATH/include/hip/hcc_detail/cuda";
|
||||
$HIPCXXFLAGS .= " -I$HSA_PATH/include";
|
||||
$HIPCXXFLAGS .= " -Wno-deprecated-register";
|
||||
|
||||
$HIPLDFLAGS .= " -lsupc++";
|
||||
$HIPLDFLAGS .= " -L$HSA_PATH/lib -L$ROCM_PATH/lib -lhsa-runtime64 -lhc_am -lhsakmt";
|
||||
$HIPLDFLAGS .= " -L$HSA_PATH/lib -L$ROCM_PATH/lib -lhsa-runtime64 -lhc_am -lhsakmt ";
|
||||
# $HIPLDFLAGS .= " -L$HCC_HOME/compiler/lib -lLLVMAMDGPUDesc -lLLVMAMDGPUUtils -lLLVMMC -lLLVMCore -lLLVMSupport ";
|
||||
|
||||
# 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.
|
||||
@@ -402,9 +412,9 @@ if ($setStdLib eq 0 and $HIP_PLATFORM eq 'hcc')
|
||||
|
||||
if ($needHipHcc) {
|
||||
if ($linkType eq 0) {
|
||||
substr($HIPLDFLAGS,0,0) = " $HIP_PATH/lib/libhip_hcc_static.a $HIP_PATH/lib/libhip_device.a " ;
|
||||
substr($HIPLDFLAGS,0,0) = " $HIP_PATH/lib/libhip_hcc_static.a $HIP_PATH/lib/libhip_device.a " ;
|
||||
} else {
|
||||
substr($HIPLDFLAGS,0,0) = " -Wl,--rpath=$HIP_PATH/lib $HIP_PATH/lib/libhip_hcc.so $HIP_PATH/lib/libhip_device.a ";
|
||||
substr($HIPLDFLAGS,0,0) = " -Wl,--rpath=$HIP_PATH/lib $HIP_PATH/lib/libhip_hcc.so $HIP_PATH/lib/libhip_device.a ";
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
+18
@@ -0,0 +1,18 @@
|
||||
#!/bin/bash
|
||||
|
||||
#usage : hipconvertinplace.sh [DIRNAME] [HIPIFY_OPTIONS]
|
||||
|
||||
#hipify "inplace" all code files in specified directory.
|
||||
# This can be quite handy when dealing with an existing CUDA code base since the script
|
||||
# preserves the existing directory structure.
|
||||
|
||||
# For each code file, this script will:
|
||||
# - If ".prehip file does not exist, copy the original code to a new file with extension ".prehip". Then Hipify the code file.
|
||||
# - If ".prehip" file exists, this is used as input to hipify.
|
||||
# (this is useful for testing improvements to the hipify toolset).
|
||||
|
||||
|
||||
SCRIPT_DIR=`dirname $0`
|
||||
SEARCH_DIR=$1
|
||||
shift
|
||||
$SCRIPT_DIR/hipify -inplace -print-stats "$@" `$SCRIPT_DIR/findcode.sh $SEARCH_DIR`
|
||||
@@ -1,18 +1,24 @@
|
||||
#!/bin/bash
|
||||
|
||||
#usage : hipconvertinplace.sh [DIRNAME] [HIPIFY_OPTIONS]
|
||||
#usage : hipconvertinplace.sh DIRNAME [hipify options] [--] [clang options]
|
||||
|
||||
#hipify "inplace" all code files in specified directory.
|
||||
#hipify "inplace" all code files in specified directory.
|
||||
# This can be quite handy when dealing with an existing CUDA code base since the script
|
||||
# preserves the existing directory structure.
|
||||
|
||||
# For each code file, this script will:
|
||||
# - If ".prehip file does not exist, copy the original code to a new file with extension ".prehip". Then Hipify the code file.
|
||||
# - If ".prehip" file exists, this is used as input to hipify.
|
||||
# (this is useful for testing improvements to the hipify toolset).
|
||||
|
||||
|
||||
SCRIPT_DIR=`dirname $0`
|
||||
SEARCH_DIR=$1
|
||||
shift
|
||||
$SCRIPT_DIR/hipify -inplace -print-stats "$@" `$SCRIPT_DIR/findcode.sh $SEARCH_DIR`
|
||||
|
||||
hipify_args=''
|
||||
while (( "$#" )); do
|
||||
shift
|
||||
if [ "$1" != "--" ]; then
|
||||
hipify_args="$hipify_args $1"
|
||||
else
|
||||
shift
|
||||
break
|
||||
fi
|
||||
done
|
||||
clang_args="$@"
|
||||
|
||||
$SCRIPT_DIR/hipify-clang -inplace -print-stats $hipify_args `$SCRIPT_DIR/findcode.sh $SEARCH_DIR` -- -x cuda $clang_args
|
||||
|
||||
@@ -1,24 +0,0 @@
|
||||
#!/bin/bash
|
||||
|
||||
#usage : hipconvertinplace.sh DIRNAME [hipify options] [--] [clang options]
|
||||
|
||||
#hipify "inplace" all code files in specified directory.
|
||||
# This can be quite handy when dealing with an existing CUDA code base since the script
|
||||
# preserves the existing directory structure.
|
||||
|
||||
SCRIPT_DIR=`dirname $0`
|
||||
SEARCH_DIR=$1
|
||||
|
||||
hipify_args=''
|
||||
while (( "$#" )); do
|
||||
shift
|
||||
if [ "$1" != "--" ]; then
|
||||
hipify_args="$hipify_args $1"
|
||||
else
|
||||
shift
|
||||
break
|
||||
fi
|
||||
done
|
||||
clang_args="$@"
|
||||
|
||||
$SCRIPT_DIR/hipify-clang -inplace -print-stats $hipify_args `$SCRIPT_DIR/findcode.sh $SEARCH_DIR` -- -x cuda $clang_args
|
||||
Executable
+12
@@ -0,0 +1,12 @@
|
||||
#!/bin/bash
|
||||
|
||||
#usage : hipexamine.sh DIRNAME [hipify.pl options]
|
||||
|
||||
# Generate HIP stats (LOC, CUDA->API conversions, missing functionality) for all the code files
|
||||
# in the specified directory.
|
||||
|
||||
|
||||
SCRIPT_DIR=`dirname $0`
|
||||
SEARCH_DIR=$1
|
||||
shift
|
||||
$SCRIPT_DIR/hipify -no-output -print-stats "$@" `$SCRIPT_DIR/findcode.sh $SEARCH_DIR`
|
||||
@@ -1,12 +1,22 @@
|
||||
#!/bin/bash
|
||||
|
||||
#usage : hipexamine.sh DIRNAME [hipify.pl options]
|
||||
|
||||
# Generate HIP stats (LOC, CUDA->API conversions, missing functionality) for all the code files
|
||||
# in the specified directory.
|
||||
#usage : hipexamine2.sh DIRNAME [hipify options] [--] [clang options]
|
||||
|
||||
# Generate CUDA->HIP conversion statistics for all the code files in the specified directory.
|
||||
|
||||
SCRIPT_DIR=`dirname $0`
|
||||
SEARCH_DIR=$1
|
||||
shift
|
||||
$SCRIPT_DIR/hipify -no-output -print-stats "$@" `$SCRIPT_DIR/findcode.sh $SEARCH_DIR`
|
||||
|
||||
hipify_args=''
|
||||
while (( "$#" )); do
|
||||
shift
|
||||
if [ "$1" != "--" ]; then
|
||||
hipify_args="$hipify_args $1"
|
||||
else
|
||||
shift
|
||||
break
|
||||
fi
|
||||
done
|
||||
clang_args="$@"
|
||||
|
||||
$SCRIPT_DIR/hipify-clang -examine $hipify_args `$SCRIPT_DIR/findcode.sh $SEARCH_DIR` -- -x cuda $clang_args
|
||||
|
||||
@@ -1,22 +0,0 @@
|
||||
#!/bin/bash
|
||||
|
||||
#usage : hipexamine2.sh DIRNAME [hipify options] [--] [clang options]
|
||||
|
||||
# Generate CUDA->HIP conversion statistics for all the code files in the specified directory.
|
||||
|
||||
SCRIPT_DIR=`dirname $0`
|
||||
SEARCH_DIR=$1
|
||||
|
||||
hipify_args=''
|
||||
while (( "$#" )); do
|
||||
shift
|
||||
if [ "$1" != "--" ]; then
|
||||
hipify_args="$hipify_args $1"
|
||||
else
|
||||
shift
|
||||
break
|
||||
fi
|
||||
done
|
||||
clang_args="$@"
|
||||
|
||||
$SCRIPT_DIR/hipify-clang -examine $hipify_args `$SCRIPT_DIR/findcode.sh $SEARCH_DIR` -- -x cuda $clang_args
|
||||
@@ -4,6 +4,7 @@
|
||||
|
||||
- [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)
|
||||
- [What is the current limitation of HIP Generic Grid Launch method?](#what-is-the-current-limitation-of-hip-generic-grid-launch-method)
|
||||
|
||||
<!-- tocstop -->
|
||||
|
||||
@@ -23,8 +24,6 @@ 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.
|
||||
@@ -43,4 +42,8 @@ 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);
|
||||
|
||||
```
|
||||
```
|
||||
|
||||
### What is the current limitation of HIP Generic Grid Launch method?
|
||||
1. __global__ functions cannot be marked as static or put in an unnamed namespace i.e. they cannot be given internal linkage (this would clash with __attribute__((weak)));
|
||||
2. using the macro based dispatch mechanism i.e. hipLaunchKernel* only works for functions that take no more than 20 arguments (this limit can be increased up to 126, and is temporary until we can enable C++14 mode and use variadic generic lambdas); no such limitation applies do dispatching directly through grid_launch.
|
||||
@@ -26,6 +26,7 @@
|
||||
- [How do I trace HIP application flow?](#how-do-i-trace-hip-application-flow)
|
||||
* [Using CodeXL markers for HIP Functions](#using-codexl-markers-for-hip-functions)
|
||||
* [Using HIP_TRACE_API](#using-hip_trace_api)
|
||||
- [How do I enable HIP Generic Grid Launch option?](#how-do-i-enable-hip-generic-grid-launch-option)
|
||||
|
||||
<!-- tocstop -->
|
||||
|
||||
@@ -235,3 +236,9 @@ Unlike CUDA, in HCC, for functions defined in the header files, the keyword of "
|
||||
Thus, if failed to define "static" keyword, you might see a lot of "symbol multiply defined!" errors at compilation.
|
||||
The workaround is to explicitly add the keyword of "static" before any functions that were defined as "__forceinline__".
|
||||
|
||||
### How do I disable HIP Generic Grid Launch option?
|
||||
Generic Grid Launch(GGL) is currently the default method for hip kernel launch.
|
||||
To disable it and use the legancy grid launch method, please either change the default value of GENERIC_GRID_LAUNCH to 0 in the following to header files and rebuild HIP:
|
||||
$HIP/include/hip/hcc_detail/hip_runtime_api.h
|
||||
$HIP/include/hip/hcc_detail/host_defines.h
|
||||
Or pass "-DGENERIC_GRID_LAUNCH=0" to hipcc at application compilation time.
|
||||
|
||||
@@ -680,7 +680,7 @@ The user can specify the target for which the binary can be generated. HIP/HCC d
|
||||
The file format for binary is `.co` which means Code Object. The following command builds the code object using `hipcc`.
|
||||
|
||||
`hipcc --genco --target-isa=[TARGET GPU] [INPUT FILE] -o [OUTPUT FILE]`
|
||||
```[TARGET GPU] = fiji/hawaii
|
||||
```[TARGET GPU] = gfx803/gfx701
|
||||
[INPUT FILE] = Name of the file containing kernels
|
||||
[OUTPUT FILE] = Name of the generated code object file```
|
||||
|
||||
|
||||
@@ -1,151 +1,119 @@
|
||||
# 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.
|
||||
- The Driver API defines a different but largely overlapping error code code space than the Runtime API, and uses a different coding convention. For example, Driver API defines `CUDA_ERROR_INVALID_VALUE` while the Runtime API defines `cudaErrorInvalidValue`
|
||||
- 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.
|
||||
- The Driver API defines a different but largely overlapping error code space than the Runtime API, and uses a different coding convention. For example, Driver API defines `CUDA_ERROR_INVALID_VALUE` while the Runtime API defines `cudaErrorInvalidValue`
|
||||
|
||||
|
||||
The Driver API offers two additional pieces of functionality not provided by the Runtime API: cuModule and cuCtx APIs.
|
||||
|
||||
### cuModule API
|
||||
The Module section of the Driver API provides additional control over how and when accelerator
|
||||
code objects are loaded.
|
||||
The Module section of the Driver API provides additional control over how and when accelerator code objects are loaded.
|
||||
For example, the driver API allows code objects to be loaded from files or memory pointers.
|
||||
Symbols for kernels or global data can be extracted from the loaded code objects.
|
||||
In contrast, the Runtime API automatically loads and (if necessary) compiles all of the
|
||||
kernels from a executable binary when run.
|
||||
In contrast, the Runtime API automatically loads and (if necessary) compiles all of the kernels from an executable binary when run.
|
||||
In this mode, NVCC must be used to compile kernel code so the automatic loading can function correctly.
|
||||
|
||||
Both Driver and Runtime APIs define a function for launching kernels (called `cuLaunchKernel` or `cudaLaunchKernel`.
|
||||
The kernel arguments and the execution configuration (grid dimensions, group dimensions, dynamic shared memory, and stream) are passed as arguments to the launch function.
|
||||
The Runtime additionally provides the `<<< >>>` syntax for launching kernels, which resembles
|
||||
a special function call and is easier to use than explicit launch API (in particular with respect to
|
||||
handling of kernel arguments).
|
||||
However, this syntax is not standard C++ and is available only when NVCC is used to compile the host code.
|
||||
The Runtime additionally provides the `<<< >>>` syntax for launching kernels, which resembles a special function call and is easier to use than explicit launch API (in particular with respect to handling of kernel arguments).
|
||||
However, this syntax is not standard C++ and is available only when NVCC is used to compile the host code.
|
||||
|
||||
The Module features are useful in an environment which generate the code objects directly, such as a new
|
||||
accelerator language front-end. Here, NVCC is not used. Instead, the environment may have a
|
||||
different kernel language or different compilation flow.
|
||||
The Module features are useful in an environment which generates the code objects directly, such as a new accelerator language front-end.
|
||||
Here, NVCC is not used. Instead, the environment may have a different kernel language or different compilation flow.
|
||||
Other environments have many kernels and do not want them to be all loaded automatically.
|
||||
The Module functions can be used to load the generated code objects and launch kernels.
|
||||
As we will see below, HIP defines a Module API which provides similar explicit control over code
|
||||
object management.
|
||||
As we will see below, HIP defines a Module API which provides similar explicit control over code object management.
|
||||
|
||||
### cuCtx API
|
||||
The Driver API defines "Context" and "Devices" as separate entities.
|
||||
Contexts contain a single device, and a device can theoretically have multiple contexts.
|
||||
Each context contains a set of streams and events specific to the context.
|
||||
Historically contexts also defined a unique address space for the GPU, though this may not longer be the case in Unified Memory platforms (since the CPU and all the devices in the same process share a single unified address space).
|
||||
The Context APIs also provide a mechanism to switch between devices, which allowed
|
||||
a single CPU thread to send commands to different GPUs. HIP as well as a recent versions
|
||||
of CUDA Runtime provide other mechanisms to accomplish this feat - for example using streams or
|
||||
`cudaSetDevice`.
|
||||
The Driver API defines "Context" and "Devices" as separate entities.
|
||||
Contexts contain a single device, and a device can theoretically have multiple contexts.
|
||||
Each context contains a set of streams and events specific to the context.
|
||||
Historically contexts also defined a unique address space for the GPU, though this may no longer be the case in Unified Memory platforms (since the CPU and all the devices in the same process share a single unified address space).
|
||||
The Context APIs also provide a mechanism to switch between devices, which allowed a single CPU thread to send commands to different GPUs.
|
||||
HIP as well as a recent versions of CUDA Runtime provide other mechanisms to accomplish this feat - for example using streams or `cudaSetDevice`.
|
||||
|
||||
The CUDA Runtime API unifies the Context API with the Device API. This simplifies the APIs
|
||||
and has little loss of functionality since each Context can contain a single device,
|
||||
and the benefits of multiple contexts has been replaced with other interfaces. HIP provides
|
||||
a context API to facilitate easy porting from existing Driver codes.
|
||||
In HIP, the Ctx functions largely provide an alternate syntax for changing the active device.
|
||||
The CUDA Runtime API unifies the Context API with the Device API. This simplifies the APIs and has little loss of functionality since each Context can contain a single device, and the benefits of multiple contexts has been replaced with other interfaces.
|
||||
HIP provides a context API to facilitate easy porting from existing Driver codes.
|
||||
In HIP, the Ctx functions largely provide an alternate syntax for changing the active device.
|
||||
Most new applications will prefer to use `hipSetDevice` or the stream APIs.
|
||||
|
||||
## HIP Module and Ctx APIs
|
||||
|
||||
Rather than present two separate APIs, HIP extends the HIP API with new APIs for Modules
|
||||
and Ctx control.
|
||||
Rather than present two separate APIs, HIP extends the HIP API with new APIs for Modules and Ctx control.
|
||||
|
||||
### hipModule API
|
||||
|
||||
Like the CUDA Driver API, the Module API provides additional control
|
||||
over how code is loaded, including options to load code from files or from in-memory pointers.
|
||||
NVCC and HCC target different architectures and use different code object formats : NVCC
|
||||
is `cubin` or `ptx` files, while the HCC path is the `hsaco` format. The external
|
||||
compilers which generate these code objects are responsible for generating and loading
|
||||
the correct code object for each platform. Notably, there is not a fat binary format that
|
||||
can contain code for both NVCC and HCC platforms. The following table summarizes the
|
||||
formats used on each platform:
|
||||
Like the CUDA Driver API, the Module API provides additional control over how code is loaded, including options to load code from files or from in-memory pointers.
|
||||
NVCC and HCC target different architectures and use different code object formats: NVCC is `cubin` or `ptx` files, while the HCC path is the `hsaco` format.
|
||||
The external compilers which generate these code objects are responsible for generating and loading the correct code object for each platform.
|
||||
Notably, there is not a fat binary format that can contain code for both NVCC and HCC platforms. The following table summarizes the formats used on each platform:
|
||||
|
||||
| Format | APIs | NVCC | HCC |
|
||||
| --- | --- | --- | --- |
|
||||
| Format | APIs | NVCC | HCC |
|
||||
| --- | --- | --- | --- |
|
||||
| Code Object | hipModuleLoad, hipModuleLoadData | .cubin or PTX text | .hsaco |
|
||||
| Fat Binary | hipModuleLoadFatBin | .fatbin | Under Development |
|
||||
| Fat Binary | hipModuleLoadFatBin | .fatbin | Under Development |
|
||||
|
||||
hipcc uses NVCC and HCC to compile host codes. Both of these may embed code objects
|
||||
into the final executable, and these code objects will be automatically loaded when
|
||||
the application starts.
|
||||
The hipModule API can be used to load additional code objects, and in this way
|
||||
provides an extended capability to the automatically loaded code objects. HCC allows
|
||||
both of these capabilities to be used together, if desired. Of course it is possible
|
||||
to create a program with no kernels and thus no automatic loading.
|
||||
`hipcc` uses NVCC and HCC to compile host codes. Both of these may embed code objects into the final executable, and these code objects will be automatically loaded when the application starts.
|
||||
The hipModule API can be used to load additional code objects, and in this way provides an extended capability to the automatically loaded code objects.
|
||||
HCC allows both of these capabilities to be used together, if desired. Of course it is possible to create a program with no kernels and thus no automatic loading.
|
||||
|
||||
|
||||
### hipCtx API
|
||||
HIP provides a `Ctx` API as a thin layer over the existing Device functions. This Ctx API
|
||||
can be used to set the current context, or to query properties of the device associated with
|
||||
the context. The current context is implicitly used by other APIs such as `hipStreamCreate`.
|
||||
HIP provides a `Ctx` API as a thin layer over the existing Device functions. This Ctx API can be used to set the current context, or to query properties of the device associated with the context.
|
||||
The current context is implicitly used by other APIs such as `hipStreamCreate`.
|
||||
|
||||
### hipify translation of CUDA Driver API
|
||||
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.
|
||||
[hipify support for translating driver API is Under Development]
|
||||
The hipify tool converts CUDA Driver APIs for streams, events, modules, devices, memory management, context, profiler 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.
|
||||
|
||||
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.
|
||||
The memory copy API requires 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`.
|
||||
The first flavor may be faster in some cases since they avoid host overhead to detect the
|
||||
different memory directions.
|
||||
The first flavor may be faster in some cases since they avoid host overhead to detect the different memory directions.
|
||||
|
||||
HIP defines a single error space, and uses camel-case for all errors (i.e. `hipErrorInvalidValue`).
|
||||
HIP defines a single error space, and uses camel-case for all errors (i.e. `hipErrorInvalidValue`).
|
||||
|
||||
### HCC Implementation Notes
|
||||
#### .hsaco
|
||||
The .hsaco format used by HCC is described in more detail [here](https://github.com/RadeonOpenCompute/ROCm-ComputeABI-Doc).
|
||||
An example and blog that show how to use the format is [here](http://gpuopen.com/rocm-with-harmony-combining-opencl-hcc-hsa-in-a-single-program). hsaco can be generated by hcc + extractkernel tool,
|
||||
cloc, the GCN assembler, or other tools.
|
||||
An example and blog that show how to use the format is [here](http://gpuopen.com/rocm-with-harmony-combining-opencl-hcc-hsa-in-a-single-program). hsaco can be generated by hcc + extractkernel tool, cloc, the GCN assembler, or other tools.
|
||||
|
||||
#### Address Spaces
|
||||
HCC defines a process-wide address space where the CPU and all devices allocate
|
||||
addresses from a single unified pool. Thus addresses may be shared between contexts, and
|
||||
unlike the original CUDA definition a new context does not create a new address space for
|
||||
the device.
|
||||
HCC defines a process-wide address space where the CPU and all devices allocate addresses from a single unified pool.
|
||||
Thus addresses may be shared between contexts, and unlike the original CUDA definition a new context does not create a new address space for the device.
|
||||
|
||||
#### Using hipModuleLaunchKernel
|
||||
`hipModuleLaunchKernel` is `cuLaunchKernel` in HIP world. It takes the same arguments as `cuLaunchKernel`. The argument `kernelParams` is not fully implemented for HCC. The workaround for it is, to use platform specific macros for each target. Or, `extra` argument can be used which works on both the platforms.
|
||||
|
||||
#### Additional Information
|
||||
- HCC allocates staging buffers (used for unpinned copies) on a per-device basis.
|
||||
- HCC creates a primary context when the HIP API is called. So in a pure driver API code, HIP/HCC will create a primary context while HIP/NVCC will have empty context stack. HIP/HCC will push primary context to context stack when it is empty. This can have subtle differences on applications which mix the runtime and driver APIs.
|
||||
- HCC creates a primary context when the HIP API is called. So in a pure driver API code, HIP/HCC will create a primary context while HIP/NVCC will have empty context stack.
|
||||
HIP/HCC will push primary context to context stack when it is empty. This can have subtle differences on applications which mix the runtime and driver APIs.
|
||||
|
||||
### NVCC Implementation Notes
|
||||
|
||||
#### Interoperation between HIP and CUDA Driver
|
||||
CUDA applications may want to mix CUDA driver code with HIP code (see example below). This
|
||||
table shows the type equivalence to enable this interaction.
|
||||
|
||||
|**HIP Type** |**CU Driver Type**| **CUDA Runtime Type** |
|
||||
| ---- | ---- | ---- |
|
||||
| hipModule | CUmodule | |
|
||||
| hipFunction | CUfunction | |
|
||||
| hipCtx_t | CUcontext | |
|
||||
| hipDevice_t | CUdevice | |
|
||||
| hipStream_t | CUstream | cudaStream_t |
|
||||
| hipEvent_t | CUevent | cudaEvent_t |
|
||||
| hipArray | CUarray | cudaArray |
|
||||
|
||||
CUDA applications may want to mix CUDA driver code with HIP code (see example below). This table shows the type equivalence to enable this interaction.
|
||||
|
||||
|**HIP Type** |**CU Driver Type**|**CUDA Runtime Type**|
|
||||
| ---- | ---- | ---- |
|
||||
| hipModule | CUmodule | |
|
||||
| hipFunction | CUfunction | |
|
||||
| hipCtx_t | CUcontext | |
|
||||
| hipDevice_t | CUdevice | |
|
||||
| hipStream_t | CUstream | cudaStream_t |
|
||||
| hipEvent_t | CUevent | cudaEvent_t |
|
||||
| hipArray | CUarray | cudaArray |
|
||||
|
||||
#### Compilation Flags
|
||||
The hipModule interface does not support the hipModuleLoadEx function, which is used to control PTX compilaton options.
|
||||
The hipModule interface does not support the `cuModuleLoadDataEx` function, which is used to control PTX compilation options.
|
||||
HCC does not use PTX and does not support the same compilation options.
|
||||
In fact, HCC code objects always contain fully compiled ISA and do not require additional compilation as part of the load step.
|
||||
Code which requires this functionally should use platform-specific coding, calling `cuModuleLoadEx`
|
||||
on the NVCC path and hipModuleLoad on the hcc path. For example:
|
||||
Code which requires this functionally should use platform-specific coding, calling `cuModuleLoadDataEx` on the NVCC path and `hipModuleLoadData` on the hcc path.
|
||||
For example:
|
||||
|
||||
```
|
||||
hipModule module;
|
||||
@@ -240,7 +208,7 @@ int main(){
|
||||
HIP_LAUNCH_PARAM_END
|
||||
};
|
||||
|
||||
hipModuleLaunchKernel(Function, 1, 1, 1, LEN, 1, 1, 0, 0, NULL, (void**)&con fig);
|
||||
hipModuleLaunchKernel(Function, 1, 1, 1, LEN, 1, 1, 0, 0, NULL, (void**)&config);
|
||||
|
||||
hipMemcpyDtoH(B, Bd, SIZE);
|
||||
for(uint32_t i=0;i<LEN;i++){
|
||||
@@ -253,5 +221,4 @@ int main(){
|
||||
|
||||
return 0;
|
||||
}
|
||||
|
||||
```
|
||||
|
||||
@@ -19,6 +19,9 @@ else()
|
||||
add_llvm_executable(hipify-clang src/Cuda2Hip.cpp)
|
||||
find_program(LIT_COMMAND lit)
|
||||
|
||||
set(CMAKE_CXX_COMPILER ${LLVM_TOOLS_BINARY_DIR}/clang++)
|
||||
set(CMAKE_C_COMPILER ${LLVM_TOOLS_BINARY_DIR}/clang)
|
||||
|
||||
# Link against LLVM and CLANG tools libraries
|
||||
target_link_libraries(hipify-clang
|
||||
clangASTMatchers
|
||||
|
||||
@@ -2271,7 +2271,7 @@ public:
|
||||
StringRef repName = found->second.hipName;
|
||||
DEBUG(dbgs() << "Identifier " << name
|
||||
<< " found as an actual argument in expansion of macro "
|
||||
<< macroName << "\n"
|
||||
<< MacroNameTok.getIdentifierInfo()->getName() << "\n"
|
||||
<< "will be replaced with: " << repName << "\n");
|
||||
size_t length = name.size();
|
||||
SourceLocation sl = tok.getLocation();
|
||||
|
||||
@@ -1,5 +1,5 @@
|
||||
/*
|
||||
Copyright (c) 2015-2017 Advanced Micro Devices, Inc. All rights reserved.
|
||||
Copyright (c) 2015 - present 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,7 +20,8 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
THE SOFTWARE.
|
||||
*/
|
||||
|
||||
#pragma once
|
||||
#ifndef HIP_INCLUDE_HIP_CHANNEL_DESCRIPTOR_H
|
||||
#define HIP_INCLUDE_HIP_CHANNEL_DESCRIPTOR_H
|
||||
|
||||
// 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
|
||||
@@ -34,3 +35,5 @@ THE SOFTWARE.
|
||||
#else
|
||||
#error("Must define exactly one of __HIP_PLATFORM_HCC__ or __HIP_PLATFORM_NVCC__");
|
||||
#endif
|
||||
|
||||
#endif
|
||||
|
||||
@@ -1,13 +1,16 @@
|
||||
/*
|
||||
Copyright (c) 2015-2017 Advanced Micro Devices, Inc. All rights reserved.
|
||||
Copyright (c) 2015 - present 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
|
||||
@@ -17,8 +20,8 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
THE SOFTWARE.
|
||||
*/
|
||||
|
||||
#ifndef HIP_DEVICE_FUNCTIONS_H
|
||||
#define HIP_DEVICE_FUNCTIONS_H
|
||||
#ifndef HIP_INCLUDE_HIP_DEVICE_FUNCTIONS_H
|
||||
#define HIP_INCLUDE_HIP_DEVICE_FUNCTIONS_H
|
||||
|
||||
#include<hip/hip_common.h>
|
||||
|
||||
|
||||
@@ -1,22 +1,27 @@
|
||||
/*
|
||||
Copyright (c) 2015-2017 Advanced Micro Devices, Inc. All rights reserved.
|
||||
Copyright (c) 2015 - present 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,
|
||||
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
|
||||
#ifndef HIP_INCLUDE_HIP_DRIVER_TYPES_H
|
||||
#define HIP_INCLUDE_HIP_DRIVER_TYPES_H
|
||||
|
||||
#include <hip/hip_common.h>
|
||||
|
||||
@@ -27,3 +32,5 @@ THE SOFTWARE.
|
||||
#else
|
||||
#error("Must define exactly one of __HIP_PLATFORM_HCC__ or __HIP_PLATFORM_NVCC__");
|
||||
#endif
|
||||
|
||||
#endif
|
||||
|
||||
@@ -1,5 +1,5 @@
|
||||
/*
|
||||
Copyright (c) 2015-2017 Advanced Micro Devices, Inc. All rights reserved.
|
||||
Copyright (c) 2015 - present 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_HCC_DETAIL_CHANNEL_DESCRIPTOR_H
|
||||
#define HIP_HCC_DETAIL_CHANNEL_DESCRIPTOR_H
|
||||
#ifndef HIP_INCLUDE_HIP_HCC_DETAIL_CHANNEL_DESCRIPTOR_H
|
||||
#define HIP_INCLUDE_HIP_HCC_DETAIL_CHANNEL_DESCRIPTOR_H
|
||||
|
||||
#include<hip/hcc_detail/driver_types.h>
|
||||
#include<hip/hcc_detail/hip_vector_types.h>
|
||||
|
||||
+6
-22
@@ -20,27 +20,11 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
THE SOFTWARE.
|
||||
*/
|
||||
|
||||
#ifndef HCC_ACC_H
|
||||
#define HCC_ACC_H
|
||||
#include "hip/hip_runtime_api.h"
|
||||
#pragma once
|
||||
|
||||
#if __cplusplus
|
||||
#ifdef __HCC__
|
||||
#include <hc.hpp>
|
||||
/**
|
||||
* @brief Return hc::accelerator associated with the specified deviceId
|
||||
* @return #hipSuccess, #hipErrorInvalidDevice
|
||||
*/
|
||||
hipError_t hipHccGetAccelerator(int deviceId, hc::accelerator *acc);
|
||||
namespace hip_impl // Documentation only.
|
||||
{
|
||||
#define requires(...)
|
||||
|
||||
/**
|
||||
* @brief Return hc::accelerator_view associated with the specified stream
|
||||
*
|
||||
* If stream is 0, the accelerator_view for the default stream is returned.
|
||||
* @return #hipSuccess
|
||||
*/
|
||||
hipError_t hipHccGetAcceleratorView(hipStream_t stream, hc::accelerator_view **av);
|
||||
#endif
|
||||
#endif
|
||||
|
||||
#endif
|
||||
#define FunctionalProcedure typename
|
||||
}
|
||||
@@ -1,13 +1,16 @@
|
||||
/*
|
||||
Copyright (c) 2015-2017 Advanced Micro Devices, Inc. All rights reserved.
|
||||
Copyright (c) 2015 - present 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
|
||||
@@ -17,8 +20,8 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
THE SOFTWARE.
|
||||
*/
|
||||
|
||||
#ifndef HIP_HCC_DETAIL_DEVICE_FUNCTIONS_H
|
||||
#define HIP_HCC_DETAIL_DEVICE_FUNCTIONS_H
|
||||
#ifndef HIP_INCLUDE_HIP_HCC_DETAIL_DEVICE_FUNCTIONS_H
|
||||
#define HIP_INCLUDE_HIP_HCC_DETAIL_DEVICE_FUNCTIONS_H
|
||||
|
||||
#include <hip/hip_runtime.h>
|
||||
#include <hip/hip_vector_types.h>
|
||||
|
||||
@@ -1,13 +1,16 @@
|
||||
/*
|
||||
Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved.
|
||||
Copyright (c) 2015 - present 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
|
||||
@@ -17,8 +20,8 @@ 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
|
||||
#ifndef HIP_INCLUDE_HIP_HCC_DETAIL_DRIVER_TYPES_H
|
||||
#define HIP_INCLUDE_HIP_HCC_DETAIL_DRIVER_TYPES_H
|
||||
|
||||
enum hipChannelFormatKind
|
||||
{
|
||||
|
||||
@@ -0,0 +1,853 @@
|
||||
/*
|
||||
Copyright (c) 2015 - present 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 "concepts.hpp"
|
||||
#include "helpers.hpp"
|
||||
|
||||
#include "hc.hpp"
|
||||
#include "hip_hcc.h"
|
||||
#include "hip_runtime.h"
|
||||
|
||||
#include <functional>
|
||||
#include <iostream>
|
||||
#include <stdexcept>
|
||||
#include <type_traits>
|
||||
#include <utility>
|
||||
|
||||
namespace hip_impl
|
||||
{
|
||||
namespace
|
||||
{
|
||||
struct New_grid_launch_tag {};
|
||||
struct Old_grid_launch_tag {};
|
||||
|
||||
template<typename C, typename D>
|
||||
class RAII_guard {
|
||||
D dtor_;
|
||||
public:
|
||||
RAII_guard() = default;
|
||||
|
||||
RAII_guard(const C& ctor, D dtor) : dtor_{std::move(dtor)}
|
||||
{
|
||||
ctor();
|
||||
}
|
||||
|
||||
RAII_guard(const RAII_guard&) = default;
|
||||
RAII_guard(RAII_guard&&) = default;
|
||||
|
||||
RAII_guard& operator=(const RAII_guard&) = default;
|
||||
RAII_guard& operator=(RAII_guard&&) = default;
|
||||
|
||||
~RAII_guard() { dtor_(); }
|
||||
};
|
||||
|
||||
template<typename C, typename D>
|
||||
RAII_guard<C, D> make_RAII_guard(const C& ctor, D dtor)
|
||||
{
|
||||
return RAII_guard<C, D>{ctor, std::move(dtor)};
|
||||
}
|
||||
|
||||
template<FunctionalProcedure F, typename... Ts>
|
||||
using is_new_grid_launch_t = typename std::conditional<
|
||||
std::is_callable<F(Ts...)>{},
|
||||
New_grid_launch_tag,
|
||||
Old_grid_launch_tag>::type;
|
||||
}
|
||||
|
||||
// TODO: - dispatch rank should be derived from the domain dimensions passed
|
||||
// in, and not always assumed to be 3;
|
||||
|
||||
template<FunctionalProcedure K, typename... Ts>
|
||||
requires(Domain<K> == {Ts...})
|
||||
inline
|
||||
void grid_launch_hip_impl_(
|
||||
New_grid_launch_tag,
|
||||
dim3 num_blocks,
|
||||
dim3 dim_blocks,
|
||||
int group_mem_bytes,
|
||||
const hc::accelerator_view& acc_v,
|
||||
K k,
|
||||
Ts&&... args)
|
||||
{
|
||||
const auto d = hc::extent<3>{
|
||||
num_blocks.z * dim_blocks.z,
|
||||
num_blocks.y * dim_blocks.y,
|
||||
num_blocks.x * dim_blocks.x}.tile_with_dynamic(
|
||||
dim_blocks.z,
|
||||
dim_blocks.y,
|
||||
dim_blocks.x,
|
||||
group_mem_bytes);
|
||||
|
||||
try {
|
||||
hc::parallel_for_each(
|
||||
acc_v,
|
||||
d,
|
||||
[=](const hc::tiled_index<3>& idx) [[hc]] {
|
||||
k(args...);
|
||||
});
|
||||
}
|
||||
catch (std::exception& ex) {
|
||||
std::cerr << "Failed in " << __FUNCTION__ << ", with exception: "
|
||||
<< ex.what() << std::endl;
|
||||
throw;
|
||||
}
|
||||
}
|
||||
|
||||
// TODO: these are workarounds, they should be removed.
|
||||
|
||||
hc::accelerator_view lock_stream_hip_(hipStream_t&, void*&);
|
||||
void unlock_stream_hip_(
|
||||
hipStream_t, void*, const char*, hc::accelerator_view*);
|
||||
|
||||
template<FunctionalProcedure K, typename... Ts>
|
||||
requires(Domain<K> == {Ts...})
|
||||
inline
|
||||
void grid_launch_hip_impl_(
|
||||
New_grid_launch_tag,
|
||||
dim3 num_blocks,
|
||||
dim3 dim_blocks,
|
||||
int group_mem_bytes,
|
||||
hipStream_t stream,
|
||||
const char* kernel_name,
|
||||
K k,
|
||||
Ts&&... args)
|
||||
{
|
||||
void* lck_stream = nullptr;
|
||||
auto acc_v = lock_stream_hip_(stream, lck_stream);
|
||||
auto stream_guard = make_RAII_guard(
|
||||
[](){ /* perhaps use a slimmed down ihipPrintKernelLaunch here */ },
|
||||
std::bind(
|
||||
unlock_stream_hip_, stream, lck_stream, kernel_name, &acc_v));
|
||||
|
||||
try {
|
||||
grid_launch_hip_impl_(
|
||||
New_grid_launch_tag{},
|
||||
std::move(num_blocks),
|
||||
std::move(dim_blocks),
|
||||
group_mem_bytes,
|
||||
acc_v,
|
||||
std::move(k),
|
||||
std::forward<Ts>(args)...);
|
||||
}
|
||||
catch (std::exception& ex) {
|
||||
std::cerr << "Failed in " << __FUNCTION__ << ", with exception: "
|
||||
<< ex.what() << std::endl;
|
||||
throw;
|
||||
}
|
||||
}
|
||||
|
||||
template<FunctionalProcedure K, typename... Ts>
|
||||
requires(Domain<K> == {hipLaunchParm, Ts...})
|
||||
inline
|
||||
void grid_launch_hip_impl_(
|
||||
Old_grid_launch_tag,
|
||||
dim3 num_blocks,
|
||||
dim3 dim_blocks,
|
||||
int group_mem_bytes,
|
||||
hipStream_t stream,
|
||||
K k,
|
||||
Ts&&... args)
|
||||
{
|
||||
grid_launch_hip_impl_(
|
||||
New_grid_launch_tag{},
|
||||
std::move(num_blocks),
|
||||
std::move(dim_blocks),
|
||||
group_mem_bytes,
|
||||
std::move(stream),
|
||||
std::move(k),
|
||||
hipLaunchParm{},
|
||||
std::forward<Ts>(args)...);
|
||||
}
|
||||
|
||||
template<FunctionalProcedure K, typename... Ts>
|
||||
requires(Domain<K> == {hipLaunchParm, Ts...})
|
||||
inline
|
||||
void grid_launch_hip_impl_(
|
||||
Old_grid_launch_tag,
|
||||
dim3 num_blocks,
|
||||
dim3 dim_blocks,
|
||||
int group_mem_bytes,
|
||||
hipStream_t stream,
|
||||
const char* kernel_name,
|
||||
K k,
|
||||
Ts&&... args)
|
||||
{
|
||||
grid_launch_hip_impl_(
|
||||
New_grid_launch_tag{},
|
||||
std::move(num_blocks),
|
||||
std::move(dim_blocks),
|
||||
group_mem_bytes,
|
||||
std::move(stream),
|
||||
kernel_name,
|
||||
std::move(k),
|
||||
hipLaunchParm{},
|
||||
std::forward<Ts>(args)...);
|
||||
}
|
||||
|
||||
template<FunctionalProcedure K, typename... Ts>
|
||||
requires(Domain<K> == {Ts...})
|
||||
inline
|
||||
std::enable_if_t<!std::is_function<K>::value> grid_launch_hip_(
|
||||
dim3 num_blocks,
|
||||
dim3 dim_blocks,
|
||||
int group_mem_bytes,
|
||||
hipStream_t stream,
|
||||
const char* kernel_name,
|
||||
K k,
|
||||
Ts&& ... args)
|
||||
{
|
||||
grid_launch_hip_impl_(
|
||||
is_new_grid_launch_t<K, Ts...>{},
|
||||
std::move(num_blocks),
|
||||
std::move(dim_blocks),
|
||||
group_mem_bytes,
|
||||
std::move(stream),
|
||||
kernel_name,
|
||||
std::move(k),
|
||||
std::forward<Ts>(args)...);
|
||||
}
|
||||
|
||||
template<FunctionalProcedure K, typename... Ts>
|
||||
requires(Domain<K> == {Ts...})
|
||||
inline
|
||||
std::enable_if_t<!std::is_function<K>::value> grid_launch_hip_(
|
||||
dim3 num_blocks,
|
||||
dim3 dim_blocks,
|
||||
int group_mem_bytes,
|
||||
hipStream_t stream,
|
||||
K k,
|
||||
Ts&& ... args)
|
||||
{
|
||||
grid_launch_hip_impl_(
|
||||
is_new_grid_launch_t<K, Ts...>{},
|
||||
std::move(num_blocks),
|
||||
std::move(dim_blocks),
|
||||
group_mem_bytes,
|
||||
std::move(stream),
|
||||
std::move(k),
|
||||
std::forward<Ts>(args)...);
|
||||
}
|
||||
|
||||
namespace
|
||||
{
|
||||
template<typename T>
|
||||
constexpr
|
||||
inline
|
||||
T&& forward_(std::remove_reference_t<T>& x) [[hc]]
|
||||
{
|
||||
return static_cast<T&&>(x);
|
||||
}
|
||||
|
||||
template<FunctionalProcedure K, K* k>
|
||||
struct Forwarder {
|
||||
template<typename... Ts>
|
||||
void operator()(Ts&&...args) const [[hc]]
|
||||
{
|
||||
k(forward_<Ts>(args)...);
|
||||
}
|
||||
};
|
||||
}
|
||||
|
||||
template<FunctionalProcedure K, K* k, typename... Ts>
|
||||
requires(Domain<K> == {Ts...})
|
||||
inline
|
||||
void grid_launch(
|
||||
New_grid_launch_tag,
|
||||
dim3 num_blocks,
|
||||
dim3 dim_blocks,
|
||||
int group_mem_bytes,
|
||||
hipStream_t stream,
|
||||
Ts&&... args)
|
||||
{
|
||||
grid_launch_hip_impl_(
|
||||
New_grid_launch_tag{},
|
||||
std::move(num_blocks),
|
||||
std::move(dim_blocks),
|
||||
group_mem_bytes,
|
||||
std::move(stream),
|
||||
Forwarder<K, k>{},
|
||||
std::forward<Ts>(args)...);
|
||||
}
|
||||
|
||||
template<FunctionalProcedure K, K* k, typename... Ts>
|
||||
requires(Domain<K> == {Ts...})
|
||||
inline
|
||||
void grid_launch(
|
||||
Old_grid_launch_tag,
|
||||
dim3 num_blocks,
|
||||
dim3 dim_blocks,
|
||||
int group_mem_bytes,
|
||||
hipStream_t stream,
|
||||
Ts&&... args)
|
||||
{
|
||||
grid_launch_hip_<K, k>(
|
||||
New_grid_launch_tag{},
|
||||
std::move(num_blocks),
|
||||
std::move(dim_blocks),
|
||||
group_mem_bytes,
|
||||
std::move(stream),
|
||||
hipLaunchParm{},
|
||||
std::forward<Ts>(args)...);
|
||||
}
|
||||
|
||||
template<FunctionalProcedure K, K* k, typename... Ts>
|
||||
requires(Domain<K> == {Ts...})
|
||||
inline
|
||||
std::enable_if_t<std::is_function<K>::value> grid_launch_hip_(
|
||||
dim3 num_blocks,
|
||||
dim3 dim_blocks,
|
||||
int group_mem_bytes,
|
||||
hipStream_t stream,
|
||||
Ts&&... args)
|
||||
{
|
||||
grid_launch_hip_<K, k>(
|
||||
is_new_grid_launch_t<K*, Ts...>{},
|
||||
std::move(num_blocks),
|
||||
std::move(dim_blocks),
|
||||
group_mem_bytes,
|
||||
std::move(stream),
|
||||
std::forward<Ts>(args)...);
|
||||
}
|
||||
|
||||
// TODO: these are temporary, they need to be completely removed once we
|
||||
// enable C++14 support and can have proper generic, variadic lambdas.
|
||||
#define make_kernel_lambda_hip_26(\
|
||||
kernel_name,\
|
||||
p0, p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13, p14, p15,\
|
||||
p16, p17, p18, p19, p20, p21, p22, p23, p24)\
|
||||
[](const std::decay_t<decltype(p0)>& _p0_,\
|
||||
const std::decay_t<decltype(p1)>& _p1_,\
|
||||
const std::decay_t<decltype(p2)>& _p2_,\
|
||||
const std::decay_t<decltype(p3)>& _p3_,\
|
||||
const std::decay_t<decltype(p4)>& _p4_,\
|
||||
const std::decay_t<decltype(p5)>& _p5_,\
|
||||
const std::decay_t<decltype(p6)>& _p6_,\
|
||||
const std::decay_t<decltype(p7)>& _p7_,\
|
||||
const std::decay_t<decltype(p8)>& _p8_,\
|
||||
const std::decay_t<decltype(p9)>& _p9_,\
|
||||
const std::decay_t<decltype(p10)>& _p10_,\
|
||||
const std::decay_t<decltype(p11)>& _p11_,\
|
||||
const std::decay_t<decltype(p12)>& _p12_,\
|
||||
const std::decay_t<decltype(p13)>& _p13_,\
|
||||
const std::decay_t<decltype(p14)>& _p14_,\
|
||||
const std::decay_t<decltype(p15)>& _p15_,\
|
||||
const std::decay_t<decltype(p16)>& _p16_,\
|
||||
const std::decay_t<decltype(p17)>& _p17_,\
|
||||
const std::decay_t<decltype(p18)>& _p18_,\
|
||||
const std::decay_t<decltype(p19)>& _p19_,\
|
||||
const std::decay_t<decltype(p20)>& _p20_,\
|
||||
const std::decay_t<decltype(p21)>& _p21_,\
|
||||
const std::decay_t<decltype(p22)>& _p22_,\
|
||||
const std::decay_t<decltype(p23)>& _p23_,\
|
||||
const std::decay_t<decltype(p24)>& _p24_) [[hc]] {\
|
||||
kernel_name(\
|
||||
_p0_, _p1_, _p2_, _p3_, _p4_, _p5_, _p6_, _p7_, _p8_,\
|
||||
_p9_, _p10_, _p11_, _p12_, _p13_, _p14_, _p15_, _p16_, _p17_,\
|
||||
_p18_, _p19_, _p20_, _p21_, _p22_, _p23_, _p24_);\
|
||||
}
|
||||
#define make_kernel_lambda_hip_25(\
|
||||
kernel_name,\
|
||||
p0, p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13, p14, p15,\
|
||||
p16, p17, p18, p19, p20, p21, p22, p23)\
|
||||
[](const std::decay_t<decltype(p0)>& _p0_,\
|
||||
const std::decay_t<decltype(p1)>& _p1_,\
|
||||
const std::decay_t<decltype(p2)>& _p2_,\
|
||||
const std::decay_t<decltype(p3)>& _p3_,\
|
||||
const std::decay_t<decltype(p4)>& _p4_,\
|
||||
const std::decay_t<decltype(p5)>& _p5_,\
|
||||
const std::decay_t<decltype(p6)>& _p6_,\
|
||||
const std::decay_t<decltype(p7)>& _p7_,\
|
||||
const std::decay_t<decltype(p8)>& _p8_,\
|
||||
const std::decay_t<decltype(p9)>& _p9_,\
|
||||
const std::decay_t<decltype(p10)>& _p10_,\
|
||||
const std::decay_t<decltype(p11)>& _p11_,\
|
||||
const std::decay_t<decltype(p12)>& _p12_,\
|
||||
const std::decay_t<decltype(p13)>& _p13_,\
|
||||
const std::decay_t<decltype(p14)>& _p14_,\
|
||||
const std::decay_t<decltype(p15)>& _p15_,\
|
||||
const std::decay_t<decltype(p16)>& _p16_,\
|
||||
const std::decay_t<decltype(p17)>& _p17_,\
|
||||
const std::decay_t<decltype(p18)>& _p18_,\
|
||||
const std::decay_t<decltype(p19)>& _p19_,\
|
||||
const std::decay_t<decltype(p20)>& _p20_,\
|
||||
const std::decay_t<decltype(p21)>& _p21_,\
|
||||
const std::decay_t<decltype(p22)>& _p22_,\
|
||||
const std::decay_t<decltype(p23)>& _p23_) [[hc]] {\
|
||||
kernel_name(\
|
||||
_p0_, _p1_, _p2_, _p3_, _p4_, _p5_, _p6_, _p7_, _p8_,\
|
||||
_p9_, _p10_, _p11_, _p12_, _p13_, _p14_, _p15_, _p16_, _p17_,\
|
||||
_p18_, _p19_, _p20_, _p21_, _p22_, _p23_);\
|
||||
}
|
||||
#define make_kernel_lambda_hip_24(\
|
||||
kernel_name,\
|
||||
p0, p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13, p14, p15,\
|
||||
p16, p17, p18, p19, p20, p21, p22)\
|
||||
[](const std::decay_t<decltype(p0)>& _p0_,\
|
||||
const std::decay_t<decltype(p1)>& _p1_,\
|
||||
const std::decay_t<decltype(p2)>& _p2_,\
|
||||
const std::decay_t<decltype(p3)>& _p3_,\
|
||||
const std::decay_t<decltype(p4)>& _p4_,\
|
||||
const std::decay_t<decltype(p5)>& _p5_,\
|
||||
const std::decay_t<decltype(p6)>& _p6_,\
|
||||
const std::decay_t<decltype(p7)>& _p7_,\
|
||||
const std::decay_t<decltype(p8)>& _p8_,\
|
||||
const std::decay_t<decltype(p9)>& _p9_,\
|
||||
const std::decay_t<decltype(p10)>& _p10_,\
|
||||
const std::decay_t<decltype(p11)>& _p11_,\
|
||||
const std::decay_t<decltype(p12)>& _p12_,\
|
||||
const std::decay_t<decltype(p13)>& _p13_,\
|
||||
const std::decay_t<decltype(p14)>& _p14_,\
|
||||
const std::decay_t<decltype(p15)>& _p15_,\
|
||||
const std::decay_t<decltype(p16)>& _p16_,\
|
||||
const std::decay_t<decltype(p17)>& _p17_,\
|
||||
const std::decay_t<decltype(p18)>& _p18_,\
|
||||
const std::decay_t<decltype(p19)>& _p19_,\
|
||||
const std::decay_t<decltype(p20)>& _p20_,\
|
||||
const std::decay_t<decltype(p21)>& _p21_,\
|
||||
const std::decay_t<decltype(p22)>& _p22_) [[hc]] {\
|
||||
kernel_name(\
|
||||
_p0_, _p1_, _p2_, _p3_, _p4_, _p5_, _p6_, _p7_, _p8_,\
|
||||
_p9_, _p10_, _p11_, _p12_, _p13_, _p14_, _p15_, _p16_, _p17_,\
|
||||
_p18_, _p19_, _p20_, _p21_, _p22_);\
|
||||
}
|
||||
#define make_kernel_lambda_hip_23(\
|
||||
kernel_name,\
|
||||
p0, p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13, p14, p15,\
|
||||
p16, p17, p18, p19, p20, p21)\
|
||||
[](const std::decay_t<decltype(p0)>& _p0_,\
|
||||
const std::decay_t<decltype(p1)>& _p1_,\
|
||||
const std::decay_t<decltype(p2)>& _p2_,\
|
||||
const std::decay_t<decltype(p3)>& _p3_,\
|
||||
const std::decay_t<decltype(p4)>& _p4_,\
|
||||
const std::decay_t<decltype(p5)>& _p5_,\
|
||||
const std::decay_t<decltype(p6)>& _p6_,\
|
||||
const std::decay_t<decltype(p7)>& _p7_,\
|
||||
const std::decay_t<decltype(p8)>& _p8_,\
|
||||
const std::decay_t<decltype(p9)>& _p9_,\
|
||||
const std::decay_t<decltype(p10)>& _p10_,\
|
||||
const std::decay_t<decltype(p11)>& _p11_,\
|
||||
const std::decay_t<decltype(p12)>& _p12_,\
|
||||
const std::decay_t<decltype(p13)>& _p13_,\
|
||||
const std::decay_t<decltype(p14)>& _p14_,\
|
||||
const std::decay_t<decltype(p15)>& _p15_,\
|
||||
const std::decay_t<decltype(p16)>& _p16_,\
|
||||
const std::decay_t<decltype(p17)>& _p17_,\
|
||||
const std::decay_t<decltype(p18)>& _p18_,\
|
||||
const std::decay_t<decltype(p19)>& _p19_,\
|
||||
const std::decay_t<decltype(p20)>& _p20_,\
|
||||
const std::decay_t<decltype(p21)>& _p21_) [[hc]] {\
|
||||
kernel_name(\
|
||||
_p0_, _p1_, _p2_, _p3_, _p4_, _p5_, _p6_, _p7_, _p8_,\
|
||||
_p9_, _p10_, _p11_, _p12_, _p13_, _p14_, _p15_, _p16_, _p17_,\
|
||||
_p18_, _p19_, _p20_, _p21_);\
|
||||
}
|
||||
#define make_kernel_lambda_hip_22(\
|
||||
kernel_name,\
|
||||
p0, p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13, p14, p15,\
|
||||
p16, p17, p18, p19, p20)\
|
||||
[](const std::decay_t<decltype(p0)>& _p0_,\
|
||||
const std::decay_t<decltype(p1)>& _p1_,\
|
||||
const std::decay_t<decltype(p2)>& _p2_,\
|
||||
const std::decay_t<decltype(p3)>& _p3_,\
|
||||
const std::decay_t<decltype(p4)>& _p4_,\
|
||||
const std::decay_t<decltype(p5)>& _p5_,\
|
||||
const std::decay_t<decltype(p6)>& _p6_,\
|
||||
const std::decay_t<decltype(p7)>& _p7_,\
|
||||
const std::decay_t<decltype(p8)>& _p8_,\
|
||||
const std::decay_t<decltype(p9)>& _p9_,\
|
||||
const std::decay_t<decltype(p10)>& _p10_,\
|
||||
const std::decay_t<decltype(p11)>& _p11_,\
|
||||
const std::decay_t<decltype(p12)>& _p12_,\
|
||||
const std::decay_t<decltype(p13)>& _p13_,\
|
||||
const std::decay_t<decltype(p14)>& _p14_,\
|
||||
const std::decay_t<decltype(p15)>& _p15_,\
|
||||
const std::decay_t<decltype(p16)>& _p16_,\
|
||||
const std::decay_t<decltype(p17)>& _p17_,\
|
||||
const std::decay_t<decltype(p18)>& _p18_,\
|
||||
const std::decay_t<decltype(p19)>& _p19_,\
|
||||
const std::decay_t<decltype(p20)>& _p20_) [[hc]] {\
|
||||
kernel_name(\
|
||||
_p0_, _p1_, _p2_, _p3_, _p4_, _p5_, _p6_, _p7_, _p8_,\
|
||||
_p9_, _p10_, _p11_, _p12_, _p13_, _p14_, _p15_, _p16_, _p17_,\
|
||||
_p18_, _p19_, _p20_);\
|
||||
}
|
||||
#define make_kernel_lambda_hip_21(\
|
||||
kernel_name,\
|
||||
p0, p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13, p14, p15,\
|
||||
p16, p17, p18, p19)\
|
||||
[](const std::decay_t<decltype(p0)>& _p0_,\
|
||||
const std::decay_t<decltype(p1)>& _p1_,\
|
||||
const std::decay_t<decltype(p2)>& _p2_,\
|
||||
const std::decay_t<decltype(p3)>& _p3_,\
|
||||
const std::decay_t<decltype(p4)>& _p4_,\
|
||||
const std::decay_t<decltype(p5)>& _p5_,\
|
||||
const std::decay_t<decltype(p6)>& _p6_,\
|
||||
const std::decay_t<decltype(p7)>& _p7_,\
|
||||
const std::decay_t<decltype(p8)>& _p8_,\
|
||||
const std::decay_t<decltype(p9)>& _p9_,\
|
||||
const std::decay_t<decltype(p10)>& _p10_,\
|
||||
const std::decay_t<decltype(p11)>& _p11_,\
|
||||
const std::decay_t<decltype(p12)>& _p12_,\
|
||||
const std::decay_t<decltype(p13)>& _p13_,\
|
||||
const std::decay_t<decltype(p14)>& _p14_,\
|
||||
const std::decay_t<decltype(p15)>& _p15_,\
|
||||
const std::decay_t<decltype(p16)>& _p16_,\
|
||||
const std::decay_t<decltype(p17)>& _p17_,\
|
||||
const std::decay_t<decltype(p18)>& _p18_,\
|
||||
const std::decay_t<decltype(p19)>& _p19_) [[hc]] {\
|
||||
kernel_name(\
|
||||
_p0_, _p1_, _p2_, _p3_, _p4_, _p5_, _p6_, _p7_, _p8_,\
|
||||
_p9_, _p10_, _p11_, _p12_, _p13_, _p14_, _p15_, _p16_, _p17_,\
|
||||
_p18_, _p19_);\
|
||||
}
|
||||
#define make_kernel_lambda_hip_20(\
|
||||
kernel_name,\
|
||||
p0, p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13, p14, p15,\
|
||||
p16, p17, p18)\
|
||||
[](const std::decay_t<decltype(p0)>& _p0_,\
|
||||
const std::decay_t<decltype(p1)>& _p1_,\
|
||||
const std::decay_t<decltype(p2)>& _p2_,\
|
||||
const std::decay_t<decltype(p3)>& _p3_,\
|
||||
const std::decay_t<decltype(p4)>& _p4_,\
|
||||
const std::decay_t<decltype(p5)>& _p5_,\
|
||||
const std::decay_t<decltype(p6)>& _p6_,\
|
||||
const std::decay_t<decltype(p7)>& _p7_,\
|
||||
const std::decay_t<decltype(p8)>& _p8_,\
|
||||
const std::decay_t<decltype(p9)>& _p9_,\
|
||||
const std::decay_t<decltype(p10)>& _p10_,\
|
||||
const std::decay_t<decltype(p11)>& _p11_,\
|
||||
const std::decay_t<decltype(p12)>& _p12_,\
|
||||
const std::decay_t<decltype(p13)>& _p13_,\
|
||||
const std::decay_t<decltype(p14)>& _p14_,\
|
||||
const std::decay_t<decltype(p15)>& _p15_,\
|
||||
const std::decay_t<decltype(p16)>& _p16_,\
|
||||
const std::decay_t<decltype(p17)>& _p17_,\
|
||||
const std::decay_t<decltype(p18)>& _p18_) [[hc]] {\
|
||||
kernel_name(\
|
||||
_p0_, _p1_, _p2_, _p3_, _p4_, _p5_, _p6_, _p7_, _p8_,\
|
||||
_p9_, _p10_, _p11_, _p12_, _p13_, _p14_, _p15_, _p16_, _p17_,\
|
||||
_p18_);\
|
||||
}
|
||||
#define make_kernel_lambda_hip_19(\
|
||||
kernel_name,\
|
||||
p0, p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13, p14, p15,\
|
||||
p16, p17)\
|
||||
[](const std::decay_t<decltype(p0)>& _p0_,\
|
||||
const std::decay_t<decltype(p1)>& _p1_,\
|
||||
const std::decay_t<decltype(p2)>& _p2_,\
|
||||
const std::decay_t<decltype(p3)>& _p3_,\
|
||||
const std::decay_t<decltype(p4)>& _p4_,\
|
||||
const std::decay_t<decltype(p5)>& _p5_,\
|
||||
const std::decay_t<decltype(p6)>& _p6_,\
|
||||
const std::decay_t<decltype(p7)>& _p7_,\
|
||||
const std::decay_t<decltype(p8)>& _p8_,\
|
||||
const std::decay_t<decltype(p9)>& _p9_,\
|
||||
const std::decay_t<decltype(p10)>& _p10_,\
|
||||
const std::decay_t<decltype(p11)>& _p11_,\
|
||||
const std::decay_t<decltype(p12)>& _p12_,\
|
||||
const std::decay_t<decltype(p13)>& _p13_,\
|
||||
const std::decay_t<decltype(p14)>& _p14_,\
|
||||
const std::decay_t<decltype(p15)>& _p15_,\
|
||||
const std::decay_t<decltype(p16)>& _p16_,\
|
||||
const std::decay_t<decltype(p17)>& _p17_) [[hc]] {\
|
||||
kernel_name(\
|
||||
_p0_, _p1_, _p2_, _p3_, _p4_, _p5_, _p6_, _p7_, _p8_,\
|
||||
_p9_, _p10_, _p11_, _p12_, _p13_, _p14_, _p15_, _p16_, _p17_);\
|
||||
}
|
||||
#define make_kernel_lambda_hip_18(\
|
||||
kernel_name,\
|
||||
p0, p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13, p14, p15,\
|
||||
p16)\
|
||||
[](const std::decay_t<decltype(p0)>& _p0_,\
|
||||
const std::decay_t<decltype(p1)>& _p1_,\
|
||||
const std::decay_t<decltype(p2)>& _p2_,\
|
||||
const std::decay_t<decltype(p3)>& _p3_,\
|
||||
const std::decay_t<decltype(p4)>& _p4_,\
|
||||
const std::decay_t<decltype(p5)>& _p5_,\
|
||||
const std::decay_t<decltype(p6)>& _p6_,\
|
||||
const std::decay_t<decltype(p7)>& _p7_,\
|
||||
const std::decay_t<decltype(p8)>& _p8_,\
|
||||
const std::decay_t<decltype(p9)>& _p9_,\
|
||||
const std::decay_t<decltype(p10)>& _p10_,\
|
||||
const std::decay_t<decltype(p11)>& _p11_,\
|
||||
const std::decay_t<decltype(p12)>& _p12_,\
|
||||
const std::decay_t<decltype(p13)>& _p13_,\
|
||||
const std::decay_t<decltype(p14)>& _p14_,\
|
||||
const std::decay_t<decltype(p15)>& _p15_,\
|
||||
const std::decay_t<decltype(p16)>& _p16_) [[hc]] {\
|
||||
kernel_name(\
|
||||
_p0_, _p1_, _p2_, _p3_, _p4_, _p5_, _p6_, _p7_, _p8_,\
|
||||
_p9_, _p10_, _p11_, _p12_, _p13_, _p14_, _p15_, _p16_);\
|
||||
}
|
||||
#define make_kernel_lambda_hip_17(\
|
||||
kernel_name,\
|
||||
p0, p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13, p14, p15)\
|
||||
[](const std::decay_t<decltype(p0)>& _p0_,\
|
||||
const std::decay_t<decltype(p1)>& _p1_,\
|
||||
const std::decay_t<decltype(p2)>& _p2_,\
|
||||
const std::decay_t<decltype(p3)>& _p3_,\
|
||||
const std::decay_t<decltype(p4)>& _p4_,\
|
||||
const std::decay_t<decltype(p5)>& _p5_,\
|
||||
const std::decay_t<decltype(p6)>& _p6_,\
|
||||
const std::decay_t<decltype(p7)>& _p7_,\
|
||||
const std::decay_t<decltype(p8)>& _p8_,\
|
||||
const std::decay_t<decltype(p9)>& _p9_,\
|
||||
const std::decay_t<decltype(p10)>& _p10_,\
|
||||
const std::decay_t<decltype(p11)>& _p11_,\
|
||||
const std::decay_t<decltype(p12)>& _p12_,\
|
||||
const std::decay_t<decltype(p13)>& _p13_,\
|
||||
const std::decay_t<decltype(p14)>& _p14_,\
|
||||
const std::decay_t<decltype(p15)>& _p15_) [[hc]] {\
|
||||
kernel_name(\
|
||||
_p0_, _p1_, _p2_, _p3_, _p4_, _p5_, _p6_, _p7_, _p8_,\
|
||||
_p9_, _p10_, _p11_, _p12_, _p13_, _p14_, _p15_);\
|
||||
}
|
||||
#define make_kernel_lambda_hip_16(\
|
||||
kernel_name,\
|
||||
p0, p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13, p14)\
|
||||
[](const std::decay_t<decltype(p0)>& _p0_,\
|
||||
const std::decay_t<decltype(p1)>& _p1_,\
|
||||
const std::decay_t<decltype(p2)>& _p2_,\
|
||||
const std::decay_t<decltype(p3)>& _p3_,\
|
||||
const std::decay_t<decltype(p4)>& _p4_,\
|
||||
const std::decay_t<decltype(p5)>& _p5_,\
|
||||
const std::decay_t<decltype(p6)>& _p6_,\
|
||||
const std::decay_t<decltype(p7)>& _p7_,\
|
||||
const std::decay_t<decltype(p8)>& _p8_,\
|
||||
const std::decay_t<decltype(p9)>& _p9_,\
|
||||
const std::decay_t<decltype(p10)>& _p10_,\
|
||||
const std::decay_t<decltype(p11)>& _p11_,\
|
||||
const std::decay_t<decltype(p12)>& _p12_,\
|
||||
const std::decay_t<decltype(p13)>& _p13_,\
|
||||
const std::decay_t<decltype(p14)>& _p14_) [[hc]] {\
|
||||
kernel_name(\
|
||||
_p0_, _p1_, _p2_, _p3_, _p4_, _p5_, _p6_, _p7_, _p8_,\
|
||||
_p9_, _p10_, _p11_, _p12_, _p13_, _p14_);\
|
||||
}
|
||||
#define make_kernel_lambda_hip_15(\
|
||||
kernel_name,\
|
||||
p0, p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13)\
|
||||
[](const std::decay_t<decltype(p0)>& _p0_,\
|
||||
const std::decay_t<decltype(p1)>& _p1_,\
|
||||
const std::decay_t<decltype(p2)>& _p2_,\
|
||||
const std::decay_t<decltype(p3)>& _p3_,\
|
||||
const std::decay_t<decltype(p4)>& _p4_,\
|
||||
const std::decay_t<decltype(p5)>& _p5_,\
|
||||
const std::decay_t<decltype(p6)>& _p6_,\
|
||||
const std::decay_t<decltype(p7)>& _p7_,\
|
||||
const std::decay_t<decltype(p8)>& _p8_,\
|
||||
const std::decay_t<decltype(p9)>& _p9_,\
|
||||
const std::decay_t<decltype(p10)>& _p10_,\
|
||||
const std::decay_t<decltype(p11)>& _p11_,\
|
||||
const std::decay_t<decltype(p12)>& _p12_,\
|
||||
const std::decay_t<decltype(p13)>& _p13_) [[hc]] {\
|
||||
kernel_name(\
|
||||
_p0_, _p1_, _p2_, _p3_, _p4_, _p5_, _p6_, _p7_, _p8_,\
|
||||
_p9_, _p10_, _p11_, _p12_, _p13_);\
|
||||
}
|
||||
#define make_kernel_lambda_hip_14(\
|
||||
kernel_name, p0, p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12)\
|
||||
[](const std::decay_t<decltype(p0)>& _p0_,\
|
||||
const std::decay_t<decltype(p1)>& _p1_,\
|
||||
const std::decay_t<decltype(p2)>& _p2_,\
|
||||
const std::decay_t<decltype(p3)>& _p3_,\
|
||||
const std::decay_t<decltype(p4)>& _p4_,\
|
||||
const std::decay_t<decltype(p5)>& _p5_,\
|
||||
const std::decay_t<decltype(p6)>& _p6_,\
|
||||
const std::decay_t<decltype(p7)>& _p7_,\
|
||||
const std::decay_t<decltype(p8)>& _p8_,\
|
||||
const std::decay_t<decltype(p9)>& _p9_,\
|
||||
const std::decay_t<decltype(p10)>& _p10_,\
|
||||
const std::decay_t<decltype(p11)>& _p11_,\
|
||||
const std::decay_t<decltype(p12)>& _p12_) [[hc]] {\
|
||||
kernel_name(\
|
||||
_p0_, _p1_, _p2_, _p3_, _p4_, _p5_, _p6_, _p7_, _p8_,\
|
||||
_p9_, _p10_, _p11_, _p12_);\
|
||||
}
|
||||
#define make_kernel_lambda_hip_13(\
|
||||
kernel_name, p0, p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11)\
|
||||
[](const std::decay_t<decltype(p0)>& _p0_,\
|
||||
const std::decay_t<decltype(p1)>& _p1_,\
|
||||
const std::decay_t<decltype(p2)>& _p2_,\
|
||||
const std::decay_t<decltype(p3)>& _p3_,\
|
||||
const std::decay_t<decltype(p4)>& _p4_,\
|
||||
const std::decay_t<decltype(p5)>& _p5_,\
|
||||
const std::decay_t<decltype(p6)>& _p6_,\
|
||||
const std::decay_t<decltype(p7)>& _p7_,\
|
||||
const std::decay_t<decltype(p8)>& _p8_,\
|
||||
const std::decay_t<decltype(p9)>& _p9_,\
|
||||
const std::decay_t<decltype(p10)>& _p10_,\
|
||||
const std::decay_t<decltype(p11)>& _p11_) [[hc]] {\
|
||||
kernel_name(\
|
||||
_p0_, _p1_, _p2_, _p3_, _p4_, _p5_, _p6_, _p7_, _p8_,\
|
||||
_p9_, _p10_, _p11_);\
|
||||
}
|
||||
#define make_kernel_lambda_hip_12(\
|
||||
kernel_name, p0, p1, p2, p3, p4, p5, p6, p7, p8, p9, p10)\
|
||||
[](const std::decay_t<decltype(p0)>& _p0_,\
|
||||
const std::decay_t<decltype(p1)>& _p1_,\
|
||||
const std::decay_t<decltype(p2)>& _p2_,\
|
||||
const std::decay_t<decltype(p3)>& _p3_,\
|
||||
const std::decay_t<decltype(p4)>& _p4_,\
|
||||
const std::decay_t<decltype(p5)>& _p5_,\
|
||||
const std::decay_t<decltype(p6)>& _p6_,\
|
||||
const std::decay_t<decltype(p7)>& _p7_,\
|
||||
const std::decay_t<decltype(p8)>& _p8_,\
|
||||
const std::decay_t<decltype(p9)>& _p9_,\
|
||||
const std::decay_t<decltype(p10)>& _p10_) [[hc]] {\
|
||||
kernel_name(\
|
||||
_p0_, _p1_, _p2_, _p3_, _p4_, _p5_, _p6_, _p7_, _p8_, _p9_,\
|
||||
_p10_);\
|
||||
}
|
||||
#define make_kernel_lambda_hip_11(\
|
||||
kernel_name, p0, p1, p2, p3, p4, p5, p6, p7, p8, p9)\
|
||||
[](const std::decay_t<decltype(p0)>& _p0_,\
|
||||
const std::decay_t<decltype(p1)>& _p1_,\
|
||||
const std::decay_t<decltype(p2)>& _p2_,\
|
||||
const std::decay_t<decltype(p3)>& _p3_,\
|
||||
const std::decay_t<decltype(p4)>& _p4_,\
|
||||
const std::decay_t<decltype(p5)>& _p5_,\
|
||||
const std::decay_t<decltype(p6)>& _p6_,\
|
||||
const std::decay_t<decltype(p7)>& _p7_,\
|
||||
const std::decay_t<decltype(p8)>& _p8_,\
|
||||
const std::decay_t<decltype(p9)>& _p9_) [[hc]] {\
|
||||
kernel_name(\
|
||||
_p0_, _p1_, _p2_, _p3_, _p4_, _p5_, _p6_, _p7_, _p8_, _p9_);\
|
||||
}
|
||||
#define make_kernel_lambda_hip_10(\
|
||||
kernel_name, p0, p1, p2, p3, p4, p5, p6, p7, p8)\
|
||||
[](const std::decay_t<decltype(p0)>& _p0_,\
|
||||
const std::decay_t<decltype(p1)>& _p1_,\
|
||||
const std::decay_t<decltype(p2)>& _p2_,\
|
||||
const std::decay_t<decltype(p3)>& _p3_,\
|
||||
const std::decay_t<decltype(p4)>& _p4_,\
|
||||
const std::decay_t<decltype(p5)>& _p5_,\
|
||||
const std::decay_t<decltype(p6)>& _p6_,\
|
||||
const std::decay_t<decltype(p7)>& _p7_,\
|
||||
const std::decay_t<decltype(p8)>& _p8_) [[hc]] {\
|
||||
kernel_name(_p0_, _p1_, _p2_, _p3_, _p4_, _p5_, _p6_, _p7_, _p8_);\
|
||||
}
|
||||
#define make_kernel_lambda_hip_9(\
|
||||
kernel_name, p0, p1, p2, p3, p4, p5, p6, p7)\
|
||||
[](const std::decay_t<decltype(p0)>& _p0_,\
|
||||
const std::decay_t<decltype(p1)>& _p1_,\
|
||||
const std::decay_t<decltype(p2)>& _p2_,\
|
||||
const std::decay_t<decltype(p3)>& _p3_,\
|
||||
const std::decay_t<decltype(p4)>& _p4_,\
|
||||
const std::decay_t<decltype(p5)>& _p5_,\
|
||||
const std::decay_t<decltype(p6)>& _p6_,\
|
||||
const std::decay_t<decltype(p7)>& _p7_) [[hc]] {\
|
||||
kernel_name(_p0_, _p1_, _p2_, _p3_, _p4_, _p5_, _p6_, _p7_);\
|
||||
}
|
||||
#define make_kernel_lambda_hip_8(kernel_name, p0, p1, p2, p3, p4, p5, p6)\
|
||||
[](const std::decay_t<decltype(p0)>& _p0_,\
|
||||
const std::decay_t<decltype(p1)>& _p1_,\
|
||||
const std::decay_t<decltype(p2)>& _p2_,\
|
||||
const std::decay_t<decltype(p3)>& _p3_,\
|
||||
const std::decay_t<decltype(p4)>& _p4_,\
|
||||
const std::decay_t<decltype(p5)>& _p5_,\
|
||||
const std::decay_t<decltype(p6)>& _p6_) [[hc]] {\
|
||||
kernel_name(_p0_, _p1_, _p2_, _p3_, _p4_, _p5_, _p6_);\
|
||||
}
|
||||
#define make_kernel_lambda_hip_7(kernel_name, p0, p1, p2, p3, p4, p5)\
|
||||
[](const std::decay_t<decltype(p0)>& _p0_,\
|
||||
const std::decay_t<decltype(p1)>& _p1_,\
|
||||
const std::decay_t<decltype(p2)>& _p2_,\
|
||||
const std::decay_t<decltype(p3)>& _p3_,\
|
||||
const std::decay_t<decltype(p4)>& _p4_,\
|
||||
const std::decay_t<decltype(p5)>& _p5_) [[hc]] {\
|
||||
kernel_name(_p0_, _p1_, _p2_, _p3_, _p4_, _p5_);\
|
||||
}
|
||||
#define make_kernel_lambda_hip_6(kernel_name, p0, p1, p2, p3, p4)\
|
||||
[](const std::decay_t<decltype(p0)>& _p0_,\
|
||||
const std::decay_t<decltype(p1)>& _p1_,\
|
||||
const std::decay_t<decltype(p2)>& _p2_,\
|
||||
const std::decay_t<decltype(p3)>& _p3_,\
|
||||
const std::decay_t<decltype(p4)>& _p4_) [[hc]] {\
|
||||
kernel_name(_p0_, _p1_, _p2_, _p3_, _p4_);\
|
||||
}
|
||||
#define make_kernel_lambda_hip_5(kernel_name, p0, p1, p2, p3)\
|
||||
[](const std::decay_t<decltype(p0)>& _p0_,\
|
||||
const std::decay_t<decltype(p1)>& _p1_,\
|
||||
const std::decay_t<decltype(p2)>& _p2_,\
|
||||
const std::decay_t<decltype(p3)>& _p3_) [[hc]] {\
|
||||
kernel_name(_p0_, _p1_, _p2_, _p3_);\
|
||||
}
|
||||
#define make_kernel_lambda_hip_4(kernel_name, p0, p1, p2)\
|
||||
[](const std::decay_t<decltype(p0)>& _p0_,\
|
||||
const std::decay_t<decltype(p1)>& _p1_,\
|
||||
const std::decay_t<decltype(p2)>& _p2_) [[hc]] {\
|
||||
kernel_name(_p0_, _p1_, _p2_);\
|
||||
}
|
||||
#define make_kernel_lambda_hip_3(kernel_name, p0, p1)\
|
||||
[](const std::decay_t<decltype(p0)>& _p0_,\
|
||||
const std::decay_t<decltype(p1)>& _p1_) [[hc]] {\
|
||||
kernel_name(_p0_, _p1_);\
|
||||
}
|
||||
#define make_kernel_lambda_hip_2(kernel_name, p0)\
|
||||
[](const std::decay_t<decltype(p0)>& _p0_) [[hc]] {\
|
||||
kernel_name(_p0_);\
|
||||
}
|
||||
#define make_kernel_lambda_hip_1(kernel_name)\
|
||||
[]() [[hc]] { return kernel_name(hipLaunchParm{}); }
|
||||
|
||||
#define make_kernel_lambda_hip_(...)\
|
||||
overload_macro_hip_(make_kernel_lambda_hip_, __VA_ARGS__)
|
||||
|
||||
#define hipLaunchKernelGGL(\
|
||||
kernel_name,\
|
||||
num_blocks,\
|
||||
dim_blocks,\
|
||||
group_mem_bytes,\
|
||||
stream,\
|
||||
...)\
|
||||
do {\
|
||||
hip_impl::grid_launch_hip_(\
|
||||
num_blocks,\
|
||||
dim_blocks,\
|
||||
group_mem_bytes,\
|
||||
stream,\
|
||||
#kernel_name,\
|
||||
make_kernel_lambda_hip_(kernel_name, __VA_ARGS__),\
|
||||
##__VA_ARGS__);\
|
||||
} while(0)
|
||||
|
||||
#define hipLaunchKernel(\
|
||||
kernel_name,\
|
||||
num_blocks,\
|
||||
dim_blocks,\
|
||||
group_mem_bytes,\
|
||||
stream,\
|
||||
...)\
|
||||
{\
|
||||
hipLaunchKernelGGL(\
|
||||
kernel_name,\
|
||||
num_blocks,\
|
||||
dim_blocks,\
|
||||
group_mem_bytes,\
|
||||
stream,\
|
||||
hipLaunchParm{},\
|
||||
##__VA_ARGS__);\
|
||||
}
|
||||
}
|
||||
@@ -0,0 +1,144 @@
|
||||
/*
|
||||
Copyright (c) 2015 - present 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 <type_traits> // For std::conditional, std::decay, std::enable_if,
|
||||
// std::false_type, std result_of and std::true_type.
|
||||
#include <utility> // For std::declval.
|
||||
|
||||
namespace std
|
||||
{ // TODO: these should be removed as soon as possible.
|
||||
#if (__cplusplus < 201406L)
|
||||
template<typename...>
|
||||
using void_t = void;
|
||||
|
||||
#if (__cplusplus < 201402L)
|
||||
template<bool cond, typename T = void>
|
||||
using enable_if_t = typename enable_if<cond, T>::type;
|
||||
template<bool cond, typename T, typename U>
|
||||
using conditional_t = typename conditional<cond, T, U>::type;
|
||||
template<typename T>
|
||||
using decay_t = typename decay<T>::type;
|
||||
template<FunctionalProcedure F, typename... Ts>
|
||||
using result_of_t = typename result_of<F(Ts...)>::type;
|
||||
template<typename T>
|
||||
using remove_reference_t = typename remove_reference<T>::type;
|
||||
template<
|
||||
FunctionalProcedure F,
|
||||
unsigned int n = 0u,
|
||||
typename = void>
|
||||
struct is_callable_impl : is_callable_impl<F, n + 1u> {};
|
||||
|
||||
// Pointer to member function, call through non-pointer.
|
||||
template<FunctionalProcedure F, typename C, typename... Ts>
|
||||
struct is_callable_impl<
|
||||
F(C, Ts...),
|
||||
0u,
|
||||
void_t<decltype((declval<C>().*declval<F>())(declval<Ts>()...))>
|
||||
> : true_type {
|
||||
};
|
||||
|
||||
// Pointer to member function, call through pointer.
|
||||
template<FunctionalProcedure F, typename C, typename... Ts>
|
||||
struct is_callable_impl<
|
||||
F(C, Ts...),
|
||||
1u,
|
||||
void_t<decltype(((*declval<C>()).*declval<F>())(declval<Ts>()...))>
|
||||
> : std::true_type {
|
||||
};
|
||||
|
||||
// Pointer to member data, call through non-pointer, no args.
|
||||
template<FunctionalProcedure F, typename C>
|
||||
struct is_callable_impl<
|
||||
F(C),
|
||||
2u,
|
||||
void_t<decltype(declval<C>().*declval<F>())>
|
||||
> : true_type {
|
||||
};
|
||||
|
||||
// Pointer to member data, call through pointer, no args.
|
||||
template<FunctionalProcedure F, typename C>
|
||||
struct is_callable_impl<
|
||||
F(C),
|
||||
3u,
|
||||
void_t<decltype(*declval<C>().*declval<F>())>
|
||||
> : true_type {
|
||||
};
|
||||
|
||||
// General call, n args.
|
||||
template<FunctionalProcedure F, typename... Ts>
|
||||
struct is_callable_impl<
|
||||
F(Ts...),
|
||||
4u,
|
||||
void_t<decltype(declval<F>()(declval<Ts>()...))>
|
||||
> : true_type {
|
||||
};
|
||||
|
||||
// Not callable.
|
||||
template<FunctionalProcedure F>
|
||||
struct is_callable_impl<F, 5u> : false_type {};
|
||||
|
||||
template<typename Call>
|
||||
struct is_callable : is_callable_impl<Call> {};
|
||||
#else
|
||||
template<typename, typename = void>
|
||||
struct is_callable_impl : false_type {};
|
||||
|
||||
template<FunctionalProcedure F, typename... Ts>
|
||||
struct is_callable_impl<
|
||||
F(Ts...),
|
||||
void_t<result_of_t<F(Ts...)>>> : true_type {};
|
||||
|
||||
template<typename F>
|
||||
struct is_callable : is_callable_impl<F> {};
|
||||
#endif
|
||||
template<typename...>
|
||||
struct disjunction : false_type {};
|
||||
template<typename B1>
|
||||
struct disjunction<B1> : B1 {};
|
||||
template<typename B1, typename... Bs>
|
||||
struct disjunction<B1, Bs...>
|
||||
: conditional_t<B1{} == true, B1, disjunction<Bs...>>
|
||||
{};
|
||||
#endif
|
||||
}
|
||||
|
||||
namespace hip_impl // Only for documentation, macros ignore namespaces.
|
||||
{
|
||||
#define count_macro_args_impl_hip_(\
|
||||
_0, _1, _2, _3, _4, _5, _6, _7, _8, _9, _10, _11, _12, _13, _14, _15,\
|
||||
_16, _17, _18, _19, _20, _21, _22, _23, _24, _25, _26, _27, _28, _29,\
|
||||
_30, _31, _n, ...)\
|
||||
_n
|
||||
#define count_macro_args_hip_(...)\
|
||||
count_macro_args_impl_hip_(\
|
||||
, ##__VA_ARGS__, 31, 30, 29, 28, 27, 26, 25, 24, 23, 22, 21, 20, 19,\
|
||||
18, 17, 16, 15, 14, 13, 12, 11, 10, 9, 8, 7, 6, 5, 4, 3, 2, 1, 0)
|
||||
|
||||
#define overloaded_macro_expand_hip_(macro, arg_cnt) macro##arg_cnt
|
||||
#define overload_macro_impl_hip_(macro, arg_cnt)\
|
||||
overloaded_macro_expand_hip_(macro, arg_cnt)
|
||||
#define overload_macro_hip_(macro, ...)\
|
||||
overload_macro_impl_hip_(macro, count_macro_args_hip_(__VA_ARGS__))\
|
||||
(__VA_ARGS__)
|
||||
}
|
||||
@@ -1,5 +1,5 @@
|
||||
/*
|
||||
Copyright (c) 2015-2017 Advanced Micro Devices, Inc. All rights reserved.
|
||||
Copyright (c) 2015 - present 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 INCLUDE_HIP_HCC_DETAIL_HIP_COMPLEX_H
|
||||
#define INCLUDE_HIP_HCC_DETAIL_HIP_COMPLEX_H
|
||||
#ifndef HIP_INCLUDE_HIP_HCC_DETAIL_HIP_COMPLEX_H
|
||||
#define HIP_INCLUDE_HIP_HCC_DETAIL_HIP_COMPLEX_H
|
||||
|
||||
#include "./hip_fp16.h"
|
||||
#include "./hip_vector_types.h"
|
||||
|
||||
@@ -1,5 +1,5 @@
|
||||
/*
|
||||
Copyright (c) 2015-2017 Advanced Micro Devices, Inc. All rights reserved.
|
||||
Copyright (c) 2015 - present 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_HCC_DETAIL_FP16_H
|
||||
#define HIP_HCC_DETAIL_FP16_H
|
||||
#ifndef HIP_INCLUDE_HIP_HCC_DETAIL_HIP_FP16_H
|
||||
#define HIP_INCLUDE_HIP_HCC_DETAIL_HIP_FP16_H
|
||||
|
||||
#include "hip/hip_runtime.h"
|
||||
|
||||
@@ -464,9 +464,6 @@ __device__ static inline __half2 h2trunc(const __half2 h) {
|
||||
return a;
|
||||
}
|
||||
|
||||
|
||||
|
||||
|
||||
#endif
|
||||
|
||||
#if __clang_major__ == 3
|
||||
|
||||
@@ -0,0 +1,62 @@
|
||||
/*
|
||||
Copyright (c) 2015 - present 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_INCLUDE_HIP_HCC_DETAIL_HIP_HCC_H
|
||||
#define HIP_INCLUDE_HIP_HCC_DETAIL_HIP_HCC_H
|
||||
|
||||
#include "hip/hip_runtime_api.h"
|
||||
|
||||
#if __cplusplus
|
||||
#ifdef __HCC__
|
||||
#include <hc.hpp>
|
||||
/**
|
||||
* @brief Return hc::accelerator associated with the specified deviceId
|
||||
* @return #hipSuccess, #hipErrorInvalidDevice
|
||||
*/
|
||||
hipError_t hipHccGetAccelerator(int deviceId, hc::accelerator *acc);
|
||||
|
||||
/**
|
||||
* @brief Return hc::accelerator_view associated with the specified stream
|
||||
*
|
||||
* If stream is 0, the accelerator_view for the default stream is returned.
|
||||
* @return #hipSuccess
|
||||
*/
|
||||
hipError_t hipHccGetAcceleratorView(hipStream_t stream, hc::accelerator_view **av);
|
||||
|
||||
|
||||
#endif // #ifdef __HCC__
|
||||
|
||||
hipError_t hipHccModuleLaunchKernel(hipFunction_t f,
|
||||
uint32_t globalWorkSizeX,
|
||||
uint32_t globalWorkSizeY,
|
||||
uint32_t globalWorkSizeZ,
|
||||
uint32_t localWorkSizeX,
|
||||
uint32_t localWorkSizeY,
|
||||
uint32_t localWorkSizeZ,
|
||||
size_t sharedMemBytes,
|
||||
hipStream_t hStream,
|
||||
void **kernelParams,
|
||||
void **extra);
|
||||
|
||||
#endif // #if __cplusplus
|
||||
|
||||
#endif //
|
||||
@@ -1,5 +1,5 @@
|
||||
/*
|
||||
Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved.
|
||||
Copyright (c) 2015 - present 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,15 +20,14 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
THE SOFTWARE.
|
||||
*/
|
||||
|
||||
#ifndef HIP_LDG_H
|
||||
#define HIP_LDG_H
|
||||
#ifndef HIP_INCLUDE_HIP_HCC_DETAIL_HIP_LDG_H
|
||||
#define HIP_INCLUDE_HIP_HCC_DETAIL_HIP_LDG_H
|
||||
|
||||
#if __HCC__
|
||||
#if defined __HCC__
|
||||
#if __hcc_workweek__ >= 16164
|
||||
#include "hip_vector_types.h"
|
||||
#include "host_defines.h"
|
||||
|
||||
|
||||
__device__ char __ldg(const char* );
|
||||
__device__ char2 __ldg(const char2* );
|
||||
__device__ char4 __ldg(const char4* );
|
||||
|
||||
@@ -1,5 +1,5 @@
|
||||
/*
|
||||
Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved.
|
||||
Copyright (c) 2015 - present 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
|
||||
@@ -26,13 +26,12 @@ THE SOFTWARE.
|
||||
*/
|
||||
|
||||
//#pragma once
|
||||
#ifndef HIP_HCC_DETAIL_RUNTIME_H
|
||||
#define HIP_HCC_DETAIL_RUNTIME_H
|
||||
#ifndef HIP_INCLUDE_HIP_HCC_DETAIL_HIP_RUNTIME_H
|
||||
#define HIP_INCLUDE_HIP_HCC_DETAIL_HIP_RUNTIME_H
|
||||
|
||||
//---
|
||||
// Top part of file can be compiled with any compiler
|
||||
|
||||
|
||||
//#include <cstring>
|
||||
#if __cplusplus
|
||||
#include <cmath>
|
||||
@@ -40,30 +39,53 @@ THE SOFTWARE.
|
||||
#include <math.h>
|
||||
#include <string.h>
|
||||
#include <stddef.h>
|
||||
#endif
|
||||
#endif//__cplusplus
|
||||
|
||||
// Define NVCC_COMPAT for CUDA compatibility
|
||||
#define NVCC_COMPAT
|
||||
#define CUDA_SUCCESS hipSuccess
|
||||
|
||||
#include <hip/hip_runtime_api.h>
|
||||
|
||||
//#include "hip/hcc_detail/hip_hcc.h"
|
||||
|
||||
#if USE_PROMOTE_FREE_HCC == 1
|
||||
#define ADDRESS_SPACE_1
|
||||
#define ADDRESS_SPACE_3
|
||||
#else
|
||||
#define ADDRESS_SPACE_1 __attribute__((address_space(1)))
|
||||
#define ADDRESS_SPACE_3 __attribute__((address_space(3)))
|
||||
#endif
|
||||
|
||||
//---
|
||||
// Remainder of this file only compiles with HCC
|
||||
#ifdef __HCC__
|
||||
#if defined __HCC__
|
||||
#include <grid_launch.h>
|
||||
//TODO-HCC-GL - change this to typedef.
|
||||
//typedef grid_launch_parm hipLaunchParm ;
|
||||
|
||||
#if defined (GRID_LAUNCH_VERSION) and (GRID_LAUNCH_VERSION >= 20)
|
||||
// Use field names for grid_launch 2.0 structure, if HCC supports GL 2.0.
|
||||
#if GENERIC_GRID_LAUNCH == 0
|
||||
#define hipLaunchParm grid_launch_parm
|
||||
#else
|
||||
namespace hip_impl
|
||||
{
|
||||
struct Empty_launch_parm {};
|
||||
}
|
||||
#define hipLaunchParm hip_impl::Empty_launch_parm
|
||||
#endif //GENERIC_GRID_LAUNCH
|
||||
|
||||
#if defined (GRID_LAUNCH_VERSION) and (GRID_LAUNCH_VERSION >= 20) || GENERIC_GRID_LAUNCH == 1
|
||||
#else // Use field names for grid_launch 2.0 structure, if HCC supports GL 2.0.
|
||||
#error (HCC must support GRID_LAUNCH_20)
|
||||
#endif
|
||||
#endif //GRID_LAUNCH_VERSION
|
||||
|
||||
#endif //HCC
|
||||
|
||||
#if GENERIC_GRID_LAUNCH==1 && defined __HCC__
|
||||
#include "grid_launch_GGL.hpp"
|
||||
#endif//GENERIC_GRID_LAUNCH
|
||||
|
||||
extern int HIP_TRACE_API;
|
||||
|
||||
//TODO-HCC-GL - change this to typedef.
|
||||
//typedef grid_launch_parm hipLaunchParm ;
|
||||
#define hipLaunchParm grid_launch_parm
|
||||
#ifdef __cplusplus
|
||||
//#include <hip/hcc_detail/hip_texture.h>
|
||||
#include <hip/hcc_detail/hip_ldg.h>
|
||||
@@ -77,6 +99,17 @@ extern int HIP_TRACE_API;
|
||||
#define __HCC_ACCELERATOR__ __KALMAR_ACCELERATOR__
|
||||
#endif
|
||||
|
||||
|
||||
|
||||
|
||||
// TODO-HCC add a dummy implementation of assert, need to replace with a proper kernel exit call.
|
||||
#if __HIP_DEVICE_COMPILE__ == 1
|
||||
#undef assert
|
||||
#define assert(COND) { if (COND) {} }
|
||||
#endif
|
||||
|
||||
|
||||
|
||||
// Feature tests:
|
||||
#if defined(__HCC_ACCELERATOR__) && (__HCC_ACCELERATOR__ != 0)
|
||||
// Device compile and not host compile:
|
||||
@@ -126,8 +159,7 @@ extern int HIP_TRACE_API;
|
||||
|
||||
// TODO - hipify-clang - change to use the function call.
|
||||
//#define warpSize hc::__wavesize()
|
||||
extern const int warpSize;
|
||||
|
||||
static constexpr int warpSize = 64;
|
||||
|
||||
#define clock_t long long int
|
||||
__device__ long long int clock64();
|
||||
@@ -247,7 +279,7 @@ __device__ float __shfl(float input, int lane, int width);
|
||||
__device__ float __shfl_up(float input, unsigned int lane_delta, int width);
|
||||
__device__ float __shfl_down(float input, unsigned int lane_delta, int width);
|
||||
__device__ float __shfl_xor(float input, int lane_mask, int width);
|
||||
#endif
|
||||
#endif //__cplusplus
|
||||
|
||||
__device__ unsigned __hip_ds_bpermute(int index, unsigned src);
|
||||
__device__ float __hip_ds_bpermutef(int index, float src);
|
||||
@@ -259,12 +291,12 @@ __device__ float __hip_ds_swizzlef(float src, int pattern);
|
||||
|
||||
__device__ int __hip_move_dpp(int src, int dpp_ctrl, int row_mask, int bank_mask, bool bound_ctrl);
|
||||
|
||||
#endif
|
||||
#endif //__HIP_ARCH_GFX803__ == 1
|
||||
|
||||
__host__ __device__ int min(int arg1, int arg2);
|
||||
__host__ __device__ int max(int arg1, int arg2);
|
||||
|
||||
__device__ __attribute__((address_space(3))) void* __get_dynamicgroupbaseptr();
|
||||
__device__ ADDRESS_SPACE_3 void* __get_dynamicgroupbaseptr();
|
||||
|
||||
|
||||
/**
|
||||
@@ -387,17 +419,18 @@ static inline __device__ void* memset(void* ptr, int val, size_t size)
|
||||
|
||||
#define __syncthreads() hc_barrier(CLK_LOCAL_MEM_FENCE)
|
||||
|
||||
#define HIP_KERNEL_NAME(...) __VA_ARGS__
|
||||
#define HIP_KERNEL_NAME(...) (__VA_ARGS__)
|
||||
#define HIP_SYMBOL(X) #X
|
||||
|
||||
#ifdef __HCC_CPP__
|
||||
#if defined __HCC_CPP__
|
||||
extern hipStream_t ihipPreLaunchKernel(hipStream_t stream, dim3 grid, dim3 block, grid_launch_parm *lp, const char *kernelNameStr);
|
||||
extern hipStream_t ihipPreLaunchKernel(hipStream_t stream, dim3 grid, size_t block, grid_launch_parm *lp, const char *kernelNameStr);
|
||||
extern hipStream_t ihipPreLaunchKernel(hipStream_t stream, size_t grid, dim3 block, grid_launch_parm *lp, const char *kernelNameStr);
|
||||
extern hipStream_t ihipPreLaunchKernel(hipStream_t stream, size_t grid, size_t block, grid_launch_parm *lp, const char *kernelNameStr);
|
||||
extern void ihipPostLaunchKernel(const char *kernelName, hipStream_t stream, grid_launch_parm &lp);
|
||||
|
||||
|
||||
#if GENERIC_GRID_LAUNCH == 0
|
||||
//#warning "Original hipLaunchKernel defined"
|
||||
// Due to multiple overloaded versions of ihipPreLaunchKernel, the numBlocks3D and blockDim3D can be either size_t or dim3 types
|
||||
#define hipLaunchKernel(_kernelName, _numBlocks3D, _blockDim3D, _groupMemBytes, _stream, ...) \
|
||||
do {\
|
||||
@@ -407,13 +440,13 @@ do {\
|
||||
_kernelName (lp, ##__VA_ARGS__);\
|
||||
ihipPostLaunchKernel(#_kernelName, trueStream, lp);\
|
||||
} while(0)
|
||||
|
||||
#endif //GENERIC_GRID_LAUNCH
|
||||
|
||||
#elif defined (__HCC_C__)
|
||||
|
||||
//TODO - develop C interface.
|
||||
|
||||
#endif
|
||||
#endif //__HCC_CPP__
|
||||
|
||||
/**
|
||||
* extern __shared__
|
||||
@@ -422,12 +455,11 @@ do {\
|
||||
// Macro to replace extern __shared__ declarations
|
||||
// to local variable definitions
|
||||
#define HIP_DYNAMIC_SHARED(type, var) \
|
||||
__attribute__((address_space(3))) type* var = \
|
||||
(__attribute__((address_space(3))) type*)__get_dynamicgroupbaseptr(); \
|
||||
ADDRESS_SPACE_3 type* var = \
|
||||
(ADDRESS_SPACE_3 type*)__get_dynamicgroupbaseptr(); \
|
||||
|
||||
#define HIP_DYNAMIC_SHARED_ATTRIBUTE __attribute__((address_space(3)))
|
||||
#define HIP_DYNAMIC_SHARED_ATTRIBUTE ADDRESS_SPACE_3
|
||||
|
||||
#endif // __HCC__
|
||||
|
||||
|
||||
/**
|
||||
@@ -451,4 +483,4 @@ do {\
|
||||
|
||||
|
||||
|
||||
#endif
|
||||
#endif//HIP_HCC_DETAIL_RUNTIME_H
|
||||
|
||||
@@ -1,5 +1,5 @@
|
||||
/*
|
||||
Copyright (c) 2015-2017 Advanced Micro Devices, Inc. All rights reserved.
|
||||
Copyright (c) 2015 - present 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,17 +21,20 @@ THE SOFTWARE.
|
||||
*/
|
||||
|
||||
//#pragma once
|
||||
#ifndef HIP_HCC_DETAIL_HIP_RUNTIME_API_H
|
||||
#define HIP_HCC_DETAIL_HIP_RUNTIME_API_H
|
||||
#ifndef HIP_INCLUDE_HIP_HCC_DETAIL_HIP_RUNTIME_API_H
|
||||
#define HIP_INCLUDE_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.
|
||||
*/
|
||||
|
||||
#include <stdint.h>
|
||||
#include <stddef.h>
|
||||
#include <iostream>
|
||||
|
||||
#ifndef GENERIC_GRID_LAUNCH
|
||||
#define GENERIC_GRID_LAUNCH 1
|
||||
#endif
|
||||
|
||||
#include <hip/hcc_detail/host_defines.h>
|
||||
#include <hip/hip_runtime_api.h>
|
||||
#include <hip/hip_texture.h>
|
||||
@@ -1153,7 +1156,7 @@ hipError_t hipMemcpyDtoDAsync(hipDeviceptr_t dst, hipDeviceptr_t src, size_t siz
|
||||
*
|
||||
* @see hipMemcpy, hipMemcpy2D, hipMemcpyToArray, hipMemcpy2DToArray, hipMemcpyFromArray, hipMemcpy2DFromArray, hipMemcpyArrayToArray, hipMemcpy2DArrayToArray, hipMemcpyFromSymbol, hipMemcpyAsync, hipMemcpy2DAsync, hipMemcpyToArrayAsync, hipMemcpy2DToArrayAsync, hipMemcpyFromArrayAsync, hipMemcpy2DFromArrayAsync, hipMemcpyToSymbolAsync, hipMemcpyFromSymbolAsync
|
||||
*/
|
||||
hipError_t hipMemcpyToSymbol(const char* symbolName, const void *src, size_t sizeBytes, size_t offset, hipMemcpyKind kind);
|
||||
hipError_t hipMemcpyToSymbol(const void* symbolName, const void *src, size_t sizeBytes, size_t offset, hipMemcpyKind kind);
|
||||
|
||||
|
||||
/**
|
||||
@@ -1173,11 +1176,11 @@ hipError_t hipMemcpyToSymbol(const char* symbolName, const void *src, size_t siz
|
||||
*
|
||||
* @see hipMemcpy, hipMemcpy2D, hipMemcpyToArray, hipMemcpy2DToArray, hipMemcpyFromArray, hipMemcpy2DFromArray, hipMemcpyArrayToArray, hipMemcpy2DArrayToArray, hipMemcpyFromSymbol, hipMemcpyAsync, hipMemcpy2DAsync, hipMemcpyToArrayAsync, hipMemcpy2DToArrayAsync, hipMemcpyFromArrayAsync, hipMemcpy2DFromArrayAsync, hipMemcpyToSymbolAsync, hipMemcpyFromSymbolAsync
|
||||
*/
|
||||
hipError_t hipMemcpyToSymbolAsync(const char* symbolName, const void *src, size_t sizeBytes, size_t offset, hipMemcpyKind kind, hipStream_t stream);
|
||||
hipError_t hipMemcpyToSymbolAsync(const void* symbolName, const void *src, size_t sizeBytes, size_t offset, hipMemcpyKind kind, hipStream_t stream);
|
||||
|
||||
hipError_t hipMemcpyFromSymbol(void *dst, const char* symbolName, size_t sizeBytes, size_t offset, hipMemcpyKind kind);
|
||||
hipError_t hipMemcpyFromSymbol(void *dst, const void* symbolName, size_t sizeBytes, size_t offset, hipMemcpyKind kind);
|
||||
|
||||
hipError_t hipMemcpyFromSymbolAsync(void *dst, const char* symbolName, size_t sizeBytes, size_t offset, hipMemcpyKind kind, hipStream_t stream);
|
||||
hipError_t hipMemcpyFromSymbolAsync(void *dst, const void* symbolName, size_t sizeBytes, size_t offset, hipMemcpyKind kind, hipStream_t stream);
|
||||
|
||||
/**
|
||||
* @brief Copy data from src to dst asynchronously.
|
||||
@@ -1693,6 +1696,66 @@ hipError_t hipCtxEnablePeerAccess (hipCtx_t peerCtx, unsigned int flags);
|
||||
*/
|
||||
hipError_t hipCtxDisablePeerAccess (hipCtx_t peerCtx);
|
||||
|
||||
/**
|
||||
* @brief Get the state of the primary context.
|
||||
*
|
||||
* @param [in] Device to get primary context flags for
|
||||
* @param [out] Pointer to store flags
|
||||
* @param [out] Pointer to store context state; 0 = inactive, 1 = active
|
||||
*
|
||||
* @returns #hipSuccess
|
||||
*
|
||||
* @see hipCtxCreate, hipCtxDestroy, hipCtxGetFlags, hipCtxPopCurrent, hipCtxGetCurrent, hipCtxSetCurrent, hipCtxPushCurrent, hipCtxSetCacheConfig, hipCtxSynchronize, hipCtxGetDevice
|
||||
*/
|
||||
hipError_t hipDevicePrimaryCtxGetState ( hipDevice_t dev, unsigned int* flags, int* active );
|
||||
|
||||
/**
|
||||
* @brief Release the primary context on the GPU.
|
||||
*
|
||||
* @param [in] Device which primary context is released
|
||||
*
|
||||
* @returns #hipSuccess
|
||||
*
|
||||
* @see hipCtxCreate, hipCtxDestroy, hipCtxGetFlags, hipCtxPopCurrent, hipCtxGetCurrent, hipCtxSetCurrent, hipCtxPushCurrent, hipCtxSetCacheConfig, hipCtxSynchronize, hipCtxGetDevice
|
||||
* @warning This function return #hipSuccess though doesn't release the primaryCtx by design on HIP/HCC path.
|
||||
*/
|
||||
hipError_t hipDevicePrimaryCtxRelease ( hipDevice_t dev);
|
||||
|
||||
/**
|
||||
* @brief Retain the primary context on the GPU.
|
||||
*
|
||||
* @param [out] Returned context handle of the new context
|
||||
* @param [in] Device which primary context is released
|
||||
*
|
||||
* @returns #hipSuccess
|
||||
*
|
||||
* @see hipCtxCreate, hipCtxDestroy, hipCtxGetFlags, hipCtxPopCurrent, hipCtxGetCurrent, hipCtxSetCurrent, hipCtxPushCurrent, hipCtxSetCacheConfig, hipCtxSynchronize, hipCtxGetDevice
|
||||
*/
|
||||
hipError_t hipDevicePrimaryCtxRetain ( hipCtx_t* pctx, hipDevice_t dev );
|
||||
|
||||
/**
|
||||
* @brief Resets the primary context on the GPU.
|
||||
*
|
||||
* @param [in] Device which primary context is reset
|
||||
*
|
||||
* @returns #hipSuccess
|
||||
*
|
||||
* @see hipCtxCreate, hipCtxDestroy, hipCtxGetFlags, hipCtxPopCurrent, hipCtxGetCurrent, hipCtxSetCurrent, hipCtxPushCurrent, hipCtxSetCacheConfig, hipCtxSynchronize, hipCtxGetDevice
|
||||
*/
|
||||
hipError_t hipDevicePrimaryCtxReset ( hipDevice_t dev );
|
||||
|
||||
/**
|
||||
* @brief Set flags for the primary context.
|
||||
*
|
||||
* @param [in] Device for which the primary context flags are set
|
||||
* @param [in] New flags for the device
|
||||
*
|
||||
* @returns #hipSuccess, #hipErrorContextAlreadyInUse
|
||||
*
|
||||
* @see hipCtxCreate, hipCtxDestroy, hipCtxGetFlags, hipCtxPopCurrent, hipCtxGetCurrent, hipCtxSetCurrent, hipCtxPushCurrent, hipCtxSetCacheConfig, hipCtxSynchronize, hipCtxGetDevice
|
||||
*/
|
||||
hipError_t hipDevicePrimaryCtxSetFlags ( hipDevice_t dev, unsigned int flags );
|
||||
|
||||
// doxygen end Context Management
|
||||
/**
|
||||
* @}
|
||||
|
||||
@@ -1,5 +1,5 @@
|
||||
/*
|
||||
Copyright (c) 2015-2017 Advanced Micro Devices, Inc. All rights reserved.
|
||||
Copyright (c) 2015 - present 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
|
||||
@@ -22,8 +22,8 @@ THE SOFTWARE.
|
||||
|
||||
//#pragma once
|
||||
|
||||
#ifndef HIP_HCC_DETAIL_TEXTURE_H
|
||||
#define HIP_HCC_DETAIL_TEXTURE_H
|
||||
#ifndef HIP_INCLUDE_HIP_HCC_DETAIL_HIP_TEXTURE_H
|
||||
#define HIP_INCLUDE_HIP_HCC_DETAIL_HIP_TEXTURE_H
|
||||
|
||||
/**
|
||||
* @file hcc_detail/hip_texture.h
|
||||
|
||||
@@ -1,5 +1,5 @@
|
||||
/*
|
||||
Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved.
|
||||
Copyright (c) 2015 - present 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
|
||||
@@ -25,8 +25,8 @@ THE SOFTWARE.
|
||||
* @brief Defines the different newt vector types for HIP runtime.
|
||||
*/
|
||||
|
||||
#ifndef HIP_VECTOR_TYPES_H
|
||||
#define HIP_VECTOR_TYPES_H
|
||||
#ifndef HIP_INCLUDE_HIP_HCC_DETAIL_HIP_VECTOR_TYPES_H
|
||||
#define HIP_INCLUDE_HIP_HCC_DETAIL_HIP_VECTOR_TYPES_H
|
||||
|
||||
#if defined (__HCC__) && (__hcc_workweek__ < 16032)
|
||||
#error("This version of HIP requires a newer version of HCC.");
|
||||
|
||||
@@ -1,5 +1,5 @@
|
||||
/*
|
||||
Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved.
|
||||
Copyright (c) 2015 - present 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
|
||||
@@ -25,8 +25,15 @@ THE SOFTWARE.
|
||||
* @brief TODO-doc
|
||||
*/
|
||||
|
||||
#ifndef HOST_DEFINES_H
|
||||
#define HOST_DEFINES_H
|
||||
#ifndef HIP_INCLUDE_HIP_HCC_DETAIL_HOST_DEFINES_H
|
||||
#define HIP_INCLUDE_HIP_HCC_DETAIL_HOST_DEFINES_H
|
||||
|
||||
#define USE_PROMOTE_FREE_HCC 1
|
||||
|
||||
// Add guard to Generic Grid Launch method
|
||||
#ifndef GENERIC_GRID_LAUNCH
|
||||
#define GENERIC_GRID_LAUNCH 1
|
||||
#endif
|
||||
|
||||
#ifdef __HCC__
|
||||
/**
|
||||
@@ -35,7 +42,14 @@ THE SOFTWARE.
|
||||
#define __host__ __attribute__((cpu))
|
||||
#define __device__ __attribute__((hc))
|
||||
|
||||
//#warning "HOST DEFINE header included"
|
||||
#if GENERIC_GRID_LAUNCH == 0
|
||||
//#warning "original global define reached"
|
||||
#define __global__ __attribute__((hc_grid_launch)) __attribute__((used))
|
||||
#else
|
||||
//#warning "GGL global define reached"
|
||||
#define __global__ [[hc]] __attribute__((weak))
|
||||
#endif //GENERIC_GRID_LAUNCH
|
||||
|
||||
#define __noinline__ __attribute__((noinline))
|
||||
#define __forceinline__ __attribute__((always_inline))
|
||||
@@ -47,7 +61,11 @@ THE SOFTWARE.
|
||||
*/
|
||||
// _restrict is supported by the compiler
|
||||
#define __shared__ tile_static
|
||||
#define __constant__ __attribute__((address_space(1)))
|
||||
#if USE_PROMOTE_FREE_HCC==1
|
||||
#define __constant__ __attribute__((hc))
|
||||
#else
|
||||
#define __constant__ ADDRESS_SPACE_1
|
||||
#endif
|
||||
|
||||
#else
|
||||
// Non-HCC compiler
|
||||
|
||||
@@ -1,13 +1,16 @@
|
||||
/*
|
||||
Copyright (c) 2015-2017 Advanced Micro Devices, Inc. All rights reserved.
|
||||
Copyright (c) 2015 - present 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
|
||||
@@ -17,8 +20,8 @@ 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
|
||||
#ifndef HIP_INCLUDE_HIP_HCC_DETAIL_MATH_FUNCTIONS_H
|
||||
#define HIP_INCLUDE_HIP_HCC_DETAIL_MATH_FUNCTIONS_H
|
||||
|
||||
#include <hip/hip_runtime.h>
|
||||
#include <hip/hip_vector_types.h>
|
||||
@@ -66,7 +69,7 @@ __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__ float lgammaf(float x);
|
||||
__device__ long long int llrintf(float x);
|
||||
__device__ long long int llroundf(float x);
|
||||
__device__ float log10f(float x);
|
||||
|
||||
@@ -1,13 +1,16 @@
|
||||
/*
|
||||
Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved.
|
||||
Copyright (c) 2015 - present 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
|
||||
@@ -18,8 +21,8 @@ THE SOFTWARE.
|
||||
*/
|
||||
|
||||
|
||||
#ifndef HIP_HCC_DETAIL_TEXTURE_TYPES_H
|
||||
#define HIP_HCC_DETAIL_TEXTURE_TYPES_H
|
||||
#ifndef HIP_INCLUDE_HIP_HCC_DETAIL_TEXTURE_TYPES_H
|
||||
#define HIP_INCLUDE_HIP_HCC_DETAIL_TEXTURE_TYPES_H
|
||||
|
||||
#include<hip/hcc_detail/driver_types.h>
|
||||
|
||||
|
||||
@@ -1,5 +1,5 @@
|
||||
/*
|
||||
Copyright (c) 2015-2017 Advanced Micro Devices, Inc. All rights reserved.
|
||||
Copyright (c) 2015 - present 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,7 +20,8 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
THE SOFTWARE.
|
||||
*/
|
||||
|
||||
#pragma once
|
||||
#ifndef HIP_INCLUDE_HIP_HIP_COMMON_H
|
||||
#define HIP_INCLUDE_HIP_HIP_COMMON_H
|
||||
|
||||
// Common code included at start of every hip file.
|
||||
// Auto enable __HIP_PLATFORM_HCC__ if compiling with HCC
|
||||
@@ -73,3 +74,5 @@ THE SOFTWARE.
|
||||
#define __HIP_ARCH_HAS_3DGRID__ (0)
|
||||
#define __HIP_ARCH_HAS_DYNAMIC_PARALLEL__ (0)
|
||||
#endif
|
||||
|
||||
#endif
|
||||
|
||||
@@ -1,5 +1,5 @@
|
||||
/*
|
||||
Copyright (c) 2015-2017 Advanced Micro Devices, Inc. All rights reserved.
|
||||
Copyright (c) 2015 - present 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,7 +20,8 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
THE SOFTWARE.
|
||||
*/
|
||||
|
||||
#pragma once
|
||||
#ifndef HIP_INCLUDE_HIP_HIP_COMPLEX_H
|
||||
#define HIP_INCLUDE_HIP_HIP_COMPLEX_H
|
||||
|
||||
#include <hip/hip_common.h>
|
||||
|
||||
@@ -31,3 +32,5 @@ THE SOFTWARE.
|
||||
#else
|
||||
#error("Must define exactly one of __HIP_PLATFORM_HCC__ or __HIP_PLATFORM_NVCC__");
|
||||
#endif
|
||||
|
||||
#endif
|
||||
|
||||
@@ -1,5 +1,5 @@
|
||||
/*
|
||||
Copyright (c) 2015-2017 Advanced Micro Devices, Inc. All rights reserved.
|
||||
Copyright (c) 2015 - present 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,7 +20,8 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
THE SOFTWARE.
|
||||
*/
|
||||
|
||||
#pragma once
|
||||
#ifdef HIP_INCLUDE_HIP_HIP_FP16_H
|
||||
#define HIP_INCLUDE_HIP_HIP_FP16_H
|
||||
|
||||
#include <hip/hip_common.h>
|
||||
|
||||
@@ -31,3 +32,5 @@ THE SOFTWARE.
|
||||
#else
|
||||
#error("Must define exactly one of __HIP_PLATFORM_HCC__ or __HIP_PLATFORM_NVCC__");
|
||||
#endif
|
||||
|
||||
#endif
|
||||
|
||||
+4
-4
@@ -1,5 +1,5 @@
|
||||
/*
|
||||
Copyright (c) 2015-2017 Advanced Micro Devices, Inc. All rights reserved.
|
||||
Copyright (c) 2015 - present 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,11 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
THE SOFTWARE.
|
||||
*/
|
||||
|
||||
#ifndef HIP_HCC_H
|
||||
#define HIP_HCC_H
|
||||
#ifndef HIP_INCLUDE_HIP_HIP_HCC_H
|
||||
#define HIP_INCLUDE_HIP_HIP_HCC_H
|
||||
|
||||
#if defined(__HIP_PLATFORM_HCC__) && !defined (__HIP_PLATFORM_NVCC__)
|
||||
#include "hip/hcc_detail/hcc_acc.h"
|
||||
#include "hip/hcc_detail/hip_hcc.h"
|
||||
#endif
|
||||
|
||||
#endif
|
||||
@@ -1,5 +1,5 @@
|
||||
/*
|
||||
Copyright (c) 2015-2017 Advanced Micro Devices, Inc. All rights reserved.
|
||||
Copyright (c) 2015 - present 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,7 +20,8 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
THE SOFTWARE.
|
||||
*/
|
||||
|
||||
#pragma once
|
||||
#ifndef HIP_INCLUDE_HIP_HIP_PROFILE_H
|
||||
#define HIP_INCLUDE_HIP_HIP_PROFILE_H
|
||||
|
||||
#if not defined (ENABLE_HIP_PROFILE)
|
||||
#define ENABLE_HIP_PROFILE 1
|
||||
@@ -36,3 +37,5 @@ THE SOFTWARE.
|
||||
#define HIP_BEGIN_MARKER(markerName, group)
|
||||
#define HIP_END_MARKER()
|
||||
#endif
|
||||
|
||||
#endif
|
||||
|
||||
@@ -1,5 +1,5 @@
|
||||
/*
|
||||
Copyright (c) 2015-2017 Advanced Micro Devices, Inc. All rights reserved.
|
||||
Copyright (c) 2015 - present 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
|
||||
@@ -33,7 +33,8 @@ THE SOFTWARE.
|
||||
//! hip_runtime.h : includes everything in hip_api.h, plus math builtins and kernel launch macros.
|
||||
//! hip_runtime_api.h : Defines HIP API. This is a C header file and does not use any C++ features.
|
||||
|
||||
#pragma once
|
||||
#ifndef HIP_INCLUDE_HIP_HIP_RUNTIME_H
|
||||
#define HIP_INCLUDE_HIP_HIP_RUNTIME_H
|
||||
|
||||
// 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
|
||||
@@ -61,3 +62,5 @@ THE SOFTWARE.
|
||||
|
||||
#include <hip/hip_runtime_api.h>
|
||||
#include <hip/hip_vector_types.h>
|
||||
|
||||
#endif
|
||||
|
||||
@@ -1,5 +1,5 @@
|
||||
/*
|
||||
Copyright (c) 2015-2017 Advanced Micro Devices, Inc. All rights reserved.
|
||||
Copyright (c) 2015 - present 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
|
||||
@@ -27,7 +27,8 @@ THE SOFTWARE.
|
||||
* This file can be compiled with a standard compiler.
|
||||
*/
|
||||
|
||||
#pragma once
|
||||
#ifndef HIP_INCLUDE_HIP_HIP_RUNTIME_API_H
|
||||
#define HIP_INCLUDE_HIP_HIP_RUNTIME_API_H
|
||||
|
||||
|
||||
#include <string.h> // for getDeviceProp
|
||||
@@ -283,3 +284,5 @@ static inline hipError_t hipHostMalloc( T** ptr, size_t size, unsigned int flags
|
||||
return hipHostMalloc((void**)ptr, size, flags);
|
||||
}
|
||||
#endif
|
||||
|
||||
#endif
|
||||
|
||||
@@ -1,5 +1,5 @@
|
||||
/*
|
||||
Copyright (c) 2015-2017 Advanced Micro Devices, Inc. All rights reserved.
|
||||
Copyright (c) 2015 - present 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,10 +20,8 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
THE SOFTWARE.
|
||||
*/
|
||||
|
||||
|
||||
|
||||
#ifndef HIP_HIP_TEXTURE_H
|
||||
#define HIP_HIP_TEXTURE_H
|
||||
#ifndef HIP_INCLUDE_HIP_HIP_TEXTURE_H
|
||||
#define HIP_INCLUDE_HIP_HIP_TEXTURE_H
|
||||
|
||||
#if defined(__HIP_PLATFORM_HCC__) && !defined (__HIP_PLATFORM_NVCC__)
|
||||
#include <hip/hcc_detail/hip_texture.h>
|
||||
|
||||
@@ -1,5 +1,5 @@
|
||||
/*
|
||||
Copyright (c) 2015-2017 Advanced Micro Devices, Inc. All rights reserved.
|
||||
Copyright (c) 2015 - present 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
|
||||
@@ -22,7 +22,8 @@ THE SOFTWARE.
|
||||
|
||||
//! hip_vector_types.h : Defines the HIP vector types.
|
||||
|
||||
#pragma once
|
||||
#ifndef HIP_INCLUDE_HIP_HIP_VECTOR_TYPES_H
|
||||
#define HIP_INCLUDE_HIP_HIP_VECTOR_TYPES_H
|
||||
|
||||
#include <hip/hip_common.h>
|
||||
|
||||
@@ -36,3 +37,5 @@ THE SOFTWARE.
|
||||
#else
|
||||
#error("Must define exactly one of __HIP_PLATFORM_HCC__ or __HIP_PLATFORM_NVCC__");
|
||||
#endif
|
||||
|
||||
#endif
|
||||
|
||||
@@ -1,5 +1,5 @@
|
||||
/*
|
||||
Copyright (c) 2015-2017 Advanced Micro Devices, Inc. All rights reserved.
|
||||
Copyright (c) 2015 - present 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,7 +20,8 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
THE SOFTWARE.
|
||||
*/
|
||||
|
||||
#pragma once
|
||||
#ifndef HIP_INCLUDE_HIP_MATH_FUNCTIONS_H
|
||||
#define HIP_INCLUDE_HIP_MATH_FUNCTIONS_H
|
||||
|
||||
// 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
|
||||
@@ -34,3 +35,5 @@ THE SOFTWARE.
|
||||
#else
|
||||
#error("Must define exactly one of __HIP_PLATFORM_HCC__ or __HIP_PLATFORM_NVCC__");
|
||||
#endif
|
||||
|
||||
#endif
|
||||
|
||||
@@ -1,5 +1,5 @@
|
||||
/*
|
||||
Copyright (c) 2015-2017 Advanced Micro Devices, Inc. All rights reserved.
|
||||
Copyright (c) 2015 - present 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,6 +20,9 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
THE SOFTWARE.
|
||||
*/
|
||||
|
||||
#pragma once
|
||||
#ifndef HIP_INCLUDE_HIP_NVCC_DETAIL_CHANNEL_DESCRIPTOR_H
|
||||
#define HIP_INCLUDE_HIP_NVCC_DETAIL_CHANNEL_DESCRIPTOR_H
|
||||
|
||||
#include"channel_descriptor.h"
|
||||
|
||||
#endif
|
||||
|
||||
@@ -1,5 +1,5 @@
|
||||
/*
|
||||
Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved.
|
||||
Copyright (c) 2015 - present 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 HIPCOMPLEX_H
|
||||
#define HIPCOMPLEX_H
|
||||
#ifndef HIP_INCLUDE_HIP_NVCC_DETAIL_HIP_COMPLEX_H
|
||||
#define HIP_INCLUDE_HIP_NVCC_DETAIL_HIP_COMPLEX_H
|
||||
|
||||
#include"cuComplex.h"
|
||||
|
||||
|
||||
@@ -1,5 +1,5 @@
|
||||
/*
|
||||
Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved.
|
||||
Copyright (c) 2015 - present 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,7 +20,8 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
THE SOFTWARE.
|
||||
*/
|
||||
|
||||
#pragma once
|
||||
#ifndef HIP_INCLUDE_HIP_NVCC_DETAIL_HIP_RUNTIME_H
|
||||
#define HIP_INCLUDE_HIP_NVCC_DETAIL_HIP_RUNTIME_H
|
||||
|
||||
#include <cuda_runtime.h>
|
||||
|
||||
@@ -105,3 +106,5 @@ kernelName<<<numblocks,numthreads,memperblock,streamId>>>(0, ##__VA_ARGS__);\
|
||||
#define HIP_DYNAMIC_SHARED_ATTRIBUTE
|
||||
|
||||
#endif
|
||||
|
||||
#endif
|
||||
|
||||
@@ -1,5 +1,5 @@
|
||||
/*
|
||||
Copyright (c) 2015-2017 Advanced Micro Devices, Inc. All rights reserved.
|
||||
Copyright (c) 2015 - present 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,7 +20,8 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
THE SOFTWARE.
|
||||
*/
|
||||
|
||||
#pragma once
|
||||
#ifndef HIP_INCLUDE_HIP_NVCC_DETAIL_HIP_RUNTIME_API_H
|
||||
#define HIP_INCLUDE_HIP_NVCC_DETAIL_HIP_RUNTIME_API_H
|
||||
|
||||
#include <cuda_runtime_api.h>
|
||||
#include <cuda.h>
|
||||
@@ -333,6 +334,16 @@ inline static hipError_t hipMemcpyToSymbolAsync(const void* symbol, const void*
|
||||
return hipCUDAErrorTohipError(cudaMemcpyToSymbolAsync(symbol, src, sizeBytes, offset, hipMemcpyKindToCudaMemcpyKind(copyType)));
|
||||
}
|
||||
|
||||
inline static hipError_t hipMemcpyFromSymbol(void *dst, const void* symbolName, size_t sizeBytes, size_t offset, hipMemcpyKind kind)
|
||||
{
|
||||
return hipCUDAErrorTohipError(cudaMemcpyFromSymbol(dst, symbolName, sizeBytes, offset, hipMemcpyKindToCudaMemcpyKind(kind)));
|
||||
}
|
||||
|
||||
inline static hipError_t hipMemcpyFromSymbolAsync(void *dst, const void* symbolName, size_t sizeBytes, size_t offset, hipMemcpyKind kind, hipStream_t stream)
|
||||
{
|
||||
return hipCUDAErrorTohipError(cudaMemcpyFromSymbolAsync(dst, symbolName, sizeBytes, offset, hipMemcpyKindToCudaMemcpyKind(kind), stream));
|
||||
}
|
||||
|
||||
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)));
|
||||
}
|
||||
@@ -670,6 +681,31 @@ inline static hipError_t hipCtxEnablePeerAccess ( hipCtx_t peerCtx, unsigned in
|
||||
return hipCUResultTohipError(cuCtxEnablePeerAccess(peerCtx, flags));
|
||||
}
|
||||
|
||||
inline static hipError_t hipDevicePrimaryCtxGetState ( hipDevice_t dev, unsigned int* flags, int* active )
|
||||
{
|
||||
return hipCUResultTohipError(cuDevicePrimaryCtxGetState(dev, flags, active));
|
||||
}
|
||||
|
||||
inline static hipError_t hipDevicePrimaryCtxRelease ( hipDevice_t dev)
|
||||
{
|
||||
return hipCUResultTohipError(cuDevicePrimaryCtxRelease(dev));
|
||||
}
|
||||
|
||||
inline static hipError_t hipDevicePrimaryCtxRetain ( hipCtx_t* pctx, hipDevice_t dev )
|
||||
{
|
||||
return hipCUResultTohipError(cuDevicePrimaryCtxRetain(pctx, dev));
|
||||
}
|
||||
|
||||
inline static hipError_t hipDevicePrimaryCtxReset ( hipDevice_t dev )
|
||||
{
|
||||
return hipCUResultTohipError(cuDevicePrimaryCtxReset(dev));
|
||||
}
|
||||
|
||||
inline static hipError_t hipDevicePrimaryCtxSetFlags ( hipDevice_t dev, unsigned int flags )
|
||||
{
|
||||
return hipCUResultTohipError(cuDevicePrimaryCtxSetFlags(dev, flags));
|
||||
}
|
||||
|
||||
inline static hipError_t hipMemGetAddressRange ( hipDeviceptr_t* pbase, size_t* psize, hipDeviceptr_t dptr )
|
||||
{
|
||||
return hipCUResultTohipError(cuMemGetAddressRange( pbase , psize , dptr));
|
||||
|
||||
@@ -1,5 +1,27 @@
|
||||
#ifndef HIP_TEXTURE_H
|
||||
#define HIP_TEXTURE_H
|
||||
/*
|
||||
Copyright (c) 2015 - present 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_INCLUDE_HIP_NVCC_DETAIL_HIP_TEXTURE_H
|
||||
#define HIP_INCLUDE_HIP_NVCC_DETAIL_HIP_TEXTURE_H
|
||||
|
||||
#include <texture_types.h>
|
||||
|
||||
|
||||
@@ -1,13 +1,16 @@
|
||||
/*
|
||||
Copyright (c) 2015-2017 Advanced Micro Devices, Inc. All rights reserved.
|
||||
Copyright (c) 2015 - present 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
|
||||
@@ -17,7 +20,8 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
THE SOFTWARE.
|
||||
*/
|
||||
|
||||
#pragma once
|
||||
#ifndef HIP_INCLUDE_HIP_TEXTURE_TYPES_H
|
||||
#define HIP_INCLUDE_HIP_TEXTURE_TYPES_H
|
||||
|
||||
#include <hip/hip_common.h>
|
||||
|
||||
@@ -28,3 +32,5 @@ THE SOFTWARE.
|
||||
#else
|
||||
#error("Must define exactly one of __HIP_PLATFORM_HCC__ or __HIP_PLATFORM_NVCC__");
|
||||
#endif
|
||||
|
||||
#endif
|
||||
|
||||
@@ -24,4 +24,3 @@ $(EXE): bit_extract.cpp
|
||||
|
||||
clean:
|
||||
rm -f *.o $(EXE)
|
||||
|
||||
|
||||
@@ -37,7 +37,7 @@ THE SOFTWARE.
|
||||
}\
|
||||
}
|
||||
|
||||
void __global__
|
||||
__global__ void
|
||||
bit_extract_kernel(hipLaunchParm lp, uint32_t *C_d, const uint32_t *A_d, size_t N)
|
||||
{
|
||||
size_t offset = (hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x);
|
||||
@@ -45,7 +45,7 @@ bit_extract_kernel(hipLaunchParm lp, uint32_t *C_d, const uint32_t *A_d, size_t
|
||||
|
||||
for (size_t i=offset; i<N; i+=stride) {
|
||||
#ifdef __HIP_PLATFORM_HCC__
|
||||
C_d[i] = hc::__bitextract_u32(A_d[i], 8, 4);
|
||||
C_d[i] = hc::__bitextract_u32(A_d[i], 8, 4);
|
||||
#else /* defined __HIP_PLATFORM_NVCC__ or other path */
|
||||
C_d[i] = ((A_d[i] & 0xf00) >> 8);
|
||||
#endif
|
||||
@@ -73,7 +73,7 @@ int main(int argc, char *argv[])
|
||||
C_h = (uint32_t*)malloc(Nbytes);
|
||||
CHECK(C_h == 0 ? hipErrorMemoryAllocation : hipSuccess );
|
||||
|
||||
for (size_t i=0; i<N; i++)
|
||||
for (size_t i=0; i<N; i++)
|
||||
{
|
||||
A_h[i] = i;
|
||||
}
|
||||
|
||||
@@ -5,14 +5,19 @@ endif
|
||||
HIPCC=$(HIP_PATH)/bin/hipcc
|
||||
HIP_PLATFORM=$(shell $(HIP_PATH)/bin/hipconfig --compiler)
|
||||
|
||||
all: vcpy_kernel.code runKernel.hip.out
|
||||
all: vcpy_kernel.code runKernel.hip.out defaultDriver.hip.out
|
||||
|
||||
runKernel.hip.out: runKernel.cpp
|
||||
$(HIPCC) $(HIPCC_FLAGS) $< -o $@
|
||||
|
||||
launchKernelHcc.hip.out: launchKernelHcc.cpp
|
||||
$(HIPCC) $(HIPCC_FLAGS) $< -o $@
|
||||
|
||||
defaultDriver.hip.out: defaultDriver.cpp
|
||||
$(HIPCC) $(HIPCC_FLAGS) $< -o $@
|
||||
|
||||
vcpy_kernel.code: vcpy_kernel.cpp
|
||||
$(HIPCC) --genco $(GENCO_FLAGS) $^ -o $@
|
||||
$(HIPCC) --genco $(GENCO_FLAGS) $^ -o $@
|
||||
|
||||
clean:
|
||||
rm -f *.code *.out
|
||||
|
||||
|
||||
@@ -0,0 +1,89 @@
|
||||
/*
|
||||
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.
|
||||
*/
|
||||
|
||||
#include "hip/hip_runtime.h"
|
||||
#include "hip/hip_runtime_api.h"
|
||||
#include<iostream>
|
||||
#include<fstream>
|
||||
#include<vector>
|
||||
|
||||
#define LEN 64
|
||||
#define SIZE LEN<<2
|
||||
|
||||
#define fileName "test.co"
|
||||
#define kernel_name "vadd"
|
||||
|
||||
int main(){
|
||||
float *A, *B, *C;
|
||||
hipDeviceptr_t Ad, Bd, Cd;
|
||||
A = new float[LEN];
|
||||
B = new float[LEN];
|
||||
C = new float[LEN];
|
||||
|
||||
for(uint32_t i=0;i<LEN;i++){
|
||||
A[i] = i*1.0f;
|
||||
B[i] = 1.0f;
|
||||
C[i] = 0.0f;
|
||||
}
|
||||
|
||||
hipInit(0);
|
||||
hipDevice_t device;
|
||||
hipCtx_t context;
|
||||
hipDeviceGet(&device, 0);
|
||||
hipCtxCreate(&context, 0, device);
|
||||
|
||||
hipMalloc((void**)&Ad, SIZE);
|
||||
hipMalloc((void**)&Bd, SIZE);
|
||||
hipMalloc((void**)&Cd, SIZE);
|
||||
|
||||
hipMemcpyHtoD(Ad, A, SIZE);
|
||||
hipMemcpyHtoD(Bd, B, SIZE);
|
||||
hipMemcpyHtoD(Cd, C, SIZE);
|
||||
|
||||
hipModule_t Module;
|
||||
hipFunction_t Function;
|
||||
hipModuleLoad(&Module, fileName);
|
||||
hipModuleGetFunction(&Function, Module, kernel_name);
|
||||
|
||||
int n = LEN;
|
||||
void * args[4] = {&Ad, &Bd, &Cd, &n};
|
||||
|
||||
hipModuleLaunchKernel(Function, 1, 1, 1, LEN, 1, 1, 0, 0, args, nullptr);
|
||||
|
||||
hipMemcpyDtoH(C, Cd, SIZE);
|
||||
int mismatchCount = 0;
|
||||
for(uint32_t i=0;i<LEN;i++){
|
||||
if (A[i] + B[i] != C[i]) {
|
||||
mismatchCount++;
|
||||
std::cout<<"error: mismatch " << A[i]<<" + "<<B[i]<<" != "<<C[i]<<std::endl;
|
||||
}
|
||||
}
|
||||
|
||||
if (mismatchCount == 0) {
|
||||
std::cout << "PASSED!\n";
|
||||
} else {
|
||||
std::cout << "FAILED!\n";
|
||||
};
|
||||
|
||||
hipCtxDestroy(context);
|
||||
return 0;
|
||||
}
|
||||
@@ -0,0 +1,112 @@
|
||||
/*
|
||||
Copyright (c) 2015 - present Advanced Micro Devices, Inc. All rights reserved.
|
||||
|
||||
Permission is hereby granted, free of charge, to any person obtaining a copy
|
||||
of this software and associated documentation files (the "Software"), to deal
|
||||
in the Software without restriction, including without limitation the rights
|
||||
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
|
||||
copies of the Software, and to permit persons to whom the Software is
|
||||
furnished to do so, subject to the following conditions:
|
||||
|
||||
The above copyright notice and this permission notice shall be included in
|
||||
all copies or substantial portions of the Software.
|
||||
|
||||
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
|
||||
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
|
||||
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
|
||||
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
|
||||
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
|
||||
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
THE SOFTWARE.
|
||||
*/
|
||||
|
||||
|
||||
|
||||
#include "hip/hip_runtime.h"
|
||||
#include "hip/hip_runtime_api.h"
|
||||
#include <iostream>
|
||||
#include <fstream>
|
||||
#include <vector>
|
||||
|
||||
#ifdef __HIP_PLATFORM_HCC__
|
||||
#include <hip/hip_hcc.h>
|
||||
#endif
|
||||
|
||||
#define LEN 64
|
||||
#define SIZE LEN<<2
|
||||
|
||||
#define fileName "vcpy_kernel.code"
|
||||
#define kernel_name "hello_world"
|
||||
|
||||
#define HIP_CHECK(status) \
|
||||
if(status != hipSuccess) {std::cout<<"Got Status: "<<status<<" at Line: "<<__LINE__<<std::endl;exit(0);}
|
||||
|
||||
int main(){
|
||||
float *A, *B;
|
||||
hipDeviceptr_t Ad, Bd;
|
||||
A = new float[LEN];
|
||||
B = new float[LEN];
|
||||
|
||||
for(uint32_t i=0;i<LEN;i++){
|
||||
A[i] = i*1.0f;
|
||||
B[i] = 0.0f;
|
||||
}
|
||||
|
||||
hipInit(0);
|
||||
hipDevice_t device;
|
||||
hipCtx_t context;
|
||||
hipDeviceGet(&device, 0);
|
||||
hipCtxCreate(&context, 0, device);
|
||||
|
||||
hipMalloc((void**)&Ad, SIZE);
|
||||
hipMalloc((void**)&Bd, SIZE);
|
||||
|
||||
hipMemcpyHtoD(Ad, A, SIZE);
|
||||
hipMemcpyHtoD(Bd, B, SIZE);
|
||||
hipModule_t Module;
|
||||
hipFunction_t Function;
|
||||
HIP_CHECK(hipModuleLoad(&Module, fileName));
|
||||
HIP_CHECK(hipModuleGetFunction(&Function, Module, kernel_name));
|
||||
|
||||
uint32_t len = LEN;
|
||||
uint32_t one = 1;
|
||||
|
||||
struct {
|
||||
void * _Ad;
|
||||
void * _Bd;
|
||||
} args;
|
||||
|
||||
args._Ad = Ad;
|
||||
args._Bd = Bd;
|
||||
|
||||
|
||||
|
||||
size_t size = sizeof(args);
|
||||
|
||||
void *config[] = {
|
||||
HIP_LAUNCH_PARAM_BUFFER_POINTER, &args,
|
||||
HIP_LAUNCH_PARAM_BUFFER_SIZE, &size,
|
||||
HIP_LAUNCH_PARAM_END
|
||||
};
|
||||
|
||||
HIP_CHECK(hipHccModuleLaunchKernel(Function, LEN, 1, 1, LEN, 1, 1, 0, 0, NULL, (void**)&config));
|
||||
|
||||
hipMemcpyDtoH(B, Bd, SIZE);
|
||||
|
||||
int mismatchCount = 0;
|
||||
for(uint32_t i=0;i<LEN;i++){
|
||||
if (A[i] != B[i]) {
|
||||
mismatchCount++;
|
||||
std::cout<<"error: mismatch " << A[i]<<" != "<<B[i]<<std::endl;
|
||||
}
|
||||
}
|
||||
|
||||
if (mismatchCount == 0) {
|
||||
std::cout << "PASSED!\n";
|
||||
} else {
|
||||
std::cout << "FAILED!\n";
|
||||
};
|
||||
|
||||
hipCtxDestroy(context);
|
||||
return 0;
|
||||
}
|
||||
@@ -1,5 +1,5 @@
|
||||
/*
|
||||
Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved.
|
||||
Copyright (c) 2015 - present 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
|
||||
@@ -22,9 +22,10 @@ THE SOFTWARE.
|
||||
|
||||
#include "hip/hip_runtime.h"
|
||||
#include "hip/hip_runtime_api.h"
|
||||
#include<iostream>
|
||||
#include<fstream>
|
||||
#include<vector>
|
||||
#include <iostream>
|
||||
#include <fstream>
|
||||
#include <vector>
|
||||
#include <hip/hip_hcc.h>
|
||||
|
||||
#define LEN 64
|
||||
#define SIZE LEN<<2
|
||||
@@ -32,6 +33,9 @@ THE SOFTWARE.
|
||||
#define fileName "vcpy_kernel.code"
|
||||
#define kernel_name "hello_world"
|
||||
|
||||
#define HIP_CHECK(status) \
|
||||
if(status != hipSuccess) {std::cout<<"Got Status: "<<status<<" at Line: "<<__LINE__<<std::endl;exit(0);}
|
||||
|
||||
int main(){
|
||||
float *A, *B;
|
||||
hipDeviceptr_t Ad, Bd;
|
||||
@@ -43,10 +47,10 @@ int main(){
|
||||
B[i] = 0.0f;
|
||||
}
|
||||
|
||||
hipInit(0);
|
||||
hipDevice_t device;
|
||||
hipCtx_t context;
|
||||
hipDeviceGet(&device, 0);
|
||||
hipInit(0);
|
||||
hipDevice_t device;
|
||||
hipCtx_t context;
|
||||
hipDeviceGet(&device, 0);
|
||||
hipCtxCreate(&context, 0, device);
|
||||
|
||||
hipMalloc((void**)&Ad, SIZE);
|
||||
@@ -56,22 +60,18 @@ int main(){
|
||||
hipMemcpyHtoD(Bd, B, SIZE);
|
||||
hipModule_t Module;
|
||||
hipFunction_t Function;
|
||||
hipModuleLoad(&Module, fileName);
|
||||
hipModuleGetFunction(&Function, Module, kernel_name);
|
||||
HIP_CHECK(hipModuleLoad(&Module, fileName));
|
||||
HIP_CHECK(hipModuleGetFunction(&Function, Module, kernel_name));
|
||||
|
||||
#ifdef __HIP_PLATFORM_HCC__
|
||||
uint32_t len = LEN;
|
||||
uint32_t one = 1;
|
||||
|
||||
struct {
|
||||
uint32_t _hidden[6];
|
||||
void * _Ad;
|
||||
void * _Bd;
|
||||
} args;
|
||||
|
||||
for (int i=0; i<6; i++) {
|
||||
args._hidden[i] = 0;
|
||||
}
|
||||
args._Ad = Ad;
|
||||
args._Bd = Bd;
|
||||
|
||||
@@ -98,9 +98,10 @@ int main(){
|
||||
HIP_LAUNCH_PARAM_END
|
||||
};
|
||||
|
||||
hipModuleLaunchKernel(Function, 1, 1, 1, LEN, 1, 1, 0, 0, NULL, (void**)&config);
|
||||
HIP_CHECK(hipModuleLaunchKernel(Function, 1, 1, 1, LEN, 1, 1, 0, 0, NULL, (void**)&config));
|
||||
|
||||
hipMemcpyDtoH(B, Bd, SIZE);
|
||||
|
||||
int mismatchCount = 0;
|
||||
for(uint32_t i=0;i<LEN;i++){
|
||||
if (A[i] != B[i]) {
|
||||
|
||||
@@ -0,0 +1,12 @@
|
||||
__kernel void memset(char in, __global int* out) {
|
||||
int tx = get_global_id(0);
|
||||
out[tx] = in;
|
||||
}
|
||||
|
||||
|
||||
__kernel void vadd(__global float *Ad, __global float *Bd, __global float *Cd, int N){
|
||||
int tx = get_global_id(0);
|
||||
if(tx < N){
|
||||
Cd[tx] = Ad[tx] + Bd[tx];
|
||||
}
|
||||
}
|
||||
دودویی (BIN)
فایل باینری نشان داده نشده است.
@@ -1,5 +1,5 @@
|
||||
/*
|
||||
Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved.
|
||||
Copyright (c) 2015 - present 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
|
||||
@@ -27,4 +27,3 @@ extern "C" __global__ void hello_world(hipLaunchParm lp, float *a, float *b)
|
||||
int tx = hipThreadIdx_x;
|
||||
b[tx] = a[tx];
|
||||
}
|
||||
|
||||
|
||||
@@ -83,7 +83,7 @@ int main(int argc, char *argv[])
|
||||
const unsigned threadsPerBlock = 256;
|
||||
|
||||
printf ("info: launch 'vector_square' kernel\n");
|
||||
hipLaunchKernel(vector_square, dim3(blocks), dim3(threadsPerBlock), 0, 0, C_d, A_d, N);
|
||||
hipLaunchKernel(vector_square, dim3(blocks), dim3(threadsPerBlock), 0, nullptr, C_d, A_d, N);
|
||||
|
||||
printf ("info: copy Device2Host\n");
|
||||
CHECK ( hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost));
|
||||
|
||||
@@ -0,0 +1,45 @@
|
||||
//===-- AMDGPUNoteType.h - AMDGPU ELF PT_NOTE section info-------*- C++ -*-===//
|
||||
//
|
||||
// The LLVM Compiler Infrastructure
|
||||
//
|
||||
// This file is distributed under the University of Illinois Open Source
|
||||
// License. See LICENSE.TXT for details.
|
||||
//
|
||||
//===----------------------------------------------------------------------===//
|
||||
//
|
||||
/// \file
|
||||
///
|
||||
/// Enums and constants for AMDGPU PT_NOTE sections.
|
||||
///
|
||||
//
|
||||
//===----------------------------------------------------------------------===//
|
||||
//
|
||||
#ifndef LLVM_LIB_TARGET_AMDGPU_AMDGPUPTNOTE_H
|
||||
#define LLVM_LIB_TARGET_AMDGPU_AMDGPUPTNOTE_H
|
||||
|
||||
namespace AMDGPU {
|
||||
|
||||
namespace ElfNote {
|
||||
|
||||
const char SectionName[] = ".note";
|
||||
|
||||
const char NoteName[] = "AMD";
|
||||
|
||||
// TODO: Move this enum to include/llvm/Support so it can be used in tools?
|
||||
enum NoteType{
|
||||
NT_AMDGPU_HSA_CODE_OBJECT_VERSION = 1,
|
||||
NT_AMDGPU_HSA_HSAIL = 2,
|
||||
NT_AMDGPU_HSA_ISA = 3,
|
||||
NT_AMDGPU_HSA_PRODUCER = 4,
|
||||
NT_AMDGPU_HSA_PRODUCER_OPTIONS = 5,
|
||||
NT_AMDGPU_HSA_EXTENSION = 6,
|
||||
NT_AMDGPU_HSA_RUNTIME_METADATA_V_1 = 7, // deprecated since 12/14/16.
|
||||
NT_AMDGPU_HSA_RUNTIME_METADATA_V_2 = 8,
|
||||
NT_AMDGPU_HSA_RUNTIME_METADATA = NT_AMDGPU_HSA_RUNTIME_METADATA_V_2,
|
||||
NT_AMDGPU_HSA_HLDEBUG_DEBUG = 101,
|
||||
NT_AMDGPU_HSA_HLDEBUG_TARGET = 102
|
||||
};
|
||||
}
|
||||
}
|
||||
|
||||
#endif // LLVM_LIB_TARGET_AMDGPU_AMDGPUNOTETYPE_H
|
||||
@@ -0,0 +1,290 @@
|
||||
//===-- AMDGPURuntimeMetadata.h - AMDGPU Runtime Metadata -------*- C++ -*-===//
|
||||
//
|
||||
// The LLVM Compiler Infrastructure
|
||||
//
|
||||
// This file is distributed under the University of Illinois Open Source
|
||||
// License. See LICENSE.TXT for details.
|
||||
//
|
||||
//===----------------------------------------------------------------------===//
|
||||
//
|
||||
/// \file
|
||||
///
|
||||
/// Enums and structure types used by runtime metadata.
|
||||
///
|
||||
/// Runtime requests certain information (metadata) about kernels to be able
|
||||
/// to execute the kernels and answer the queries about the kernels.
|
||||
/// The metadata is represented as a note element in the .note ELF section of a
|
||||
/// binary (code object). The desc field of the note element is a YAML string
|
||||
/// consisting of key-value pairs. Each key is a string. Each value can be
|
||||
/// an integer, a string, or an YAML sequence. There are 3 levels of YAML maps.
|
||||
/// At the beginning of the YAML string is the module level YAML map. A
|
||||
/// kernel-level YAML map is in the amd.Kernels sequence. A
|
||||
/// kernel-argument-level map is in the amd.Args sequence.
|
||||
///
|
||||
/// The format should be kept backward compatible. New enum values and bit
|
||||
/// fields should be appended at the end. It is suggested to bump up the
|
||||
/// revision number whenever the format changes and document the change
|
||||
/// in the revision in this header.
|
||||
///
|
||||
//
|
||||
//===----------------------------------------------------------------------===//
|
||||
//
|
||||
#ifndef LLVM_LIB_TARGET_AMDGPU_AMDGPURUNTIMEMETADATA_H
|
||||
#define LLVM_LIB_TARGET_AMDGPU_AMDGPURUNTIMEMETADATA_H
|
||||
|
||||
#include <cstdint>
|
||||
#include <vector>
|
||||
#include <string>
|
||||
|
||||
namespace AMDGPU {
|
||||
namespace RuntimeMD {
|
||||
|
||||
// Version and revision of runtime metadata
|
||||
const unsigned char MDVersion = 2;
|
||||
const unsigned char MDRevision = 1;
|
||||
|
||||
// Name of keys for runtime metadata.
|
||||
namespace KeyName {
|
||||
|
||||
// Runtime metadata version
|
||||
const char MDVersion[] = "amd.MDVersion";
|
||||
|
||||
// Instruction set architecture information
|
||||
const char IsaInfo[] = "amd.IsaInfo";
|
||||
// Wavefront size
|
||||
const char IsaInfoWavefrontSize[] = "amd.IsaInfoWavefrontSize";
|
||||
// Local memory size in bytes
|
||||
const char IsaInfoLocalMemorySize[] = "amd.IsaInfoLocalMemorySize";
|
||||
// Number of execution units per compute unit
|
||||
const char IsaInfoEUsPerCU[] = "amd.IsaInfoEUsPerCU";
|
||||
// Maximum number of waves per execution unit
|
||||
const char IsaInfoMaxWavesPerEU[] = "amd.IsaInfoMaxWavesPerEU";
|
||||
// Maximum flat work group size
|
||||
const char IsaInfoMaxFlatWorkGroupSize[] = "amd.IsaInfoMaxFlatWorkGroupSize";
|
||||
// SGPR allocation granularity
|
||||
const char IsaInfoSGPRAllocGranule[] = "amd.IsaInfoSGPRAllocGranule";
|
||||
// Total number of SGPRs
|
||||
const char IsaInfoTotalNumSGPRs[] = "amd.IsaInfoTotalNumSGPRs";
|
||||
// Addressable number of SGPRs
|
||||
const char IsaInfoAddressableNumSGPRs[] = "amd.IsaInfoAddressableNumSGPRs";
|
||||
// VGPR allocation granularity
|
||||
const char IsaInfoVGPRAllocGranule[] = "amd.IsaInfoVGPRAllocGranule";
|
||||
// Total number of VGPRs
|
||||
const char IsaInfoTotalNumVGPRs[] = "amd.IsaInfoTotalNumVGPRs";
|
||||
// Addressable number of VGPRs
|
||||
const char IsaInfoAddressableNumVGPRs[] = "amd.IsaInfoAddressableNumVGPRs";
|
||||
|
||||
// Language
|
||||
const char Language[] = "amd.Language";
|
||||
// Language version
|
||||
const char LanguageVersion[] = "amd.LanguageVersion";
|
||||
|
||||
// Kernels
|
||||
const char Kernels[] = "amd.Kernels";
|
||||
// Kernel name
|
||||
const char KernelName[] = "amd.KernelName";
|
||||
// Kernel arguments
|
||||
const char Args[] = "amd.Args";
|
||||
// Kernel argument size in bytes
|
||||
const char ArgSize[] = "amd.ArgSize";
|
||||
// Kernel argument alignment
|
||||
const char ArgAlign[] = "amd.ArgAlign";
|
||||
// Kernel argument type name
|
||||
const char ArgTypeName[] = "amd.ArgTypeName";
|
||||
// Kernel argument name
|
||||
const char ArgName[] = "amd.ArgName";
|
||||
// Kernel argument kind
|
||||
const char ArgKind[] = "amd.ArgKind";
|
||||
// Kernel argument value type
|
||||
const char ArgValueType[] = "amd.ArgValueType";
|
||||
// Kernel argument address qualifier
|
||||
const char ArgAddrQual[] = "amd.ArgAddrQual";
|
||||
// Kernel argument access qualifier
|
||||
const char ArgAccQual[] = "amd.ArgAccQual";
|
||||
// Kernel argument is const qualified
|
||||
const char ArgIsConst[] = "amd.ArgIsConst";
|
||||
// Kernel argument is restrict qualified
|
||||
const char ArgIsRestrict[] = "amd.ArgIsRestrict";
|
||||
// Kernel argument is volatile qualified
|
||||
const char ArgIsVolatile[] = "amd.ArgIsVolatile";
|
||||
// Kernel argument is pipe qualified
|
||||
const char ArgIsPipe[] = "amd.ArgIsPipe";
|
||||
// Required work group size
|
||||
const char ReqdWorkGroupSize[] = "amd.ReqdWorkGroupSize";
|
||||
// Work group size hint
|
||||
const char WorkGroupSizeHint[] = "amd.WorkGroupSizeHint";
|
||||
// Vector type hint
|
||||
const char VecTypeHint[] = "amd.VecTypeHint";
|
||||
// Kernel index for device enqueue
|
||||
const char KernelIndex[] = "amd.KernelIndex";
|
||||
// No partial work groups
|
||||
const char NoPartialWorkGroups[] = "amd.NoPartialWorkGroups";
|
||||
// Prinf function call information
|
||||
const char PrintfInfo[] = "amd.PrintfInfo";
|
||||
// The actual kernel argument access qualifier
|
||||
const char ArgActualAcc[] = "amd.ArgActualAcc";
|
||||
// Alignment of pointee type
|
||||
const char ArgPointeeAlign[] = "amd.ArgPointeeAlign";
|
||||
|
||||
} // end namespace KeyName
|
||||
|
||||
namespace KernelArg {
|
||||
|
||||
enum Kind : uint8_t {
|
||||
ByValue = 0,
|
||||
GlobalBuffer = 1,
|
||||
DynamicSharedPointer = 2,
|
||||
Sampler = 3,
|
||||
Image = 4,
|
||||
Pipe = 5,
|
||||
Queue = 6,
|
||||
HiddenGlobalOffsetX = 7,
|
||||
HiddenGlobalOffsetY = 8,
|
||||
HiddenGlobalOffsetZ = 9,
|
||||
HiddenNone = 10,
|
||||
HiddenPrintfBuffer = 11,
|
||||
HiddenDefaultQueue = 12,
|
||||
HiddenCompletionAction = 13,
|
||||
};
|
||||
|
||||
enum ValueType : uint16_t {
|
||||
Struct = 0,
|
||||
I8 = 1,
|
||||
U8 = 2,
|
||||
I16 = 3,
|
||||
U16 = 4,
|
||||
F16 = 5,
|
||||
I32 = 6,
|
||||
U32 = 7,
|
||||
F32 = 8,
|
||||
I64 = 9,
|
||||
U64 = 10,
|
||||
F64 = 11,
|
||||
};
|
||||
|
||||
// Avoid using 'None' since it conflicts with a macro in X11 header file.
|
||||
enum AccessQualifer : uint8_t {
|
||||
AccNone = 0,
|
||||
ReadOnly = 1,
|
||||
WriteOnly = 2,
|
||||
ReadWrite = 3,
|
||||
};
|
||||
|
||||
enum AddressSpaceQualifer : uint8_t {
|
||||
Private = 0,
|
||||
Global = 1,
|
||||
Constant = 2,
|
||||
Local = 3,
|
||||
Generic = 4,
|
||||
Region = 5,
|
||||
};
|
||||
|
||||
} // end namespace KernelArg
|
||||
|
||||
// Invalid values are used to indicate an optional key should not be emitted.
|
||||
const uint8_t INVALID_ADDR_QUAL = 0xff;
|
||||
const uint8_t INVALID_ACC_QUAL = 0xff;
|
||||
const uint32_t INVALID_KERNEL_INDEX = ~0U;
|
||||
|
||||
namespace KernelArg {
|
||||
|
||||
// In-memory representation of kernel argument information.
|
||||
struct Metadata {
|
||||
uint32_t Size = 0;
|
||||
uint32_t Align = 0;
|
||||
uint32_t PointeeAlign = 0;
|
||||
uint8_t Kind = 0;
|
||||
uint16_t ValueType = 0;
|
||||
std::string TypeName;
|
||||
std::string Name;
|
||||
uint8_t AddrQual = INVALID_ADDR_QUAL;
|
||||
uint8_t AccQual = INVALID_ACC_QUAL;
|
||||
uint8_t IsVolatile = 0;
|
||||
uint8_t IsConst = 0;
|
||||
uint8_t IsRestrict = 0;
|
||||
uint8_t IsPipe = 0;
|
||||
|
||||
Metadata() = default;
|
||||
};
|
||||
|
||||
} // end namespace KernelArg
|
||||
|
||||
namespace Kernel {
|
||||
|
||||
// In-memory representation of kernel information.
|
||||
struct Metadata {
|
||||
std::string Name;
|
||||
std::string Language;
|
||||
std::vector<uint8_t> LanguageVersion;
|
||||
std::vector<uint32_t> ReqdWorkGroupSize;
|
||||
std::vector<uint32_t> WorkGroupSizeHint;
|
||||
std::string VecTypeHint;
|
||||
uint32_t KernelIndex = INVALID_KERNEL_INDEX;
|
||||
uint8_t NoPartialWorkGroups = 0;
|
||||
std::vector<KernelArg::Metadata> Args;
|
||||
|
||||
Metadata() = default;
|
||||
};
|
||||
|
||||
} // end namespace Kernel
|
||||
|
||||
namespace IsaInfo {
|
||||
|
||||
/// \brief In-memory representation of instruction set architecture
|
||||
/// information.
|
||||
struct Metadata {
|
||||
/// \brief Wavefront size.
|
||||
unsigned WavefrontSize = 0;
|
||||
/// \brief Local memory size in bytes.
|
||||
unsigned LocalMemorySize = 0;
|
||||
/// \brief Number of execution units per compute unit.
|
||||
unsigned EUsPerCU = 0;
|
||||
/// \brief Maximum number of waves per execution unit.
|
||||
unsigned MaxWavesPerEU = 0;
|
||||
/// \brief Maximum flat work group size.
|
||||
unsigned MaxFlatWorkGroupSize = 0;
|
||||
/// \brief SGPR allocation granularity.
|
||||
unsigned SGPRAllocGranule = 0;
|
||||
/// \brief Total number of SGPRs.
|
||||
unsigned TotalNumSGPRs = 0;
|
||||
/// \brief Addressable number of SGPRs.
|
||||
unsigned AddressableNumSGPRs = 0;
|
||||
/// \brief VGPR allocation granularity.
|
||||
unsigned VGPRAllocGranule = 0;
|
||||
/// \brief Total number of VGPRs.
|
||||
unsigned TotalNumVGPRs = 0;
|
||||
/// \brief Addressable number of VGPRs.
|
||||
unsigned AddressableNumVGPRs = 0;
|
||||
|
||||
Metadata() = default;
|
||||
};
|
||||
|
||||
} // end namespace IsaInfo
|
||||
|
||||
namespace Program {
|
||||
|
||||
// In-memory representation of program information.
|
||||
struct Metadata {
|
||||
std::vector<uint8_t> MDVersionSeq;
|
||||
IsaInfo::Metadata IsaInfo;
|
||||
std::vector<std::string> PrintfInfo;
|
||||
std::vector<Kernel::Metadata> Kernels;
|
||||
|
||||
explicit Metadata() = default;
|
||||
|
||||
// Construct from an YAML string.
|
||||
explicit Metadata(const std::string &YAML);
|
||||
|
||||
// Convert to YAML string.
|
||||
std::string toYAML();
|
||||
|
||||
// Convert from YAML string.
|
||||
static Metadata fromYAML(const std::string &S);
|
||||
};
|
||||
|
||||
} //end namespace Program
|
||||
|
||||
} // end namespace RuntimeMD
|
||||
} // end namespace AMDGPU
|
||||
|
||||
#endif // LLVM_LIB_TARGET_AMDGPU_AMDGPURUNTIMEMETADATA_H
|
||||
@@ -41,8 +41,8 @@ struct holder32Bit {
|
||||
};
|
||||
} __attribute__((aligned(4)));
|
||||
|
||||
struct holder64Bit hold64;
|
||||
struct holder32Bit hold32;
|
||||
__device__ struct holder64Bit hold64;
|
||||
__device__ struct holder32Bit hold32;
|
||||
|
||||
__device__ float __double2float_rd(double x)
|
||||
{
|
||||
|
||||
@@ -34,8 +34,8 @@ THE SOFTWARE.
|
||||
This is the best place to put them because the device
|
||||
global variables need to be initialized at the start.
|
||||
*/
|
||||
__attribute__((address_space(1))) char gpuHeap[SIZE_OF_HEAP];
|
||||
__attribute__((address_space(1))) uint32_t gpuFlags[NUM_PAGES];
|
||||
__device__ ADDRESS_SPACE_1 char gpuHeap[SIZE_OF_HEAP];
|
||||
__device__ ADDRESS_SPACE_1 uint32_t gpuFlags[NUM_PAGES];
|
||||
|
||||
__device__ void *__hip_hc_malloc(size_t size)
|
||||
{
|
||||
@@ -835,10 +835,6 @@ __device__ float __hip_ynf(int n, float x)
|
||||
return b0;
|
||||
}
|
||||
|
||||
|
||||
|
||||
const int warpSize = 64;
|
||||
|
||||
__device__ long long int clock64() { return (long long int)hc::__cycle_u64(); };
|
||||
__device__ clock_t clock() { return (clock_t)hc::__cycle_u64(); };
|
||||
|
||||
@@ -1083,7 +1079,7 @@ __host__ __device__ int max(int arg1, int arg2)
|
||||
return (int)(hc::precise_math::fmax((float)arg1, (float)arg2));
|
||||
}
|
||||
|
||||
__device__ __attribute__((address_space(3))) void* __get_dynamicgroupbaseptr()
|
||||
__device__ ADDRESS_SPACE_3 void* __get_dynamicgroupbaseptr()
|
||||
{
|
||||
return hc::get_dynamic_group_segment_base_pointer();
|
||||
}
|
||||
|
||||
@@ -1,4 +1,26 @@
|
||||
#include "hip_hcc.h"
|
||||
/*
|
||||
Copyright (c) 2015 - present Advanced Micro Devices, Inc. All rights reserved.
|
||||
|
||||
Permission is hereby granted, free of charge, to any person obtaining a copy
|
||||
of this software and associated documentation files (the "Software"), to deal
|
||||
in the Software without restriction, including without limitation the rights
|
||||
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
|
||||
copies of the Software, and to permit persons to whom the Software is
|
||||
furnished to do so, subject to the following conditions:
|
||||
|
||||
The above copyright notice and this permission notice shall be included in
|
||||
all copies or substantial portions of the Software.
|
||||
|
||||
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
|
||||
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
|
||||
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
|
||||
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
|
||||
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
|
||||
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
THE SOFTWARE.
|
||||
*/
|
||||
|
||||
#include "hip_hcc_internal.h"
|
||||
#include "trace_helper.h"
|
||||
#include "env.h"
|
||||
|
||||
|
||||
@@ -0,0 +1,60 @@
|
||||
/*
|
||||
Copyright (c) 2015 - present Advanced Micro Devices, Inc. All rights reserved.
|
||||
|
||||
Permission is hereby granted, free of charge, to any person obtaining a copy
|
||||
of this software and associated documentation files (the "Software"), to deal
|
||||
in the Software without restriction, including without limitation the rights
|
||||
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
|
||||
copies of the Software, and to permit persons to whom the Software is
|
||||
furnished to do so, subject to the following conditions:
|
||||
|
||||
The above copyright notice and this permission notice shall be included in
|
||||
all copies or substantial portions of the Software.
|
||||
|
||||
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
|
||||
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
|
||||
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
|
||||
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
|
||||
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
|
||||
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
THE SOFTWARE.
|
||||
*/
|
||||
|
||||
#include "hip/hcc_detail/grid_launch_GGL.hpp"
|
||||
|
||||
// Internal header, do not percolate upwards.
|
||||
#include "hip_hcc_internal.h"
|
||||
#include "hc.hpp"
|
||||
#include "trace_helper.h"
|
||||
|
||||
namespace hip_impl
|
||||
{
|
||||
hc::accelerator_view lock_stream_hip_(
|
||||
hipStream_t& stream, void*& locked_stream)
|
||||
{ // This allocated but does not take ownership of locked_stream. If it is
|
||||
// not deleted elsewhere it will leak.
|
||||
using L = decltype(stream->lockopen_preKernelCommand());
|
||||
|
||||
HIP_INIT();
|
||||
|
||||
stream = ihipSyncAndResolveStream(stream);
|
||||
locked_stream = new L{stream->lockopen_preKernelCommand()};
|
||||
return (*static_cast<L*>(locked_stream))->_av;
|
||||
}
|
||||
|
||||
void unlock_stream_hip_(
|
||||
hipStream_t stream,
|
||||
void* locked_stream,
|
||||
const char* kernel_name,
|
||||
hc::accelerator_view* acc_v)
|
||||
{ // Precondition: acc_v is the accelerator_view associated with stream
|
||||
// which is guarded by locked_stream;
|
||||
// locked_stream is deletable.
|
||||
using L = decltype(stream->lockopen_preKernelCommand());
|
||||
|
||||
stream->lockclose_postKernelCommand(kernel_name, acc_v);
|
||||
|
||||
delete static_cast<L*>(locked_stream);
|
||||
locked_stream = nullptr;
|
||||
}
|
||||
}
|
||||
@@ -1,5 +1,5 @@
|
||||
/*
|
||||
Copyright (c) 2015-2017 Advanced Micro Devices, Inc. All rights reserved.
|
||||
Copyright (c) 2015 - present 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
|
||||
@@ -26,7 +26,7 @@ THE SOFTWARE.
|
||||
#include <stack>
|
||||
|
||||
#include "hip/hip_runtime.h"
|
||||
#include "hip_hcc.h"
|
||||
#include "hip_hcc_internal.h"
|
||||
#include "trace_helper.h"
|
||||
|
||||
// Stack of contexts
|
||||
@@ -283,3 +283,79 @@ hipError_t hipCtxGetFlags ( unsigned int* flags )
|
||||
*flags = tempCtx->_ctxFlags;
|
||||
return ihipLogStatus(e);
|
||||
}
|
||||
|
||||
hipError_t hipDevicePrimaryCtxGetState ( hipDevice_t dev, unsigned int* flags, int* active )
|
||||
{
|
||||
HIP_INIT_API(dev, flags, active);
|
||||
hipError_t e = hipSuccess;
|
||||
auto deviceHandle = ihipGetDevice(dev);
|
||||
|
||||
if (deviceHandle == NULL) {
|
||||
e = hipErrorInvalidDevice;
|
||||
}
|
||||
|
||||
ihipCtx_t* tempCtx;
|
||||
tempCtx = ihipGetTlsDefaultCtx();
|
||||
ihipCtx_t* primaryCtx = deviceHandle->_primaryCtx;
|
||||
if(tempCtx == primaryCtx) {
|
||||
*active = 1;
|
||||
*flags = tempCtx->_ctxFlags;
|
||||
} else {
|
||||
*active = 0;
|
||||
*flags = primaryCtx->_ctxFlags;
|
||||
}
|
||||
return ihipLogStatus(e);
|
||||
}
|
||||
|
||||
hipError_t hipDevicePrimaryCtxRelease ( hipDevice_t dev)
|
||||
{
|
||||
HIP_INIT_API(dev);
|
||||
hipError_t e = hipSuccess;
|
||||
auto deviceHandle = ihipGetDevice(dev);
|
||||
|
||||
if (deviceHandle == NULL) {
|
||||
e = hipErrorInvalidDevice;
|
||||
}
|
||||
return ihipLogStatus(e);
|
||||
}
|
||||
|
||||
hipError_t hipDevicePrimaryCtxRetain ( hipCtx_t* pctx, hipDevice_t dev )
|
||||
{
|
||||
HIP_INIT_API(pctx, dev);
|
||||
hipError_t e = hipSuccess;
|
||||
auto deviceHandle = ihipGetDevice(dev);
|
||||
|
||||
if (deviceHandle == NULL) {
|
||||
e = hipErrorInvalidDevice;
|
||||
}
|
||||
*pctx = deviceHandle->_primaryCtx;
|
||||
return ihipLogStatus(e);
|
||||
}
|
||||
|
||||
hipError_t hipDevicePrimaryCtxReset ( hipDevice_t dev )
|
||||
{
|
||||
HIP_INIT_API(dev);
|
||||
hipError_t e = hipSuccess;
|
||||
auto deviceHandle = ihipGetDevice(dev);
|
||||
|
||||
if (deviceHandle == NULL) {
|
||||
e = hipErrorInvalidDevice;
|
||||
}
|
||||
ihipCtx_t* primaryCtx = deviceHandle->_primaryCtx;
|
||||
primaryCtx->locked_reset();
|
||||
return ihipLogStatus(e);
|
||||
}
|
||||
|
||||
hipError_t hipDevicePrimaryCtxSetFlags ( hipDevice_t dev, unsigned int flags )
|
||||
{
|
||||
HIP_INIT_API(dev, flags);
|
||||
hipError_t e = hipSuccess;
|
||||
auto deviceHandle = ihipGetDevice(dev);
|
||||
|
||||
if (deviceHandle == NULL) {
|
||||
e = hipErrorInvalidDevice;
|
||||
} else {
|
||||
e = hipErrorContextAlreadyInUse;
|
||||
}
|
||||
return ihipLogStatus(e);
|
||||
}
|
||||
|
||||
@@ -1,5 +1,5 @@
|
||||
/*
|
||||
Copyright (c) 2015-2017 Advanced Micro Devices, Inc. All rights reserved.
|
||||
Copyright (c) 2015 - present 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,7 +21,7 @@ THE SOFTWARE.
|
||||
*/
|
||||
|
||||
#include "hip/hip_runtime.h"
|
||||
#include "hip_hcc.h"
|
||||
#include "hip_hcc_internal.h"
|
||||
#include "trace_helper.h"
|
||||
#include "device_util.h"
|
||||
|
||||
|
||||
@@ -1,5 +1,5 @@
|
||||
/*
|
||||
Copyright (c) 2015-2017 Advanced Micro Devices, Inc. All rights reserved.
|
||||
Copyright (c) 2015 - present 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,7 +21,7 @@ THE SOFTWARE.
|
||||
*/
|
||||
|
||||
#include "hip/hip_runtime.h"
|
||||
#include "hip_hcc.h"
|
||||
#include "hip_hcc_internal.h"
|
||||
#include "trace_helper.h"
|
||||
|
||||
//-------------------------------------------------------------------------------------------------
|
||||
|
||||
@@ -1,5 +1,5 @@
|
||||
/*
|
||||
Copyright (c) 2015-2017 Advanced Micro Devices, Inc. All rights reserved.
|
||||
Copyright (c) 2015 - present 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,7 +21,7 @@ THE SOFTWARE.
|
||||
*/
|
||||
|
||||
#include "hip/hip_runtime.h"
|
||||
#include "hip_hcc.h"
|
||||
#include "hip_hcc_internal.h"
|
||||
#include "trace_helper.h"
|
||||
|
||||
//-------------------------------------------------------------------------------------------------
|
||||
|
||||
@@ -1,5 +1,5 @@
|
||||
/*
|
||||
Copyright (c) 2015-2017 Advanced Micro Devices, Inc. All rights reserved.
|
||||
Copyright (c) 2015 - present 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,7 +20,7 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
THE SOFTWARE.
|
||||
*/
|
||||
|
||||
#include"hip/hip_fp16.h"
|
||||
#include"hip/hcc_detail/hip_fp16.h"
|
||||
|
||||
struct hipHalfHolder{
|
||||
union {
|
||||
@@ -31,7 +31,7 @@ struct hipHalfHolder{
|
||||
|
||||
#define HINF 65504
|
||||
|
||||
static struct hipHalfHolder __hInfValue = {HINF};
|
||||
__device__ static struct hipHalfHolder __hInfValue = {HINF};
|
||||
|
||||
__device__ __half __hadd(__half a, __half b) {
|
||||
return a + b;
|
||||
|
||||
@@ -1,5 +1,5 @@
|
||||
/*
|
||||
Copyright (c) 2015-2017 Advanced Micro Devices, Inc. All rights reserved.
|
||||
Copyright (c) 2015 - present 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
|
||||
@@ -43,7 +43,7 @@ THE SOFTWARE.
|
||||
#include "hsa/hsa_ext_amd.h"
|
||||
|
||||
#include "hip/hip_runtime.h"
|
||||
#include "hip_hcc.h"
|
||||
#include "hip_hcc_internal.h"
|
||||
#include "trace_helper.h"
|
||||
#include "env.h"
|
||||
|
||||
@@ -762,7 +762,7 @@ hipError_t ihipDevice_t::initProperties(hipDeviceProp_t* prop)
|
||||
err = hsa_agent_get_info(_hsaAgent, (hsa_agent_info_t)HSA_AMD_AGENT_INFO_PRODUCT_NAME, &(prop->name));
|
||||
char archName[256];
|
||||
err = hsa_agent_get_info(_hsaAgent, HSA_AGENT_INFO_NAME, &archName);
|
||||
|
||||
|
||||
if(strcmp(archName,"gfx701")==0){
|
||||
prop->gcnArch = 701;
|
||||
}
|
||||
@@ -1556,7 +1556,6 @@ void ihipPostLaunchKernel(const char *kernelName, hipStream_t stream, grid_launc
|
||||
MARKER_END();
|
||||
}
|
||||
|
||||
|
||||
//=================================================================================================
|
||||
// HIP API Implementation
|
||||
//
|
||||
@@ -1806,7 +1805,7 @@ void ihipStream_t::resolveHcMemcpyDirection(unsigned hipMemKind,
|
||||
void printPointerInfo(unsigned dbFlag, const char *tag, const void *ptr, const hc::AmPointerInfo &ptrInfo)
|
||||
{
|
||||
tprintf (dbFlag, " %s=%p baseHost=%p baseDev=%p sz=%zu home_dev=%d tracked=%d isDevMem=%d registered=%d\n",
|
||||
tag, ptr,
|
||||
tag, ptr,
|
||||
ptrInfo._hostPointer, ptrInfo._devicePointer, ptrInfo._sizeBytes,
|
||||
ptrInfo._appId, ptrInfo._sizeBytes != 0, ptrInfo._isInDeviceMem, !ptrInfo._isAmManaged);
|
||||
}
|
||||
@@ -1873,31 +1872,47 @@ void ihipStream_t::addSymbolPtrToTracker(hc::accelerator& acc, void* ptr, size_t
|
||||
hc::am_memtracker_add(ptr, ptrInfo);
|
||||
}
|
||||
|
||||
void ihipStream_t::lockedSymbolCopySync(hc::accelerator &acc, void* dst, void* src, size_t sizeBytes, unsigned kind)
|
||||
void ihipStream_t::lockedSymbolCopySync(hc::accelerator &acc, void* dst, void* src, size_t sizeBytes, size_t offset, unsigned kind)
|
||||
{
|
||||
if(kind == hipMemcpyHostToHost){
|
||||
acc.memcpy_symbol(dst, (void*)src, sizeBytes, Kalmar::hcMemcpyHostToHost);
|
||||
acc.memcpy_symbol(dst, (void*)src, sizeBytes, offset, Kalmar::hcMemcpyHostToHost);
|
||||
}
|
||||
if(kind == hipMemcpyHostToDevice){
|
||||
acc.memcpy_symbol(dst, (void*)src, sizeBytes);
|
||||
acc.memcpy_symbol(dst, (void*)src, sizeBytes, offset);
|
||||
}
|
||||
if(kind == hipMemcpyDeviceToDevice){
|
||||
acc.memcpy_symbol(dst, (void*)src, sizeBytes, Kalmar::hcMemcpyDeviceToDevice);
|
||||
acc.memcpy_symbol(dst, (void*)src, sizeBytes, offset, Kalmar::hcMemcpyDeviceToDevice);
|
||||
}
|
||||
if(kind == hipMemcpyDeviceToHost){
|
||||
acc.memcpy_symbol(dst, (void*)src, sizeBytes, Kalmar::hcMemcpyDeviceToHost);
|
||||
acc.memcpy_symbol((void*) src, (void*)dst, sizeBytes, offset, Kalmar::hcMemcpyDeviceToHost);
|
||||
}
|
||||
}
|
||||
|
||||
void ihipStream_t::lockedSymbolCopyAsync(hc::accelerator &acc, void* dst, void* src, size_t sizeBytes, unsigned kind)
|
||||
void ihipStream_t::lockedSymbolCopyAsync(hc::accelerator &acc, void* dst, void* src, size_t sizeBytes, size_t offset, unsigned kind)
|
||||
{
|
||||
if(kind == hipMemcpyHostToDevice) {
|
||||
addSymbolPtrToTracker(acc, dst, sizeBytes);
|
||||
locked_getAv()->copy_async((void*)src, dst, sizeBytes);
|
||||
hc::AmPointerInfo srcPtrInfo(NULL, NULL, 0, acc, 0, 0);
|
||||
bool srcTracked = (hc::am_memtracker_getinfo(&srcPtrInfo, src) == AM_SUCCESS);
|
||||
if(srcTracked) {
|
||||
addSymbolPtrToTracker(acc, dst, sizeBytes);
|
||||
locked_getAv()->copy_async((void*)src, dst, sizeBytes);
|
||||
} else {
|
||||
LockedAccessor_StreamCrit_t crit(_criticalData);
|
||||
this->wait(crit);
|
||||
acc.memcpy_symbol(dst, (void*)src, sizeBytes, offset);
|
||||
}
|
||||
}
|
||||
if(kind == hipMemcpyDeviceToHost) {
|
||||
addSymbolPtrToTracker(acc, src, sizeBytes);
|
||||
locked_getAv()->copy_async((void*)src, dst, sizeBytes);
|
||||
hc::AmPointerInfo dstPtrInfo(NULL, NULL, 0, acc, 0, 0);
|
||||
bool dstTracked = (hc::am_memtracker_getinfo(&dstPtrInfo, dst) == AM_SUCCESS);
|
||||
if(dstTracked) {
|
||||
addSymbolPtrToTracker(acc, src, sizeBytes);
|
||||
locked_getAv()->copy_async((void*)src, dst, sizeBytes);
|
||||
} else {
|
||||
LockedAccessor_StreamCrit_t crit(_criticalData);
|
||||
this->wait(crit);
|
||||
acc.memcpy_symbol((void*)src, (void*)dst, sizeBytes, offset, Kalmar::hcMemcpyDeviceToHost);
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
@@ -1,5 +1,5 @@
|
||||
/*
|
||||
Copyright (c) 2015-2017 Advanced Micro Devices, Inc. All rights reserved.
|
||||
Copyright (c) 2015 - present 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_HCC_H
|
||||
#define HIP_HCC_H
|
||||
#ifndef HIP_SRC_HIP_HCC_INTERNAL_H
|
||||
#define HIP_SRC_HIP_HCC_INTERNAL_H
|
||||
|
||||
#include <hc.hpp>
|
||||
#include <hsa/hsa.h>
|
||||
@@ -503,8 +503,8 @@ public:
|
||||
void locked_copySync (void* dst, const void* src, size_t sizeBytes, unsigned kind, bool resolveOn = true);
|
||||
void locked_copyAsync(void* dst, const void* src, size_t sizeBytes, unsigned kind);
|
||||
|
||||
void lockedSymbolCopySync(hc::accelerator &acc, void *dst, void* src, size_t sizeBytes, unsigned kind);
|
||||
void lockedSymbolCopyAsync(hc::accelerator &acc, void *dst, void* src, size_t sizeBytes, unsigned kind);
|
||||
void lockedSymbolCopySync(hc::accelerator &acc, void *dst, void* src, size_t sizeBytes, size_t offset, unsigned kind);
|
||||
void lockedSymbolCopyAsync(hc::accelerator &acc, void *dst, void* src, size_t sizeBytes, size_t offset, unsigned kind);
|
||||
|
||||
//---
|
||||
// Member functions that begin with locked_ are thread-safe accessors - these acquire / release the critical mutex.
|
||||
@@ -1,5 +1,5 @@
|
||||
/*
|
||||
Copyright (c) 2015-2017 Advanced Micro Devices, Inc. All rights reserved.
|
||||
Copyright (c) 2015 - present 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
|
||||
|
||||
@@ -1,5 +1,5 @@
|
||||
/*
|
||||
Copyright (c) 2015-2017 Advanced Micro Devices, Inc. All rights reserved.
|
||||
Copyright (c) 2015 - present 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
|
||||
@@ -25,7 +25,7 @@ THE SOFTWARE.
|
||||
#include "hsa/hsa_ext_amd.h"
|
||||
|
||||
#include "hip/hip_runtime.h"
|
||||
#include "hip_hcc.h"
|
||||
#include "hip_hcc_internal.h"
|
||||
#include "trace_helper.h"
|
||||
#include "hip/hcc_detail/hip_texture.h"
|
||||
#include <hc_am.hpp>
|
||||
@@ -261,11 +261,11 @@ hipError_t hipHostMalloc(void** ptr, size_t sizeBytes, unsigned int flags)
|
||||
auto device = ctx->getWriteableDevice();
|
||||
unsigned amFlags = HIP_COHERENT_HOST_ALLOC ? amHostCoherent : amHostPinned;
|
||||
|
||||
*ptr = hip_internal::allocAndSharePtr(HIP_COHERENT_HOST_ALLOC ? "finegrained_host":"pinned_host",
|
||||
*ptr = hip_internal::allocAndSharePtr(HIP_COHERENT_HOST_ALLOC ? "finegrained_host":"pinned_host",
|
||||
sizeBytes, ctx, amFlags, flags);
|
||||
if(sizeBytes && (*ptr == NULL)){
|
||||
hip_status = hipErrorMemoryAllocation;
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
@@ -314,7 +314,7 @@ hipError_t hipMallocPitch(void** ptr, size_t* pitch, size_t width, size_t height
|
||||
|
||||
if (sizeBytes && (*ptr == NULL)) {
|
||||
hip_status = hipErrorMemoryAllocation;
|
||||
}
|
||||
}
|
||||
} else {
|
||||
hip_status = hipErrorMemoryAllocation;
|
||||
}
|
||||
@@ -372,7 +372,7 @@ hipError_t hipMallocArray(hipArray** array, const hipChannelFormatDesc* desc,
|
||||
*ptr = hip_internal::allocAndSharePtr("device_array", allocSize, ctx, am_flags, 0);
|
||||
if (size && (*ptr == NULL)) {
|
||||
hip_status = hipErrorMemoryAllocation;
|
||||
}
|
||||
}
|
||||
|
||||
} else {
|
||||
hip_status = hipErrorMemoryAllocation;
|
||||
@@ -472,7 +472,7 @@ hipError_t hipHostUnregister(void *hostPtr)
|
||||
return ihipLogStatus(hip_status);
|
||||
}
|
||||
|
||||
hipError_t hipMemcpyToSymbol(const char* symbolName, const void *src, size_t count, size_t offset, hipMemcpyKind kind)
|
||||
hipError_t hipMemcpyToSymbol(const void* symbolName, const void *src, size_t count, size_t offset, hipMemcpyKind kind)
|
||||
{
|
||||
HIP_INIT_CMD_API(symbolName, src, count, offset, kind);
|
||||
|
||||
@@ -485,7 +485,7 @@ hipError_t hipMemcpyToSymbol(const char* symbolName, const void *src, size_t cou
|
||||
|
||||
hc::accelerator acc = ctx->getDevice()->_acc;
|
||||
|
||||
void *dst = acc.get_symbol_address(symbolName);
|
||||
void *dst = acc.get_symbol_address((const char*) symbolName);
|
||||
tprintf(DB_MEM, " symbol '%s' resolved to address:%p\n", symbolName, dst);
|
||||
|
||||
if(dst == nullptr)
|
||||
@@ -497,7 +497,7 @@ hipError_t hipMemcpyToSymbol(const char* symbolName, const void *src, size_t cou
|
||||
|
||||
if(kind == hipMemcpyHostToDevice || kind == hipMemcpyDeviceToHost || kind == hipMemcpyDeviceToDevice || kind == hipMemcpyHostToHost)
|
||||
{
|
||||
stream->lockedSymbolCopySync(acc, dst, (void*)src, count + offset, kind);
|
||||
stream->lockedSymbolCopySync(acc, dst, (void*)src, count, offset, kind);
|
||||
// acc.memcpy_symbol(dst, (void*)src, count+offset);
|
||||
} else {
|
||||
return ihipLogStatus(hipErrorInvalidValue);
|
||||
@@ -507,7 +507,7 @@ hipError_t hipMemcpyToSymbol(const char* symbolName, const void *src, size_t cou
|
||||
}
|
||||
|
||||
|
||||
hipError_t hipMemcpyFromSymbol(void* dst, const char* symbolName, size_t count, size_t offset, hipMemcpyKind kind)
|
||||
hipError_t hipMemcpyFromSymbol(void* dst, const void* symbolName, size_t count, size_t offset, hipMemcpyKind kind)
|
||||
{
|
||||
HIP_INIT_CMD_API(symbolName, dst, count, offset, kind);
|
||||
|
||||
@@ -520,7 +520,7 @@ hipError_t hipMemcpyFromSymbol(void* dst, const char* symbolName, size_t count,
|
||||
|
||||
hc::accelerator acc = ctx->getDevice()->_acc;
|
||||
|
||||
void *src = acc.get_symbol_address(symbolName);
|
||||
void *src = acc.get_symbol_address((const char*) symbolName);
|
||||
tprintf(DB_MEM, " symbol '%s' resolved to address:%p\n", symbolName, dst);
|
||||
|
||||
if(dst == nullptr)
|
||||
@@ -532,7 +532,7 @@ hipError_t hipMemcpyFromSymbol(void* dst, const char* symbolName, size_t count,
|
||||
|
||||
if(kind == hipMemcpyHostToDevice || kind == hipMemcpyDeviceToHost || kind == hipMemcpyDeviceToDevice || kind == hipMemcpyHostToHost)
|
||||
{
|
||||
stream->lockedSymbolCopySync(acc, dst, (void*)src, count + offset, kind);
|
||||
stream->lockedSymbolCopySync(acc, dst, (void*)src, count, offset, kind);
|
||||
}
|
||||
else {
|
||||
return ihipLogStatus(hipErrorInvalidValue);
|
||||
@@ -542,7 +542,7 @@ hipError_t hipMemcpyFromSymbol(void* dst, const char* symbolName, size_t count,
|
||||
}
|
||||
|
||||
|
||||
hipError_t hipMemcpyToSymbolAsync(const char* symbolName, const void *src, size_t count, size_t offset, hipMemcpyKind kind, hipStream_t stream)
|
||||
hipError_t hipMemcpyToSymbolAsync(const void* symbolName, const void *src, size_t count, size_t offset, hipMemcpyKind kind, hipStream_t stream)
|
||||
{
|
||||
HIP_INIT_CMD_API(symbolName, src, count, offset, kind, stream);
|
||||
|
||||
@@ -557,7 +557,7 @@ hipError_t hipMemcpyToSymbolAsync(const char* symbolName, const void *src, size_
|
||||
|
||||
hc::accelerator acc = ctx->getDevice()->_acc;
|
||||
|
||||
void *dst = acc.get_symbol_address(symbolName);
|
||||
void *dst = acc.get_symbol_address((const char*) symbolName);
|
||||
tprintf(DB_MEM, " symbol '%s' resolved to address:%p\n", symbolName, dst);
|
||||
|
||||
if(dst == nullptr)
|
||||
@@ -567,7 +567,7 @@ hipError_t hipMemcpyToSymbolAsync(const char* symbolName, const void *src, size_
|
||||
|
||||
if (stream) {
|
||||
try {
|
||||
stream->lockedSymbolCopyAsync(acc, dst, (void*)src, count + offset, kind);
|
||||
stream->lockedSymbolCopyAsync(acc, dst, (void*)src, count, offset, kind);
|
||||
}
|
||||
catch (ihipException ex) {
|
||||
e = ex._code;
|
||||
@@ -580,7 +580,7 @@ hipError_t hipMemcpyToSymbolAsync(const char* symbolName, const void *src, size_
|
||||
}
|
||||
|
||||
|
||||
hipError_t hipMemcpyFromSymbolAsync(void* dst, const char* symbolName, size_t count, size_t offset, hipMemcpyKind kind, hipStream_t stream)
|
||||
hipError_t hipMemcpyFromSymbolAsync(void* dst, const void* symbolName, size_t count, size_t offset, hipMemcpyKind kind, hipStream_t stream)
|
||||
{
|
||||
HIP_INIT_CMD_API(symbolName, dst, count, offset, kind, stream);
|
||||
|
||||
@@ -595,7 +595,7 @@ hipError_t hipMemcpyFromSymbolAsync(void* dst, const char* symbolName, size_t co
|
||||
|
||||
hc::accelerator acc = ctx->getDevice()->_acc;
|
||||
|
||||
void *src = acc.get_symbol_address(symbolName);
|
||||
void *src = acc.get_symbol_address((const char*) symbolName);
|
||||
tprintf(DB_MEM, " symbol '%s' resolved to address:%p\n", symbolName, src);
|
||||
|
||||
if(src == nullptr || dst == nullptr)
|
||||
@@ -603,9 +603,10 @@ hipError_t hipMemcpyFromSymbolAsync(void* dst, const char* symbolName, size_t co
|
||||
return ihipLogStatus(hipErrorInvalidSymbol);
|
||||
}
|
||||
|
||||
stream = ihipSyncAndResolveStream(stream);
|
||||
if (stream) {
|
||||
try {
|
||||
stream->lockedSymbolCopyAsync(acc, dst, src, count + offset, kind);
|
||||
stream->lockedSymbolCopyAsync(acc, dst, src, count, offset, kind);
|
||||
}
|
||||
catch (ihipException ex) {
|
||||
e = ex._code;
|
||||
|
||||
@@ -1,5 +1,5 @@
|
||||
/*
|
||||
Copyright (c) 2015-2017 Advanced Micro Devices, Inc. All rights reserved.
|
||||
Copyright (c) 2015 - present 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
|
||||
@@ -23,18 +23,45 @@ THE SOFTWARE.
|
||||
#include <fstream>
|
||||
#include <stdio.h>
|
||||
#include <stdlib.h>
|
||||
#include <libelf.h>
|
||||
#include <elf.h>
|
||||
#include <gelf.h>
|
||||
#include <map>
|
||||
|
||||
#include "hsa/hsa.h"
|
||||
#include "hsa/hsa_ext_amd.h"
|
||||
#include "hsa/amd_hsa_kernel_code.h"
|
||||
|
||||
#include "hip/hip_runtime.h"
|
||||
#include "hip_hcc.h"
|
||||
#include "hip_hcc_internal.h"
|
||||
#include "trace_helper.h"
|
||||
|
||||
//TODO Use Pool APIs from HCC to get memory regions.
|
||||
|
||||
#include <cassert>
|
||||
inline uint64_t alignTo(uint64_t Value, uint64_t Align, uint64_t Skew = 0) {
|
||||
assert(Align != 0u && "Align can't be 0.");
|
||||
Skew %= Align;
|
||||
return (Value + Align - 1 - Skew) / Align * Align + Skew;
|
||||
}
|
||||
|
||||
struct ihipKernArgInfo{
|
||||
std::vector<uint32_t> Size;
|
||||
std::vector<uint32_t> Align;
|
||||
std::vector<std::string> ArgType;
|
||||
std::vector<std::string> ArgName;
|
||||
uint32_t totalSize;
|
||||
};
|
||||
|
||||
std::map<std::string,struct ihipKernArgInfo> kernelArguments;
|
||||
|
||||
struct MyElfNote {
|
||||
uint32_t n_namesz = 0;
|
||||
uint32_t n_descsz = 0;
|
||||
uint32_t n_type = 0;
|
||||
|
||||
MyElfNote() = default;
|
||||
};
|
||||
|
||||
struct ihipModuleSymbol_t{
|
||||
uint64_t _object; // The kernel object.
|
||||
@@ -172,9 +199,57 @@ hipError_t hipModuleLoad(hipModule_t *module, const char *fname){
|
||||
(*module)->ptr = p;
|
||||
(*module)->size = size;
|
||||
in.seekg(0, std::ios::beg);
|
||||
std::copy(std::istreambuf_iterator<char>(in),
|
||||
std::istreambuf_iterator<char>(), ptr);
|
||||
std::copy(std::istreambuf_iterator<char>(in), std::istreambuf_iterator<char>(), ptr);
|
||||
/*
|
||||
Elf *e = elf_memory((char*)p, size);
|
||||
if(elf_kind(e) != ELF_K_ELF){
|
||||
return ihipLogStatus(hipErrorInvalidValue);
|
||||
}
|
||||
size_t numpHdrs;
|
||||
if(elf_getphdrnum(e, &numpHdrs) != 0){
|
||||
return ihipLogStatus(hipErrorInvalidValue);
|
||||
}
|
||||
for(size_t i=0;i<numpHdrs;i++){
|
||||
GElf_Phdr pHdr;
|
||||
if(gelf_getphdr(e, i, &pHdr) != &pHdr){
|
||||
continue;
|
||||
}
|
||||
if(pHdr.p_type == PT_NOTE && pHdr.p_align >= sizeof(int)){
|
||||
char* ptr = (char*) p + pHdr.p_offset;
|
||||
char* segmentEnd = ptr + pHdr.p_filesz;
|
||||
while(ptr < segmentEnd){
|
||||
MyElfNote *note = (MyElfNote*) ptr;
|
||||
char *name = (char*) ¬e[1];
|
||||
char *desc = name + alignTo(note->n_namesz, sizeof(int));
|
||||
if (note->n_type == 8) {
|
||||
std::string metadatastr((const char*)desc, (size_t)note->n_descsz);
|
||||
AMDGPU::RuntimeMD::Program::Metadata meta;
|
||||
meta = meta.fromYAML(metadatastr);
|
||||
for(int i=0;i<meta.Kernels.size();i++){
|
||||
struct ihipKernArgInfo pl;
|
||||
size_t totalSizeOfArgs = 0;
|
||||
for(int j=0;j<meta.Kernels[i].Args.size();j++){
|
||||
if(meta.Kernels[i].Args[j].Kind < 7) {
|
||||
pl.Size.push_back(meta.Kernels[i].Args[j].Size);
|
||||
pl.Align.push_back(meta.Kernels[i].Args[j].Align);
|
||||
pl.ArgType.push_back(meta.Kernels[i].Args[j].TypeName);
|
||||
pl.ArgName.push_back(meta.Kernels[i].Args[j].Name);
|
||||
totalSizeOfArgs += meta.Kernels[i].Args[j].Align;
|
||||
}
|
||||
}
|
||||
pl.totalSize = totalSizeOfArgs;
|
||||
kernelArguments[meta.Kernels[i].Name] = pl;
|
||||
}
|
||||
break;
|
||||
}
|
||||
|
||||
ptr += sizeof(*note)
|
||||
+ alignTo(note->n_namesz, sizeof(int))
|
||||
+ alignTo(note->n_descsz, sizeof(int));
|
||||
}
|
||||
}
|
||||
}
|
||||
*/
|
||||
status = hsa_code_object_deserialize(ptr, size, NULL, &(*module)->object);
|
||||
|
||||
if(status != HSA_STATUS_SUCCESS){
|
||||
@@ -288,16 +363,12 @@ hipError_t hipModuleGetFunction(hipFunction_t *hfunc, hipModule_t hmod,
|
||||
}
|
||||
|
||||
|
||||
hipError_t hipModuleLaunchKernel(hipFunction_t f,
|
||||
uint32_t gridDimX, uint32_t gridDimY, uint32_t gridDimZ,
|
||||
uint32_t blockDimX, uint32_t blockDimY, uint32_t blockDimZ,
|
||||
uint32_t sharedMemBytes, hipStream_t hStream,
|
||||
hipError_t ihipModuleLaunchKernel(hipFunction_t f,
|
||||
uint32_t globalWorkSizeX, uint32_t globalWorkSizeY, uint32_t globalWorkSizeZ,
|
||||
uint32_t localWorkSizeX, uint32_t localWorkSizeY, uint32_t localWorkSizeZ,
|
||||
size_t sharedMemBytes, hipStream_t hStream,
|
||||
void **kernelParams, void **extra)
|
||||
{
|
||||
HIP_INIT_API(f, gridDimX, gridDimY, gridDimZ,
|
||||
blockDimX, blockDimY, blockDimZ,
|
||||
sharedMemBytes, hStream,
|
||||
kernelParams, extra);
|
||||
|
||||
auto ctx = ihipGetTlsDefaultCtx();
|
||||
hipError_t ret = hipSuccess;
|
||||
@@ -313,13 +384,26 @@ hipError_t hipModuleLaunchKernel(hipFunction_t f,
|
||||
void *config[5] = {0};
|
||||
size_t kernArgSize;
|
||||
|
||||
if(extra != NULL){
|
||||
if(kernelParams != NULL){
|
||||
std::string name = f->_name;
|
||||
struct ihipKernArgInfo pl = kernelArguments[name];
|
||||
char* argBuf = (char*)malloc(pl.totalSize);
|
||||
memset(argBuf, 0, pl.totalSize);
|
||||
int index = 0;
|
||||
for(int i=0;i<pl.Size.size();i++){
|
||||
memcpy(argBuf + index, kernelParams[i], pl.Size[i]);
|
||||
index += pl.Align[i];
|
||||
}
|
||||
config[1] = (void*)argBuf;
|
||||
kernArgSize = pl.totalSize;
|
||||
} else if(extra != NULL){
|
||||
memcpy(config, extra, sizeof(size_t)*5);
|
||||
if(config[0] == HIP_LAUNCH_PARAM_BUFFER_POINTER && config[2] == HIP_LAUNCH_PARAM_BUFFER_SIZE && config[4] == HIP_LAUNCH_PARAM_END){
|
||||
kernArgSize = *(size_t*)(config[3]);
|
||||
} else {
|
||||
return ihipLogStatus(hipErrorNotInitialized);
|
||||
}
|
||||
|
||||
}else{
|
||||
return ihipLogStatus(hipErrorInvalidValue);
|
||||
}
|
||||
@@ -331,7 +415,7 @@ hipError_t hipModuleLaunchKernel(hipFunction_t f,
|
||||
*/
|
||||
grid_launch_parm lp;
|
||||
lp.dynamic_group_mem_bytes = sharedMemBytes; // TODO - this should be part of preLaunchKernel.
|
||||
hStream = ihipPreLaunchKernel(hStream, dim3(gridDimX, gridDimY, gridDimZ), dim3(blockDimX, blockDimY, blockDimZ), &lp, f->_name.c_str());
|
||||
hStream = ihipPreLaunchKernel(hStream, dim3(globalWorkSizeX, globalWorkSizeY, globalWorkSizeZ), dim3(localWorkSizeX, localWorkSizeY, localWorkSizeZ), &lp, f->_name.c_str());
|
||||
|
||||
|
||||
hsa_kernel_dispatch_packet_t aql;
|
||||
@@ -341,12 +425,12 @@ hipError_t hipModuleLaunchKernel(hipFunction_t f,
|
||||
//aql.completion_signal._handle = 0;
|
||||
//aql.kernarg_address = 0;
|
||||
|
||||
aql.workgroup_size_x = blockDimX;
|
||||
aql.workgroup_size_y = blockDimY;
|
||||
aql.workgroup_size_z = blockDimZ;
|
||||
aql.grid_size_x = blockDimX * gridDimX;
|
||||
aql.grid_size_y = blockDimY * gridDimY;
|
||||
aql.grid_size_z = blockDimZ * gridDimZ;
|
||||
aql.workgroup_size_x = localWorkSizeX;
|
||||
aql.workgroup_size_y = localWorkSizeY;
|
||||
aql.workgroup_size_z = localWorkSizeZ;
|
||||
aql.grid_size_x = globalWorkSizeX;
|
||||
aql.grid_size_y = globalWorkSizeY;
|
||||
aql.grid_size_z = globalWorkSizeZ;
|
||||
aql.group_segment_size = f->_groupSegmentSize + sharedMemBytes;
|
||||
aql.private_segment_size = f->_privateSegmentSize;
|
||||
aql.kernel_object = f->_object;
|
||||
@@ -364,13 +448,46 @@ hipError_t hipModuleLaunchKernel(hipFunction_t f,
|
||||
|
||||
lp.av->dispatch_hsa_kernel(&aql, config[1] /* kernarg*/, kernArgSize, nullptr/*completion_future*/);
|
||||
|
||||
|
||||
if(kernelParams != NULL){
|
||||
free(config[1]);
|
||||
}
|
||||
ihipPostLaunchKernel(f->_name.c_str(), hStream, lp);
|
||||
}
|
||||
|
||||
return ihipLogStatus(ret);
|
||||
return ret;
|
||||
}
|
||||
|
||||
hipError_t hipModuleLaunchKernel(hipFunction_t f,
|
||||
uint32_t gridDimX, uint32_t gridDimY, uint32_t gridDimZ,
|
||||
uint32_t blockDimX, uint32_t blockDimY, uint32_t blockDimZ,
|
||||
uint32_t sharedMemBytes, hipStream_t hStream,
|
||||
void **kernelParams, void **extra)
|
||||
{
|
||||
HIP_INIT_API(f, gridDimX, gridDimY, gridDimZ,
|
||||
blockDimX, blockDimY, blockDimZ,
|
||||
sharedMemBytes, hStream,
|
||||
kernelParams, extra);
|
||||
return ihipLogStatus(ihipModuleLaunchKernel(f,
|
||||
blockDimX * gridDimX, blockDimY * gridDimY, gridDimZ * blockDimZ,
|
||||
blockDimX, blockDimY, blockDimZ,
|
||||
sharedMemBytes, hStream, kernelParams, extra));
|
||||
}
|
||||
|
||||
|
||||
hipError_t hipHccModuleLaunchKernel(hipFunction_t f,
|
||||
uint32_t globalWorkSizeX, uint32_t globalWorkSizeY, uint32_t globalWorkSizeZ,
|
||||
uint32_t localWorkSizeX, uint32_t localWorkSizeY, uint32_t localWorkSizeZ,
|
||||
size_t sharedMemBytes, hipStream_t hStream,
|
||||
void **kernelParams, void **extra)
|
||||
{
|
||||
HIP_INIT_API(f, globalWorkSizeX, globalWorkSizeY, globalWorkSizeZ,
|
||||
localWorkSizeX, localWorkSizeY, localWorkSizeZ,
|
||||
sharedMemBytes, hStream,
|
||||
kernelParams, extra);
|
||||
return ihipLogStatus(ihipModuleLaunchKernel(f, globalWorkSizeX, globalWorkSizeY, globalWorkSizeZ,
|
||||
localWorkSizeX, localWorkSizeY, localWorkSizeZ,
|
||||
sharedMemBytes, hStream, kernelParams, extra));
|
||||
}
|
||||
|
||||
hipError_t hipModuleGetGlobal(hipDeviceptr_t *dptr, size_t *bytes,
|
||||
hipModule_t hmod, const char* name)
|
||||
|
||||
@@ -1,5 +1,5 @@
|
||||
/*
|
||||
Copyright (c) 2015-2017 Advanced Micro Devices, Inc. All rights reserved.
|
||||
Copyright (c) 2015 - present 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
|
||||
@@ -23,7 +23,7 @@ THE SOFTWARE.
|
||||
#include <hc_am.hpp>
|
||||
|
||||
#include "hip/hip_runtime.h"
|
||||
#include "hip_hcc.h"
|
||||
#include "hip_hcc_internal.h"
|
||||
#include "trace_helper.h"
|
||||
|
||||
|
||||
|
||||
@@ -1,5 +1,5 @@
|
||||
/*
|
||||
Copyright (c) 2015-2017 Advanced Micro Devices, Inc. All rights reserved.
|
||||
Copyright (c) 2015 - present 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,7 +21,7 @@ THE SOFTWARE.
|
||||
*/
|
||||
|
||||
#include "hip/hip_runtime.h"
|
||||
#include "hip_hcc.h"
|
||||
#include "hip_hcc_internal.h"
|
||||
#include "trace_helper.h"
|
||||
|
||||
|
||||
|
||||
@@ -1,5 +1,5 @@
|
||||
/*
|
||||
Copyright (c) 2015-2017 Advanced Micro Devices, Inc. All rights reserved.
|
||||
Copyright (c) 2015 - present 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_UTIL_H
|
||||
#define HIP_UTIL_H
|
||||
#ifndef HIP_INCLUDE_HCC_DETAIL_HIP_UTIL_H
|
||||
#define HIP_INCLUDE_HCC_DETAIL_HIP_UTIL_H
|
||||
|
||||
#include <assert.h>
|
||||
#include <stdint.h>
|
||||
|
||||
@@ -188,8 +188,12 @@ __device__ float ldexpf(float x, int exp)
|
||||
}
|
||||
__device__ float lgammaf(float x)
|
||||
{
|
||||
int sign;
|
||||
return hc::precise_math::lgammaf(x, &sign);
|
||||
float val = 0.0f;
|
||||
float y = x - 1;
|
||||
while(y > 0){
|
||||
val += logf(y--);
|
||||
}
|
||||
return val;
|
||||
}
|
||||
__device__ long long int llrintf(float x)
|
||||
{
|
||||
@@ -202,7 +206,8 @@ __device__ long long int llroundf(float x)
|
||||
int y = hc::precise_math::roundf(x);
|
||||
long long int z = y;
|
||||
return z;
|
||||
}__device__ float log10f(float x)
|
||||
}
|
||||
__device__ float log10f(float x)
|
||||
{
|
||||
return hc::precise_math::log10f(x);
|
||||
}
|
||||
@@ -569,8 +574,12 @@ __device__ double ldexp(double x, int exp)
|
||||
}
|
||||
__device__ double lgamma(double x)
|
||||
{
|
||||
int sign;
|
||||
return hc::precise_math::lgamma(x, &sign);
|
||||
double val = 0.0;
|
||||
double y = x - 1;
|
||||
while(y > 0){
|
||||
val += log(y--);
|
||||
}
|
||||
return val;
|
||||
}
|
||||
__device__ long long int llrint(double x)
|
||||
{
|
||||
|
||||
@@ -28,7 +28,6 @@ THE SOFTWARE.
|
||||
#include <iostream>
|
||||
#include <iomanip>
|
||||
#include <string>
|
||||
|
||||
//---
|
||||
// Helper functions to convert HIP function arguments into strings.
|
||||
// Handles POD data types as well as enumerations (ie hipMemcpyKind).
|
||||
@@ -71,7 +70,6 @@ inline std::string ToString(hipEvent_t v)
|
||||
ss << v;
|
||||
return ss.str();
|
||||
};
|
||||
|
||||
// hipStream_t
|
||||
template <>
|
||||
inline std::string ToString(hipStream_t v)
|
||||
|
||||
@@ -59,6 +59,8 @@ __device__ void integer_intrinsics()
|
||||
__umulhi((unsigned int)1, (unsigned int)2);
|
||||
__urhadd((unsigned int)1, (unsigned int)2);
|
||||
__usad((unsigned int)1, (unsigned int)2, 0);
|
||||
|
||||
assert(1);
|
||||
}
|
||||
|
||||
__global__ void compileIntegerIntrinsics(hipLaunchParm lp, int ignored)
|
||||
|
||||
@@ -31,9 +31,10 @@ THE SOFTWARE.
|
||||
#define NUM 1024
|
||||
#define SIZE 1024*4
|
||||
|
||||
// TODO - collapse:
|
||||
#ifdef __HIP_PLATFORM_HCC__
|
||||
__attribute__((address_space(1))) int globalIn[NUM];
|
||||
__attribute__((address_space(1))) int globalOut[NUM];
|
||||
__device__ ADDRESS_SPACE_1 int globalIn[NUM];
|
||||
__device__ ADDRESS_SPACE_1 int globalOut[NUM];
|
||||
#endif
|
||||
|
||||
#ifdef __HIP_PLATFORM_NVCC__
|
||||
@@ -92,8 +93,29 @@ int main()
|
||||
hipMemcpyFromSymbol(C, HIP_SYMBOL(globalOut), SIZE, 0, hipMemcpyDeviceToHost);
|
||||
for(unsigned i=0;i<NUM;i++) {
|
||||
assert(A[i] == B[i]);
|
||||
// assert(A[i] == C[i]);
|
||||
assert(A[i] == C[i]);
|
||||
}
|
||||
|
||||
for(unsigned i=0;i<NUM;i++) {
|
||||
A[i] = -3*i;
|
||||
B[i] = 0;
|
||||
}
|
||||
|
||||
hipMemcpyToSymbolAsync(HIP_SYMBOL(globalIn), A, SIZE, 0, hipMemcpyHostToDevice, stream);
|
||||
hipStreamSynchronize(stream);
|
||||
hipLaunchKernel(Assign, dim3(1,1,1), dim3(NUM,1,1), 0, 0, Ad);
|
||||
hipMemcpy(B, Ad, SIZE, hipMemcpyDeviceToHost);
|
||||
hipMemcpyFromSymbolAsync(C, HIP_SYMBOL(globalOut), SIZE, 0, hipMemcpyDeviceToHost, stream);
|
||||
hipStreamSynchronize(stream);
|
||||
for(unsigned i=0;i<NUM;i++) {
|
||||
assert(A[i] == B[i]);
|
||||
assert(A[i] == C[i]);
|
||||
}
|
||||
hipHostFree(Am);
|
||||
hipHostFree(Cm);
|
||||
hipFree(Ad);
|
||||
delete[] A;
|
||||
delete[] B;
|
||||
delete[] C;
|
||||
passed();
|
||||
}
|
||||
|
||||
@@ -3874,7 +3874,22 @@ int main() {
|
||||
assert(sizeof(float2) == 8);
|
||||
assert(sizeof(float3) == 12);
|
||||
assert(sizeof(float4) == 16);
|
||||
bool *ptr;
|
||||
hipLaunchKernel(CheckVectorTypes, dim3(1,1,1), dim3(1,1,1), 0, 0, ptr);
|
||||
passed();
|
||||
|
||||
bool* ptr = nullptr;
|
||||
if (hipMalloc(&ptr, sizeof(bool)) != HIP_SUCCESS) return EXIT_FAILURE;
|
||||
std::unique_ptr<bool, decltype(hipFree)*> correct{ptr, hipFree};
|
||||
hipLaunchKernel(
|
||||
CheckVectorTypes, dim3(1,1,1), dim3(1,1,1), 0, 0, correct.get());
|
||||
bool passed = false;
|
||||
if (hipMemcpyDtoH(&passed, correct.get(), sizeof(bool)) != HIP_SUCCESS) {
|
||||
return EXIT_FAILURE;
|
||||
}
|
||||
|
||||
if (passed == true){
|
||||
std::cout << "PASSED" << std::endl;
|
||||
return 0;
|
||||
}
|
||||
else
|
||||
return EXIT_FAILURE;
|
||||
}
|
||||
|
||||
|
||||
@@ -30,7 +30,7 @@ THE SOFTWARE.
|
||||
#include <stdio.h>
|
||||
#include <iostream>
|
||||
#include "hip/hip_runtime.h"
|
||||
#include "hip/hcc.h"
|
||||
#include "hip/hip_hcc.h"
|
||||
#include "test_common.h"
|
||||
|
||||
#define CHECK(error) \
|
||||
@@ -61,4 +61,3 @@ int main(int argc, char *argv[])
|
||||
passed();
|
||||
|
||||
};
|
||||
|
||||
|
||||
مرجع در شماره جدید
Block a user