AI/rocm-systems
2
0
Fork 0
Koodi Ongelmat Pull-pyynnöt Actions Paketit Projektit Julkaisut Wiki Toiminta

Merge 'master' into 'amd-master'

Selaa lähdekoodia 
Change-Id: I09eb56a50f07d30b6bd0e03e5ee72e92d373c904


[ROCm/clr commit: e27b0c8129]
This commit is contained in:
Jenkins
2020-01-09 06:06:55 -05:00
vanhempi 544ad9a8f1 d814f1e292
commit 80f40374b0
50 muutettua tiedostoa jossa 1441 lisäystä ja 875 poistoa
Download Patch File Download Diff File Expand all files Collapse all files
projects/clr/hipamd/.gitignore
+3 -3
Näytä tiedosto
@@ -3,9 +3,6 @@
*.o
*.exe
*.swp
*.Po
hip-amdinternal
HIP-Examples
lib
packages
build
@@ -14,5 +11,8 @@ bin/hipBusBandwidth
bin/hipDispatchLatency
bin/hipify-clang
include/hip/hcc_detail/hip_prof_str.h
include/hip/hip_version.h
tags
samples/1_Utils/hipInfo/hipInfo
samples/1_Utils/hipBusBandwidth/hipBusBandwidth
samples/1_Utils/hipDispatchLatency/hipDispatchLatency
projects/clr/hipamd/CMakeLists.txt
+13
Näytä tiedosto
@@ -343,6 +343,19 @@ endif()
# Generate .hipVersion
file(WRITE "${PROJECT_BINARY_DIR}/.hipVersion" ${_versionInfo})
# Generate hip_version.h
set(_versionInfoHeader
"// Auto-generated by cmake\n
#ifndef HIP_VERSION_H
#define HIP_VERSION_H\n
#define HIP_VERSION_MAJOR ${HIP_VERSION_MAJOR}
#define HIP_VERSION_MINOR ${HIP_VERSION_MINOR}
#define HIP_VERSION_PATCH ${HIP_VERSION_GITDATE}
#define HIP_VERSION (HIP_VERSION_MAJOR * 100 + HIP_VERSION_MINOR)\n
#endif\n
")
file(WRITE "${CMAKE_CURRENT_SOURCE_DIR}/include/hip/hip_version.h" ${_versionInfoHeader})
# Build doxygen documentation
find_program(DOXYGEN_EXE doxygen)
if(DOXYGEN_EXE)
projects/clr/hipamd/INSTALL.md
+32 -9
Näytä tiedosto
@@ -42,20 +42,43 @@ apt-get install hip-hcc
* Install the [rocm](http://gpuopen.com/getting-started-with-boltzmann-components-platforms-installation/) packages. ROCm will install some of the necessary components, including the kernel driver, HSA runtime, etc.
* Build LLVM/clang/lld by using the following repository and branch and following the general LLVM/clang build procedure. It is recommended to use -DCMAKE_INSTALL_PREFIX=/opt/rocm/llvm with cmake so that LLVM/clang/lld are installed to the default path expected by hipcc.
* Build HIP-Clang
```
git clone https://github.com/llvm/llvm-project.git
cd llvm-project/llvm/tools
ln -s clang ../../clang
ln -s lld ../../lld
cd ../..
mkdir -p build && cd build
cmake -DCMAKE_INSTALL_PREFIX=/opt/rocm/llvm -DCMAKE_BUILD_TYPE=Release -DLLVM_ENABLE_ASSERTIONS=1 -DLLVM_TARGETS_TO_BUILD="AMDGPU;X86" ../llvm
make -j
sudo make install
```
* LLVM: https://github.com/RadeonOpenCompute/llvm.git amd-common branch
* clang: https://github.com/RadeonOpenCompute/clang amd-common branch
* lld: https://github.com/RadeonOpenCompute/lld amd-common branch
* Build Rocm device library
* Checkout https://github.com/RadeonOpenCompute/ROCm-Device-Libs.git master branch and build it with clang built from the last step.
```
export PATH=/opt/rocm/llvm/bin:$PATH
git clone -b master https://github.com/RadeonOpenCompute/ROCm-Device-Libs.git
cd ROCm-Device-Libs
mkdir -p build && cd build
CC=clang CXX=clang++ cmake -DLLVM_DIR=/opt/rocm/llvm -DCMAKE_BUILD_TYPE=Release -DLLVM_ENABLE_WERROR=1 -DLLVM_ENABLE_ASSERTIONS=1 ..
make -j
sudo make install
```
* Build HIP
* Checkout https://github.com/ROCm-Developer-Tools/HIP.git master branch and build it with HCC installed with ROCm packages. Please use -DHIP_COMPILER=clang with cmake to enable hip-clang.
```
git clone -b master https://github.com/ROCm-Developer-Tools/HIP.git
cd HIP
mkdir -p build && cd build
cmake -DCMAKE_INSTALL_PREFIX=/opt/rocm/hip -DHIP_COMPILER=clang -DCMAKE_BUILD_TYPE=Release ..
make -j
sudo make install
```
* Default paths and environment variables:
* By default HIP looks for HSA in /opt/rocm/hsa (can be overridden by setting HSA_PATH environment variable)
projects/clr/hipamd/Jenkinsfile
Vendored
+4 -4
Näytä tiedosto
@@ -295,13 +295,13 @@ def docker_upload_dockerhub( String local_org, String image_name, String remote_
String build_config = 'Release'
String job_name = env.JOB_NAME.toLowerCase( )
// The following launches 3 builds in parallel: rocm-head, rocm-2.9.x and cuda-10.x
parallel rocm_2_9:
// The following launches 3 builds in parallel: rocm-head, rocm-3.0.x and cuda-10.x
parallel rocm_3_0:
{
node('hip-rocm')
{
String hcc_ver = 'rocm-2.9.x'
String from_image = 'ci_test_nodes/rocm-2.9.x/ubuntu-16.04:latest'
String hcc_ver = 'rocm-3.0.x'
String from_image = 'ci_test_nodes/rocm-3.0.x/ubuntu-16.04:latest'
String inside_args = '--device=/dev/kfd --device=/dev/dri --group-add=video'
// Checkout source code, dependencies and version files
projects/clr/hipamd/bin/hipcc
+1 -3
Näytä tiedosto
@@ -119,8 +119,6 @@ sub delete_temp_dirs {
#HIP_PLATFORM controls whether to use NVCC or HCC for compilation:
$HIP_PLATFORM= `$HIP_PATH/bin/hipconfig --platform` // "hcc";
$HIP_VERSION= `$HIP_PATH/bin/hipconfig --version`;
($HIP_VERSION_MAJOR, $HIP_VERSION_MINOR, $HIP_VERSION_PATCH) = split(/\./, $HIP_VERSION);
my ($HIP_VERSION_GITDATE, $HIP_VERSION_GITHASH) = split(/-/, $HIP_VERSION_PATCH);
$HIP_COMPILER= $hipConfig{'HIP_COMPILER'};
$HIP_RUNTIME= $hipConfig{'HIP_RUNTIME'};
@@ -341,7 +339,7 @@ if ($HIP_PLATFORM eq "clang") {
}
# Add paths to common HIP includes:
$HIPCXXFLAGS .= " -isystem $HIP_INCLUDE_PATH -DHIP_VERSION_MAJOR=$HIP_VERSION_MAJOR -DHIP_VERSION_MINOR=$HIP_VERSION_MINOR -DHIP_VERSION_PATCH=$HIP_VERSION_GITDATE" ;
$HIPCXXFLAGS .= " -isystem $HIP_INCLUDE_PATH" ;
my $compileOnly = 0;
my $needCXXFLAGS = 0; # need to add CXX flags to compile step
projects/clr/hipamd/bin/hipconfig
+1 -1
Näytä tiedosto
@@ -1,7 +1,7 @@
#!/usr/bin/perl -w
$HIP_BASE_VERSION_MAJOR = "3";
$HIP_BASE_VERSION_MINOR = "0";
$HIP_BASE_VERSION_MINOR = "1";
# Need perl > 5.10 to use logic-defined or
use 5.006; use v5.10.1;
projects/clr/hipamd/bin/hipify-perl
+122 -10
Näytä tiedosto
@@ -199,9 +199,9 @@ sub simpleSubstitutions {
$ft{'memory'} += s/\bcuIpcGetMemHandle\b/hipIpcGetMemHandle/g;
$ft{'memory'} += s/\bcuIpcOpenMemHandle\b/hipIpcOpenMemHandle/g;
$ft{'memory'} += s/\bcuMemAlloc\b/hipMalloc/g;
$ft{'memory'} += s/\bcuMemAllocHost\b/hipMemAllocHost/g;
$ft{'memory'} += s/\bcuMemAllocHost_v2\b/hipMemAllocHost/g;
$ft{'memory'} += s/\bcuMemAllocManaged\b/hipMemAllocManaged/g;
$ft{'memory'} += s/\bcuMemAllocHost\b/hipHostMalloc/g;
$ft{'memory'} += s/\bcuMemAllocHost_v2\b/hipHostMalloc/g;
$ft{'memory'} += s/\bcuMemAllocManaged\b/hipMallocManaged/g;
$ft{'memory'} += s/\bcuMemAllocPitch\b/hipMemAllocPitch/g;
$ft{'memory'} += s/\bcuMemAllocPitch_v2\b/hipMemAllocPitch/g;
$ft{'memory'} += s/\bcuMemAlloc_v2\b/hipMalloc/g;
@@ -215,7 +215,7 @@ sub simpleSubstitutions {
$ft{'memory'} += s/\bcuMemHostAlloc\b/hipHostMalloc/g;
$ft{'memory'} += s/\bcuMemHostGetDevicePointer\b/hipHostGetDevicePointer/g;
$ft{'memory'} += s/\bcuMemHostGetDevicePointer_v2\b/hipHostGetDevicePointer/g;
$ft{'memory'} += s/\bcuMemHostGetFlags\b/hipMemHostGetFlags/g;
$ft{'memory'} += s/\bcuMemHostGetFlags\b/hipHostGetFlags/g;
$ft{'memory'} += s/\bcuMemHostRegister\b/hipHostRegister/g;
$ft{'memory'} += s/\bcuMemHostRegister_v2\b/hipHostRegister/g;
$ft{'memory'} += s/\bcuMemHostUnregister\b/hipHostUnregister/g;
@@ -287,6 +287,8 @@ sub simpleSubstitutions {
$ft{'memory'} += s/\bcudaMemset\b/hipMemset/g;
$ft{'memory'} += s/\bcudaMemset2D\b/hipMemset2D/g;
$ft{'memory'} += s/\bcudaMemset2DAsync\b/hipMemset2DAsync/g;
$ft{'memory'} += s/\bcudaMemset3D\b/hipMemset3D/g;
$ft{'memory'} += s/\bcudaMemset3DAsync\b/hipMemset3DAsync/g;
$ft{'memory'} += s/\bcudaMemsetAsync\b/hipMemsetAsync/g;
$ft{'memory'} += s/\bmake_cudaExtent\b/make_hipExtent/g;
$ft{'memory'} += s/\bmake_cudaPitchedPtr\b/make_hipPitchedPtr/g;
@@ -407,8 +409,33 @@ sub simpleSubstitutions {
$ft{'complex'} += s/\bmake_cuComplex\b/make_hipComplex/g;
$ft{'complex'} += s/\bmake_cuDoubleComplex\b/make_hipDoubleComplex/g;
$ft{'complex'} += s/\bmake_cuFloatComplex\b/make_hipFloatComplex/g;
$ft{'library'} += s/\bcublasCaxpy\b/hipblasCaxpy/g;
$ft{'library'} += s/\bcublasCaxpy_v2\b/hipblasCaxpy/g;
$ft{'library'} += s/\bcublasCcopy\b/hipblasCcopy/g;
$ft{'library'} += s/\bcublasCcopy_v2\b/hipblasCcopy/g;
$ft{'library'} += s/\bcublasCdotc\b/hipblasCdotc/g;
$ft{'library'} += s/\bcublasCdotc_v2\b/hipblasCdotc/g;
$ft{'library'} += s/\bcublasCdotu\b/hipblasCdotu/g;
$ft{'library'} += s/\bcublasCdotu_v2\b/hipblasCdotu/g;
$ft{'library'} += s/\bcublasCgemm\b/hipblasCgemm/g;
$ft{'library'} += s/\bcublasCgemmBatched\b/hipblasCgemmBatched/g;
$ft{'library'} += s/\bcublasCgemmStridedBatched\b/hipblasCgemmStridedBatched/g;
$ft{'library'} += s/\bcublasCgemv\b/hipblasCgemv/g;
$ft{'library'} += s/\bcublasCgemv_v2\b/hipblasCgemv/g;
$ft{'library'} += s/\bcublasCreate\b/hipblasCreate/g;
$ft{'library'} += s/\bcublasCreate_v2\b/hipblasCreate/g;
$ft{'library'} += s/\bcublasCrot\b/hipblasCrot/g;
$ft{'library'} += s/\bcublasCrot_v2\b/hipblasCrot/g;
$ft{'library'} += s/\bcublasCrotg\b/hipblasCrotg/g;
$ft{'library'} += s/\bcublasCrotg_v2\b/hipblasCrotg/g;
$ft{'library'} += s/\bcublasCscal\b/hipblasCscal/g;
$ft{'library'} += s/\bcublasCscal_v2\b/hipblasCscal/g;
$ft{'library'} += s/\bcublasCsrot\b/hipblasCsrot/g;
$ft{'library'} += s/\bcublasCsrot_v2\b/hipblasCsrot/g;
$ft{'library'} += s/\bcublasCsscal\b/hipblasCsscal/g;
$ft{'library'} += s/\bcublasCsscal_v2\b/hipblasCsscal/g;
$ft{'library'} += s/\bcublasCswap\b/hipblasCswap/g;
$ft{'library'} += s/\bcublasCswap_v2\b/hipblasCswap/g;
$ft{'library'} += s/\bcublasDasum\b/hipblasDasum/g;
$ft{'library'} += s/\bcublasDasum_v2\b/hipblasDasum/g;
$ft{'library'} += s/\bcublasDaxpy\b/hipblasDaxpy/g;
@@ -430,10 +457,28 @@ sub simpleSubstitutions {
$ft{'library'} += s/\bcublasDger_v2\b/hipblasDger/g;
$ft{'library'} += s/\bcublasDnrm2\b/hipblasDnrm2/g;
$ft{'library'} += s/\bcublasDnrm2_v2\b/hipblasDnrm2/g;
$ft{'library'} += s/\bcublasDrot\b/hipblasDrot/g;
$ft{'library'} += s/\bcublasDrot_v2\b/hipblasDrot/g;
$ft{'library'} += s/\bcublasDrotg\b/hipblasDrotg/g;
$ft{'library'} += s/\bcublasDrotg_v2\b/hipblasDrotg/g;
$ft{'library'} += s/\bcublasDrotm\b/hipblasDrotm/g;
$ft{'library'} += s/\bcublasDrotm_v2\b/hipblasDrotm/g;
$ft{'library'} += s/\bcublasDrotmg\b/hipblasDrotmg/g;
$ft{'library'} += s/\bcublasDrotmg_v2\b/hipblasDrotmg/g;
$ft{'library'} += s/\bcublasDscal\b/hipblasDscal/g;
$ft{'library'} += s/\bcublasDscal_v2\b/hipblasDscal/g;
$ft{'library'} += s/\bcublasDswap\b/hipblasDswap/g;
$ft{'library'} += s/\bcublasDswap_v2\b/hipblasDswap/g;
$ft{'library'} += s/\bcublasDsyr\b/hipblasDsyr/g;
$ft{'library'} += s/\bcublasDsyr_v2\b/hipblasDsyr/g;
$ft{'library'} += s/\bcublasDtrsm\b/hipblasDtrsm/g;
$ft{'library'} += s/\bcublasDtrsm_v2\b/hipblasDtrsm/g;
$ft{'library'} += s/\bcublasDtrsv\b/hipblasDtrsv/g;
$ft{'library'} += s/\bcublasDtrsv_v2\b/hipblasDtrsv/g;
$ft{'library'} += s/\bcublasDzasum\b/hipblasDzasum/g;
$ft{'library'} += s/\bcublasDzasum_v2\b/hipblasDzasum/g;
$ft{'library'} += s/\bcublasDznrm2\b/hipblasDznrm2/g;
$ft{'library'} += s/\bcublasDznrm2_v2\b/hipblasDznrm2/g;
$ft{'library'} += s/\bcublasGemmEx\b/hipblasGemmEx/g;
$ft{'library'} += s/\bcublasGetMatrix\b/hipblasGetMatrix/g;
$ft{'library'} += s/\bcublasGetPointerMode\b/hipblasGetPointerMode/g;
@@ -442,14 +487,32 @@ sub simpleSubstitutions {
$ft{'library'} += s/\bcublasGetStream_v2\b/hipblasGetStream/g;
$ft{'library'} += s/\bcublasGetVector\b/hipblasGetVector/g;
$ft{'library'} += s/\bcublasHgemm\b/hipblasHgemm/g;
$ft{'library'} += s/\bcublasHgemmBatched\b/hipblasHgemmBatched/g;
$ft{'library'} += s/\bcublasHgemmStridedBatched\b/hipblasHgemmStridedBatched/g;
$ft{'library'} += s/\bcublasIcamax\b/hipblasIcamax/g;
$ft{'library'} += s/\bcublasIcamax_v2\b/hipblasIcamax/g;
$ft{'library'} += s/\bcublasIcamin\b/hipblasIcamin/g;
$ft{'library'} += s/\bcublasIcamin_v2\b/hipblasIcamin/g;
$ft{'library'} += s/\bcublasIdamax\b/hipblasIdamax/g;
$ft{'library'} += s/\bcublasIdamax_v2\b/hipblasIdamax/g;
$ft{'library'} += s/\bcublasIdamin\b/hipblasIdamin/g;
$ft{'library'} += s/\bcublasIdamin_v2\b/hipblasIdamin/g;
$ft{'library'} += s/\bcublasIsamax\b/hipblasIsamax/g;
$ft{'library'} += s/\bcublasIsamax_v2\b/hipblasIsamax/g;
$ft{'library'} += s/\bcublasIsamin\b/hipblasIsamin/g;
$ft{'library'} += s/\bcublasIsamin_v2\b/hipblasIsamin/g;
$ft{'library'} += s/\bcublasIzamax\b/hipblasIzamax/g;
$ft{'library'} += s/\bcublasIzamax_v2\b/hipblasIzamax/g;
$ft{'library'} += s/\bcublasIzamin\b/hipblasIzamin/g;
$ft{'library'} += s/\bcublasIzamin_v2\b/hipblasIzamin/g;
$ft{'library'} += s/\bcublasSasum\b/hipblasSasum/g;
$ft{'library'} += s/\bcublasSasum_v2\b/hipblasSasum/g;
$ft{'library'} += s/\bcublasSaxpy\b/hipblasSaxpy/g;
$ft{'library'} += s/\bcublasSaxpy_v2\b/hipblasSaxpy/g;
$ft{'library'} += s/\bcublasScasum\b/hipblasScasum/g;
$ft{'library'} += s/\bcublasScasum_v2\b/hipblasScasum/g;
$ft{'library'} += s/\bcublasScnrm2\b/hipblasScnrm2/g;
$ft{'library'} += s/\bcublasScnrm2_v2\b/hipblasScnrm2/g;
$ft{'library'} += s/\bcublasScopy\b/hipblasScopy/g;
$ft{'library'} += s/\bcublasScopy_v2\b/hipblasScopy/g;
$ft{'library'} += s/\bcublasSdot\b/hipblasSdot/g;
@@ -466,15 +529,56 @@ sub simpleSubstitutions {
$ft{'library'} += s/\bcublasSgemmStridedBatched\b/hipblasSgemmStridedBatched/g;
$ft{'library'} += s/\bcublasSgemm_v2\b/hipblasSgemm/g;
$ft{'library'} += s/\bcublasSgemv\b/hipblasSgemv/g;
$ft{'library'} += s/\bcublasSgemvBatched\b/hipblasSgemvBatched/g;
$ft{'library'} += s/\bcublasSgemv_v2\b/hipblasSgemv/g;
$ft{'library'} += s/\bcublasSger\b/hipblasSger/g;
$ft{'library'} += s/\bcublasSger_v2\b/hipblasSger/g;
$ft{'library'} += s/\bcublasSnrm2\b/hipblasSnrm2/g;
$ft{'library'} += s/\bcublasSnrm2_v2\b/hipblasSnrm2/g;
$ft{'library'} += s/\bcublasSrot\b/hipblasSrot/g;
$ft{'library'} += s/\bcublasSrot_v2\b/hipblasSrot/g;
$ft{'library'} += s/\bcublasSrotg\b/hipblasSrotg/g;
$ft{'library'} += s/\bcublasSrotg_v2\b/hipblasSrotg/g;
$ft{'library'} += s/\bcublasSrotm\b/hipblasSrotm/g;
$ft{'library'} += s/\bcublasSrotm_v2\b/hipblasSrotm/g;
$ft{'library'} += s/\bcublasSrotmg\b/hipblasSrotmg/g;
$ft{'library'} += s/\bcublasSrotmg_v2\b/hipblasSrotmg/g;
$ft{'library'} += s/\bcublasSscal\b/hipblasSscal/g;
$ft{'library'} += s/\bcublasSscal_v2\b/hipblasSscal/g;
$ft{'library'} += s/\bcublasSswap\b/hipblasSswap/g;
$ft{'library'} += s/\bcublasSswap_v2\b/hipblasSswap/g;
$ft{'library'} += s/\bcublasSsyr\b/hipblasSsyr/g;
$ft{'library'} += s/\bcublasSsyr_v2\b/hipblasSsyr/g;
$ft{'library'} += s/\bcublasStrsm\b/hipblasStrsm/g;
$ft{'library'} += s/\bcublasStrsm_v2\b/hipblasStrsm/g;
$ft{'library'} += s/\bcublasStrsv\b/hipblasStrsv/g;
$ft{'library'} += s/\bcublasStrsv_v2\b/hipblasStrsv/g;
$ft{'library'} += s/\bcublasZaxpy\b/hipblasZaxpy/g;
$ft{'library'} += s/\bcublasZaxpy_v2\b/hipblasZaxpy/g;
$ft{'library'} += s/\bcublasZcopy\b/hipblasZcopy/g;
$ft{'library'} += s/\bcublasZcopy_v2\b/hipblasZcopy/g;
$ft{'library'} += s/\bcublasZdotc\b/hipblasZdotc/g;
$ft{'library'} += s/\bcublasZdotc_v2\b/hipblasZdotc/g;
$ft{'library'} += s/\bcublasZdotu\b/hipblasZdotu/g;
$ft{'library'} += s/\bcublasZdotu_v2\b/hipblasZdotu/g;
$ft{'library'} += s/\bcublasZdrot\b/hipblasZdrot/g;
$ft{'library'} += s/\bcublasZdrot_v2\b/hipblasZdrot/g;
$ft{'library'} += s/\bcublasZdscal\b/hipblasZdscal/g;
$ft{'library'} += s/\bcublasZdscal_v2\b/hipblasZdscal/g;
$ft{'library'} += s/\bcublasZgemm\b/hipblasZgemm/g;
$ft{'library'} += s/\bcublasZgemmBatched\b/hipblasZgemmBatched/g;
$ft{'library'} += s/\bcublasZgemmStridedBatched\b/hipblasZgemmStridedBatched/g;
$ft{'library'} += s/\bcublasZgemm_v2\b/hipblasZgemm/g;
$ft{'library'} += s/\bcublasZgemv\b/hipblasZgemv/g;
$ft{'library'} += s/\bcublasZgemv_v2\b/hipblasZgemv/g;
$ft{'library'} += s/\bcublasZrot\b/hipblasZrot/g;
$ft{'library'} += s/\bcublasZrot_v2\b/hipblasZrot/g;
$ft{'library'} += s/\bcublasZrotg\b/hipblasZrotg/g;
$ft{'library'} += s/\bcublasZrotg_v2\b/hipblasZrotg/g;
$ft{'library'} += s/\bcublasZscal\b/hipblasZscal/g;
$ft{'library'} += s/\bcublasZscal_v2\b/hipblasZscal/g;
$ft{'library'} += s/\bcublasZswap\b/hipblasZswap/g;
$ft{'library'} += s/\bcublasZswap_v2\b/hipblasZswap/g;
$ft{'library'} += s/\bcuda_stream\b/hip_stream/g;
$ft{'library'} += s/\bcudnnActivationBackward\b/hipdnnActivationBackward/g;
$ft{'library'} += s/\bcudnnActivationForward\b/hipdnnActivationForward/g;
@@ -1011,7 +1115,11 @@ sub simpleSubstitutions {
$ft{'numeric_literal'} += s/\bCUBLAS_STATUS_SUCCESS\b/HIPBLAS_STATUS_SUCCESS/g;
$ft{'numeric_literal'} += s/\bCUDA_C_16F\b/HIPBLAS_C_16F/g;
$ft{'numeric_literal'} += s/\bCUDA_C_32F\b/HIPBLAS_C_32F/g;
$ft{'numeric_literal'} += s/\bCUDA_C_32I\b/HIPBLAS_C_32I/g;
$ft{'numeric_literal'} += s/\bCUDA_C_32U\b/HIPBLAS_C_32U/g;
$ft{'numeric_literal'} += s/\bCUDA_C_64F\b/HIPBLAS_C_64F/g;
$ft{'numeric_literal'} += s/\bCUDA_C_8I\b/HIPBLAS_C_8I/g;
$ft{'numeric_literal'} += s/\bCUDA_C_8U\b/HIPBLAS_C_8U/g;
$ft{'numeric_literal'} += s/\bCUDA_ERROR_ALREADY_ACQUIRED\b/hipErrorAlreadyAcquired/g;
$ft{'numeric_literal'} += s/\bCUDA_ERROR_ALREADY_MAPPED\b/hipErrorAlreadyMapped/g;
$ft{'numeric_literal'} += s/\bCUDA_ERROR_ARRAY_IS_MAPPED\b/hipErrorArrayIsMapped/g;
@@ -1062,7 +1170,11 @@ sub simpleSubstitutions {
$ft{'numeric_literal'} += s/\bCUDA_ERROR_UNSUPPORTED_LIMIT\b/hipErrorUnsupportedLimit/g;
$ft{'numeric_literal'} += s/\bCUDA_R_16F\b/HIPBLAS_R_16F/g;
$ft{'numeric_literal'} += s/\bCUDA_R_32F\b/HIPBLAS_R_32F/g;
$ft{'numeric_literal'} += s/\bCUDA_R_32I\b/HIPBLAS_R_32I/g;
$ft{'numeric_literal'} += s/\bCUDA_R_32U\b/HIPBLAS_R_32U/g;
$ft{'numeric_literal'} += s/\bCUDA_R_64F\b/HIPBLAS_R_64F/g;
$ft{'numeric_literal'} += s/\bCUDA_R_8I\b/HIPBLAS_R_8I/g;
$ft{'numeric_literal'} += s/\bCUDA_R_8U\b/HIPBLAS_R_8U/g;
$ft{'numeric_literal'} += s/\bCUDA_SUCCESS\b/hipSuccess/g;
$ft{'numeric_literal'} += s/\bCUDNN_16BIT_INDICES\b/HIPDNN_16BIT_INDICES/g;
$ft{'numeric_literal'} += s/\bCUDNN_32BIT_INDICES\b/HIPDNN_32BIT_INDICES/g;
@@ -1487,7 +1599,7 @@ sub simpleSubstitutions {
$ft{'numeric_literal'} += s/\bcudaErrorHostMemoryAlreadyRegistered\b/hipErrorHostMemoryAlreadyRegistered/g;
$ft{'numeric_literal'} += s/\bcudaErrorHostMemoryNotRegistered\b/hipErrorHostMemoryNotRegistered/g;
$ft{'numeric_literal'} += s/\bcudaErrorIllegalAddress\b/hipErrorIllegalAddress/g;
$ft{'numeric_literal'} += s/\bcudaErrorInitializationError\b/hipErrorInitializationError/g;
$ft{'numeric_literal'} += s/\bcudaErrorInitializationError\b/hipErrorNotInitialized/g;
$ft{'numeric_literal'} += s/\bcudaErrorInsufficientDriver\b/hipErrorInsufficientDriver/g;
$ft{'numeric_literal'} += s/\bcudaErrorInvalidConfiguration\b/hipErrorInvalidConfiguration/g;
$ft{'numeric_literal'} += s/\bcudaErrorInvalidDevice\b/hipErrorInvalidDevice/g;
@@ -1497,15 +1609,15 @@ sub simpleSubstitutions {
$ft{'numeric_literal'} += s/\bcudaErrorInvalidKernelImage\b/hipErrorInvalidImage/g;
$ft{'numeric_literal'} += s/\bcudaErrorInvalidMemcpyDirection\b/hipErrorInvalidMemcpyDirection/g;
$ft{'numeric_literal'} += s/\bcudaErrorInvalidPtx\b/hipErrorInvalidKernelFile/g;
$ft{'numeric_literal'} += s/\bcudaErrorInvalidResourceHandle\b/hipErrorInvalidResourceHandle/g;
$ft{'numeric_literal'} += s/\bcudaErrorInvalidResourceHandle\b/hipErrorInvalidHandle/g;
$ft{'numeric_literal'} += s/\bcudaErrorInvalidSource\b/hipErrorInvalidSource/g;
$ft{'numeric_literal'} += s/\bcudaErrorInvalidSymbol\b/hipErrorInvalidSymbol/g;
$ft{'numeric_literal'} += s/\bcudaErrorInvalidValue\b/hipErrorInvalidValue/g;
$ft{'numeric_literal'} += s/\bcudaErrorLaunchFailure\b/hipErrorLaunchFailure/g;
$ft{'numeric_literal'} += s/\bcudaErrorLaunchOutOfResources\b/hipErrorLaunchOutOfResources/g;
$ft{'numeric_literal'} += s/\bcudaErrorLaunchTimeout\b/hipErrorLaunchTimeOut/g;
$ft{'numeric_literal'} += s/\bcudaErrorMapBufferObjectFailed\b/hipErrorMapBufferObjectFailed/g;
$ft{'numeric_literal'} += s/\bcudaErrorMemoryAllocation\b/hipErrorMemoryAllocation/g;
$ft{'numeric_literal'} += s/\bcudaErrorMapBufferObjectFailed\b/hipErrorMapFailed/g;
$ft{'numeric_literal'} += s/\bcudaErrorMemoryAllocation\b/hipErrorOutOfMemory/g;
$ft{'numeric_literal'} += s/\bcudaErrorMissingConfiguration\b/hipErrorMissingConfiguration/g;
$ft{'numeric_literal'} += s/\bcudaErrorNoDevice\b/hipErrorNoDevice/g;
$ft{'numeric_literal'} += s/\bcudaErrorNoKernelImageForDevice\b/hipErrorNoBinaryForGpu/g;
@@ -1596,7 +1708,7 @@ sub simpleSubstitutions {
$ft{'define'} += s/\bCU_LAUNCH_PARAM_END\b/HIP_LAUNCH_PARAM_END/g;
$ft{'define'} += s/\bCU_MEMHOSTALLOC_DEVICEMAP\b/hipHostMallocMapped/g;
$ft{'define'} += s/\bCU_MEMHOSTALLOC_PORTABLE\b/hipHostMallocPortable/g;
$ft{'define'} += s/\bCU_MEMHOSTALLOC_WRITECOMBINED\b/hipHostAllocWriteCombined/g;
$ft{'define'} += s/\bCU_MEMHOSTALLOC_WRITECOMBINED\b/hipHostMallocWriteCombined/g;
$ft{'define'} += s/\bCU_MEMHOSTREGISTER_DEVICEMAP\b/hipHostRegisterMapped/g;
$ft{'define'} += s/\bCU_MEMHOSTREGISTER_IOMEMORY\b/hipHostRegisterIoMemory/g;
$ft{'define'} += s/\bCU_MEMHOSTREGISTER_PORTABLE\b/hipHostRegisterPortable/g;
@@ -1626,7 +1738,7 @@ sub simpleSubstitutions {
$ft{'define'} += s/\bcudaHostAllocDefault\b/hipHostMallocDefault/g;
$ft{'define'} += s/\bcudaHostAllocMapped\b/hipHostMallocMapped/g;
$ft{'define'} += s/\bcudaHostAllocPortable\b/hipHostMallocPortable/g;
$ft{'define'} += s/\bcudaHostAllocWriteCombined\b/hipHostAllocWriteCombined/g;
$ft{'define'} += s/\bcudaHostAllocWriteCombined\b/hipHostMallocWriteCombined/g;
$ft{'define'} += s/\bcudaHostRegisterDefault\b/hipHostRegisterDefault/g;
$ft{'define'} += s/\bcudaHostRegisterIoMemory\b/hipHostRegisterIoMemory/g;
$ft{'define'} += s/\bcudaHostRegisterMapped\b/hipHostRegisterMapped/g;
projects/clr/hipamd/docs/markdown/CUBLAS_API_supported_by_HIP.md
+101 -98
Näytä tiedosto
@@ -137,128 +137,131 @@
|`cublasSnrm2_v2` |`hipblasSnrm2` |
|`cublasDnrm2` |`hipblasDnrm2` |
|`cublasDnrm2_v2` |`hipblasDnrm2` |
|`cublasScnrm2` | |
|`cublasScnrm2_v2` | |
|`cublasDznrm2` | |
|`cublasDznrm2_v2` | |
|`cublasScnrm2` |`hipblasScnrm2` |
|`cublasScnrm2_v2` |`hipblasScnrm2` |
|`cublasDznrm2` |`hipblasDznrm2` |
|`cublasDznrm2_v2` |`hipblasDznrm2` |
|`cublasDotEx` | | 8.0 |
|`cublasDotcEx` | | 8.0 |
|`cublasSdot` |`hipblasSdot` |
|`cublasSdot_v2` |`hipblasSdot` |
|`cublasDdot` |`hipblasDdot` |
|`cublasDdot_v2` |`hipblasDdot` |
|`cublasCdotu` | |
|`cublasCdotu_v2` | |
|`cublasCdotc` | |
|`cublasCdotc_v2` | |
|`cublasZdotu` | |
|`cublasZdotu_v2` | |
|`cublasZdotc` | |
|`cublasZdotc_v2` | |
|`cublasCdotu` |`hipblasCdotu` |
|`cublasCdotu_v2` |`hipblasCdotu` |
|`cublasCdotc` |`hipblasCdotc` |
|`cublasCdotc_v2` |`hipblasCdotc` |
|`cublasZdotu` |`hipblasZdotu` |
|`cublasZdotu_v2` |`hipblasZdotu` |
|`cublasZdotc` |`hipblasZdotc` |
|`cublasZdotc_v2` |`hipblasZdotc` |
|`cublasScalEx` | | 8.0 |
|`cublasSscal` |`hipblasSscal` |
|`cublasSscal_v2` |`hipblasSscal` |
|`cublasDscal` |`hipblasDscal` |
|`cublasDscal_v2` |`hipblasDscal` |
|`cublasCscal` | |
|`cublasCscal_v2` | |
|`cublasCsscal` | |
|`cublasCsscal_v2` | |
|`cublasZscal` | |
|`cublasZscal_v2` | |
|`cublasZdscal` | |
|`cublasZdscal_v2` | |
|`cublasCscal` |`hipblasCscal` |
|`cublasCscal_v2` |`hipblasCscal` |
|`cublasCsscal` |`hipblasCsscal` |
|`cublasCsscal_v2` |`hipblasCsscal` |
|`cublasZscal` |`hipblasZscal` |
|`cublasZscal_v2` |`hipblasZscal` |
|`cublasZdscal` |`hipblasZdscal` |
|`cublasZdscal_v2` |`hipblasZdscal` |
|`cublasAxpyEx` | | 8.0 |
|`cublasSaxpy` |`hipblasSaxpy` |
|`cublasSaxpy_v2` |`hipblasSaxpy` |
|`cublasDaxpy` |`hipblasDaxpy` |
|`cublasDaxpy_v2` |`hipblasDaxpy` |
|`cublasCaxpy` | |
|`cublasCaxpy_v2` | |
|`cublasZaxpy` | |
|`cublasZaxpy_v2` | |
|`cublasCaxpy` |`hipblasCaxpy` |
|`cublasCaxpy_v2` |`hipblasCaxpy` |
|`cublasZaxpy` |`hipblasZaxpy` |
|`cublasZaxpy_v2` |`hipblasZaxpy` |
|`cublasScopy` |`hipblasScopy` |
|`cublasScopy_v2` |`hipblasScopy` |
|`cublasDcopy` |`hipblasDcopy` |
|`cublasDcopy_v2` |`hipblasDcopy` |
|`cublasCcopy` | |
|`cublasCcopy` |`hipblasCcopy` |
|`cublasCopyEx` | | 10.1 |
|`cublasCcopy_v2` | |
|`cublasZcopy` | |
|`cublasZcopy_v2` | |
|`cublasSswap` | |
|`cublasSswap_v2` | |
|`cublasDswap` | |
|`cublasDswap_v2` | |
|`cublasCswap` | |
|`cublasCswap_v2` | |
|`cublasZswap` | |
|`cublasZswap_v2` | |
|`cublasCcopy_v2` |`hipblasCcopy` |
|`cublasZcopy` |`hipblasZcopy` |
|`cublasZcopy_v2` |`hipblasZcopy` |
|`cublasSswap` |`hipblasSswap` |
|`cublasSswap_v2` |`hipblasSswap` |
|`cublasDswap` |`hipblasDswap` |
|`cublasDswap_v2` |`hipblasDswap` |
|`cublasCswap` |`hipblasCswap` |
|`cublasCswap_v2` |`hipblasCswap` |
|`cublasZswap` |`hipblasZswap` |
|`cublasZswap_v2` |`hipblasZswap` |
|`cublasIamaxEx` | | 10.1 |
|`cublasIsamax` |`hipblasIsamax` |
|`cublasIsamax_v2` |`hipblasIsamax` |
|`cublasIdamax` |`hipblasIdamax` |
|`cublasIdamax_v2` |`hipblasIdamax` |
|`cublasIcamax` | |
|`cublasIcamax_v2` | |
|`cublasIzamax` | |
|`cublasIzamax_v2` | |
|`cublasIcamax` |`hipblasIcamax` |
|`cublasIcamax_v2` |`hipblasIcamax` |
|`cublasIzamax` |`hipblasIzamax` |
|`cublasIzamax_v2` |`hipblasIzamax` |
|`cublasIaminEx` | | 10.1 |
|`cublasIsamin` | |
|`cublasIsamin_v2` | |
|`cublasIdamin` | |
|`cublasIdamin_v2` | |
|`cublasIcamin` | |
|`cublasIcamin_v2` | |
|`cublasIzamin` | |
|`cublasIzamin_v2` | |
|`cublasIsamin` |`hipblasIsamin` |
|`cublasIsamin_v2` |`hipblasIsamin` |
|`cublasIdamin` |`hipblasIdamin` |
|`cublasIdamin_v2` |`hipblasIdamin` |
|`cublasIcamin` |`hipblasIcamin` |
|`cublasIcamin_v2` |`hipblasIcamin` |
|`cublasIzamin` |`hipblasIzamin` |
|`cublasIzamin_v2` |`hipblasIzamin` |
|`cublasAsumEx` | | 10.1 |
|`cublasSasum` |`hipblasSasum` |
|`cublasSasum_v2` |`hipblasSasum` |
|`cublasDasum` |`hipblasDasum` |
|`cublasDasum_v2` |`hipblasDasum` |
|`cublasScasum` | |
|`cublasScasum_v2` | |
|`cublasDzasum` | |
|`cublasDzasum_v2` | |
|`cublasScasum` |`hipblasScasum` |
|`cublasScasum_v2` |`hipblasScasum` |
|`cublasDzasum` |`hipblasDzasum` |
|`cublasDzasum_v2` |`hipblasDzasum` |
|`cublasRotEx` | | 10.1 |
|`cublasSrot` | |
|`cublasSrot_v2` | |
|`cublasDrot` | |
|`cublasDrot_v2` | |
|`cublasCrot` | |
|`cublasCrot_v2` | |
|`cublasZrot` | |
|`cublasZrot_v2` | |
|`cublasSrot` |`hipblasSrot` |
|`cublasSrot_v2` |`hipblasSrot` |
|`cublasDrot` |`hipblasDrot` |
|`cublasDrot_v2` |`hipblasDrot` |
|`cublasCrot` |`hipblasCrot` |
|`cublasCrot_v2` |`hipblasCrot` |
|`cublasCsrot` |`hipblasCsrot` |
|`cublasCsrot_v2` |`hipblasCsrot` |
|`cublasZrot` |`hipblasZrot` |
|`cublasZrot_v2` |`hipblasZrot` |
|`cublasRotgEx` | | 10.1 |
|`cublasZdrot` | |
|`cublasZdrot_v2` | |
|`cublasSrotg` | |
|`cublasSrotg_v2` | |
|`cublasDrotg` | |
|`cublasDrotg_v2` | |
|`cublasCrotg` | |
|`cublasCrotg_v2` | |
|`cublasZrotg` | |
|`cublasZrotg_v2` | |
|`cublasZdrot` |`hipblasZdrot` |
|`cublasZdrot_v2` |`hipblasZdrot` |
|`cublasSrotg` |`hipblasSrotg` |
|`cublasSrotg_v2` |`hipblasSrotg` |
|`cublasDrotg` |`hipblasDrotg` |
|`cublasDrotg_v2` |`hipblasDrotg` |
|`cublasCrotg` |`hipblasCrotg` |
|`cublasCrotg_v2` |`hipblasCrotg` |
|`cublasZrotg` |`hipblasZrotg` |
|`cublasZrotg_v2` |`hipblasZrotg` |
|`cublasRotmEx` | | 10.1 |
|`cublasSrotm` | |
|`cublasSrotm_v2` | |
|`cublasDrotm` | |
|`cublasDrotm_v2` | |
|`cublasSrotm` |`hipblasSrotm` |
|`cublasSrotm_v2` |`hipblasSrotm` |
|`cublasDrotm` |`hipblasDrotm` |
|`cublasDrotm_v2` |`hipblasDrotm` |
|`cublasRotmgEx` | | 10.1 |
|`cublasSrotmg` | |
|`cublasSrotmg_v2` | |
|`cublasDrotmg` | |
|`cublasDrotmg_v2` | |
|`cublasSrotmg` |`hipblasSrotmg` |
|`cublasSrotmg_v2` |`hipblasSrotmg` |
|`cublasDrotmg` |`hipblasDrotmg` |
|`cublasDrotmg_v2` |`hipblasDrotmg` |
|`cublasSgemv` |`hipblasSgemv` |
|`cublasSgemv_v2` |`hipblasSgemv` |
|`cublasSgemvBatched` |`hipblasSgemvBatched` |
|`cublasDgemv` |`hipblasDgemv` |
|`cublasDgemv_v2` |`hipblasDgemv` |
|`cublasCgemv` | |
|`cublasCgemv_v2` | |
|`cublasZgemv` | |
|`cublasZgemv_v2` | |
|`cublasCgemv` |`hipblasCgemv` |
|`cublasCgemv_v2` |`hipblasCgemv` |
|`cublasZgemv` |`hipblasZgemv` |
|`cublasZgemv_v2` |`hipblasZgemv` |
|`cublasSgbmv` | |
|`cublasSgbmv_v2` | |
|`cublasDgbmv` | |
@@ -291,10 +294,10 @@
|`cublasCtpmv_v2` | |
|`cublasZtpmv` | |
|`cublasZtpmv_v2` | |
|`cublasStrsv` | |
|`cublasStrsv_v2` | |
|`cublasDtrsv` | |
|`cublasDtrsv_v2` | |
|`cublasStrsv` |`hipblasStrsv` |
|`cublasStrsv_v2` |`hipblasStrsv` |
|`cublasDtrsv` |`hipblasDtrsv` |
|`cublasDtrsv_v2` |`hipblasDtrsv` |
|`cublasCtrsv` | |
|`cublasCtrsv_v2` | |
|`cublasZtrsv` | |
@@ -355,10 +358,10 @@
|`cublasZgeru_v2` | |
|`cublasZgerc` | |
|`cublasZgerc_v2` | |
|`cublasSsyr` | |
|`cublasSsyr_v2` | |
|`cublasDsyr` | |
|`cublasDsyr_v2` | |
|`cublasSsyr` |`hipblasSsyr` |
|`cublasSsyr_v2` |`hipblasSsyr` |
|`cublasDsyr` |`hipblasDsyr` |
|`cublasDsyr_v2` |`hipblasDsyr` |
|`cublasCsyr` | |
|`cublasCsyr_v2` | |
|`cublasZsyr` | |
@@ -403,8 +406,8 @@
|`cublasCgemm_v2` | |
|`cublasCgemm3m` | | 8.0 |
|`cublasCgemm3mEx` | | 8.0 |
|`cublasZgemm` | |
|`cublasZgemm_v2` | |
|`cublasZgemm` |`hipblasZgemm` |
|`cublasZgemm_v2` |`hipblasZgemm` |
|`cublasZgemm3m` | | 8.0 |
|`cublasHgemm` |`hipblasHgemm` | 7.5 |
|`cublasSgemmEx` | | 7.5 |
@@ -473,20 +476,20 @@
|`cublasCtrmm_v2` | |
|`cublasZtrmm` | |
|`cublasZtrmm_v2` | |
|`cublasHgemmBatched` | | 9.0 |
|`cublasHgemmBatched` |`hipblasHgemmBatched` | 9.0 |
|`cublasSgemmBatched` |`hipblasSgemmBatched` |
|`cublasDgemmBatched` |`hipblasDgemmBatched` |
|`cublasCgemmBatched` | |
|`cublasCgemmBatched` |`hipblasCgemmBatched` |
|`cublasCgemm3mBatched` | | 8.0 |
|`cublasZgemmBatched` | |
|`cublasZgemmBatched` |`hipblasZgemmBatched` |
|`cublasGemmBatchedEx` | | 9.1 |
|`cublasGemmStridedBatchedEx` | | 9.1 |
|`cublasSgemmStridedBatched` |`hipblasSgemmStridedBatched` | 8.0 |
|`cublasDgemmStridedBatched` |`hipblasDgemmStridedBatched` | 8.0 |
|`cublasCgemmStridedBatched` | | 8.0 |
|`cublasCgemmStridedBatched` |`hipblasCgemmStridedBatched` | 8.0 |
|`cublasCgemm3mStridedBatched` | | 8.0 |
|`cublasZgemmStridedBatched` | | 8.0 |
|`cublasHgemmStridedBatched` | | 8.0 |
|`cublasZgemmStridedBatched` |`hipblasZgemmStridedBatched` | 8.0 |
|`cublasHgemmStridedBatched` |`hipblasHgemmStridedBatched` | 8.0 |
|`cublasSgeam` |`hipblasSgeam` |
|`cublasDgeam` |`hipblasDgeam` |
|`cublasCgeam` | |
projects/clr/hipamd/docs/markdown/CUDA_Driver_API_functions_supported_by_HIP.md
+4 -4
Näytä tiedosto
@@ -561,7 +561,7 @@
| define |`CU_LAUNCH_PARAM_END` |`HIP_LAUNCH_PARAM_END` |
| define |`CU_MEMHOSTALLOC_DEVICEMAP` |`hipHostMallocMapped` |
| define |`CU_MEMHOSTALLOC_PORTABLE` |`hipHostMallocPortable` |
| define |`CU_MEMHOSTALLOC_WRITECOMBINED` |`hipHostAllocWriteCombined` |
| define |`CU_MEMHOSTALLOC_WRITECOMBINED` |`hipHostMallocWriteCombined` |
| define |`CU_MEMHOSTREGISTER_DEVICEMAP` |`hipHostRegisterMapped` |
| define |`CU_MEMHOSTREGISTER_IOMEMORY` |`hipHostRegisterIoMemory` | 7.5 |
| define |`CU_MEMHOSTREGISTER_PORTABLE` |`hipHostRegisterPortable` |
@@ -916,8 +916,8 @@
| `cuIpcOpenEventHandle` | |
| `cuIpcOpenMemHandle` | `hipIpcOpenMemHandle` |
| `cuMemAlloc` | `hipMalloc` |
| `cuMemAllocHost` | `hipMemAllocHost` |
| `cuMemAllocManaged` | `hipMemAllocManaged` |
| `cuMemAllocHost` | `hipHostMalloc` |
| `cuMemAllocManaged` | `hipMallocManaged` |
| `cuMemAllocPitch` | `hipMemAllocPitch` |
| `cuMemcpy` | |
| `cuMemcpy2D` | `hipMemcpyParam2D` |
@@ -944,7 +944,7 @@
| `cuMemcpyPeer` | |
| `cuMemcpyPeerAsync` | |
| `cuMemFree` | `hipFree` |
| `cuMemFreeHost` | `hipFreeHost` |
| `cuMemFreeHost` | `hipHostFree` |
| `cuMemGetAddressRange` | `hipMemGetAddressRange` |
| `cuMemGetInfo` | `hipMemGetInfo` |
| `cuMemHostAlloc` | `hipHostMalloc` |
projects/clr/hipamd/docs/markdown/CUDA_Runtime_API_functions_supported_by_HIP.md
+37 -37
Näytä tiedosto
@@ -166,7 +166,7 @@
| `cudaMemcpy2DToArray` | `hipMemcpy2DToArray` |
| `cudaMemcpy2DToArrayAsync` | |
| `cudaMemcpy3D` | `hipMemcpy3D` |
| `cudaMemcpy3DAsync` | |
| `cudaMemcpy3DAsync` | `hipMemcpy3DAsync` |
| `cudaMemcpy3DPeer` | |
| `cudaMemcpy3DPeerAsync` | |
| `cudaMemcpyAsync` | `hipMemcpyAsync` |
@@ -179,8 +179,8 @@
| `cudaMemset` | `hipMemset` |
| `cudaMemset2D` | `hipMemset2D` |
| `cudaMemset2DAsync` | `hipMemset2DAsync` |
| `cudaMemset3D` | |
| `cudaMemset3DAsync` | |
| `cudaMemset3D` | `hipMemset3D` |
| `cudaMemset3DAsync` | `hipMemset3DAsync` |
| `cudaMemsetAsync` | `hipMemsetAsync` |
| `make_cudaExtent` | `make_hipExtent` |
| `make_cudaPitchedPtr` | `make_hipPitchedPtr` |
@@ -716,23 +716,23 @@
| enum |***`cudaError`*** | |***`hipError_t`*** |
| typedef |***`cudaError_t`*** | |***`hipError_t`*** |
| 0 |*`cudaSuccess`* | |*`hipSuccess`* |
| 1 |*`cudaErrorInvalidValue`* | |*`hipErrorInvalidValue`* | 1011 |
| 2 |*`cudaErrorMemoryAllocation`* | |*`hipErrorMemoryAllocation`* | 1002 |
| 3 |*`cudaErrorInitializationError`* | |*`hipErrorInitializationError`* | 1003 |
| 1 |*`cudaErrorInvalidValue`* | |*`hipErrorInvalidValue`* |
| 2 |*`cudaErrorMemoryAllocation`* | |*`hipErrorOutOfMemory`* |
| 3 |*`cudaErrorInitializationError`* | |*`hipErrorNotInitialized`* |
| 4 |*`cudaErrorCudartUnloading`* | |*`hipErrorDeinitialized`* |
| 5 |*`cudaErrorProfilerDisabled`* | |*`hipErrorProfilerDisabled`* |
| 6 |*`cudaErrorProfilerNotInitialized`* | |*`hipErrorProfilerNotInitialized`* |
| 7 |*`cudaErrorProfilerAlreadyStarted`* | |*`hipErrorProfilerAlreadyStarted`* |
| 8 |*`cudaErrorProfilerAlreadyStopped`* | |*`hipErrorProfilerAlreadyStopped`* |
| 9 |*`cudaErrorInvalidConfiguration`* | |*`hipErrorInvalidConfiguration`* | 1009 |
| 9 |*`cudaErrorInvalidConfiguration`* | |*`hipErrorInvalidConfiguration`* |
| 12 |*`cudaErrorInvalidPitchValue`* | | |
| 13 |*`cudaErrorInvalidSymbol`* | |*`hipErrorInvalidSymbol`* | 701 |
| 13 |*`cudaErrorInvalidSymbol`* | |*`hipErrorInvalidSymbol`* |
| 16 |*`cudaErrorInvalidHostPointer`* | | |
| 17 |*`cudaErrorInvalidDevicePointer`* | |*`hipErrorInvalidDevicePointer`* | 1017 |
| 17 |*`cudaErrorInvalidDevicePointer`* | |*`hipErrorInvalidDevicePointer`* |
| 18 |*`cudaErrorInvalidTexture`* | | |
| 19 |*`cudaErrorInvalidTextureBinding`* | | |
| 20 |*`cudaErrorInvalidChannelDescriptor`* | | |
| 21 |*`cudaErrorInvalidMemcpyDirection`* | |*`hipErrorInvalidMemcpyDirection`* | 1021 |
| 21 |*`cudaErrorInvalidMemcpyDirection`* | |*`hipErrorInvalidMemcpyDirection`* |
| 22 |*`cudaErrorAddressOfConstant`* | | |
| 23 |*`cudaErrorTextureFetchFailed`* | | |
| 24 |*`cudaErrorTextureNotBound`* | | |
@@ -749,20 +749,20 @@
| 45 |*`cudaErrorDuplicateSurfaceName`* | | |
| 46 |*`cudaErrorDevicesUnavailable`* | | |
| 49 |*`cudaErrorIncompatibleDriverContext`* | | |
| 52 |*`cudaErrorMissingConfiguration`* | |*`hipErrorMissingConfiguration`* | 1001 |
| 53 |*`cudaErrorPriorLaunchFailure`* | |*`hipErrorPriorLaunchFailure`* | 1005 |
| 52 |*`cudaErrorMissingConfiguration`* | |*`hipErrorMissingConfiguration`* |
| 53 |*`cudaErrorPriorLaunchFailure`* | |*`hipErrorPriorLaunchFailure`* |
| 65 |*`cudaErrorLaunchMaxDepthExceeded`* | | |
| 66 |*`cudaErrorLaunchFileScopedTex`* | | |
| 67 |*`cudaErrorLaunchFileScopedSurf`* | | |
| 68 |*`cudaErrorSyncDepthExceeded`* | | |
| 69 |*`cudaErrorLaunchPendingCountExceeded`* | | |
| 98 |*`cudaErrorInvalidDeviceFunction`* | |*`hipErrorInvalidDeviceFunction`* | 1008 |
| 100 |*`cudaErrorNoDevice`* | |*`hipErrorNoDevice`* | 1038 |
| 101 |*`cudaErrorInvalidDevice`* | |*`hipErrorInvalidDevice`* | 1010 |
| 98 |*`cudaErrorInvalidDeviceFunction`* | |*`hipErrorInvalidDeviceFunction`* |
| 100 |*`cudaErrorNoDevice`* | |*`hipErrorNoDevice`* |
| 101 |*`cudaErrorInvalidDevice`* | |*`hipErrorInvalidDevice`* |
| 127 |*`cudaErrorStartupFailure`* | 10.0 | |
| 200 |*`cudaErrorInvalidKernelImage`* | |*`hipErrorInvalidImage`* |
| 201 |*`cudaErrorDeviceUninitilialized`* | |*`hipErrorInvalidContext`* |
| 205 |*`cudaErrorMapBufferObjectFailed`* | |*`hipErrorMapBufferObjectFailed`* | 1071 |
| 205 |*`cudaErrorMapBufferObjectFailed`* | |*`hipErrorMapFailed`* |
| 206 |*`cudaErrorUnmapBufferObjectFailed`* | |*`hipErrorUnmapFailed`* |
| 209 |*`cudaErrorNoKernelImageForDevice`* | |*`hipErrorNoBinaryForGpu`* |
| 214 |*`cudaErrorECCUncorrectable`* | |*`hipErrorECCNotCorrectable`* |
@@ -778,31 +778,31 @@
| 302 |*`cudaErrorSharedObjectSymbolNotFound`* | |*`hipErrorSharedObjectSymbolNotFound`* |
| 303 |*`cudaErrorSharedObjectInitFailed`* | |*`hipErrorSharedObjectInitFailed`* |
| 304 |*`cudaErrorOperatingSystem`* | |*`hipErrorOperatingSystem`* |
| 400 |*`cudaErrorInvalidResourceHandle`* | |*`hipErrorInvalidResourceHandle`* | 1033 |
| 400 |*`cudaErrorInvalidResourceHandle`* | |*`hipErrorInvalidHandle`* |
| 401 |*`cudaErrorIllegalState`* | 10.0 | |
| 500 |*`cudaErrorSymbolNotFound`* | 10.1 |*`hipErrorNotFound`* |
| 600 |*`cudaErrorNotReady`* | |*`hipErrorNotReady`* | 1034 |
| 600 |*`cudaErrorNotReady`* | |*`hipErrorNotReady`* |
| 700 |*`cudaErrorIllegalAddress`* | |*`hipErrorIllegalAddress`* |
| 701 |*`cudaErrorLaunchOutOfResources`* | |*`hipErrorLaunchOutOfResources`* | 1007 |
| 702 |*`cudaErrorLaunchTimeout`* | |*`hipErrorLaunchTimeOut`* | 1006 |
| 701 |*`cudaErrorLaunchOutOfResources`* | |*`hipErrorLaunchOutOfResources`* |
| 702 |*`cudaErrorLaunchTimeout`* | |*`hipErrorLaunchTimeOut`* |
| 703 |*`cudaErrorLaunchIncompatibleTexturing`* | | |
| 704 |*`cudaErrorPeerAccessAlreadyEnabled`* | |*`hipErrorPeerAccessAlreadyEnabled`* | 1050 |
| 705 |*`cudaErrorPeerAccessNotEnabled`* | |*`hipErrorPeerAccessNotEnabled`* | 1051 |
| 708 |*`cudaErrorSetOnActiveProcess`* | |*`hipErrorSetOnActiveProcess`* | 305 |
| 704 |*`cudaErrorPeerAccessAlreadyEnabled`* | |*`hipErrorPeerAccessAlreadyEnabled`* |
| 705 |*`cudaErrorPeerAccessNotEnabled`* | |*`hipErrorPeerAccessNotEnabled`* |
| 708 |*`cudaErrorSetOnActiveProcess`* | |*`hipErrorSetOnActiveProcess`* |
| 709 |*`cudaErrorContextIsDestroyed`* | | |
| 710 |*`cudaErrorAssert`* | |*`hipErrorAssert`* | 1081 |
| 710 |*`cudaErrorAssert`* | |*`hipErrorAssert`* |
| 711 |*`cudaErrorTooManyPeers`* | | |
| 712 |*`cudaErrorHostMemoryAlreadyRegistered`* | |*`hipErrorHostMemoryAlreadyRegistered`* | 1061 |
| 713 |*`cudaErrorHostMemoryNotRegistered`* | |*`hipErrorHostMemoryNotRegistered`* | 1062 |
| 712 |*`cudaErrorHostMemoryAlreadyRegistered`* | |*`hipErrorHostMemoryAlreadyRegistered`* |
| 713 |*`cudaErrorHostMemoryNotRegistered`* | |*`hipErrorHostMemoryNotRegistered`* |
| 714 |*`cudaErrorHardwareStackError`* | | |
| 715 |*`cudaErrorIllegalInstruction`* | | |
| 716 |*`cudaErrorMisalignedAddress`* | | |
| 717 |*`cudaErrorInvalidAddressSpace`* | | |
| 718 |*`cudaErrorInvalidPc`* | | |
| 719 |*`cudaErrorLaunchFailure`* | |*`hipErrorLaunchFailure`* | 1004 |
| 719 |*`cudaErrorLaunchFailure`* | |*`hipErrorLaunchFailure`* |
| 720 |*`cudaErrorCooperativeLaunchTooLarge`* | 9.0 | |
| 800 |*`cudaErrorNotPermitted`* | | |
| 801 |*`cudaErrorNotSupported`* | |*`hipErrorNotSupported`* | 1082 |
| 801 |*`cudaErrorNotSupported`* | |*`hipErrorNotSupported`* |
| 802 |*`cudaErrorSystemNotReady`* | 10.0 | |
| 803 |*`cudaErrorSystemDriverMismatch`* | 10.0 | |
| 804 |*`cudaErrorCompatNotSupportedOnDevice`* | 10.0 | |
@@ -817,7 +817,7 @@
| 908 |*`cudaErrorStreamCaptureWrongThread`* | 10.1 | |
| 909 |*`cudaErrorTimeout`* | 10.2 | |
| 910 |*`cudaErrorGraphExecUpdateFailure`* | 10.2 | |
| 999 |*`cudaErrorUnknown`* | |*`hipErrorUnknown`* | 1030 |
| 999 |*`cudaErrorUnknown`* | |*`hipErrorUnknown`* |
| 10000 |*`cudaErrorApiFailureBase`* | | |
| enum |***`cudaFuncCache`*** | |***`hipFuncCache_t`*** |
| 0 |*`cudaFuncCachePreferNone`* | |*`hipFuncCachePreferNone`* |
@@ -1067,14 +1067,14 @@
| 4 |*`CUDA_C_32F`* | 8.0 |*`HIPBLAS_C_32F`* | 154 |
| 1 |*`CUDA_R_64F`* | 8.0 |*`HIPBLAS_R_64F`* | 152 |
| 5 |*`CUDA_C_64F`* | 8.0 |*`HIPBLAS_C_64F`* | 155 |
| 3 |*`CUDA_R_8I`* | 8.0 | |
| 7 |*`CUDA_C_8I`* | 8.0 | |
| 8 |*`CUDA_R_8U`* | 8.0 | |
| 9 |*`CUDA_C_8U`* | 8.0 | |
| 10 |*`CUDA_R_32I`* | 8.0 | |
| 11 |*`CUDA_C_32I`* | 8.0 | |
| 12 |*`CUDA_R_32U`* | 8.0 | |
| 13 |*`CUDA_C_32U`* | 8.0 | |
| 3 |*`CUDA_R_8I`* | 8.0 |*`HIPBLAS_R_8I`* | 160 |
| 7 |*`CUDA_C_8I`* | 8.0 |*`HIPBLAS_C_8I`* | 164 |
| 8 |*`CUDA_R_8U`* | 8.0 |*`HIPBLAS_R_8U`* | 161 |
| 9 |*`CUDA_C_8U`* | 8.0 |*`HIPBLAS_C_8U`* | 165 |
| 10 |*`CUDA_R_32I`* | 8.0 |*`HIPBLAS_R_32I`* | 162 |
| 11 |*`CUDA_C_32I`* | 8.0 |*`HIPBLAS_C_32I`* | 166 |
| 12 |*`CUDA_R_32U`* | 8.0 |*`HIPBLAS_R_32U`* | 163 |
| 13 |*`CUDA_C_32U`* | 8.0 |*`HIPBLAS_C_32U`* | 167 |
| struct |`cudaExternalMemoryBufferDesc` | 10.0 | |
| struct |`cudaExternalMemoryHandleDesc` | 10.0 | |
| struct |`cudaExternalMemoryMipmappedArrayDesc` | 10.0 | |
projects/clr/hipamd/docs/markdown/hip_kernel_language.md
+2 -2
Näytä tiedosto
@@ -119,7 +119,7 @@ __global__ MyKernel(hipLaunchParm lp, float *A, float *B, float *C, size_t N)
...
}
MyKernel<<<dim3(gridDim), dim3(gridDim), 0, 0>>> (a,b,c,n);
MyKernel<<<dim3(gridDim), dim3(groupDim), 0, 0>>> (a,b,c,n);
// Alternatively, kernel can be launched by
// hipLaunchKernel(MyKernel, dim3(gridDim), dim3(groupDim), 0/*dynamicShared*/, 0/*stream), a, b, c, n);
@@ -153,7 +153,7 @@ void callMyKernel()
unsigned N = 1000000;
const unsigned blockSize = 256;
MyKernel<<<dim3(gridDim), dim3(gridDim), 0, 0>>> (a,b,c,n);
MyKernel<<<dim3(gridDim), dim3(groupDim), 0, 0>>> (a,b,c,n);
// Alternatively, kernel can be launched by
// hipLaunchKernel(MyKernel, dim3(N/blockSize), dim3(blockSize), 0, 0, a,b,c,N);
}
projects/clr/hipamd/docs/markdown/hip_programming_guide.md
+11
Näytä tiedosto
@@ -115,4 +115,15 @@ allocated.
In HCC and HIP-Clang, long double type is 80-bit extended precision format for x86_64, which is not supported by AMDGPU. HCC and HIP-Clang treat long double type as IEEE double type for AMDGPU. Using long double type in HIP source code will not cause issue as long as data of long double type is not transferred between host and device. However, long double type should not be used as kernel argument type.
## FMA and contractions
By default HIP-Clang assumes -ffp-contract=fast and HCC assumes -ffp-contract=off.
For x86_64, FMA is off by default since the generic x86_64 target does not
support FMA by default. To turn on FMA on x86_64, either use -mfma or -march=native
on CPU's supporting FMA.
When contractions are enabled and the CPU has not enabled FMA instructions, the
GPU can produce different numerical results than the CPU for expressions that
can be contracted. Tolerance should be used for floating point comparsions.
## [Supported Clang Options](clang_options.md)
projects/clr/hipamd/hipify-clang/CMakeLists.txt
+4
Näytä tiedosto
@@ -58,6 +58,10 @@ if (LLVM_PACKAGE_VERSION VERSION_GREATER "6.0.1")
target_link_libraries(hipify-clang PRIVATE clangToolingInclusions)
endif()
if (LLVM_PACKAGE_VERSION VERSION_GREATER "9.0.1")
target_link_libraries(hipify-clang PRIVATE LLVMFrontendOpenMP)
endif()
if (MSVC)
target_link_libraries(hipify-clang PRIVATE version)
target_compile_options(hipify-clang PRIVATE "/Od /GR- /EHs- /EHc-")
projects/clr/hipamd/hipify-clang/README.md
+52 -51
Näytä tiedosto
@@ -34,36 +34,37 @@
`hipify-clang` requires:
1. [**LLVM+CLANG**](http://releases.llvm.org) of at least version [3.8.0](http://releases.llvm.org/download.html#3.8.0); the latest stable and recommended release: [**9.0.0**](http://releases.llvm.org/download.html#9.0.0).
1. [**LLVM+CLANG**](http://releases.llvm.org) of at least version [3.8.0](http://releases.llvm.org/download.html#3.8.0); the latest stable and recommended release: [**9.0.1**](http://releases.llvm.org/download.html#9.0.1).
2. [**CUDA**](https://developer.nvidia.com/cuda-downloads) of at least version [7.0](https://developer.nvidia.com/cuda-toolkit-70), the latest supported version is [**10.1 Update 2**](https://developer.nvidia.com/cuda-downloads).
2. [**CUDA**](https://developer.nvidia.com/cuda-downloads) of at least version [7.0](https://developer.nvidia.com/cuda-toolkit-70), the latest supported version is [**10.1 Update 2**](https://developer.nvidia.com/cuda-10.1-download-archive-base).
| **LLVM release version** | **CUDA latest supported version** | **Windows** | **Linux** |
|:----------------------------------------------------------:|:-------------------------------------------------------------------:|:------------:|:---------:|
| [3.8.0](http://releases.llvm.org/download.html#3.8.0) | [7.5](https://developer.nvidia.com/cuda-75-downloads-archive) | + | + |
| [3.8.1](http://releases.llvm.org/download.html#3.8.1) | [7.5](https://developer.nvidia.com/cuda-75-downloads-archive) | + | + |
| [3.9.0](http://releases.llvm.org/download.html#3.9.0) | [7.5](https://developer.nvidia.com/cuda-75-downloads-archive) | + | + |
| [3.9.1](http://releases.llvm.org/download.html#3.9.1) | [7.5](https://developer.nvidia.com/cuda-75-downloads-archive) | + | + |
| [4.0.0](http://releases.llvm.org/download.html#4.0.0) | [8.0](https://developer.nvidia.com/cuda-80-ga2-download-archive) | + | + |
| [4.0.1](http://releases.llvm.org/download.html#4.0.1) | [8.0](https://developer.nvidia.com/cuda-80-ga2-download-archive) | + | + |
| [5.0.0](http://releases.llvm.org/download.html#5.0.0) | [8.0](https://developer.nvidia.com/cuda-80-ga2-download-archive) | + | + |
| [5.0.1](http://releases.llvm.org/download.html#5.0.1) | [8.0](https://developer.nvidia.com/cuda-80-ga2-download-archive) | + | + |
| [5.0.2](http://releases.llvm.org/download.html#5.0.2) | [8.0](https://developer.nvidia.com/cuda-80-ga2-download-archive) | + | + |
| [6.0.0](http://releases.llvm.org/download.html#6.0.0) | [9.0](https://developer.nvidia.com/cuda-90-download-archive) | + | + |
| [6.0.1](http://releases.llvm.org/download.html#6.0.1) | [9.0](https://developer.nvidia.com/cuda-90-download-archive) | + | + |
| [7.0.0](http://releases.llvm.org/download.html#7.0.0) | [9.2](https://developer.nvidia.com/cuda-92-download-archive) | - <br/> not working due to <br/> the clang's bug [38811](https://bugs.llvm.org/show_bug.cgi?id=38811) <br/>+<br/>[patch](patches/patch_for_clang_7.0.0_bug_38811.zip)*</br> | - <br/> not working due to <br/> the clang's bug [36384](https://bugs.llvm.org/show_bug.cgi?id=36384) |
| [7.0.1](http://releases.llvm.org/download.html#7.0.1) | [9.2](https://developer.nvidia.com/cuda-92-download-archive) | - <br/> not working due to <br/> the clang's bug [38811](https://bugs.llvm.org/show_bug.cgi?id=38811) <br/>+<br/>[patch](patches/patch_for_clang_7.0.1_bug_38811.zip)*</br> | - <br/> not working due to <br/> the clang's bug [36384](https://bugs.llvm.org/show_bug.cgi?id=36384) |
| [7.1.0](http://releases.llvm.org/download.html#7.1.0) | [9.2](https://developer.nvidia.com/cuda-92-download-archive) | - <br/> not working due to <br/> the clang's bug [38811](https://bugs.llvm.org/show_bug.cgi?id=38811) <br/>+<br/>[patch](patches/patch_for_clang_7.1.0_bug_38811.zip)*</br> | - <br/> not working due to <br/> the clang's bug [36384](https://bugs.llvm.org/show_bug.cgi?id=36384) |
| [8.0.0](http://releases.llvm.org/download.html#8.0.0) | [10.0](https://developer.nvidia.com/cuda-10.0-download-archive) | - <br/> not working due to <br/> the clang's bug [38811](https://bugs.llvm.org/show_bug.cgi?id=38811) <br/>+<br/>[patch](patches/patch_for_clang_8.0.0_bug_38811.zip)*</br> | + |
| [8.0.1](http://releases.llvm.org/download.html#8.0.1) | [10.0](https://developer.nvidia.com/cuda-10.0-download-archive) | - <br/> not working due to <br/> the clang's bug [38811](https://bugs.llvm.org/show_bug.cgi?id=38811) <br/>+<br/>[patch](patches/patch_for_clang_8.0.1_bug_38811.zip)*</br> | + |
| [**9.0.0**](http://releases.llvm.org/download.html#9.0.0) | [**10.1**](https://developer.nvidia.com/cuda-downloads) | + <br/> **LATEST STABLE RELEASE** | + <br/> **LATEST STABLE RELEASE** |
| **LLVM release version** | **CUDA latest supported version** | **Windows** | **Linux** |
|:----------------------------------------------------------:|:------------------------------------------------------------------------:|:-----------:|:---------:|
| [3.8.0](http://releases.llvm.org/download.html#3.8.0) | [7.5](https://developer.nvidia.com/cuda-75-downloads-archive) | + | + |
| [3.8.1](http://releases.llvm.org/download.html#3.8.1) | [7.5](https://developer.nvidia.com/cuda-75-downloads-archive) | + | + |
| [3.9.0](http://releases.llvm.org/download.html#3.9.0) | [7.5](https://developer.nvidia.com/cuda-75-downloads-archive) | + | + |
| [3.9.1](http://releases.llvm.org/download.html#3.9.1) | [7.5](https://developer.nvidia.com/cuda-75-downloads-archive) | + | + |
| [4.0.0](http://releases.llvm.org/download.html#4.0.0) | [8.0](https://developer.nvidia.com/cuda-80-ga2-download-archive) | + | + |
| [4.0.1](http://releases.llvm.org/download.html#4.0.1) | [8.0](https://developer.nvidia.com/cuda-80-ga2-download-archive) | + | + |
| [5.0.0](http://releases.llvm.org/download.html#5.0.0) | [8.0](https://developer.nvidia.com/cuda-80-ga2-download-archive) | + | + |
| [5.0.1](http://releases.llvm.org/download.html#5.0.1) | [8.0](https://developer.nvidia.com/cuda-80-ga2-download-archive) | + | + |
| [5.0.2](http://releases.llvm.org/download.html#5.0.2) | [8.0](https://developer.nvidia.com/cuda-80-ga2-download-archive) | + | + |
| [6.0.0](http://releases.llvm.org/download.html#6.0.0) | [9.0](https://developer.nvidia.com/cuda-90-download-archive) | + | + |
| [6.0.1](http://releases.llvm.org/download.html#6.0.1) | [9.0](https://developer.nvidia.com/cuda-90-download-archive) | + | + |
| [7.0.0](http://releases.llvm.org/download.html#7.0.0) | [9.2](https://developer.nvidia.com/cuda-92-download-archive) | - <br/> not working due to <br/> the clang's bug [38811](https://bugs.llvm.org/show_bug.cgi?id=38811) <br/>+<br/>[patch](patches/patch_for_clang_7.0.0_bug_38811.zip)*</br> | - <br/> not working due to <br/> the clang's bug [36384](https://bugs.llvm.org/show_bug.cgi?id=36384) |
| [7.0.1](http://releases.llvm.org/download.html#7.0.1) | [9.2](https://developer.nvidia.com/cuda-92-download-archive) | - <br/> not working due to <br/> the clang's bug [38811](https://bugs.llvm.org/show_bug.cgi?id=38811) <br/>+<br/>[patch](patches/patch_for_clang_7.0.1_bug_38811.zip)*</br> | - <br/> not working due to <br/> the clang's bug [36384](https://bugs.llvm.org/show_bug.cgi?id=36384) |
| [7.1.0](http://releases.llvm.org/download.html#7.1.0) | [9.2](https://developer.nvidia.com/cuda-92-download-archive) | - <br/> not working due to <br/> the clang's bug [38811](https://bugs.llvm.org/show_bug.cgi?id=38811) <br/>+<br/>[patch](patches/patch_for_clang_7.1.0_bug_38811.zip)*</br> | - <br/> not working due to <br/> the clang's bug [36384](https://bugs.llvm.org/show_bug.cgi?id=36384) |
| [8.0.0](http://releases.llvm.org/download.html#8.0.0) | [10.0](https://developer.nvidia.com/cuda-10.0-download-archive) | - <br/> not working due to <br/> the clang's bug [38811](https://bugs.llvm.org/show_bug.cgi?id=38811) <br/>+<br/>[patch](patches/patch_for_clang_8.0.0_bug_38811.zip)*</br> | + |
| [8.0.1](http://releases.llvm.org/download.html#8.0.1) | [10.0](https://developer.nvidia.com/cuda-10.0-download-archive) | - <br/> not working due to <br/> the clang's bug [38811](https://bugs.llvm.org/show_bug.cgi?id=38811) <br/>+<br/>[patch](patches/patch_for_clang_8.0.1_bug_38811.zip)*</br> | + |
| [9.0.0](http://releases.llvm.org/download.html#9.0.0) | [10.1](https://developer.nvidia.com/cuda-10.1-download-archive-base) | + | + |
| [**9.0.1**](http://releases.llvm.org/download.html#9.0.1) | [**10.1**](https://developer.nvidia.com/cuda-10.1-download-archive-base) | + <br/> **LATEST STABLE RELEASE** | + <br/> **LATEST STABLE RELEASE** |
`*` Download the patch and unpack it into your LLVM distributive directory; a few header files will be overwritten; rebuilding of LLVM is not needed.
In most cases, you can get a suitable version of LLVM+CLANG with your package manager.
Failing that or having multiple versions of LLVM, you can [download a release archive](http://releases.llvm.org/), build or install it, and set
[CMAKE_PREFIX_PATH](https://cmake.org/cmake/help/v3.5/variable/CMAKE_PREFIX_PATH.html) so `cmake` can find it; for instance: `-DCMAKE_PREFIX_PATH=f:\LLVM\9.0.0\dist`
[CMAKE_PREFIX_PATH](https://cmake.org/cmake/help/v3.5/variable/CMAKE_PREFIX_PATH.html) so `cmake` can find it; for instance: `-DCMAKE_PREFIX_PATH=f:\LLVM\9.0.1\dist`
## <a name="build-and-install"></a> Build and install
@@ -97,7 +98,7 @@ The binary can then be found at `./dist/bin/hipify-clang`.
**LLVM+CLANG should be built from sources, pre-built binaries are not exhaustive for testing.**
To run it:
1. download [`LLVM`](http://releases.llvm.org/9.0.0/llvm-9.0.0.src.tar.xz)+[`CLANG`](http://releases.llvm.org/9.0.0/cfe-9.0.0.src.tar.xz) sources;
1. download [`LLVM`](http://releases.llvm.org/9.0.1/llvm-9.0.1.src.tar.xz)+[`CLANG`](http://releases.llvm.org/9.0.1/cfe-9.0.1.src.tar.xz) sources;
2. build [`LLVM+CLANG`](http://llvm.org/docs/CMake.html):
```shell
cd llvm
@@ -123,7 +124,7 @@ To run it:
-A x64 \
-DCMAKE_INSTALL_PREFIX=../dist \
-DLLVM_SOURCE_DIR=../llvm \
-DLLVM_TARGETS_TO_BUILD="X86;NVPTX" \
-DLLVM_TARGETS_TO_BUILD="NVPTX" \
-DCMAKE_BUILD_TYPE=Release \
-Thost=x64 \
../llvm
@@ -164,21 +165,21 @@ To run it:
* Install `lit` into `python`:
- ***Linux***: `python /srv/git/LLVM/9.0.0/llvm/utils/lit/setup.py install`
- ***Linux***: `python /srv/git/LLVM/9.0.1/llvm/utils/lit/setup.py install`
- ***Windows***: `python f:/LLVM/9.0.0/llvm/utils/lit/setup.py install`
- ***Windows***: `python f:/LLVM/9.0.1/llvm/utils/lit/setup.py install`
* Starting with LLVM 6.0.1 path to `llvm-lit` python script should be specified by the `LLVM_EXTERNAL_LIT` option:
- ***Linux***: `-DLLVM_EXTERNAL_LIT=/srv/git/LLVM/9.0.0/build/bin/llvm-lit`
- ***Linux***: `-DLLVM_EXTERNAL_LIT=/srv/git/LLVM/9.0.1/build/bin/llvm-lit`
- ***Windows***: `-DLLVM_EXTERNAL_LIT=f:/LLVM/9.0.0/build/Release/bin/llvm-lit.py`
- ***Windows***: `-DLLVM_EXTERNAL_LIT=f:/LLVM/9.0.1/build/Release/bin/llvm-lit.py`
* `FileCheck`:
- ***Linux***: copy from `/srv/git/LLVM/9.0.0/build/bin/` to `CMAKE_INSTALL_PREFIX/dist/bin`
- ***Linux***: copy from `/srv/git/LLVM/9.0.1/build/bin/` to `CMAKE_INSTALL_PREFIX/dist/bin`
- ***Windows***: copy from `f:/LLVM/9.0.0/build/Release/bin` to `CMAKE_INSTALL_PREFIX/dist/bin`
- ***Windows***: copy from `f:/LLVM/9.0.1/build/Release/bin` to `CMAKE_INSTALL_PREFIX/dist/bin`
- Or specify the path to `FileCheck` in `CMAKE_INSTALL_PREFIX` option
@@ -200,7 +201,7 @@ On Linux the following configurations are tested:
Ubuntu 14: LLVM 5.0.0 - 6.0.1, CUDA 7.0 - 9.0, cudnn-5.0.5 - cudnn-7.6.5.32
Ubuntu 16-18: LLVM 8.0.0 - 9.0.0, CUDA 8.0 - 10.1, cudnn-5.1.10 - cudnn-7.6.5.32
Ubuntu 16-18: LLVM 8.0.0 - 9.0.1, CUDA 8.0 - 10.1 Update 2, cudnn-5.1.10 - cudnn-7.6.5.32
Minimum build system requirements for the above configurations:
@@ -213,11 +214,11 @@ cmake
-DHIPIFY_CLANG_TESTS=1 \
-DCMAKE_BUILD_TYPE=Release \
-DCMAKE_INSTALL_PREFIX=../dist \
-DCMAKE_PREFIX_PATH=/srv/git/LLVM/9.0.0/dist \
-DCMAKE_PREFIX_PATH=/srv/git/LLVM/9.0.1/dist \
-DCUDA_TOOLKIT_ROOT_DIR=/usr/local/cuda-10.1 \
-DCUDA_DNN_ROOT_DIR=/srv/CUDNN/cudnn-10.1-v7.6.5.32 \
-DCUDA_CUB_ROOT_DIR=/srv/git/CUB \
-DLLVM_EXTERNAL_LIT=/srv/git/LLVM/9.0.0/build/bin/llvm-lit \
-DLLVM_EXTERNAL_LIT=/srv/git/LLVM/9.0.1/build/bin/llvm-lit \
..
```
*A corresponding successful output:*
@@ -236,14 +237,14 @@ cmake
-- Detecting CXX compiler ABI info - done
-- Detecting CXX compile features
-- Detecting CXX compile features - done
-- Found LLVM 9.0.0:
-- - CMake module path: /srv/git/LLVM/9.0.0/dist/lib/cmake/llvm
-- - Include path : /srv/git/LLVM/9.0.0/dist/include
-- - Binary path : /srv/git/LLVM/9.0.0/dist/bin
-- Found LLVM 9.0.1:
-- - CMake module path: /srv/git/LLVM/9.0.1/dist/lib/cmake/llvm
-- - Include path : /srv/git/LLVM/9.0.1/dist/include
-- - Binary path : /srv/git/LLVM/9.0.1/dist/bin
-- Linker detection: GNU ld
-- Found PythonInterp: /usr/bin/python2.7 (found suitable version "2.7.12", minimum required is "2.7")
-- Found lit: /usr/local/bin/lit
-- Found FileCheck: /srv/git/LLVM/9.0.0/dist/bin/FileCheck
-- Found FileCheck: /srv/git/LLVM/9.0.1/dist/bin/FileCheck
-- Looking for pthread.h
-- Looking for pthread.h - found
-- Looking for pthread_create
@@ -266,7 +267,7 @@ make test-hipify
Running HIPify regression tests
========================================
CUDA 10.1 - will be used for testing
LLVM 9.0.0 - will be used for testing
LLVM 9.0.1 - will be used for testing
x86_64 - Platform architecture
Linux 5.2.0 - Platform OS
64 - hipify-clang binary bitness
@@ -352,11 +353,11 @@ LLVM 5.0.0 - 5.0.2, CUDA 8.0, cudnn 5.1.10 - 7.1.4.18
LLVM 6.0.0 - 6.0.1, CUDA 9.0, cudnn 7.0.5.15 - 7.6.5.32
LLVM 7.0.0 - 9.0.0, CUDA 7.5 - 10.1, cudnn 7.0.5.15 - 7.6.5.32
LLVM 7.0.0 - 9.0.1, CUDA 7.5 - 10.1 Update 2, cudnn 7.0.5.15 - 7.6.5.32
Build system requirements for the latest configuration LLVM 9.0.0/CUDA 10.1 Update 2:
Build system requirements for the latest configuration LLVM 9.0.1/CUDA 10.1 Update 2:
Python 3.6.0 - 3.8.0, cmake 3.5.1 - 3.16.0, Visual Studio 2017 (15.5.2) - 2019 (16.3.8).
Python 3.6.0 - 3.8.1, cmake 3.5.1 - 3.16.2, Visual Studio 2017 (15.5.2) - 2019 (16.4.2).
Here is an example of building `hipify-clang` with testing support on `Windows 10` by `Visual Studio 16 2019`:
@@ -367,24 +368,24 @@ cmake
-DHIPIFY_CLANG_TESTS=1 \
-DCMAKE_BUILD_TYPE=Release \
-DCMAKE_INSTALL_PREFIX=../dist \
-DCMAKE_PREFIX_PATH=f:/LLVM/9.0.0/dist \
-DCMAKE_PREFIX_PATH=f:/LLVM/9.0.1/dist \
-DCUDA_TOOLKIT_ROOT_DIR="c:/Program Files/NVIDIA GPU Computing Toolkit/CUDA/v10.1" \
-DCUDA_SDK_ROOT_DIR="c:/ProgramData/NVIDIA Corporation/CUDA Samples/v10.1" \
-DCUDA_DNN_ROOT_DIR=f:/CUDNN/cudnn-10.1-windows10-x64-v7.6.5.32 \
-DCUDA_CUB_ROOT_DIR=f:/GIT/cub \
-DLLVM_EXTERNAL_LIT=f:/LLVM/9.0.0/build/Release/bin/llvm-lit.py \
-DLLVM_EXTERNAL_LIT=f:/LLVM/9.0.1/build/Release/bin/llvm-lit.py \
-Thost=x64
..
```
*A corresponding successful output:*
```shell
-- Found LLVM 9.0.0:
-- - CMake module path: F:/LLVM/9.0.0/dist/lib/cmake/llvm
-- - Include path : F:/LLVM/9.0.0/dist/include
-- - Binary path : F:/LLVM/9.0.0/dist/bin
-- Found PythonInterp: C:/Program Files/Python38/python.exe (found suitable version "3.8.0", minimum required is "3.6")
-- Found lit: C:/Program Files/Python36/Scripts/lit.exe
-- Found FileCheck: F:/LLVM/9.0.0/dist/bin/FileCheck.exe
-- Found LLVM 9.0.1:
-- - CMake module path: F:/LLVM/9.0.1/dist/lib/cmake/llvm
-- - Include path : F:/LLVM/9.0.1/dist/include
-- - Binary path : F:/LLVM/9.0.1/dist/bin
-- Found PythonInterp: C:/Program Files/Python38/python.exe (found suitable version "3.8.1", minimum required is "3.6")
-- Found lit: C:/Program Files/Python38/Scripts/lit.exe
-- Found FileCheck: F:/LLVM/9.0.1/dist/bin/FileCheck.exe
-- Found CUDA: C:/Program Files/NVIDIA GPU Computing Toolkit/CUDA/v10.1 (found version "10.1")
-- Configuring done
-- Generating done
projects/clr/hipamd/hipify-clang/src/CUDA2HIP_BLAS_API_functions.cpp
+112 -129
Näytä tiedosto
@@ -64,108 +64,91 @@ const std::map<llvm::StringRef, hipCounter> CUDA_BLAS_FUNCTION_MAP{
// NRM2
{"cublasSnrm2", {"hipblasSnrm2", "rocblas_snrm2", CONV_LIB_FUNC, API_BLAS}},
{"cublasDnrm2", {"hipblasDnrm2", "rocblas_dnrm2", CONV_LIB_FUNC, API_BLAS}},
{"cublasScnrm2", {"hipblasScnrm2", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}},
{"cublasDznrm2", {"hipblasDznrm2", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}},
{"cublasScnrm2", {"hipblasScnrm2", "rocblas_scnrm2", CONV_LIB_FUNC, API_BLAS}},
{"cublasDznrm2", {"hipblasDznrm2", "rocblas_dznrm2", CONV_LIB_FUNC, API_BLAS}},
{"cublasNrm2Ex", {"hipblasNrm2Ex", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}},
// DOT
{"cublasSdot", {"hipblasSdot", "rocblas_sdot", CONV_LIB_FUNC, API_BLAS}},
// NOTE: there is no such a function in CUDA
{"cublasSdotBatched", {"hipblasSdotBatched", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}},
{"cublasDdot", {"hipblasDdot", "rocblas_ddot", CONV_LIB_FUNC, API_BLAS}},
// NOTE: there is no such a function in CUDA
{"cublasDdotBatched", {"hipblasDdotBatched", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}},
{"cublasCdotu", {"hipblasCdotu", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}},
{"cublasCdotc", {"hipblasCdotc", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}},
{"cublasZdotu", {"hipblasZdotu", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}},
{"cublasZdotc", {"hipblasZdotc", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}},
{"cublasCdotu", {"hipblasCdotu", "rocblas_cdotu", CONV_LIB_FUNC, API_BLAS}},
{"cublasCdotc", {"hipblasCdotc", "rocblas_cdotc", CONV_LIB_FUNC, API_BLAS}},
{"cublasZdotu", {"hipblasZdotu", "rocblas_zdotu", CONV_LIB_FUNC, API_BLAS}},
{"cublasZdotc", {"hipblasZdotc", "rocblas_zdotc", CONV_LIB_FUNC, API_BLAS}},
// SCAL
{"cublasSscal", {"hipblasSscal", "rocblas_sscal", CONV_LIB_FUNC, API_BLAS}},
// NOTE: there is no such a function in CUDA
{"cublasSscalBatched", {"hipblasSscalBatched", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}},
{"cublasDscal", {"hipblasDscal", "rocblas_dscal", CONV_LIB_FUNC, API_BLAS}},
// NOTE: there is no such a function in CUDA
{"cublasDscalBatched", {"hipblasDscalBatched", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}},
{"cublasCscal", {"hipblasCscal", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}},
{"cublasCsscal", {"hipblasCsscal", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}},
{"cublasZscal", {"hipblasZscal", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}},
{"cublasZdscal", {"hipblasZdscal", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}},
{"cublasCscal", {"hipblasCscal", "rocblas_cscal", CONV_LIB_FUNC, API_BLAS}},
{"cublasCsscal", {"hipblasCsscal", "rocblas_csscal", CONV_LIB_FUNC, API_BLAS}},
{"cublasZscal", {"hipblasZscal", "rocblas_zscal", CONV_LIB_FUNC, API_BLAS}},
{"cublasZdscal", {"hipblasZdscal", "rocblas_zdscal", CONV_LIB_FUNC, API_BLAS}},
// AXPY
{"cublasSaxpy", {"hipblasSaxpy", "rocblas_saxpy", CONV_LIB_FUNC, API_BLAS}},
{"cublasSaxpyBatched", {"hipblasSaxpyBatched", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}},
{"cublasDaxpy", {"hipblasDaxpy", "rocblas_daxpy", CONV_LIB_FUNC, API_BLAS}},
{"cublasCaxpy", {"hipblasCaxpy", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}},
{"cublasZaxpy", {"hipblasZaxpy", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}},
{"cublasCaxpy", {"hipblasCaxpy", "rocblas_caxpy", CONV_LIB_FUNC, API_BLAS}},
{"cublasZaxpy", {"hipblasZaxpy", "rocblas_zaxpy", CONV_LIB_FUNC, API_BLAS}},
// COPY
{"cublasScopy", {"hipblasScopy", "rocblas_scopy", CONV_LIB_FUNC, API_BLAS}},
// NOTE: there is no such a function in CUDA
{"cublasScopyBatched", {"hipblasScopyBatched", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}},
{"cublasDcopy", {"hipblasDcopy", "rocblas_dcopy", CONV_LIB_FUNC, API_BLAS}},
// NOTE: there is no such a function in CUDA
{"cublasDcopyBatched", {"hipblasDcopyBatched", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}},
{"cublasCcopy", {"hipblasCcopy", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}},
{"cublasZcopy", {"hipblasZcopy", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}},
{"cublasCcopy", {"hipblasCcopy", "rocblas_ccopy", CONV_LIB_FUNC, API_BLAS}},
{"cublasZcopy", {"hipblasZcopy", "rocblas_zcopy", CONV_LIB_FUNC, API_BLAS}},
// SWAP
{"cublasSswap", {"hipblasSswap", "rocblas_sswap", CONV_LIB_FUNC, API_BLAS, HIP_UNSUPPORTED}},
{"cublasDswap", {"hipblasDswap", "rocblas_dswap", CONV_LIB_FUNC, API_BLAS, HIP_UNSUPPORTED}},
{"cublasCswap", {"hipblasCswap", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}},
{"cublasZswap", {"hipblasZswap", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}},
{"cublasSswap", {"hipblasSswap", "rocblas_sswap", CONV_LIB_FUNC, API_BLAS}},
{"cublasDswap", {"hipblasDswap", "rocblas_dswap", CONV_LIB_FUNC, API_BLAS}},
{"cublasCswap", {"hipblasCswap", "rocblas_cswap", CONV_LIB_FUNC, API_BLAS}},
{"cublasZswap", {"hipblasZswap", "rocblas_zswap", CONV_LIB_FUNC, API_BLAS}},
// AMAX
{"cublasIsamax", {"hipblasIsamax", "rocblas_isamax", CONV_LIB_FUNC, API_BLAS}},
{"cublasIdamax", {"hipblasIdamax", "rocblas_idamax", CONV_LIB_FUNC, API_BLAS}},
{"cublasIcamax", {"hipblasIcamax", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}},
{"cublasIzamax", {"hipblasIzamax", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}},
{"cublasIcamax", {"hipblasIcamax", "rocblas_icamax", CONV_LIB_FUNC, API_BLAS}},
{"cublasIzamax", {"hipblasIzamax", "rocblas_izamax", CONV_LIB_FUNC, API_BLAS}},
// AMIN
{"cublasIsamin", {"hipblasIsamin", "rocblas_isamin", CONV_LIB_FUNC, API_BLAS, HIP_UNSUPPORTED}},
{"cublasIdamin", {"hipblasIdamin", "rocblas_idamin", CONV_LIB_FUNC, API_BLAS, HIP_UNSUPPORTED}},
{"cublasIcamin", {"hipblasIcamin", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}},
{"cublasIzamin", {"hipblasIzamin", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}},
{"cublasIsamin", {"hipblasIsamin", "rocblas_isamin", CONV_LIB_FUNC, API_BLAS}},
{"cublasIdamin", {"hipblasIdamin", "rocblas_idamin", CONV_LIB_FUNC, API_BLAS}},
{"cublasIcamin", {"hipblasIcamin", "rocblas_icamin", CONV_LIB_FUNC, API_BLAS}},
{"cublasIzamin", {"hipblasIzamin", "rocblas_izamin", CONV_LIB_FUNC, API_BLAS}},
// ASUM
{"cublasSasum", {"hipblasSasum", "rocblas_sasum", CONV_LIB_FUNC, API_BLAS}},
// NOTE: there is no such a function in CUDA
{"cublasSasumBatched", {"hipblasSasumBatched", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}},
{"cublasDasum", {"hipblasDasum", "rocblas_dasum", CONV_LIB_FUNC, API_BLAS}},
// NOTE: there is no such a function in CUDA
{"cublasDasumBatched", {"hipblasDasumBatched", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}},
{"cublasScasum", {"hipblasScasum", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}},
{"cublasDzasum", {"hipblasDzasum", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}},
{"cublasScasum", {"hipblasScasum", "rocblas_scasum", CONV_LIB_FUNC, API_BLAS}},
{"cublasDzasum", {"hipblasDzasum", "rocblas_dzasum", CONV_LIB_FUNC, API_BLAS}},
// ROT
{"cublasSrot", {"hipblasSrot", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}},
{"cublasDrot", {"hipblasDrot", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}},
{"cublasCrot", {"hipblasCrot", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}},
{"cublasCsrot", {"hipblasCsrot", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}},
{"cublasZrot", {"hipblasZrot", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}},
{"cublasZdrot", {"hipblasZdrot", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}},
{"cublasSrot", {"hipblasSrot", "rocblas_srot", CONV_LIB_FUNC, API_BLAS}},
{"cublasDrot", {"hipblasDrot", "rocblas_drot", CONV_LIB_FUNC, API_BLAS}},
{"cublasCrot", {"hipblasCrot", "rocblas_crot", CONV_LIB_FUNC, API_BLAS}},
{"cublasCsrot", {"hipblasCsrot", "rocblas_csrot", CONV_LIB_FUNC, API_BLAS}},
{"cublasZrot", {"hipblasZrot", "rocblas_zrot", CONV_LIB_FUNC, API_BLAS}},
{"cublasZdrot", {"hipblasZdrot", "rocblas_zdrot", CONV_LIB_FUNC, API_BLAS}},
// ROTG
{"cublasSrotg", {"hipblasSrotg", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}},
{"cublasDrotg", {"hipblasDrotg", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}},
{"cublasCrotg", {"hipblasCrotg", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}},
{"cublasZrotg", {"hipblasZrotg", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}},
{"cublasSrotg", {"hipblasSrotg", "rocblas_srotg", CONV_LIB_FUNC, API_BLAS}},
{"cublasDrotg", {"hipblasDrotg", "rocblas_drotg", CONV_LIB_FUNC, API_BLAS}},
{"cublasCrotg", {"hipblasCrotg", "rocblas_crotg", CONV_LIB_FUNC, API_BLAS}},
{"cublasZrotg", {"hipblasZrotg", "rocblas_zrotg", CONV_LIB_FUNC, API_BLAS}},
// ROTM
{"cublasSrotm", {"hipblasSrotm", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}},
{"cublasDrotm", {"hipblasDrotm", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}},
{"cublasSrotm", {"hipblasSrotm", "rocblas_srotm", CONV_LIB_FUNC, API_BLAS}},
{"cublasDrotm", {"hipblasDrotm", "rocblas_drotm", CONV_LIB_FUNC, API_BLAS}},
// ROTMG
{"cublasSrotmg", {"hipblasSrotmg", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}},
{"cublasDrotmg", {"hipblasDrotmg", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}},
{"cublasSrotmg", {"hipblasSrotmg", "rocblas_srotmg", CONV_LIB_FUNC, API_BLAS}},
{"cublasDrotmg", {"hipblasDrotmg", "rocblas_drotmg", CONV_LIB_FUNC, API_BLAS}},
// GEMV
{"cublasSgemv", {"hipblasSgemv", "rocblas_sgemv", CONV_LIB_FUNC, API_BLAS}},
// NOTE: there is no such a function in CUDA
{"cublasSgemvBatched", {"hipblasSgemvBatched", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}},
{"cublasSgemvBatched", {"hipblasSgemvBatched", "rocblas_sgemv_batched", CONV_LIB_FUNC, API_BLAS}},
{"cublasDgemv", {"hipblasDgemv", "rocblas_dgemv", CONV_LIB_FUNC, API_BLAS}},
{"cublasCgemv", {"hipblasCgemv", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}},
{"cublasZgemv", {"hipblasZgemv", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}},
{"cublasCgemv", {"hipblasCgemv", "rocblas_cgemv", CONV_LIB_FUNC, API_BLAS}},
{"cublasZgemv", {"hipblasZgemv", "rocblas_zgemv", CONV_LIB_FUNC, API_BLAS}},
// GBMV
{"cublasSgbmv", {"hipblasSgbmv", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}},
@@ -192,8 +175,8 @@ const std::map<llvm::StringRef, hipCounter> CUDA_BLAS_FUNCTION_MAP{
{"cublasZtpmv", {"hipblasZtpmv", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}},
// TRSV
{"cublasStrsv", {"hipblasStrsv", "rocblas_strsv", CONV_LIB_FUNC, API_BLAS, HIP_UNSUPPORTED}},
{"cublasDtrsv", {"hipblasDtrsv", "rocblas_dtrsv", CONV_LIB_FUNC, API_BLAS, HIP_UNSUPPORTED}},
{"cublasStrsv", {"hipblasStrsv", "rocblas_strsv", CONV_LIB_FUNC, API_BLAS}},
{"cublasDtrsv", {"hipblasDtrsv", "rocblas_dtrsv", CONV_LIB_FUNC, API_BLAS}},
{"cublasCtrsv", {"hipblasCtrsv", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}},
{"cublasZtrsv", {"hipblasZtrsv", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}},
@@ -238,8 +221,8 @@ const std::map<llvm::StringRef, hipCounter> CUDA_BLAS_FUNCTION_MAP{
{"cublasZgerc", {"hipblasZgerc", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}},
// SYR/HER
{"cublasSsyr", {"hipblasSsyr", "rocblas_ssyr", CONV_LIB_FUNC, API_BLAS, HIP_UNSUPPORTED}},
{"cublasDsyr", {"hipblasDsyr", "rocblas_dsyr", CONV_LIB_FUNC, API_BLAS, HIP_UNSUPPORTED}},
{"cublasSsyr", {"hipblasSsyr", "rocblas_ssyr", CONV_LIB_FUNC, API_BLAS}},
{"cublasDsyr", {"hipblasDsyr", "rocblas_dsyr", CONV_LIB_FUNC, API_BLAS}},
{"cublasCsyr", {"hipblasCsyr", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}},
{"cublasZsyr", {"hipblasZsyr", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}},
{"cublasCher", {"hipblasCher", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}},
@@ -269,23 +252,23 @@ const std::map<llvm::StringRef, hipCounter> CUDA_BLAS_FUNCTION_MAP{
// GEMM
{"cublasSgemm", {"hipblasSgemm", "rocblas_sgemm", CONV_LIB_FUNC, API_BLAS}},
{"cublasDgemm", {"hipblasDgemm", "rocblas_dgemm", CONV_LIB_FUNC, API_BLAS}},
{"cublasCgemm", {"hipblasCgemm", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}},
{"cublasZgemm", {"hipblasZgemm", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}},
{"cublasCgemm", {"hipblasCgemm", "rocblas_cgemm", CONV_LIB_FUNC, API_BLAS}},
{"cublasZgemm", {"hipblasZgemm", "rocblas_zgemm", CONV_LIB_FUNC, API_BLAS}},
{"cublasHgemm", {"hipblasHgemm", "rocblas_hgemm", CONV_LIB_FUNC, API_BLAS}},
// BATCH GEMM
{"cublasSgemmBatched", {"hipblasSgemmBatched", "", CONV_LIB_FUNC, API_BLAS, ROC_UNSUPPORTED}},
{"cublasDgemmBatched", {"hipblasDgemmBatched", "", CONV_LIB_FUNC, API_BLAS, ROC_UNSUPPORTED}},
{"cublasHgemmBatched", {"hipblasHgemmBatched", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}},
{"cublasSgemmBatched", {"hipblasSgemmBatched", "rocblas_sgemm_batched", CONV_LIB_FUNC, API_BLAS}},
{"cublasDgemmBatched", {"hipblasDgemmBatched", "rocblas_dgemm_batched", CONV_LIB_FUNC, API_BLAS}},
{"cublasHgemmBatched", {"hipblasHgemmBatched", "rocblas_hgemm_batched", CONV_LIB_FUNC, API_BLAS}},
{"cublasSgemmStridedBatched", {"hipblasSgemmStridedBatched", "rocblas_sgemm_strided_batched", CONV_LIB_FUNC, API_BLAS}},
{"cublasDgemmStridedBatched", {"hipblasDgemmStridedBatched", "rocblas_dgemm_strided_batched", CONV_LIB_FUNC, API_BLAS}},
{"cublasCgemmBatched", {"hipblasCgemmBatched", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}},
{"cublasCgemmBatched", {"hipblasCgemmBatched", "rocblas_cgemm_batched", CONV_LIB_FUNC, API_BLAS}},
{"cublasCgemm3mBatched", {"hipblasCgemm3mBatched", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}},
{"cublasZgemmBatched", {"hipblasZgemmBatched", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}},
{"cublasCgemmStridedBatched", {"hipblasCgemmStridedBatched", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}},
{"cublasZgemmBatched", {"hipblasZgemmBatched", "rocblas_zgemm_batched", CONV_LIB_FUNC, API_BLAS}},
{"cublasCgemmStridedBatched", {"hipblasCgemmStridedBatched", "rocblas_cgemm_strided_batched", CONV_LIB_FUNC, API_BLAS}},
{"cublasCgemm3mStridedBatched", {"hipblasCgemm3mStridedBatched", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}},
{"cublasZgemmStridedBatched", {"hipblasZgemmStridedBatched", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}},
{"cublasHgemmStridedBatched", {"hipblasHgemmStridedBatched", "rocblas_hgemm_strided_batched", CONV_LIB_FUNC, API_BLAS, HIP_UNSUPPORTED}},
{"cublasZgemmStridedBatched", {"hipblasZgemmStridedBatched", "rocblas_zgemm_strided_batched", CONV_LIB_FUNC, API_BLAS}},
{"cublasHgemmStridedBatched", {"hipblasHgemmStridedBatched", "rocblas_hgemm_strided_batched", CONV_LIB_FUNC, API_BLAS}},
// SYRK
{"cublasSsyrk", {"hipblasSsyrk", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}},
@@ -334,8 +317,8 @@ const std::map<llvm::StringRef, hipCounter> CUDA_BLAS_FUNCTION_MAP{
{"cublasZtrsm", {"hipblasZtrsm", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}},
// TRMM
{"cublasStrmm", {"hipblasStrmm", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}},
{"cublasDtrmm", {"hipblasDtrmm", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}},
{"cublasStrmm", {"hipblasStrmm", "rocblas_strmm", CONV_LIB_FUNC, API_BLAS, HIP_UNSUPPORTED}},
{"cublasDtrmm", {"hipblasDtrmm", "rocblas_dtrmm", CONV_LIB_FUNC, API_BLAS, HIP_UNSUPPORTED}},
{"cublasCtrmm", {"hipblasCtrmm", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}},
{"cublasZtrmm", {"hipblasZtrmm", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}},
@@ -420,8 +403,8 @@ const std::map<llvm::StringRef, hipCounter> CUDA_BLAS_FUNCTION_MAP{
// GEMV
{"cublasSgemv_v2", {"hipblasSgemv", "rocblas_sgemv", CONV_LIB_FUNC, API_BLAS}},
{"cublasDgemv_v2", {"hipblasDgemv", "rocblas_dgemv", CONV_LIB_FUNC, API_BLAS}},
{"cublasCgemv_v2", {"hipblasCgemv", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}},
{"cublasZgemv_v2", {"hipblasZgemv", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}},
{"cublasCgemv_v2", {"hipblasCgemv", "rocblas_cgemv", CONV_LIB_FUNC, API_BLAS}},
{"cublasZgemv_v2", {"hipblasZgemv", "rocblas_zgemv", CONV_LIB_FUNC, API_BLAS}},
// GBMV
{"cublasSgbmv_v2", {"hipblasSgbmv", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}},
@@ -448,8 +431,8 @@ const std::map<llvm::StringRef, hipCounter> CUDA_BLAS_FUNCTION_MAP{
{"cublasZtpmv_v2", {"hipblasZtpmv", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}},
// TRSV
{"cublasStrsv_v2", {"hipblasStrsv", "rocblas_strsv", CONV_LIB_FUNC, API_BLAS, HIP_UNSUPPORTED}},
{"cublasDtrsv_v2", {"hipblasDtrsv", "rocblas_dtrsv", CONV_LIB_FUNC, API_BLAS, HIP_UNSUPPORTED}},
{"cublasStrsv_v2", {"hipblasStrsv", "rocblas_strsv", CONV_LIB_FUNC, API_BLAS}},
{"cublasDtrsv_v2", {"hipblasDtrsv", "rocblas_dtrsv", CONV_LIB_FUNC, API_BLAS}},
{"cublasCtrsv_v2", {"hipblasCtrsv", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}},
{"cublasZtrsv_v2", {"hipblasZtrsv", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}},
@@ -494,8 +477,8 @@ const std::map<llvm::StringRef, hipCounter> CUDA_BLAS_FUNCTION_MAP{
{"cublasZgerc_v2", {"hipblasZgerc", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}},
// SYR/HER
{"cublasSsyr_v2", {"hipblasSsyr", "rocblas_ssyr", CONV_LIB_FUNC, API_BLAS, HIP_UNSUPPORTED}},
{"cublasDsyr_v2", {"hipblasDsyr", "rocblas_dsyr", CONV_LIB_FUNC, API_BLAS, HIP_UNSUPPORTED}},
{"cublasSsyr_v2", {"hipblasSsyr", "rocblas_ssyr", CONV_LIB_FUNC, API_BLAS}},
{"cublasDsyr_v2", {"hipblasDsyr", "rocblas_dsyr", CONV_LIB_FUNC, API_BLAS}},
{"cublasCsyr_v2", {"hipblasCsyr", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}},
{"cublasZsyr_v2", {"hipblasZsyr", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}},
{"cublasCher_v2", {"hipblasCher", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}},
@@ -528,7 +511,7 @@ const std::map<llvm::StringRef, hipCounter> CUDA_BLAS_FUNCTION_MAP{
{"cublasCgemm_v2", {"hipblasCgemm", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}},
{"cublasCgemm3m", {"hipblasCgemm3m", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}},
{"cublasCgemm3mEx", {"hipblasCgemm3mEx", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}},
{"cublasZgemm_v2", {"hipblasZgemm", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}},
{"cublasZgemm_v2", {"hipblasZgemm", "rocblas_zgemm", CONV_LIB_FUNC, API_BLAS}},
{"cublasZgemm3m", {"hipblasZgemm3m", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}},
//IO in FP16 / FP32, computation in float
@@ -586,16 +569,16 @@ const std::map<llvm::StringRef, hipCounter> CUDA_BLAS_FUNCTION_MAP{
{"cublasZtrsm_v2", {"hipblasZtrsm", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}},
// TRMM
{"cublasStrmm_v2", {"hipblasStrmm", "", CONV_LIB_FUNC, API_BLAS, HIP_UNSUPPORTED}},
{"cublasDtrmm_v2", {"hipblasDtrmm", "", CONV_LIB_FUNC, API_BLAS, HIP_UNSUPPORTED}},
{"cublasCtrmm_v2", {"hipblasCtrmm", "", CONV_LIB_FUNC, API_BLAS, HIP_UNSUPPORTED}},
{"cublasZtrmm_v2", {"hipblasZtrmm", "", CONV_LIB_FUNC, API_BLAS, HIP_UNSUPPORTED}},
{"cublasStrmm_v2", {"hipblasStrmm", "rocblas_strmm", CONV_LIB_FUNC, API_BLAS, HIP_UNSUPPORTED}},
{"cublasDtrmm_v2", {"hipblasDtrmm", "rocblas_dtrmm", CONV_LIB_FUNC, API_BLAS, HIP_UNSUPPORTED}},
{"cublasCtrmm_v2", {"hipblasCtrmm", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}},
{"cublasZtrmm_v2", {"hipblasZtrmm", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}},
// NRM2
{"cublasSnrm2_v2", {"hipblasSnrm2", "rocblas_snrm2", CONV_LIB_FUNC, API_BLAS}},
{"cublasDnrm2_v2", {"hipblasDnrm2", "rocblas_dnrm2", CONV_LIB_FUNC, API_BLAS}},
{"cublasScnrm2_v2", {"hipblasScnrm2", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}},
{"cublasDznrm2_v2", {"hipblasDznrm2", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}},
{"cublasScnrm2_v2", {"hipblasScnrm2", "rocblas_scnrm2", CONV_LIB_FUNC, API_BLAS}},
{"cublasDznrm2_v2", {"hipblasDznrm2", "rocblas_dznrm2", CONV_LIB_FUNC, API_BLAS}},
// DOT
{"cublasDotEx", {"hipblasDotEx", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}},
@@ -604,85 +587,85 @@ const std::map<llvm::StringRef, hipCounter> CUDA_BLAS_FUNCTION_MAP{
{"cublasSdot_v2", {"hipblasSdot", "rocblas_sdot", CONV_LIB_FUNC, API_BLAS}},
{"cublasDdot_v2", {"hipblasDdot", "rocblas_ddot", CONV_LIB_FUNC, API_BLAS}},
{"cublasCdotu_v2", {"hipblasCdotu", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}},
{"cublasCdotc_v2", {"hipblasCdotc", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}},
{"cublasZdotu_v2", {"hipblasZdotu", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}},
{"cublasZdotc_v2", {"hipblasZdotc", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}},
{"cublasCdotu_v2", {"hipblasCdotu", "rocblas_cdotu", CONV_LIB_FUNC, API_BLAS}},
{"cublasCdotc_v2", {"hipblasCdotc", "rocblas_cdotc", CONV_LIB_FUNC, API_BLAS,}},
{"cublasZdotu_v2", {"hipblasZdotu", "rocblas_zdotu", CONV_LIB_FUNC, API_BLAS}},
{"cublasZdotc_v2", {"hipblasZdotc", "rocblas_zdotc", CONV_LIB_FUNC, API_BLAS}},
// SCAL
{"cublasScalEx", {"hipblasScalEx", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}},
{"cublasSscal_v2", {"hipblasSscal", "rocblas_sscal", CONV_LIB_FUNC, API_BLAS}},
{"cublasDscal_v2", {"hipblasDscal", "rocblas_dscal", CONV_LIB_FUNC, API_BLAS}},
{"cublasCscal_v2", {"hipblasCscal", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}},
{"cublasCsscal_v2", {"hipblasCsscal", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}},
{"cublasZscal_v2", {"hipblasZscal", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}},
{"cublasZdscal_v2", {"hipblasZdscal", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}},
{"cublasCscal_v2", {"hipblasCscal", "rocblas_cscal", CONV_LIB_FUNC, API_BLAS}},
{"cublasCsscal_v2", {"hipblasCsscal", "rocblas_csscal", CONV_LIB_FUNC, API_BLAS}},
{"cublasZscal_v2", {"hipblasZscal", "rocblas_zscal", CONV_LIB_FUNC, API_BLAS}},
{"cublasZdscal_v2", {"hipblasZdscal", "rocblas_zdscal", CONV_LIB_FUNC, API_BLAS}},
// AXPY
{"cublasAxpyEx", {"hipblasAxpyEx", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}},
{"cublasSaxpy_v2", {"hipblasSaxpy", "rocblas_saxpy", CONV_LIB_FUNC, API_BLAS}},
{"cublasDaxpy_v2", {"hipblasDaxpy", "rocblas_daxpy", CONV_LIB_FUNC, API_BLAS}},
{"cublasCaxpy_v2", {"hipblasCaxpy", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}},
{"cublasZaxpy_v2", {"hipblasZaxpy", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}},
{"cublasCaxpy_v2", {"hipblasCaxpy", "rocblas_caxpy", CONV_LIB_FUNC, API_BLAS}},
{"cublasZaxpy_v2", {"hipblasZaxpy", "rocblas_zaxpy", CONV_LIB_FUNC, API_BLAS}},
// COPY
{"cublasCopyEx", {"hipblasCopyEx", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}},
{"cublasScopy_v2", {"hipblasScopy", "rocblas_scopy", CONV_LIB_FUNC, API_BLAS}},
{"cublasDcopy_v2", {"hipblasDcopy", "rocblas_dcopy", CONV_LIB_FUNC, API_BLAS}},
{"cublasCcopy_v2", {"hipblasCcopy", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}},
{"cublasZcopy_v2", {"hipblasZcopy", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}},
{"cublasCcopy_v2", {"hipblasCcopy", "rocblas_ccopy", CONV_LIB_FUNC, API_BLAS}},
{"cublasZcopy_v2", {"hipblasZcopy", "rocblas_zcopy", CONV_LIB_FUNC, API_BLAS}},
// SWAP
{"cublasSwapEx", {"hipblasSwapEx", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}},
{"cublasSswap_v2", {"hipblasSswap", "rocblas_sswap", CONV_LIB_FUNC, API_BLAS, HIP_UNSUPPORTED}},
{"cublasDswap_v2", {"hipblasDswap", "rocblas_dswap", CONV_LIB_FUNC, API_BLAS, HIP_UNSUPPORTED}},
{"cublasCswap_v2", {"hipblasCswap", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}},
{"cublasZswap_v2", {"hipblasZswap", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}},
{"cublasSswap_v2", {"hipblasSswap", "rocblas_sswap", CONV_LIB_FUNC, API_BLAS}},
{"cublasDswap_v2", {"hipblasDswap", "rocblas_dswap", CONV_LIB_FUNC, API_BLAS}},
{"cublasCswap_v2", {"hipblasCswap", "rocblas_cswap", CONV_LIB_FUNC, API_BLAS}},
{"cublasZswap_v2", {"hipblasZswap", "rocblas_zswap", CONV_LIB_FUNC, API_BLAS}},
// AMAX
{"cublasIamaxEx", {"hipblasIamaxEx", "", CONV_LIB_FUNC, API_BLAS, HIP_UNSUPPORTED}},
{"cublasIamaxEx", {"hipblasIamaxEx", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}},
{"cublasIsamax_v2", {"hipblasIsamax", "rocblas_isamax", CONV_LIB_FUNC, API_BLAS}},
{"cublasIdamax_v2", {"hipblasIdamax", "rocblas_idamax", CONV_LIB_FUNC, API_BLAS}},
{"cublasIcamax_v2", {"hipblasIcamax", "", CONV_LIB_FUNC, API_BLAS, HIP_UNSUPPORTED}},
{"cublasIzamax_v2", {"hipblasIzamax", "", CONV_LIB_FUNC, API_BLAS, HIP_UNSUPPORTED}},
{"cublasIcamax_v2", {"hipblasIcamax", "rocblas_icamax", CONV_LIB_FUNC, API_BLAS}},
{"cublasIzamax_v2", {"hipblasIzamax", "rocblas_izamax", CONV_LIB_FUNC, API_BLAS}},
// AMIN
{"cublasIaminEx", {"hipblasIaminEx", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}},
{"cublasIsamin_v2", {"hipblasIsamin", "rocblas_isamin", CONV_LIB_FUNC, API_BLAS, HIP_UNSUPPORTED}},
{"cublasIdamin_v2", {"hipblasIdamin", "rocblas_idamin", CONV_LIB_FUNC, API_BLAS, HIP_UNSUPPORTED}},
{"cublasIcamin_v2", {"hipblasIcamin", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}},
{"cublasIzamin_v2", {"hipblasIzamin", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}},
{"cublasIsamin_v2", {"hipblasIsamin", "rocblas_isamin", CONV_LIB_FUNC, API_BLAS}},
{"cublasIdamin_v2", {"hipblasIdamin", "rocblas_idamin", CONV_LIB_FUNC, API_BLAS}},
{"cublasIcamin_v2", {"hipblasIcamin", "rocblas_icamin", CONV_LIB_FUNC, API_BLAS}},
{"cublasIzamin_v2", {"hipblasIzamin", "rocblas_izamin", CONV_LIB_FUNC, API_BLAS}},
// ASUM
{"cublasAsumEx", {"hipblasAsumEx", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}},
{"cublasSasum_v2", {"hipblasSasum", "rocblas_sasum", CONV_LIB_FUNC, API_BLAS}},
{"cublasDasum_v2", {"hipblasDasum", "rocblas_dasum", CONV_LIB_FUNC, API_BLAS}},
{"cublasScasum_v2", {"hipblasScasum", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}},
{"cublasDzasum_v2", {"hipblasDzasum", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}},
{"cublasScasum_v2", {"hipblasScasum", "rocblas_scasum", CONV_LIB_FUNC, API_BLAS}},
{"cublasDzasum_v2", {"hipblasDzasum", "rocblas_dzasum", CONV_LIB_FUNC, API_BLAS}},
// ROT
{"cublasRotEx", {"hipblasRotEx", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}},
{"cublasSrot_v2", {"hipblasSrot", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}},
{"cublasDrot_v2", {"hipblasDrot", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}},
{"cublasCrot_v2", {"hipblasCrot", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}},
{"cublasCsrot_v2", {"hipblasCsrot", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}},
{"cublasZrot_v2", {"hipblasZrot", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}},
{"cublasZdrot_v2", {"hipblasZdrot", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}},
{"cublasSrot_v2", {"hipblasSrot", "rocblas_srot", CONV_LIB_FUNC, API_BLAS}},
{"cublasDrot_v2", {"hipblasDrot", "rocblas_drot", CONV_LIB_FUNC, API_BLAS}},
{"cublasCrot_v2", {"hipblasCrot", "rocblas_crot", CONV_LIB_FUNC, API_BLAS}},
{"cublasCsrot_v2", {"hipblasCsrot", "rocblas_csrot", CONV_LIB_FUNC, API_BLAS}},
{"cublasZrot_v2", {"hipblasZrot", "rocblas_zrot", CONV_LIB_FUNC, API_BLAS}},
{"cublasZdrot_v2", {"hipblasZdrot", "rocblas_zdrot", CONV_LIB_FUNC, API_BLAS}},
// ROTG
{"cublasRotgEx", {"hipblasRotgEx", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}},
{"cublasSrotg_v2", {"hipblasSrotg", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}},
{"cublasDrotg_v2", {"hipblasDrotg", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}},
{"cublasCrotg_v2", {"hipblasCrotg", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}},
{"cublasZrotg_v2", {"hipblasZrotg", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}},
{"cublasSrotg_v2", {"hipblasSrotg", "rocblas_srotg", CONV_LIB_FUNC, API_BLAS}},
{"cublasDrotg_v2", {"hipblasDrotg", "rocblas_drotg", CONV_LIB_FUNC, API_BLAS}},
{"cublasCrotg_v2", {"hipblasCrotg", "rocblas_crotg", CONV_LIB_FUNC, API_BLAS}},
{"cublasZrotg_v2", {"hipblasZrotg", "rocblas_zrotg", CONV_LIB_FUNC, API_BLAS}},
// ROTM
{"cublasRotmEx", {"hipblasRotmEx", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}},
{"cublasSrotm_v2", {"hipblasSrotm", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}},
{"cublasDrotm_v2", {"hipblasDrotm", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}},
{"cublasRotmEx", {"hipblasRotmEx", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}},
{"cublasSrotm_v2", {"hipblasSrotm", "rocblas_srotm", CONV_LIB_FUNC, API_BLAS}},
{"cublasDrotm_v2", {"hipblasDrotm", "rocblas_drotm", CONV_LIB_FUNC, API_BLAS}},
// ROTMG
{"cublasRotmgEx", {"hipblasRotmgEx", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}},
{"cublasSrotmg_v2", {"hipblasSrotmg", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}},
{"cublasDrotmg_v2", {"hipblasDrotmg", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}},
{"cublasSrotmg_v2", {"hipblasSrotmg", "rocblas_srotmg", CONV_LIB_FUNC, API_BLAS}},
{"cublasDrotmg_v2", {"hipblasDrotmg", "rocblas_drotmg", CONV_LIB_FUNC, API_BLAS}},
};
projects/clr/hipamd/hipify-clang/src/CUDA2HIP_BLAS_API_types.cpp
+8 -8
Näytä tiedosto
@@ -143,14 +143,14 @@ const std::map<llvm::StringRef, hipCounter> CUDA_BLAS_TYPE_NAME_MAP{
{"CUDA_C_32F", {"HIPBLAS_C_32F", "rocblas_datatype_f32_c", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 4 // 154
{"CUDA_R_64F", {"HIPBLAS_R_64F", "rocblas_datatype_f64_r", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 1 // 152
{"CUDA_C_64F", {"HIPBLAS_C_64F", "rocblas_datatype_f64_c", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 5 // 155
{"CUDA_R_8I", {"HIPBLAS_R_8I", "rocblas_datatype_i8_r", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 3 // 160
{"CUDA_C_8I", {"HIPBLAS_C_8I", "rocblas_datatype_i8_c", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 7 // 164
{"CUDA_R_8U", {"HIPBLAS_R_8U", "rocblas_datatype_u8_r", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 8 // 161
{"CUDA_C_8U", {"HIPBLAS_C_8U", "rocblas_datatype_u8_c", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 9 // 165
{"CUDA_R_32I", {"HIPBLAS_R_32I", "rocblas_datatype_i32_r", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 10 // 162
{"CUDA_C_32I", {"HIPBLAS_C_32I", "rocblas_datatype_i32_c", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 11 // 166
{"CUDA_R_32U", {"HIPBLAS_R_32U", "rocblas_datatype_u32_r", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 12 // 163
{"CUDA_C_32U", {"HIPBLAS_C_32U", "rocblas_datatype_u32_c", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 13 // 167
{"CUDA_R_8I", {"HIPBLAS_R_8I", "rocblas_datatype_i8_r", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 3 // 160
{"CUDA_C_8I", {"HIPBLAS_C_8I", "rocblas_datatype_i8_c", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 7 // 164
{"CUDA_R_8U", {"HIPBLAS_R_8U", "rocblas_datatype_u8_r", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 8 // 161
{"CUDA_C_8U", {"HIPBLAS_C_8U", "rocblas_datatype_u8_c", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 9 // 165
{"CUDA_R_32I", {"HIPBLAS_R_32I", "rocblas_datatype_i32_r", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 10 // 162
{"CUDA_C_32I", {"HIPBLAS_C_32I", "rocblas_datatype_i32_c", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 11 // 166
{"CUDA_R_32U", {"HIPBLAS_R_32U", "rocblas_datatype_u32_r", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 12 // 163
{"CUDA_C_32U", {"HIPBLAS_C_32U", "rocblas_datatype_u32_c", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 13 // 167
{"cublasHandle_t", {"hipblasHandle_t", "rocblas_handle", CONV_TYPE, API_BLAS}},
// TODO: dereferencing: typedef struct cublasContext *cublasHandle_t;
projects/clr/hipamd/hipify-clang/src/CUDA2HIP_Driver_API_functions.cpp
+5 -5
Näytä tiedosto
@@ -164,10 +164,10 @@ const std::map<llvm::StringRef, hipCounter> CUDA_DRIVER_FUNCTION_MAP{
{"cuMemAlloc", {"hipMalloc", "", CONV_MEMORY, API_DRIVER}},
{"cuMemAlloc_v2", {"hipMalloc", "", CONV_MEMORY, API_DRIVER}},
// cudaHostAlloc
{"cuMemAllocHost", {"hipMemAllocHost", "", CONV_MEMORY, API_DRIVER}},
{"cuMemAllocHost_v2", {"hipMemAllocHost", "", CONV_MEMORY, API_DRIVER}},
{"cuMemAllocHost", {"hipHostMalloc", "", CONV_MEMORY, API_DRIVER}},
{"cuMemAllocHost_v2", {"hipHostMalloc", "", CONV_MEMORY, API_DRIVER}},
// cudaMallocManaged
{"cuMemAllocManaged", {"hipMemAllocManaged", "", CONV_MEMORY, API_DRIVER}},
{"cuMemAllocManaged", {"hipMallocManaged", "", CONV_MEMORY, API_DRIVER}},
// no analogue
// NOTE: Not equal to cudaMallocPitch due to different signatures
{"cuMemAllocPitch", {"hipMemAllocPitch", "", CONV_MEMORY, API_DRIVER}},
@@ -266,7 +266,7 @@ const std::map<llvm::StringRef, hipCounter> CUDA_DRIVER_FUNCTION_MAP{
{"cuMemHostGetDevicePointer", {"hipHostGetDevicePointer", "", CONV_MEMORY, API_DRIVER}},
{"cuMemHostGetDevicePointer_v2", {"hipHostGetDevicePointer", "", CONV_MEMORY, API_DRIVER}},
// cudaHostGetFlags
{"cuMemHostGetFlags", {"hipMemHostGetFlags", "", CONV_MEMORY, API_DRIVER}},
{"cuMemHostGetFlags", {"hipHostGetFlags", "", CONV_MEMORY, API_DRIVER}},
// cudaHostRegister
{"cuMemHostRegister", {"hipHostRegister", "", CONV_MEMORY, API_DRIVER}},
{"cuMemHostRegister_v2", {"hipHostRegister", "", CONV_MEMORY, API_DRIVER}},
@@ -340,7 +340,7 @@ const std::map<llvm::StringRef, hipCounter> CUDA_DRIVER_FUNCTION_MAP{
{"cuPointerGetAttribute", {"hipPointerGetAttribute", "", CONV_ADDRESSING, API_DRIVER, HIP_UNSUPPORTED}},
// no analogue
// NOTE: Not equal to cudaPointerGetAttributes due to different signatures
{"cuPointerGetAttributes", {"hipPointerGetAttributes", "", CONV_ADDRESSING, API_DRIVER, HIP_UNSUPPORTED}},
{"cuPointerGetAttributes", {"hipPointerGetAttributes_", "", CONV_ADDRESSING, API_DRIVER, HIP_UNSUPPORTED}},
// no analogue
{"cuPointerSetAttribute", {"hipPointerSetAttribute", "", CONV_ADDRESSING, API_DRIVER, HIP_UNSUPPORTED}},
projects/clr/hipamd/hipify-clang/src/CUDA2HIP_Driver_API_types.cpp
+1 -41
Näytä tiedosto
@@ -1166,16 +1166,10 @@ const std::map<llvm::StringRef, hipCounter> CUDA_DRIVER_TYPE_NAME_MAP{
// CUresult enum values
// cudaSuccess
{"CUDA_SUCCESS", {"hipSuccess", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 0
// cudaErrorInvalidValue = 1, CUDA_ERROR_INVALID_VALUE = 1, hipErrorInvalidValue = 1011
// TODO [HIP]: make hipErrorInvalidValue = 1
// cudaErrorInvalidValue
{"CUDA_ERROR_INVALID_VALUE", {"hipErrorInvalidValue", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 1
// cudaErrorMemoryAllocation = 2, CUDA_ERROR_OUT_OF_MEMORY = 2, hipErrorOutOfMemory = 2, hipErrorMemoryAllocation = 1002
// TODO [HIP]: remove hipErrorMemoryAllocation
// cudaErrorMemoryAllocation
{"CUDA_ERROR_OUT_OF_MEMORY", {"hipErrorOutOfMemory", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 2
// cudaErrorInitializationError = 3, CUDA_ERROR_NOT_INITIALIZED = 3, hipErrorNotInitialized = 3, hipErrorInitializationError = 1003
// TODO [HIP]: remove hipErrorInitializationError
// cudaErrorInitializationError
{"CUDA_ERROR_NOT_INITIALIZED", {"hipErrorNotInitialized", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 3
// cudaErrorCudartUnloading
@@ -1191,12 +1185,8 @@ const std::map<llvm::StringRef, hipCounter> CUDA_DRIVER_TYPE_NAME_MAP{
// cudaErrorProfilerAlreadyStopped
// NOTE: Deprecated since CUDA 5.0
{"CUDA_ERROR_PROFILER_ALREADY_STOPPED", {"hipErrorProfilerAlreadyStopped", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 8
// cudaErrorNoDevice = 100, CUDA_ERROR_NO_DEVICE = 100, hipErrorNoDevice = 1038
// TODO [HIP]: make hipErrorNoDevice = 100
// cudaErrorNoDevice
{"CUDA_ERROR_NO_DEVICE", {"hipErrorNoDevice", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 100
// cudaErrorInvalidDevice = 101, CUDA_ERROR_INVALID_DEVICE = 101, hipErrorInvalidDevice = 1010
// TODO [HIP]: make hipErrorInvalidDevice = 101
// cudaErrorInvalidDevice
{"CUDA_ERROR_INVALID_DEVICE", {"hipErrorInvalidDevice", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 101
// cudaErrorInvalidKernelImage
@@ -1206,9 +1196,6 @@ const std::map<llvm::StringRef, hipCounter> CUDA_DRIVER_TYPE_NAME_MAP{
// no analogue
// NOTE: Deprecated since CUDA 3.2
{"CUDA_ERROR_CONTEXT_ALREADY_CURRENT", {"hipErrorContextAlreadyCurrent", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 202
// cudaErrorMapBufferObjectFailed = 205, CUDA_ERROR_MAP_FAILED = 205, hipErrorMapFailed = 205, hipErrorMapBufferObjectFailed = 1071
// TODO [HIP]: remove hipErrorMapBufferObjectFailed
// TODO [HIPIFY]: rename hipErrorMapBufferObjectFailed to hipErrorMapFailed
// cudaErrorMapBufferObjectFailed
{"CUDA_ERROR_MAP_FAILED", {"hipErrorMapFailed", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 205
// cudaErrorUnmapBufferObjectFailed
@@ -1253,57 +1240,36 @@ const std::map<llvm::StringRef, hipCounter> CUDA_DRIVER_TYPE_NAME_MAP{
{"CUDA_ERROR_SHARED_OBJECT_INIT_FAILED", {"hipErrorSharedObjectInitFailed", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 303
// cudaErrorOperatingSystem
{"CUDA_ERROR_OPERATING_SYSTEM", {"hipErrorOperatingSystem", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 304
// cudaErrorInvalidResourceHandle = 400, CUDA_ERROR_INVALID_HANDLE = 400, hipErrorInvalidHandle = 400, hipErrorInvalidResourceHandle = 1033
// TODO [HIP]: remove hipErrorInvalidResourceHandle
// TODO [HIPIFY]: rename hipErrorInvalidResourceHandle to hipErrorInvalidHandle
// cudaErrorInvalidResourceHandle
{"CUDA_ERROR_INVALID_HANDLE", {"hipErrorInvalidHandle", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 400
// cudaErrorIllegalState
{"CUDA_ERROR_ILLEGAL_STATE", {"hipErrorIllegalState", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 401
// cudaErrorSymbolNotFound
{"CUDA_ERROR_NOT_FOUND", {"hipErrorNotFound", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 500
// cudaErrorNotReady = 600, CUDA_ERROR_NOT_READY = 600, hipErrorNotReady = 1034
// TODO [HIP]: make hipErrorNotReady = 600
// cudaErrorNotReady
{"CUDA_ERROR_NOT_READY", {"hipErrorNotReady", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 600
// cudaErrorIllegalAddress
{"CUDA_ERROR_ILLEGAL_ADDRESS", {"hipErrorIllegalAddress", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 700
// cudaErrorLaunchOutOfResources = 701, CUDA_ERROR_LAUNCH_OUT_OF_RESOURCES = 701, hipErrorLaunchOutOfResources = 1007
// TODO [HIP]: make hipErrorLaunchOutOfResources = 701
// cudaErrorLaunchOutOfResources
{"CUDA_ERROR_LAUNCH_OUT_OF_RESOURCES", {"hipErrorLaunchOutOfResources", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 701
// cudaErrorLaunchTimeout = 702, CUDA_ERROR_LAUNCH_TIMEOUT = 702, hipErrorLaunchTimeOut = 1006
// TODO [HIP]: make hipErrorLaunchTimeOut = 702
// cudaErrorLaunchTimeout
{"CUDA_ERROR_LAUNCH_TIMEOUT", {"hipErrorLaunchTimeOut", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 702
// cudaErrorLaunchIncompatibleTexturing
{"CUDA_ERROR_LAUNCH_INCOMPATIBLE_TEXTURING", {"hipErrorLaunchIncompatibleTexturing", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 703
// cudaErrorPeerAccessAlreadyEnabled = 704, CUDA_ERROR_PEER_ACCESS_ALREADY_ENABLED = 704, hipErrorPeerAccessAlreadyEnabled = 1050
// TODO [HIP]: make hipErrorPeerAccessAlreadyEnabled = 704
// cudaErrorPeerAccessAlreadyEnabled
{"CUDA_ERROR_PEER_ACCESS_ALREADY_ENABLED", {"hipErrorPeerAccessAlreadyEnabled", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 704
// cudaErrorPeerAccessNotEnabled = 705, CUDA_ERROR_PEER_ACCESS_NOT_ENABLED = 705, hipErrorPeerAccessNotEnabled = 1051
// TODO [HIP]: make hipErrorPeerAccessNotEnabled = 705
// cudaErrorPeerAccessNotEnabled
{"CUDA_ERROR_PEER_ACCESS_NOT_ENABLED", {"hipErrorPeerAccessNotEnabled", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 705
// cudaErrorSetOnActiveProcess = 708, CUDA_ERROR_PRIMARY_CONTEXT_ACTIVE = 708, hipErrorSetOnActiveProcess = 305
// TODO [HIP]: make hipErrorSetOnActiveProcess = 708
// cudaErrorSetOnActiveProcess
{"CUDA_ERROR_PRIMARY_CONTEXT_ACTIVE", {"hipErrorSetOnActiveProcess", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 708
// cudaErrorContextIsDestroyed
{"CUDA_ERROR_CONTEXT_IS_DESTROYED", {"hipErrorContextIsDestroyed", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 709
// cudaErrorAssert = 710, CUDA_ERROR_ASSERT = 710, hipErrorAssert = 1081
// TODO [HIP]: make hipErrorAssert = 710
// cudaErrorAssert
{"CUDA_ERROR_ASSERT", {"hipErrorAssert", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 710
// cudaErrorTooManyPeers
{"CUDA_ERROR_TOO_MANY_PEERS", {"hipErrorTooManyPeers", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 711
// cudaErrorHostMemoryAlreadyRegistered = 712, CUDA_ERROR_HOST_MEMORY_ALREADY_REGISTERED = 712, hipErrorHostMemoryAlreadyRegistered = 1061
// TODO [HIP]: make hipErrorHostMemoryAlreadyRegistered = 712
// cudaErrorHostMemoryAlreadyRegistered
{"CUDA_ERROR_HOST_MEMORY_ALREADY_REGISTERED", {"hipErrorHostMemoryAlreadyRegistered", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 712
// cudaErrorHostMemoryNotRegistered = 713, CUDA_ERROR_HOST_MEMORY_NOT_REGISTERED = 713, hipErrorHostMemoryNotRegistered = 1062
// TODO [HIP]: make hipErrorHostMemoryNotRegistered = 713
// cudaErrorHostMemoryNotRegistered
{"CUDA_ERROR_HOST_MEMORY_NOT_REGISTERED", {"hipErrorHostMemoryNotRegistered", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 713
// cudaErrorHardwareStackError
@@ -1316,14 +1282,10 @@ const std::map<llvm::StringRef, hipCounter> CUDA_DRIVER_TYPE_NAME_MAP{
{"CUDA_ERROR_INVALID_ADDRESS_SPACE", {"hipErrorInvalidAddressSpace", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 717
// cudaErrorInvalidPc
{"CUDA_ERROR_INVALID_PC", {"hipErrorInvalidPc", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 718
// cudaErrorLaunchFailure = 719, CUDA_ERROR_LAUNCH_FAILED = 719, hipErrorLaunchFailure = 1004
// TODO [HIP]: make hipErrorSetOnActiveProcess = 719
// cudaErrorLaunchFailure
{"CUDA_ERROR_LAUNCH_FAILED", {"hipErrorLaunchFailure", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 719
// cudaErrorNotPermitted
{"CUDA_ERROR_NOT_PERMITTED", {"hipErrorNotPermitted", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 800
// cudaErrorNotSupported = 801, CUDA_ERROR_NOT_SUPPORTED = 801, hipErrorNotSupported = 1082
// TODO [HIP]: make hipErrorNotSupported = 801
// cudaErrorNotSupported
{"CUDA_ERROR_NOT_SUPPORTED", {"hipErrorNotSupported", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 801
// cudaErrorSystemNotReady
@@ -1354,8 +1316,6 @@ const std::map<llvm::StringRef, hipCounter> CUDA_DRIVER_TYPE_NAME_MAP{
{"CUDA_ERROR_TIMEOUT", {"hipErrorTimeout", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 909
// cudaErrorGraphExecUpdateFailure
{"CUDA_ERROR_GRAPH_EXEC_UPDATE_FAILURE", {"hipErrorGraphExecUpdateFailure", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 910
// cudaErrorUnknown = 999, CUDA_ERROR_UNKNOWN = 999, hipErrorUnknown = 1030
// TODO [HIP]: make hipErrorUnknown = 999
// cudaErrorUnknown
{"CUDA_ERROR_UNKNOWN", {"hipErrorUnknown", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 999
@@ -1606,7 +1566,7 @@ const std::map<llvm::StringRef, hipCounter> CUDA_DRIVER_TYPE_NAME_MAP{
// cudaHostAllocMapped
{"CU_MEMHOSTALLOC_DEVICEMAP", {"hipHostMallocMapped", "", CONV_DEFINE, API_DRIVER}}, // 0x02
// cudaHostAllocWriteCombined
{"CU_MEMHOSTALLOC_WRITECOMBINED", {"hipHostAllocWriteCombined", "", CONV_DEFINE, API_DRIVER}}, // 0x04
{"CU_MEMHOSTALLOC_WRITECOMBINED", {"hipHostMallocWriteCombined", "", CONV_DEFINE, API_DRIVER}}, // 0x04
// cudaHostRegisterPortable
{"CU_MEMHOSTREGISTER_PORTABLE", {"hipHostRegisterPortable", "", CONV_DEFINE, API_DRIVER}}, // 0x01
// cudaHostRegisterMapped
projects/clr/hipamd/hipify-clang/src/CUDA2HIP_Runtime_API_functions.cpp
+3 -4
Näytä tiedosto
@@ -254,7 +254,6 @@ const std::map<llvm::StringRef, hipCounter> CUDA_RUNTIME_FUNCTION_MAP{
// no analogue
{"cudaGetSymbolSize", {"hipGetSymbolSize", "", CONV_MEMORY, API_RUNTIME}},
// cuMemHostAlloc
// NOTE: hipHostAlloc deprecated - use hipHostMalloc instead
{"cudaHostAlloc", {"hipHostMalloc", "", CONV_MEMORY, API_RUNTIME}},
// cuMemHostGetDevicePointer
{"cudaHostGetDevicePointer", {"hipHostGetDevicePointer", "", CONV_MEMORY, API_RUNTIME}},
@@ -340,16 +339,16 @@ const std::map<llvm::StringRef, hipCounter> CUDA_RUNTIME_FUNCTION_MAP{
{"cudaMemRangeGetAttribute", {"hipMemRangeGetAttribute", "", CONV_MEMORY, API_RUNTIME, HIP_UNSUPPORTED}},
// cuMemRangeGetAttributes
{"cudaMemRangeGetAttributes", {"hipMemRangeGetAttributes", "", CONV_MEMORY, API_RUNTIME, HIP_UNSUPPORTED}},
// cuMemsetD32
// cuMemsetD32 - hipMemsetD32
{"cudaMemset", {"hipMemset", "", CONV_MEMORY, API_RUNTIME}},
// no analogue
{"cudaMemset2D", {"hipMemset2D", "", CONV_MEMORY, API_RUNTIME}},
// no analogue
{"cudaMemset2DAsync", {"hipMemset2DAsync", "", CONV_MEMORY, API_RUNTIME}},
// no analogue
{"cudaMemset3D", {"hipMemset3D", "", CONV_MEMORY, API_RUNTIME, HIP_UNSUPPORTED}},
{"cudaMemset3D", {"hipMemset3D", "", CONV_MEMORY, API_RUNTIME}},
// no analogue
{"cudaMemset3DAsync", {"hipMemset3DAsync", "", CONV_MEMORY, API_RUNTIME, HIP_UNSUPPORTED}},
{"cudaMemset3DAsync", {"hipMemset3DAsync", "", CONV_MEMORY, API_RUNTIME}},
// cuMemsetD32Async
{"cudaMemsetAsync", {"hipMemsetAsync", "", CONV_MEMORY, API_RUNTIME}},
// no analogue
projects/clr/hipamd/hipify-clang/src/CUDA2HIP_Runtime_API_types.cpp
+5 -61
Näytä tiedosto
@@ -597,20 +597,12 @@ const std::map<llvm::StringRef, hipCounter> CUDA_RUNTIME_TYPE_NAME_MAP {
// cudaError enum values
// CUDA_SUCCESS
{"cudaSuccess", {"hipSuccess", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 0
// cudaErrorInvalidValue = 1, CUDA_ERROR_INVALID_VALUE = 1, hipErrorInvalidValue = 1011
// TODO [HIP]: make hipErrorInvalidValue = 1
// CUDA_ERROR_INVALID_VALUE
{"cudaErrorInvalidValue", {"hipErrorInvalidValue", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 1
// cudaErrorMemoryAllocation = 2, CUDA_ERROR_OUT_OF_MEMORY = 2, hipErrorOutOfMemory = 2, hipErrorMemoryAllocation = 1002
// TODO [HIP]: remove hipErrorMemoryAllocation
// TODO [HIPIFY]: rename hipErrorMemoryAllocation to hipErrorOutOfMemory
// CUDA_ERROR_OUT_OF_MEMORY
{"cudaErrorMemoryAllocation", {"hipErrorMemoryAllocation", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 2
// cudaErrorInitializationError = 3, CUDA_ERROR_NOT_INITIALIZED = 3, hipErrorNotInitialized = 3, hipErrorInitializationError = 1003
// TODO [HIP]: remove hipErrorInitializationError
// TODO [HIPIFY]: rename hipErrorInitializationError to hipErrorNotInitialized
{"cudaErrorMemoryAllocation", {"hipErrorOutOfMemory", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 2
// CUDA_ERROR_NOT_INITIALIZED
{"cudaErrorInitializationError", {"hipErrorInitializationError", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 3
{"cudaErrorInitializationError", {"hipErrorNotInitialized", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 3
// CUDA_ERROR_DEINITIALIZED
{"cudaErrorCudartUnloading", {"hipErrorDeinitialized", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 4
// CUDA_ERROR_PROFILER_DISABLED
@@ -624,21 +616,15 @@ const std::map<llvm::StringRef, hipCounter> CUDA_RUNTIME_TYPE_NAME_MAP {
// Deprecated since CUDA 5.0
// CUDA_ERROR_PROFILER_ALREADY_STOPPED
{"cudaErrorProfilerAlreadyStopped", {"hipErrorProfilerAlreadyStopped", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 8
// cudaErrorInvalidConfiguration = 9, hipErrorInvalidConfiguration = 1009
// TODO [HIP]: make hipErrorInvalidConfiguration = 9
// no analogue
{"cudaErrorInvalidConfiguration", {"hipErrorInvalidConfiguration", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 9
// no analogue
{"cudaErrorInvalidPitchValue", {"hipErrorInvalidPitchValue", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 12
// cudaErrorInvalidSymbol = 13, hipErrorInvalidSymbol = 701, but also cudaErrorLaunchOutOfResources = CUDA_ERROR_LAUNCH_OUT_OF_RESOURCES = 701
// TODO [HIP]: make hipErrorInvalidSymbol = 13
// no analogue
{"cudaErrorInvalidSymbol", {"hipErrorInvalidSymbol", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 13
// Deprecated since CUDA 10.1
// no analogue
{"cudaErrorInvalidHostPointer", {"hipErrorInvalidHostPointer", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 16
// cudaErrorInvalidDevicePointer = 17, hipErrorInvalidDevicePointer = 1017
// TODO [HIP]: make hipErrorInvalidDevicePointer = 17
// Deprecated since CUDA 10.1
// no analogue
{"cudaErrorInvalidDevicePointer", {"hipErrorInvalidDevicePointer", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 17
@@ -648,8 +634,6 @@ const std::map<llvm::StringRef, hipCounter> CUDA_RUNTIME_TYPE_NAME_MAP {
{"cudaErrorInvalidTextureBinding", {"hipErrorInvalidTextureBinding", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 19
// no analogue
{"cudaErrorInvalidChannelDescriptor", {"hipErrorInvalidChannelDescriptor", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 20
// cudaErrorInvalidMemcpyDirection = 21, hipErrorInvalidMemcpyDirection = 1021
// TODO [HIP]: make hipErrorInvalidMemcpyDirection = 21
// no analogue
{"cudaErrorInvalidMemcpyDirection", {"hipErrorInvalidMemcpyDirection", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 21
// no analogue
@@ -686,12 +670,8 @@ const std::map<llvm::StringRef, hipCounter> CUDA_RUNTIME_TYPE_NAME_MAP {
{"cudaErrorDevicesUnavailable", {"hipErrorDevicesUnavailable", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 46
// no analogue
{"cudaErrorIncompatibleDriverContext", {"hipErrorIncompatibleDriverContext", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 49
// cudaErrorMissingConfiguration = 52, hipErrorMissingConfiguration = 1001
// TODO [HIP]: make hipErrorMissingConfiguration = 52
// no analogue
{"cudaErrorMissingConfiguration", {"hipErrorMissingConfiguration", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 52
// cudaErrorPriorLaunchFailure = 53, hipErrorPriorLaunchFailure = 1005
// TODO [HIP]: make hipErrorPriorLaunchFailure = 53
// no analogue
{"cudaErrorPriorLaunchFailure", {"hipErrorPriorLaunchFailure", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 53
// no analogue
@@ -704,16 +684,10 @@ const std::map<llvm::StringRef, hipCounter> CUDA_RUNTIME_TYPE_NAME_MAP {
{"cudaErrorSyncDepthExceeded", {"hipErrorSyncDepthExceeded", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 68
// no analogue
{"cudaErrorLaunchPendingCountExceeded", {"hipErrorLaunchPendingCountExceeded", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 69
// cudaErrorInvalidDeviceFunction = 98, hipErrorInvalidDeviceFunction = 1008
// TODO [HIP]: make hipErrorInvalidDeviceFunction = 98
// no analogue
{"cudaErrorInvalidDeviceFunction", {"hipErrorInvalidDeviceFunction", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 98
// cudaErrorNoDevice = 100, CUDA_ERROR_NO_DEVICE = 100, hipErrorNoDevice = 1038
// TODO [HIP]: make hipErrorNoDevice = 100
// CUDA_ERROR_NO_DEVICE
{"cudaErrorNoDevice", {"hipErrorNoDevice", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 100
// cudaErrorInvalidDevice = 101, CUDA_ERROR_INVALID_DEVICE = 101, hipErrorInvalidDevice = 1010
// TODO [HIP]: make hipErrorInvalidDevice = 101
// CUDA_ERROR_INVALID_DEVICE
{"cudaErrorInvalidDevice", {"hipErrorInvalidDevice", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 101
// no analogue
@@ -725,11 +699,8 @@ const std::map<llvm::StringRef, hipCounter> CUDA_RUNTIME_TYPE_NAME_MAP {
{"cudaErrorDeviceUninitialized", {"hipErrorInvalidContext", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 201
// CUDA_ERROR_INVALID_CONTEXT
{"cudaErrorDeviceUninitilialized", {"hipErrorInvalidContext", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 201
// cudaErrorMapBufferObjectFailed = 205, CUDA_ERROR_MAP_FAILED = 205, hipErrorMapFailed = 205, hipErrorMapBufferObjectFailed = 1071
// TODO [HIP]: remove hipErrorMapBufferObjectFailed
// TODO [HIPIFY]: rename hipErrorMapBufferObjectFailed to hipErrorMapFailed
// CUDA_ERROR_MAP_FAILED
{"cudaErrorMapBufferObjectFailed", {"hipErrorMapBufferObjectFailed", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 205
{"cudaErrorMapBufferObjectFailed", {"hipErrorMapFailed", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 205
// CUDA_ERROR_UNMAP_FAILED
{"cudaErrorUnmapBufferObjectFailed", {"hipErrorUnmapFailed", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 206
// CUDA_ERROR_ARRAY_IS_MAPPED
@@ -772,57 +743,36 @@ const std::map<llvm::StringRef, hipCounter> CUDA_RUNTIME_TYPE_NAME_MAP {
{"cudaErrorSharedObjectInitFailed", {"hipErrorSharedObjectInitFailed", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 303
// CUDA_ERROR_OPERATING_SYSTEM
{"cudaErrorOperatingSystem", {"hipErrorOperatingSystem", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 304
// cudaErrorInvalidResourceHandle = 400, CUDA_ERROR_INVALID_HANDLE = 400, hipErrorInvalidHandle = 400, hipErrorInvalidResourceHandle = 1033
// TODO [HIP]: remove hipErrorInvalidResourceHandle
// TODO [HIPIFY]: rename hipErrorInvalidResourceHandle to hipErrorInvalidHandle
// CUDA_ERROR_INVALID_HANDLE
{"cudaErrorInvalidResourceHandle", {"hipErrorInvalidResourceHandle", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 400
{"cudaErrorInvalidResourceHandle", {"hipErrorInvalidHandle", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 400
// CUDA_ERROR_ILLEGAL_STATE
{"cudaErrorIllegalState", {"hipErrorIllegalState", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 401
// CUDA_ERROR_NOT_FOUND
{"cudaErrorSymbolNotFound", {"hipErrorNotFound", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 500
// cudaErrorNotReady = 600, CUDA_ERROR_NOT_READY = 600, hipErrorNotReady = 1034
// TODO [HIP]: make hipErrorNotReady = 600
// CUDA_ERROR_NOT_READY
{"cudaErrorNotReady", {"hipErrorNotReady", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 600
// CUDA_ERROR_ILLEGAL_ADDRESS
{"cudaErrorIllegalAddress", {"hipErrorIllegalAddress", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 700
// cudaErrorLaunchOutOfResources = 701, CUDA_ERROR_LAUNCH_OUT_OF_RESOURCES = 701, hipErrorLaunchOutOfResources = 1007
// TODO [HIP]: make hipErrorLaunchOutOfResources = 701
// CUDA_ERROR_LAUNCH_OUT_OF_RESOURCES
{"cudaErrorLaunchOutOfResources", {"hipErrorLaunchOutOfResources", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 701
// cudaErrorLaunchTimeout = 702, CUDA_ERROR_LAUNCH_TIMEOUT = 702, hipErrorLaunchTimeOut = 1006
// TODO [HIP]: make hipErrorLaunchTimeOut = 702
// CUDA_ERROR_LAUNCH_TIMEOUT
{"cudaErrorLaunchTimeout", {"hipErrorLaunchTimeOut", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 702
// CUDA_ERROR_LAUNCH_INCOMPATIBLE_TEXTURING
{"cudaErrorLaunchIncompatibleTexturing", {"hipErrorLaunchIncompatibleTexturing", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 703
// cudaErrorPeerAccessAlreadyEnabled = 704, CUDA_ERROR_PEER_ACCESS_ALREADY_ENABLED = 704, hipErrorPeerAccessAlreadyEnabled = 1050
// TODO [HIP]: make hipErrorPeerAccessAlreadyEnabled = 704
// CUDA_ERROR_PEER_ACCESS_ALREADY_ENABLED
{"cudaErrorPeerAccessAlreadyEnabled", {"hipErrorPeerAccessAlreadyEnabled", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 704
// cudaErrorPeerAccessNotEnabled = 705, CUDA_ERROR_PEER_ACCESS_NOT_ENABLED = 705, hipErrorPeerAccessNotEnabled = 1051
// TODO [HIP]: make hipErrorPeerAccessNotEnabled = 705
// CUDA_ERROR_PEER_ACCESS_NOT_ENABLED
{"cudaErrorPeerAccessNotEnabled", {"hipErrorPeerAccessNotEnabled", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 705
// cudaErrorSetOnActiveProcess = 708, CUDA_ERROR_PRIMARY_CONTEXT_ACTIVE = 708, hipErrorSetOnActiveProcess = 305
// TODO [HIP]: make hipErrorSetOnActiveProcess = 708
// CUDA_ERROR_PRIMARY_CONTEXT_ACTIVE
{"cudaErrorSetOnActiveProcess", {"hipErrorSetOnActiveProcess", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 708
// CUDA_ERROR_CONTEXT_IS_DESTROYED
{"cudaErrorContextIsDestroyed", {"hipErrorContextIsDestroyed", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 709
// cudaErrorAssert = 710, CUDA_ERROR_ASSERT = 710, hipErrorAssert = 1081
// TODO [HIP]: make hipErrorAssert = 710
// CUDA_ERROR_ASSERT
{"cudaErrorAssert", {"hipErrorAssert", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 710
// CUDA_ERROR_TOO_MANY_PEERS
{"cudaErrorTooManyPeers", {"hipErrorTooManyPeers", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 711
// cudaErrorHostMemoryAlreadyRegistered = 712, CUDA_ERROR_HOST_MEMORY_ALREADY_REGISTERED = 712, hipErrorHostMemoryAlreadyRegistered = 1061
// TODO [HIP]: make hipErrorHostMemoryAlreadyRegistered = 712
// CUDA_ERROR_HOST_MEMORY_ALREADY_REGISTERED
{"cudaErrorHostMemoryAlreadyRegistered", {"hipErrorHostMemoryAlreadyRegistered", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 712
// cudaErrorHostMemoryNotRegistered = 713, CUDA_ERROR_HOST_MEMORY_NOT_REGISTERED = 713, hipErrorHostMemoryNotRegistered = 1062
// TODO [HIP]: make hipErrorHostMemoryNotRegistered = 713
// CUDA_ERROR_HOST_MEMORY_NOT_REGISTERED
{"cudaErrorHostMemoryNotRegistered", {"hipErrorHostMemoryNotRegistered", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 713
// CUDA_ERROR_HARDWARE_STACK_ERROR
@@ -835,16 +785,12 @@ const std::map<llvm::StringRef, hipCounter> CUDA_RUNTIME_TYPE_NAME_MAP {
{"cudaErrorInvalidAddressSpace", {"hipErrorInvalidAddressSpace", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 717
// CUDA_ERROR_INVALID_PC
{"cudaErrorInvalidPc", {"hipErrorInvalidPc", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 718
// cudaErrorLaunchFailure = 719, CUDA_ERROR_LAUNCH_FAILED = 719, hipErrorLaunchFailure = 1004
// TODO [HIP]: make hipErrorSetOnActiveProcess = 719
// CUDA_ERROR_LAUNCH_FAILED
{"cudaErrorLaunchFailure", {"hipErrorLaunchFailure", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 719
// no analogue
{"cudaErrorCooperativeLaunchTooLarge", {"hipErrorCooperativeLaunchTooLarge", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 720
// CUDA_ERROR_NOT_PERMITTED
{"cudaErrorNotPermitted", {"hipErrorNotPermitted", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 800
// cudaErrorNotSupported = 801, CUDA_ERROR_NOT_SUPPORTED = 801, hipErrorNotSupported = 1082
// TODO [HIP]: make hipErrorNotSupported = 801
// CUDA_ERROR_NOT_SUPPORTED
{"cudaErrorNotSupported", {"hipErrorNotSupported", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 801
// CUDA_ERROR_SYSTEM_NOT_READY
@@ -875,8 +821,6 @@ const std::map<llvm::StringRef, hipCounter> CUDA_RUNTIME_TYPE_NAME_MAP {
{"cudaErrorTimeout", {"hipErrorTimeout", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 909
// CUDA_ERROR_GRAPH_EXEC_UPDATE_FAILURE
{"cudaErrorGraphExecUpdateFailure", {"hipErrorGraphExecUpdateFailure", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 910
// cudaErrorUnknown = 999, CUDA_ERROR_UNKNOWN = 999, hipErrorUnknown = 1030
// TODO [HIP]: make hipErrorUnknown = 999
// CUDA_ERROR_UNKNOWN
{"cudaErrorUnknown", {"hipErrorUnknown", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 999
// Deprecated since CUDA 4.1
@@ -1436,7 +1380,7 @@ const std::map<llvm::StringRef, hipCounter> CUDA_RUNTIME_TYPE_NAME_MAP {
// CU_MEMHOSTALLOC_DEVICEMAP
{"cudaHostAllocMapped", {"hipHostMallocMapped", "", CONV_DEFINE, API_RUNTIME}}, // 0x02
// CU_MEMHOSTALLOC_WRITECOMBINED
{"cudaHostAllocWriteCombined", {"hipHostAllocWriteCombined", "", CONV_DEFINE, API_RUNTIME}}, // 0x04
{"cudaHostAllocWriteCombined", {"hipHostMallocWriteCombined", "", CONV_DEFINE, API_RUNTIME}}, // 0x04
// no analogue
{"cudaHostRegisterDefault", {"hipHostRegisterDefault", "", CONV_DEFINE, API_RUNTIME}}, // 0x00
// CU_MEMHOSTREGISTER_PORTABLE
projects/clr/hipamd/hipify-clang/src/main.cpp
+14 -6
Näytä tiedosto
@@ -158,7 +158,7 @@ void appendArgumentsAdjusters(ct::RefactoringTool &Tool, const std::string &sSou
Tool.appendArgumentsAdjuster(ct::getClangSyntaxOnlyAdjuster());
}
int main(int argc, const char **argv) {
void addSeparatorIfNeeded(int& argc, const char** argv) {
std::vector<const char*> new_argv(argv, argv + argc);
if (std::find(new_argv.begin(), new_argv.end(), std::string("--")) == new_argv.end()) {
new_argv.push_back("--");
@@ -166,6 +166,18 @@ int main(int argc, const char **argv) {
argv = new_argv.data();
argc++;
}
}
bool generatePython() {
bool bToRoc = TranslateToRoc;
TranslateToRoc = true;
bool bToPython = python::generate(GeneratePython);
TranslateToRoc = bToRoc;
return bToPython;
}
int main(int argc, const char **argv) {
addSeparatorIfNeeded(argc, argv);
llcompat::PrintStackTraceOnErrorSignal();
ct::CommonOptionsParser OptionsParser(argc, argv, ToolTemplateCategory, llvm::cl::ZeroOrMore);
if (!llcompat::CheckCompatibility()) {
@@ -180,11 +192,7 @@ int main(int argc, const char **argv) {
llvm::errs() << "\n" << sHipify << sError << "hipify-perl generating failed" << "\n";
return 1;
}
bool bToRoc = TranslateToRoc;
TranslateToRoc = true;
bool bToPython = python::generate(GeneratePython);
TranslateToRoc = bToRoc;
if (!bToPython) {
if (!generatePython()) {
llvm::errs() << "\n" << sHipify << sError << "hipify-python generating failed" << "\n";
return 1;
}
projects/clr/hipamd/include/hip/hcc_detail/functional_grid_launch.hpp
+30
Näytä tiedosto
@@ -37,6 +37,15 @@ THE SOFTWARE.
hipError_t ihipExtLaunchMultiKernelMultiDevice(hipLaunchParams* launchParamsList, int numDevices,
unsigned int flags, hip_impl::program_state& ps);
hipError_t ihipLaunchCooperativeKernel(const void* f, dim3 gridDim, dim3 blockDimX, void** kernelParams,
unsigned int sharedMemBytes, hipStream_t stream, hip_impl::program_state& ps);
hipError_t ihipLaunchCooperativeKernelMultiDevice(hipLaunchParams* launchParamsList, int numDevices,
unsigned int flags, hip_impl::program_state& ps);
#pragma GCC visibility push(hidden)
namespace hip_impl {
@@ -190,4 +199,25 @@ hipError_t hipExtLaunchMultiKernelMultiDevice(hipLaunchParams* launchParamsList,
}
template <typename F>
inline
__attribute__((visibility("hidden")))
hipError_t hipLaunchCooperativeKernel(F f, dim3 gridDim, dim3 blockDimX, void** kernelParams,
unsigned int sharedMemBytes, hipStream_t stream) {
hip_impl::hip_init();
auto& ps = hip_impl::get_program_state();
return ihipLaunchCooperativeKernel(reinterpret_cast<void*>(f), gridDim, blockDimX, kernelParams, sharedMemBytes, stream, ps);
}
inline
__attribute__((visibility("hidden")))
hipError_t hipLaunchCooperativeKernelMultiDevice(hipLaunchParams* launchParamsList, int numDevices,
unsigned int flags) {
hip_impl::hip_init();
auto& ps = hip_impl::get_program_state();
return ihipLaunchCooperativeKernelMultiDevice(launchParamsList, numDevices, flags, ps);
}
#pragma GCC visibility pop
projects/clr/hipamd/include/hip/hcc_detail/grid_launch.h
-2
Näytä tiedosto
@@ -48,8 +48,6 @@ typedef struct grid_launch_parm
//! Value of packet fences to apply to launch.
//! The correspond to the value of bits 9:14 in the AQL packet,
//! see HSA_PACKET_HEADER_ACQUIRE_FENCE_SCOPE and hsa_fence_scope_t.
//! Set to -1 for conservative defaults.
//! Placeholder, is not used to control packet dispatch yet
unsigned int launch_fence;
//! Pointer to the accelerator_view where the kernel should execute.
projects/clr/hipamd/include/hip/hcc_detail/hip_runtime.h
+2 -2
Näytä tiedosto
@@ -506,7 +506,7 @@ hc_get_workitem_absolute_id(int dim)
#endif
// Support std::complex.
#ifndef _OPENMP
#if !_OPENMP || __HIP_ENABLE_CUDA_WRAPPER_FOR_OPENMP__
#pragma push_macro("__CUDA__")
#define __CUDA__
#include <__clang_cuda_math_forward_declares.h>
@@ -516,7 +516,7 @@ hc_get_workitem_absolute_id(int dim)
#include <cuda_wrappers/new>
#undef __CUDA__
#pragma pop_macro("__CUDA__")
#endif // ndef _OPENMP
#endif // !_OPENMP || __HIP_ENABLE_CUDA_WRAPPER_FOR_OPENMP__
#endif // defined(__clang__) && defined(__HIP__)
projects/clr/hipamd/include/hip/hcc_detail/hip_runtime_api.h
+56 -50
Näytä tiedosto
@@ -444,7 +444,7 @@ hipError_t hipGetDeviceProperties(hipDeviceProp_t* prop, int deviceId);
*
* @param [in] cacheConfig
*
* @returns #hipSuccess, #hipErrorInitializationError
* @returns #hipSuccess, #hipErrorNotInitialized
* Note: AMD devices and some Nvidia GPUS do not support reconfigurable cache. This hint is ignored
* on those architectures.
*
@@ -457,7 +457,7 @@ hipError_t hipDeviceSetCacheConfig(hipFuncCache_t cacheConfig);
*
* @param [in] cacheConfig
*
* @returns #hipSuccess, #hipErrorInitializationError
* @returns #hipSuccess, #hipErrorNotInitialized
* Note: AMD devices and some Nvidia GPUS do not support reconfigurable cache. This hint is ignored
* on those architectures.
*
@@ -482,7 +482,7 @@ hipError_t hipDeviceGetLimit(size_t* pValue, enum hipLimit_t limit);
*
* @param [in] config;
*
* @returns #hipSuccess, #hipErrorInitializationError
* @returns #hipSuccess, #hipErrorNotInitialized
* Note: AMD devices and some Nvidia GPUS do not support reconfigurable cache. This hint is ignored
* on those architectures.
*
@@ -494,7 +494,7 @@ hipError_t hipFuncSetCacheConfig(const void* func, hipFuncCache_t config);
*
* @param [out] pConfig
*
* @returns #hipSuccess, #hipErrorInvalidValue, #hipErrorInitializationError
* @returns #hipSuccess, #hipErrorInvalidValue, #hipErrorNotInitialized
*
* Note: AMD devices and some Nvidia GPUS do not support shared cache banking, and the hint is
* ignored on those architectures.
@@ -508,7 +508,7 @@ hipError_t hipDeviceGetSharedMemConfig(hipSharedMemConfig* pConfig);
*
* @param [in] config
*
* @returns #hipSuccess, #hipErrorInvalidValue, #hipErrorInitializationError
* @returns #hipSuccess, #hipErrorInvalidValue, #hipErrorNotInitialized
*
* Note: AMD devices and some Nvidia GPUS do not support shared cache banking, and the hint is
* ignored on those architectures.
@@ -727,7 +727,7 @@ hipError_t hipDeviceGetStreamPriorityRange(int* leastPriority, int* greatestPrio
*
* @param[in, out] stream Valid pointer to hipStream_t. This function writes the memory with the
* newly created stream.
* @return #hipSuccess #hipErrorInvalidResourceHandle
* @return #hipSuccess #hipErrorInvalidHandle
*
* Destroys the specified stream.
*
@@ -749,7 +749,7 @@ hipError_t hipStreamDestroy(hipStream_t stream);
*
* @param[in] stream stream to query
*
* @return #hipSuccess, #hipErrorNotReady, #hipErrorInvalidResourceHandle
* @return #hipSuccess, #hipErrorNotReady, #hipErrorInvalidHandle
*
* This is thread-safe and returns a snapshot of the current state of the queue. However, if other
* host threads are sending work to the stream, the status may change immediately after the function
@@ -766,7 +766,7 @@ hipError_t hipStreamQuery(hipStream_t stream);
*
* @param[in] stream stream identifier.
*
* @return #hipSuccess, #hipErrorInvalidResourceHandle
* @return #hipSuccess, #hipErrorInvalidHandle
*
* This command is host-synchronous : the host will block until the specified stream is empty.
*
@@ -790,7 +790,7 @@ hipError_t hipStreamSynchronize(hipStream_t stream);
* @param[in] event event to wait on
* @param[in] flags control operation [must be 0]
*
* @return #hipSuccess, #hipErrorInvalidResourceHandle
* @return #hipSuccess, #hipErrorInvalidHandle
*
* This function inserts a wait operation into the specified stream.
* All future work submitted to @p stream will wait until @p event reports completion before
@@ -810,9 +810,9 @@ hipError_t hipStreamWaitEvent(hipStream_t stream, hipEvent_t event, unsigned int
*
* @param[in] stream stream to be queried
* @param[in,out] flags Pointer to an unsigned integer in which the stream's flags are returned
* @return #hipSuccess, #hipErrorInvalidValue, #hipErrorInvalidResourceHandle
* @return #hipSuccess, #hipErrorInvalidValue, #hipErrorInvalidHandle
*
* @returns #hipSuccess #hipErrorInvalidValue #hipErrorInvalidResourceHandle
* @returns #hipSuccess #hipErrorInvalidValue #hipErrorInvalidHandle
*
* Return flags associated with this stream in *@p flags.
*
@@ -826,9 +826,9 @@ hipError_t hipStreamGetFlags(hipStream_t stream, unsigned int* flags);
*
* @param[in] stream stream to be queried
* @param[in,out] priority Pointer to an unsigned integer in which the stream's priority is returned
* @return #hipSuccess, #hipErrorInvalidValue, #hipErrorInvalidResourceHandle
* @return #hipSuccess, #hipErrorInvalidValue, #hipErrorInvalidHandle
*
* @returns #hipSuccess #hipErrorInvalidValue #hipErrorInvalidResourceHandle
* @returns #hipSuccess #hipErrorInvalidValue #hipErrorInvalidHandle
*
* Query the priority of a stream. The priority is returned in in priority.
*
@@ -851,7 +851,7 @@ typedef void (*hipStreamCallback_t)(hipStream_t stream, hipError_t status, void*
* @param[in] callback - The function to call once preceding stream operations are complete
* @param[in] userData - User specified data to be passed to the callback function
* @param[in] flags - Reserved for future use, must be 0
* @return #hipSuccess, #hipErrorInvalidResourceHandle, #hipErrorNotSupported
* @return #hipSuccess, #hipErrorInvalidHandle, #hipErrorNotSupported
*
* @see hipStreamCreate, hipStreamCreateWithFlags, hipStreamQuery, hipStreamSynchronize,
* hipStreamWaitEvent, hipStreamDestroy, hipStreamCreateWithPriority
@@ -893,8 +893,8 @@ hipError_t hipStreamAddCallback(hipStream_t stream, hipStreamCallback_t callback
* @warning On HCC platform, hipEventInterprocess support is under development. Use of this flag
will return an error.
*
* @returns #hipSuccess, #hipErrorInitializationError, #hipErrorInvalidValue,
#hipErrorLaunchFailure, #hipErrorMemoryAllocation
* @returns #hipSuccess, #hipErrorNotInitialized, #hipErrorInvalidValue,
#hipErrorLaunchFailure, #hipErrorOutOfMemory
*
* @see hipEventCreate, hipEventSynchronize, hipEventDestroy, hipEventElapsedTime
*/
@@ -906,8 +906,8 @@ hipError_t hipEventCreateWithFlags(hipEvent_t* event, unsigned flags);
*
* @param[in,out] event Returns the newly created event.
*
* @returns #hipSuccess, #hipErrorInitializationError, #hipErrorInvalidValue,
* #hipErrorLaunchFailure, #hipErrorMemoryAllocation
* @returns #hipSuccess, #hipErrorNotInitialized, #hipErrorInvalidValue,
* #hipErrorLaunchFailure, #hipErrorOutOfMemory
*
* @see hipEventCreateWithFlags, hipEventRecord, hipEventQuery, hipEventSynchronize,
* hipEventDestroy, hipEventElapsedTime
@@ -920,8 +920,8 @@ hipError_t hipEventCreate(hipEvent_t* event);
*
* @param[in] event event to record.
* @param[in] stream stream in which to record event.
* @returns #hipSuccess, #hipErrorInvalidValue, #hipErrorInitializationError,
* #hipErrorInvalidResourceHandle, #hipErrorLaunchFailure
* @returns #hipSuccess, #hipErrorInvalidValue, #hipErrorNotInitialized,
* #hipErrorInvalidHandle, #hipErrorLaunchFailure
*
* hipEventQuery() or hipEventSynchronize() must be used to determine when the event
* transitions from "recording" (after hipEventRecord() is called) to "recorded"
@@ -952,7 +952,7 @@ hipError_t hipEventRecord(hipEvent_t event, hipStream_t stream);
* @brief Destroy the specified event.
*
* @param[in] event Event to destroy.
* @returns #hipSuccess, #hipErrorInitializationError, #hipErrorInvalidValue,
* @returns #hipSuccess, #hipErrorNotInitialized, #hipErrorInvalidValue,
* #hipErrorLaunchFailure
*
* Releases memory associated with the event. If the event is recording but has not completed
@@ -978,8 +978,8 @@ hipError_t hipEventDestroy(hipEvent_t event);
* TODO-hcc - This function needs to support hipEventBlockingSync parameter.
*
* @param[in] event Event on which to wait.
* @returns #hipSuccess, #hipErrorInvalidValue, #hipErrorInitializationError,
* #hipErrorInvalidResourceHandle, #hipErrorLaunchFailure
* @returns #hipSuccess, #hipErrorInvalidValue, #hipErrorNotInitialized,
* #hipErrorInvalidHandle, #hipErrorLaunchFailure
*
* @see hipEventCreate, hipEventCreateWithFlags, hipEventQuery, hipEventDestroy, hipEventRecord,
* hipEventElapsedTime
@@ -993,8 +993,8 @@ hipError_t hipEventSynchronize(hipEvent_t event);
* @param[out] ms : Return time between start and stop in ms.
* @param[in] start : Start event.
* @param[in] stop : Stop event.
* @returns #hipSuccess, #hipErrorInvalidValue, #hipErrorNotReady, #hipErrorInvalidResourceHandle,
* #hipErrorInitializationError, #hipErrorLaunchFailure
* @returns #hipSuccess, #hipErrorInvalidValue, #hipErrorNotReady, #hipErrorInvalidHandle,
* #hipErrorNotInitialized, #hipErrorLaunchFailure
*
* Computes the elapsed time between two events. Time is computed in ms, with
* a resolution of approximately 1 us.
@@ -1007,7 +1007,7 @@ hipError_t hipEventSynchronize(hipEvent_t event);
* commands in that stream have completed executing. Thus the time that
* the event recorded may be significantly after the host calls hipEventRecord().
*
* If hipEventRecord() has not been called on either event, then #hipErrorInvalidResourceHandle is
* If hipEventRecord() has not been called on either event, then #hipErrorInvalidHandle is
* returned. If hipEventRecord() has been called on both events, but the timestamp has not yet been
* recorded on one or both events (that is, hipEventQuery() would return #hipErrorNotReady on at
* least one of the events), then #hipErrorNotReady is returned.
@@ -1022,8 +1022,8 @@ hipError_t hipEventElapsedTime(float* ms, hipEvent_t start, hipEvent_t stop);
* @brief Query event status
*
* @param[in] event Event to query.
* @returns #hipSuccess, #hipErrorNotReady, #hipErrorInvalidResourceHandle, #hipErrorInvalidValue,
* #hipErrorInitializationError, #hipErrorLaunchFailure
* @returns #hipSuccess, #hipErrorNotReady, #hipErrorInvalidHandle, #hipErrorInvalidValue,
* #hipErrorNotInitialized, #hipErrorLaunchFailure
*
* Query the status of the specified event. This function will return #hipErrorNotReady if all
* commands in the appropriate stream (specified to hipEventRecord()) have completed. If that work
@@ -1077,7 +1077,7 @@ hipError_t hipPointerGetAttributes(hipPointerAttribute_t* attributes, const void
*
* If size is 0, no memory is allocated, *ptr returns nullptr, and hipSuccess is returned.
*
* @return #hipSuccess, #hipErrorMemoryAllocation, #hipErrorInvalidValue (bad context, null *ptr)
* @return #hipSuccess, #hipErrorOutOfMemory, #hipErrorInvalidValue (bad context, null *ptr)
*
* @see hipMallocPitch, hipFree, hipMallocArray, hipFreeArray, hipMalloc3D, hipMalloc3DArray,
* hipHostFree, hipHostMalloc
@@ -1093,7 +1093,7 @@ hipError_t hipMalloc(void** ptr, size_t size);
*
* If size is 0, no memory is allocated, *ptr returns nullptr, and hipSuccess is returned.
*
* @return #hipSuccess, #hipErrorMemoryAllocation, #hipErrorInvalidValue (bad context, null *ptr)
* @return #hipSuccess, #hipErrorOutOfMemory, #hipErrorInvalidValue (bad context, null *ptr)
*
* @see hipMallocPitch, hipFree, hipMallocArray, hipFreeArray, hipMalloc3D, hipMalloc3DArray,
* hipHostFree, hipHostMalloc
@@ -1108,7 +1108,7 @@ hipError_t hipExtMallocWithFlags(void** ptr, size_t sizeBytes, unsigned int flag
*
* If size is 0, no memory is allocated, *ptr returns nullptr, and hipSuccess is returned.
*
* @return #hipSuccess, #hipErrorMemoryAllocation
* @return #hipSuccess, #hipErrorOutOfMemory
*
* @deprecated use hipHostMalloc() instead
*/
@@ -1123,7 +1123,7 @@ hipError_t hipMallocHost(void** ptr, size_t size);
*
* If size is 0, no memory is allocated, *ptr returns nullptr, and hipSuccess is returned.
*
* @return #hipSuccess, #hipErrorMemoryAllocation
* @return #hipSuccess, #hipErrorOutOfMemory
*
* @deprecated use hipHostMalloc() instead
*/
@@ -1139,7 +1139,7 @@ hipError_t hipMemAllocHost(void** ptr, size_t size);
*
* If size is 0, no memory is allocated, *ptr returns nullptr, and hipSuccess is returned.
*
* @return #hipSuccess, #hipErrorMemoryAllocation
* @return #hipSuccess, #hipErrorOutOfMemory
*
* @see hipSetDeviceFlags, hipHostFree
*/
@@ -1152,7 +1152,7 @@ hipError_t hipHostMalloc(void** ptr, size_t size, unsigned int flags);
* @param[in] size Requested memory size
* @param[in] flags must be either hipMemAttachGlobal/hipMemAttachHost
*
* @return #hipSuccess, #hipErrorMemoryAllocation
* @return #hipSuccess, #hipErrorOutOfMemory
*/
hipError_t hipMallocManaged(void** devPtr, size_t size, unsigned int flags __dparm(0));
@@ -1165,7 +1165,7 @@ hipError_t hipMallocManaged(void** devPtr, size_t size, unsigned int flags __dpa
*
* If size is 0, no memory is allocated, *ptr returns nullptr, and hipSuccess is returned.
*
* @return #hipSuccess, #hipErrorMemoryAllocation
* @return #hipSuccess, #hipErrorOutOfMemory
*
* @deprecated use hipHostMalloc() instead
*/
@@ -1179,7 +1179,7 @@ hipError_t hipHostAlloc(void** ptr, size_t size, unsigned int flags);
* @param[in] hstPtr Host Pointer allocated through hipHostMalloc
* @param[in] flags Flags to be passed for extension
*
* @return #hipSuccess, #hipErrorInvalidValue, #hipErrorMemoryAllocation
* @return #hipSuccess, #hipErrorInvalidValue, #hipErrorOutOfMemory
*
* @see hipSetDeviceFlags, hipHostMalloc
*/
@@ -1228,7 +1228,7 @@ hipError_t hipHostGetFlags(unsigned int* flagsPtr, void* hostPtr);
* typically one of the writes will "win" and overwrite data from the other registered memory
* region.
*
* @return #hipSuccess, #hipErrorMemoryAllocation
* @return #hipSuccess, #hipErrorOutOfMemory
*
* @see hipHostUnregister, hipHostGetFlags, hipHostGetDevicePointer
*/
@@ -1361,6 +1361,9 @@ hipError_t hipHostFree(void* ptr);
*/
hipError_t hipMemcpy(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind kind);
// TODO: Add description
hipError_t hipMemcpyWithStream(void* dst, const void* src, size_t sizeBytes,
hipMemcpyKind kind, hipStream_t stream);
/**
* @brief Copy data from Host to Device
*
@@ -1909,7 +1912,7 @@ hipError_t hipMemPtrGetInfo(void* ptr, size_t* size);
* @param[in] width Requested array allocation width
* @param[in] height Requested array allocation height
* @param[in] flags Requested properties of allocated array
* @return #hipSuccess, #hipErrorMemoryAllocation
* @return #hipSuccess, #hipErrorOutOfMemory
*
* @see hipMalloc, hipMallocPitch, hipFree, hipFreeArray, hipHostMalloc, hipHostFree
*/
@@ -1925,7 +1928,7 @@ hipError_t hipMalloc3D(hipPitchedPtr* pitchedDevPtr, hipExtent extent);
* @brief Frees an array on the device.
*
* @param[in] array Pointer to array to free
* @return #hipSuccess, #hipErrorInvalidValue, #hipErrorInitializationError
* @return #hipSuccess, #hipErrorInvalidValue, #hipErrorNotInitialized
*
* @see hipMalloc, hipMallocPitch, hipFree, hipMallocArray, hipHostMalloc, hipHostFree
*/
@@ -1938,7 +1941,7 @@ hipError_t hipFreeArray(hipArray* array);
* @param[in] desc Requested channel format
* @param[in] extent Requested array allocation width, height and depth
* @param[in] flags Requested properties of allocated array
* @return #hipSuccess, #hipErrorMemoryAllocation
* @return #hipSuccess, #hipErrorOutOfMemory
*
* @see hipMalloc, hipMallocPitch, hipFree, hipFreeArray, hipHostMalloc, hipHostFree
*/
@@ -2886,6 +2889,7 @@ hipError_t hipModuleLaunchKernel(hipFunction_t f, unsigned int gridDimX, unsigne
void** kernelParams, void** extra);
#if __HIP_VDI__ && !defined(__HCC__)
/**
* @brief launches kernel f with launch parameters and shared memory on stream with arguments passed
* to kernelparams or extra, where thread blocks can cooperate and synchronize as they execute
@@ -2918,6 +2922,8 @@ hipError_t hipLaunchCooperativeKernel(const void* f, dim3 gridDim, dim3 blockDim
hipError_t hipLaunchCooperativeKernelMultiDevice(hipLaunchParams* launchParamsList,
int numDevices, unsigned int flags);
#endif
/**
* @brief determine the grid and block sizes to achieves maximum occupancy for a kernel
*
@@ -2956,7 +2962,7 @@ hipError_t hipOccupancyMaxActiveBlocksPerMultiprocessor(
hipError_t hipOccupancyMaxActiveBlocksPerMultiprocessorWithFlags(
uint32_t* numBlocks, hipFunction_t f, uint32_t blockSize, size_t dynSharedMemPerBlk, unsigned int flags);
#if __HIP_VDI__
#if __HIP_VDI__ && !defined(__HCC__)
/**
* @brief Launches kernels on multiple devices and guarantees all specified kernels are dispatched
* on respective streams before enqueuing any other work on the specified streams from any other threads
@@ -3036,9 +3042,9 @@ hipError_t hipProfilerStop();
*
* @returns
* hipSuccess,
* hipErrorInvalidResourceHandle,
* hipErrorMemoryAllocation,
* hipErrorMapBufferObjectFailed,
* hipErrorInvalidHandle,
* hipErrorOutOfMemory,
* hipErrorMapFailed,
*
*/
hipError_t hipIpcGetMemHandle(hipIpcMemHandle_t* handle, void* devPtr);
@@ -3071,8 +3077,8 @@ hipError_t hipIpcGetMemHandle(hipIpcMemHandle_t* handle, void* devPtr);
*
* @returns
* hipSuccess,
* hipErrorMapBufferObjectFailed,
* hipErrorInvalidResourceHandle,
* hipErrorMapFailed,
* hipErrorInvalidHandle,
* hipErrorTooManyPeers
*
* @note No guarantees are made about the address returned in @p *devPtr.
@@ -3095,8 +3101,8 @@ hipError_t hipIpcOpenMemHandle(void** devPtr, hipIpcMemHandle_t handle, unsigned
*
* @returns
* hipSuccess,
* hipErrorMapBufferObjectFailed,
* hipErrorInvalidResourceHandle,
* hipErrorMapFailed,
* hipErrorInvalidHandle,
*
*/
hipError_t hipIpcCloseMemHandle(void* devPtr);
@@ -3379,7 +3385,7 @@ hipError_t hipBindTextureToMipmappedArray(const texture<T, dim, readMode>& tex,
return hipSuccess;
}
#if __HIP_VDI__ && !defined(__HCC__)
template <class T>
inline hipError_t hipLaunchCooperativeKernel(T f, dim3 gridDim, dim3 blockDim,
void** kernelParams, unsigned int sharedMemBytes, hipStream_t stream) {
@@ -3393,7 +3399,7 @@ inline hipError_t hipLaunchCooperativeKernelMultiDevice(hipLaunchParams* launchP
return hipLaunchCooperativeKernelMultiDevice(launchParamsList, numDevices, flags);
}
#if defined(__clang__) && defined(__HIP__)
template <class T>
inline hipError_t hipExtLaunchMultiKernelMultiDevice(hipLaunchParams* launchParamsList,
unsigned int numDevices, unsigned int flags = 0) {
projects/clr/hipamd/include/hip/hcc_detail/hip_vector_types.h
+5 -3
Näytä tiedosto
@@ -34,12 +34,12 @@ THE SOFTWARE.
#include "hip/hcc_detail/host_defines.h"
#if !defined(_MSC_VER)
#if !defined(_MSC_VER)|| __clang__
#if __has_attribute(ext_vector_type)
#define __NATIVE_VECTOR__(n, ...) __attribute__((ext_vector_type(n)))
#else
#define __NATIVE_VECTOR__(n, ...) [n]
#endif
#endif
#if defined(__cplusplus)
#include <array>
@@ -109,7 +109,8 @@ THE SOFTWARE.
operator T() const noexcept { return data[idx]; }
__host__ __device__
operator T() const volatile noexcept { return data[idx]; }
#ifdef __HIP_ENABLE_VECTOR_SCALAR_ACCESSORY_ENUM_CONVERSION__
// The conversions to enum are fairly ghastly, but unfortunately used in
// some pre-existing, difficult to modify, code.
template<
@@ -130,6 +131,7 @@ THE SOFTWARE.
T, typename std::enable_if<std::is_enum<U>::value, std::underlying_type<U>>::type::type>{}>::type* = nullptr>
__host__ __device__
operator U() const volatile noexcept { return static_cast<U>(data[idx]); }
#endif
__host__ __device__
operator T&() noexcept {
projects/clr/hipamd/include/hip/hcc_detail/hiprtc.h
+5 -4
Näytä tiedosto
@@ -84,7 +84,7 @@ hiprtcResult hiprtcGetCodeSize(hiprtcProgram prog, std::size_t* codeSizeRet);
namespace hip_impl
{
std::string demangle(const char* mangled_expression);
char* demangle(const char* mangled_expression);
}
#if defined(HIPRTC_GET_TYPE_NAME)
@@ -102,10 +102,11 @@ namespace hip_impl
{
if (!result) return HIPRTC_ERROR_INVALID_INPUT;
*result = hip_impl::demangle(typeid(T).name());
char * res= hip_impl::demangle(typeid(T).name());
result->assign(res == nullptr ? "" : res);
std::free(res);
return (result->empty()) ? HIPRTC_ERROR_INTERNAL_ERROR :
HIPRTC_SUCCESS;
}
#endif
#endif
#endif
projects/clr/hipamd/include/hip/hcc_detail/texture_functions.h
+14 -12
Näytä tiedosto
@@ -157,8 +157,6 @@ union TData {
#define TEXTURE_RETURN_UINT_XYZW return make_uint4(texel.u.x, texel.u.y, texel.u.z, texel.u.w);
#define HIP_AD_FORMAT_NOT_INITIALIZED 0
#define TEXTURE_RETURN_FLOAT return (texFormatToSize[texRef.format] == 1)? texel.f.x : (float)texel.u.x/texFormatToSize[texRef.format];
#define TEXTURE_RETURN_FLOAT_X return (texFormatToSize[texRef.format] == 1)? make_float1(texel.f.x) : make_float1((float)texel.u.x/texFormatToSize[texRef.format]);
@@ -169,16 +167,20 @@ union TData {
extern "C" {
__device__ __constant__ static int texFormatToSize[] = {
[HIP_AD_FORMAT_NOT_INITIALIZED] = 1 ,
[HIP_AD_FORMAT_UNSIGNED_INT8] = UCHAR_MAX ,
[HIP_AD_FORMAT_UNSIGNED_INT16]= USHRT_MAX,
[HIP_AD_FORMAT_UNSIGNED_INT32]= 1 ,
[HIP_AD_FORMAT_SIGNED_INT8] = SCHAR_MAX,
[HIP_AD_FORMAT_SIGNED_INT16] = SHRT_MAX,
[HIP_AD_FORMAT_SIGNED_INT32] = 1 ,
[HIP_AD_FORMAT_HALF] = 1 ,
[HIP_AD_FORMAT_FLOAT] = 1
// this is really a sparse array with only valid values being the ones indexed by the enum hipArray_Format(e.g. texFormatToSize[HIP_AD_FORMAT_UNSIGNED_INT8] = UCHAR_MAX)
__device__ __constant__ static int texFormatToSize[] = {
1, /* HIP_AD_FORMAT_NOT_INITIALIZED */
UCHAR_MAX, /* HIP_AD_FORMAT_UNSIGNED_INT8 */
USHRT_MAX, /* HIP_AD_FORMAT_UNSIGNED_INT16 */
1, /* HIP_AD_FORMAT_UNSIGNED_INT32 */
1,1,1,1, /* Invalid values */
SCHAR_MAX, /* HIP_AD_FORMAT_SIGNED_INT8 */
SHRT_MAX, /* HIP_AD_FORMAT_SIGNED_INT16 */
1, /* HIP_AD_FORMAT_SIGNED_INT32 */
1,1,1,1,1, /* Invalid values */
1, /* HIP_AD_FORMAT_HALF */
1,1,1,1,1,1,1,1,1,1,1,1,1,1,1, /* Invalid values */
1 /* HIP_AD_FORMAT_FLOAT */
};
__device__
projects/clr/hipamd/include/hip/hip_runtime.h
+1 -1
Näytä tiedosto
@@ -49,7 +49,7 @@ THE SOFTWARE.
#include <thread>
#endif
#include <hip/hip_version.h>
#include <hip/hip_common.h>
#if defined(__HIP_PLATFORM_HCC__) && !defined(__HIP_PLATFORM_NVCC__)
projects/clr/hipamd/include/hip/hip_runtime_api.h
+38 -38
Näytä tiedosto
@@ -183,18 +183,35 @@ typedef struct hipPointerAttribute_t {
typedef enum __HIP_NODISCARD hipError_t {
hipSuccess = 0, ///< Successful completion.
hipErrorInvalidValue = 1, ///< One or more of the parameters passed to the API call is NULL
///< or not in an acceptable range.
hipErrorOutOfMemory = 2,
// Deprecated
hipErrorMemoryAllocation = 2, ///< Memory allocation error.
hipErrorNotInitialized = 3,
// Deprecated
hipErrorInitializationError = 3,
hipErrorDeinitialized = 4,
hipErrorProfilerDisabled = 5,
hipErrorProfilerNotInitialized = 6,
hipErrorProfilerAlreadyStarted = 7,
hipErrorProfilerAlreadyStopped = 8,
hipErrorInvalidConfiguration = 9,
hipErrorInvalidSymbol = 13,
hipErrorInvalidDevicePointer = 17, ///< Invalid Device Pointer
hipErrorInvalidMemcpyDirection = 21, ///< Invalid memory copy direction
hipErrorInsufficientDriver = 35,
hipErrorMissingConfiguration = 52,
hipErrorPriorLaunchFailure = 53,
hipErrorInvalidDeviceFunction = 98,
hipErrorNoDevice = 100, ///< Call to hipGetDeviceCount returned 0 devices
hipErrorInvalidDevice = 101, ///< DeviceID must be in range 0...#compute-devices.
hipErrorInvalidImage = 200,
hipErrorInvalidContext = 201, ///< Produced when input context is invalid.
hipErrorContextAlreadyCurrent = 202,
hipErrorMapFailed = 205,
// Deprecated
hipErrorMapBufferObjectFailed = 205, ///< Produced when the IPC memory attach failed from ROCr.
hipErrorUnmapFailed = 206,
hipErrorArrayIsMapped = 207,
hipErrorAlreadyMapped = 208,
@@ -214,53 +231,36 @@ typedef enum __HIP_NODISCARD hipError_t {
hipErrorSharedObjectSymbolNotFound = 302,
hipErrorSharedObjectInitFailed = 303,
hipErrorOperatingSystem = 304,
hipErrorSetOnActiveProcess = 305,
hipErrorInvalidHandle = 400,
// Deprecated
hipErrorInvalidResourceHandle = 400, ///< Resource handle (hipEvent_t or hipStream_t) invalid.
hipErrorNotFound = 500,
hipErrorNotReady = 600, ///< Indicates that asynchronous operations enqueued earlier are not
///< ready. This is not actually an error, but is used to distinguish
///< from hipSuccess (which indicates completion). APIs that return
///< this error include hipEventQuery and hipStreamQuery.
hipErrorIllegalAddress = 700,
hipErrorInvalidSymbol = 701,
// Runtime Error Codes start here.
hipErrorMissingConfiguration = 1001,
hipErrorMemoryAllocation = 1002, ///< Memory allocation error.
hipErrorInitializationError = 1003, ///< TODO comment from hipErrorInitializationError
hipErrorLaunchFailure =
1004, ///< An exception occurred on the device while executing a kernel.
hipErrorPriorLaunchFailure = 1005,
hipErrorLaunchTimeOut = 1006,
hipErrorLaunchOutOfResources = 1007, ///< Out of resources error.
hipErrorInvalidDeviceFunction = 1008,
hipErrorInvalidConfiguration = 1009,
hipErrorInvalidDevice = 1010, ///< DeviceID must be in range 0...#compute-devices.
hipErrorInvalidValue = 1011, ///< One or more of the parameters passed to the API call is NULL
///< or not in an acceptable range.
hipErrorInvalidDevicePointer = 1017, ///< Invalid Device Pointer
hipErrorInvalidMemcpyDirection = 1021, ///< Invalid memory copy direction
hipErrorUnknown = 1030, ///< Unknown error.
hipErrorInvalidResourceHandle = 1033, ///< Resource handle (hipEvent_t or hipStream_t) invalid.
hipErrorNotReady = 1034, ///< Indicates that asynchronous operations enqueued earlier are not
///< ready. This is not actually an error, but is used to distinguish
///< from hipSuccess (which indicates completion). APIs that return
///< this error include hipEventQuery and hipStreamQuery.
hipErrorNoDevice = 1038, ///< Call to hipGetDeviceCount returned 0 devices
hipErrorLaunchOutOfResources = 701, ///< Out of resources error.
hipErrorLaunchTimeOut = 702,
hipErrorPeerAccessAlreadyEnabled =
1050, ///< Peer access was already enabled from the current device.
704, ///< Peer access was already enabled from the current device.
hipErrorPeerAccessNotEnabled =
1051, ///< Peer access was never enabled from the current device.
705, ///< Peer access was never enabled from the current device.
hipErrorSetOnActiveProcess = 708,
hipErrorAssert = 710, ///< Produced when the kernel calls assert.
hipErrorHostMemoryAlreadyRegistered =
712, ///< Produced when trying to lock a page-locked memory.
hipErrorHostMemoryNotRegistered =
713, ///< Produced when trying to unlock a non-page-locked memory.
hipErrorLaunchFailure =
719, ///< An exception occurred on the device while executing a kernel.
hipErrorNotSupported = 801, ///< Produced when the hip API is not supported/implemented
hipErrorUnknown = 999, //< Unknown error.
// HSA Runtime Error Codes start here.
hipErrorRuntimeMemory = 1052, ///< HSA runtime memory call returned error. Typically not seen
///< in production systems.
hipErrorRuntimeOther = 1053, ///< HSA runtime call other than memory returned error. Typically
///< not seen in production systems.
hipErrorHostMemoryAlreadyRegistered =
1061, ///< Produced when trying to lock a page-locked memory.
hipErrorHostMemoryNotRegistered =
1062, ///< Produced when trying to unlock a non-page-locked memory.
hipErrorMapBufferObjectFailed =
1071, ///< Produced when the IPC memory attach failed from ROCr.
hipErrorAssert =
1081, ///< Produced when the kernel calls assert.
hipErrorNotSupported =
1082, ///< Produced when the hip API is not supported/implemented
hipErrorTbd ///< Marker that more error codes are needed.
} hipError_t;
projects/clr/hipamd/include/hip/nvcc_detail/hip_runtime_api.h
+30 -8
Näytä tiedosto
@@ -262,9 +262,9 @@ inline static hipError_t hipCUDAErrorTohipError(cudaError_t cuError) {
case cudaErrorMissingConfiguration:
return hipErrorMissingConfiguration;
case cudaErrorMemoryAllocation:
return hipErrorMemoryAllocation;
return hipErrorOutOfMemory;
case cudaErrorInitializationError:
return hipErrorInitializationError;
return hipErrorNotInitialized;
case cudaErrorLaunchFailure:
return hipErrorLaunchFailure;
case cudaErrorPriorLaunchFailure:
@@ -286,7 +286,7 @@ inline static hipError_t hipCUDAErrorTohipError(cudaError_t cuError) {
case cudaErrorUnknown:
return hipErrorUnknown;
case cudaErrorInvalidResourceHandle:
return hipErrorInvalidResourceHandle;
return hipErrorInvalidHandle;
case cudaErrorNotReady:
return hipErrorNotReady;
case cudaErrorNoDevice:
@@ -300,7 +300,7 @@ inline static hipError_t hipCUDAErrorTohipError(cudaError_t cuError) {
case cudaErrorHostMemoryNotRegistered:
return hipErrorHostMemoryNotRegistered;
case cudaErrorMapBufferObjectFailed:
return hipErrorMapBufferObjectFailed;
return hipErrorMapFailed;
case cudaErrorAssert:
return hipErrorAssert;
case cudaErrorNotSupported:
@@ -315,7 +315,7 @@ inline static hipError_t hipCUResultTohipError(CUresult cuError) { // TODO Popu
case CUDA_SUCCESS:
return hipSuccess;
case CUDA_ERROR_OUT_OF_MEMORY:
return hipErrorMemoryAllocation;
return hipErrorOutOfMemory;
case CUDA_ERROR_INVALID_VALUE:
return hipErrorInvalidValue;
case CUDA_ERROR_INVALID_DEVICE:
@@ -328,6 +328,10 @@ inline static hipError_t hipCUResultTohipError(CUresult cuError) { // TODO Popu
return hipErrorInvalidContext;
case CUDA_ERROR_NOT_INITIALIZED:
return hipErrorNotInitialized;
case CUDA_ERROR_INVALID_HANDLE:
return hipErrorInvalidHandle;
case CUDA_ERROR_MAP_FAILED:
return hipErrorMapFailed;
default:
return hipErrorUnknown; // Note - translated error.
}
@@ -338,13 +342,13 @@ inline static cudaError_t hipErrorToCudaError(hipError_t hError) {
switch (hError) {
case hipSuccess:
return cudaSuccess;
case hipErrorMemoryAllocation:
case hipErrorOutOfMemory:
return cudaErrorMemoryAllocation;
case hipErrorLaunchOutOfResources:
return cudaErrorLaunchOutOfResources;
case hipErrorInvalidValue:
return cudaErrorInvalidValue;
case hipErrorInvalidResourceHandle:
case hipErrorInvalidHandle:
return cudaErrorInvalidResourceHandle;
case hipErrorInvalidDevice:
return cudaErrorInvalidDevice;
@@ -352,7 +356,7 @@ inline static cudaError_t hipErrorToCudaError(hipError_t hError) {
return cudaErrorInvalidMemcpyDirection;
case hipErrorInvalidDevicePointer:
return cudaErrorInvalidDevicePointer;
case hipErrorInitializationError:
case hipErrorNotInitialized:
return cudaErrorInitializationError;
case hipErrorNoDevice:
return cudaErrorNoDevice;
@@ -611,6 +615,18 @@ inline static hipError_t hipMemcpy(void* dst, const void* src, size_t sizeBytes,
}
inline hipError_t hipMemcpyWithStream(void* dst, const void* src,
size_t sizeBytes, hipMemcpyKind copyKind,
hipStream_t stream) {
cudaError_t error = cudaMemcpyAsync(dst, src, sizeBytes,
hipMemcpyKindToCudaMemcpyKind(copyKind),
stream);
if (error != cudaSuccess) return hipCUDAErrorTohipError(error);
return hipCUDAErrorTohipError(cudaStreamSynchronize(stream));
}
inline static hipError_t hipMemcpyAsync(void* dst, const void* src, size_t sizeBytes,
hipMemcpyKind copyKind, hipStream_t stream __dparm(0)) {
return hipCUDAErrorTohipError(
@@ -1367,6 +1383,12 @@ inline static hipError_t hipBindTexture(size_t* offset, struct textureReference*
return hipCUDAErrorTohipError(cudaBindTexture(offset, tex, devPtr, desc, size));
}
inline static hipError_t hipBindTexture2D(size_t* offset, struct textureReference* tex, const void* devPtr,
const hipChannelFormatDesc* desc, size_t width, size_t height,
size_t pitch) {
return hipCUDAErrorTohipError(cudaBindTexture2D(offset, tex, devPtr, desc, width, height, pitch));
}
inline static hipChannelFormatDesc hipCreateChannelDesc(int x, int y, int z, int w,
hipChannelFormatKind f) {
return cudaCreateChannelDesc(x, y, z, w, hipChannelFormatKindToCudaChannelFormatKind(f));
projects/clr/hipamd/samples/0_Intro/bit_extract/bit_extract.cpp
+2 -2
Näytä tiedosto
@@ -63,9 +63,9 @@ int main(int argc, char* argv[]) {
printf("info: allocate host mem (%6.2f MB)\n", 2 * Nbytes / 1024.0 / 1024.0);
A_h = (uint32_t*)malloc(Nbytes);
CHECK(A_h == 0 ? hipErrorMemoryAllocation : hipSuccess);
CHECK(A_h == 0 ? hipErrorOutOfMemory : hipSuccess);
C_h = (uint32_t*)malloc(Nbytes);
CHECK(C_h == 0 ? hipErrorMemoryAllocation : hipSuccess);
CHECK(C_h == 0 ? hipErrorOutOfMemory : hipSuccess);
for (size_t i = 0; i < N; i++) {
A_h[i] = i;
projects/clr/hipamd/samples/0_Intro/square/square.hipref.cpp
+2 -2
Näytä tiedosto
@@ -62,9 +62,9 @@ int main(int argc, char* argv[]) {
#endif
printf("info: allocate host mem (%6.2f MB)\n", 2 * Nbytes / 1024.0 / 1024.0);
A_h = (float*)malloc(Nbytes);
CHECK(A_h == 0 ? hipErrorMemoryAllocation : hipSuccess);
CHECK(A_h == 0 ? hipErrorOutOfMemory : hipSuccess);
C_h = (float*)malloc(Nbytes);
CHECK(C_h == 0 ? hipErrorMemoryAllocation : hipSuccess);
CHECK(C_h == 0 ? hipErrorOutOfMemory : hipSuccess);
// Fill with Phi + i
for (size_t i = 0; i < N; i++) {
A_h[i] = 1.618f + i;
projects/clr/hipamd/src/hip_event.cpp
+9 -9
Näytä tiedosto
@@ -103,10 +103,10 @@ hipError_t hipEventCreate(hipEvent_t* event) {
hipError_t hipEventRecord(hipEvent_t event, hipStream_t stream) {
HIP_INIT_SPECIAL_API(hipEventRecord, TRACE_SYNC, event, stream);
if (!event) return ihipLogStatus(hipErrorInvalidResourceHandle);
if (!event) return ihipLogStatus(hipErrorInvalidHandle);
stream = ihipSyncAndResolveStream(stream);
LockedAccessor_EventCrit_t eCrit(event->criticalData());
if (eCrit->_eventData._state == hipEventStatusUnitialized) return ihipLogStatus(hipErrorInvalidResourceHandle);
if (eCrit->_eventData._state == hipEventStatusUnitialized) return ihipLogStatus(hipErrorInvalidHandle);
if (HIP_SYNC_NULL_STREAM && stream->isDefaultStream()) {
// TODO-HIP_SYNC_NULL_STREAM : can remove this code when HIP_SYNC_NULL_STREAM = 0
// If default stream , then wait on all queues.
@@ -136,7 +136,7 @@ hipError_t hipEventDestroy(hipEvent_t event) {
return ihipLogStatus(hipSuccess);
} else {
return ihipLogStatus(hipErrorInvalidResourceHandle);
return ihipLogStatus(hipErrorInvalidHandle);
}
}
@@ -152,7 +152,7 @@ hipError_t hipEventSynchronize(hipEvent_t event) {
auto ecd = event->locked_copyCrit();
if (ecd._state == hipEventStatusUnitialized) {
return ihipLogStatus(hipErrorInvalidResourceHandle);
return ihipLogStatus(hipErrorInvalidHandle);
} else if (ecd._state == hipEventStatusCreated) {
// Created but not actually recorded on any device:
return ihipLogStatus(hipSuccess);
@@ -167,7 +167,7 @@ hipError_t hipEventSynchronize(hipEvent_t event) {
return ihipLogStatus(hipSuccess);
}
} else {
return ihipLogStatus(hipErrorInvalidResourceHandle);
return ihipLogStatus(hipErrorInvalidHandle);
}
}
@@ -175,7 +175,7 @@ hipError_t hipEventElapsedTime(float* ms, hipEvent_t start, hipEvent_t stop) {
HIP_INIT_API(hipEventElapsedTime, ms, start, stop);
if (ms == nullptr) return ihipLogStatus(hipErrorInvalidValue);
if ((start == nullptr) || (stop == nullptr)) return ihipLogStatus(hipErrorInvalidResourceHandle);
if ((start == nullptr) || (stop == nullptr)) return ihipLogStatus(hipErrorInvalidHandle);
*ms = 0.0f;
auto startEcd = start->locked_copyCrit();
@@ -187,8 +187,8 @@ hipError_t hipEventElapsedTime(float* ms, hipEvent_t start, hipEvent_t stop) {
(stop->_flags & hipEventDisableTiming) ||
(stopEcd._state == hipEventStatusUnitialized) ||
(stopEcd._state == hipEventStatusCreated)) {
// Both events must be at least recorded else return hipErrorInvalidResourceHandle
return ihipLogStatus(hipErrorInvalidResourceHandle);
// Both events must be at least recorded else return hipErrorInvalidHandle
return ihipLogStatus(hipErrorInvalidHandle);
}
// Refresh status, if still recording...
@@ -222,7 +222,7 @@ hipError_t hipEventElapsedTime(float* ms, hipEvent_t start, hipEvent_t stop) {
hipError_t hipEventQuery(hipEvent_t event) {
HIP_INIT_SPECIAL_API(hipEventQuery, TRACE_QUERY, event);
if (!event) return ihipLogStatus(hipErrorInvalidResourceHandle);
if (!event) return ihipLogStatus(hipErrorInvalidHandle);
if (!(event->_flags & hipEventReleaseToSystem)) {
tprintf(DB_WARN,
projects/clr/hipamd/src/hip_hcc.cpp
+48 -23
Näytä tiedosto
@@ -304,18 +304,28 @@ void ihipStream_t::wait(LockedAccessor_StreamCrit_t& crit) {
//---
// Wait for all kernel and data copy commands in this stream to complete.
void ihipStream_t::locked_wait() {
inline void ihipStream_t::locked_wait(bool& waited) {
// create a marker while holding stream lock,
// but release lock prior to waiting on the marker
hc::completion_future marker;
{
LockedAccessor_StreamCrit_t crit(_criticalData);
// skipping marker since stream is empty
if (crit->_av.get_is_empty()) return;
if (crit->_av.get_is_empty()) {
waited = false;
return;
}
marker = crit->_av.create_marker(hc::no_scope);
}
marker.wait(waitMode());
waited = true;
return;
};
void ihipStream_t::locked_wait() {
bool waited;
locked_wait(waited);
};
// Causes current stream to wait for specified event to complete:
@@ -1040,6 +1050,7 @@ void ihipCtx_t::locked_syncDefaultStream(bool waitOnSelf, bool syncHost) {
// Vector of ops sent to each stream that will complete before ops sent to null stream:
std::vector<hc::completion_future> depOps;
bool last_stream_waited = false;
for (auto streamI = crit->const_streams().begin(); streamI != crit->const_streams().end();
streamI++) {
ihipStream_t* stream = *streamI;
@@ -1051,7 +1062,8 @@ void ihipCtx_t::locked_syncDefaultStream(bool waitOnSelf, bool syncHost) {
if (HIP_SYNC_NULL_STREAM) {
if (waitThisStream) {
stream->locked_wait();
last_stream_waited = false;
stream->locked_wait(last_stream_waited);
}
} else {
if (waitThisStream) {
@@ -1082,6 +1094,14 @@ void ihipCtx_t::locked_syncDefaultStream(bool waitOnSelf, bool syncHost) {
defaultCf.wait(); // TODO - account for active or blocking here.
}
}
else if ( (HIP_SYNC_NULL_STREAM && !last_stream_waited) ||
(!HIP_SYNC_NULL_STREAM && depOps.empty()) ) {
// This catches all the conditions where the printf buffer
// need to be explicitly flushed
if (syncHost) {
Kalmar::getContext()->flushPrintfBuffer();
}
}
tprintf(DB_SYNC, " syncDefaultStream depOps=%zu\n", depOps.size());
}
@@ -1101,9 +1121,18 @@ void ihipCtx_t::locked_waitAllStreams() {
LockedAccessor_CtxCrit_t crit(_criticalData);
tprintf(DB_SYNC, "waitAllStream\n");
bool need_printf_flush = false;
for (auto streamI = crit->const_streams().begin(); streamI != crit->const_streams().end();
streamI++) {
(*streamI)->locked_wait();
bool waited;
(*streamI)->locked_wait(waited);
need_printf_flush = !waited;
}
// When synchronizing with the last stream, if we didn't do an explicit wait,
// then we need to an extra flush to the printf buffer
if (need_printf_flush) {
Kalmar::getContext()->flushPrintfBuffer();
}
}
@@ -1492,7 +1521,11 @@ hipError_t ihipStreamSynchronize(TlsData *tls, hipStream_t stream) {
ctx->locked_syncDefaultStream(true /*waitOnSelf*/, true /*syncToHost*/);
} else {
// note this does not synchornize with the NULL stream:
stream->locked_wait();
bool waited;
stream->locked_wait(waited);
if (!waited) {
Kalmar::getContext()->flushPrintfBuffer();
}
e = hipSuccess;
}
@@ -1619,16 +1652,18 @@ hipStream_t ihipPreLaunchKernel(hipStream_t stream, dim3 grid, dim3 block, grid_
lp->group_dim.y = block.y;
lp->group_dim.z = block.z;
lp->barrier_bit = barrier_bit_queue_default;
lp->launch_fence = -1;
if (!lockAcquired) {
auto crit = stream->lockopen_preKernelCommand();
lp->av = &(crit->_av);
} else {
// this stream is already locked (e.g., call from hipExtLaunchMultiKernelMultiDevice)
lp->av = &(stream->criticalData()._av);
}
if (!lockAcquired) stream->lockopen_preKernelCommand();
auto &crit = stream->criticalData();
lp->av = &(crit._av);
lp->cf = nullptr;
auto acq = (HCC_OPT_FLUSH && !crit._last_op_was_a_copy) ?
HSA_FENCE_SCOPE_AGENT : HSA_FENCE_SCOPE_SYSTEM;
auto rel = HCC_OPT_FLUSH ?
HSA_FENCE_SCOPE_AGENT : HSA_FENCE_SCOPE_SYSTEM;
lp->launch_fence = (acq << HSA_PACKET_HEADER_SCACQUIRE_FENCE_SCOPE) |
(rel << HSA_PACKET_HEADER_SCRELEASE_FENCE_SCOPE);
crit._last_op_was_a_copy = false;
ihipPrintKernelLaunch(kernelNameStr, lp, stream);
return (stream);
@@ -1756,13 +1791,8 @@ const char* ihipErrorString(hipError_t hip_error) {
return "hipErrorIllegalAddress";
case hipErrorInvalidSymbol:
return "hipErrorInvalidSymbol";
case hipErrorMissingConfiguration:
return "hipErrorMissingConfiguration";
case hipErrorMemoryAllocation:
return "hipErrorMemoryAllocation";
case hipErrorInitializationError:
return "hipErrorInitializationError";
case hipErrorLaunchFailure:
return "hipErrorLaunchFailure";
case hipErrorPriorLaunchFailure:
@@ -1785,15 +1815,12 @@ const char* ihipErrorString(hipError_t hip_error) {
return "hipErrorInvalidMemcpyDirection";
case hipErrorUnknown:
return "hipErrorUnknown";
case hipErrorInvalidResourceHandle:
return "hipErrorInvalidResourceHandle";
case hipErrorNotReady:
return "hipErrorNotReady";
case hipErrorNoDevice:
return "hipErrorNoDevice";
case hipErrorPeerAccessAlreadyEnabled:
return "hipErrorPeerAccessAlreadyEnabled";
case hipErrorPeerAccessNotEnabled:
return "hipErrorPeerAccessNotEnabled";
case hipErrorRuntimeMemory:
@@ -1804,8 +1831,6 @@ const char* ihipErrorString(hipError_t hip_error) {
return "hipErrorHostMemoryAlreadyRegistered";
case hipErrorHostMemoryNotRegistered:
return "hipErrorHostMemoryNotRegistered";
case hipErrorMapBufferObjectFailed:
return "hipErrorMapBufferObjectFailed";
case hipErrorAssert:
return "hipErrorAssert";
case hipErrorNotSupported:
projects/clr/hipamd/src/hip_hcc_internal.h
+5 -7
Näytä tiedosto
@@ -495,9 +495,10 @@ struct LockedBase {
template <typename MUTEX_TYPE>
class ihipStreamCriticalBase_t : public LockedBase<MUTEX_TYPE> {
public:
public:
ihipStreamCriticalBase_t(ihipStream_t* parentStream, hc::accelerator_view av)
: _av(av), _parent(parentStream){};
: _parent{parentStream}, _av{av}, _last_op_was_a_copy{false}
{}
~ihipStreamCriticalBase_t() {}
@@ -519,12 +520,9 @@ class ihipStreamCriticalBase_t : public LockedBase<MUTEX_TYPE> {
return gotLock ? this : nullptr;
};
public:
ihipStream_t* _parent;
hc::accelerator_view _av;
private:
bool _last_op_was_a_copy;
};
@@ -572,7 +570,7 @@ class ihipStream_t {
LockedAccessor_StreamCrit_t lockopen_preKernelCommand();
void lockclose_postKernelCommand(const char* kernelName, hc::accelerator_view* av, bool unlockNotNeeded = 0);
void locked_wait(bool& waited);
void locked_wait();
hc::accelerator_view* locked_getAv() {
projects/clr/hipamd/src/hip_memory.cpp
+484 -149
Näytä tiedosto
@@ -27,6 +27,7 @@ THE SOFTWARE.
#include "hip_hcc_internal.h"
#include "trace_helper.h"
#include <functional>
#include <fstream>
__device__ char __hip_device_heap[__HIP_SIZE_OF_HEAP];
@@ -35,23 +36,372 @@ __device__ uint32_t __hip_device_page_flag[__HIP_NUM_PAGES];
// Internal HIP APIS:
namespace hip_internal {
hipError_t memcpyAsync(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind kind,
hipStream_t stream) {
hipError_t e = hipSuccess;
namespace {
inline
const char* hsa_to_string(hsa_status_t err) noexcept
{
const char* r{};
// Return success if number of bytes to copy is 0
if (sizeBytes == 0) return e;
if (hsa_status_string(err, &r) == HSA_STATUS_SUCCESS) return r;
return "Unknown.";
}
template<std::size_t m, std::size_t n>
inline
void throwing_result_check(hsa_status_t res, const char (&file)[m],
const char (&function)[n], int line) {
if (res == HSA_STATUS_SUCCESS) return;
if (res == HSA_STATUS_INFO_BREAK) return;
throw std::runtime_error{"Failed in file " + (file +
(", in function \"" + (function +
("\", on line " + std::to_string(line))))) +
", with error: " + hsa_to_string(res)};
}
inline
hsa_agent_t cpu_agent() {
hsa_agent_t r{};
throwing_result_check(hsa_iterate_agents([](hsa_agent_t x, void* pr) {
hsa_device_type_t t{};
hsa_agent_get_info(x, HSA_AGENT_INFO_DEVICE, &t);
if (t != HSA_DEVICE_TYPE_CPU) return HSA_STATUS_SUCCESS;
*static_cast<hsa_agent_t *>(pr) = x;
return HSA_STATUS_INFO_BREAK;
}, &r), __FILE__, __func__, __LINE__);
return r;
}
inline
hsa_device_type_t type(hsa_agent_t x)
{
hsa_device_type_t r{};
throwing_result_check(hsa_agent_get_info(x, HSA_AGENT_INFO_DEVICE, &r),
__FILE__, __func__, __LINE__);
return r;
}
const auto is_large_BAR{[](){
std::unique_ptr<void, void (*)(void*)> hsa{
(hsa_init() == HSA_STATUS_SUCCESS)
? reinterpret_cast<void*>(UINT64_MAX) : nullptr,
[](void* p) { if (p) hsa_shut_down(); }};
if (!hsa) return false;
bool r{true};
throwing_result_check(hsa_iterate_agents([](hsa_agent_t x, void* pr) {
if (x.handle == cpu_agent().handle) return HSA_STATUS_SUCCESS;
throwing_result_check(
hsa_agent_iterate_regions(x, [](hsa_region_t y, void* p) {
hsa_region_segment_t seg{};
throwing_result_check(
hsa_region_get_info(y, HSA_REGION_INFO_SEGMENT, &seg),
__FILE__, __func__, __LINE__);
if (seg != HSA_REGION_SEGMENT_GLOBAL) {
return HSA_STATUS_SUCCESS;
}
uint32_t flags{};
throwing_result_check(hsa_region_get_info(
y, HSA_REGION_INFO_GLOBAL_FLAGS, &flags),
__FILE__, __func__, __LINE__);
if (flags & HSA_REGION_GLOBAL_FLAG_COARSE_GRAINED) {
hsa_amd_memory_pool_access_t tmp{};
throwing_result_check(
hsa_amd_agent_memory_pool_get_info(
cpu_agent(),
hsa_amd_memory_pool_t{y.handle},
HSA_AMD_AGENT_MEMORY_POOL_INFO_ACCESS,
&tmp),
__FILE__, __func__, __LINE__);
*static_cast<bool*>(p) &=
tmp != HSA_AMD_MEMORY_POOL_ACCESS_NEVER_ALLOWED;
}
return HSA_STATUS_SUCCESS;
}, pr), __FILE__, __func__, __LINE__);
return HSA_STATUS_SUCCESS;
}, &r), __FILE__, __func__, __LINE__);
return r;
}()};
inline
hsa_amd_pointer_info_t info(const void* p)
{
hsa_amd_pointer_info_t r{sizeof(hsa_amd_pointer_info_t)};
throwing_result_check(
hsa_amd_pointer_info(
const_cast<void*>(p), &r, nullptr, nullptr, nullptr),
__FILE__, __func__, __LINE__);
r.size = is_large_BAR || (type(r.agentOwner) == HSA_DEVICE_TYPE_CPU) ?
UINT32_MAX : sizeof(hsa_amd_pointer_info_t);
return r;
}
constexpr size_t staging_sz{4 * 1024 * 1024}; // 2 Pages.
thread_local const std::unique_ptr<void, void (*)(void *)> staging_buffer{
[]() {
hsa_region_t r{};
throwing_result_check(hsa_agent_iterate_regions(
cpu_agent(), [](hsa_region_t x, void *p) {
hsa_region_segment_t seg{};
throwing_result_check(
hsa_region_get_info(x, HSA_REGION_INFO_SEGMENT, &seg),
__FILE__, __func__, __LINE__);
if (seg != HSA_REGION_SEGMENT_GLOBAL) return HSA_STATUS_SUCCESS;
uint32_t flags{};
throwing_result_check(hsa_region_get_info(
x, HSA_REGION_INFO_GLOBAL_FLAGS, &flags),
__FILE__, __func__, __LINE__);
if (flags & HSA_REGION_GLOBAL_FLAG_COARSE_GRAINED) {
*static_cast<hsa_region_t *>(p) = x;
return HSA_STATUS_INFO_BREAK;
}
return HSA_STATUS_SUCCESS;
}, &r), __FILE__, __func__, __LINE__);
void *tp{};
throwing_result_check(hsa_memory_allocate(r, staging_sz, &tp),
__FILE__, __func__, __LINE__);
return tp;
}(),
[](void *ptr) { hsa_memory_free(ptr); }};
thread_local hsa_signal_t copy_signal{[]() {
hsa_agent_t cpu{cpu_agent()};
hsa_signal_t sgn{};
throwing_result_check(hsa_signal_create(1, 1, &cpu, &sgn),
__FILE__, __func__, __LINE__);
return sgn;
}()};
} // Unnamed namespace.
inline
void do_copy(void* __restrict dst, const void* __restrict src, std::size_t n,
hsa_agent_t da, hsa_agent_t sa) {
hsa_signal_silent_store_relaxed(copy_signal, 1);
throwing_result_check(
hsa_amd_memory_async_copy(dst, da, src, sa, n, 0, nullptr, copy_signal),
__FILE__, __func__, __LINE__);
while (hsa_signal_wait_relaxed(copy_signal, HSA_SIGNAL_CONDITION_EQ, 0,
UINT64_MAX, HSA_WAIT_STATE_ACTIVE));
}
inline
void do_std_memcpy(
void* __restrict dst, const void* __restrict src, std::size_t n) {
std::memcpy(dst, src, n);
return std::atomic_thread_fence(std::memory_order_seq_cst);
}
inline
void d2h_copy(void* __restrict dst, const void* __restrict src, size_t n,
hsa_amd_pointer_info_t si) {
// TODO: characterise direct largeBAR reads from agent-allocated memory.
// if (si.size == UINT32_MAX) {
// return do_std_memcpy(dst, src, n);
// }
const auto di{info(dst)};
if (di.type == HSA_EXT_POINTER_TYPE_LOCKED) {
dst = static_cast<char*>(di.agentBaseAddress) +
(static_cast<char*>(dst) -
static_cast<char*>(di.hostBaseAddress));
do_copy(dst, src, n, si.agentOwner, si.agentOwner);
}
else if (n <= staging_sz) {
do_copy(staging_buffer.get(), src, n, si.agentOwner, si.agentOwner);
std::memcpy(dst, staging_buffer.get(), n);
}
else {
std::unique_ptr<void, void (*)(void*)> lck{
dst, [](void* p) { hsa_amd_memory_unlock(p); }};
throwing_result_check(hsa_amd_memory_lock(dst, n, &si.agentOwner, 1,
const_cast<void**>(&dst)),
__FILE__, __func__, __LINE__);
do_copy(dst, src, n, si.agentOwner, si.agentOwner);
}
}
inline
void h2d_copy(void* __restrict dst, const void* __restrict src, size_t n,
hsa_amd_pointer_info_t di) {
if (di.size == UINT32_MAX) {
return do_std_memcpy(dst, src, n);
}
const auto si{info(const_cast<void*>(src))};
if (si.type == HSA_EXT_POINTER_TYPE_LOCKED) {
src = static_cast<char*>(si.agentBaseAddress) +
(static_cast<const char*>(src) -
static_cast<char*>(si.hostBaseAddress));
do_copy(dst, src, n, di.agentOwner, di.agentOwner);
}
else if (n <= staging_sz) {
std::memcpy(staging_buffer.get(), src, n);
do_copy(dst, staging_buffer.get(), n, di.agentOwner, di.agentOwner);
}
else {
std::unique_ptr<void, void (*)(void*)> lck{
const_cast<void*>(src), [](void* p) { hsa_amd_memory_unlock(p); }};
throwing_result_check(hsa_amd_memory_lock(const_cast<void*>(src), n,
&di.agentOwner, 1,
const_cast<void**>(&src)),
__FILE__, __func__, __LINE__);
do_copy(dst, src, n, di.agentOwner, di.agentOwner);
}
}
inline
void generic_copy(void* __restrict dst, const void* __restrict src, size_t n,
hsa_amd_pointer_info_t di, hsa_amd_pointer_info_t si) {
if (di.size == UINT32_MAX && si.size == UINT32_MAX) {
return do_std_memcpy(dst, src, n);
}
std::unique_ptr<void, void (*)(void*)> lck0{
nullptr, [](void* p) { hsa_amd_memory_unlock(p); }};
std::unique_ptr<void, void (*)(void*)> lck1{nullptr, lck0.get_deleter()};
switch (si.type) {
case HSA_EXT_POINTER_TYPE_HSA:
if (di.type == HSA_EXT_POINTER_TYPE_HSA) {
hsa_memory_copy(dst, src, n);
return; // TODO: do_copy(dst, src, n, di.agentOwner, si.agentOwner);
}
if (di.type == HSA_EXT_POINTER_TYPE_UNKNOWN ||
di.type == HSA_EXT_POINTER_TYPE_LOCKED) {
return d2h_copy(dst, src, n, si);
}
break;
case HSA_EXT_POINTER_TYPE_LOCKED:
if (di.type == HSA_EXT_POINTER_TYPE_UNKNOWN) {
std::memcpy(dst, si.hostBaseAddress, n);
return;
}
if (di.type == HSA_EXT_POINTER_TYPE_LOCKED) {
std::memcpy(di.hostBaseAddress, si.hostBaseAddress, n);
return;
}
src = si.agentBaseAddress;
si.agentOwner = di.agentOwner;
break;
case HSA_EXT_POINTER_TYPE_UNKNOWN:
if (di.type == HSA_EXT_POINTER_TYPE_UNKNOWN) {
std::memcpy(dst, src, n);
return;
}
if (di.type == HSA_EXT_POINTER_TYPE_LOCKED) {
std::memcpy(di.hostBaseAddress, src, n);
return;
}
return h2d_copy(dst, src, n, di);
default: do_copy(dst, src, n, di.agentOwner, si.agentOwner); break;
}
}
inline
void memcpy_impl(void* __restrict dst, const void* __restrict src, size_t n,
hipMemcpyKind k) noexcept {
switch (k) {
case hipMemcpyHostToHost: std::memcpy(dst, src, n); break;
case hipMemcpyHostToDevice:
return is_large_BAR ? do_std_memcpy(dst, src, n)
: h2d_copy(dst, src, n, info(dst));
case hipMemcpyDeviceToHost:
// TODO: characterise direct largeBAR reads from agent-allocated memory.
return /*is_large_BAR ? do_std_memcpy(dst, src, n)
: */d2h_copy(dst, src, n, info(src));
case hipMemcpyDeviceToDevice: hsa_memory_copy(dst, src, n); break;
default: return generic_copy(dst, src, n, info(dst), info(src));
}
}
hipError_t memcpyAsync(void* dst, const void* src, size_t sizeBytes,
hipMemcpyKind kind, hipStream_t stream) {
if (sizeBytes == 0) return hipSuccess;
if (!dst || !src) return hipErrorInvalidValue;
if (!(stream = ihipSyncAndResolveStream(stream))) {
return hipErrorInvalidValue;
}
try {
stream = ihipSyncAndResolveStream(stream);
if (!stream) return hipErrorInvalidValue;
stream->locked_copyAsync(dst, src, sizeBytes, kind);
}
catch (ihipException& ex) {
e = ex._code;
catch (const ihipException& ex) {
return ex._code;
}
catch (const std::exception& ex) {
std::cerr << ex.what() << std::endl;
throw;
}
catch (...) {
return hipErrorUnknown;
}
return hipSuccess;
}
hipError_t memcpySync(void* dst, const void* src, size_t sizeBytes,
hipMemcpyKind kind, hipStream_t stream) {
if (sizeBytes == 0) return hipSuccess;
if (!dst || !src) return hipErrorInvalidValue;
try {
stream = ihipSyncAndResolveStream(stream);
if (!stream) return hipErrorInvalidValue;
LockedAccessor_StreamCrit_t cs{stream->criticalData()};
cs->_av.wait();
memcpy_impl(dst, src, sizeBytes, kind);
cs->_last_op_was_a_copy = true;
}
catch (const ihipException& ex) {
return ex._code;
}
catch (const std::exception& ex) {
std::cerr << ex.what() << std::endl;
throw;
}
catch (...) {
return hipErrorUnknown;
@@ -199,7 +549,7 @@ hipError_t ihipHostMalloc(TlsData *tls, void** ptr, size_t sizeBytes, unsigned i
true /*shareWithAll*/, amFlags, flags, 0);
if (sizeBytes && (*ptr == NULL)) {
hip_status = hipErrorMemoryAllocation;
hip_status = hipErrorOutOfMemory;
}
}
}
@@ -328,7 +678,7 @@ hipError_t hipHostGetDevicePointer(void** devicePointer, void* hostPointer, unsi
tprintf(DB_MEM, " host_ptr=%p returned device_pointer=%p\n", hostPointer,
*devicePointer);
} else {
e = hipErrorMemoryAllocation;
e = hipErrorOutOfMemory;
}
}
return ihipLogStatus(e);
@@ -354,7 +704,7 @@ hipError_t hipMalloc(void** ptr, size_t sizeBytes) {
0 /*amFlags*/, 0 /*hipFlags*/, 0);
if (sizeBytes && (*ptr == NULL)) {
hip_status = hipErrorMemoryAllocation;
hip_status = hipErrorOutOfMemory;
}
}
@@ -389,11 +739,11 @@ hipError_t hipExtMallocWithFlags(void** ptr, size_t sizeBytes, unsigned int flag
amFlags /*amFlags*/, 0 /*hipFlags*/, 0);
if (sizeBytes && (*ptr == NULL)) {
hip_status = hipErrorMemoryAllocation;
hip_status = hipErrorOutOfMemory;
}
}
#else
hipError_t hip_status = hipErrorMemoryAllocation;
hipError_t hip_status = hipErrorOutOfMemory;
#endif
return ihipLogStatus(hip_status);
@@ -436,7 +786,7 @@ hipError_t allocImage(TlsData* tls,hsa_ext_image_geometry_t geometry, int width,
hc::accelerator acc = ctx->getDevice()->_acc;
hsa_agent_t* agent = static_cast<hsa_agent_t*>(acc.get_hsa_agent());
if (!agent)
return hipErrorInvalidResourceHandle;
return hipErrorInvalidHandle;
size_t allocGranularity = 0;
hsa_amd_memory_pool_t* allocRegion = static_cast<hsa_amd_memory_pool_t*>(acc.get_hsa_am_region());
hsa_amd_memory_pool_get_info(*allocRegion, HSA_AMD_MEMORY_POOL_INFO_RUNTIME_ALLOC_GRANULE, &allocGranularity);
@@ -461,12 +811,12 @@ hipError_t allocImage(TlsData* tls,hsa_ext_image_geometry_t geometry, int width,
*ptr = hip_internal::allocAndSharePtr("device_array", imageInfo.size, ctx,
false /*shareWithAll*/, am_flags, 0, alignment);
if (*ptr == NULL) {
return hipErrorMemoryAllocation;
return hipErrorOutOfMemory;
}
return hipSuccess;
}
else {
return hipErrorMemoryAllocation;
return hipErrorOutOfMemory;
}
}
@@ -485,7 +835,7 @@ hipError_t ihipMallocPitch(TlsData* tls, void** ptr, size_t* pitch, size_t width
HSA_EXT_IMAGE_CHANNEL_TYPE_UNSIGNED_INT32,ptr,imageInfo);
if(hip_status == hipSuccess)
*pitch = imageInfo.size/(height == 0 ? 1:height)/(depth == 0 ? 1:depth);
*pitch = imageInfo.size/(height == 0 ? 1 : height)/(depth == 0 ? 1 : depth);
return hip_status;
}
@@ -565,7 +915,7 @@ hipError_t GetImageInfo(hsa_ext_image_geometry_t geometry,int width, int height,
hc::accelerator acc;
hsa_agent_t* agent = static_cast<hsa_agent_t*>(acc.get_hsa_agent());
if (!agent)
return hipErrorInvalidResourceHandle;
return hipErrorInvalidHandle;
hsa_status_t status =
hsa_ext_image_data_get_info_with_layout(*agent, &imageDescriptor, permission, HSA_EXT_IMAGE_DATA_LAYOUT_LINEAR, 0, 0, &imageInfo);
if(HSA_STATUS_SUCCESS != status){
@@ -575,6 +925,31 @@ hipError_t GetImageInfo(hsa_ext_image_geometry_t geometry,int width, int height,
return hipSuccess;
}
hipError_t GetImageInfo(hsa_ext_image_geometry_t geometry,size_t width, size_t height, size_t depth, hsa_ext_image_channel_order_t channelOrder, hsa_ext_image_channel_type_t channelType, hsa_ext_image_data_info_t &imageInfo,size_t array_size __dparm(0))
{
hsa_ext_image_descriptor_t imageDescriptor;
imageDescriptor.geometry = geometry;
imageDescriptor.width = width;
imageDescriptor.height = height;
imageDescriptor.depth = depth;
imageDescriptor.array_size = array_size;
imageDescriptor.format.channel_order = channelOrder;
imageDescriptor.format.channel_type = channelType;
// Get the current device agent.
hc::accelerator acc;
hsa_agent_t* agent = static_cast<hsa_agent_t*>(acc.get_hsa_agent());
if (!agent)
return hipErrorInvalidResourceHandle;
hsa_access_permission_t permission = HSA_ACCESS_PERMISSION_RW;
hsa_status_t status =
hsa_ext_image_data_get_info_with_layout(*agent, &imageDescriptor, permission, HSA_EXT_IMAGE_DATA_LAYOUT_LINEAR, 0, 0, &imageInfo);
if(HSA_STATUS_SUCCESS != status){
return hipErrorRuntimeOther;
}
return hipSuccess;
}
hipError_t ihipArrayToImageFormat(hipArray_Format format,hsa_ext_image_channel_type_t &channelType) {
switch (format) {
case HIP_AD_FORMAT_UNSIGNED_INT8:
@@ -720,6 +1095,7 @@ hipError_t hipArray3DCreate(hipArray** array, const HIP_ARRAY3D_DESCRIPTOR* pAll
case hipArrayDefault:
case hipArrayCubemap:
default:
array[0]->type = hipArrayCubemap;
hip_status = allocImage(tls,HSA_EXT_IMAGE_GEOMETRY_3D,pAllocateArray->Width,pAllocateArray->Height,
pAllocateArray->Depth,channelOrder,channelType,ptr,imageInfo);
array[0]->textureType = hipTextureType3D;
@@ -763,6 +1139,7 @@ hipError_t hipMalloc3DArray(hipArray** array, const struct hipChannelFormatDesc*
case hipArrayDefault:
case hipArrayCubemap:
default:
array[0]->type = hipArrayCubemap;
hip_status = allocImage(tls,HSA_EXT_IMAGE_GEOMETRY_3D,extent.width,extent.height,extent.depth,channelOrder,channelType,ptr,imageInfo);
array[0]->textureType = hipTextureType3D;
break;
@@ -860,7 +1237,7 @@ hipError_t hipHostRegister(void* hostPtr, size_t sizeBytes, unsigned int flags)
if (am_status == AM_SUCCESS) {
hip_status = hipSuccess;
} else {
hip_status = hipErrorMemoryAllocation;
hip_status = hipErrorOutOfMemory;
}
} else {
hip_status = hipErrorInvalidValue;
@@ -900,16 +1277,14 @@ hipError_t hipMemcpyToSymbol(void* dst, const void* src, size_t count,
return ihipLogStatus(hipErrorInvalidSymbol);
}
hipStream_t stream = ihipSyncAndResolveStream(hipStreamNull);
if (kind == hipMemcpyHostToDevice || kind == hipMemcpyDefault ||
kind == hipMemcpyDeviceToDevice || kind == hipMemcpyHostToHost) {
stream->locked_copySync((char*)dst+offset, (void*)src, count, kind, false);
} else {
return ihipLogStatus(hipErrorInvalidValue);
if (kind == hipMemcpyDeviceToHost || kind == hipMemcpyHostToHost) {
return ihipLogStatus(hipErrorInvalidMemcpyDirection);
} else if (kind == hipMemcpyDeviceToDevice) {
return ihipLogStatus(hipErrorInvalidValue);
}
return ihipLogStatus(hipSuccess);
return ihipLogStatus(hip_internal::memcpySync(static_cast<char*>(dst)+offset, src, count, kind,
hipStreamNull));
}
hipError_t hipMemcpyFromSymbol(void* dst, const void* src, size_t count,
@@ -920,20 +1295,18 @@ hipError_t hipMemcpyFromSymbol(void* dst, const void* src, size_t count,
tprintf(DB_MEM, " symbol '%s' resolved to address:%p\n", symbol_name, dst);
if (dst == nullptr) {
if (src == nullptr || dst == nullptr) {
return ihipLogStatus(hipErrorInvalidSymbol);
}
hipStream_t stream = ihipSyncAndResolveStream(hipStreamNull);
if (kind == hipMemcpyDefault || kind == hipMemcpyDeviceToHost ||
kind == hipMemcpyDeviceToDevice || kind == hipMemcpyHostToHost) {
stream->locked_copySync((void*)dst, (char*)src+offset, count, kind, false);
} else {
return ihipLogStatus(hipErrorInvalidValue);
if (kind == hipMemcpyHostToDevice || kind == hipMemcpyHostToHost) {
return ihipLogStatus(hipErrorInvalidMemcpyDirection);
} else if (kind == hipMemcpyDeviceToDevice) {
return ihipLogStatus(hipErrorInvalidValue);
}
return ihipLogStatus(hipSuccess);
return ihipLogStatus(hip_internal::memcpySync(dst, static_cast<const char*>(src)+offset, count, kind,
hipStreamNull));
}
@@ -948,11 +1321,17 @@ hipError_t hipMemcpyToSymbolAsync(void* dst, const void* src, size_t count,
if (dst == nullptr) {
return ihipLogStatus(hipErrorInvalidSymbol);
}
if (kind == hipMemcpyDeviceToHost || kind == hipMemcpyHostToHost) {
return ihipLogStatus(hipErrorInvalidMemcpyDirection);
} else if (kind == hipMemcpyDeviceToDevice) {
return ihipLogStatus(hipErrorInvalidValue);
}
hipError_t e = hipSuccess;
if (stream) {
try {
hip_internal::memcpyAsync((char*)dst+offset, src, count, kind, stream);
hip_internal::memcpyAsync(static_cast<char*>(dst)+offset, src, count, kind, stream);
} catch (ihipException& ex) {
e = ex._code;
}
@@ -974,12 +1353,18 @@ hipError_t hipMemcpyFromSymbolAsync(void* dst, const void* src, size_t count,
if (src == nullptr || dst == nullptr) {
return ihipLogStatus(hipErrorInvalidSymbol);
}
if (kind == hipMemcpyHostToDevice || kind == hipMemcpyHostToHost) {
return ihipLogStatus(hipErrorInvalidMemcpyDirection);
} else if (kind == hipMemcpyDeviceToDevice) {
return ihipLogStatus(hipErrorInvalidValue);
}
hipError_t e = hipSuccess;
stream = ihipSyncAndResolveStream(stream);
if (stream) {
try {
hip_internal::memcpyAsync(dst, (char*)src+offset, count, kind, stream);
hip_internal::memcpyAsync(dst, static_cast<const char*>(src)+offset, count, kind, stream);
} catch (ihipException& ex) {
e = ex._code;
}
@@ -995,120 +1380,49 @@ hipError_t hipMemcpyFromSymbolAsync(void* dst, const void* src, size_t count,
hipError_t hipMemcpy(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind kind) {
HIP_INIT_SPECIAL_API(hipMemcpy, (TRACE_MCMD), dst, src, sizeBytes, kind);
hipError_t e = hipSuccess;
// Return success if number of bytes to copy is 0
if (sizeBytes == 0) return ihipLogStatus(e);
hipStream_t stream = ihipSyncAndResolveStream(hipStreamNull);
hc::completion_future marker;
if(dst==NULL || src==NULL)
{
e=hipErrorInvalidValue;
return ihipLogStatus(e);
}
try {
stream->locked_copySync(dst, src, sizeBytes, kind);
} catch (ihipException& ex) {
e = ex._code;
}
return ihipLogStatus(e);
return ihipLogStatus(hip_internal::memcpySync(dst, src, sizeBytes, kind,
hipStreamNull));
}
hipError_t hipMemcpyHtoD(hipDeviceptr_t dst, void* src, size_t sizeBytes) {
HIP_INIT_SPECIAL_API(hipMemcpyHtoD, (TRACE_MCMD), dst, src, sizeBytes);
hipError_t e = hipSuccess;
if (sizeBytes == 0) return ihipLogStatus(e);
if(dst==NULL || src==NULL){
return ihipLogStatus(hipErrorInvalidValue);
}
hipStream_t stream = ihipSyncAndResolveStream(hipStreamNull);
hc::completion_future marker;
try {
stream->locked_copySync((void*)dst, (void*)src, sizeBytes, hipMemcpyHostToDevice, false);
} catch (ihipException& ex) {
e = ex._code;
}
return ihipLogStatus(e);
return ihipLogStatus(hip_internal::memcpySync(dst, src, sizeBytes,
hipMemcpyHostToDevice,
hipStreamNull));
}
hipError_t hipMemcpyDtoH(void* dst, hipDeviceptr_t src, size_t sizeBytes) {
HIP_INIT_SPECIAL_API(hipMemcpyDtoH, (TRACE_MCMD), dst, src, sizeBytes);
hipError_t e = hipSuccess;
if (sizeBytes == 0) return ihipLogStatus(e);
if(dst==NULL || src==NULL){
return ihipLogStatus(hipErrorInvalidValue);
}
hipStream_t stream = ihipSyncAndResolveStream(hipStreamNull);
hc::completion_future marker;
try {
stream->locked_copySync((void*)dst, (void*)src, sizeBytes, hipMemcpyDeviceToHost, false);
} catch (ihipException& ex) {
e = ex._code;
}
return ihipLogStatus(e);
return ihipLogStatus(hip_internal::memcpySync(dst, src, sizeBytes,
hipMemcpyDeviceToHost,
hipStreamNull));
}
hipError_t hipMemcpyDtoD(hipDeviceptr_t dst, hipDeviceptr_t src, size_t sizeBytes) {
HIP_INIT_SPECIAL_API(hipMemcpyDtoD, (TRACE_MCMD), dst, src, sizeBytes);
hipError_t e = hipSuccess;
if (sizeBytes == 0) return ihipLogStatus(e);
if(dst==NULL || src==NULL){
return ihipLogStatus(hipErrorInvalidValue);
}
hipStream_t stream = ihipSyncAndResolveStream(hipStreamNull);
hc::completion_future marker;
try {
stream->locked_copySync((void*)dst, (void*)src, sizeBytes, hipMemcpyDeviceToDevice, false);
} catch (ihipException& ex) {
e = ex._code;
}
return ihipLogStatus(e);
return ihipLogStatus(hip_internal::memcpySync(dst, src, sizeBytes,
hipMemcpyDeviceToDevice,
hipStreamNull));
}
hipError_t hipMemcpyHtoH(void* dst, void* src, size_t sizeBytes) {
HIP_INIT_SPECIAL_API(hipMemcpyHtoH, (TRACE_MCMD), dst, src, sizeBytes);
hipError_t e = hipSuccess;
if (sizeBytes == 0) return ihipLogStatus(e);
return ihipLogStatus(hip_internal::memcpySync(dst, src, sizeBytes,
hipMemcpyHostToHost,
hipStreamNull));
}
if(dst==NULL || src==NULL){
return ihipLogStatus(hipErrorInvalidValue);
}
hipError_t hipMemcpyWithStream(void* dst, const void* src, size_t sizeBytes,
hipMemcpyKind kind, hipStream_t stream) {
HIP_INIT_SPECIAL_API(hipMemcpyWithStream, (TRACE_MCMD), dst, src, sizeBytes,
kind, stream);
hipStream_t stream = ihipSyncAndResolveStream(hipStreamNull);
hc::completion_future marker;
try {
stream->locked_copySync((void*)dst, (void*)src, sizeBytes, hipMemcpyHostToHost, false);
} catch (ihipException& ex) {
e = ex._code;
}
return ihipLogStatus(e);
return ihipLogStatus(hip_internal::memcpySync(dst, src, sizeBytes, kind,
stream));
}
hipError_t hipMemcpyAsync(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind kind,
@@ -1291,25 +1605,46 @@ hipError_t ihipMemcpy3D(const struct hipMemcpy3DParms* p, hipStream_t stream, bo
ySize = p->srcPtr.ysize;
desc = p->dstArray->desc;
dstPtr = p->dstArray->data;
hsa_ext_image_data_info_t imageInfo;
if(hipTextureType2DLayered == p->dstArray->textureType)
GetImageInfo(HSA_EXT_IMAGE_GEOMETRY_2DA, width, height, 0, desc, imageInfo, depth);
else
GetImageInfo(HSA_EXT_IMAGE_GEOMETRY_3D, width, height, depth, desc, imageInfo);
dstPitch = imageInfo.size/(height == 0 ? 1 : height)/(depth == 0 ? 1 : depth);
} else {
depth = p->Depth;
height = p->Height;
widthInBytes = p->WidthInBytes;
width = p->dstArray->width;
desc = hipCreateChannelDesc(32, 0, 0, 0, hipChannelFormatKindSigned);
hsa_ext_image_channel_order_t channelOrder;
switch(p->dstArray->NumChannels) {
case 2:
channelOrder = HSA_EXT_IMAGE_CHANNEL_ORDER_RG;
break;
case 3:
channelOrder = HSA_EXT_IMAGE_CHANNEL_ORDER_RGB;
break;
case 4:
channelOrder = HSA_EXT_IMAGE_CHANNEL_ORDER_RGBA;
break;
case 1:
default:
channelOrder = HSA_EXT_IMAGE_CHANNEL_ORDER_R;
break;
}
hsa_ext_image_channel_type_t channelType;
e = ihipArrayToImageFormat(p->dstArray->Format,channelType);
srcPitch = p->srcPitch;
srcPtr = (void*)p->srcHost;
ySize = p->srcHeight;
dstPtr = p->dstArray->data;
hsa_ext_image_data_info_t imageInfo;
if(hipTextureType2DLayered == p->dstArray->textureType)
GetImageInfo(HSA_EXT_IMAGE_GEOMETRY_2DA, width, height, 0, channelOrder, channelType, imageInfo, depth);
else
GetImageInfo(HSA_EXT_IMAGE_GEOMETRY_3D, width, height, depth, channelOrder, channelType, imageInfo);
dstPitch = imageInfo.size/(height == 0 ? 1 : height)/(depth == 0 ? 1 : depth);
}
hsa_ext_image_data_info_t imageInfo;
if(hipTextureType2DLayered == p->dstArray->textureType)
GetImageInfo(HSA_EXT_IMAGE_GEOMETRY_2DA, width, height, 0, desc, imageInfo, depth);
else
GetImageInfo(HSA_EXT_IMAGE_GEOMETRY_3D, width, height, depth, desc, imageInfo);
dstPitch = imageInfo.size/(height == 0 ? 1:height)/(depth == 0 ? 1:depth);
} else {
// Non array destination
depth = p->extent.depth;
@@ -2083,7 +2418,7 @@ hipError_t hipIpcGetMemHandle(hipIpcMemHandle_t* handle, void* devPtr) {
size_t psize = 0u;
hc::accelerator acc;
if ((handle == NULL) || (devPtr == NULL)) {
hipStatus = hipErrorInvalidResourceHandle;
hipStatus = hipErrorInvalidHandle;
} else {
#if (__hcc_workweek__ >= 17332)
hc::AmPointerInfo amPointerInfo(NULL, NULL, NULL, 0, acc, 0, 0);
@@ -2094,7 +2429,7 @@ hipError_t hipIpcGetMemHandle(hipIpcMemHandle_t* handle, void* devPtr) {
if (status == AM_SUCCESS) {
psize = (size_t)amPointerInfo._sizeBytes;
} else {
hipStatus = hipErrorInvalidResourceHandle;
hipStatus = hipErrorInvalidHandle;
}
ihipIpcMemHandle_t* iHandle = (ihipIpcMemHandle_t*)handle;
// Save the size of the pointer to hipIpcMemHandle
@@ -2104,7 +2439,7 @@ hipError_t hipIpcGetMemHandle(hipIpcMemHandle_t* handle, void* devPtr) {
// Create HSA ipc memory
hsa_status_t hsa_status =
hsa_amd_ipc_memory_create(devPtr, psize, (hsa_amd_ipc_memory_t*)&(iHandle->ipc_handle));
if (hsa_status != HSA_STATUS_SUCCESS) hipStatus = hipErrorMemoryAllocation;
if (hsa_status != HSA_STATUS_SUCCESS) hipStatus = hipErrorOutOfMemory;
#else
hipStatus = hipErrorRuntimeOther;
#endif
@@ -2123,7 +2458,7 @@ hipError_t hipIpcOpenMemHandle(void** devPtr, hipIpcMemHandle_t handle, unsigned
hc::accelerator acc;
hsa_agent_t* agent = static_cast<hsa_agent_t*>(acc.get_hsa_agent());
if (!agent)
return ihipLogStatus(hipErrorInvalidResourceHandle);
return ihipLogStatus(hipErrorInvalidHandle);
ihipIpcMemHandle_t* iHandle = (ihipIpcMemHandle_t*)&handle;
// Attach ipc memory
@@ -2140,7 +2475,7 @@ hipError_t hipIpcOpenMemHandle(void** devPtr, hipIpcMemHandle_t handle, unsigned
hc::AmPointerInfo ampi(NULL, *devPtr, *devPtr, sizeof(*devPtr), acc, true, true);
am_status_t am_status = hc::am_memtracker_add(*devPtr,ampi);
if (am_status != AM_SUCCESS)
return ihipLogStatus(hipErrorMapBufferObjectFailed);
return ihipLogStatus(hipErrorMapFailed);
#if USE_APP_PTR_FOR_CTX
am_status = hc::am_memtracker_update(*devPtr, device->_deviceId, 0, ctx);
@@ -2148,7 +2483,7 @@ hipError_t hipIpcOpenMemHandle(void** devPtr, hipIpcMemHandle_t handle, unsigned
am_status = hc::am_memtracker_update(*devPtr, device->_deviceId, 0);
#endif
if(am_status != AM_SUCCESS)
return ihipLogStatus(hipErrorMapBufferObjectFailed);
return ihipLogStatus(hipErrorMapFailed);
}
#else
hipStatus = hipErrorRuntimeOther;
@@ -2168,7 +2503,7 @@ hipError_t hipIpcCloseMemHandle(void* devPtr) {
return ihipLogStatus(hipErrorInvalidValue);
if (hsa_amd_ipc_memory_detach(devPtr) != HSA_STATUS_SUCCESS)
return ihipLogStatus(hipErrorInvalidResourceHandle);
return ihipLogStatus(hipErrorInvalidHandle);
#else
hipStatus = hipErrorRuntimeOther;
#endif
projects/clr/hipamd/src/hip_module.cpp
+36 -40
Näytä tiedosto
@@ -28,6 +28,7 @@ THE SOFTWARE.
#include "hip/hip_ext.h"
#include "program_state.inl"
#include "trace_helper.h"
#include "hc_am.hpp"
#include <hsa/amd_hsa_kernel_code.h>
#include <hsa/hsa.h>
@@ -239,14 +240,7 @@ hipError_t ihipModuleLaunchKernel(TlsData *tls, hipFunction_t f, uint32_t global
aql.header |= (1 << HSA_PACKET_HEADER_BARRIER);
}
if (HCC_OPT_FLUSH) {
aql.header |= (HSA_FENCE_SCOPE_AGENT << HSA_PACKET_HEADER_ACQUIRE_FENCE_SCOPE) |
(HSA_FENCE_SCOPE_AGENT << HSA_PACKET_HEADER_RELEASE_FENCE_SCOPE);
} else {
aql.header |= (HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_ACQUIRE_FENCE_SCOPE) |
(HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_RELEASE_FENCE_SCOPE);
};
aql.header |= lp.launch_fence;
hc::completion_future cf;
@@ -389,46 +383,45 @@ __global__ void init_gws(uint nwm1) {
}
}
hipError_t hipLaunchCooperativeKernel(const void* f, dim3 gridDim,
__attribute__((visibility("default")))
hipError_t ihipLaunchCooperativeKernel(const void* f, dim3 gridDim,
dim3 blockDimX, void** kernelParams, unsigned int sharedMemBytes,
hipStream_t stream) {
hipStream_t stream, hip_impl::program_state& ps) {
HIP_INIT_API(hipLaunchCooperativeKernel, f, gridDim, blockDimX, kernelParams, sharedMemBytes, stream);
hipError_t result;
if ((f == nullptr) || (stream == nullptr) || (kernelParams == nullptr)) {
return ihipLogStatus(hipErrorNotInitialized);
return hipErrorNotInitialized;
}
if (!stream->getDevice()->_props.cooperativeLaunch) {
return ihipLogStatus(hipErrorInvalidConfiguration);
return hipErrorInvalidConfiguration;
}
// Prepare the kernel descriptor for initializing the GWS
hipFunction_t gwsKD = hip_impl::get_program_state().kernel_descriptor(
hipFunction_t gwsKD = ps.kernel_descriptor(
reinterpret_cast<std::uintptr_t>(&init_gws),
hip_impl::target_agent(stream));
if (gwsKD == nullptr) {
return ihipLogStatus(hipErrorInvalidValue);
return hipErrorInvalidValue;
}
hip_impl::kernargs_size_align gwsKargs =
hip_impl::get_program_state().get_kernargs_size_align(
hip_impl::kernargs_size_align gwsKargs = ps.get_kernargs_size_align(
reinterpret_cast<std::uintptr_t>(&init_gws));
gwsKD->_kernarg_layout = *reinterpret_cast<const std::vector<
std::pair<std::size_t, std::size_t>>*>(gwsKargs.getHandle());
// Prepare the kernel descriptor for the main kernel
hipFunction_t kd = hip_impl::get_program_state().kernel_descriptor(
hipFunction_t kd = ps.kernel_descriptor(
reinterpret_cast<std::uintptr_t>(f),
hip_impl::target_agent(stream));
if (kd == nullptr) {
return ihipLogStatus(hipErrorInvalidValue);
return hipErrorInvalidValue;
}
hip_impl::kernargs_size_align kargs =
hip_impl::get_program_state().get_kernargs_size_align(
ps.get_kernargs_size_align(
reinterpret_cast<std::uintptr_t>(f));
kd->_kernarg_layout = *reinterpret_cast<const std::vector<
@@ -445,6 +438,7 @@ hipError_t hipLaunchCooperativeKernel(const void* f, dim3 gridDim,
streamCrit->_av.acquire_locked_hsa_queue();
#endif
GET_TLS();
// launch the init_gws kernel to initialize the GWS
result = ihipModuleLaunchKernel(tls, gwsKD, 1, 1, 1, 1, 1, 1,
0, stream, gwsKernelParam, nullptr, nullptr, nullptr, 0, true);
@@ -455,7 +449,7 @@ hipError_t hipLaunchCooperativeKernel(const void* f, dim3 gridDim,
stream->criticalData()._av.release_locked_hsa_queue();
#endif
return ihipLogStatus(hipErrorLaunchFailure);
return hipErrorLaunchFailure;
}
size_t impCoopArg = 1;
@@ -476,29 +470,30 @@ hipError_t hipLaunchCooperativeKernel(const void* f, dim3 gridDim,
stream->criticalData()._av.release_locked_hsa_queue();
#endif
return ihipLogStatus(result);
return result;
}
hipError_t hipLaunchCooperativeKernelMultiDevice(hipLaunchParams* launchParamsList,
int numDevices, unsigned int flags) {
__attribute__((visibility("default")))
hipError_t ihipLaunchCooperativeKernelMultiDevice(hipLaunchParams* launchParamsList,
int numDevices, unsigned int flags, hip_impl::program_state& ps) {
HIP_INIT_API(hipLaunchCooperativeKernelMultiDevice, launchParamsList, numDevices, flags);
hipError_t result;
if (numDevices > g_deviceCnt || launchParamsList == nullptr || numDevices > MAX_COOPERATIVE_GPUs) {
return ihipLogStatus(hipErrorInvalidValue);
return hipErrorInvalidValue;
}
for (int i = 0; i < numDevices; ++i) {
if (!launchParamsList[i].stream->getDevice()->_props.cooperativeMultiDeviceLaunch) {
return ihipLogStatus(hipErrorInvalidConfiguration);
return hipErrorInvalidConfiguration;
}
}
hipFunction_t* gwsKds = reinterpret_cast<hipFunction_t*>(malloc(sizeof(hipFunction_t) * numDevices));
hipFunction_t* kds = reinterpret_cast<hipFunction_t*>(malloc(sizeof(hipFunction_t) * numDevices));
if (kds == nullptr || gwsKds == nullptr) {
return ihipLogStatus(hipErrorNotInitialized);
return hipErrorNotInitialized;
}
// prepare all kernel descriptors for initializing the GWS and the main kernels per device
@@ -507,30 +502,30 @@ hipError_t hipLaunchCooperativeKernelMultiDevice(hipLaunchParams* launchParamsLi
if (lp.stream == nullptr) {
free(gwsKds);
free(kds);
return ihipLogStatus(hipErrorNotInitialized);
return hipErrorNotInitialized;
}
gwsKds[i] = hip_impl::get_program_state().kernel_descriptor(reinterpret_cast<std::uintptr_t>(&init_gws),
gwsKds[i] = ps.kernel_descriptor(reinterpret_cast<std::uintptr_t>(&init_gws),
hip_impl::target_agent(lp.stream));
if (gwsKds[i] == nullptr) {
free(gwsKds);
free(kds);
return ihipLogStatus(hipErrorInvalidValue);
return hipErrorInvalidValue;
}
hip_impl::kernargs_size_align gwsKargs = hip_impl::get_program_state().get_kernargs_size_align(
hip_impl::kernargs_size_align gwsKargs = ps.get_kernargs_size_align(
reinterpret_cast<std::uintptr_t>(&init_gws));
gwsKds[i]->_kernarg_layout = *reinterpret_cast<const std::vector<std::pair<std::size_t, std::size_t>>*>(
gwsKargs.getHandle());
kds[i] = hip_impl::get_program_state().kernel_descriptor(reinterpret_cast<std::uintptr_t>(lp.func),
kds[i] = ps.kernel_descriptor(reinterpret_cast<std::uintptr_t>(lp.func),
hip_impl::target_agent(lp.stream));
if (kds[i] == nullptr) {
free(gwsKds);
free(kds);
return ihipLogStatus(hipErrorInvalidValue);
return hipErrorInvalidValue;
}
hip_impl::kernargs_size_align kargs = hip_impl::get_program_state().get_kernargs_size_align(
hip_impl::kernargs_size_align kargs = ps.get_kernargs_size_align(
reinterpret_cast<std::uintptr_t>(lp.func));
kds[i]->_kernarg_layout = *reinterpret_cast<const std::vector<std::pair<std::size_t, std::size_t>>*>(
kargs.getHandle());
@@ -539,9 +534,10 @@ hipError_t hipLaunchCooperativeKernelMultiDevice(hipLaunchParams* launchParamsLi
mg_sync *mg_sync_ptr = 0;
mg_info *mg_info_ptr[MAX_COOPERATIVE_GPUs] = {0};
GET_TLS();
result = hip_internal::ihipHostMalloc(tls, (void **)&mg_sync_ptr, sizeof(mg_sync), hipHostMallocDefault);
if (result != hipSuccess) {
return ihipLogStatus(hipErrorInvalidValue);
return hipErrorInvalidValue;
}
mg_sync_ptr->w0 = 0;
mg_sync_ptr->w1 = 0;
@@ -554,7 +550,7 @@ hipError_t hipLaunchCooperativeKernelMultiDevice(hipLaunchParams* launchParamsLi
for (int j = 0; j < i; ++j) {
hip_internal::ihipHostFree(tls, mg_info_ptr[j]);
}
return ihipLogStatus(hipErrorInvalidValue);
return hipErrorInvalidValue;
}
// calculate the sum of sizes of all grids
const hipLaunchParams& lp = launchParamsList[i];
@@ -593,7 +589,7 @@ hipError_t hipLaunchCooperativeKernelMultiDevice(hipLaunchParams* launchParamsLi
hip_internal::ihipHostFree(tls, mg_info_ptr[j]);
}
return ihipLogStatus(hipErrorLaunchFailure);
return hipErrorLaunchFailure;
}
}
@@ -636,7 +632,7 @@ hipError_t hipLaunchCooperativeKernelMultiDevice(hipLaunchParams* launchParamsLi
hip_internal::ihipHostFree(tls, mg_info_ptr[j]);
}
return ihipLogStatus(hipErrorLaunchFailure);
return hipErrorLaunchFailure;
}
}
@@ -657,7 +653,7 @@ hipError_t hipLaunchCooperativeKernelMultiDevice(hipLaunchParams* launchParamsLi
hip_internal::ihipHostFree(tls, mg_info_ptr[j]);
}
return ihipLogStatus(result);
return result;
}
namespace hip_impl {
projects/clr/hipamd/src/hip_stream.cpp
+6 -6
Näytä tiedosto
@@ -133,7 +133,7 @@ hipError_t hipStreamWaitEvent(hipStream_t stream, hipEvent_t event, unsigned int
hipError_t e = hipSuccess;
if (event == nullptr) {
e = hipErrorInvalidResourceHandle;
e = hipErrorInvalidHandle;
} else {
auto ecd = event->locked_copyCrit();
@@ -189,7 +189,7 @@ hipError_t hipStreamSynchronize(hipStream_t stream) {
//---
/**
* @return #hipSuccess, #hipErrorInvalidResourceHandle
* @return #hipSuccess, #hipErrorInvalidHandle
*/
hipError_t hipStreamDestroy(hipStream_t stream) {
HIP_INIT_API(hipStreamDestroy, stream);
@@ -199,7 +199,7 @@ hipError_t hipStreamDestroy(hipStream_t stream) {
//--- Drain the stream:
if (stream == NULL) {
if (!HIP_FORCE_NULL_STREAM) {
e = hipErrorInvalidResourceHandle;
e = hipErrorInvalidHandle;
}
} else {
stream->locked_wait();
@@ -210,7 +210,7 @@ hipError_t hipStreamDestroy(hipStream_t stream) {
ctx->locked_removeStream(stream);
delete stream;
} else {
e = hipErrorInvalidResourceHandle;
e = hipErrorInvalidHandle;
}
}
@@ -225,7 +225,7 @@ hipError_t hipStreamGetFlags(hipStream_t stream, unsigned int* flags) {
if (flags == NULL) {
return ihipLogStatus(hipErrorInvalidValue);
} else if (stream == hipStreamNull) {
return ihipLogStatus(hipErrorInvalidResourceHandle);
return ihipLogStatus(hipErrorInvalidHandle);
} else {
*flags = stream->_flags;
return ihipLogStatus(hipSuccess);
@@ -240,7 +240,7 @@ hipError_t hipStreamGetPriority(hipStream_t stream, int* priority) {
if (priority == NULL) {
return ihipLogStatus(hipErrorInvalidValue);
} else if (stream == hipStreamNull) {
return ihipLogStatus(hipErrorInvalidResourceHandle);
return ihipLogStatus(hipErrorInvalidHandle);
} else {
#if defined(__HCC__) && (__hcc_major__ < 3) && (__hcc_minor__ < 3)
*priority = 0;
projects/clr/hipamd/src/hiprtc.cpp
+9 -9
Näytä tiedosto
@@ -143,7 +143,9 @@ struct _hiprtcProgram {
{
using namespace std;
name = hip_impl::demangle(name.c_str());
char* demangled = hip_impl::demangle(name.c_str());
name.assign(demangled == nullptr ? "" : demangled);
free(demangled);
if (name.empty()) return name;
@@ -392,18 +394,16 @@ namespace
namespace hip_impl
{
inline
std::string demangle(const char* x)
char* demangle(const char* x)
{
if (!x) return {};
if (!x) return nullptr;
int s{};
std::unique_ptr<char, decltype(std::free)*> tmp{
abi::__cxa_demangle(x, nullptr, nullptr, &s), std::free};
char* tmp = abi::__cxa_demangle(x, nullptr, nullptr, &s);
if (s != 0) return {};
if (s != 0) return nullptr;
return tmp.get();
return tmp;
}
} // Namespace hip_impl.
@@ -598,4 +598,4 @@ hiprtcResult hiprtcGetCodeSize(hiprtcProgram p, std::size_t* sz)
*sz = p->elf.size();
return HIPRTC_SUCCESS;
}
}
projects/clr/hipamd/tests/hipify-clang/unit_tests/samples/square.cu
+2 -2
Näytä tiedosto
@@ -67,11 +67,11 @@ int main(int argc, char *argv[])
printf ("info: allocate host mem (%6.2f MB)\n", 2*Nbytes/1024.0/1024.0);
A_h = (float*)malloc(Nbytes);
// CHECK: CHECK(A_h == 0 ? hipErrorMemoryAllocation : hipSuccess );
// CHECK: CHECK(A_h == 0 ? hipErrorOutOfMemory : hipSuccess );
CHECK(A_h == 0 ? cudaErrorMemoryAllocation : cudaSuccess );
C_h = (float*)malloc(Nbytes);
// CHECK: CHECK(C_h == 0 ? hipErrorMemoryAllocation : hipSuccess );
// CHECK: CHECK(C_h == 0 ? hipErrorOutOfMemory : hipSuccess );
CHECK(C_h == 0 ? cudaErrorMemoryAllocation : cudaSuccess );
// Fill with Phi + i
for (size_t i=0; i<N; i++)
projects/clr/hipamd/tests/src/deviceLib/hipVersion.cpp
+42
Näytä tiedosto
@@ -0,0 +1,42 @@
/*
Copyright (c) 2015-present Advanced Micro Devices, Inc. All rights reserved.
Permission is hereby granted, free of charge, to any person obtaining a copy
of this software and associated documentation files (the "Software"), to deal
in the Software without restriction, including without limitation the rights
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
copies of the Software, and to permit persons to whom the Software is
furnished to do so, subject to the following conditions:
The above copyright notice and this permission notice shall be included in
all copies or substantial portions of the Software.
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
/* HIT_START
* BUILD: %t %s ../test_common.cpp
* TEST: %t
* HIT_END
*/
#include "test_common.h"
/*
* Note : Compile time test hence always returns success while run.
* Intension is to make sure apilcation can access hip version.
*/
int main()
{
std::cout<<"Hip major version : "<<HIP_VERSION_MAJOR<<std::endl;
std::cout<<"Hip minor version : "<<HIP_VERSION_MINOR<<std::endl;
std::cout<<"Hip patch version : "<<HIP_VERSION_PATCH<<std::endl;
std::cout<<"Hip patch version : "<<HIP_VERSION<<std::endl;
passed();
}
projects/clr/hipamd/tests/src/dynamicLoading/complex_loading_behavior.cpp
+6 -6
Näytä tiedosto
@@ -62,11 +62,11 @@ int launch_local_kernel() {
hipDeviceProp_t props;
CHECK(hipGetDeviceProperties(&props, device /*deviceID*/));
A_h = (float*)malloc(Nbytes);
CHECK(A_h == 0 ? hipErrorMemoryAllocation : hipSuccess);
CHECK(A_h == 0 ? hipErrorOutOfMemory : hipSuccess);
B_h = (float*)malloc(Nbytes);
CHECK(B_h == 0 ? hipErrorMemoryAllocation : hipSuccess);
CHECK(B_h == 0 ? hipErrorOutOfMemory : hipSuccess);
C_h = (float*)malloc(Nbytes);
CHECK(C_h == 0 ? hipErrorMemoryAllocation : hipSuccess);
CHECK(C_h == 0 ? hipErrorOutOfMemory : hipSuccess);
// Fill with Phi + i
for (size_t i = 0; i < N; i++) {
A_h[i] = 1.618f + i;
@@ -174,11 +174,11 @@ extern "C" int foo() {
hipDeviceProp_t props;
CHECK(hipGetDeviceProperties(&props, device /*deviceID*/));
A_h = (float*)malloc(Nbytes);
CHECK(A_h == 0 ? hipErrorMemoryAllocation : hipSuccess);
CHECK(A_h == 0 ? hipErrorOutOfMemory : hipSuccess);
B_h = (float*)malloc(Nbytes);
CHECK(B_h == 0 ? hipErrorMemoryAllocation : hipSuccess);
CHECK(B_h == 0 ? hipErrorOutOfMemory : hipSuccess);
C_h = (float*)malloc(Nbytes);
CHECK(C_h == 0 ? hipErrorMemoryAllocation : hipSuccess);
CHECK(C_h == 0 ? hipErrorOutOfMemory : hipSuccess);
// Fill with Phi + i
for (size_t i = 0; i < N; i++) {
A_h[i] = 1.618f + i;
projects/clr/hipamd/tests/src/runtimeApi/event/record_event.cpp
+6 -6
Näytä tiedosto
@@ -133,14 +133,14 @@ void test(unsigned testMask, int* C_d, int* C_h, int64_t numElements, hipStream_
{
// Check some error conditions for incomplete events:
HIPCHECK_API(hipEventElapsedTime(&t, timingDisabled, stop), hipErrorInvalidResourceHandle);
HIPCHECK_API(hipEventElapsedTime(&t, start, timingDisabled), hipErrorInvalidResourceHandle);
HIPCHECK_API(hipEventElapsedTime(&t, timingDisabled, stop), hipErrorInvalidHandle);
HIPCHECK_API(hipEventElapsedTime(&t, start, timingDisabled), hipErrorInvalidHandle);
HIPCHECK_API(hipEventElapsedTime(&t, neverCreated, stop), hipErrorInvalidResourceHandle);
HIPCHECK_API(hipEventElapsedTime(&t, start, neverCreated), hipErrorInvalidResourceHandle);
HIPCHECK_API(hipEventElapsedTime(&t, neverCreated, stop), hipErrorInvalidHandle);
HIPCHECK_API(hipEventElapsedTime(&t, start, neverCreated), hipErrorInvalidHandle);
HIPCHECK_API(hipEventElapsedTime(&t, neverRecorded, stop), hipErrorInvalidResourceHandle);
HIPCHECK_API(hipEventElapsedTime(&t, start, neverRecorded), hipErrorInvalidResourceHandle);
HIPCHECK_API(hipEventElapsedTime(&t, neverRecorded, stop), hipErrorInvalidHandle);
HIPCHECK_API(hipEventElapsedTime(&t, start, neverRecorded), hipErrorInvalidHandle);
}
HIPCHECK(hipEventDestroy(start));
projects/clr/hipamd/tests/src/runtimeApi/memory/hipMemcpy.cpp
+6 -2
Näytä tiedosto
@@ -292,8 +292,12 @@ void memcpytest2_get_host_memory(size_t& free, size_t& total) {
MEMORYSTATUSEX status;
status.dwLength = sizeof(status);
GlobalMemoryStatusEx(&status);
free = status.ullAvailPhys;
total = status.ullTotalPhys;
// Windows doesn't allow allocating more than half of system memory to the gpu.
// Since the runtime also needs space for its internal allocations,
// we should not try to allocate more than 40% of reported system memory,
// otherwise we can run into OOM issues.
free = static_cast<size_t>(0.4 * status.ullAvailPhys);
total = static_cast<size_t>(0.4 * status.ullTotalPhys);
}
#else
struct sysinfo memInfo;
projects/clr/hipamd/tests/src/runtimeApi/memory/hipMemcpyWithStream.cpp
+57
Näytä tiedosto
@@ -0,0 +1,57 @@
/*
Copyright (c) 2015-present Advanced Micro Devices, Inc. All rights reserved.
Permission is hereby granted, free of charge, to any person obtaining a copy
of this software and associated documentation files (the "Software"), to deal
in the Software without restriction, including without limitation the rights
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
copies of the Software, and to permit persons to whom the Software is
furnished to do so, subject to the following conditions:
The above copyright notice and this permission notice shall be included in
all copies or substantial portions of the Software.
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
/*
* Conformance test for checking functionality of
* hipError_t hipMemcpyPeer(void* dst, int dstDeviceId, const void* src, int srcDeviceId, size_t
* sizeBytes);
*/
/* HIT_START
* BUILD: %t %s ../../test_common.cpp
* TEST: %t
* HIT_END
*/
#include "test_common.h"
int main() {
size_t Nbytes = N * sizeof(int);
int numDevices = 0;
int *A_d, *B_d, *C_d;
int *A_h, *B_h, *C_h;
unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N);
HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N, false);
hipStream_t stream;
HIPCHECK(hipStreamCreate(&stream));
HIPCHECK(hipMemcpyWithStream(A_d, A_h, Nbytes, hipMemcpyHostToDevice, stream));
HIPCHECK(hipMemcpyWithStream(B_d, B_h, Nbytes, hipMemcpyHostToDevice,stream));
hipLaunchKernelGGL(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock), 0, stream,
static_cast<const int*>(A_d), static_cast<const int*>(B_d), C_d, N);
HIPCHECK(hipStreamSynchronize(stream));
HIPCHECK(hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost));
HipTest::checkVectorADD(A_h, B_h, C_h, N);
HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false);
HIPCHECK(hipStreamDestroy(stream));
passed();
}
projects/clr/hipamd/tests/src/runtimeApi/module/hipExtLaunchMultiKernelMultiDevice.cpp
+2 -2
Näytä tiedosto
@@ -67,9 +67,9 @@ int main(int argc, char *argv[])
printf ("info: allocate host mem (%6.2f MB)\n", 2*Nbytes/1024.0/1024.0);
A_h = (float*)malloc(Nbytes);
HIPCHECK(A_h == 0 ? hipErrorMemoryAllocation : hipSuccess );
HIPCHECK(A_h == 0 ? hipErrorOutOfMemory : hipSuccess );
C_h = (float*)malloc(Nbytes);
HIPCHECK(C_h == 0 ? hipErrorMemoryAllocation : hipSuccess );
HIPCHECK(C_h == 0 ? hipErrorOutOfMemory : hipSuccess );
// Fill with Phi + i
for (size_t i = 0; i < N; i++)
{
projects/clr/hipamd/tests/src/runtimeApi/stream/hipStreamAddCallback.cpp
+2 -2
Näytä tiedosto
@@ -60,9 +60,9 @@ int main(int argc, char* argv[]) {
size_t Nbytes = N * sizeof(float);
A_h = (float*)malloc(Nbytes);
HIPCHECK(A_h == 0 ? hipErrorMemoryAllocation : hipSuccess);
HIPCHECK(A_h == 0 ? hipErrorOutOfMemory : hipSuccess);
C_h = (float*)malloc(Nbytes);
HIPCHECK(C_h == 0 ? hipErrorMemoryAllocation : hipSuccess);
HIPCHECK(C_h == 0 ? hipErrorOutOfMemory : hipSuccess);
// Fill with Phi + i
for (size_t i = 0; i < N; i++) {
projects/clr/hipamd/tests/src/runtimeApi/stream/hipStreamSync2.cpp
+1 -1
Näytä tiedosto
@@ -201,7 +201,7 @@ void runTests(int64_t numElements) {
int main(int argc, char* argv[]) {
// Can' destroy the default stream:// TODO - move to another test
HIPCHECK_API(hipStreamDestroy(0), hipErrorInvalidResourceHandle);
HIPCHECK_API(hipStreamDestroy(0), hipErrorInvalidHandle);
HipTest::parseStandardArguments(argc, argv, true /*failOnUndefinedArg*/);
projects/clr/hipamd/tests/src/texture/hipBindTex2DPitch.cpp
-11
Näytä tiedosto
@@ -58,19 +58,8 @@ int main (void)
SIZE_W*sizeof(TYPE_t), SIZE_H, hipMemcpyHostToDevice));
tex.normalized = false;
#if defined(__HIP_PLATFORM_NVCC__)
cudaError_t status = cudaBindTexture2D(&tex_ofs, &tex, devPtrA, &tex.channelDesc,
SIZE_W, SIZE_H, devPitchA);
if (status != cudaSuccess) {
printf("%serror: '%s'(%d) at %s:%d%s\n", KRED, cudaGetErrorString(status),
status,__FILE__, __LINE__, KNRM);
failed("API returned error code.");
}
#else
HIPCHECK(hipBindTexture2D(&tex_ofs, &tex, devPtrA, &tex.channelDesc,
SIZE_W, SIZE_H, devPitchA));
#endif
HIPCHECK(hipMalloc((void**)&devPtrB, SIZE_W*sizeof(TYPE_t)*SIZE_H)) ;
hipLaunchKernelGGL(texture2dCopyKernel, dim3(4,4,1), dim3(32,32,1), 0, 0, devPtrB);