Merge 'master' into 'amd-master'
Change-Id: I7e307defaf54b3e95d97e8910bb61faba1021d3e
Tá an tiomantas seo le fáil i:
Jenkins
+48
-6
@@ -24,6 +24,7 @@ use Cwd 'abs_path';
|
||||
# CUDA_PATH : Path to CUDA SDK (default /usr/local/cuda). Used on NVIDIA platforms only.
|
||||
# HCC_HOME : Path to HCC SDK (default /opt/rocm/hcc). Used on AMD platforms only.
|
||||
# HSA_PATH : Path to HSA dir (default /opt/rocm/hsa). Used on AMD platforms only.
|
||||
# HIP_VDI_HOME : Path to HIP/VDI directory. Used on AMD platforms only.
|
||||
|
||||
if(scalar @ARGV == 0){
|
||||
print "No Arguments passed, exiting ...\n";
|
||||
@@ -53,6 +54,7 @@ $verbose = $ENV{'HIPCC_VERBOSE'} // 0;
|
||||
# Verbose: 0x1=commands, 0x2=paths, 0x4=hipcc args
|
||||
|
||||
$HIP_PATH=$ENV{'HIP_PATH'} // dirname (dirname $0); # use parent directory of hipcc
|
||||
$HIP_VDI_HOME=$ENV{'HIP_VDI_HOME'};
|
||||
$HIP_CLANG_PATH=$ENV{'HIP_CLANG_PATH'};
|
||||
$DEVICE_LIB_PATH=$ENV{'DEVICE_LIB_PATH'};
|
||||
|
||||
@@ -88,6 +90,18 @@ $HIP_PLATFORM= `$HIP_PATH/bin/hipconfig --platform` // "hcc";
|
||||
$HIP_VERSION= `$HIP_PATH/bin/hipconfig --version`;
|
||||
($HIP_VERSION_MAJOR, $HIP_VERSION_MINOR, $HIP_VERSION_PATCH) = split(/\./, $HIP_VERSION);
|
||||
|
||||
if (defined $HIP_VDI_HOME) {
|
||||
if (!defined $HIP_CLANG_PATH) {
|
||||
$HIP_CLANG_PATH = "$HIP_VDI_HOME/bin/x86_64";
|
||||
}
|
||||
if (!defined $DEVICE_LIB_PATH) {
|
||||
$DEVICE_LIB_PATH = "$HIP_VDI_HOME/lib/x86_64/bitcode";
|
||||
}
|
||||
$HIP_CLANG_INCLUDE_PATH = "$HIP_VDI_HOME/include/clang";
|
||||
$HIP_INCLUDE_PATH = "$HIP_VDI_HOME/include";
|
||||
$HIP_LIB_PATH = "$HIP_VDI_HOME/lib/x86_64";
|
||||
}
|
||||
|
||||
if (defined $HIP_CLANG_PATH) {
|
||||
$HIP_PLATFORM = "clang"
|
||||
}
|
||||
@@ -109,9 +123,6 @@ $target_gfx906 = 0;
|
||||
$default_amdgpu_target = 1;
|
||||
|
||||
if ($HIP_PLATFORM eq "clang") {
|
||||
if ($verbose & 0x2) {
|
||||
print ("HIP_CLANG_PATH=$HIP_CLANG_PATH\n");
|
||||
}
|
||||
$ROCM_PATH=$ENV{'ROCM_PATH'} // "/opt/rocm";
|
||||
$HIPCC="$HIP_CLANG_PATH/clang++";
|
||||
|
||||
@@ -125,9 +136,30 @@ if ($HIP_PLATFORM eq "clang") {
|
||||
$HIP_CLANG_VERSION=~/.*clang version ([^ ]+).*/;
|
||||
$HIP_CLANG_VERSION=$1;
|
||||
|
||||
$HIPCXXFLAGS .= " -std=c++11 -isystem $HIP_CLANG_PATH/../lib/clang/$HIP_CLANG_VERSION/include -I$HIP_PATH/include";
|
||||
$HIPLDFLAGS .= " --hip-link --hip-device-lib-path=$DEVICE_LIB_PATH -L$HIP_PATH/lib -lhip_hcc";
|
||||
if (! defined $HIP_CLANG_INCLUDE_PATH) {
|
||||
$HIP_CLANG_INCLUDE_PATH = "$HIP_CLANG_PATH/../lib/clang/$HIP_CLANG_VERSION/include";
|
||||
}
|
||||
if (! defined $HIP_INCLUDE_PATH) {
|
||||
$HIP_INCLUDE_PATH = "$HIP_PATH/include";
|
||||
}
|
||||
if (! defined $HIP_LIB_PATH) {
|
||||
$HIP_LIB_PATH = "$HIP_PATH/lib";
|
||||
}
|
||||
if ($verbose & 0x2) {
|
||||
if (defined $HIP_VDI_HOME) {
|
||||
print ("HIP_VDI_HOME=$HIP_VDI_HOME\n");
|
||||
}
|
||||
print ("HIP_CLANG_PATH=$HIP_CLANG_PATH\n");
|
||||
print ("HIP_CLANG_INCLUDE_PATH=$HIP_CLANG_INCLUDE_PATH\n");
|
||||
print ("HIP_INCLUDE_PATH=$HIP_INCLUDE_PATH\n");
|
||||
print ("HIP_LIB_PATH=$HIP_LIB_PATH\n");
|
||||
print ("DEVICE_LIB_PATH=$DEVICE_LIB_PATH\n");
|
||||
}
|
||||
|
||||
$HIPCXXFLAGS .= " -std=c++11 -isystem $HIP_CLANG_INCLUDE_PATH";
|
||||
$HIPLDFLAGS .= " --hip-device-lib-path=$DEVICE_LIB_PATH -L$HIP_LIB_PATH -Wl,--rpath=$HIP_LIB_PATH -lhip_hcc";
|
||||
} elsif ($HIP_PLATFORM eq "hcc") {
|
||||
$HIP_INCLUDE_PATH = "$HIP_PATH/include";
|
||||
$HSA_PATH=$ENV{'HSA_PATH'} // "/opt/rocm/hsa";
|
||||
|
||||
$HCC_HOME=$ENV{'HCC_HOME'} // $hipConfig{'HCC_HOME'} // "/opt/rocm/hcc";
|
||||
@@ -209,6 +241,7 @@ if ($HIP_PLATFORM eq "clang") {
|
||||
}
|
||||
|
||||
} elsif ($HIP_PLATFORM eq "nvcc") {
|
||||
$HIP_INCLUDE_PATH = "$HIP_PATH/include";
|
||||
if ($verbose & 0x2) {
|
||||
print ("CUDA_PATH=$CUDA_PATH\n");
|
||||
}
|
||||
@@ -225,7 +258,7 @@ if ($HIP_PLATFORM eq "clang") {
|
||||
}
|
||||
|
||||
# Add paths to common HIP includes:
|
||||
$HIPCXXFLAGS .= " -I$HIP_PATH/include -DHIP_VERSION_MAJOR=$HIP_VERSION_MAJOR -DHIP_VERSION_MINOR=$HIP_VERSION_MINOR -DHIP_VERSION_PATCH=$HIP_VERSION_PATCH" ;
|
||||
$HIPCXXFLAGS .= " -I$HIP_INCLUDE_PATH -DHIP_VERSION_MAJOR=$HIP_VERSION_MAJOR -DHIP_VERSION_MINOR=$HIP_VERSION_MINOR -DHIP_VERSION_PATCH=$HIP_VERSION_PATCH" ;
|
||||
|
||||
my $compileOnly = 0;
|
||||
my $needCXXFLAGS = 0; # need to add CXX flags to compile step
|
||||
@@ -602,6 +635,15 @@ if ($buildDeps and $HIP_PLATFORM eq 'nvcc') {
|
||||
$HIPCXXFLAGS .= " -M -D__CUDACC__";
|
||||
}
|
||||
|
||||
if ($buildDeps and $HIP_PLATFORM eq 'clang') {
|
||||
$HIPCXXFLAGS .= " --cuda-host-only";
|
||||
}
|
||||
|
||||
# Add --hip-link only if there are no source files.
|
||||
if (!$needCXXFLAGS and $HIP_PLATFORM eq 'clang') {
|
||||
$HIPLDFLAGS .= " --hip-link";
|
||||
}
|
||||
|
||||
if ($setStdLib eq 0 and $HIP_PLATFORM eq 'hcc')
|
||||
{
|
||||
$HIPCXXFLAGS .= $HCC_WA_FLAGS;
|
||||
|
||||
@@ -23,6 +23,27 @@ option(HIP_HOST_COMPILATION_CPP "Host code compilation mode" ON)
|
||||
option(HIP_VERBOSE_BUILD "Print out the commands run while compiling the HIP source file. With the Makefile generator this defaults to VERBOSE variable specified on the command line, but can be forced on with this option." OFF)
|
||||
mark_as_advanced(HIP_HOST_COMPILATION_CPP)
|
||||
|
||||
###############################################################################
|
||||
# Set HIP CMAKE Flags
|
||||
###############################################################################
|
||||
# Copy the invocation styles from CXX to HIP
|
||||
set(CMAKE_HIP_ARCHIVE_CREATE ${CMAKE_CXX_ARCHIVE_CREATE})
|
||||
set(CMAKE_HIP_ARCHIVE_APPEND ${CMAKE_CXX_ARCHIVE_APPEND})
|
||||
set(CMAKE_HIP_ARCHIVE_FINISH ${CMAKE_CXX_ARCHIVE_FINISH})
|
||||
set(CMAKE_SHARED_LIBRARY_SONAME_HIP_FLAG ${CMAKE_SHARED_LIBRARY_SONAME_CXX_FLAG})
|
||||
set(CMAKE_SHARED_LIBRARY_CREATE_HIP_FLAGS ${CMAKE_SHARED_LIBRARY_CREATE_CXX_FLAGS})
|
||||
set(CMAKE_SHARED_LIBRARY_HIP_FLAGS ${CMAKE_SHARED_LIBRARY_CXX_FLAGS})
|
||||
set(CMAKE_SHARED_LIBRARY_LINK_HIP_FLAGS ${CMAKE_SHARED_LIBRARY_LINK_CXX_FLAGS})
|
||||
set(CMAKE_SHARED_LIBRARY_RUNTIME_HIP_FLAG ${CMAKE_SHARED_LIBRARY_RUNTIME_CXX_FLAG})
|
||||
set(CMAKE_SHARED_LIBRARY_RUNTIME_HIP_FLAG_SEP ${CMAKE_SHARED_LIBRARY_RUNTIME_CXX_FLAG_SEP})
|
||||
set(CMAKE_SHARED_LIBRARY_LINK_STATIC_HIP_FLAGS ${CMAKE_SHARED_LIBRARY_LINK_STATIC_CXX_FLAGS})
|
||||
set(CMAKE_SHARED_LIBRARY_LINK_DYNAMIC_HIP_FLAGS ${CMAKE_SHARED_LIBRARY_LINK_DYNAMIC_CXX_FLAGS})
|
||||
|
||||
# Set the CMake Flags to use the HCC Compilier.
|
||||
set(CMAKE_HIP_CREATE_SHARED_LIBRARY "${HIP_HIPCC_CMAKE_LINKER_HELPER} ${HCC_PATH} <CMAKE_SHARED_LIBRARY_CXX_FLAGS> <LANGUAGE_COMPILE_FLAGS> <LINK_FLAGS> <CMAKE_SHARED_LIBRARY_CREATE_CXX_FLAGS> <SONAME_FLAG><TARGET_SONAME> -o <TARGET> <OBJECTS> <LINK_LIBRARIES>")
|
||||
set(CMAKE_HIP_CREATE_SHARED_MODULE "${HIP_HIPCC_CMAKE_LINKER_HELPER} ${HCC_PATH} <CMAKE_CXX_LINK_FLAGS> <LINK_FLAGS> <OBJECTS> <SONAME_FLAG><TARGET_SONAME> -o <TARGET> <LINK_LIBRARIES> -shared" )
|
||||
set(CMAKE_HIP_LINK_EXECUTABLE "${HIP_HIPCC_CMAKE_LINKER_HELPER} ${HCC_PATH} <FLAGS> <CMAKE_CXX_LINK_FLAGS> <LINK_FLAGS> <OBJECTS> -o <TARGET> <LINK_LIBRARIES>")
|
||||
|
||||
###############################################################################
|
||||
# FIND: HIP and associated helper binaries
|
||||
###############################################################################
|
||||
@@ -370,6 +391,11 @@ macro(HIP_PREPARE_TARGET_COMMANDS _target _format _generated_files _source_files
|
||||
# Initialize list of includes with those specified by the user. Append with
|
||||
# ones specified to cmake directly.
|
||||
set(HIP_HIPCC_INCLUDE_ARGS ${HIP_HIPCC_INCLUDE_ARGS_USER})
|
||||
|
||||
# Add the include directories
|
||||
set(include_directories_generator "$<TARGET_PROPERTY:${_target},INCLUDE_DIRECTORIES>")
|
||||
list(APPEND HIP_HIPCC_INCLUDE_ARGS "$<$<BOOL:${include_directories_generator}>:-I$<JOIN:${include_directories_generator}, -I>>")
|
||||
|
||||
get_directory_property(_hip_include_directories INCLUDE_DIRECTORIES)
|
||||
list(REMOVE_DUPLICATES _hip_include_directories)
|
||||
if(_hip_include_directories)
|
||||
@@ -383,6 +409,10 @@ macro(HIP_PREPARE_TARGET_COMMANDS _target _format _generated_files _source_files
|
||||
HIP_PARSE_HIPCC_OPTIONS(HIP_HCC_FLAGS ${_hcc_options})
|
||||
HIP_PARSE_HIPCC_OPTIONS(HIP_NVCC_FLAGS ${_nvcc_options})
|
||||
|
||||
# Add the compile definitions
|
||||
set(compile_definition_generator "$<TARGET_PROPERTY:${_target},COMPILE_DEFINITIONS>")
|
||||
list(APPEND HIP_HIPCC_FLAGS "$<$<BOOL:${compile_definition_generator}>:-D$<JOIN:${compile_definition_generator}, -D>>")
|
||||
|
||||
# Check if we are building shared library.
|
||||
set(_hip_build_shared_libs FALSE)
|
||||
list(FIND _hip_cmake_options SHARED _hip_found_SHARED)
|
||||
@@ -478,7 +508,7 @@ macro(HIP_PREPARE_TARGET_COMMANDS _target _format _generated_files _source_files
|
||||
set(verbose_output ON)
|
||||
else()
|
||||
set(verbose_output OFF)
|
||||
endif()
|
||||
endif()
|
||||
|
||||
# Create up the comment string
|
||||
file(RELATIVE_PATH generated_file_relative_path "${CMAKE_BINARY_DIR}" "${generated_file}")
|
||||
|
||||
@@ -0,0 +1,379 @@
|
||||
# CUDNN API supported by HIP
|
||||
|
||||
## **1. CUDNN Data types**
|
||||
|
||||
| **type** | **CUDA** | **HIP** |
|
||||
|-------------:|---------------------------------------------------------------|------------------------------------------------------------|
|
||||
| define |`CUDNN_VERSION` |`HIPDNN_VERSION` |
|
||||
| struct |`cudnnContext` | |
|
||||
| struct* |`cudnnHandle_t` |`hipdnnHandle_t` |
|
||||
| enum |***`cudnnStatus_t`*** |***`hipdnnStatus_t`*** |
|
||||
| 0 |*`CUDNN_STATUS_SUCCESS`* |*`HIPDNN_STATUS_SUCCESS`* |
|
||||
| 1 |*`CUDNN_STATUS_NOT_INITIALIZED`* |*`HIPDNN_STATUS_NOT_INITIALIZED`* |
|
||||
| 2 |*`CUDNN_STATUS_ALLOC_FAILED`* |*`HIPDNN_STATUS_ALLOC_FAILED`* |
|
||||
| 3 |*`CUDNN_STATUS_BAD_PARAM`* |*`HIPDNN_STATUS_BAD_PARAM`* |
|
||||
| 4 |*`CUDNN_STATUS_INTERNAL_ERROR`* |*`HIPDNN_STATUS_INTERNAL_ERROR`* |
|
||||
| 5 |*`CUDNN_STATUS_INVALID_VALUE`* |*`HIPDNN_STATUS_INVALID_VALUE`* |
|
||||
| 6 |*`CUDNN_STATUS_ARCH_MISMATCH`* |*`HIPDNN_STATUS_ARCH_MISMATCH`* |
|
||||
| 7 |*`CUDNN_STATUS_MAPPING_ERROR`* |*`HIPDNN_STATUS_MAPPING_ERROR`* |
|
||||
| 8 |*`CUDNN_STATUS_EXECUTION_FAILED`* |*`HIPDNN_STATUS_EXECUTION_FAILED`* |
|
||||
| 9 |*`CUDNN_STATUS_NOT_SUPPORTED`* |*`HIPDNN_STATUS_NOT_SUPPORTED`* |
|
||||
| 10 |*`CUDNN_STATUS_LICENSE_ERROR`* |*`HIPDNN_STATUS_LICENSE_ERROR`* |
|
||||
| 11 |*`CUDNN_STATUS_RUNTIME_PREREQUISITE_MISSING`* |*`HIPDNN_STATUS_RUNTIME_PREREQUISITE_MISSING`* |
|
||||
| 12 |*`CUDNN_STATUS_RUNTIME_IN_PROGRESS`* | |
|
||||
| 13 |*`CUDNN_STATUS_RUNTIME_FP_OVERFLOW`* | |
|
||||
| struct |`cudnnRuntimeTag_t` | |
|
||||
| enum |***`cudnnErrQueryMode_t`*** | |
|
||||
| 0 |*`CUDNN_ERRQUERY_RAWCODE`* | |
|
||||
| 1 |*`CUDNN_ERRQUERY_NONBLOCKING`* | |
|
||||
| 2 |*`CUDNN_ERRQUERY_BLOCKING`* | |
|
||||
| enum |***`libraryPropertyType_t`*** | |
|
||||
| struct |`cudnnTensorStruct` | |
|
||||
| struct* |`cudnnTensorDescriptor_t` |`hipdnnTensorDescriptor_t` |
|
||||
| struct |`cudnnConvolutionStruct` | |
|
||||
| struct* |`cudnnConvolutionDescriptor_t` |`hipdnnConvolutionDescriptor_t` |
|
||||
| struct |`cudnnPoolingStruct` | |
|
||||
| struct* |`cudnnPoolingDescriptor_t` |`hipdnnPoolingDescriptor_t` |
|
||||
| struct |`cudnnFilterStruct` | |
|
||||
| struct* |`cudnnFilterDescriptor_t` |`hipdnnFilterDescriptor_t` |
|
||||
| struct |`cudnnLRNStruct` | |
|
||||
| struct* |`cudnnLRNDescriptor_t` |`hipdnnLRNDescriptor_t` |
|
||||
| struct |`cudnnActivationStruct` | |
|
||||
| struct* |`cudnnActivationDescriptor_t` |`hipdnnActivationDescriptor_t` |
|
||||
| struct |`cudnnSpatialTransformerStruct` | |
|
||||
| struct* |`cudnnSpatialTransformerDescriptor_t` | |
|
||||
| struct |`cudnnOpTensorStruct` | |
|
||||
| struct* |`cudnnOpTensorDescriptor_t` |`hipdnnOpTensorDescriptor_t` |
|
||||
| struct |`cudnnReduceTensorStruct` | |
|
||||
| struct* |`cudnnReduceTensorDescriptor_t` |`hipdnnReduceTensorDescriptor_t` |
|
||||
| struct |`cudnnCTCLossStruct` | |
|
||||
| struct* |`cudnnCTCLossDescriptor_t` | |
|
||||
| enum |***`cudnnDataType_t`*** |***`hipdnnDataType_t`*** |
|
||||
| 0 |*`CUDNN_DATA_FLOAT`* |*`HIPDNN_DATA_FLOAT`* |
|
||||
| 1 |*`CUDNN_DATA_DOUBLE`* |*`HIPDNN_DATA_DOUBLE`* |
|
||||
| 2 |*`CUDNN_DATA_HALF`* |*`HIPDNN_DATA_HALF`* |
|
||||
| 3 |*`CUDNN_DATA_INT8`* |*`HIPDNN_DATA_INT8`* |
|
||||
| 4 |*`CUDNN_DATA_INT32`* |*`HIPDNN_DATA_INT32`* |
|
||||
| 5 |*`CUDNN_DATA_INT8x4`* |*`HIPDNN_DATA_INT8x4`* |
|
||||
| 6 |*`CUDNN_DATA_UINT8`* |*`HIPDNN_DATA_UINT8`* |
|
||||
| 7 |*`CUDNN_DATA_UINT8x4`* |*`HIPDNN_DATA_UINT8x4`* |
|
||||
| enum |***`cudnnMathType_t`*** |***`hipdnnMathType_t`*** |
|
||||
| 0 |*`CUDNN_DEFAULT_MATH`* |*`HIPDNN_DEFAULT_MATH`* |
|
||||
| 1 |*`CUDNN_TENSOR_OP_MATH`* |*`HIPDNN_TENSOR_OP_MATH`* |
|
||||
| enum |***`cudnnNanPropagation_t`*** |***`hipdnnNanPropagation_t`*** |
|
||||
| 0 |*`CUDNN_NOT_PROPAGATE_NAN`* |*`HIPDNN_NOT_PROPAGATE_NAN`* |
|
||||
| 1 |*`CUDNN_PROPAGATE_NAN`* |*`HIPDNN_PROPAGATE_NAN`* |
|
||||
| enum |***`cudnnDeterminism_t`*** | |
|
||||
| 0 |*`CUDNN_NON_DETERMINISTIC`* | |
|
||||
| 1 |*`CUDNN_DETERMINISTIC`* | |
|
||||
| define |`CUDNN_DIM_MAX` | |
|
||||
| enum |***`cudnnTensorFormat_t`*** |***`hipdnnTensorFormat_t`*** |
|
||||
| 0 |*`CUDNN_TENSOR_NCHW`* |*`HIPDNN_TENSOR_NCHW`* |
|
||||
| 1 |*`CUDNN_TENSOR_NHWC`* |*`HIPDNN_TENSOR_NHWC`* |
|
||||
| 2 |*`CUDNN_TENSOR_NCHW_VECT_C`* |*`HIPDNN_TENSOR_NCHW_VECT_C`* |
|
||||
| enum |***`cudnnOpTensorOp_t`*** |***`hipdnnOpTensorOp_t`*** |
|
||||
| 0 |*`CUDNN_OP_TENSOR_ADD`* |*`HIPDNN_OP_TENSOR_ADD`* |
|
||||
| 1 |*`CUDNN_OP_TENSOR_MUL`* |*`HIPDNN_OP_TENSOR_MUL`* |
|
||||
| 2 |*`CUDNN_OP_TENSOR_MIN`* |*`HIPDNN_OP_TENSOR_MIN`* |
|
||||
| 3 |*`CUDNN_OP_TENSOR_MAX`* |*`HIPDNN_OP_TENSOR_MAX`* |
|
||||
| 4 |*`CUDNN_OP_TENSOR_SQRT`* |*`HIPDNN_OP_TENSOR_SQRT`* |
|
||||
| 5 |*`CUDNN_OP_TENSOR_NOT`* | |
|
||||
| enum |***`cudnnReduceTensorOp_t`*** |***`hipdnnReduceTensorOp_t`*** |
|
||||
| 0 |*`CUDNN_REDUCE_TENSOR_ADD`* |*`HIPDNN_REDUCE_TENSOR_ADD`* |
|
||||
| 1 |*`CUDNN_REDUCE_TENSOR_MUL`* |*`HIPDNN_REDUCE_TENSOR_MUL`* |
|
||||
| 2 |*`CUDNN_REDUCE_TENSOR_MIN`* |*`HIPDNN_REDUCE_TENSOR_MIN`* |
|
||||
| 3 |*`CUDNN_REDUCE_TENSOR_MAX`* |*`HIPDNN_REDUCE_TENSOR_MAX`* |
|
||||
| 4 |*`CUDNN_REDUCE_TENSOR_AMAX`* |*`HIPDNN_REDUCE_TENSOR_AMAX`* |
|
||||
| 5 |*`CUDNN_REDUCE_TENSOR_AVG`* |*`HIPDNN_REDUCE_TENSOR_AVG`* |
|
||||
| 6 |*`CUDNN_REDUCE_TENSOR_NORM1`* |*`HIPDNN_REDUCE_TENSOR_NORM1`* |
|
||||
| 7 |*`CUDNN_REDUCE_TENSOR_NORM2`* |*`HIPDNN_REDUCE_TENSOR_NORM2`* |
|
||||
| 8 |*`CUDNN_REDUCE_TENSOR_MUL_NO_ZEROS`* |*`HIPDNN_REDUCE_TENSOR_MUL_NO_ZEROS`* |
|
||||
| enum |***`cudnnReduceTensorIndices_t`*** |***`hipdnnReduceTensorIndices_t`*** |
|
||||
| 0 |*`CUDNN_REDUCE_TENSOR_NO_INDICES`* |*`HIPDNN_REDUCE_TENSOR_NO_INDICES`* |
|
||||
| 1 |*`CUDNN_REDUCE_TENSOR_FLATTENED_INDICES`* |*`HIPDNN_REDUCE_TENSOR_FLATTENED_INDICES`* |
|
||||
| enum |***`cudnnIndicesType_t`*** |***`hipdnnIndicesType_t`*** |
|
||||
| 0 |*`CUDNN_32BIT_INDICES`* |*`HIPDNN_32BIT_INDICES`* |
|
||||
| 1 |*`CUDNN_64BIT_INDICES`* |*`HIPDNN_64BIT_INDICES`* |
|
||||
| 2 |*`CUDNN_16BIT_INDICES`* |*`HIPDNN_16BIT_INDICES`* |
|
||||
| 3 |*`CUDNN_8BIT_INDICES`* |*`HIPDNN_8BIT_INDICES`* |
|
||||
| enum |***`cudnnConvolutionMode_t`*** |***`hipdnnConvolutionMode_t`*** |
|
||||
| 0 |*`CUDNN_CONVOLUTION`* |*`HIPDNN_CONVOLUTION`* |
|
||||
| 1 |*`CUDNN_CROSS_CORRELATION`* |*`HIPDNN_CROSS_CORRELATION`* |
|
||||
| enum |***`cudnnConvolutionFwdPreference_t`*** |***`hipdnnConvolutionFwdPreference_t`*** |
|
||||
| 0 |*`CUDNN_CONVOLUTION_FWD_NO_WORKSPACE`* |*`HIPDNN_CONVOLUTION_FWD_NO_WORKSPACE`* |
|
||||
| 1 |*`CUDNN_CONVOLUTION_FWD_PREFER_FASTEST`* |*`HIPDNN_CONVOLUTION_FWD_PREFER_FASTEST`* |
|
||||
| 2 |*`CUDNN_CONVOLUTION_FWD_SPECIFY_WORKSPACE_LIMIT`* |*`HIPDNN_CONVOLUTION_FWD_SPECIFY_WORKSPACE_LIMIT`* |
|
||||
| enum |***`cudnnConvolutionFwdAlgo_t`*** |***`hipdnnConvolutionFwdAlgo_t`*** |
|
||||
| 0 |*`CUDNN_CONVOLUTION_FWD_ALGO_IMPLICIT_GEMM`* |*`HIPDNN_CONVOLUTION_FWD_ALGO_IMPLICIT_GEMM`* |
|
||||
| 1 |*`CUDNN_CONVOLUTION_FWD_ALGO_IMPLICIT_PRECOMP_GEMM`* |*`HIPDNN_CONVOLUTION_FWD_ALGO_IMPLICIT_PRECOMP_GEMM`* |
|
||||
| 2 |*`CUDNN_CONVOLUTION_FWD_ALGO_GEMM`* |*`HIPDNN_CONVOLUTION_FWD_ALGO_GEMM`* |
|
||||
| 3 |*`CUDNN_CONVOLUTION_FWD_ALGO_DIRECT`* |*`HIPDNN_CONVOLUTION_FWD_ALGO_DIRECT`* |
|
||||
| 4 |*`CUDNN_CONVOLUTION_FWD_ALGO_FFT`* |*`HIPDNN_CONVOLUTION_FWD_ALGO_FFT`* |
|
||||
| 5 |*`CUDNN_CONVOLUTION_FWD_ALGO_FFT_TILING`* |*`HIPDNN_CONVOLUTION_FWD_ALGO_FFT_TILING`* |
|
||||
| 6 |*`CUDNN_CONVOLUTION_FWD_ALGO_WINOGRAD`* |*`HIPDNN_CONVOLUTION_FWD_ALGO_WINOGRAD`* |
|
||||
| 7 |*`CUDNN_CONVOLUTION_FWD_ALGO_WINOGRAD_NONFUSED`* |*`HIPDNN_CONVOLUTION_FWD_ALGO_WINOGRAD_NONFUSED`* |
|
||||
| 8 |*`CUDNN_CONVOLUTION_FWD_ALGO_COUNT`* |*`HIPDNN_CONVOLUTION_FWD_ALGO_COUNT`* |
|
||||
| struct |`cudnnConvolutionFwdAlgoPerf_t` |`hipdnnConvolutionFwdAlgoPerf_t` |
|
||||
| enum |***`cudnnConvolutionBwdFilterPreference_t`*** |***`hipdnnConvolutionBwdFilterPreference_t`*** |
|
||||
| 0 |*`CUDNN_CONVOLUTION_BWD_FILTER_NO_WORKSPACE`* |*`HIPDNN_CONVOLUTION_BWD_FILTER_NO_WORKSPACE`* |
|
||||
| 1 |*`CUDNN_CONVOLUTION_BWD_FILTER_PREFER_FASTEST`* |*`HIPDNN_CONVOLUTION_BWD_FILTER_PREFER_FASTEST`* |
|
||||
| 2 |*`CUDNN_CONVOLUTION_BWD_FILTER_SPECIFY_WORKSPACE_LIMIT`* |*`HIPDNN_CONVOLUTION_BWD_FILTER_SPECIFY_WORKSPACE_LIMIT`* |
|
||||
| enum |***`cudnnConvolutionBwdFilterAlgo_t`*** |***`hipdnnConvolutionBwdFilterAlgo_t`*** |
|
||||
| 0 |*`CUDNN_CONVOLUTION_BWD_FILTER_ALGO_0`* |*`HIPDNN_CONVOLUTION_BWD_FILTER_ALGO_0`* |
|
||||
| 1 |*`CUDNN_CONVOLUTION_BWD_FILTER_ALGO_1`* |*`HIPDNN_CONVOLUTION_BWD_FILTER_ALGO_1`* |
|
||||
| 2 |*`CUDNN_CONVOLUTION_BWD_FILTER_ALGO_FFT`* |*`HIPDNN_CONVOLUTION_BWD_FILTER_ALGO_FFT`* |
|
||||
| 3 |*`CUDNN_CONVOLUTION_BWD_FILTER_ALGO_3`* |*`HIPDNN_CONVOLUTION_BWD_FILTER_ALGO_3`* |
|
||||
| 4 |*`CUDNN_CONVOLUTION_BWD_FILTER_ALGO_WINOGRAD`* |*`HIPDNN_CONVOLUTION_BWD_FILTER_ALGO_WINOGRAD`* |
|
||||
| 5 |*`CUDNN_CONVOLUTION_BWD_FILTER_ALGO_WINOGRAD_NONFUSED`* |*`HIPDNN_CONVOLUTION_BWD_FILTER_ALGO_WINOGRAD_NONFUSED`* |
|
||||
| 6 |*`CUDNN_CONVOLUTION_BWD_FILTER_ALGO_FFT_TILING`* |*`HIPDNN_CONVOLUTION_BWD_FILTER_ALGO_FFT_TILING`* |
|
||||
| 7 |*`CUDNN_CONVOLUTION_BWD_FILTER_ALGO_COUNT`* |*`HIPDNN_CONVOLUTION_BWD_FILTER_ALGO_COUNT`* |
|
||||
| struct |`cudnnConvolutionBwdDataAlgoPerf_t` |`hipdnnConvolutionBwdDataAlgoPerf_t` |
|
||||
| enum |***`cudnnSoftmaxAlgorithm_t`*** |***`hipdnnSoftmaxAlgorithm_t`*** |
|
||||
| 0 |*`CUDNN_SOFTMAX_FAST`* |*`HIPDNN_SOFTMAX_FAST`* |
|
||||
| 1 |*`CUDNN_SOFTMAX_ACCURATE`* |*`HIPDNN_SOFTMAX_ACCURATE`* |
|
||||
| 2 |*`CUDNN_SOFTMAX_LOG`* |*`HIPDNN_SOFTMAX_LOG`* |
|
||||
| enum |***`cudnnSoftmaxMode_t`*** |***`hipdnnSoftmaxMode_t`*** |
|
||||
| 0 |*`CUDNN_SOFTMAX_MODE_INSTANCE`* |*`HIPDNN_SOFTMAX_MODE_INSTANCE`* |
|
||||
| 1 |*`CUDNN_SOFTMAX_MODE_CHANNEL`* |*`HIPDNN_SOFTMAX_MODE_CHANNEL`* |
|
||||
| enum |***`cudnnPoolingMode_t`*** |***`hipdnnPoolingMode_t`*** |
|
||||
| 0 |*`CUDNN_POOLING_MAX`* |*`HIPDNN_POOLING_MAX`* |
|
||||
| 1 |*`CUDNN_POOLING_AVERAGE_COUNT_INCLUDE_PADDING`* |*`HIPDNN_POOLING_AVERAGE_COUNT_INCLUDE_PADDING`* |
|
||||
| 2 |*`CUDNN_POOLING_AVERAGE_COUNT_EXCLUDE_PADDING`* |*`HIPDNN_POOLING_AVERAGE_COUNT_EXCLUDE_PADDING`* |
|
||||
| 3 |*`CUDNN_POOLING_MAX_DETERMINISTIC`* |*`HIPDNN_POOLING_MAX_DETERMINISTIC`* |
|
||||
| enum |***`cudnnActivationMode_t`*** |***`hipdnnActivationMode_t`*** |
|
||||
| 0 |*`CUDNN_ACTIVATION_SIGMOID`* |*`HIPDNN_ACTIVATION_SIGMOID`* |
|
||||
| 1 |*`CUDNN_ACTIVATION_RELU`* |*`HIPDNN_ACTIVATION_RELU`* |
|
||||
| 2 |*`CUDNN_ACTIVATION_TANH`* |*`HIPDNN_ACTIVATION_TANH`* |
|
||||
| 3 |*`CUDNN_ACTIVATION_CLIPPED_RELU`* |*`HIPDNN_ACTIVATION_CLIPPED_RELU`* |
|
||||
| 4 |*`CUDNN_ACTIVATION_ELU`* |*`HIPDNN_ACTIVATION_ELU`* |
|
||||
| 5 |*`CUDNN_ACTIVATION_IDENTITY`* |*`HIPDNN_ACTIVATION_PATHTRU`* |
|
||||
| define |`CUDNN_LRN_MIN_N` | |
|
||||
| define |`CUDNN_LRN_MAX_N` | |
|
||||
| define |`CUDNN_LRN_MIN_K` | |
|
||||
| define |`CUDNN_LRN_MIN_BETA` | |
|
||||
| enum |***`cudnnLRNMode_t`*** |***`hipdnnLRNMode_t`*** |
|
||||
| 0 |*`CUDNN_LRN_CROSS_CHANNEL_DIM1`* |*`HIPDNN_LRN_CROSS_CHANNEL`* |
|
||||
| enum |***`cudnnDivNormMode_t`*** | |
|
||||
| 0 |*`CUDNN_DIVNORM_PRECOMPUTED_MEANS`* | |
|
||||
| enum |***`cudnnBatchNormMode_t`*** |***`hipdnnBatchNormMode_t`*** |
|
||||
| 0 |*`CUDNN_BATCHNORM_PER_ACTIVATION`* |*`HIPDNN_BATCHNORM_PER_ACTIVATION`* |
|
||||
| 1 |*`CUDNN_BATCHNORM_SPATIAL`* |*`HIPDNN_BATCHNORM_SPATIAL`* |
|
||||
| 2 |*`CUDNN_BATCHNORM_SPATIAL_PERSISTENT`* |*`HIPDNN_BATCHNORM_SPATIAL_PERSISTENT`* |
|
||||
| define |`CUDNN_BN_MIN_EPSILON` |`HIPDNN_BN_MIN_EPSILON` |
|
||||
| enum |***`cudnnSamplerType_t`*** | |
|
||||
| 0 |*`CUDNN_SAMPLER_BILINEAR`* | |
|
||||
| struct |`cudnnDropoutStruct` | |
|
||||
| struct* |`cudnnDropoutDescriptor_t` |`hipdnnDropoutDescriptor_t` |
|
||||
| enum |***`cudnnRNNMode_t`*** |***`hipdnnRNNMode_t`*** |
|
||||
| 0 |*`CUDNN_RNN_RELU`* |*`HIPDNN_RNN_RELU`* |
|
||||
| 1 |*`CUDNN_RNN_TANH`* |*`HIPDNN_RNN_TANH`* |
|
||||
| 2 |*`CUDNN_LSTM`* |*`HIPDNN_LSTM`* |
|
||||
| 3 |*`CUDNN_GRU`* |*`HIPDNN_GRU`* |
|
||||
| enum |***`cudnnDirectionMode_t`*** |***`hipdnnDirectionMode_t`*** |
|
||||
| 0 |*`CUDNN_UNIDIRECTIONAL`* |*`HIPDNN_UNIDIRECTIONAL`* |
|
||||
| 1 |*`CUDNN_BIDIRECTIONAL`* |*`HIPDNN_BIDIRECTIONAL`* |
|
||||
| enum |***`cudnnRNNAlgo_t`*** |***`hipdnnRNNAlgo_t`*** |
|
||||
| 0 |*`CUDNN_RNN_ALGO_STANDARD`* |*`HIPDNN_RNN_ALGO_STANDARD`* |
|
||||
| 1 |*`CUDNN_RNN_ALGO_PERSIST_STATIC`* |*`HIPDNN_RNN_ALGO_PERSIST_STATIC`* |
|
||||
| 2 |*`CUDNN_RNN_ALGO_PERSIST_DYNAMIC`* |*`HIPDNN_RNN_ALGO_PERSIST_DYNAMIC`* |
|
||||
| 3 |*`CUDNN_RNN_ALGO_COUNT`* | |
|
||||
| struct |`cudnnAlgorithmStruct` | |
|
||||
| struct* |`cudnnAlgorithmDescriptor_t` | |
|
||||
| struct |`cudnnAlgorithmPerformanceStruct` | |
|
||||
| struct* |`cudnnAlgorithmPerformance_t` | |
|
||||
| struct |`cudnnRNNStruct` | |
|
||||
| struct* |`cudnnRNNDescriptor_t` |`hipdnnRNNDescriptor_t` |
|
||||
| struct |`cudnnPersistentRNNPlan` | |
|
||||
| struct* |`cudnnPersistentRNNPlan_t` |`hipdnnPersistentRNNPlan_t` |
|
||||
| enum |***`cudnnCTCLossAlgo_t`*** | |
|
||||
| 0 |*`CUDNN_CTC_LOSS_ALGO_DETERMINISTIC`* | |
|
||||
| 1 |*`CUDNN_CTC_LOSS_ALGO_NON_DETERMINISTIC`* | |
|
||||
| struct |`cudnnAlgorithm_t` | |
|
||||
| enum |***`cudnnSeverity_t`*** | |
|
||||
| 0 |*`CUDNN_SEV_FATAL`* | |
|
||||
| 1 |*`CUDNN_SEV_ERROR`* | |
|
||||
| 2 |*`CUDNN_SEV_WARNING`* | |
|
||||
| 3 |*`CUDNN_SEV_INFO`* | |
|
||||
| define |`CUDNN_SEV_ERROR_EN` | |
|
||||
| define |`CUDNN_SEV_WARNING_EN` | |
|
||||
| define |`CUDNN_SEV_INFO_EN` | |
|
||||
| struct |`cudnnDebug_t` | |
|
||||
| struct |`cudnnCallback_t` | |
|
||||
|
||||
## **2. CUDNN API functions**
|
||||
|
||||
| **CUDA** | **HIP** |
|
||||
|-----------------------------------------------------------|-------------------------------------------------|
|
||||
|`cudnnGetVersion` |`hipdnnGetVersion` |
|
||||
|`cudnnGetCudartVersion` | |
|
||||
|`cudnnGetErrorString` |`hipdnnGetErrorString` |
|
||||
|`cudnnQueryRuntimeError` | |
|
||||
|`cudnnGetProperty` | |
|
||||
|`cudnnCreate` |`hipdnnCreate` |
|
||||
|`cudnnDestroy` |`hipdnnDestroy` |
|
||||
|`cudnnSetStream` |`hipdnnSetStream` |
|
||||
|`cudnnSetStream` |`hipdnnGetStream` |
|
||||
|`cudnnCreateTensorDescriptor` |`hipdnnCreateTensorDescriptor` |
|
||||
|`cudnnSetTensor4dDescriptor` |`hipdnnSetTensor4dDescriptor` |
|
||||
|`cudnnSetTensor4dDescriptorEx` | |
|
||||
|`cudnnGetTensor4dDescriptor` |`hipdnnGetTensor4dDescriptor` |
|
||||
|`cudnnSetTensorNdDescriptor` |`hipdnnSetTensorNdDescriptor` |
|
||||
|`cudnnSetTensorNdDescriptorEx` | |
|
||||
|`cudnnGetTensorNdDescriptor` |`hipdnnGetTensorNdDescriptor` |
|
||||
|`cudnnGetTensorSizeInBytes` | |
|
||||
|`cudnnDestroyTensorDescriptor` |`hipdnnDestroyTensorDescriptor` |
|
||||
|`cudnnTransformTensor` | |
|
||||
|`cudnnAddTensor` |`hipdnnAddTensor` |
|
||||
|`cudnnCreateOpTensorDescriptor` |`hipdnnCreateOpTensorDescriptor` |
|
||||
|`cudnnSetOpTensorDescriptor` |`hipdnnSetOpTensorDescriptor` |
|
||||
|`cudnnGetOpTensorDescriptor` |`hipdnnGetOpTensorDescriptor` |
|
||||
|`cudnnDestroyOpTensorDescriptor` |`hipdnnDestroyOpTensorDescriptor` |
|
||||
|`cudnnOpTensor` |`hipdnnOpTensor` |
|
||||
|`cudnnCreateReduceTensorDescriptor` |`hipdnnCreateReduceTensorDescriptor` |
|
||||
|`cudnnSetReduceTensorDescriptor` |`hipdnnSetReduceTensorDescriptor` |
|
||||
|`cudnnGetReduceTensorDescriptor` |`hipdnnGetReduceTensorDescriptor` |
|
||||
|`cudnnDestroyReduceTensorDescriptor` |`hipdnnDestroyReduceTensorDescriptor` |
|
||||
|`cudnnGetReductionIndicesSize` | |
|
||||
|`cudnnGetReductionWorkspaceSize` |`hipdnnGetReductionWorkspaceSize` |
|
||||
|`cudnnReduceTensor` |`hipdnnReduceTensor` |
|
||||
|`cudnnSetTensor` |`hipdnnSetTensor` |
|
||||
|`cudnnScaleTensor` |`hipdnnScaleTensor` |
|
||||
|`cudnnCreateFilterDescriptor` |`hipdnnCreateFilterDescriptor` |
|
||||
|`cudnnSetFilter4dDescriptor` | |
|
||||
|`cudnnGetFilter4dDescriptor` | |
|
||||
|`cudnnSetFilterNdDescriptor` |`hipdnnSetFilterNdDescriptor` |
|
||||
|`cudnnGetFilterNdDescriptor` |`hipdnnGetFilterNdDescriptor` |
|
||||
|`cudnnDestroyFilterDescriptor` |`hipdnnDestroyFilterDescriptor` |
|
||||
|`cudnnCreateConvolutionDescriptor` |`hipdnnCreateConvolutionDescriptor` |
|
||||
|`cudnnSetConvolutionMathType` |`hipdnnSetConvolutionMathType` |
|
||||
|`cudnnGetConvolutionMathType` | |
|
||||
|`cudnnSetConvolutionGroupCount` | |
|
||||
|`cudnnGetConvolutionGroupCount` | |
|
||||
|`cudnnSetConvolution2dDescriptor` |`hipdnnSetConvolution2dDescriptor` |
|
||||
|`cudnnGetConvolution2dDescriptor` |`hipdnnGetConvolution2dDescriptor` |
|
||||
|`cudnnGetConvolution2dForwardOutputDim` |`hipdnnGetConvolution2dForwardOutputDim` |
|
||||
|`cudnnSetConvolutionNdDescriptor` |`hipdnnSetConvolutionNdDescriptor` |
|
||||
|`cudnnGetConvolutionNdDescriptor` | |
|
||||
|`cudnnGetConvolutionNdForwardOutputDim` | |
|
||||
|`cudnnDestroyConvolutionDescriptor` | |
|
||||
|`cudnnGetConvolutionForwardAlgorithmMaxCount` | |
|
||||
|`cudnnFindConvolutionForwardAlgorithm` |`hipdnnFindConvolutionForwardAlgorithm` |
|
||||
|`cudnnFindConvolutionForwardAlgorithmEx` |`hipdnnFindConvolutionForwardAlgorithmEx` |
|
||||
|`cudnnGetConvolutionForwardAlgorithm` |`hipdnnGetConvolutionForwardAlgorithm` |
|
||||
|`cudnnGetConvolutionForwardAlgorithm_v7` | |
|
||||
|`cudnnGetConvolutionForwardWorkspaceSize` |`hipdnnGetConvolutionForwardWorkspaceSize` |
|
||||
|`cudnnConvolutionForward` |`hipdnnConvolutionForward` |
|
||||
|`cudnnConvolutionBiasActivationForward` | |
|
||||
|`cudnnConvolutionBackwardBias` |`hipdnnConvolutionBackwardBias` |
|
||||
|`cudnnGetConvolutionBackwardFilterAlgorithmMaxCount` | |
|
||||
|`cudnnFindConvolutionBackwardFilterAlgorithm` |`hipdnnFindConvolutionBackwardFilterAlgorithm` |
|
||||
|`cudnnFindConvolutionBackwardFilterAlgorithmEx` |`hipdnnFindConvolutionBackwardFilterAlgorithmEx` |
|
||||
|`cudnnGetConvolutionBackwardFilterAlgorithm` |`hipdnnGetConvolutionBackwardFilterAlgorithm` |
|
||||
|`cudnnGetConvolutionBackwardFilterAlgorithm_v7` | |
|
||||
|`cudnnGetConvolutionBackwardFilterWorkspaceSize` |`hipdnnGetConvolutionBackwardFilterWorkspaceSize`|
|
||||
|`cudnnConvolutionBackwardFilter` |`hipdnnConvolutionBackwardFilter` |
|
||||
|`cudnnGetConvolutionBackwardDataAlgorithmMaxCount` | |
|
||||
|`cudnnFindConvolutionBackwardDataAlgorithm` |`hipdnnFindConvolutionBackwardDataAlgorithm` |
|
||||
|`cudnnFindConvolutionBackwardDataAlgorithmEx` |`hipdnnFindConvolutionBackwardDataAlgorithmEx` |
|
||||
|`cudnnGetConvolutionBackwardDataAlgorithm` |`hipdnnGetConvolutionBackwardDataAlgorithm` |
|
||||
|`cudnnGetConvolutionBackwardDataAlgorithm_v7` | |
|
||||
|`cudnnGetConvolutionBackwardDataWorkspaceSize` |`hipdnnGetConvolutionBackwardDataWorkspaceSize` |
|
||||
|`cudnnConvolutionBackwardData` |`hipdnnConvolutionBackwardData` |
|
||||
|`cudnnIm2Col` | |
|
||||
|`cudnnSoftmaxForward` |`hipdnnSoftmaxForward` |
|
||||
|`cudnnSoftmaxBackward` |`hipdnnSoftmaxBackward` |
|
||||
|`cudnnCreatePoolingDescriptor` |`hipdnnCreatePoolingDescriptor` |
|
||||
|`cudnnSetPooling2dDescriptor` |`hipdnnSetPooling2dDescriptor` |
|
||||
|`cudnnGetPooling2dDescriptor` |`hipdnnGetPooling2dDescriptor` |
|
||||
|`cudnnSetPoolingNdDescriptor` |`hipdnnSetPoolingNdDescriptor` |
|
||||
|`cudnnGetPoolingNdDescriptor` | |
|
||||
|`cudnnGetPoolingNdForwardOutputDim` | |
|
||||
|`cudnnGetPooling2dForwardOutputDim` |`hipdnnGetPooling2dForwardOutputDim` |
|
||||
|`cudnnDestroyPoolingDescriptor` |`hipdnnDestroyPoolingDescriptor` |
|
||||
|`cudnnPoolingForward` |`hipdnnPoolingForward` |
|
||||
|`cudnnPoolingBackward` |`hipdnnPoolingBackward` |
|
||||
|`cudnnCreateActivationDescriptor` |`hipdnnCreateActivationDescriptor` |
|
||||
|`cudnnSetActivationDescriptor` |`hipdnnSetActivationDescriptor` |
|
||||
|`cudnnGetActivationDescriptor` |`hipdnnGetActivationDescriptor` |
|
||||
|`cudnnDestroyActivationDescriptor` |`hipdnnDestroyActivationDescriptor` |
|
||||
|`cudnnActivationForward` |`hipdnnActivationForward` |
|
||||
|`cudnnActivationBackward` |`hipdnnActivationBackward` |
|
||||
|`cudnnCreateLRNDescriptor` |`hipdnnCreateLRNDescriptor` |
|
||||
|`cudnnSetLRNDescriptor` |`hipdnnSetLRNDescriptor` |
|
||||
|`cudnnGetLRNDescriptor` |`hipdnnGetLRNDescriptor` |
|
||||
|`cudnnDestroyLRNDescriptor` |`hipdnnDestroyLRNDescriptor` |
|
||||
|`cudnnLRNCrossChannelForward` |`hipdnnLRNCrossChannelForward` |
|
||||
|`cudnnLRNCrossChannelBackward` |`hipdnnLRNCrossChannelBackward` |
|
||||
|`cudnnDivisiveNormalizationForward` | |
|
||||
|`cudnnDivisiveNormalizationBackward` | |
|
||||
|`cudnnDeriveBNTensorDescriptor` |`hipdnnDeriveBNTensorDescriptor` |
|
||||
|`cudnnBatchNormalizationForwardTraining` |`hipdnnBatchNormalizationForwardTraining` |
|
||||
|`cudnnBatchNormalizationForwardInference` |`hipdnnBatchNormalizationForwardInference` |
|
||||
|`cudnnBatchNormalizationBackward` |`hipdnnBatchNormalizationBackward` |
|
||||
|`cudnnCreateSpatialTransformerDescriptor` | |
|
||||
|`cudnnSetSpatialTransformerNdDescriptor` | |
|
||||
|`cudnnDestroySpatialTransformerDescriptor` | |
|
||||
|`cudnnSpatialTfGridGeneratorForward` | |
|
||||
|`cudnnSpatialTfGridGeneratorBackward` | |
|
||||
|`cudnnSpatialTfSamplerForward` | |
|
||||
|`cudnnSpatialTfSamplerBackward` | |
|
||||
|`cudnnCreateDropoutDescriptor` |`hipdnnCreateDropoutDescriptor` |
|
||||
|`cudnnDestroyDropoutDescriptor` |`hipdnnDestroyDropoutDescriptor` |
|
||||
|`cudnnDropoutGetStatesSize` |`hipdnnDropoutGetStatesSize` |
|
||||
|`cudnnDropoutGetReserveSpaceSize` | |
|
||||
|`cudnnSetDropoutDescriptor` |`hipdnnSetDropoutDescriptor` |
|
||||
|`cudnnGetDropoutDescriptor` | |
|
||||
|`cudnnRestoreDropoutDescriptor` | |
|
||||
|`cudnnDropoutForward` | |
|
||||
|`cudnnDropoutBackward` | |
|
||||
|`cudnnCreateRNNDescriptor` |`hipdnnCreateRNNDescriptor` |
|
||||
|`cudnnDestroyRNNDescriptor` |`hipdnnDestroyRNNDescriptor` |
|
||||
|`cudnnGetRNNForwardInferenceAlgorithmMaxCount` | |
|
||||
|`cudnnFindRNNForwardInferenceAlgorithmEx` | |
|
||||
|`cudnnGetRNNForwardTrainingAlgorithmMaxCount` | |
|
||||
|`cudnnFindRNNForwardTrainingAlgorithmEx` | |
|
||||
|`cudnnGetRNNBackwardDataAlgorithmMaxCount` | |
|
||||
|`cudnnFindRNNBackwardDataAlgorithmEx` | |
|
||||
|`cudnnGetRNNBackwardWeightsAlgorithmMaxCount` | |
|
||||
|`cudnnFindRNNBackwardWeightsAlgorithmEx` | |
|
||||
|`cudnnCreatePersistentRNNPlan` |`hipdnnCreatePersistentRNNPlan` |
|
||||
|`cudnnSetPersistentRNNPlan` |`hipdnnSetPersistentRNNPlan` |
|
||||
|`cudnnDestroyPersistentRNNPlan` |`hipdnnDestroyPersistentRNNPlan` |
|
||||
|`cudnnSetRNNDescriptor` |`hipdnnSetRNNDescriptor` |
|
||||
|`cudnnGetRNNDescriptor` | |
|
||||
|`cudnnSetRNNProjectionLayers` | |
|
||||
|`cudnnGetRNNProjectionLayers` | |
|
||||
|`cudnnSetRNNAlgorithmDescriptor` | |
|
||||
|`cudnnSetRNNMatrixMathType` | |
|
||||
|`cudnnGetRNNMatrixMathType` | |
|
||||
|`cudnnGetRNNWorkspaceSize` |`hipdnnGetRNNWorkspaceSize` |
|
||||
|`cudnnGetRNNTrainingReserveSize` |`hipdnnGetRNNTrainingReserveSize` |
|
||||
|`cudnnGetRNNParamsSize` |`hipdnnGetRNNParamsSize` |
|
||||
|`cudnnGetRNNLinLayerMatrixParams` |`hipdnnGetRNNLinLayerMatrixParams` |
|
||||
|`cudnnGetRNNLinLayerBiasParams` |`hipdnnGetRNNLinLayerBiasParams` |
|
||||
|`cudnnRNNForwardInference` |`hipdnnRNNForwardInference` |
|
||||
|`cudnnRNNForwardTraining` |`hipdnnRNNForwardTraining` |
|
||||
|`cudnnRNNBackwardData` |`hipdnnRNNBackwardData` |
|
||||
|`cudnnRNNBackwardWeights` |`hipdnnRNNBackwardWeights` |
|
||||
|`cudnnCreateCTCLossDescriptor` | |
|
||||
|`cudnnSetCTCLossDescriptor` | |
|
||||
|`cudnnGetCTCLossDescriptor` | |
|
||||
|`cudnnDestroyCTCLossDescriptor` | |
|
||||
|`cudnnCTCLoss` | |
|
||||
|`cudnnGetCTCLossWorkspaceSize` | |
|
||||
|`cudnnCreateAlgorithmDescriptor` | |
|
||||
|`cudnnSetAlgorithmDescriptor` | |
|
||||
|`cudnnGetAlgorithmDescriptor` | |
|
||||
|`cudnnCopyAlgorithmDescriptor` | |
|
||||
|`cudnnDestroyAlgorithmDescriptor` | |
|
||||
|`cudnnCreateAlgorithmPerformance` | |
|
||||
|`cudnnSetAlgorithmPerformance` | |
|
||||
|`cudnnGetAlgorithmPerformance` | |
|
||||
|`cudnnDestroyAlgorithmPerformance` | |
|
||||
|`cudnnGetAlgorithmSpaceSize` | |
|
||||
|`cudnnSaveAlgorithm` | |
|
||||
|`cudnnRestoreAlgorithm` | |
|
||||
|`cudnnSetRNNDescriptor_v5` |`hipdnnSetRNNDescriptor_v5` |
|
||||
|`cudnnSetRNNDescriptor_v6` |`hipdnnSetRNNDescriptor_v6` |
|
||||
|`cudnnSetCallback` | |
|
||||
|`cudnnGetCallback` | |
|
||||
@@ -1,245 +1,253 @@
|
||||
# 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:
|
||||
- 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 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.
|
||||
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 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 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.
|
||||
|
||||
### 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 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.
|
||||
|
||||
Most new applications will prefer to use `hipSetDevice` or the stream APIs , therefore HIP has marked hipCtx APIs as **deprecated**. Support for these APIs may not be available in future releases. For more details on deprecated APIs please refer [HIP deprecated APIs](https://github.com/ROCm-Developer-Tools/HIP/tree/master/docs/markdown/hip_deprecated_api_list.md).
|
||||
|
||||
## HIP Module and Ctx APIs
|
||||
|
||||
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:
|
||||
|
||||
| Format | APIs | NVCC | HCC |
|
||||
| --- | --- | --- | --- |
|
||||
| Code Object | hipModuleLoad, hipModuleLoadData | .cubin or PTX text | .hsaco |
|
||||
| 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.
|
||||
|
||||
|
||||
### 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`.
|
||||
|
||||
### hipify translation of CUDA Driver API
|
||||
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 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.
|
||||
|
||||
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.
|
||||
|
||||
#### 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.
|
||||
|
||||
#### 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.
|
||||
|
||||
### 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_t | CUmodule | |
|
||||
| hipFunction_t | CUfunction | |
|
||||
| hipCtx_t | CUcontext | |
|
||||
| hipDevice_t | CUdevice | |
|
||||
| hipStream_t | CUstream | cudaStream_t |
|
||||
| hipEvent_t | CUevent | cudaEvent_t |
|
||||
| hipArray | CUarray | cudaArray |
|
||||
|
||||
#### Compilation Options
|
||||
The `hipModule_t` interface does not support `cuModuleLoadDataEx` function, which is used to control PTX compilation options.
|
||||
HCC does not use PTX and does not support these compilation options.
|
||||
In fact, HCC code objects always contain fully compiled ISA and do not require additional compilation as a part of the load step.
|
||||
The corresponding HIP function `hipModuleLoadDataEx` behaves as `hipModuleLoadData` on HCC path (compilation options are not used) and as `cuModuleLoadDataEx` on NVCC path.
|
||||
For example (CUDA):
|
||||
```
|
||||
CUmodule module;
|
||||
void *imagePtr = ...; // Somehow populate data pointer with code object
|
||||
|
||||
const int numOptions = 1;
|
||||
CUJit_option options[numOptions];
|
||||
void * optionValues[numOptions];
|
||||
|
||||
options[0] = CU_JIT_MAX_REGISTERS;
|
||||
unsigned maxRegs = 15;
|
||||
optionValues[0] = (void*)(&maxRegs);
|
||||
|
||||
cuModuleLoadDataEx(module, imagePtr, numOptions, options, optionValues);
|
||||
|
||||
CUfunction k;
|
||||
cuModuleGetFunction(&k, module, "myKernel");
|
||||
```
|
||||
HIP:
|
||||
```
|
||||
hipModule_t module;
|
||||
void *imagePtr = ...; // Somehow populate data pointer with code object
|
||||
|
||||
const int numOptions = 1;
|
||||
hipJitOption options[numOptions];
|
||||
void * optionValues[numOptions];
|
||||
|
||||
options[0] = hipJitOptionMaxRegisters;
|
||||
unsigned maxRegs = 15;
|
||||
optionValues[0] = (void*)(&maxRegs);
|
||||
|
||||
// hipModuleLoadData(module, imagePtr) will be called on HCC path, JIT options will not be used, and
|
||||
// cupModuleLoadDataEx(module, imagePtr, numOptions, options, optionValues) will be called on NVCC path
|
||||
hipModuleLoadDataEx(module, imagePtr, numOptions, options, optionValues);
|
||||
|
||||
hipFunction_t k;
|
||||
hipModuleGetFunction(&k, module, "myKernel");
|
||||
```
|
||||
|
||||
The below sample shows how to use `hipModuleGetFunction`.
|
||||
|
||||
```
|
||||
#include<hip_runtime.h>
|
||||
#include<hip_runtime_api.h>
|
||||
#include<iostream>
|
||||
#include<fstream>
|
||||
#include<vector>
|
||||
|
||||
#define LEN 64
|
||||
#define SIZE LEN<<2
|
||||
|
||||
#ifdef __HIP_PLATFORM_HCC__
|
||||
#define fileName "vcpy_isa.co"
|
||||
#endif
|
||||
|
||||
#ifdef __HIP_PLATFORM_NVCC__
|
||||
#define fileName "vcpy_isa.ptx"
|
||||
#endif
|
||||
|
||||
#define kernel_name "hello_world"
|
||||
|
||||
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;
|
||||
std::cout<<A[i] << " "<<B[i]<<std::endl;
|
||||
}
|
||||
|
||||
|
||||
#ifdef __HIP_PLATFORM_NVCC__
|
||||
hipInit(0);
|
||||
hipDevice_t device;
|
||||
hipCtx_t context;
|
||||
hipDeviceGet(&device, 0);
|
||||
hipCtxCreate(&context, 0, device);
|
||||
#endif
|
||||
|
||||
hipMalloc((void**)&Ad, SIZE);
|
||||
hipMalloc((void**)&Bd, SIZE);
|
||||
|
||||
hipMemcpyHtoD(Ad, A, SIZE);
|
||||
hipMemcpyHtoD(Bd, B, SIZE);
|
||||
hipModule_t Module;
|
||||
hipFunction_t Function;
|
||||
hipModuleLoad(&Module, fileName);
|
||||
hipModuleGetFunction(&Function, Module, kernel_name);
|
||||
|
||||
std::vector<void*>argBuffer(2);
|
||||
memcpy(&argBuffer[0], &Ad, sizeof(void*));
|
||||
memcpy(&argBuffer[1], &Bd, sizeof(void*));
|
||||
|
||||
size_t size = argBuffer.size()*sizeof(void*);
|
||||
|
||||
void *config[] = {
|
||||
HIP_LAUNCH_PARAM_BUFFER_POINTER, &argBuffer[0],
|
||||
HIP_LAUNCH_PARAM_BUFFER_SIZE, &size,
|
||||
HIP_LAUNCH_PARAM_END
|
||||
};
|
||||
|
||||
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++){
|
||||
std::cout<<A[i]<<" - "<<B[i]<<std::endl;
|
||||
}
|
||||
|
||||
#ifdef __HIP_PLATFORM_NVCC__
|
||||
hipCtxDetach(context);
|
||||
#endif
|
||||
|
||||
return 0;
|
||||
}
|
||||
```
|
||||
|
||||
## HIP Module and Texture Driver API
|
||||
|
||||
HIP supports texture driver APIs however texture reference should be declared in host scope. Following code explains the use of texture reference for __HIP_PLATFORM_HCC__ platform.
|
||||
|
||||
```
|
||||
// Code to generate code object
|
||||
|
||||
# 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:
|
||||
- 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 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.
|
||||
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 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 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.
|
||||
|
||||
### 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 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.
|
||||
|
||||
Most new applications will prefer to use `hipSetDevice` or the stream APIs , therefore HIP has marked hipCtx APIs as **deprecated**. Support for these APIs may not be available in future releases. For more details on deprecated APIs please refer [HIP deprecated APIs](https://github.com/ROCm-Developer-Tools/HIP/tree/master/docs/markdown/hip_deprecated_api_list.md).
|
||||
|
||||
## HIP Module and Ctx APIs
|
||||
|
||||
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:
|
||||
|
||||
| Format | APIs | NVCC | HCC | HIP-CLANG |
|
||||
| --- | --- | --- | --- | ---
|
||||
| Code Object | hipModuleLoad, hipModuleLoadData | .cubin or PTX text | .hsaco | .hsaco |
|
||||
| Fat Binary | hipModuleLoadFatBin | .fatbin | Under Development | .hip_fatbin |
|
||||
|
||||
`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`.
|
||||
|
||||
### hipify translation of CUDA Driver API
|
||||
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 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.
|
||||
|
||||
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.
|
||||
|
||||
#### 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.
|
||||
|
||||
#### 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.
|
||||
|
||||
### hip-clang Implementation Notes
|
||||
#### .hip_fatbin
|
||||
hip-clang links device code from different translation units together. For each device target, a code object is generated. Code objects for different device targets are bundled by clang-offload-bundler as one fatbinary, which is embeded as a global symbol `__hip_fatbin` in the .hip_fatbin section of the ELF file of the executable or shared object.
|
||||
|
||||
#### Initialization and Termination Functions
|
||||
hip-clang generates initializatiion and termination functions for each translation unit for host code compilation. The initialization functions call `__hipRegisterFatBinary` to register the fatbinary embeded in the ELF file. They also call `__hipRegisterFunction` and `__hipRegisterVar` to register kernel functions and device side global variables. The termination functions call `__hipUnregisterFatBinary`.
|
||||
hip-clang emits a global variable `__hip_gpubin_handle` of void** type with linkonce linkage and inital value 0 for each host translation unit. Each initialization function checks `__hip_gpubin_handle` and register the fatbinary only if `__hip_gpubin_handle` is 0 and saves the return value of `__hip_gpubin_handle` to `__hip_gpubin_handle`. This is to guarantee that the fatbinary is only registered once. Similar check is done in the termination functions.
|
||||
|
||||
### 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_t | CUmodule | |
|
||||
| hipFunction_t | CUfunction | |
|
||||
| hipCtx_t | CUcontext | |
|
||||
| hipDevice_t | CUdevice | |
|
||||
| hipStream_t | CUstream | cudaStream_t |
|
||||
| hipEvent_t | CUevent | cudaEvent_t |
|
||||
| hipArray | CUarray | cudaArray |
|
||||
|
||||
#### Compilation Options
|
||||
The `hipModule_t` interface does not support `cuModuleLoadDataEx` function, which is used to control PTX compilation options.
|
||||
HCC does not use PTX and does not support these compilation options.
|
||||
In fact, HCC code objects always contain fully compiled ISA and do not require additional compilation as a part of the load step.
|
||||
The corresponding HIP function `hipModuleLoadDataEx` behaves as `hipModuleLoadData` on HCC path (compilation options are not used) and as `cuModuleLoadDataEx` on NVCC path.
|
||||
For example (CUDA):
|
||||
```
|
||||
CUmodule module;
|
||||
void *imagePtr = ...; // Somehow populate data pointer with code object
|
||||
|
||||
const int numOptions = 1;
|
||||
CUJit_option options[numOptions];
|
||||
void * optionValues[numOptions];
|
||||
|
||||
options[0] = CU_JIT_MAX_REGISTERS;
|
||||
unsigned maxRegs = 15;
|
||||
optionValues[0] = (void*)(&maxRegs);
|
||||
|
||||
cuModuleLoadDataEx(module, imagePtr, numOptions, options, optionValues);
|
||||
|
||||
CUfunction k;
|
||||
cuModuleGetFunction(&k, module, "myKernel");
|
||||
```
|
||||
HIP:
|
||||
```
|
||||
hipModule_t module;
|
||||
void *imagePtr = ...; // Somehow populate data pointer with code object
|
||||
|
||||
const int numOptions = 1;
|
||||
hipJitOption options[numOptions];
|
||||
void * optionValues[numOptions];
|
||||
|
||||
options[0] = hipJitOptionMaxRegisters;
|
||||
unsigned maxRegs = 15;
|
||||
optionValues[0] = (void*)(&maxRegs);
|
||||
|
||||
// hipModuleLoadData(module, imagePtr) will be called on HCC path, JIT options will not be used, and
|
||||
// cupModuleLoadDataEx(module, imagePtr, numOptions, options, optionValues) will be called on NVCC path
|
||||
hipModuleLoadDataEx(module, imagePtr, numOptions, options, optionValues);
|
||||
|
||||
hipFunction_t k;
|
||||
hipModuleGetFunction(&k, module, "myKernel");
|
||||
```
|
||||
|
||||
The below sample shows how to use `hipModuleGetFunction`.
|
||||
|
||||
```
|
||||
#include<hip_runtime.h>
|
||||
#include<hip_runtime_api.h>
|
||||
#include<iostream>
|
||||
#include<fstream>
|
||||
#include<vector>
|
||||
|
||||
#define LEN 64
|
||||
#define SIZE LEN<<2
|
||||
|
||||
#ifdef __HIP_PLATFORM_HCC__
|
||||
#define fileName "vcpy_isa.co"
|
||||
#endif
|
||||
|
||||
#ifdef __HIP_PLATFORM_NVCC__
|
||||
#define fileName "vcpy_isa.ptx"
|
||||
#endif
|
||||
|
||||
#define kernel_name "hello_world"
|
||||
|
||||
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;
|
||||
std::cout<<A[i] << " "<<B[i]<<std::endl;
|
||||
}
|
||||
|
||||
|
||||
#ifdef __HIP_PLATFORM_NVCC__
|
||||
hipInit(0);
|
||||
hipDevice_t device;
|
||||
hipCtx_t context;
|
||||
hipDeviceGet(&device, 0);
|
||||
hipCtxCreate(&context, 0, device);
|
||||
#endif
|
||||
|
||||
hipMalloc((void**)&Ad, SIZE);
|
||||
hipMalloc((void**)&Bd, SIZE);
|
||||
|
||||
hipMemcpyHtoD(Ad, A, SIZE);
|
||||
hipMemcpyHtoD(Bd, B, SIZE);
|
||||
hipModule_t Module;
|
||||
hipFunction_t Function;
|
||||
hipModuleLoad(&Module, fileName);
|
||||
hipModuleGetFunction(&Function, Module, kernel_name);
|
||||
|
||||
std::vector<void*>argBuffer(2);
|
||||
memcpy(&argBuffer[0], &Ad, sizeof(void*));
|
||||
memcpy(&argBuffer[1], &Bd, sizeof(void*));
|
||||
|
||||
size_t size = argBuffer.size()*sizeof(void*);
|
||||
|
||||
void *config[] = {
|
||||
HIP_LAUNCH_PARAM_BUFFER_POINTER, &argBuffer[0],
|
||||
HIP_LAUNCH_PARAM_BUFFER_SIZE, &size,
|
||||
HIP_LAUNCH_PARAM_END
|
||||
};
|
||||
|
||||
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++){
|
||||
std::cout<<A[i]<<" - "<<B[i]<<std::endl;
|
||||
}
|
||||
|
||||
#ifdef __HIP_PLATFORM_NVCC__
|
||||
hipCtxDetach(context);
|
||||
#endif
|
||||
|
||||
return 0;
|
||||
}
|
||||
```
|
||||
|
||||
## HIP Module and Texture Driver API
|
||||
|
||||
HIP supports texture driver APIs however texture reference should be declared in host scope. Following code explains the use of texture reference for __HIP_PLATFORM_HCC__ platform.
|
||||
|
||||
```
|
||||
// Code to generate code object
|
||||
|
||||
#include "hip/hip_runtime.h"
|
||||
extern texture<float, 2, hipReadModeElementType> tex;
|
||||
|
||||
@@ -251,26 +259,26 @@ __global__ void tex2dKernel(hipLaunchParm lp, float* outputData,
|
||||
int y = hipBlockIdx_y*hipBlockDim_y + hipThreadIdx_y;
|
||||
outputData[y*width + x] = tex2D(tex, x, y);
|
||||
}
|
||||
|
||||
```
|
||||
```
|
||||
// Host code:
|
||||
|
||||
texture<float, 2, hipReadModeElementType> tex;
|
||||
|
||||
void myFunc ()
|
||||
{
|
||||
// ...
|
||||
|
||||
textureReference* texref;
|
||||
hipModuleGetTexRef(&texref, Module1, "tex");
|
||||
hipTexRefSetAddressMode(texref, 0, hipAddressModeWrap);
|
||||
hipTexRefSetAddressMode(texref, 1, hipAddressModeWrap);
|
||||
hipTexRefSetFilterMode(texref, hipFilterModePoint);
|
||||
hipTexRefSetFlags(texref, 0);
|
||||
hipTexRefSetFormat(texref, HIP_AD_FORMAT_FLOAT, 1);
|
||||
hipTexRefSetArray(texref, array, HIP_TRSA_OVERRIDE_FORMAT);
|
||||
|
||||
// ...
|
||||
}
|
||||
```
|
||||
|
||||
```
|
||||
```
|
||||
// Host code:
|
||||
|
||||
texture<float, 2, hipReadModeElementType> tex;
|
||||
|
||||
void myFunc ()
|
||||
{
|
||||
// ...
|
||||
|
||||
textureReference* texref;
|
||||
hipModuleGetTexRef(&texref, Module1, "tex");
|
||||
hipTexRefSetAddressMode(texref, 0, hipAddressModeWrap);
|
||||
hipTexRefSetAddressMode(texref, 1, hipAddressModeWrap);
|
||||
hipTexRefSetFilterMode(texref, hipFilterModePoint);
|
||||
hipTexRefSetFlags(texref, 0);
|
||||
hipTexRefSetFormat(texref, HIP_AD_FORMAT_FLOAT, 1);
|
||||
hipTexRefSetArray(texref, array, HIP_TRSA_OVERRIDE_FORMAT);
|
||||
|
||||
// ...
|
||||
}
|
||||
```
|
||||
|
||||
@@ -2920,64 +2920,60 @@ const std::map<llvm::StringRef, hipCounter> CUDA_IDENTIFIER_MAP{
|
||||
// unchanged function names: skipahead, skipahead_sequence, skipahead_subsequence
|
||||
|
||||
///////////////////////////// cuDNN /////////////////////////////
|
||||
// defines
|
||||
{"CUDNN_VERSION", {"HIPDNN_VERSION", CONV_NUMERIC_LITERAL, API_DNN}}, // 7000
|
||||
{"CUDNN_DIM_MAX", {"HIPDNN_DIM_MAX", CONV_NUMERIC_LITERAL, API_DNN, HIP_UNSUPPORTED}}, // 8
|
||||
{"CUDNN_LRN_MIN_N", {"HIPDNN_LRN_MIN_N", CONV_NUMERIC_LITERAL, API_DNN, HIP_UNSUPPORTED}}, // 1
|
||||
{"CUDNN_LRN_MAX_N", {"HIPDNN_LRN_MAX_N", CONV_NUMERIC_LITERAL, API_DNN, HIP_UNSUPPORTED}}, // 16
|
||||
{"CUDNN_LRN_MIN_K", {"HIPDNN_LRN_MIN_K", CONV_NUMERIC_LITERAL, API_DNN, HIP_UNSUPPORTED}}, // 1e-5
|
||||
{"CUDNN_LRN_MIN_BETA", {"HIPDNN_LRN_MIN_BETA", CONV_NUMERIC_LITERAL, API_DNN, HIP_UNSUPPORTED}}, // 0.01
|
||||
{"CUDNN_BN_MIN_EPSILON", {"HIPDNN_BN_MIN_EPSILON", CONV_NUMERIC_LITERAL, API_DNN}}, // 1e-5
|
||||
{"CUDNN_SEV_ERROR_EN", {"HIPDNN_SEV_ERROR_EN", CONV_NUMERIC_LITERAL, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"CUDNN_SEV_WARNING_EN", {"HIPDNN_SEV_WARNING_EN", CONV_NUMERIC_LITERAL, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"CUDNN_SEV_INFO_EN", {"HIPDNN_SEV_INFO_EN", CONV_NUMERIC_LITERAL, API_DNN, HIP_UNSUPPORTED}},
|
||||
|
||||
{"cudnnContext", {"hipdnnContext", CONV_TYPE, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnHandle_t", {"hipdnnHandle_t", CONV_TYPE, API_DNN}},
|
||||
{"cudnnStatus_t", {"hipdnnStatus_t", CONV_TYPE, API_DNN}},
|
||||
{"CUDNN_STATUS_SUCCESS", {"HIPDNN_STATUS_SUCCESS", CONV_NUMERIC_LITERAL, API_DNN}}, // 0
|
||||
{"CUDNN_STATUS_NOT_INITIALIZED", {"HIPDNN_STATUS_NOT_INITIALIZED", CONV_NUMERIC_LITERAL, API_DNN}}, // 1
|
||||
{"CUDNN_STATUS_ALLOC_FAILED", {"HIPDNN_STATUS_ALLOC_FAILED", CONV_NUMERIC_LITERAL, API_DNN}}, // 2
|
||||
{"CUDNN_STATUS_BAD_PARAM", {"HIPDNN_STATUS_BAD_PARAM", CONV_NUMERIC_LITERAL, API_DNN}}, // 3
|
||||
{"CUDNN_STATUS_INTERNAL_ERROR", {"HIPDNN_STATUS_INTERNAL_ERROR", CONV_NUMERIC_LITERAL, API_DNN}}, // 4
|
||||
{"CUDNN_STATUS_INVALID_VALUE", {"HIPDNN_STATUS_INVALID_VALUE", CONV_NUMERIC_LITERAL, API_DNN}}, // 5
|
||||
{"CUDNN_STATUS_ARCH_MISMATCH", {"HIPDNN_STATUS_ARCH_MISMATCH", CONV_NUMERIC_LITERAL, API_DNN}}, // 6
|
||||
{"CUDNN_STATUS_MAPPING_ERROR", {"HIPDNN_STATUS_MAPPING_ERROR", CONV_NUMERIC_LITERAL, API_DNN}}, // 7
|
||||
{"CUDNN_STATUS_EXECUTION_FAILED", {"HIPDNN_STATUS_EXECUTION_FAILED", CONV_NUMERIC_LITERAL, API_DNN}}, // 8
|
||||
{"CUDNN_STATUS_NOT_SUPPORTED", {"HIPDNN_STATUS_NOT_SUPPORTED", CONV_NUMERIC_LITERAL, API_DNN}}, // 9
|
||||
{"CUDNN_STATUS_LICENSE_ERROR", {"HIPDNN_STATUS_LICENSE_ERROR", CONV_NUMERIC_LITERAL, API_DNN}}, // 10
|
||||
{"CUDNN_STATUS_RUNTIME_PREREQUISITE_MISSING", {"HIPDNN_STATUS_RUNTIME_PREREQUISITE_MISSING", CONV_NUMERIC_LITERAL, API_DNN}}, // 11
|
||||
{"CUDNN_STATUS_RUNTIME_IN_PROGRESS", {"HIPDNN_STATUS_RUNTIME_IN_PROGRESS", CONV_NUMERIC_LITERAL, API_DNN, HIP_UNSUPPORTED}}, // 12
|
||||
{"CUDNN_STATUS_RUNTIME_FP_OVERFLOW", {"HIPDNN_STATUS_RUNTIME_FP_OVERFLOW", CONV_NUMERIC_LITERAL, API_DNN, HIP_UNSUPPORTED}}, // 13
|
||||
{"cudnnRuntimeTag_t", {"hipdnnRuntimeTag_t", CONV_TYPE, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnTensorDescriptor_t", {"hipdnnTensorDescriptor_t", CONV_TYPE, API_DNN}},
|
||||
{"cudnnConvolutionDescriptor_t", {"hipdnnConvolutionDescriptor_t", CONV_TYPE, API_DNN}},
|
||||
{"cudnnConvolutionMode_t", {"hipdnnConvolutionMode_t", CONV_TYPE, API_DNN}},
|
||||
{"CUDNN_CONVOLUTION", {"HIPDNN_CONVOLUTION", CONV_NUMERIC_LITERAL, API_DNN}}, // 0
|
||||
{"CUDNN_CROSS_CORRELATION", {"HIPDNN_CROSS_CORRELATION", CONV_NUMERIC_LITERAL, API_DNN}}, // 1
|
||||
{"cudnnTensorFormat_t", {"hipdnnTensorFormat_t", CONV_TYPE, API_DNN}},
|
||||
{"CUDNN_TENSOR_NCHW", {"HIPDNN_TENSOR_NCHW", CONV_NUMERIC_LITERAL, API_DNN}}, // 0
|
||||
{"CUDNN_TENSOR_NHWC", {"HIPDNN_TENSOR_NHWC", CONV_NUMERIC_LITERAL, API_DNN}}, // 1
|
||||
{"CUDNN_TENSOR_NCHW_VECT_C", {"HIPDNN_TENSOR_NCHW_VECT_C", CONV_NUMERIC_LITERAL, API_DNN}}, // 2
|
||||
{"cudnnDataType_t", {"hipdnnDataType_t", CONV_TYPE, API_DNN}},
|
||||
{"CUDNN_DATA_FLOAT", {"HIPDNN_DATA_FLOAT", CONV_NUMERIC_LITERAL, API_DNN}}, // 0
|
||||
{"CUDNN_DATA_DOUBLE", {"HIPDNN_DATA_DOUBLE", CONV_NUMERIC_LITERAL, API_DNN}}, // 1
|
||||
{"CUDNN_DATA_HALF", {"HIPDNN_DATA_HALF", CONV_NUMERIC_LITERAL, API_DNN}}, // 2
|
||||
{"CUDNN_DATA_INT8", {"HIPDNN_DATA_INT8", CONV_NUMERIC_LITERAL, API_DNN}}, // 3
|
||||
{"CUDNN_DATA_INT32", {"HIPDNN_DATA_INT32", CONV_NUMERIC_LITERAL, API_DNN}}, // 4
|
||||
{"CUDNN_DATA_INT8x4", {"HIPDNN_DATA_INT8x4", CONV_NUMERIC_LITERAL, API_DNN}}, // 5
|
||||
{"CUDNN_DATA_UINT8", {"HIPDNN_DATA_UINT8", CONV_NUMERIC_LITERAL, API_DNN, HIP_UNSUPPORTED}}, // 6
|
||||
{"CUDNN_DATA_UINT8x4", {"HIPDNN_DATA_UINT8x4", CONV_NUMERIC_LITERAL, API_DNN, HIP_UNSUPPORTED}}, // 7
|
||||
{"cudnnErrQueryMode_t", {"hipdnnErrQueryMode_t", CONV_TYPE, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"CUDNN_ERRQUERY_RAWCODE", {"HIPDNN_ERRQUERY_RAWCODE", CONV_NUMERIC_LITERAL, API_DNN, HIP_UNSUPPORTED}}, // 0
|
||||
{"CUDNN_ERRQUERY_NONBLOCKING", {"HIPDNN_ERRQUERY_NONBLOCKING", CONV_NUMERIC_LITERAL, API_DNN, HIP_UNSUPPORTED}}, // 1
|
||||
{"CUDNN_ERRQUERY_BLOCKING", {"HIPDNN_ERRQUERY_BLOCKING", CONV_NUMERIC_LITERAL, API_DNN, HIP_UNSUPPORTED}}, // 2
|
||||
{"cudnnSeverity_t", {"hipdnnSeverity_t", CONV_TYPE, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"CUDNN_SEV_FATAL", {"HIPDNN_SEV_FATAL", CONV_NUMERIC_LITERAL, API_DNN, HIP_UNSUPPORTED}}, // 0
|
||||
{"CUDNN_SEV_ERROR", {"HIPDNN_SEV_ERROR", CONV_NUMERIC_LITERAL, API_DNN, HIP_UNSUPPORTED}}, // 1
|
||||
{"CUDNN_SEV_WARNING", {"HIPDNN_SEV_WARNING", CONV_NUMERIC_LITERAL, API_DNN, HIP_UNSUPPORTED}}, // 2
|
||||
{"CUDNN_SEV_INFO", {"HIPDNN_SEV_INFO", CONV_NUMERIC_LITERAL, API_DNN, HIP_UNSUPPORTED}}, // 3
|
||||
// cuDNN defines
|
||||
{"CUDNN_VERSION", {"HIPDNN_VERSION", CONV_NUMERIC_LITERAL, API_DNN}}, // 7000
|
||||
{"CUDNN_DIM_MAX", {"HIPDNN_DIM_MAX", CONV_NUMERIC_LITERAL, API_DNN, HIP_UNSUPPORTED}}, // 8
|
||||
{"CUDNN_LRN_MIN_N", {"HIPDNN_LRN_MIN_N", CONV_NUMERIC_LITERAL, API_DNN, HIP_UNSUPPORTED}}, // 1
|
||||
{"CUDNN_LRN_MAX_N", {"HIPDNN_LRN_MAX_N", CONV_NUMERIC_LITERAL, API_DNN, HIP_UNSUPPORTED}}, // 16
|
||||
{"CUDNN_LRN_MIN_K", {"HIPDNN_LRN_MIN_K", CONV_NUMERIC_LITERAL, API_DNN, HIP_UNSUPPORTED}}, // 1e-5
|
||||
{"CUDNN_LRN_MIN_BETA", {"HIPDNN_LRN_MIN_BETA", CONV_NUMERIC_LITERAL, API_DNN, HIP_UNSUPPORTED}}, // 0.01
|
||||
{"CUDNN_BN_MIN_EPSILON", {"HIPDNN_BN_MIN_EPSILON", CONV_NUMERIC_LITERAL, API_DNN}}, // 1e-5
|
||||
{"CUDNN_SEV_ERROR_EN", {"HIPDNN_SEV_ERROR_EN", CONV_NUMERIC_LITERAL, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"CUDNN_SEV_WARNING_EN", {"HIPDNN_SEV_WARNING_EN", CONV_NUMERIC_LITERAL, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"CUDNN_SEV_INFO_EN", {"HIPDNN_SEV_INFO_EN", CONV_NUMERIC_LITERAL, API_DNN, HIP_UNSUPPORTED}},
|
||||
|
||||
// cuDNN enums
|
||||
{"cudnnStatus_t", {"hipdnnStatus_t", CONV_TYPE, API_DNN}},
|
||||
{"CUDNN_STATUS_SUCCESS", {"HIPDNN_STATUS_SUCCESS", CONV_NUMERIC_LITERAL, API_DNN}}, // 0
|
||||
{"CUDNN_STATUS_NOT_INITIALIZED", {"HIPDNN_STATUS_NOT_INITIALIZED", CONV_NUMERIC_LITERAL, API_DNN}}, // 1
|
||||
{"CUDNN_STATUS_ALLOC_FAILED", {"HIPDNN_STATUS_ALLOC_FAILED", CONV_NUMERIC_LITERAL, API_DNN}}, // 2
|
||||
{"CUDNN_STATUS_BAD_PARAM", {"HIPDNN_STATUS_BAD_PARAM", CONV_NUMERIC_LITERAL, API_DNN}}, // 3
|
||||
{"CUDNN_STATUS_INTERNAL_ERROR", {"HIPDNN_STATUS_INTERNAL_ERROR", CONV_NUMERIC_LITERAL, API_DNN}}, // 4
|
||||
{"CUDNN_STATUS_INVALID_VALUE", {"HIPDNN_STATUS_INVALID_VALUE", CONV_NUMERIC_LITERAL, API_DNN}}, // 5
|
||||
{"CUDNN_STATUS_ARCH_MISMATCH", {"HIPDNN_STATUS_ARCH_MISMATCH", CONV_NUMERIC_LITERAL, API_DNN}}, // 6
|
||||
{"CUDNN_STATUS_MAPPING_ERROR", {"HIPDNN_STATUS_MAPPING_ERROR", CONV_NUMERIC_LITERAL, API_DNN}}, // 7
|
||||
{"CUDNN_STATUS_EXECUTION_FAILED", {"HIPDNN_STATUS_EXECUTION_FAILED", CONV_NUMERIC_LITERAL, API_DNN}}, // 8
|
||||
{"CUDNN_STATUS_NOT_SUPPORTED", {"HIPDNN_STATUS_NOT_SUPPORTED", CONV_NUMERIC_LITERAL, API_DNN}}, // 9
|
||||
{"CUDNN_STATUS_LICENSE_ERROR", {"HIPDNN_STATUS_LICENSE_ERROR", CONV_NUMERIC_LITERAL, API_DNN}}, // 10
|
||||
{"CUDNN_STATUS_RUNTIME_PREREQUISITE_MISSING", {"HIPDNN_STATUS_RUNTIME_PREREQUISITE_MISSING", CONV_NUMERIC_LITERAL, API_DNN}}, // 11
|
||||
{"CUDNN_STATUS_RUNTIME_IN_PROGRESS", {"HIPDNN_STATUS_RUNTIME_IN_PROGRESS", CONV_NUMERIC_LITERAL, API_DNN, HIP_UNSUPPORTED}}, // 12
|
||||
{"CUDNN_STATUS_RUNTIME_FP_OVERFLOW", {"HIPDNN_STATUS_RUNTIME_FP_OVERFLOW", CONV_NUMERIC_LITERAL, API_DNN, HIP_UNSUPPORTED}}, // 13
|
||||
{"cudnnRuntimeTag_t", {"hipdnnRuntimeTag_t", CONV_TYPE, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnConvolutionMode_t", {"hipdnnConvolutionMode_t", CONV_TYPE, API_DNN}},
|
||||
{"CUDNN_CONVOLUTION", {"HIPDNN_CONVOLUTION", CONV_NUMERIC_LITERAL, API_DNN}}, // 0
|
||||
{"CUDNN_CROSS_CORRELATION", {"HIPDNN_CROSS_CORRELATION", CONV_NUMERIC_LITERAL, API_DNN}}, // 1
|
||||
{"cudnnTensorFormat_t", {"hipdnnTensorFormat_t", CONV_TYPE, API_DNN}},
|
||||
{"CUDNN_TENSOR_NCHW", {"HIPDNN_TENSOR_NCHW", CONV_NUMERIC_LITERAL, API_DNN}}, // 0
|
||||
{"CUDNN_TENSOR_NHWC", {"HIPDNN_TENSOR_NHWC", CONV_NUMERIC_LITERAL, API_DNN}}, // 1
|
||||
{"CUDNN_TENSOR_NCHW_VECT_C", {"HIPDNN_TENSOR_NCHW_VECT_C", CONV_NUMERIC_LITERAL, API_DNN}}, // 2
|
||||
{"cudnnDataType_t", {"hipdnnDataType_t", CONV_TYPE, API_DNN}},
|
||||
{"CUDNN_DATA_FLOAT", {"HIPDNN_DATA_FLOAT", CONV_NUMERIC_LITERAL, API_DNN}}, // 0
|
||||
{"CUDNN_DATA_DOUBLE", {"HIPDNN_DATA_DOUBLE", CONV_NUMERIC_LITERAL, API_DNN}}, // 1
|
||||
{"CUDNN_DATA_HALF", {"HIPDNN_DATA_HALF", CONV_NUMERIC_LITERAL, API_DNN}}, // 2
|
||||
{"CUDNN_DATA_INT8", {"HIPDNN_DATA_INT8", CONV_NUMERIC_LITERAL, API_DNN}}, // 3
|
||||
{"CUDNN_DATA_INT32", {"HIPDNN_DATA_INT32", CONV_NUMERIC_LITERAL, API_DNN}}, // 4
|
||||
{"CUDNN_DATA_INT8x4", {"HIPDNN_DATA_INT8x4", CONV_NUMERIC_LITERAL, API_DNN}}, // 5
|
||||
{"CUDNN_DATA_UINT8", {"HIPDNN_DATA_UINT8", CONV_NUMERIC_LITERAL, API_DNN, HIP_UNSUPPORTED}}, // 6
|
||||
{"CUDNN_DATA_UINT8x4", {"HIPDNN_DATA_UINT8x4", CONV_NUMERIC_LITERAL, API_DNN, HIP_UNSUPPORTED}}, // 7
|
||||
{"cudnnErrQueryMode_t", {"hipdnnErrQueryMode_t", CONV_TYPE, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"CUDNN_ERRQUERY_RAWCODE", {"HIPDNN_ERRQUERY_RAWCODE", CONV_NUMERIC_LITERAL, API_DNN, HIP_UNSUPPORTED}}, // 0
|
||||
{"CUDNN_ERRQUERY_NONBLOCKING", {"HIPDNN_ERRQUERY_NONBLOCKING", CONV_NUMERIC_LITERAL, API_DNN, HIP_UNSUPPORTED}}, // 1
|
||||
{"CUDNN_ERRQUERY_BLOCKING", {"HIPDNN_ERRQUERY_BLOCKING", CONV_NUMERIC_LITERAL, API_DNN, HIP_UNSUPPORTED}}, // 2
|
||||
{"cudnnSeverity_t", {"hipdnnSeverity_t", CONV_TYPE, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"CUDNN_SEV_FATAL", {"HIPDNN_SEV_FATAL", CONV_NUMERIC_LITERAL, API_DNN, HIP_UNSUPPORTED}}, // 0
|
||||
{"CUDNN_SEV_ERROR", {"HIPDNN_SEV_ERROR", CONV_NUMERIC_LITERAL, API_DNN, HIP_UNSUPPORTED}}, // 1
|
||||
{"CUDNN_SEV_WARNING", {"HIPDNN_SEV_WARNING", CONV_NUMERIC_LITERAL, API_DNN, HIP_UNSUPPORTED}}, // 2
|
||||
{"CUDNN_SEV_INFO", {"HIPDNN_SEV_INFO", CONV_NUMERIC_LITERAL, API_DNN, HIP_UNSUPPORTED}}, // 3
|
||||
{"cudnnConvolutionFwdAlgo_t", {"hipdnnConvolutionFwdAlgo_t", CONV_TYPE, API_DNN}},
|
||||
{"CUDNN_CONVOLUTION_FWD_ALGO_IMPLICIT_GEMM", {"HIPDNN_CONVOLUTION_FWD_ALGO_IMPLICIT_GEMM", CONV_NUMERIC_LITERAL, API_DNN}}, // 0
|
||||
{"CUDNN_CONVOLUTION_FWD_ALGO_IMPLICIT_PRECOMP_GEMM", {"HIPDNN_CONVOLUTION_FWD_ALGO_IMPLICIT_PRECOMP_GEMM", CONV_NUMERIC_LITERAL, API_DNN}}, // 1
|
||||
@@ -3000,42 +2996,6 @@ const std::map<llvm::StringRef, hipCounter> CUDA_IDENTIFIER_MAP{
|
||||
{"cudnnCTCLossAlgo_t", {"hipdnnCTCLossAlgo_t", CONV_TYPE, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"CUDNN_CTC_LOSS_ALGO_DETERMINISTIC", {"HIPDNN_CTC_LOSS_ALGO_DETERMINISTIC", CONV_NUMERIC_LITERAL, API_DNN, HIP_UNSUPPORTED}}, // 0
|
||||
{"CUDNN_CTC_LOSS_ALGO_NON_DETERMINISTIC", {"HIPDNN_CTC_LOSS_ALGO_NON_DETERMINISTIC", CONV_NUMERIC_LITERAL, API_DNN, HIP_UNSUPPORTED}}, // 1
|
||||
|
||||
{"cudnnFilterDescriptor_t", {"hipdnnFilterDescriptor_t", CONV_TYPE, API_DNN}},
|
||||
{"cudnnDropoutDescriptor_t", {"hipdnnDropoutDescriptor_t", CONV_TYPE, API_DNN}},
|
||||
{"cudnnConvolutionFwdAlgoPerf_t", {"hipdnnConvolutionFwdAlgoPerf_t", CONV_TYPE, API_DNN}},
|
||||
{"cudnnConvolutionBwdFilterAlgoPerf_t", {"hipdnnConvolutionBwdFilterAlgoPerf_t", CONV_TYPE, API_DNN}},
|
||||
{"cudnnRNNDescriptor_t", {"hipdnnRNNDescriptor_t", CONV_TYPE, API_DNN}},
|
||||
{"cudnnPersistentRNNPlan", {"hipdnnPersistentRNNPlan", CONV_TYPE, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnPersistentRNNPlan_t", {"hipdnnPersistentRNNPlan_t", CONV_TYPE, API_DNN}},
|
||||
{"cudnnTensorStruct", {"hipdnnTensorStruct", CONV_TYPE, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnConvolutionStruct", {"hipdnnConvolutionStruct", CONV_TYPE, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnPoolingStruct", {"hipdnnPoolingStruct", CONV_TYPE, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnPoolingDescriptor_t", {"hipdnnPoolingDescriptor_t", CONV_TYPE, API_DNN}},
|
||||
{"cudnnFilterStruct", {"hipdnnFilterStruct", CONV_TYPE, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnLRNDescriptor_t", {"hipdnnLRNDescriptor_t", CONV_TYPE, API_DNN}},
|
||||
{"cudnnLRNStruct", {"hipdnnLRNStruct", CONV_TYPE, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnActivationDescriptor_t", {"hipdnnActivationDescriptor_t", CONV_TYPE, API_DNN}},
|
||||
{"cudnnActivationStruct", {"hipdnnActivationStruct", CONV_TYPE, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnSpatialTransformerDescriptor_t", {"hipdnnSpatialTransformerDescriptor_t", CONV_TYPE, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnSpatialTransformerStruct", {"hipdnnSpatialTransformerStruct", CONV_TYPE, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnOpTensorDescriptor_t", {"hipdnnOpTensorDescriptor_t", CONV_TYPE, API_DNN}},
|
||||
{"cudnnOpTensorStruct", {"hipdnnOpTensorStruct", CONV_TYPE, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnReduceTensorDescriptor_t", {"hipdnnReduceTensorDescriptor_t", CONV_TYPE, API_DNN}},
|
||||
{"cudnnReduceTensorStruct", {"hipdnnReduceTensorStruct", CONV_TYPE, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnCTCLossDescriptor_t", {"hipdnnCTCLossDescriptor_t", CONV_TYPE, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnCTCLossStruct", {"hipdnnCTCLossStruct", CONV_TYPE, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnConvolutionBwdDataAlgoPerf_t", {"hipdnnConvolutionBwdDataAlgoPerf_t", CONV_TYPE, API_DNN}},
|
||||
{"cudnnAlgorithmDescriptor_t", {"hipdnnAlgorithmDescriptor_t", CONV_TYPE, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnAlgorithmStruct", {"hipdnnAlgorithmStruct", CONV_TYPE, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnAlgorithmPerformance_t", {"hipdnnAlgorithmPerformance_t", CONV_TYPE, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnAlgorithmPerformanceStruct", {"hipdnnAlgorithmPerformanceStruct", CONV_TYPE, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnRNNStruct", {"hipdnnRNNStruct", CONV_TYPE, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnRNNDescriptor_t", {"hipdnnRNNDescriptor_t", CONV_TYPE, API_DNN}},
|
||||
{"cudnnAlgorithm_t", {"hipdnnAlgorithm_t", CONV_TYPE, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnCallback_t", {"hipdnnCallback_t", CONV_TYPE, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnDebug_t", {"hipdnnDebug_t", CONV_TYPE, API_DNN, HIP_UNSUPPORTED}},
|
||||
|
||||
{"cudnnLRNMode_t", {"hipdnnLRNMode_t", CONV_TYPE, API_DNN}},
|
||||
{"CUDNN_LRN_CROSS_CHANNEL_DIM1", {"HIPDNN_LRN_CROSS_CHANNEL", CONV_NUMERIC_LITERAL, API_DNN}}, // 0 vs 1
|
||||
{"cudnnRNNInputMode_t", {"hipdnnRNNInputMode_t", CONV_TYPE, API_DNN}},
|
||||
@@ -3136,180 +3096,250 @@ const std::map<llvm::StringRef, hipCounter> CUDA_IDENTIFIER_MAP{
|
||||
{"cudnnSamplerType_t", {"hipdnnSamplerType_t", CONV_TYPE, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"CUDNN_SAMPLER_BILINEAR", {"HIPDNN_SAMPLER_BILINEAR", CONV_NUMERIC_LITERAL, API_DNN, HIP_UNSUPPORTED}}, // 0
|
||||
|
||||
{"cudnnGetVersion", {"hipdnnGetVersion", CONV_VERSION, API_DNN}},
|
||||
{"cudnnGetCudartVersion", {"hipdnnGetCudartVersion", CONV_VERSION, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnQueryRuntimeError", {"hipdnnQueryRuntimeError", CONV_VERSION, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnGetProperty", {"hipdnnGetProperty", CONV_VERSION, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnGetErrorString", {"hipdnnGetErrorString", CONV_ERROR, API_DNN}},
|
||||
// cuDNN types
|
||||
{"cudnnContext", {"hipdnnContext", CONV_TYPE, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnHandle_t", {"hipdnnHandle_t", CONV_TYPE, API_DNN}},
|
||||
{"cudnnTensorStruct", {"hipdnnTensorStruct", CONV_TYPE, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnTensorDescriptor_t", {"hipdnnTensorDescriptor_t", CONV_TYPE, API_DNN}},
|
||||
{"cudnnConvolutionStruct", {"hipdnnConvolutionStruct", CONV_TYPE, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnConvolutionDescriptor_t", {"hipdnnConvolutionDescriptor_t", CONV_TYPE, API_DNN}},
|
||||
{"cudnnPoolingStruct", {"hipdnnPoolingStruct", CONV_TYPE, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnPoolingDescriptor_t", {"hipdnnPoolingDescriptor_t", CONV_TYPE, API_DNN}},
|
||||
{"cudnnFilterStruct", {"hipdnnFilterStruct", CONV_TYPE, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnFilterDescriptor_t", {"hipdnnFilterDescriptor_t", CONV_TYPE, API_DNN}},
|
||||
{"cudnnLRNStruct", {"hipdnnLRNStruct", CONV_TYPE, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnLRNDescriptor_t", {"hipdnnLRNDescriptor_t", CONV_TYPE, API_DNN}},
|
||||
{"cudnnActivationStruct", {"hipdnnActivationStruct", CONV_TYPE, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnActivationDescriptor_t", {"hipdnnActivationDescriptor_t", CONV_TYPE, API_DNN}},
|
||||
{"cudnnSpatialTransformerStruct", {"hipdnnSpatialTransformerStruct", CONV_TYPE, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnSpatialTransformerDescriptor_t", {"hipdnnSpatialTransformerDescriptor_t", CONV_TYPE, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnOpTensorStruct", {"hipdnnOpTensorStruct", CONV_TYPE, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnOpTensorDescriptor_t", {"hipdnnOpTensorDescriptor_t", CONV_TYPE, API_DNN}},
|
||||
{"cudnnReduceTensorStruct", {"hipdnnReduceTensorStruct", CONV_TYPE, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnReduceTensorDescriptor_t", {"hipdnnReduceTensorDescriptor_t", CONV_TYPE, API_DNN}},
|
||||
{"cudnnCTCLossStruct", {"hipdnnCTCLossStruct", CONV_TYPE, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnCTCLossDescriptor_t", {"hipdnnCTCLossDescriptor_t", CONV_TYPE, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnConvolutionFwdAlgoPerf_t", {"hipdnnConvolutionFwdAlgoPerf_t", CONV_TYPE, API_DNN}},
|
||||
{"cudnnConvolutionBwdFilterAlgoPerf_t", {"hipdnnConvolutionBwdFilterAlgoPerf_t", CONV_TYPE, API_DNN}},
|
||||
{"cudnnConvolutionBwdDataAlgoPerf_t", {"hipdnnConvolutionBwdDataAlgoPerf_t", CONV_TYPE, API_DNN}},
|
||||
{"cudnnDropoutStruct", {"hipdnnDropoutStruct", CONV_TYPE, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnDropoutDescriptor_t", {"hipdnnDropoutDescriptor_t", CONV_TYPE, API_DNN}},
|
||||
{"cudnnAlgorithmStruct", {"hipdnnAlgorithmStruct", CONV_TYPE, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnAlgorithmDescriptor_t", {"hipdnnAlgorithmDescriptor_t", CONV_TYPE, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnAlgorithmPerformanceStruct", {"hipdnnAlgorithmPerformanceStruct", CONV_TYPE, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnAlgorithmPerformance_t", {"hipdnnAlgorithmPerformance_t", CONV_TYPE, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnRNNStruct", {"hipdnnRNNStruct", CONV_TYPE, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnRNNDescriptor_t", {"hipdnnRNNDescriptor_t", CONV_TYPE, API_DNN}},
|
||||
{"cudnnPersistentRNNPlan", {"hipdnnPersistentRNNPlan", CONV_TYPE, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnPersistentRNNPlan_t", {"hipdnnPersistentRNNPlan_t", CONV_TYPE, API_DNN}},
|
||||
{"cudnnAlgorithm_t", {"hipdnnAlgorithm_t", CONV_TYPE, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnDebug_t", {"hipdnnDebug_t", CONV_TYPE, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnCallback_t", {"hipdnnCallback_t", CONV_TYPE, API_DNN, HIP_UNSUPPORTED}},
|
||||
|
||||
{"cudnnCreate", {"hipdnnCreate", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnCreateTensorDescriptor", {"hipdnnCreateTensorDescriptor", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnCreateDropoutDescriptor", {"hipdnnCreateDropoutDescriptor", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnCreateReduceTensorDescriptor", {"hipdnnCreateReduceTensorDescriptor", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnSetReduceTensorDescriptor", {"hipdnnSetReduceTensorDescriptor", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnGetReduceTensorDescriptor", {"hipdnnGetReduceTensorDescriptor", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnGetReductionIndicesSize", {"hipdnnGetReductionIndicesSize", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnGetReductionWorkspaceSize", {"hipdnnGetReductionWorkspaceSize", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnCreateOpTensorDescriptor", {"hipdnnCreateOpTensorDescriptor", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnSetOpTensorDescriptor", {"hipdnnSetOpTensorDescriptor", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnGetOpTensorDescriptor", {"hipdnnGetOpTensorDescriptor", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnCreateRNNDescriptor", {"hipdnnCreateRNNDescriptor", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnSetStream", {"hipdnnSetStream", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnGetStream", {"hipdnnGetStream", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnSetRNNDescriptor_v5", {"hipdnnSetRNNDescriptor_v5", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnSetRNNDescriptor_v6", {"hipdnnSetRNNDescriptor_v6", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnSetRNNDescriptor", {"hipdnnSetRNNDescriptor", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnDropoutGetStatesSize", {"hipdnnDropoutGetStatesSize", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnDropoutGetReserveSpaceSize", {"hipdnnDropoutGetReserveSpaceSize", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnTransformTensor", {"hipdnnTransformTensor", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnSetTensor4dDescriptor", {"hipdnnSetTensor4dDescriptor", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnGetTensor4dDescriptor", {"hipdnnGetTensor4dDescriptor", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnAddTensor", {"hipdnnAddTensor", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnOpTensor", {"hipdnnOpTensor", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnGetTensorSizeInBytes", {"hipdnnGetTensorSizeInBytes", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnSetTensor4dDescriptorEx", {"hipdnnSetTensor4dDescriptorEx", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnSetTensorNdDescriptor", {"hipdnnSetTensorNdDescriptor", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnGetTensorNdDescriptor", {"hipdnnGetTensorNdDescriptor", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnSetTensorNdDescriptorEx", {"hipdnnSetTensorNdDescriptorEx", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnFindConvolutionForwardAlgorithm", {"hipdnnFindConvolutionForwardAlgorithm", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnFindConvolutionForwardAlgorithmEx", {"hipdnnFindConvolutionForwardAlgorithmEx", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnConvolutionBackwardFilter", {"hipdnnConvolutionBackwardFilter", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnConvolutionBackwardData", {"hipdnnConvolutionBackwardData", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnFindConvolutionBackwardFilterAlgorithm", {"hipdnnFindConvolutionBackwardFilterAlgorithm", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnFindConvolutionBackwardFilterAlgorithmEx", {"hipdnnFindConvolutionBackwardFilterAlgorithmEx", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnGetConvolutionBackwardFilterAlgorithm", {"hipdnnGetConvolutionBackwardFilterAlgorithm", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnGetConvolutionBackwardFilterAlgorithm_v7", {"hipdnnGetConvolutionBackwardFilterAlgorithm_v7", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnGetConvolutionBackwardFilterWorkspaceSize",{"hipdnnGetConvolutionBackwardFilterWorkspaceSize",CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnGetConvolutionBackwardDataWorkspaceSize", {"hipdnnGetConvolutionBackwardDataWorkspaceSize", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnGetConvolutionBackwardDataAlgorithm", {"hipdnnGetConvolutionBackwardDataAlgorithm", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnGetConvolutionBackwardDataAlgorithm_v7", {"hipdnnGetConvolutionBackwardDataAlgorithm_v7", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnGetConvolutionBackwardDataAlgorithmMaxCount", {"hipdnnGetConvolutionBackwardDataAlgorithmMaxCount", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnGetConvolutionForwardAlgorithmMaxCount", {"hipdnnGetConvolutionForwardAlgorithmMaxCount", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnGetRNNLinLayerMatrixParams", {"hipdnnGetRNNLinLayerMatrixParams", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnGetRNNLinLayerBiasParams", {"hipdnnGetRNNLinLayerBiasParams", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnSetRNNProjectionLayers", {"hipdnnSetRNNProjectionLayers", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnGetRNNProjectionLayers", {"hipdnnGetRNNProjectionLayers", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnSetRNNAlgorithmDescriptor", {"hipdnnSetRNNAlgorithmDescriptor", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnGetRNNDescriptor", {"hipdnnGetRNNDescriptor", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnGetFilterNdDescriptor", {"hipdnnGetFilterNdDescriptor", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnFindConvolutionBackwardDataAlgorithm", {"hipdnnFindConvolutionBackwardDataAlgorithm", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnFindConvolutionBackwardDataAlgorithmEx", {"hipdnnFindConvolutionBackwardDataAlgorithmEx", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnSetDropoutDescriptor", {"hipdnnSetDropoutDescriptor", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnRestoreDropoutDescriptor", {"hipdnnRestoreDropoutDescriptor", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnGetDropoutDescriptor", {"hipdnnGetDropoutDescriptor", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnGetDropoutDescriptor", {"hipdnnGetDropoutDescriptor", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnSetConvolution2dDescriptor", {"hipdnnSetConvolution2dDescriptor", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnGetConvolution2dDescriptor", {"hipdnnGetConvolution2dDescriptor", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnSetConvolutionMathType", {"hipdnnSetConvolutionMathType", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnDropoutForward", {"hipdnnDropoutForward", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnDropoutBackward", {"hipdnnDropoutBackward", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnGetConvolutionMathType", {"hipdnnGetConvolutionMathType", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnSetConvolutionGroupCount", {"hipdnnSetConvolutionGroupCount", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnGetConvolutionGroupCount", {"hipdnnGetConvolutionGroupCount", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnGetConvolution2dForwardOutputDim", {"hipdnnGetConvolution2dForwardOutputDim", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnSetConvolutionNdDescriptor", {"hipdnnSetConvolutionNdDescriptor", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnGetConvolutionNdDescriptor", {"hipdnnGetConvolutionNdDescriptor", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnGetConvolutionNdForwardOutputDim", {"hipdnnGetConvolutionNdForwardOutputDim", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnCreateFilterDescriptor", {"hipdnnCreateFilterDescriptor", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnCreatePersistentRNNPlan", {"hipdnnCreatePersistentRNNPlan", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnSetPersistentRNNPlan", {"hipdnnSetPersistentRNNPlan", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnRNNForwardInference", {"hipdnnRNNForwardInference", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnRNNBackwardWeights", {"hipdnnRNNBackwardWeights", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnGetRNNParamsSize", {"hipdnnGetRNNParamsSize", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnGetRNNWorkspaceSize", {"hipdnnGetRNNWorkspaceSize", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnGetRNNTrainingReserveSize", {"hipdnnGetRNNTrainingReserveSize", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnSetFilterNdDescriptor", {"hipdnnSetFilterNdDescriptor", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnRNNForwardTraining", {"hipdnnRNNForwardTraining", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnRNNBackwardData", {"hipdnnRNNBackwardData", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnSetFilter4dDescriptor", {"hipdnnSetFilter4dDescriptor", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnSetRNNMatrixMathType", {"hipdnnSetRNNMatrixMathType", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnGetRNNMatrixMathType", {"hipdnnGetRNNMatrixMathType", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnGetRNNForwardInferenceAlgorithmMaxCount", {"hipdnnGetRNNForwardInferenceAlgorithmMaxCount", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnFindRNNForwardInferenceAlgorithmEx", {"hipdnnFindRNNForwardInferenceAlgorithmEx", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnGetRNNForwardTrainingAlgorithmMaxCount", {"hipdnnGetRNNForwardTrainingAlgorithmMaxCount", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnFindRNNForwardTrainingAlgorithmEx", {"hipdnnFindRNNForwardTrainingAlgorithmEx", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnGetRNNBackwardDataAlgorithmMaxCount", {"hipdnnGetRNNBackwardDataAlgorithmMaxCount", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnFindRNNBackwardDataAlgorithmEx", {"hipdnnFindRNNBackwardDataAlgorithmEx", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnGetRNNBackwardWeightsAlgorithmMaxCount", {"hipdnnGetRNNBackwardWeightsAlgorithmMaxCount", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnFindRNNBackwardWeightsAlgorithmEx", {"hipdnnFindRNNBackwardWeightsAlgorithmEx", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnCreateConvolutionDescriptor", {"hipdnnCreateConvolutionDescriptor", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnGetConvolutionForwardAlgorithm", {"hipdnnGetConvolutionForwardAlgorithm", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnGetConvolutionForwardAlgorithm_v7", {"hipdnnGetConvolutionForwardAlgorithm_v7", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnConvolutionForward", {"hipdnnConvolutionForward", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnGetConvolutionForwardWorkspaceSize", {"hipdnnGetConvolutionForwardWorkspaceSize", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnConvolutionBiasActivationForward", {"hipdnnConvolutionBiasActivationForward", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnGetConvolutionBackwardFilterAlgorithmMaxCount", {"hipdnnGetConvolutionBackwardFilterAlgorithmMaxCount", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnConvolutionBackwardBias", {"hipdnnConvolutionBackwardBias", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnReduceTensor", {"hipdnnReduceTensor", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnSetTensor", {"hipdnnSetTensor", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnScaleTensor", {"hipdnnScaleTensor", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnIm2Col", {"hipdnnIm2Col", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnDestroyTensorDescriptor", {"hipdnnDestroyTensorDescriptor", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnDestroyOpTensorDescriptor", {"hipdnnDestroyOpTensorDescriptor", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnDestroyConvolutionDescriptor", {"hipdnnDestroyConvolutionDescriptor", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnDestroyDropoutDescriptor", {"hipdnnDestroyDropoutDescriptor", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnDestroyFilterDescriptor", {"hipdnnDestroyFilterDescriptor", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnDestroyRNNDescriptor", {"hipdnnDestroyRNNDescriptor", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnDestroyReduceTensorDescriptor", {"hipdnnDestroyReduceTensorDescriptor", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnDestroyPersistentRNNPlan", {"hipdnnDestroyPersistentRNNPlan", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnDestroy", {"hipdnnDestroy", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnSoftmaxForward", {"hipdnnSoftmaxForward", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnSoftmaxBackward", {"hipdnnSoftmaxBackward", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnCreatePoolingDescriptor", {"hipdnnCreatePoolingDescriptor", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnSetPooling2dDescriptor", {"hipdnnSetPooling2dDescriptor", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnGetPooling2dDescriptor", {"hipdnnGetPooling2dDescriptor", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnSetPoolingNdDescriptor", {"hipdnnSetPoolingNdDescriptor", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnGetPoolingNdDescriptor", {"hipdnnGetPoolingNdDescriptor", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnGetPoolingNdForwardOutputDim", {"hipdnnGetPoolingNdForwardOutputDim", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnGetPooling2dForwardOutputDim", {"hipdnnGetPooling2dForwardOutputDim", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnDestroyPoolingDescriptor", {"hipdnnDestroyPoolingDescriptor", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnPoolingForward", {"hipdnnPoolingForward", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnPoolingBackward", {"hipdnnPoolingBackward", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnCreateActivationDescriptor", {"hipdnnCreateActivationDescriptor", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnSetActivationDescriptor", {"hipdnnSetActivationDescriptor", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnGetActivationDescriptor", {"hipdnnGetActivationDescriptor", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnDestroyActivationDescriptor", {"hipdnnDestroyActivationDescriptor", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnActivationForward", {"hipdnnActivationForward", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnActivationBackward", {"hipdnnActivationBackward", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnCreateLRNDescriptor", {"hipdnnCreateLRNDescriptor", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnSetLRNDescriptor", {"hipdnnSetLRNDescriptor", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnGetLRNDescriptor", {"hipdnnGetLRNDescriptor", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnDestroyLRNDescriptor", {"hipdnnDestroyLRNDescriptor", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnLRNCrossChannelForward", {"hipdnnLRNCrossChannelForward", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnLRNCrossChannelBackward", {"hipdnnLRNCrossChannelBackward", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnDivisiveNormalizationForward", {"hipdnnDivisiveNormalizationForward", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnDivisiveNormalizationBackward", {"hipdnnDivisiveNormalizationBackward", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnDeriveBNTensorDescriptor", {"hipdnnDeriveBNTensorDescriptor", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnBatchNormalizationForwardTraining", {"hipdnnBatchNormalizationForwardTraining", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnBatchNormalizationForwardInference", {"hipdnnBatchNormalizationForwardInference", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnBatchNormalizationBackward", {"hipdnnBatchNormalizationBackward", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnCreateSpatialTransformerDescriptor", {"hipdnnCreateSpatialTransformerDescriptor", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnSetSpatialTransformerNdDescriptor", {"hipdnnSetSpatialTransformerNdDescriptor", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnDestroySpatialTransformerDescriptor", {"hipdnnDestroySpatialTransformerDescriptor", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnSpatialTfGridGeneratorForward", {"hipdnnSpatialTfGridGeneratorForward", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnSpatialTfGridGeneratorBackward", {"hipdnnSpatialTfGridGeneratorBackward", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnSpatialTfSamplerForward", {"hipdnnSpatialTfSamplerForward", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnSpatialTfSamplerBackward", {"hipdnnSpatialTfSamplerBackward", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnCreateCTCLossDescriptor", {"hipdnnCreateCTCLossDescriptor", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnSetCTCLossDescriptor", {"hipdnnSetCTCLossDescriptor", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnGetCTCLossDescriptor", {"hipdnnGetCTCLossDescriptor", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnDestroyCTCLossDescriptor", {"hipdnnDestroyCTCLossDescriptor", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnCTCLoss", {"hipdnnCTCLoss", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnGetCTCLossWorkspaceSize", {"hipdnnGetCTCLossWorkspaceSize", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnCreateAlgorithmDescriptor", {"hipdnnCreateAlgorithmDescriptor", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnSetAlgorithmDescriptor", {"hipdnnSetAlgorithmDescriptor", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnGetAlgorithmDescriptor", {"hipdnnGetAlgorithmDescriptor", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnCopyAlgorithmDescriptor", {"hipdnnCopyAlgorithmDescriptor", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnDestroyAlgorithmDescriptor", {"hipdnnDestroyAlgorithmDescriptor", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnCreateAlgorithmPerformance", {"hipdnnCreateAlgorithmPerformance", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnSetAlgorithmPerformance", {"hipdnnSetAlgorithmPerformance", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnGetAlgorithmPerformance", {"hipdnnGetAlgorithmPerformance", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnDestroyAlgorithmPerformance", {"hipdnnDestroyAlgorithmPerformance", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnGetAlgorithmSpaceSize", {"hipdnnGetAlgorithmSpaceSize", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnSaveAlgorithm", {"hipdnnSaveAlgorithm", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnRestoreAlgorithm", {"hipdnnRestoreAlgorithm", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnSetCallback", {"hipdnnSetCallback", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnGetCallback", {"hipdnnGetCallback", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
|
||||
///////////////////////////// cuDNN functions /////////////////////////////
|
||||
{"cudnnGetVersion", {"hipdnnGetVersion", CONV_VERSION, API_DNN}},
|
||||
{"cudnnGetCudartVersion", {"hipdnnGetCudartVersion", CONV_VERSION, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnQueryRuntimeError", {"hipdnnQueryRuntimeError", CONV_VERSION, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnGetProperty", {"hipdnnGetProperty", CONV_VERSION, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnGetErrorString", {"hipdnnGetErrorString", CONV_ERROR, API_DNN}},
|
||||
{"cudnnIm2Col", {"hipdnnIm2Col", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnCreate", {"hipdnnCreate", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnDestroy", {"hipdnnDestroy", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnSetStream", {"hipdnnSetStream", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnGetStream", {"hipdnnGetStream", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnSetCallback", {"hipdnnSetCallback", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnGetCallback", {"hipdnnGetCallback", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
|
||||
|
||||
// cuDNN Tensor functions
|
||||
{"cudnnCreateTensorDescriptor", {"hipdnnCreateTensorDescriptor", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnSetTensor4dDescriptor", {"hipdnnSetTensor4dDescriptor", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnSetTensor4dDescriptorEx", {"hipdnnSetTensor4dDescriptorEx", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnGetTensor4dDescriptor", {"hipdnnGetTensor4dDescriptor", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnSetTensorNdDescriptor", {"hipdnnSetTensorNdDescriptor", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnSetTensorNdDescriptorEx", {"hipdnnSetTensorNdDescriptorEx", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnGetTensorNdDescriptor", {"hipdnnGetTensorNdDescriptor", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnGetTensorSizeInBytes", {"hipdnnGetTensorSizeInBytes", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnDestroyTensorDescriptor", {"hipdnnDestroyTensorDescriptor", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnTransformTensor", {"hipdnnTransformTensor", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnAddTensor", {"hipdnnAddTensor", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnCreateOpTensorDescriptor", {"hipdnnCreateOpTensorDescriptor", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnSetOpTensorDescriptor", {"hipdnnSetOpTensorDescriptor", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnGetOpTensorDescriptor", {"hipdnnGetOpTensorDescriptor", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnDestroyOpTensorDescriptor", {"hipdnnDestroyOpTensorDescriptor", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnOpTensor", {"hipdnnOpTensor", CONV_MATH_FUNC, API_DNN}},
|
||||
|
||||
// cuDNN Reduce Tensor functions
|
||||
{"cudnnCreateReduceTensorDescriptor", {"hipdnnCreateReduceTensorDescriptor", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnSetReduceTensorDescriptor", {"hipdnnSetReduceTensorDescriptor", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnGetReduceTensorDescriptor", {"hipdnnGetReduceTensorDescriptor", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnDestroyReduceTensorDescriptor", {"hipdnnDestroyReduceTensorDescriptor", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnGetReductionIndicesSize", {"hipdnnGetReductionIndicesSize", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnGetReductionWorkspaceSize", {"hipdnnGetReductionWorkspaceSize", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnReduceTensor", {"hipdnnReduceTensor", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnSetTensor", {"hipdnnSetTensor", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnScaleTensor", {"hipdnnScaleTensor", CONV_MATH_FUNC, API_DNN}},
|
||||
|
||||
// cuDNN Filter functions
|
||||
{"cudnnCreateFilterDescriptor", {"hipdnnCreateFilterDescriptor", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnSetFilter4dDescriptor", {"hipdnnSetFilter4dDescriptor", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnGetFilter4dDescriptor", {"hipdnnGetFilter4dDescriptor", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnSetFilterNdDescriptor", {"hipdnnSetFilterNdDescriptor", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnGetFilterNdDescriptor", {"hipdnnGetFilterNdDescriptor", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnDestroyFilterDescriptor", {"hipdnnDestroyFilterDescriptor", CONV_MATH_FUNC, API_DNN}},
|
||||
|
||||
// cuDNN Convolution functions
|
||||
{"cudnnCreateConvolutionDescriptor", {"hipdnnCreateConvolutionDescriptor", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnSetConvolutionMathType", {"hipdnnSetConvolutionMathType", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnGetConvolutionMathType", {"hipdnnGetConvolutionMathType", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnSetConvolutionGroupCount", {"hipdnnSetConvolutionGroupCount", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnGetConvolutionGroupCount", {"hipdnnGetConvolutionGroupCount", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnSetConvolution2dDescriptor", {"hipdnnSetConvolution2dDescriptor", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnGetConvolution2dDescriptor", {"hipdnnGetConvolution2dDescriptor", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnGetConvolution2dForwardOutputDim", {"hipdnnGetConvolution2dForwardOutputDim", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnSetConvolutionNdDescriptor", {"hipdnnSetConvolutionNdDescriptor", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnGetConvolutionNdDescriptor", {"hipdnnGetConvolutionNdDescriptor", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnGetConvolutionNdForwardOutputDim", {"hipdnnGetConvolutionNdForwardOutputDim", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnDestroyConvolutionDescriptor", {"hipdnnDestroyConvolutionDescriptor", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnGetConvolutionForwardAlgorithmMaxCount", {"hipdnnGetConvolutionForwardAlgorithmMaxCount", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnFindConvolutionForwardAlgorithm", {"hipdnnFindConvolutionForwardAlgorithm", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnFindConvolutionForwardAlgorithmEx", {"hipdnnFindConvolutionForwardAlgorithmEx", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnGetConvolutionForwardAlgorithm", {"hipdnnGetConvolutionForwardAlgorithm", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnGetConvolutionForwardAlgorithm_v7", {"hipdnnGetConvolutionForwardAlgorithm_v7", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnGetConvolutionForwardWorkspaceSize", {"hipdnnGetConvolutionForwardWorkspaceSize", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnConvolutionForward", {"hipdnnConvolutionForward", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnConvolutionBiasActivationForward", {"hipdnnConvolutionBiasActivationForward", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnConvolutionBackwardBias", {"hipdnnConvolutionBackwardBias", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnGetConvolutionBackwardFilterAlgorithmMaxCount", {"hipdnnGetConvolutionBackwardFilterAlgorithmMaxCount", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnFindConvolutionBackwardFilterAlgorithm", {"hipdnnFindConvolutionBackwardFilterAlgorithm", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnFindConvolutionBackwardFilterAlgorithmEx", {"hipdnnFindConvolutionBackwardFilterAlgorithmEx", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnGetConvolutionBackwardFilterAlgorithm", {"hipdnnGetConvolutionBackwardFilterAlgorithm", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnGetConvolutionBackwardFilterAlgorithm_v7", {"hipdnnGetConvolutionBackwardFilterAlgorithm_v7", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnGetConvolutionBackwardFilterWorkspaceSize", {"hipdnnGetConvolutionBackwardFilterWorkspaceSize", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnConvolutionBackwardFilter", {"hipdnnConvolutionBackwardFilter", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnGetConvolutionBackwardDataAlgorithmMaxCount", {"hipdnnGetConvolutionBackwardDataAlgorithmMaxCount", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnFindConvolutionBackwardDataAlgorithm", {"hipdnnFindConvolutionBackwardDataAlgorithm", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnFindConvolutionBackwardDataAlgorithmEx", {"hipdnnFindConvolutionBackwardDataAlgorithmEx", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnGetConvolutionBackwardDataAlgorithm", {"hipdnnGetConvolutionBackwardDataAlgorithm", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnGetConvolutionBackwardDataAlgorithm_v7", {"hipdnnGetConvolutionBackwardDataAlgorithm_v7", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnGetConvolutionBackwardDataWorkspaceSize", {"hipdnnGetConvolutionBackwardDataWorkspaceSize", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnConvolutionBackwardData", {"hipdnnConvolutionBackwardData", CONV_MATH_FUNC, API_DNN}},
|
||||
|
||||
// cuDNN Sortmax functions
|
||||
{"cudnnSoftmaxForward", {"hipdnnSoftmaxForward", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnSoftmaxBackward", {"hipdnnSoftmaxBackward", CONV_MATH_FUNC, API_DNN}},
|
||||
|
||||
// cuDNN Pooling functions
|
||||
{"cudnnCreatePoolingDescriptor", {"hipdnnCreatePoolingDescriptor", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnSetPooling2dDescriptor", {"hipdnnSetPooling2dDescriptor", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnGetPooling2dDescriptor", {"hipdnnGetPooling2dDescriptor", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnSetPoolingNdDescriptor", {"hipdnnSetPoolingNdDescriptor", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnGetPoolingNdDescriptor", {"hipdnnGetPoolingNdDescriptor", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnGetPoolingNdForwardOutputDim", {"hipdnnGetPoolingNdForwardOutputDim", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnGetPooling2dForwardOutputDim", {"hipdnnGetPooling2dForwardOutputDim", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnDestroyPoolingDescriptor", {"hipdnnDestroyPoolingDescriptor", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnPoolingForward", {"hipdnnPoolingForward", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnPoolingBackward", {"hipdnnPoolingBackward", CONV_MATH_FUNC, API_DNN}},
|
||||
|
||||
// cuDNN Activation functions
|
||||
{"cudnnCreateActivationDescriptor", {"hipdnnCreateActivationDescriptor", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnSetActivationDescriptor", {"hipdnnSetActivationDescriptor", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnGetActivationDescriptor", {"hipdnnGetActivationDescriptor", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnDestroyActivationDescriptor", {"hipdnnDestroyActivationDescriptor", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnActivationForward", {"hipdnnActivationForward", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnActivationBackward", {"hipdnnActivationBackward", CONV_MATH_FUNC, API_DNN}},
|
||||
|
||||
// cuDNN LRN functions
|
||||
{"cudnnCreateLRNDescriptor", {"hipdnnCreateLRNDescriptor", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnSetLRNDescriptor", {"hipdnnSetLRNDescriptor", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnGetLRNDescriptor", {"hipdnnGetLRNDescriptor", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnDestroyLRNDescriptor", {"hipdnnDestroyLRNDescriptor", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnLRNCrossChannelForward", {"hipdnnLRNCrossChannelForward", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnLRNCrossChannelBackward", {"hipdnnLRNCrossChannelBackward", CONV_MATH_FUNC, API_DNN}},
|
||||
|
||||
// cuDNN Divisive Normalization functions
|
||||
{"cudnnDivisiveNormalizationForward", {"hipdnnDivisiveNormalizationForward", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnDivisiveNormalizationBackward", {"hipdnnDivisiveNormalizationBackward", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
|
||||
|
||||
// cuDNN Batch Normalization functions
|
||||
{"cudnnDeriveBNTensorDescriptor", {"hipdnnDeriveBNTensorDescriptor", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnBatchNormalizationForwardTraining", {"hipdnnBatchNormalizationForwardTraining", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnBatchNormalizationForwardInference", {"hipdnnBatchNormalizationForwardInference", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnBatchNormalizationBackward", {"hipdnnBatchNormalizationBackward", CONV_MATH_FUNC, API_DNN}},
|
||||
|
||||
// cuDNN Spatial Transformer functions
|
||||
{"cudnnCreateSpatialTransformerDescriptor", {"hipdnnCreateSpatialTransformerDescriptor", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnSetSpatialTransformerNdDescriptor", {"hipdnnSetSpatialTransformerNdDescriptor", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnDestroySpatialTransformerDescriptor", {"hipdnnDestroySpatialTransformerDescriptor", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnSpatialTfGridGeneratorForward", {"hipdnnSpatialTfGridGeneratorForward", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnSpatialTfGridGeneratorBackward", {"hipdnnSpatialTfGridGeneratorBackward", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnSpatialTfSamplerForward", {"hipdnnSpatialTfSamplerForward", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnSpatialTfSamplerBackward", {"hipdnnSpatialTfSamplerBackward", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
|
||||
|
||||
// cuDNN Dropout functions
|
||||
{"cudnnCreateDropoutDescriptor", {"hipdnnCreateDropoutDescriptor", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnDestroyDropoutDescriptor", {"hipdnnDestroyDropoutDescriptor", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnDropoutGetStatesSize", {"hipdnnDropoutGetStatesSize", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnDropoutGetReserveSpaceSize", {"hipdnnDropoutGetReserveSpaceSize", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnSetDropoutDescriptor", {"hipdnnSetDropoutDescriptor", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnGetDropoutDescriptor", {"hipdnnGetDropoutDescriptor", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnRestoreDropoutDescriptor", {"hipdnnRestoreDropoutDescriptor", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnDropoutForward", {"hipdnnDropoutForward", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnDropoutBackward", {"hipdnnDropoutBackward", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
|
||||
|
||||
// cuDNN RNN functions
|
||||
{"cudnnCreateRNNDescriptor", {"hipdnnCreateRNNDescriptor", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnDestroyRNNDescriptor", {"hipdnnDestroyRNNDescriptor", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnGetRNNForwardInferenceAlgorithmMaxCount", {"hipdnnGetRNNForwardInferenceAlgorithmMaxCount", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnFindRNNForwardInferenceAlgorithmEx", {"hipdnnFindRNNForwardInferenceAlgorithmEx", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnGetRNNForwardTrainingAlgorithmMaxCount", {"hipdnnGetRNNForwardTrainingAlgorithmMaxCount", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnFindRNNForwardTrainingAlgorithmEx", {"hipdnnFindRNNForwardTrainingAlgorithmEx", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnGetRNNBackwardDataAlgorithmMaxCount", {"hipdnnGetRNNBackwardDataAlgorithmMaxCount", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnFindRNNBackwardDataAlgorithmEx", {"hipdnnFindRNNBackwardDataAlgorithmEx", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnGetRNNBackwardWeightsAlgorithmMaxCount", {"hipdnnGetRNNBackwardWeightsAlgorithmMaxCount", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnFindRNNBackwardWeightsAlgorithmEx", {"hipdnnFindRNNBackwardWeightsAlgorithmEx", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnCreatePersistentRNNPlan", {"hipdnnCreatePersistentRNNPlan", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnSetPersistentRNNPlan", {"hipdnnSetPersistentRNNPlan", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnDestroyPersistentRNNPlan", {"hipdnnDestroyPersistentRNNPlan", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnSetRNNDescriptor", {"hipdnnSetRNNDescriptor", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnGetRNNDescriptor", {"hipdnnGetRNNDescriptor", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnSetRNNProjectionLayers", {"hipdnnSetRNNProjectionLayers", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnGetRNNProjectionLayers", {"hipdnnGetRNNProjectionLayers", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnSetRNNAlgorithmDescriptor", {"hipdnnSetRNNAlgorithmDescriptor", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnSetRNNMatrixMathType", {"hipdnnSetRNNMatrixMathType", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnGetRNNMatrixMathType", {"hipdnnGetRNNMatrixMathType", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnGetRNNWorkspaceSize", {"hipdnnGetRNNWorkspaceSize", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnGetRNNTrainingReserveSize", {"hipdnnGetRNNTrainingReserveSize", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnGetRNNParamsSize", {"hipdnnGetRNNParamsSize", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnGetRNNLinLayerMatrixParams", {"hipdnnGetRNNLinLayerMatrixParams", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnGetRNNLinLayerBiasParams", {"hipdnnGetRNNLinLayerBiasParams", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnRNNForwardInference", {"hipdnnRNNForwardInference", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnRNNForwardTraining", {"hipdnnRNNForwardTraining", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnRNNBackwardData", {"hipdnnRNNBackwardData", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnRNNBackwardWeights", {"hipdnnRNNBackwardWeights", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnSetRNNDescriptor_v5", {"hipdnnSetRNNDescriptor_v5", CONV_MATH_FUNC, API_DNN}},
|
||||
{"cudnnSetRNNDescriptor_v6", {"hipdnnSetRNNDescriptor_v6", CONV_MATH_FUNC, API_DNN}},
|
||||
|
||||
// cuDNN Connectionist Temporal Classification loss functions
|
||||
{"cudnnCreateCTCLossDescriptor", {"hipdnnCreateCTCLossDescriptor", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnSetCTCLossDescriptor", {"hipdnnSetCTCLossDescriptor", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnGetCTCLossDescriptor", {"hipdnnGetCTCLossDescriptor", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnDestroyCTCLossDescriptor", {"hipdnnDestroyCTCLossDescriptor", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnCTCLoss", {"hipdnnCTCLoss", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnGetCTCLossWorkspaceSize", {"hipdnnGetCTCLossWorkspaceSize", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
|
||||
|
||||
// cuDNN Algorithm functions
|
||||
{"cudnnCreateAlgorithmDescriptor", {"hipdnnCreateAlgorithmDescriptor", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnSetAlgorithmDescriptor", {"hipdnnSetAlgorithmDescriptor", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnGetAlgorithmDescriptor", {"hipdnnGetAlgorithmDescriptor", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnCopyAlgorithmDescriptor", {"hipdnnCopyAlgorithmDescriptor", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnDestroyAlgorithmDescriptor", {"hipdnnDestroyAlgorithmDescriptor", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnCreateAlgorithmPerformance", {"hipdnnCreateAlgorithmPerformance", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnSetAlgorithmPerformance", {"hipdnnSetAlgorithmPerformance", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnGetAlgorithmPerformance", {"hipdnnGetAlgorithmPerformance", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnDestroyAlgorithmPerformance", {"hipdnnDestroyAlgorithmPerformance", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnGetAlgorithmSpaceSize", {"hipdnnGetAlgorithmSpaceSize", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnSaveAlgorithm", {"hipdnnSaveAlgorithm", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnRestoreAlgorithm", {"hipdnnRestoreAlgorithm", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
|
||||
};
|
||||
|
||||
const std::map<llvm::StringRef, hipCounter>& CUDA_RENAMES_MAP() {
|
||||
|
||||
@@ -108,10 +108,10 @@ extern int HIP_TRACE_API;
|
||||
#endif
|
||||
#include <hip/hcc_detail/hip_atomic.h>
|
||||
#include <hip/hcc_detail/host_defines.h>
|
||||
#include <hip/hcc_detail/math_functions.h>
|
||||
#include <hip/hcc_detail/device_functions.h>
|
||||
#include <hip/hcc_detail/surface_functions.h>
|
||||
#if __HCC__
|
||||
#include <hip/hcc_detail/math_functions.h>
|
||||
#include <hip/hcc_detail/texture_functions.h>
|
||||
#endif // __HCC__
|
||||
|
||||
@@ -438,8 +438,6 @@ extern const __device__ __attribute__((weak)) __hip_builtin_gridDim_t gridDim;
|
||||
#define hipGridDim_y gridDim.y
|
||||
#define hipGridDim_z gridDim.z
|
||||
|
||||
#include <hip/hcc_detail/math_functions.h>
|
||||
|
||||
// Support std::complex.
|
||||
#pragma push_macro("__CUDA__")
|
||||
#define __CUDA__
|
||||
@@ -450,6 +448,8 @@ extern const __device__ __attribute__((weak)) __hip_builtin_gridDim_t gridDim;
|
||||
#undef __CUDA__
|
||||
#pragma pop_macro("__CUDA__")
|
||||
|
||||
#include <hip/hcc_detail/math_functions.h>
|
||||
|
||||
#endif
|
||||
|
||||
#endif // HIP_HCC_DETAIL_RUNTIME_H
|
||||
|
||||
@@ -123,7 +123,7 @@ THE SOFTWARE.
|
||||
typename std::enable_if<
|
||||
(rank > 1) && sizeof...(Us) == rank>::type* = nullptr>
|
||||
__host__ __device__
|
||||
HIP_vector_type(Us... xs) noexcept { data = Native_vec_{xs...}; }
|
||||
HIP_vector_type(Us... xs) noexcept { data = Native_vec_{static_cast<T>(xs)...}; }
|
||||
__host__ __device__
|
||||
HIP_vector_type(const HIP_vector_type&) = default;
|
||||
__host__ __device__
|
||||
|
||||
@@ -0,0 +1,159 @@
|
||||
// RUN: %run_test hipify "%s" "%t" %cuda_args
|
||||
|
||||
// CHECK: #include <hip/hip_runtime.h>
|
||||
#include <stdio.h>
|
||||
// CHECK: #include <hipDNN.h>
|
||||
#include <cudnn.h>
|
||||
|
||||
/**
|
||||
*
|
||||
* Author: Jon Gauthier <jon@gauthiers.net>
|
||||
* February 2015
|
||||
*
|
||||
. * Adopted for CUDA/CUDNN 9.0
|
||||
*/
|
||||
|
||||
void printMatrix(const double *mat, int m, int n) {
|
||||
for (int j = 0; j < n; j++) {
|
||||
for (int i = 0; i < m; i++) {
|
||||
printf("%f\n", mat[j * m + i]);
|
||||
}
|
||||
printf("\n\n");
|
||||
}
|
||||
}
|
||||
|
||||
double *makeDiffData(int m, int c) {
|
||||
double *diff = (double *) calloc(m * c, sizeof(double));
|
||||
for (int j = 0; j < m; j++) {
|
||||
int class_ = rand() % c;
|
||||
printf("%d class: %d\n", j, class_);
|
||||
for (int i = 0; i < c; i++)
|
||||
diff[j * c + i] = class_ == i ? -c / (double) m : 0;
|
||||
}
|
||||
|
||||
return diff;
|
||||
}
|
||||
|
||||
int main() {
|
||||
int m = 5, c = 4, numChannels = 1;
|
||||
|
||||
double *fcLayer = (double *) malloc(m * c * sizeof(double));
|
||||
for (int i = 0; i < m; i++) {
|
||||
double def = rand() % 25;
|
||||
for (int c_idx = 0; c_idx < c; c_idx++) {
|
||||
int offset = i * c + c_idx;
|
||||
fcLayer[offset] = def;
|
||||
}
|
||||
}
|
||||
printf("FC LAYER:\n");
|
||||
printMatrix(fcLayer, c, m);
|
||||
|
||||
double *d_fcLayer;
|
||||
// CHECK: hipMalloc((void**) &d_fcLayer, m * c * sizeof(double));
|
||||
cudaMalloc((void**) &d_fcLayer, m * c * sizeof(double));
|
||||
// CHECK: hipMemcpy(d_fcLayer, fcLayer, m * c * sizeof(double), hipMemcpyHostToDevice);
|
||||
cudaMemcpy(d_fcLayer, fcLayer, m * c * sizeof(double), cudaMemcpyHostToDevice);
|
||||
|
||||
double *d_softmaxData;
|
||||
// CHECK: hipMalloc((void**) &d_softmaxData, m * c * sizeof(double));
|
||||
cudaMalloc((void**) &d_softmaxData, m * c * sizeof(double));
|
||||
|
||||
// CHECK: hipdnnHandle_t handle;
|
||||
cudnnHandle_t handle;
|
||||
// CHECK: hipdnnCreate(&handle);
|
||||
cudnnCreate(&handle);
|
||||
|
||||
float one = 1;
|
||||
float zero = 0;
|
||||
|
||||
// softmaxForward(n, c, h, w, dstData, &srcData);
|
||||
// CHECK: hipdnnTensorDescriptor_t srcTensorDesc, sftTensorDesc;
|
||||
// CHECK: hipdnnCreateTensorDescriptor(&srcTensorDesc);
|
||||
// CHECK: hipdnnCreateTensorDescriptor(&sftTensorDesc);
|
||||
cudnnTensorDescriptor_t srcTensorDesc, sftTensorDesc;
|
||||
cudnnCreateTensorDescriptor(&srcTensorDesc);
|
||||
cudnnCreateTensorDescriptor(&sftTensorDesc);
|
||||
// CHECK: hipdnnSetTensor4dDescriptor(srcTensorDesc, HIPDNN_TENSOR_NCHW, HIPDNN_DATA_DOUBLE,
|
||||
cudnnSetTensor4dDescriptor(srcTensorDesc, CUDNN_TENSOR_NCHW, CUDNN_DATA_DOUBLE,
|
||||
m, c, 1, 1);
|
||||
// CHECK: hipdnnSetTensor4dDescriptor(sftTensorDesc, HIPDNN_TENSOR_NCHW, HIPDNN_DATA_DOUBLE,
|
||||
cudnnSetTensor4dDescriptor(sftTensorDesc, CUDNN_TENSOR_NCHW, CUDNN_DATA_DOUBLE,
|
||||
m, c, 1, 1);
|
||||
// CHECK: hipdnnSoftmaxForward(handle, HIPDNN_SOFTMAX_ACCURATE, HIPDNN_SOFTMAX_MODE_CHANNEL, &one,
|
||||
cudnnSoftmaxForward(handle, CUDNN_SOFTMAX_ACCURATE, CUDNN_SOFTMAX_MODE_CHANNEL, &one,
|
||||
srcTensorDesc, d_fcLayer, &zero, sftTensorDesc, d_softmaxData);
|
||||
// CHECK: hipDeviceSynchronize();
|
||||
cudaDeviceSynchronize();
|
||||
|
||||
// Copy back
|
||||
double *result = (double *) malloc(m * c * sizeof(double));
|
||||
// CHECK: hipMemcpy(result, d_softmaxData, m * c * sizeof(double), hipMemcpyDeviceToHost);
|
||||
// CHECK: hipDeviceSynchronize();
|
||||
cudaMemcpy(result, d_softmaxData, m * c * sizeof(double), cudaMemcpyDeviceToHost);
|
||||
cudaDeviceSynchronize();
|
||||
|
||||
// Log
|
||||
printf("SOFTMAX:\n");
|
||||
printMatrix(result, c, m);
|
||||
|
||||
// Try backward
|
||||
// CHECK: hipdnnTensorDescriptor_t diffTensorDesc;
|
||||
// CHECK: hipdnnCreateTensorDescriptor(&diffTensorDesc);
|
||||
// CHECK: hipdnnSetTensor4dDescriptor(diffTensorDesc, HIPDNN_TENSOR_NCHW, HIPDNN_DATA_DOUBLE,
|
||||
cudnnTensorDescriptor_t diffTensorDesc;
|
||||
cudnnCreateTensorDescriptor(&diffTensorDesc);
|
||||
cudnnSetTensor4dDescriptor(diffTensorDesc, CUDNN_TENSOR_NCHW, CUDNN_DATA_DOUBLE,
|
||||
m, c, 1, 1);
|
||||
|
||||
double *d_gradData;
|
||||
// CHECK: hipMalloc((void**) &d_gradData, m * c * sizeof(double));
|
||||
cudaMalloc((void**) &d_gradData, m * c * sizeof(double));
|
||||
|
||||
double *diffData = makeDiffData(m, c);
|
||||
double *d_diffData;
|
||||
// CHECK: hipMalloc((void**) &d_diffData, m * c * sizeof(double));
|
||||
// CHECK: hipMemcpy(d_diffData, diffData, m * c * sizeof(double), hipMemcpyHostToDevice);
|
||||
// CHECK: hipDeviceSynchronize();
|
||||
cudaMalloc((void**) &d_diffData, m * c * sizeof(double));
|
||||
cudaMemcpy(d_diffData, diffData, m * c * sizeof(double), cudaMemcpyHostToDevice);
|
||||
cudaDeviceSynchronize();
|
||||
// CHECK: hipdnnSoftmaxBackward(handle, HIPDNN_SOFTMAX_ACCURATE, HIPDNN_SOFTMAX_MODE_CHANNEL,
|
||||
cudnnSoftmaxBackward(handle, CUDNN_SOFTMAX_ACCURATE, CUDNN_SOFTMAX_MODE_CHANNEL,
|
||||
&one, srcTensorDesc, d_softmaxData, diffTensorDesc, d_diffData, &zero, sftTensorDesc, d_gradData);
|
||||
// CHECK: hipDeviceSynchronize();
|
||||
cudaDeviceSynchronize();
|
||||
|
||||
// Copy back
|
||||
double *result_backward = (double *) malloc(m * c * sizeof(double));
|
||||
// CHECK: hipMemcpy(result_backward, d_gradData, m * c * sizeof(double), hipMemcpyDeviceToHost);
|
||||
// CHECK: hipDeviceSynchronize();
|
||||
cudaMemcpy(result_backward, d_gradData, m * c * sizeof(double), cudaMemcpyDeviceToHost);
|
||||
cudaDeviceSynchronize();
|
||||
|
||||
// Log
|
||||
printf("GRADIENT:\n");
|
||||
printMatrix(result_backward, c, m);
|
||||
|
||||
// Destruct
|
||||
free(result);
|
||||
free(diffData);
|
||||
free(result_backward);
|
||||
free(fcLayer);
|
||||
|
||||
// CHECK: hipdnnDestroyTensorDescriptor(srcTensorDesc);
|
||||
// CHECK: hipdnnDestroyTensorDescriptor(sftTensorDesc);
|
||||
// CHECK: hipdnnDestroyTensorDescriptor(diffTensorDesc);
|
||||
// CHECK: hipFree(d_fcLayer);
|
||||
// CHECK: hipFree(d_softmaxData);
|
||||
// CHECK: hipFree(d_gradData);
|
||||
// CHECK: hipFree(d_diffData);
|
||||
// CHECK: hipdnnDestroy(handle);
|
||||
cudnnDestroyTensorDescriptor(srcTensorDesc);
|
||||
cudnnDestroyTensorDescriptor(sftTensorDesc);
|
||||
cudnnDestroyTensorDescriptor(diffTensorDesc);
|
||||
cudaFree(d_fcLayer);
|
||||
cudaFree(d_softmaxData);
|
||||
cudaFree(d_gradData);
|
||||
cudaFree(d_diffData);
|
||||
cudnnDestroy(handle);
|
||||
}
|
||||
@@ -23,66 +23,913 @@ THE SOFTWARE.
|
||||
* HIT_END
|
||||
*/
|
||||
|
||||
#include <cstdint>
|
||||
#include "hip/hip_runtime.h"
|
||||
#include "test_common.h"
|
||||
#include "hip/hip_runtime_api.h"
|
||||
#include <iostream>
|
||||
|
||||
__global__ void vAdd(hipLaunchParm lp, float* a) {}
|
||||
// Memory alignment is broken
|
||||
// Update: with latest changes the aligment is working fine, hence enabled
|
||||
#define ENABLE_ALIGNMENT_TEST_SMALL_BAR 1
|
||||
|
||||
// Packed member atribute broken
|
||||
#define ENABLE_PACKED_TEST 0
|
||||
|
||||
// Update: with latest changes struct class object
|
||||
// from device is working fine, hence enabled
|
||||
#define ENABLE_CLASS_OBJ_ACCESS 1
|
||||
|
||||
// accessing dynamic/heap memory from device is broken
|
||||
#define ENABLE_HEAP_MEMORY_ACCESS 0
|
||||
|
||||
// Update: with latest changes it's working hence enabled
|
||||
#define ENABLE_USER_STL 1
|
||||
|
||||
// Update: with latest changes it's working hence enabled
|
||||
#define ENABLE_OUT_OF_ORDER_INITIALIZATION 1
|
||||
|
||||
// Direct initialization of struct broken,
|
||||
// ip_d9 is a pointer, uint_t*, hipLaunchKernelStruct_h9 = {'c', ip_d9};
|
||||
#define ENABLE_DECLARE_INITIALIZATION_POINTER 0
|
||||
|
||||
// Bit fields are broken
|
||||
#define ENABLE_BIT_FIELDS 0
|
||||
|
||||
static const int BLOCK_DIM_SIZE = 1024;
|
||||
|
||||
// allocate memory on device and host for result validation
|
||||
static bool *result_d, *result_h;
|
||||
static hipError_t hipMallocError = hipMalloc((void**)&result_d,
|
||||
BLOCK_DIM_SIZE*sizeof(bool));
|
||||
static hipError_t hipHostMallocError = hipHostMalloc((void**)&result_h,
|
||||
BLOCK_DIM_SIZE*sizeof(bool));
|
||||
static hipError_t hipMemsetError = hipMemset(result_d,
|
||||
false, BLOCK_DIM_SIZE);
|
||||
|
||||
static void ResultValidation() {
|
||||
hipMemcpy(result_h, result_d, BLOCK_DIM_SIZE*sizeof(bool),
|
||||
hipMemcpyDeviceToHost);
|
||||
|
||||
for (int k = 0; k < BLOCK_DIM_SIZE; ++k) {
|
||||
HIPASSERT(result_h[k] == true);
|
||||
}
|
||||
return;
|
||||
}
|
||||
|
||||
// Segregating the reset part as it was causing a problem when i put inside
|
||||
// ResultValidation() function, the memory was not reset correctly for the
|
||||
// tests which were disabled.
|
||||
static void ResetValidationMem() {
|
||||
// reset the memory to false to reuse it.
|
||||
hipMemset(result_d, false, BLOCK_DIM_SIZE);
|
||||
hipMemset(result_h, false, BLOCK_DIM_SIZE);
|
||||
return;
|
||||
}
|
||||
|
||||
// This test is to verify Struct with variables
|
||||
// support, read from device.
|
||||
typedef struct hipLaunchKernelStruct1 {
|
||||
int li; // local int
|
||||
float lf; // local float
|
||||
bool result; // local bool
|
||||
} hipLaunchKernelStruct_t1;
|
||||
|
||||
// This test is to verify struct with padding, read from device
|
||||
typedef struct hipLaunchKernelStruct2 {
|
||||
char c1;
|
||||
long l1;
|
||||
char c2;
|
||||
long l2;
|
||||
bool result;
|
||||
} hipLaunchKernelStruct_t2;
|
||||
|
||||
// This test is to verify struct with padding, read from device
|
||||
typedef struct hipLaunchKernelStruct3 {
|
||||
char bf1;
|
||||
char bf2;
|
||||
long l1;
|
||||
char bf3;
|
||||
bool result;
|
||||
} hipLaunchKernelStruct_t3;
|
||||
|
||||
// This test is to verify empty struct
|
||||
typedef struct hipLaunchKernelStruct4 {
|
||||
// empty struct, size will be verified from device side,size 1Byte
|
||||
} hipLaunchKernelStruct_t4;
|
||||
|
||||
// This test is to verify struct with pointer member variable.
|
||||
typedef struct hipLaunchKernelStruct5 {
|
||||
char c1;
|
||||
char* cp; // char pointer
|
||||
} hipLaunchKernelStruct_t5;
|
||||
|
||||
|
||||
// This test is to verify struct with aligned(8),
|
||||
// right now it's broken on hcc & hip-clang
|
||||
typedef struct hipLaunchKernelStruct6 {
|
||||
char c1;
|
||||
short int si;
|
||||
} __attribute__((aligned(8))) hipLaunchKernelStruct_t6;
|
||||
|
||||
// This test is to verify struct with aligned(16),
|
||||
// right now it's brokenon hcc & hip-clang
|
||||
typedef struct hipLaunchKernelStruct7 {
|
||||
char c1;
|
||||
short int si;
|
||||
} __attribute__((aligned(16))) hipLaunchKernelStruct_t7;
|
||||
|
||||
// This test is to verify struct with packed & aligned,
|
||||
// size should be 4Bytes right now it's broken on hcc & hip-clang
|
||||
typedef struct hipLaunchKernelStruct8 {
|
||||
char c1;
|
||||
short int si;
|
||||
bool b;
|
||||
}__attribute__((packed, aligned(4))) hipLaunchKernelStruct_t8;
|
||||
|
||||
// This test is to verify struct with packed, no alignment as Sam suggested
|
||||
// size should be 4Bytes, right now it's broken on hcc & hip-clang
|
||||
typedef struct hipLaunchKernelStruct8A {
|
||||
char c1;
|
||||
short int si;
|
||||
bool b;
|
||||
}__attribute__((packed)) hipLaunchKernelStruct_t8A;
|
||||
|
||||
// This test is to verify struct with alignment, no packing as Sam suggested
|
||||
// size should be 8Bytes as no packing, right now it's broken on hcc & hip-clang
|
||||
typedef struct hipLaunchKernelStruct8B {
|
||||
char c1;
|
||||
short int si;
|
||||
bool b;
|
||||
}__attribute__((aligned(8))) hipLaunchKernelStruct_t8B;
|
||||
|
||||
// This test is to verify const struct object
|
||||
typedef struct hipLaunchKernelStruct9 {
|
||||
char c1;
|
||||
uint32_t* ip; // uint pointer
|
||||
} hipLaunchKernelStruct_t9;
|
||||
|
||||
// This test is to verify struct with stdint types, uintN_t
|
||||
typedef struct hipLaunchKernelStruct10 {
|
||||
uint64_t u64;
|
||||
uint32_t u32;
|
||||
uint8_t u8;
|
||||
} hipLaunchKernelStruct_t10;
|
||||
|
||||
// This test is to verify struct with volatile member
|
||||
typedef struct hipLaunchKernelStruct11 {
|
||||
int i1;
|
||||
volatile unsigned int vint;
|
||||
} hipLaunchKernelStruct_t11;
|
||||
|
||||
// This test is to verify struct with simple class object
|
||||
class base {
|
||||
public:
|
||||
int i = 0;
|
||||
base() {}
|
||||
};
|
||||
typedef struct hipLaunchKernelStruct12 {
|
||||
base b;
|
||||
char c1;
|
||||
} hipLaunchKernelStruct_t12;
|
||||
|
||||
// This test is to verify struct with __device__ func() attribute
|
||||
typedef struct hipLaunchKernelStruct13 {
|
||||
int i1;
|
||||
__device__ int getvalue() { return i1; }
|
||||
} hipLaunchKernelStruct_t13;
|
||||
|
||||
// This test is to verify struct with array variable,
|
||||
// write to from device
|
||||
typedef struct hipLaunchKernelStruct14 {
|
||||
int readint;
|
||||
int writeint[BLOCK_DIM_SIZE]; // will write to this from device
|
||||
} hipLaunchKernelStruct_t14;
|
||||
|
||||
// This test is to verify struct with dynamic memory, new int
|
||||
// the heap memory will be accessed from device
|
||||
typedef struct hipLaunchKernelStruct15 {
|
||||
char c1;
|
||||
int* heapmem; // allocated using hipMalloc()
|
||||
} hipLaunchKernelStruct_t15;
|
||||
|
||||
// This test is to verify simple template struct
|
||||
template<typename T>
|
||||
struct hipLaunchKernelStruct_t16 {
|
||||
T t1;
|
||||
};
|
||||
|
||||
// This test is to verify simple explicity template struct
|
||||
template<typename T> struct hipLaunchKernelStruct_t17 {};
|
||||
template<> // explicit template
|
||||
struct hipLaunchKernelStruct_t17<int> {
|
||||
int t1;
|
||||
};
|
||||
|
||||
// This test is to verity write to struct memory using __device__ func()
|
||||
typedef struct hipLaunchKernelStruct18 {
|
||||
char c1;
|
||||
__device__ void setChar(char c) { c1 = c; }
|
||||
__device__ int getChar() { return c1; }
|
||||
} hipLaunchKernelStruct_t18;
|
||||
|
||||
// This test is to verity user defined STL, simple stack implementation
|
||||
typedef struct stackNode {
|
||||
int data;
|
||||
stackNode* nextNode = NULL;
|
||||
} stackNode_t;
|
||||
typedef struct hipLaunchKernelStruct19 {
|
||||
stackNode_t* stack = NULL;
|
||||
unsigned int size_ = 0;
|
||||
void pushMe(int value) { // not a device function, setting from host
|
||||
stackNode_t* newNode;
|
||||
hipMalloc((void**)&newNode, sizeof(stackNode_t));
|
||||
hipMemset(&newNode->data, value, sizeof(stackNode_t));
|
||||
//newNode->data = value;
|
||||
++size_;
|
||||
if (stack == NULL) {
|
||||
stack = newNode;
|
||||
return;
|
||||
}
|
||||
stackNode_t* currentHead = stack;
|
||||
stack = newNode;
|
||||
stack->nextNode = currentHead;
|
||||
return;
|
||||
}
|
||||
__device__ void popMe() {
|
||||
stackNode_t* currentHead = stack;
|
||||
stack = stack->nextNode;
|
||||
--size_;
|
||||
// delete currentHead; // no idea why delete not working
|
||||
return;
|
||||
}
|
||||
int stackSize() {
|
||||
return size_;
|
||||
}
|
||||
} hipLaunchKernelStruct_t19;
|
||||
|
||||
// This test is to verify out of order initalizer of struct elements
|
||||
// and access in-order, from device.
|
||||
typedef struct hipLaunchKernelStruct20 {
|
||||
char name;
|
||||
int age;
|
||||
int rank;
|
||||
} hipLaunchKernelStruct_t20;
|
||||
|
||||
// This test is to verify bit fields operations
|
||||
// the size should be 1Bytes
|
||||
typedef struct hipLaunchKernelStruct21 {
|
||||
int i : 3; // limiting bits to 3
|
||||
int j : 2; // limiting bits to 2
|
||||
} hipLaunchKernelStruct_t21;
|
||||
|
||||
// Passing struct to a hipLaunchKernelGGL(),
|
||||
// read and write into the same struct
|
||||
__global__ void hipLaunchKernelStructFunc1(
|
||||
hipLaunchKernelStruct_t1 hipLaunchKernelStruct_,
|
||||
bool* result_d1) {
|
||||
int x = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
|
||||
// set the result to true if the condition met
|
||||
result_d1[x] = ((hipLaunchKernelStruct_.li == 1)
|
||||
&& (hipLaunchKernelStruct_.lf == 1.0)
|
||||
&& (hipLaunchKernelStruct_.result == false));
|
||||
}
|
||||
|
||||
// Passing struct to a hipLaunchKernelGGL(), checks padding,
|
||||
// read and write into the same struct
|
||||
__global__ void hipLaunchKernelStructFunc2(
|
||||
hipLaunchKernelStruct_t2 hipLaunchKernelStruct_,
|
||||
bool* result_d2) {
|
||||
int x = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
|
||||
// set the result to true if the condition met
|
||||
result_d2[x] = ((hipLaunchKernelStruct_.c1 == 'a')
|
||||
&& (hipLaunchKernelStruct_.l1 == 1.0)
|
||||
&& (hipLaunchKernelStruct_.c2 == 'b')
|
||||
&& (hipLaunchKernelStruct_.l2 == 2.0) );
|
||||
}
|
||||
|
||||
// Passing struct to a hipLaunchKernelGGL(), checks padding,
|
||||
// read and write into the same struct
|
||||
__global__ void hipLaunchKernelStructFunc3(
|
||||
hipLaunchKernelStruct_t3 hipLaunchKernelStruct_,
|
||||
bool* result_d3) {
|
||||
int x = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
|
||||
// set the result to true if the condition met
|
||||
result_d3[x] = ((hipLaunchKernelStruct_.bf1 == 1)
|
||||
&& (hipLaunchKernelStruct_.bf2 == 1)
|
||||
&& (hipLaunchKernelStruct_.l1 == 1.0)
|
||||
&& (hipLaunchKernelStruct_.bf3 == 1) );
|
||||
}
|
||||
|
||||
// Passing empty struct to a hipLaunchKernelGGL(),
|
||||
// check the size of 1Byte, set result_d4 to true if condition met
|
||||
__global__ void hipLaunchKernelStructFunc4(
|
||||
hipLaunchKernelStruct_t4 hipLaunchKernelStruct_,
|
||||
bool* result_d4) {
|
||||
int x = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
|
||||
// set the result to true if the condition met
|
||||
result_d4[x] = (sizeof(hipLaunchKernelStruct_) == 1);
|
||||
}
|
||||
|
||||
// Passing struct with pointer object to a hipLaunchKernelGGL()
|
||||
__global__ void hipLaunchKernelStructFunc5(
|
||||
hipLaunchKernelStruct_t5 hipLaunchKernelStruct_,
|
||||
bool* result_d5) {
|
||||
int x = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
|
||||
// set the result to true if the condition met
|
||||
result_d5[x] = ((hipLaunchKernelStruct_.c1 == 'c')
|
||||
&& (*hipLaunchKernelStruct_.cp == 'p'));
|
||||
}
|
||||
|
||||
// Passing struct which is aligned to 8Byte to a hipLaunchKernelGGL(),
|
||||
// set the result_d6 to true if condition met
|
||||
__global__ void hipLaunchKernelStructFunc6(
|
||||
hipLaunchKernelStruct_t6 hipLaunchKernelStruct_,
|
||||
bool* result_d6) {
|
||||
int x = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
|
||||
// set the result to true if the condition met
|
||||
// get the address of the struct
|
||||
// size_t(p)%8 will be 0 if aligned to 8Byte address space
|
||||
int *p = (int*)(&hipLaunchKernelStruct_);
|
||||
result_d6[x] = ((hipLaunchKernelStruct_.c1 == 'c')
|
||||
&& (hipLaunchKernelStruct_.si == 1)
|
||||
&& ((size_t(p))%8 ==0));
|
||||
}
|
||||
|
||||
// Passing struct which is aligned to 16Byte,
|
||||
// set the result_d7 to true if condition met
|
||||
__global__ void hipLaunchKernelStructFunc7(
|
||||
hipLaunchKernelStruct_t7 hipLaunchKernelStruct_,
|
||||
bool* result_d7) {
|
||||
int x = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
|
||||
// set the result to true if the condition met
|
||||
// get the address of the struct
|
||||
// size_t(p)%16 will be 0 if aligned to 16Byte address space
|
||||
int *p = (int*)(&hipLaunchKernelStruct_);
|
||||
result_d7[x] = ((hipLaunchKernelStruct_.c1 == 'c')
|
||||
&& (hipLaunchKernelStruct_.si == 1)
|
||||
&& ((size_t(p))%16 ==0) );
|
||||
}
|
||||
|
||||
// Passing struct which is packed & aligned to 4Byte,
|
||||
// set the result_d8 to true if condition met
|
||||
__global__ void hipLaunchKernelStructFunc8(
|
||||
hipLaunchKernelStruct_t8 hipLaunchKernelStruct_,
|
||||
bool* result_d8) {
|
||||
int x = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
|
||||
// set the result to true if the condition met
|
||||
// get the address of the xth element, struct[x],
|
||||
// size_t(p)%4 will be 0 if aligned to 4Byte address space
|
||||
int *p = (int*)(&hipLaunchKernelStruct_);
|
||||
result_d8[x] = ((hipLaunchKernelStruct_.c1 == 'c')
|
||||
&& (hipLaunchKernelStruct_.si == 1)
|
||||
&& ((size_t(p))%4 ==0)
|
||||
&& (sizeof(hipLaunchKernelStruct_) == 4));
|
||||
}
|
||||
|
||||
// Passing struct which is packed only, as Sam suggested, should be 4Bytes
|
||||
// set the result_d8A to true if condition met
|
||||
__global__ void hipLaunchKernelStructFunc8A(
|
||||
hipLaunchKernelStruct_t8A hipLaunchKernelStruct_,
|
||||
bool* result_d8A) {
|
||||
int x = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
|
||||
// set the result to true if the condition met
|
||||
// this is packed struct
|
||||
// the address will not be aglined in this case hence condition removed
|
||||
// only sizeof(hipLaunchKernelStruct_) will be valided
|
||||
result_d8A[x] = ((hipLaunchKernelStruct_.c1 == 'c')
|
||||
&& (hipLaunchKernelStruct_.si == 1)
|
||||
&& (sizeof(hipLaunchKernelStruct_) == 4));
|
||||
}
|
||||
|
||||
// Passing struct which is aligned(4) only, as Sam suggested
|
||||
// , size should be 8Bytes, set the result_d8B to true if condition met
|
||||
__global__ void hipLaunchKernelStructFunc8B(
|
||||
hipLaunchKernelStruct_t8B hipLaunchKernelStruct_,
|
||||
bool* result_d8B) {
|
||||
int x = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
|
||||
// set the result to true if the condition met
|
||||
// get the address of the xth element, struct[x],
|
||||
// size_t(p)%4 will be 0 if aligned to 4Byte address space
|
||||
int *p = (int*)(&hipLaunchKernelStruct_);
|
||||
result_d8B[x] = ((hipLaunchKernelStruct_.c1 == 'c')
|
||||
&& (hipLaunchKernelStruct_.si == 1)
|
||||
&& ((size_t(p))%8 == 0)
|
||||
&& (sizeof(hipLaunchKernelStruct_) == 8));
|
||||
}
|
||||
|
||||
// Passing struct with uint pointer object to a hipLaunchKernelGGL()
|
||||
__global__ void hipLaunchKernelStructFunc9(
|
||||
const hipLaunchKernelStruct_t9 hipLaunchKernelStruct_,
|
||||
bool* result_d9) {
|
||||
int x = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
|
||||
// set the result to true if the condition met
|
||||
result_d9[x] = ((hipLaunchKernelStruct_.c1 == 'c')
|
||||
&& (*hipLaunchKernelStruct_.ip == 1));
|
||||
}
|
||||
|
||||
// Passing struct with stdint types object, uintN_t, to a hipLaunchKernelGGL()
|
||||
__global__ void hipLaunchKernelStructFunc10(
|
||||
hipLaunchKernelStruct_t10 hipLaunchKernelStruct_,
|
||||
bool* result_d10) {
|
||||
int x = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
|
||||
// set the result to true if the condition met
|
||||
result_d10[x] = ((hipLaunchKernelStruct_.u64 == UINT64_MAX)
|
||||
&& (hipLaunchKernelStruct_.u32 == 1)
|
||||
&& (hipLaunchKernelStruct_.u8 == UINT8_MAX));
|
||||
}
|
||||
|
||||
// Passing struct with volatile member, to a hipLaunchKernelGGL()
|
||||
__global__ void hipLaunchKernelStructFunc11(
|
||||
hipLaunchKernelStruct_t11 hipLaunchKernelStruct_,
|
||||
bool* result_d11) {
|
||||
int x = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
|
||||
// set the result to true if the condition met
|
||||
result_d11[x] = ((hipLaunchKernelStruct_.i1 == 1)
|
||||
&& (hipLaunchKernelStruct_.vint == 0));
|
||||
}
|
||||
|
||||
// Passing struct with simple class obj, to a hipLaunchKernelGGL()
|
||||
__global__ void hipLaunchKernelStructFunc12(
|
||||
hipLaunchKernelStruct_t12 hipLaunchKernelStruct_,
|
||||
bool* result_d12) {
|
||||
int x = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
|
||||
// set the result to true if the condition met
|
||||
result_d12[x] = ((hipLaunchKernelStruct_.c1 == 'c')
|
||||
&& (hipLaunchKernelStruct_.b.i == 0));
|
||||
}
|
||||
|
||||
// Passing struct with simple __device__ func(), to a hipLaunchKernelGGL()
|
||||
__global__ void hipLaunchKernelStructFunc13(
|
||||
hipLaunchKernelStruct_t13 hipLaunchKernelStruct_,
|
||||
bool* result_d13) {
|
||||
int x = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
|
||||
// set the result to true if the condition met
|
||||
result_d13[x] = ((hipLaunchKernelStruct_.i1 == 1)
|
||||
&& (hipLaunchKernelStruct_.getvalue() == 1));
|
||||
}
|
||||
|
||||
// Passing struct with array variable, write to from device
|
||||
__global__ void hipLaunchKernelStructFunc14(
|
||||
hipLaunchKernelStruct_t14 hipLaunchKernelStruct_,
|
||||
bool* result_d14) {
|
||||
int x = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
|
||||
hipLaunchKernelStruct_.writeint[x] = 1;
|
||||
// set the result to true if the condition met
|
||||
result_d14[x] = ((hipLaunchKernelStruct_.readint == 1)
|
||||
&& (hipLaunchKernelStruct_.writeint[x] == 1));
|
||||
}
|
||||
|
||||
// Passing struct with struct with dynamic memory, new int
|
||||
// the heap memory will be accessed from device
|
||||
__global__ void hipLaunchKernelStructFunc15(
|
||||
hipLaunchKernelStruct_t15 hipLaunchKernelStruct_,
|
||||
bool* result_d15) {
|
||||
int x = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
|
||||
// set the result to true if the condition met
|
||||
result_d15[x] = ((hipLaunchKernelStruct_.c1 == 'c')
|
||||
&& (hipLaunchKernelStruct_.heapmem[x] == 1));
|
||||
}
|
||||
|
||||
// Passing simple template struct
|
||||
__global__ void hipLaunchKernelStructFunc16(
|
||||
hipLaunchKernelStruct_t16<char> hipLaunchKernelStruct_,
|
||||
bool* result_d16) {
|
||||
int x = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
|
||||
// set the result to true if the condition met
|
||||
result_d16[x] = (hipLaunchKernelStruct_.t1 == 'c');
|
||||
}
|
||||
|
||||
// Passing simple explicit template struct
|
||||
__global__ void hipLaunchKernelStructFunc17(
|
||||
hipLaunchKernelStruct_t17<int> hipLaunchKernelStruct_,
|
||||
bool* result_d17) {
|
||||
int x = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
|
||||
// set the result to true if the condition met
|
||||
result_d17[x] = (hipLaunchKernelStruct_.t1 == 1);
|
||||
}
|
||||
|
||||
// Passing struct and write to struct memory using __device__ func()
|
||||
__global__ void hipLaunchKernelStructFunc18(
|
||||
hipLaunchKernelStruct_t18 hipLaunchKernelStruct_,
|
||||
bool* result_d18) {
|
||||
int x = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
|
||||
hipLaunchKernelStruct_.setChar('c');
|
||||
// set the result to true if the condition met
|
||||
result_d18[x] = (hipLaunchKernelStruct_.getChar() == 'c');
|
||||
}
|
||||
|
||||
// Passing simple user defined stack implemenration, using __device__ func()
|
||||
__global__ void hipLaunchKernelStructFunc19(
|
||||
hipLaunchKernelStruct_t19 hipLaunchKernelStruct_) {
|
||||
int x = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
|
||||
// stack should be empty after the kernel execustion, verify on host side
|
||||
hipLaunchKernelStruct_.popMe();
|
||||
}
|
||||
|
||||
// Passing out of order initalized struct, access in-order
|
||||
__global__ void hipLaunchKernelStructFunc20(
|
||||
hipLaunchKernelStruct_t20 hipLaunchKernelStruct_,
|
||||
bool* result_d20) {
|
||||
int x = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
|
||||
// accessing struct members in order
|
||||
result_d20[x] = (hipLaunchKernelStruct_.name == 'A'
|
||||
// strcmp(hipLaunchKernelStruct_.name, "AMD") -> strcmp is not broken
|
||||
&& hipLaunchKernelStruct_.age == 42
|
||||
&& hipLaunchKernelStruct_.rank == 2);
|
||||
}
|
||||
|
||||
// Passing struct with bit fields
|
||||
__global__ void hipLaunchKernelStructFunc21(
|
||||
hipLaunchKernelStruct_t21 hipLaunchKernelStruct_,
|
||||
bool* result_d21) {
|
||||
int x = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
|
||||
// accessing struct members in order
|
||||
result_d21[x] = (hipLaunchKernelStruct_.i == 2
|
||||
&& hipLaunchKernelStruct_.j == 0
|
||||
&& (sizeof(hipLaunchKernelStruct_) == 1));
|
||||
}
|
||||
|
||||
__global__ void vAdd(float* a) {}
|
||||
|
||||
//---
|
||||
// Some wrapper macro for testing:
|
||||
#define WRAP(...) __VA_ARGS__
|
||||
|
||||
#include <sys/time.h>
|
||||
#define GPU_PRINT_TIME(cmd, elapsed, quiet) \
|
||||
do { \
|
||||
struct timeval start, stop; \
|
||||
float elapsed; \
|
||||
gettimeofday(&start, NULL); \
|
||||
hipDeviceSynchronize(); \
|
||||
cmd; \
|
||||
hipDeviceSynchronize(); \
|
||||
gettimeofday(&stop, NULL); \
|
||||
#define GPU_PRINT_TIME(cmd, elapsed, quiet) \
|
||||
do { \
|
||||
struct timeval start, stop; \
|
||||
float elapsed; \
|
||||
gettimeofday(&start, NULL); \
|
||||
hipDeviceSynchronize(); \
|
||||
cmd; \
|
||||
hipDeviceSynchronize(); \
|
||||
gettimeofday(&stop, NULL); \
|
||||
} while (0);
|
||||
|
||||
|
||||
#define MY_LAUNCH(command, doTrace, msg) \
|
||||
{ \
|
||||
if (doTrace) printf("TRACE: %s %s\n", msg, #command); \
|
||||
command; \
|
||||
#define MY_LAUNCH(command, doTrace, msg) \
|
||||
{ \
|
||||
if (doTrace) printf("TRACE: %s %s\n", msg, #command); \
|
||||
command; \
|
||||
}
|
||||
|
||||
|
||||
#define MY_LAUNCH_WITH_PAREN(command, doTrace, msg) \
|
||||
{ \
|
||||
if (doTrace) printf("TRACE: %s %s\n", msg, #command); \
|
||||
(command); \
|
||||
#define MY_LAUNCH_WITH_PAREN(command, doTrace, msg) \
|
||||
{ \
|
||||
if (doTrace) printf("TRACE: %s %s\n", msg, #command); \
|
||||
(command); \
|
||||
}
|
||||
|
||||
|
||||
int main() {
|
||||
// Validating memory & initial value, for result_d, result_h
|
||||
HIPASSERT(hipMallocError == hipSuccess);
|
||||
HIPASSERT(hipHostMallocError == hipSuccess);
|
||||
HIPASSERT(hipMemsetError == hipSuccess);
|
||||
|
||||
// Test: Passing Struct type, check access from device.
|
||||
ResetValidationMem();
|
||||
hipLaunchKernelStruct_t1 hipLaunchKernelStruct_h1;
|
||||
hipLaunchKernelStruct_h1.li = 1;
|
||||
hipLaunchKernelStruct_h1.lf = 1.0;
|
||||
hipLaunchKernelStruct_h1.result = false;
|
||||
hipLaunchKernelGGL(HIP_KERNEL_NAME(hipLaunchKernelStructFunc1),
|
||||
dim3(BLOCK_DIM_SIZE),
|
||||
dim3(1), 0, 0, hipLaunchKernelStruct_h1,
|
||||
result_d);
|
||||
ResultValidation();
|
||||
|
||||
// Test: Passing Struct type, checks padding
|
||||
ResetValidationMem();
|
||||
hipLaunchKernelStruct_t2 hipLaunchKernelStruct_h2;
|
||||
hipLaunchKernelStruct_h2.c1 = 'a';
|
||||
hipLaunchKernelStruct_h2.l1 = 1.0;
|
||||
hipLaunchKernelStruct_h2.c2 = 'b';
|
||||
hipLaunchKernelStruct_h2.l2 = 2.0;
|
||||
hipLaunchKernelStruct_h2.result = false;
|
||||
hipLaunchKernelGGL(HIP_KERNEL_NAME(hipLaunchKernelStructFunc2),
|
||||
dim3(BLOCK_DIM_SIZE),
|
||||
dim3(1), 0, 0, hipLaunchKernelStruct_h2,
|
||||
result_d);
|
||||
ResultValidation();
|
||||
|
||||
// Test: Passing Struct type, checks padding, assigning integer to a char
|
||||
ResetValidationMem();
|
||||
hipLaunchKernelStruct_t3 hipLaunchKernelStruct_h3;
|
||||
hipLaunchKernelStruct_h3.bf1 = 1;
|
||||
hipLaunchKernelStruct_h3.bf2 = 1;
|
||||
hipLaunchKernelStruct_h3.l1 = 1.0;
|
||||
hipLaunchKernelStruct_h3.bf3 = 1;
|
||||
hipLaunchKernelStruct_h3.result = false;
|
||||
// initialize to false, will be set to
|
||||
// true if the struct size is 1Byte, from device size
|
||||
hipLaunchKernelGGL(HIP_KERNEL_NAME(hipLaunchKernelStructFunc3),
|
||||
dim3(BLOCK_DIM_SIZE),
|
||||
dim3(1), 0, 0, hipLaunchKernelStruct_h3,
|
||||
result_d);
|
||||
ResultValidation();
|
||||
|
||||
// Test: Passing empty struct
|
||||
ResetValidationMem();
|
||||
hipLaunchKernelStruct_t4 hipLaunchKernelStruct_h4;
|
||||
hipLaunchKernelGGL(HIP_KERNEL_NAME(hipLaunchKernelStructFunc4),
|
||||
dim3(BLOCK_DIM_SIZE),
|
||||
dim3(1), 0, 0, hipLaunchKernelStruct_h4,
|
||||
result_d);
|
||||
ResultValidation();
|
||||
|
||||
// Test: Passing struct with pointer object to a hipLaunchKernelGGL()
|
||||
ResetValidationMem();
|
||||
hipLaunchKernelStruct_t5 hipLaunchKernelStruct_h5;
|
||||
char* cp_d5; // This is passed as pointer to struct member
|
||||
// allocating memory for char pointer on device
|
||||
HIPCHECK(hipMalloc((void**)&cp_d5, sizeof(char)));
|
||||
HIPCHECK(hipMemset(cp_d5, 'p', sizeof(char)));
|
||||
hipLaunchKernelStruct_h5.c1 = 'c';
|
||||
hipLaunchKernelStruct_h5.cp = cp_d5;
|
||||
hipLaunchKernelGGL(HIP_KERNEL_NAME(hipLaunchKernelStructFunc5),
|
||||
dim3(BLOCK_DIM_SIZE),
|
||||
dim3(1), 0, 0, hipLaunchKernelStruct_h5,
|
||||
result_d);
|
||||
ResultValidation();
|
||||
|
||||
// Test: Passing struct with aligned(8)
|
||||
ResetValidationMem();
|
||||
hipLaunchKernelStruct_t6 hipLaunchKernelStruct_h6;
|
||||
hipLaunchKernelStruct_h6.c1 = 'c';
|
||||
hipLaunchKernelStruct_h6.si = 1;
|
||||
hipLaunchKernelGGL(HIP_KERNEL_NAME(hipLaunchKernelStructFunc6),
|
||||
dim3(BLOCK_DIM_SIZE),
|
||||
dim3(1), 0, 0, hipLaunchKernelStruct_h6,
|
||||
result_d);
|
||||
// alignment is broken hence disabled the validation part
|
||||
#if ENABLE_ALIGNMENT_TEST_SMALL_BAR
|
||||
ResultValidation();
|
||||
#endif
|
||||
|
||||
|
||||
// Test: Passing struct with aligned(16)
|
||||
ResetValidationMem();
|
||||
hipLaunchKernelStruct_t7 hipLaunchKernelStruct_h7;
|
||||
hipLaunchKernelStruct_h7.c1 = 'c';
|
||||
hipLaunchKernelStruct_h7.si = 1;
|
||||
#if ENABLE_ALIGNMENT_TEST_SMALL_BAR // This is broken on small bar
|
||||
hipLaunchKernelGGL(HIP_KERNEL_NAME(hipLaunchKernelStructFunc7),
|
||||
dim3(BLOCK_DIM_SIZE),
|
||||
dim3(1), 0, 0, hipLaunchKernelStruct_h7,
|
||||
result_d);
|
||||
ResultValidation();
|
||||
#endif
|
||||
|
||||
// Test: Passing struct with packed aligned to 4Bytes
|
||||
ResetValidationMem();
|
||||
hipLaunchKernelStruct_t8 hipLaunchKernelStruct_h8;
|
||||
hipLaunchKernelStruct_h8.c1 = 'c';
|
||||
hipLaunchKernelStruct_h8.si = 1;
|
||||
hipLaunchKernelGGL(HIP_KERNEL_NAME(hipLaunchKernelStructFunc8),
|
||||
dim3(BLOCK_DIM_SIZE),
|
||||
dim3(1), 0, 0, hipLaunchKernelStruct_h8,
|
||||
result_d);
|
||||
// packed member broken on large and small bar setup.
|
||||
#if ENABLE_PACKED_TEST
|
||||
ResultValidation();
|
||||
#endif
|
||||
|
||||
// Test: Passing struct with packed to 4Bytes
|
||||
ResetValidationMem();
|
||||
hipLaunchKernelStruct_t8A hipLaunchKernelStruct_h8A;
|
||||
hipLaunchKernelStruct_h8A.c1 = 'c';
|
||||
hipLaunchKernelStruct_h8A.si = 1;
|
||||
hipLaunchKernelGGL(HIP_KERNEL_NAME(hipLaunchKernelStructFunc8A),
|
||||
dim3(BLOCK_DIM_SIZE),
|
||||
dim3(1), 0, 0, hipLaunchKernelStruct_h8A,
|
||||
result_d);
|
||||
// packed member broken on large and small bar setup.
|
||||
#if ENABLE_PACKED_TEST
|
||||
ResultValidation();
|
||||
#endif
|
||||
|
||||
// Test: Passing struct with aligned(4) to 4Bytes, size is 8Bytes
|
||||
ResetValidationMem();
|
||||
hipLaunchKernelStruct_t8B hipLaunchKernelStruct_h8B;
|
||||
hipLaunchKernelStruct_h8B.c1 = 'c';
|
||||
hipLaunchKernelStruct_h8B.si = 1;
|
||||
hipLaunchKernelGGL(HIP_KERNEL_NAME(hipLaunchKernelStructFunc8B),
|
||||
dim3(BLOCK_DIM_SIZE),
|
||||
dim3(1), 0, 0, hipLaunchKernelStruct_h8B,
|
||||
result_d);
|
||||
// alignment is broken hence disabled the validation part
|
||||
#if ENABLE_ALIGNMENT_TEST_SMALL_BAR
|
||||
ResultValidation();
|
||||
#endif
|
||||
|
||||
// Test: Passing const struct object to a hipLaunchKernelGGL()
|
||||
ResetValidationMem();
|
||||
uint32_t* ip_d9;
|
||||
// allocating memory for char pointer on device
|
||||
HIPCHECK(hipMalloc((void**)&ip_d9, sizeof(uint32_t)));
|
||||
HIPCHECK(hipMemset(ip_d9, 1, sizeof(uint32_t)));
|
||||
// ip_d9 passed as pointer to struct member, struct.ip = &ip_d9
|
||||
const hipLaunchKernelStruct_t9 hipLaunchKernelStruct_h9 = {'c', ip_d9};
|
||||
hipLaunchKernelGGL(HIP_KERNEL_NAME(hipLaunchKernelStructFunc9),
|
||||
dim3(BLOCK_DIM_SIZE),
|
||||
dim3(1), 0, 0, hipLaunchKernelStruct_h9,
|
||||
result_d);
|
||||
#if ENABLE_DECLARE_INITIALIZATION_POINTER
|
||||
ResultValidation();
|
||||
#endif
|
||||
|
||||
|
||||
// Test: Passing struct with uintN_t as member variables
|
||||
ResetValidationMem();
|
||||
hipLaunchKernelStruct_t10 hipLaunchKernelStruct_h10;
|
||||
hipLaunchKernelStruct_h10.u64 = UINT64_MAX;
|
||||
hipLaunchKernelStruct_h10.u32 = 1;
|
||||
hipLaunchKernelStruct_h10.u8 = UINT8_MAX;
|
||||
hipLaunchKernelGGL(HIP_KERNEL_NAME(hipLaunchKernelStructFunc10),
|
||||
dim3(BLOCK_DIM_SIZE),
|
||||
dim3(1), 0, 0, hipLaunchKernelStruct_h10,
|
||||
result_d);
|
||||
ResultValidation();
|
||||
|
||||
|
||||
// Test: Passing struct with uintN_t as member variables
|
||||
ResetValidationMem();
|
||||
hipLaunchKernelStruct_t11 hipLaunchKernelStruct_h11;
|
||||
hipLaunchKernelStruct_h11.i1 = 1;
|
||||
hipLaunchKernelStruct_h11.vint = 0;
|
||||
hipLaunchKernelGGL(HIP_KERNEL_NAME(hipLaunchKernelStructFunc11),
|
||||
dim3(BLOCK_DIM_SIZE),
|
||||
dim3(1), 0, 0, hipLaunchKernelStruct_h11,
|
||||
result_d);
|
||||
ResultValidation();
|
||||
|
||||
// Test: Passing struct with simple class object
|
||||
ResetValidationMem();
|
||||
hipLaunchKernelStruct_t12 hipLaunchKernelStruct_h12;
|
||||
hipLaunchKernelStruct_h12.c1 = 'c';
|
||||
hipLaunchKernelGGL(HIP_KERNEL_NAME(hipLaunchKernelStructFunc12),
|
||||
dim3(BLOCK_DIM_SIZE),
|
||||
dim3(1), 0, 0, hipLaunchKernelStruct_h12,
|
||||
result_d);
|
||||
#if ENABLE_CLASS_OBJ_ACCESS // access class obj from device broken
|
||||
// Validation part of the struct, hipLaunchKernelStructFunc12
|
||||
ResultValidation();
|
||||
#endif
|
||||
|
||||
// Test: Passing struct with simple __device__ func()
|
||||
ResetValidationMem();
|
||||
hipLaunchKernelStruct_t13 hipLaunchKernelStruct_h13;
|
||||
hipLaunchKernelStruct_h13.i1 = 1;
|
||||
hipLaunchKernelGGL(HIP_KERNEL_NAME(hipLaunchKernelStructFunc13),
|
||||
dim3(BLOCK_DIM_SIZE),
|
||||
dim3(1), 0, 0, hipLaunchKernelStruct_h13,
|
||||
result_d);
|
||||
ResultValidation();
|
||||
|
||||
// Test: Passing struct with array variable, write to from device
|
||||
ResetValidationMem();
|
||||
hipLaunchKernelStruct_t14 hipLaunchKernelStruct_h14;
|
||||
hipLaunchKernelStruct_h14.readint = 1;
|
||||
hipLaunchKernelGGL(HIP_KERNEL_NAME(hipLaunchKernelStructFunc14),
|
||||
dim3(BLOCK_DIM_SIZE),
|
||||
dim3(1), 0, 0, hipLaunchKernelStruct_h14,
|
||||
result_d);
|
||||
ResultValidation();
|
||||
|
||||
|
||||
// Test: Passing struct with heap memory, read to from device
|
||||
ResetValidationMem();
|
||||
hipLaunchKernelStruct_t15 hipLaunchKernelStruct_h15;
|
||||
hipLaunchKernelStruct_h15.c1 = 'c';
|
||||
|
||||
#if ENABLE_HEAP_MEMORY_ACCESS // causing page fault here,
|
||||
// on small bar set
|
||||
HIPCHECK(hipMalloc(&hipLaunchKernelStruct_h15.heapmem,
|
||||
BLOCK_DIM_SIZE*sizeof(int)));
|
||||
HIPCHECK(hipMemset(&hipLaunchKernelStruct_h15.heapmem,
|
||||
0, BLOCK_DIM_SIZE));
|
||||
hipLaunchKernelGGL(HIP_KERNEL_NAME(hipLaunchKernelStructFunc15),
|
||||
dim3(BLOCK_DIM_SIZE),
|
||||
dim3(1), 0, 0, hipLaunchKernelStruct_h15,
|
||||
result_d);
|
||||
ResultValidation();
|
||||
#endif
|
||||
|
||||
// Test: Passing simple template struct
|
||||
ResetValidationMem();
|
||||
hipLaunchKernelStruct_t16<char> hipLaunchKernelStruct_h16;
|
||||
hipLaunchKernelStruct_h16.t1 = 'c';
|
||||
hipLaunchKernelGGL(HIP_KERNEL_NAME(hipLaunchKernelStructFunc16),
|
||||
dim3(BLOCK_DIM_SIZE),
|
||||
dim3(1), 0, 0, hipLaunchKernelStruct_h16,
|
||||
result_d);
|
||||
ResultValidation();
|
||||
|
||||
// Test: Passing simple explicit template struct
|
||||
ResetValidationMem();
|
||||
hipLaunchKernelStruct_t17<int> hipLaunchKernelStruct_h17;
|
||||
hipLaunchKernelStruct_h17.t1 = 1;
|
||||
hipLaunchKernelGGL(HIP_KERNEL_NAME(hipLaunchKernelStructFunc17),
|
||||
dim3(BLOCK_DIM_SIZE),
|
||||
dim3(1), 0, 0, hipLaunchKernelStruct_h17,
|
||||
result_d);
|
||||
ResultValidation();
|
||||
|
||||
// Test: Passing struct with simple __device__ func() to struct memory
|
||||
ResetValidationMem();
|
||||
hipLaunchKernelStruct_t18 hipLaunchKernelStruct_h18;
|
||||
hipLaunchKernelGGL(HIP_KERNEL_NAME(hipLaunchKernelStructFunc18),
|
||||
dim3(BLOCK_DIM_SIZE),
|
||||
dim3(1), 0, 0, hipLaunchKernelStruct_h18,
|
||||
result_d);
|
||||
ResultValidation();
|
||||
|
||||
// Test: Passing user defined stack,
|
||||
ResetValidationMem();
|
||||
hipLaunchKernelStruct_t19 hipLaunchKernelStruct_h19;
|
||||
hipLaunchKernelGGL(HIP_KERNEL_NAME(hipLaunchKernelStructFunc19),
|
||||
dim3(BLOCK_DIM_SIZE),
|
||||
dim3(1), 0, 0, hipLaunchKernelStruct_h19);
|
||||
#if ENABLE_USER_STL
|
||||
// Validation part of the struct, hipLaunchKernelStructFunc19
|
||||
HIPASSERT(hipLaunchKernelStruct_h19.stackSize() == 0);
|
||||
#endif
|
||||
|
||||
// Test: Passing struct which is initiazed out of order
|
||||
// accessing same elements in order from device
|
||||
ResetValidationMem();
|
||||
hipLaunchKernelStruct_t20 hipLaunchKernelStruct_h20 =
|
||||
// out of order initalization
|
||||
{.name = 'A', .rank = 2, .age = 42};
|
||||
bool *result_d20, *result_h20;
|
||||
#if ENABLE_OUT_OF_ORDER_INITIALIZATION
|
||||
hipLaunchKernelGGL(HIP_KERNEL_NAME(hipLaunchKernelStructFunc20),
|
||||
dim3(BLOCK_DIM_SIZE),
|
||||
dim3(1), 0, 0, hipLaunchKernelStruct_h20, result_d);
|
||||
ResultValidation();
|
||||
#endif
|
||||
|
||||
// Test: Passing struct with bit fields operation
|
||||
// accessing same elements in order from device
|
||||
ResetValidationMem();
|
||||
hipLaunchKernelStruct_t21 hipLaunchKernelStruct_h21 =
|
||||
// out of order initalization
|
||||
{2,0};
|
||||
bool *result_d21, *result_h21;
|
||||
hipLaunchKernelGGL(HIP_KERNEL_NAME(hipLaunchKernelStructFunc21),
|
||||
dim3(BLOCK_DIM_SIZE),
|
||||
dim3(1), 0, 0, hipLaunchKernelStruct_h21, result_d);
|
||||
#if ENABLE_BIT_FIELDS
|
||||
ResultValidation();
|
||||
#endif
|
||||
|
||||
// Test: Passing the different hipLaunchParm options:
|
||||
float* Ad;
|
||||
hipMalloc((void**)&Ad, 1024);
|
||||
hipLaunchKernelGGL(HIP_KERNEL_NAME(vAdd), size_t(1024), 1, 0, 0, Ad);
|
||||
hipLaunchKernelGGL(HIP_KERNEL_NAME(vAdd), 1024, dim3(1), 0, 0, Ad);
|
||||
hipLaunchKernelGGL(HIP_KERNEL_NAME(vAdd), dim3(1024), 1, 0, 0, Ad);
|
||||
hipLaunchKernelGGL(HIP_KERNEL_NAME(vAdd), dim3(1024), dim3(1), 0, 0, Ad);
|
||||
|
||||
// Test the different hipLaunchParm options:
|
||||
hipLaunchKernel(vAdd, size_t(1024), 1, 0, 0, Ad);
|
||||
hipLaunchKernel(vAdd, 1024, dim3(1), 0, 0, Ad);
|
||||
hipLaunchKernel(vAdd, dim3(1024), 1, 0, 0, Ad);
|
||||
hipLaunchKernel(vAdd, dim3(1024), dim3(1), 0, 0, Ad);
|
||||
|
||||
// Test case with hipLaunchKernel inside another macro:
|
||||
// Test: Passing hipLaunchKernel inside another macro:
|
||||
float e0;
|
||||
GPU_PRINT_TIME(hipLaunchKernel(vAdd, dim3(1024), dim3(1), 0, 0, Ad), e0, j);
|
||||
GPU_PRINT_TIME(WRAP(hipLaunchKernel(vAdd, dim3(1024), dim3(1), 0, 0, Ad)), e0, j);
|
||||
GPU_PRINT_TIME(hipLaunchKernelGGL(vAdd, dim3(1024),
|
||||
dim3(1), 0, 0, Ad), e0, j);
|
||||
GPU_PRINT_TIME(WRAP(hipLaunchKernelGGL(vAdd, dim3(1024),
|
||||
dim3(1), 0, 0, Ad)), e0, j);
|
||||
|
||||
#ifdef EXTRA_PARENS_1
|
||||
// Don't wrap hipLaunchKernel in extra set of parens:
|
||||
GPU_PRINT_TIME((hipLaunchKernel(vAdd, dim3(1024), dim3(1), 0, 0, Ad)), e0, j);
|
||||
GPU_PRINT_TIME((hipLaunchKernelGGL(vAdd, dim3(1024),
|
||||
dim3(1), 0, 0, Ad)), e0, j);
|
||||
#endif
|
||||
|
||||
MY_LAUNCH(hipLaunchKernel(vAdd, dim3(1024), dim3(1), 0, 0, Ad), true, "firstCall");
|
||||
MY_LAUNCH(hipLaunchKernelGGL(vAdd, dim3(1024), dim3(1),
|
||||
0, 0, Ad), true, "firstCall");
|
||||
|
||||
float* A;
|
||||
float e1;
|
||||
@@ -90,8 +937,12 @@ int main() {
|
||||
|
||||
#ifdef EXTRA_PARENS_2
|
||||
// MY_LAUNCH_WITH_PAREN wraps cmd in () which can cause issues.
|
||||
MY_LAUNCH_WITH_PAREN(hipLaunchKernel(vAdd, dim3(1024), dim3(1), 0, 0, Ad), true, "firstCall");
|
||||
MY_LAUNCH_WITH_PAREN(hipLaunchKernelGGL(vAdd, dim3(1024),
|
||||
dim3(1), 0, 0, Ad), true, "firstCall");
|
||||
#endif
|
||||
|
||||
HIPCHECK(hipHostFree(result_h));
|
||||
HIPCHECK(hipFree(result_d));
|
||||
|
||||
passed();
|
||||
}
|
||||
|
||||
@@ -0,0 +1,449 @@
|
||||
/*
|
||||
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.
|
||||
*/
|
||||
|
||||
/* HIT_START
|
||||
* BUILD: %t %s ../test_common.cpp HIPCC_OPTIONS -O3
|
||||
* RUN: %t
|
||||
* HIT_END
|
||||
*/
|
||||
|
||||
|
||||
#include "../test_common.h"
|
||||
|
||||
#define test_passed(test_name) printf("%s %s PASSED!%s\n", KGRN, #test_name, KNRM);
|
||||
|
||||
class HipFunctorTests {
|
||||
public:
|
||||
// Test that a class functor can be passed to hiplaunchparam
|
||||
// and can be used in kernel
|
||||
void TestForSimpleClassFunctor(void);
|
||||
// Test that a templated class functor can be passed to hiplaunchparam
|
||||
// and can be used in kernel
|
||||
void TestForClassTemplateFunctor(void);
|
||||
// Test that a class functor object ptr can be passed to hiplaunchparam
|
||||
// and can be used in kernel
|
||||
void TestForClassObjPtrFunctor(void);
|
||||
// Test that a class object containing functor can be passed to hiplaunchparam
|
||||
// and can be used in kernel
|
||||
void TestForFunctorContainInClassObj(void);
|
||||
// Test that a stuct functor can be passed to hiplaunchparam
|
||||
// and can be used in kernel
|
||||
void TestForSimpleStructFunctor(void);
|
||||
// Test that a stuct functor object ptr can be passed to hiplaunchparam
|
||||
// and can be used in kernel
|
||||
void TestForStructObjPtrFunctor(void);
|
||||
// Test that a templated struct functor can be passed to hiplaunchparam
|
||||
// and can be used in kernel
|
||||
void TestForStructTemplateFunctor(void);
|
||||
// Test that a struct object containing functor can be passed to hiplaunchparam
|
||||
// and can be used in kernel
|
||||
void TestForFunctorContainInStructObj(void);
|
||||
};
|
||||
|
||||
|
||||
|
||||
|
||||
static const int BLOCK_DIM_SIZE = 1024;
|
||||
static const int THREADS_PER_BLOCK = 1;
|
||||
|
||||
// class functor tests
|
||||
|
||||
// Simple doubler Functor
|
||||
class DoublerFunctor{
|
||||
public:
|
||||
__device__ int operator()(int x) { return x * 2;}
|
||||
};
|
||||
|
||||
|
||||
|
||||
|
||||
// simple doubler functor passed to kernel
|
||||
__global__ void DoublerFunctorKernel(
|
||||
DoublerFunctor doubler_,
|
||||
bool* deviceResult) {
|
||||
int x = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
int result = doubler_(5);
|
||||
deviceResult[x] = (result == 10);
|
||||
}
|
||||
|
||||
void HipFunctorTests::TestForSimpleClassFunctor(void) {
|
||||
DoublerFunctor doubler;
|
||||
bool *deviceResults, *hostResults;
|
||||
HIPCHECK(hipMalloc(&deviceResults, BLOCK_DIM_SIZE*sizeof(bool)));
|
||||
HIPCHECK(hipHostMalloc(&hostResults, BLOCK_DIM_SIZE*sizeof(bool)));
|
||||
for (int k = 0; k < BLOCK_DIM_SIZE; ++k) {
|
||||
// initialize to false, will be set to
|
||||
// true if the functor is called in device code
|
||||
hostResults[k] = false;
|
||||
}
|
||||
|
||||
|
||||
HIPCHECK(hipMemcpy(deviceResults, hostResults, BLOCK_DIM_SIZE*sizeof(bool),
|
||||
hipMemcpyHostToDevice));
|
||||
hipLaunchKernelGGL(DoublerFunctorKernel, dim3(BLOCK_DIM_SIZE),
|
||||
dim3(THREADS_PER_BLOCK), 0, 0, doubler, deviceResults);
|
||||
|
||||
// Validation part of TestForSimpleClassFunctor
|
||||
HIPCHECK(hipMemcpy(hostResults, deviceResults, BLOCK_DIM_SIZE*sizeof(bool),
|
||||
hipMemcpyDeviceToHost));
|
||||
for (int k = 0; k < BLOCK_DIM_SIZE; ++k)
|
||||
HIPASSERT(hostResults[k] == true);
|
||||
HIPCHECK(hipHostFree(hostResults));
|
||||
HIPCHECK(hipFree(deviceResults));
|
||||
}
|
||||
|
||||
|
||||
|
||||
|
||||
// pointer functor passed to kernel
|
||||
__global__ void PtrDoublerFunctorKernel(
|
||||
DoublerFunctor *doubler_,
|
||||
bool* deviceResult) {
|
||||
int x = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
int result = (*doubler_)(5);
|
||||
deviceResult[x] = (result == 10);
|
||||
}
|
||||
|
||||
void HipFunctorTests::TestForClassObjPtrFunctor(void) {
|
||||
DoublerFunctor *ptrdoubler;
|
||||
bool *deviceResults, *hostResults;
|
||||
HIPCHECK(hipMalloc(&deviceResults, BLOCK_DIM_SIZE*sizeof(bool)));
|
||||
HIPCHECK(hipHostMalloc(&hostResults, BLOCK_DIM_SIZE*sizeof(bool)));
|
||||
for (int k = 0; k < BLOCK_DIM_SIZE; ++k) {
|
||||
// initialize to false, will be set to
|
||||
// true if the functor is called in device code
|
||||
hostResults[k] = false;
|
||||
}
|
||||
|
||||
|
||||
HIPCHECK(hipMemcpy(deviceResults, hostResults, BLOCK_DIM_SIZE*sizeof(bool),
|
||||
hipMemcpyHostToDevice));
|
||||
hipLaunchKernelGGL(PtrDoublerFunctorKernel, dim3(BLOCK_DIM_SIZE),
|
||||
dim3(THREADS_PER_BLOCK), 0, 0, ptrdoubler, deviceResults);
|
||||
|
||||
// Validation part of TestForClassObjPtrFunctor
|
||||
HIPCHECK(hipMemcpy(hostResults, deviceResults, BLOCK_DIM_SIZE*sizeof(bool),
|
||||
hipMemcpyDeviceToHost));
|
||||
for (int k = 0; k < BLOCK_DIM_SIZE; ++k)
|
||||
HIPASSERT(hostResults[k] == true);
|
||||
HIPCHECK(hipHostFree(hostResults));
|
||||
HIPCHECK(hipFree(deviceResults));
|
||||
delete ptrdoubler;
|
||||
}
|
||||
|
||||
class compare {
|
||||
public:
|
||||
template<typename T1, typename T2>
|
||||
__device__ bool operator()(const T1& v1, const T2& v2) {
|
||||
return v1 > v2;
|
||||
}
|
||||
};
|
||||
|
||||
|
||||
|
||||
|
||||
// template functor passed to kernel
|
||||
__global__ void TemplateFunctorKernel(
|
||||
compare compare_,
|
||||
bool* deviceResult) {
|
||||
int x = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
deviceResult[x] = compare_(2.2, 2.1);
|
||||
deviceResult[x] = compare_(2, 1);
|
||||
deviceResult[x] = compare_('b', 'a');
|
||||
}
|
||||
|
||||
void HipFunctorTests::TestForClassTemplateFunctor(void) {
|
||||
compare comparefunctor;
|
||||
bool *deviceResults, *hostResults;
|
||||
HIPCHECK(hipMalloc(&deviceResults, BLOCK_DIM_SIZE*sizeof(bool)));
|
||||
HIPCHECK(hipHostMalloc(&hostResults, BLOCK_DIM_SIZE*sizeof(bool)));
|
||||
for (int k = 0; k < BLOCK_DIM_SIZE; ++k) {
|
||||
// initialize to false, will be set to
|
||||
// true if the functor is called in device code
|
||||
hostResults[k] = false;
|
||||
}
|
||||
|
||||
HIPCHECK(hipMemcpy(deviceResults, hostResults, BLOCK_DIM_SIZE*sizeof(bool),
|
||||
hipMemcpyHostToDevice));
|
||||
hipLaunchKernelGGL(TemplateFunctorKernel, dim3(BLOCK_DIM_SIZE),
|
||||
dim3(THREADS_PER_BLOCK), 0, 0, comparefunctor, deviceResults);
|
||||
|
||||
// Validation part of TestForClassTemplateFunctor
|
||||
HIPCHECK(hipMemcpy(hostResults, deviceResults, BLOCK_DIM_SIZE*sizeof(bool),
|
||||
hipMemcpyDeviceToHost));
|
||||
for (int k = 0; k < BLOCK_DIM_SIZE; ++k)
|
||||
HIPASSERT(hostResults[k] == true);
|
||||
HIPCHECK(hipHostFree(hostResults));
|
||||
HIPCHECK(hipFree(deviceResults));
|
||||
}
|
||||
|
||||
|
||||
// Doubler calculator
|
||||
class DoublerCalculator {
|
||||
public:
|
||||
int a, result;
|
||||
// fucntor contained in class object
|
||||
DoublerFunctor doubler;
|
||||
};
|
||||
|
||||
|
||||
|
||||
// doubler functor conatined in class obj passed to kernel
|
||||
__global__ void DoublerCalculatorFunctorKernel(
|
||||
DoublerCalculator doubler_,
|
||||
bool* deviceResult) {
|
||||
int x = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
int result = doubler_.doubler(doubler_.a);
|
||||
deviceResult[x] = (doubler_.result == result);
|
||||
}
|
||||
|
||||
void HipFunctorTests::TestForFunctorContainInClassObj(void) {
|
||||
DoublerCalculator Doubler;
|
||||
bool *deviceResults, *hostResults;
|
||||
HIPCHECK(hipMalloc(&deviceResults, BLOCK_DIM_SIZE*sizeof(bool)));
|
||||
HIPCHECK(hipHostMalloc(&hostResults, BLOCK_DIM_SIZE*sizeof(bool)));
|
||||
for (int k = 0; k < BLOCK_DIM_SIZE; ++k) {
|
||||
// initialize to false, will be set to
|
||||
// true if the functor is called in device code
|
||||
hostResults[k] = false;
|
||||
}
|
||||
|
||||
Doubler.a = 5;
|
||||
Doubler.result = 10;
|
||||
// pass comparefunctor to hipLaunchParm
|
||||
|
||||
HIPCHECK(hipMemcpy(deviceResults, hostResults, BLOCK_DIM_SIZE*sizeof(bool),
|
||||
hipMemcpyHostToDevice));
|
||||
hipLaunchKernelGGL(DoublerCalculatorFunctorKernel, dim3(BLOCK_DIM_SIZE),
|
||||
dim3(THREADS_PER_BLOCK), 0, 0, Doubler, deviceResults);
|
||||
|
||||
// Validation part of TestForStructTemplateFunctor
|
||||
HIPCHECK(hipMemcpy(hostResults, deviceResults, BLOCK_DIM_SIZE*sizeof(bool),
|
||||
hipMemcpyDeviceToHost));
|
||||
for (int k = 0; k < BLOCK_DIM_SIZE; ++k)
|
||||
HIPASSERT(hostResults[k] == true);
|
||||
HIPCHECK(hipHostFree(hostResults));
|
||||
HIPCHECK(hipFree(deviceResults));
|
||||
}
|
||||
|
||||
|
||||
// Struct functor tests
|
||||
|
||||
// Simple doubler Functor
|
||||
struct sDoublerFunctor {
|
||||
public:
|
||||
__device__ int operator()(int x) { return x * 2;}
|
||||
};
|
||||
|
||||
|
||||
|
||||
|
||||
// simple sturct doubler functor passed to kernel
|
||||
__global__ void structDoublerFunctorKernel(
|
||||
sDoublerFunctor doubler_,
|
||||
bool* deviceResult) {
|
||||
int x = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
int result = doubler_(5);
|
||||
deviceResult[x] = (result == 10);
|
||||
}
|
||||
|
||||
void HipFunctorTests::TestForSimpleStructFunctor(void) {
|
||||
sDoublerFunctor doubler;
|
||||
bool *deviceResults, *hostResults;
|
||||
HIPCHECK(hipMalloc(&deviceResults, BLOCK_DIM_SIZE*sizeof(bool)));
|
||||
HIPCHECK(hipHostMalloc(&hostResults, BLOCK_DIM_SIZE*sizeof(bool)));
|
||||
for (int k = 0; k < BLOCK_DIM_SIZE; ++k) {
|
||||
// initialize to false, will be set to
|
||||
// true if the functor is called in device code
|
||||
hostResults[k] = false;
|
||||
}
|
||||
|
||||
HIPCHECK(hipMemcpy(deviceResults, hostResults, BLOCK_DIM_SIZE*sizeof(bool),
|
||||
hipMemcpyHostToDevice));
|
||||
hipLaunchKernelGGL(structDoublerFunctorKernel, dim3(BLOCK_DIM_SIZE),
|
||||
dim3(THREADS_PER_BLOCK), 0, 0, doubler, deviceResults);
|
||||
|
||||
// Validation part of TestForSimpleStructFunctor
|
||||
HIPCHECK(hipMemcpy(hostResults, deviceResults, BLOCK_DIM_SIZE*sizeof(bool),
|
||||
hipMemcpyDeviceToHost));
|
||||
for (int k = 0; k < BLOCK_DIM_SIZE; ++k)
|
||||
HIPASSERT(hostResults[k] == true);
|
||||
HIPCHECK(hipHostFree(hostResults));
|
||||
HIPCHECK(hipFree(deviceResults));
|
||||
}
|
||||
|
||||
// ptr functor passed to kernel
|
||||
__global__ void structPtrDoublerFunctorKernel(
|
||||
sDoublerFunctor *doubler_,
|
||||
bool* deviceResult) {
|
||||
int x = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
int result = (*doubler_)(5);
|
||||
deviceResult[x] = (result == 10);
|
||||
}
|
||||
|
||||
void HipFunctorTests::TestForStructObjPtrFunctor(void) {
|
||||
sDoublerFunctor *ptrdoubler;
|
||||
bool *deviceResults, *hostResults;
|
||||
HIPCHECK(hipMalloc(&deviceResults, BLOCK_DIM_SIZE*sizeof(bool)));
|
||||
HIPCHECK(hipHostMalloc(&hostResults, BLOCK_DIM_SIZE*sizeof(bool)));
|
||||
for (int k = 0; k < BLOCK_DIM_SIZE; ++k) {
|
||||
// initialize to false, will be set to
|
||||
// true if the functor is called in device code
|
||||
hostResults[k] = false;
|
||||
}
|
||||
|
||||
|
||||
HIPCHECK(hipMemcpy(deviceResults, hostResults, BLOCK_DIM_SIZE*sizeof(bool),
|
||||
hipMemcpyHostToDevice));
|
||||
hipLaunchKernelGGL(structPtrDoublerFunctorKernel, dim3(BLOCK_DIM_SIZE),
|
||||
dim3(THREADS_PER_BLOCK), 0, 0, ptrdoubler, deviceResults);
|
||||
|
||||
// Validation part of TestForStructObjPtrFunctor
|
||||
HIPCHECK(hipMemcpy(hostResults, deviceResults, BLOCK_DIM_SIZE*sizeof(bool),
|
||||
hipMemcpyDeviceToHost));
|
||||
for (int k = 0; k < BLOCK_DIM_SIZE; ++k)
|
||||
HIPASSERT(hostResults[k] == true);
|
||||
HIPCHECK(hipHostFree(hostResults));
|
||||
HIPCHECK(hipFree(deviceResults));
|
||||
delete ptrdoubler;
|
||||
}
|
||||
|
||||
struct sCompare {
|
||||
public:
|
||||
template< typename T1, typename T2 >
|
||||
__device__ bool operator()(const T1& v1, const T2& v2) {
|
||||
return v1 > v2;
|
||||
}
|
||||
};
|
||||
|
||||
|
||||
|
||||
|
||||
// template functor passed to kernel
|
||||
__global__ void structTemplateFunctorKernel(
|
||||
sCompare compare_,
|
||||
bool* deviceResult) {
|
||||
int x = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
deviceResult[x] = compare_(2.2, 2.1);
|
||||
deviceResult[x] = compare_(2, 1);
|
||||
deviceResult[x] = compare_('b', 'a');
|
||||
}
|
||||
|
||||
void HipFunctorTests::TestForStructTemplateFunctor(void) {
|
||||
sCompare comparefunctor;
|
||||
bool *deviceResults, *hostResults;
|
||||
HIPCHECK(hipMalloc(&deviceResults, BLOCK_DIM_SIZE*sizeof(bool)));
|
||||
HIPCHECK(hipHostMalloc(&hostResults, BLOCK_DIM_SIZE*sizeof(bool)));
|
||||
for (int k = 0; k < BLOCK_DIM_SIZE; ++k) {
|
||||
// initialize to false, will be set to
|
||||
// true if the functor is called in device code
|
||||
hostResults[k] = false;
|
||||
}
|
||||
|
||||
HIPCHECK(hipMemcpy(deviceResults, hostResults, BLOCK_DIM_SIZE*sizeof(bool),
|
||||
hipMemcpyHostToDevice));
|
||||
|
||||
// pass comparefunctor to hipLaunchKernelGGL
|
||||
hipLaunchKernelGGL(structTemplateFunctorKernel, dim3(BLOCK_DIM_SIZE),
|
||||
dim3(THREADS_PER_BLOCK), 0, 0, comparefunctor, deviceResults);
|
||||
|
||||
// Validation part of TestForStructTemplateFunctor
|
||||
HIPCHECK(hipMemcpy(hostResults, deviceResults, BLOCK_DIM_SIZE*sizeof(bool),
|
||||
hipMemcpyDeviceToHost));
|
||||
for (int k = 0; k < BLOCK_DIM_SIZE; ++k)
|
||||
HIPASSERT(hostResults[k] == true);
|
||||
HIPCHECK(hipHostFree(hostResults));
|
||||
HIPCHECK(hipFree(deviceResults));
|
||||
}
|
||||
|
||||
// Doubler calculator struct
|
||||
struct sDoublerCalculator {
|
||||
public:
|
||||
int a, result;
|
||||
// fucntor contained in class object
|
||||
DoublerFunctor doubler;
|
||||
};
|
||||
|
||||
|
||||
|
||||
// doubler functor contained in struct passed to kernel
|
||||
__global__ void DoublerCalculatorFunctorKernel(
|
||||
sDoublerCalculator doubler_,
|
||||
bool* deviceResult) {
|
||||
int x = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
int result = doubler_.doubler(doubler_.a);
|
||||
deviceResult[x] = (doubler_.result == result);
|
||||
}
|
||||
|
||||
void HipFunctorTests::TestForFunctorContainInStructObj(void) {
|
||||
sDoublerCalculator Doubler;
|
||||
bool *deviceResults, *hostResults;
|
||||
HIPCHECK(hipMalloc(&deviceResults, BLOCK_DIM_SIZE*sizeof(bool)));
|
||||
HIPCHECK(hipHostMalloc(&hostResults, BLOCK_DIM_SIZE*sizeof(bool)));
|
||||
for (int k = 0; k < BLOCK_DIM_SIZE; ++k) {
|
||||
// initialize to false, will be set to
|
||||
// true if the functor is called in device code
|
||||
hostResults[k] = false;
|
||||
}
|
||||
|
||||
Doubler.a = 5;
|
||||
Doubler.result = 10;
|
||||
HIPCHECK(hipMemcpy(deviceResults, hostResults, BLOCK_DIM_SIZE*sizeof(bool),
|
||||
hipMemcpyHostToDevice));
|
||||
|
||||
|
||||
// pass comparefunctor to hipLaunchKernelGGL
|
||||
hipLaunchKernelGGL(DoublerCalculatorFunctorKernel, dim3(BLOCK_DIM_SIZE),
|
||||
dim3(THREADS_PER_BLOCK), 0, 0, Doubler, deviceResults);
|
||||
|
||||
// Validation part of TestForStructTemplateFunctor
|
||||
HIPCHECK(hipMemcpy(hostResults, deviceResults, BLOCK_DIM_SIZE*sizeof(bool),
|
||||
hipMemcpyDeviceToHost));
|
||||
for (int k = 0; k < BLOCK_DIM_SIZE; ++k)
|
||||
HIPASSERT(hostResults[k] == true);
|
||||
HIPCHECK(hipHostFree(hostResults));
|
||||
HIPCHECK(hipFree(deviceResults));
|
||||
}
|
||||
|
||||
int main() {
|
||||
HipFunctorTests FunctorTests;
|
||||
FunctorTests.TestForSimpleClassFunctor();
|
||||
test_passed(TestForSimpleClassFunctor);
|
||||
|
||||
FunctorTests.TestForClassObjPtrFunctor();
|
||||
test_passed(TestForClassObjPtrFunctor);
|
||||
|
||||
FunctorTests.TestForClassTemplateFunctor();
|
||||
test_passed(TestForClassTemplateFunctor);
|
||||
|
||||
FunctorTests.TestForSimpleStructFunctor();
|
||||
test_passed(TestForSimpleStructFunctor);
|
||||
|
||||
FunctorTests.TestForStructObjPtrFunctor();
|
||||
test_passed(TestForStructObjPtrFunctor);
|
||||
|
||||
FunctorTests.TestForStructTemplateFunctor();
|
||||
test_passed(TestForStructTemplateFunctor);
|
||||
|
||||
FunctorTests.TestForFunctorContainInClassObj();
|
||||
test_passed(TestForFunctorContainInClassObj);
|
||||
|
||||
FunctorTests.TestForFunctorContainInStructObj();
|
||||
test_passed(TestForFunctorContainInStructObj);
|
||||
}
|
||||
@@ -108,6 +108,7 @@ int main(int argc, char* argv[]) {
|
||||
HIPCHECK(hipHostUnregister(A));
|
||||
|
||||
free(A);
|
||||
delete [] Ad;
|
||||
}
|
||||
|
||||
|
||||
@@ -144,6 +145,8 @@ int main(int argc, char* argv[]) {
|
||||
|
||||
|
||||
free(A);
|
||||
free(Bh);
|
||||
hipFree(Bd);
|
||||
}
|
||||
|
||||
|
||||
|
||||
@@ -41,7 +41,7 @@ void simpleNegTest() {
|
||||
|
||||
|
||||
// Not sure what happens here, the memory must be pinned.
|
||||
e = hipMemcpyAsync(A_malloc, A_d, Nbytes, hipMemcpyHostToDevice, NULL);
|
||||
e = hipMemcpyAsync(A_malloc, A_d, Nbytes, hipMemcpyDeviceToHost, NULL);
|
||||
|
||||
printf(" async memcpy of A_malloc to A_d. Result=%d\n", e);
|
||||
// HIPASSERT (e==hipErrorInvalidValue);
|
||||
|
||||
@@ -1,98 +1,94 @@
|
||||
/*
|
||||
Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved.
|
||||
|
||||
Permission is hereby granted, free of charge, to any person obtaining a copy
|
||||
of this software and associated documentation files (the "Software"), to deal
|
||||
in the Software without restriction, including without limitation the rights
|
||||
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
|
||||
copies of the Software, and to permit persons to whom the Software is
|
||||
furnished to do so, subject to the following conditions:
|
||||
|
||||
The above copyright notice and this permission notice shall be included in
|
||||
all copies or substantial portions of the Software.
|
||||
|
||||
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
|
||||
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
|
||||
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
|
||||
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
|
||||
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
|
||||
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
THE SOFTWARE.
|
||||
*/
|
||||
// Simple test for memset.
|
||||
// Also serves as a template for other tests.
|
||||
|
||||
/* HIT_START
|
||||
* BUILD: %t %s ../../test_common.cpp
|
||||
* //Small copy
|
||||
* RUN: %t -N 10 --memsetval 0x42
|
||||
* HIT_END
|
||||
*/
|
||||
|
||||
#include "hip/hip_runtime.h"
|
||||
#include "test_common.h"
|
||||
|
||||
bool testhipMemset3D(int memsetval,int p_gpuDevice)
|
||||
{
|
||||
size_t numH = 256;
|
||||
size_t numW = 256;
|
||||
size_t depth = 1;
|
||||
size_t pitch_A;
|
||||
size_t width = numW * sizeof(char);
|
||||
size_t sizeElements = width * numH * depth;
|
||||
size_t elements = numW* numH* depth;
|
||||
|
||||
|
||||
printf ("testhipMemset3D memsetval=%2x device=%d\n", memsetval, p_gpuDevice);
|
||||
char *A_d;
|
||||
char *A_h;
|
||||
bool testResult = true;
|
||||
hipExtent extent = make_hipExtent(width, numH, depth);
|
||||
hipPitchedPtr devPitchedPtr;
|
||||
|
||||
HIPCHECK(hipMalloc3D(&devPitchedPtr, extent));
|
||||
A_h = (char*)malloc(sizeElements);
|
||||
HIPASSERT(A_h != NULL);
|
||||
for (size_t i=0; i<elements; i++) {
|
||||
A_h[i] = 1;
|
||||
}
|
||||
hipStream_t stream;
|
||||
//HIPCHECK(hipStreamCreate(&stream));
|
||||
HIPCHECK ( hipMemset3D( devPitchedPtr, memsetval, extent) );
|
||||
//HIPCHECK ( hipMemcpy3D(A_h, width, A_d, pitch_A, numW, numH, hipMemcpyDeviceToHost));
|
||||
hipMemcpy3DParms myparms = {0};
|
||||
myparms.srcPos = make_hipPos(0,0,0);
|
||||
myparms.dstPos = make_hipPos(0,0,0);
|
||||
myparms.dstPtr = make_hipPitchedPtr(A_h, width , numW, numH);
|
||||
myparms.srcPtr = devPitchedPtr;
|
||||
//myparms.dstArray = cuArray;
|
||||
myparms.extent = extent;
|
||||
#ifdef __HIP_PLATFORM_NVCC__
|
||||
myparms.kind = hipMemcpyKindToCudaMemcpyKind(hipMemcpyDeviceToHost);
|
||||
#else
|
||||
myparms.kind = hipMemcpyDeviceToHost;
|
||||
#endif
|
||||
hipMemcpy3D(&myparms);
|
||||
|
||||
for (int i=0; i<elements; i++) {
|
||||
if (A_h[i] != memsetval) {
|
||||
testResult = false;
|
||||
printf("mismatch at index:%d computed:%02x, memsetval:%02x\n", i, (int)A_h[i], (int)memsetval);
|
||||
break;
|
||||
}
|
||||
}
|
||||
hipFree(A_d);
|
||||
free(A_h);
|
||||
return testResult;
|
||||
}
|
||||
|
||||
int main(int argc, char *argv[])
|
||||
{
|
||||
HipTest::parseStandardArguments(argc, argv, true);
|
||||
bool testResult = false;
|
||||
HIPCHECK(hipSetDevice(p_gpuDevice));
|
||||
testResult = testhipMemset3D(memsetval, p_gpuDevice);
|
||||
|
||||
passed();
|
||||
|
||||
}
|
||||
/*
|
||||
Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved.
|
||||
|
||||
Permission is hereby granted, free of charge, to any person obtaining a copy
|
||||
of this software and associated documentation files (the "Software"), to deal
|
||||
in the Software without restriction, including without limitation the rights
|
||||
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
|
||||
copies of the Software, and to permit persons to whom the Software is
|
||||
furnished to do so, subject to the following conditions:
|
||||
|
||||
The above copyright notice and this permission notice shall be included in
|
||||
all copies or substantial portions of the Software.
|
||||
|
||||
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
|
||||
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
|
||||
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
|
||||
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
|
||||
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
|
||||
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
THE SOFTWARE.
|
||||
*/
|
||||
// Simple test for memset.
|
||||
// Also serves as a template for other tests.
|
||||
|
||||
/* HIT_START
|
||||
* BUILD: %t %s ../../test_common.cpp
|
||||
* //Small copy
|
||||
* RUN: %t -N 10 --memsetval 0x42
|
||||
* HIT_END
|
||||
*/
|
||||
|
||||
#include "hip/hip_runtime.h"
|
||||
#include "test_common.h"
|
||||
|
||||
bool testhipMemset3D(int memsetval,int p_gpuDevice)
|
||||
{
|
||||
size_t numH = 256;
|
||||
size_t numW = 256;
|
||||
size_t depth = 10;
|
||||
size_t width = numW * sizeof(char);
|
||||
size_t sizeElements = width * numH * depth;
|
||||
size_t elements = numW* numH* depth;
|
||||
|
||||
|
||||
printf ("testhipMemset3D memsetval=%2x device=%d\n", memsetval, p_gpuDevice);
|
||||
char *A_h;
|
||||
bool testResult = true;
|
||||
hipExtent extent = make_hipExtent(width, numH, depth);
|
||||
hipPitchedPtr devPitchedPtr;
|
||||
|
||||
HIPCHECK(hipMalloc3D(&devPitchedPtr, extent));
|
||||
A_h = (char*)malloc(sizeElements);
|
||||
HIPASSERT(A_h != NULL);
|
||||
for (size_t i=0; i<elements; i++) {
|
||||
A_h[i] = 1;
|
||||
}
|
||||
HIPCHECK ( hipMemset3D( devPitchedPtr, memsetval, extent) );
|
||||
hipMemcpy3DParms myparms = {0};
|
||||
myparms.srcPos = make_hipPos(0,0,0);
|
||||
myparms.dstPos = make_hipPos(0,0,0);
|
||||
myparms.dstPtr = make_hipPitchedPtr(A_h, width , numW, numH);
|
||||
myparms.srcPtr = devPitchedPtr;
|
||||
myparms.extent = extent;
|
||||
#ifdef __HIP_PLATFORM_NVCC__
|
||||
myparms.kind = hipMemcpyKindToCudaMemcpyKind(hipMemcpyDeviceToHost);
|
||||
#else
|
||||
myparms.kind = hipMemcpyDeviceToHost;
|
||||
#endif
|
||||
HIPCHECK(hipMemcpy3D(&myparms));
|
||||
|
||||
for (int i=0; i<elements; i++) {
|
||||
if (A_h[i] != memsetval) {
|
||||
testResult = false;
|
||||
printf("mismatch at index:%d computed:%02x, memsetval:%02x\n", i, (int)A_h[i], (int)memsetval);
|
||||
break;
|
||||
}
|
||||
}
|
||||
HIPCHECK(hipFree(devPitchedPtr.ptr));
|
||||
free(A_h);
|
||||
return testResult;
|
||||
}
|
||||
|
||||
int main(int argc, char *argv[])
|
||||
{
|
||||
HipTest::parseStandardArguments(argc, argv, true);
|
||||
bool testResult = false;
|
||||
HIPCHECK(hipSetDevice(p_gpuDevice));
|
||||
testResult = testhipMemset3D(memsetval, p_gpuDevice);
|
||||
if (testResult) {
|
||||
passed();
|
||||
} else {
|
||||
exit(EXIT_FAILURE);
|
||||
}
|
||||
}
|
||||
|
||||
Tagairt in Eagrán Nua
Cuir bac ar úsáideoir