AI/rocm-systems
2
0
Vork 0
Code Kwesties Pull-aanvragen Actions Paketten Projecten Publicaties Wiki Activiteit

Merge 'master' into 'amd-master'

Bladeren bron 
Change-Id: Ifc6f69fb7048d2dca686c696d11f76eecee19e42


[ROCm/hip commit: 59c54da4f0]
This commit is contained in:
Jenkins
2019-02-08 04:09:44 -06:00
bovenliggende ab5d981551 59ccae8101
commit d2dbe30ac2
23 gewijzigde bestanden met toevoegingen van 959 en 662 verwijderingen
Download Patch-bestand Download Diff-bestand Expand all files Collapse all files
projects/hip/CMakeLists.txt
+10
Bestand weergeven
@@ -71,6 +71,16 @@ if(NOT (HIP_COMPILER STREQUAL "hcc" OR HIP_COMPILER STREQUAL "clang"))
message(FATAL_ERROR "Must use HIP_COMPILER as hcc or clang")
endif()
message(STATUS "HIP Compiler: " ${HIP_COMPILER})
add_to_config(_buildInfo HIP_COMPILER)
# Determine HIP_RUNTIME
# Either HCC or VDI; default is HCC
if(NOT DEFINED ENV{HIP_RUNTIME})
set(HIP_RUNTIME "HCC" CACHE STRING "HIP Runtime")
else()
set(HIP_RUNTIME $ENV{HIP_RUNTIME} CACHE STRING "HIP Runtime")
endif()
add_to_config(_buildInfo HIP_RUNTIME)
# If HIP_PLATFORM is hcc, we need HCC_HOME and HSA_PATH to be defined
projects/hip/Jenkinsfile
Vendored
+1 -46
Bestand weergeven
@@ -295,7 +295,7 @@ 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 4 builds in parallel: rocm-head, rocm-2.1.x, rocm-2.0.x and cuda-9.x
// The following launches 3 builds in parallel: rocm-head, rocm-2.1.x and cuda-9.x
parallel rocm_2_1:
{
node('hip-rocm')
@@ -341,51 +341,6 @@ parallel rocm_2_1:
*/
}
},
rocm_2_0:
{
node('hip-rocm')
{
String hcc_ver = 'rocm-2.0.x'
String from_image = 'ci_test_nodes/rocm-2.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
String source_hip_rel = checkout_and_version( hcc_ver )
// Create/reuse a docker image that represents the hip build environment
def hip_build_image = docker_build_image( hcc_ver, 'hip', '', source_hip_rel, from_image )
// Print system information for the log
hip_build_image.inside( inside_args )
{
sh """#!/usr/bin/env bash
set -x
/opt/rocm/bin/rocm_agent_enumerator -t ALL
/opt/rocm/bin/hcc --version
"""
}
// Conctruct a binary directory path based on build config
String build_hip_rel = build_directory_rel( build_config );
// Build hip inside of the build environment
docker_build_inside_image( hip_build_image, inside_args, hcc_ver, '', build_config, source_hip_rel, build_hip_rel )
// Clean docker build image
docker_clean_images( 'hip', docker_build_image_name( ) )
// After a successful build, upload a docker image of the results
/*
String hip_image_name = docker_upload_artifactory( hcc_ver, job_name, from_image, source_hip_rel, build_hip_rel )
if( params.push_image_to_docker_hub )
{
docker_upload_dockerhub( job_name, hip_image_name, 'rocm' )
docker_clean_images( 'rocm', hip_image_name )
}
docker_clean_images( job_name, hip_image_name )
*/
}
},
rocm_head:
{
node('hip-rocm')
projects/hip/bin/hipcc
+34 -10
Bestand weergeven
@@ -62,10 +62,16 @@ $HIP_CLANG_PATH=$ENV{'HIP_CLANG_PATH'};
$DEVICE_LIB_PATH=$ENV{'DEVICE_LIB_PATH'};
$HIP_CLANG_HCC_COMPAT_MODE=$ENV{'HIP_CLANG_HCC_COMPAT_MODE'}; # HCC compatibility mode
if (defined $HIP_VDI_HOME) {
$HIP_INFO_PATH= "$HIP_VDI_HOME/lib/.hipInfo";
} else {
$HIP_INFO_PATH= "$HIP_PATH/lib/.hipInfo";
}
#---
# Read .hipInfo
my %hipConfig = ();
parse_config_file("$HIP_PATH/lib/.hipInfo", \%hipConfig);
parse_config_file("$HIP_INFO_PATH", \%hipConfig);
#---
# Temporary directories
@@ -93,6 +99,13 @@ sub delete_temp_dirs {
$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);
$HIP_COMPILER= $hipConfig{'HIP_COMPILER'};
$HIP_RUNTIME= $hipConfig{'HIP_RUNTIME'};
# If using VDI runtime, need to find HIP_VDI_HOME
if ($HIP_RUNTIME eq "VDI" and !defined $HIP_VDI_HOME) {
$HIP_VDI_HOME = "/opt/rocm/hip"
}
if (defined $HIP_VDI_HOME) {
my $bits = "";
@@ -110,8 +123,14 @@ if (defined $HIP_VDI_HOME) {
$HIP_LIB_PATH = "$HIP_VDI_HOME/lib" . $bits;
}
if (defined $HIP_CLANG_PATH) {
$HIP_PLATFORM = "clang"
if ($HIP_COMPILER eq "clang") {
$HIP_PLATFORM = "clang";
if (!defined $HIP_CLANG_PATH) {
$HIP_CLANG_PATH = "/opt/rocm/llvm/bin";
}
if (!defined $DEVICE_LIB_PATH) {
$DEVICE_LIB_PATH = "/opt/rocm/lib";
}
}
if ($verbose & 0x2) {
@@ -193,7 +212,7 @@ if ($HIP_PLATFORM eq "clang") {
# HCC* may be used to compile src/hip_hcc.o (and also feed the HIPCXXFLAGS below)
$HCC = "$HCC_HOME/bin/hcc";
$HCCFLAGS = "-hc -D__HIPCC__ -I$HCC_HOME/include ";
$HCCFLAGS = "-hc -D__HIPCC__ -isystem $HCC_HOME/include ";
$HIPCC=$HCC;
$HIPCXXFLAGS = $HCCFLAGS;
@@ -215,7 +234,7 @@ if ($HIP_PLATFORM eq "clang") {
# That's because HCC already uses libstdc++ by default if a newer g++/libstdc++ is available
# Cent OS 7 and RHEL 7.4 cannot use libstdc++ for compilation, default to libc++
if (${GCC_CUR_VER} eq ${GPP_CUR_VER} and $GPP_VER_FIELDS[0] < 5 and ($HOST_OSNAME ne "\"centos\"") and ($HOST_OSNAME ne "\"rhel\"")) {
$HCC_WA_FLAGS .= " -stdlib=libstdc++ -I/usr/include/x86_64-linux-gnu -I/usr/include/x86_64-linux-gnu/c++/${GCC_CUR_VER} -I/usr/include/c++/${GCC_CUR_VER} ";
$HCC_WA_FLAGS .= " -stdlib=libstdc++ -isystem /usr/include/x86_64-linux-gnu -isystem /usr/include/x86_64-linux-gnu/c++/${GCC_CUR_VER} -isystem /usr/include/c++/${GCC_CUR_VER} ";
# Add C++ libs for GCC.
$HIPLDFLAGS .= " -lstdc++";
}
@@ -228,8 +247,8 @@ if ($HIP_PLATFORM eq "clang") {
$setStdLib = 1;
}
$HIPCXXFLAGS .= " -I$HIP_PATH/include/hip/hcc_detail/cuda";
$HIPCXXFLAGS .= " -I$HSA_PATH/include";
$HIPCXXFLAGS .= " -isystem $HIP_PATH/include/hip/hcc_detail/cuda";
$HIPCXXFLAGS .= " -isystem $HSA_PATH/include";
$HIPCXXFLAGS .= " -Wno-deprecated-register";
$HIPLDFLAGS .= " -L$HSA_PATH/lib -L$ROCM_PATH/lib -lhsa-runtime64 -lhc_am ";
@@ -240,7 +259,7 @@ if ($HIP_PLATFORM eq "clang") {
if ($HIP_ATP_MARKER) {
$marker_inc_path = "$marker_path/include";
if (-e $marker_inc_path) {
$HIPCXXFLAGS .= " -I$marker_inc_path";
$HIPCXXFLAGS .= " -isystem $marker_inc_path";
}
}
@@ -265,7 +284,7 @@ if ($HIP_PLATFORM eq "clang") {
$HIPCC="$CUDA_PATH/bin/nvcc";
$HIPCXXFLAGS .= " -Wno-deprecated-gpu-targets ";
$HIPCXXFLAGS .= " -I$CUDA_PATH/include";
$HIPCXXFLAGS .= " -isystem $CUDA_PATH/include";
$HIPLDFLAGS = " -Wno-deprecated-gpu-targets -lcuda -lcudart -L$CUDA_PATH/lib64";
} else {
@@ -274,7 +293,7 @@ if ($HIP_PLATFORM eq "clang") {
}
# Add paths to common HIP includes:
$HIPCXXFLAGS .= " -I$HIP_INCLUDE_PATH -DHIP_VERSION_MAJOR=$HIP_VERSION_MAJOR -DHIP_VERSION_MINOR=$HIP_VERSION_MINOR -DHIP_VERSION_PATCH=$HIP_VERSION_PATCH" ;
$HIPCXXFLAGS .= " -isystem $HIP_INCLUDE_PATH -DHIP_VERSION_MAJOR=$HIP_VERSION_MAJOR -DHIP_VERSION_MINOR=$HIP_VERSION_MINOR -DHIP_VERSION_PATCH=$HIP_VERSION_PATCH" ;
my $compileOnly = 0;
my $needCXXFLAGS = 0; # need to add CXX flags to compile step
@@ -750,6 +769,11 @@ if ($HIP_PLATFORM eq "clang") {
$HIPLDFLAGS .= " -lgcc_s -lgcc -lpthread -lm";
}
# Windows specific options
if ($^O eq 'MSWin32' and $HIP_PLATFORM eq "clang") {
$HIPCXXFLAGS .= " -std=c++14 -fms-extensions -fms-compatibility"
}
if ($HIPCC_COMPILE_FLAGS_APPEND) {
$HIPCXXFLAGS .= " $HIPCC_COMPILE_FLAGS_APPEND";
}
projects/hip/bin/hipify-perl
+4 -3
Bestand weergeven
@@ -269,7 +269,8 @@ while (@ARGV) {
$ft{'mem'} += s/\bcudaMemcpyDeviceToHost\b/hipMemcpyDeviceToHost/g;
$ft{'mem'} += s/\bcudaMemcpyDeviceToDevice\b/hipMemcpyDeviceToDevice/g;
$ft{'mem'} += s/\bcudaMemcpyDefault\b/hipMemcpyDefault/g;
$ft{'mem'} += s/\bcudaMemcpyToSymbol\b/hipMemcpyToSymbol/g;
$ft{'mem'} += s/\bcudaMemcpyToSymbol\s*\(\s*(\w+)\b/hipMemcpyToSymbol\(HIP_SYMBOL\($1\)/g;
$ft{'mem'} += s/\bcudaMemcpyFromSymbol\s*\(\s*(.+?)\s*,\s*(.+?)\b/hipMemcpyFromSymbol\($1, HIP_SYMBOL\($2\)/g;
$ft{'mem'} += s/\bcudaMemset\b/hipMemset/g;
$ft{'mem'} += s/\bcudaMemsetAsync\b/hipMemsetAsync/g;
@@ -287,8 +288,8 @@ while (@ARGV) {
$ft{'mem'} += s/\bcudaMemcpy2DToArray\b/hipMemcpy2DToArray/g;
$ft{'mem'} += s/\bcudaMemcpyToArray\b/hipMemcpyToArray/g;
$ft{'mem'} += s/\bcudaGetSymbolAddress\b/hipGetSymbolAddress/g;
$ft{'mem'} += s/\bcudaGetSymbolSize\b/hipGetSymbolSize/g;
$ft{'mem'} += s/\bcudaGetSymbolAddress\s*\(\s*(.+?)\s*,\s*(.+?)\b/hipGetSymbolAddress\($1, HIP_SYMBOL\($2\)/g;
$ft{'mem'} += s/\bcudaGetSymbolSize\s*\(\s*&(\w+)\s*,\s*(.+?)\b/hipGetSymbolSize(&$1, HIP_SYMBOL\($2\)/g;
#--------
# Memory management:
projects/hip/docs/markdown/CUBLAS_API_supported_by_HIP.md
+8 -8
Bestand weergeven
@@ -99,13 +99,13 @@
|`cublasGetVersion_v2` | |
|`cublasGetProperty` | |
|`cublasGetStream` |`hipblasGetStream` |
|`cublasGetStream_v2` |`hipblasGetStream_v2` |
|`cublasGetStream_v2` |`hipblasGetStream` |
|`cublasSetStream` |`hipblasSetStream` |
|`cublasSetStream_v2` |`hipblasSetStream_v2` |
|`cublasSetStream_v2` |`hipblasSetStream` |
|`cublasGetPointerMode` |`hipblasGetPointerMode` |
|`cublasGetPointerMode_v2` |`hipblasGetPointerMode_v2` |
|`cublasGetPointerMode_v2` |`hipblasGetPointerMode` |
|`cublasSetPointerMode` |`hipblasSetPointerMode` |
|`cublasSetPointerMode_v2` |`hipblasSetPointerMode_v2` |
|`cublasSetPointerMode_v2` |`hipblasSetPointerMode` |
|`cublasGetAtomicsMode` | |
|`cublasSetAtomicsMode` | |
|`cublasGetMathMode` | |
@@ -382,8 +382,8 @@
|`cublasSgemm_v2` |`hipblasSgemm` |
|`cublasDgemm` |`hipblasDgemm` |
|`cublasDgemm_v2` |`hipblasDgemm` |
|`cublasCgemm` |`hipblasCgemm` |
|`cublasCgemm_v2` |`hipblasCgemm` |
|`cublasCgemm` | |
|`cublasCgemm_v2` | |
|`cublasCgemm3m` | |
|`cublasCgemm3mEx` | |
|`cublasZgemm` | |
@@ -464,8 +464,8 @@
|`cublasZgemmBatched` | |
|`cublasGemmBatchedEx` | |
|`cublasGemmStridedBatchedEx` | |
|`cublasSgemmStridedBatched` | |
|`cublasDgemmStridedBatched` | |
|`cublasSgemmStridedBatched` |`hipblasSgemmStridedBatched` |
|`cublasDgemmStridedBatched` |`hipblasDgemmStridedBatched` |
|`cublasCgemmStridedBatched` | |
|`cublasCgemm3mStridedBatched` | |
|`cublasZgemmStridedBatched` | |
projects/hip/docs/markdown/hip_porting_guide.md
+2
Bestand weergeven
@@ -418,6 +418,8 @@ HIP support for hipMemcpyToSymbol is complete. This feature allows a kernel
to define a device-side data symbol which can be accessed on the host side. The symbol
can be in __constant or device space.
Note that the symbol name needs to be encased in the HIP_SYMBOL macro, as shown in the code example below. This also applies to hipMemcpyFromSymbol, hipGetSymbolAddress, and hipGetSymbolSize.
For example:
Device Code:
projects/hip/docs/markdown/hip_programming_guide.md
+5 -1
Bestand weergeven
@@ -109,4 +109,8 @@ freed.
The memory are allocated in pages. Users can define macro
`__HIP_SIZE_OF_PAGE` for controlling the page size in bytes and macro
`__HIP_NUM_PAGES` for controlling the total number of pages that can be
allocated.
allocated.
## Use of Long Double Type
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.
projects/hip/hipify-clang/src/CUDA2HIP.cpp
+2 -2
Bestand weergeven
@@ -37,8 +37,8 @@ const std::map <llvm::StringRef, hipCounter> CUDA_INCLUDE_MAP{
// cuComplex includes
{"cuComplex.h", {"hip/hip_complex.h", "", CONV_INCLUDE_CUDA_MAIN_H, API_COMPLEX}},
// cuBLAS includes
{"cublas.h", {"hipblas.h", "", CONV_INCLUDE_CUDA_MAIN_H, API_BLAS}},
{"cublas_v2.h", {"hipblas.h", "", CONV_INCLUDE_CUDA_MAIN_H, API_BLAS}},
{"cublas.h", {"hipblas.h", "rocblas.h", CONV_INCLUDE_CUDA_MAIN_H, API_BLAS}},
{"cublas_v2.h", {"hipblas.h", "rocblas.h", CONV_INCLUDE_CUDA_MAIN_H, API_BLAS}},
// cuRAND includes
{"curand.h", {"hiprand.h", "", CONV_INCLUDE_CUDA_MAIN_H, API_RAND}},
{"curand_kernel.h", {"hiprand_kernel.h", "", CONV_INCLUDE, API_RAND}},
projects/hip/hipify-clang/src/CUDA2HIP_BLAS_API_functions.cpp
+454 -454
Bestand weergeven
Diff onderdrukt omdat het te groot bestand Laad Diff
projects/hip/hipify-clang/src/CUDA2HIP_BLAS_API_types.cpp
+98 -103
Bestand weergeven
@@ -24,130 +24,125 @@ THE SOFTWARE.
// Map of all functions
const std::map<llvm::StringRef, hipCounter> CUDA_BLAS_TYPE_NAME_MAP{
////////////////////// Blas Data Types
// Blas operations
{"cublasOperation_t", {"hipblasOperation_t", "", CONV_TYPE, API_BLAS}},
{"CUBLAS_OP_N", {"HIPBLAS_OP_N", "", CONV_NUMERIC_LITERAL, API_BLAS}},
{"CUBLAS_OP_T", {"HIPBLAS_OP_T", "", CONV_NUMERIC_LITERAL, API_BLAS}},
{"CUBLAS_OP_C", {"HIPBLAS_OP_C", "", CONV_NUMERIC_LITERAL, API_BLAS}},
{"cublasOperation_t", {"hipblasOperation_t", "rocblas_operation", CONV_TYPE, API_BLAS}},
{"CUBLAS_OP_N", {"HIPBLAS_OP_N", "rocblas_operation_none", CONV_NUMERIC_LITERAL, API_BLAS}},
{"CUBLAS_OP_T", {"HIPBLAS_OP_T", "rocblas_operation_transpose", CONV_NUMERIC_LITERAL, API_BLAS}},
{"CUBLAS_OP_C", {"HIPBLAS_OP_C", "rocblas_operation_conjugate_transpose", CONV_NUMERIC_LITERAL, API_BLAS}},
// Blas statuses
{"cublasStatus", {"hipblasStatus_t", "", CONV_TYPE, API_BLAS}},
{"cublasStatus_t", {"hipblasStatus_t", "", CONV_TYPE, API_BLAS}},
{"CUBLAS_STATUS_SUCCESS", {"HIPBLAS_STATUS_SUCCESS", "", CONV_NUMERIC_LITERAL, API_BLAS}},
{"CUBLAS_STATUS_NOT_INITIALIZED", {"HIPBLAS_STATUS_NOT_INITIALIZED", "", CONV_NUMERIC_LITERAL, API_BLAS}},
{"CUBLAS_STATUS_ALLOC_FAILED", {"HIPBLAS_STATUS_ALLOC_FAILED", "", CONV_NUMERIC_LITERAL, API_BLAS}},
{"CUBLAS_STATUS_INVALID_VALUE", {"HIPBLAS_STATUS_INVALID_VALUE", "", CONV_NUMERIC_LITERAL, API_BLAS}},
{"CUBLAS_STATUS_MAPPING_ERROR", {"HIPBLAS_STATUS_MAPPING_ERROR", "", CONV_NUMERIC_LITERAL, API_BLAS}},
{"CUBLAS_STATUS_EXECUTION_FAILED", {"HIPBLAS_STATUS_EXECUTION_FAILED", "", CONV_NUMERIC_LITERAL, API_BLAS}},
{"CUBLAS_STATUS_INTERNAL_ERROR", {"HIPBLAS_STATUS_INTERNAL_ERROR", "", CONV_NUMERIC_LITERAL, API_BLAS}},
{"CUBLAS_STATUS_NOT_SUPPORTED", {"HIPBLAS_STATUS_NOT_SUPPORTED", "", CONV_NUMERIC_LITERAL, API_BLAS}},
{"CUBLAS_STATUS_ARCH_MISMATCH", {"HIPBLAS_STATUS_ARCH_MISMATCH", "", CONV_NUMERIC_LITERAL, API_BLAS}},
{"CUBLAS_STATUS_LICENSE_ERROR", {"HIPBLAS_STATUS_LICENSE_ERROR", "", CONV_NUMERIC_LITERAL, API_BLAS, HIP_UNSUPPORTED}},
{"cublasStatus", {"hipblasStatus_t", "rocblas_status", CONV_TYPE, API_BLAS}},
{"cublasStatus_t", {"hipblasStatus_t", "rocblas_status", CONV_TYPE, API_BLAS}},
{"CUBLAS_STATUS_SUCCESS", {"HIPBLAS_STATUS_SUCCESS", "rocblas_status_success", CONV_NUMERIC_LITERAL, API_BLAS}},
{"CUBLAS_STATUS_NOT_INITIALIZED", {"HIPBLAS_STATUS_NOT_INITIALIZED", "rocblas_status_invalid_handle", CONV_NUMERIC_LITERAL, API_BLAS}},
{"CUBLAS_STATUS_ALLOC_FAILED", {"HIPBLAS_STATUS_ALLOC_FAILED", "rocblas_status_memory_error", CONV_NUMERIC_LITERAL, API_BLAS}},
{"CUBLAS_STATUS_INVALID_VALUE", {"HIPBLAS_STATUS_INVALID_VALUE", "rocblas_status_invalid_pointer", CONV_NUMERIC_LITERAL, API_BLAS}},
{"CUBLAS_STATUS_MAPPING_ERROR", {"HIPBLAS_STATUS_MAPPING_ERROR", "rocblas_status_internal_error", CONV_NUMERIC_LITERAL, API_BLAS}},
{"CUBLAS_STATUS_EXECUTION_FAILED", {"HIPBLAS_STATUS_EXECUTION_FAILED", "rocblas_status_internal_error", CONV_NUMERIC_LITERAL, API_BLAS}},
{"CUBLAS_STATUS_INTERNAL_ERROR", {"HIPBLAS_STATUS_INTERNAL_ERROR", "rocblas_status_internal_error", CONV_NUMERIC_LITERAL, API_BLAS}},
{"CUBLAS_STATUS_NOT_SUPPORTED", {"HIPBLAS_STATUS_NOT_SUPPORTED", "rocblas_status_not_implemented", CONV_NUMERIC_LITERAL, API_BLAS}},
{"CUBLAS_STATUS_ARCH_MISMATCH", {"HIPBLAS_STATUS_ARCH_MISMATCH", "rocblas_status_not_implemented", CONV_NUMERIC_LITERAL, API_BLAS}},
{"CUBLAS_STATUS_LICENSE_ERROR", {"HIPBLAS_STATUS_LICENSE_ERROR", "rocblas_status_not_implemented", CONV_NUMERIC_LITERAL, API_BLAS, UNSUPPORTED}},
// Blas Fill Modes
{"cublasFillMode_t", {"hipblasFillMode_t", "", CONV_TYPE, API_BLAS}},
{"CUBLAS_FILL_MODE_LOWER", {"HIPBLAS_FILL_MODE_LOWER", "", CONV_NUMERIC_LITERAL, API_BLAS}},
{"CUBLAS_FILL_MODE_UPPER", {"HIPBLAS_FILL_MODE_UPPER", "", CONV_NUMERIC_LITERAL, API_BLAS}},
{"cublasFillMode_t", {"hipblasFillMode_t", "rocblas_fill", CONV_TYPE, API_BLAS}},
{"CUBLAS_FILL_MODE_LOWER", {"HIPBLAS_FILL_MODE_LOWER", "rocblas_fill_lower", CONV_NUMERIC_LITERAL, API_BLAS}},
{"CUBLAS_FILL_MODE_UPPER", {"HIPBLAS_FILL_MODE_UPPER", "rocblas_fill_upper", CONV_NUMERIC_LITERAL, API_BLAS}},
// Blas Diag Types
{"cublasDiagType_t", {"hipblasDiagType_t", "", CONV_TYPE, API_BLAS}},
{"CUBLAS_DIAG_NON_UNIT", {"HIPBLAS_DIAG_NON_UNIT", "", CONV_NUMERIC_LITERAL, API_BLAS}},
{"CUBLAS_DIAG_UNIT", {"HIPBLAS_DIAG_UNIT", "", CONV_NUMERIC_LITERAL, API_BLAS}},
{"cublasDiagType_t", {"hipblasDiagType_t", "rocblas_diagonal", CONV_TYPE, API_BLAS}},
{"CUBLAS_DIAG_NON_UNIT", {"HIPBLAS_DIAG_NON_UNIT", "rocblas_diagonal_non_unit", CONV_NUMERIC_LITERAL, API_BLAS}},
{"CUBLAS_DIAG_UNIT", {"HIPBLAS_DIAG_UNIT", "rocblas_diagonal_unit", CONV_NUMERIC_LITERAL, API_BLAS}},
// Blas Side Modes
{"cublasSideMode_t", {"hipblasSideMode_t", "", CONV_TYPE, API_BLAS}},
{"CUBLAS_SIDE_LEFT", {"HIPBLAS_SIDE_LEFT", "", CONV_NUMERIC_LITERAL, API_BLAS}},
{"CUBLAS_SIDE_RIGHT", {"HIPBLAS_SIDE_RIGHT", "", CONV_NUMERIC_LITERAL, API_BLAS}},
{"cublasSideMode_t", {"hipblasSideMode_t", "rocblas_side", CONV_TYPE, API_BLAS}},
{"CUBLAS_SIDE_LEFT", {"HIPBLAS_SIDE_LEFT", "rocblas_side_left", CONV_NUMERIC_LITERAL, API_BLAS}},
{"CUBLAS_SIDE_RIGHT", {"HIPBLAS_SIDE_RIGHT", "rocblas_side_right", CONV_NUMERIC_LITERAL, API_BLAS}},
// Blas Pointer Modes
{"cublasPointerMode_t", {"hipblasPointerMode_t", "", CONV_TYPE, API_BLAS}},
{"CUBLAS_POINTER_MODE_HOST", {"HIPBLAS_POINTER_MODE_HOST", "", CONV_NUMERIC_LITERAL, API_BLAS}},
{"CUBLAS_POINTER_MODE_DEVICE", {"HIPBLAS_POINTER_MODE_DEVICE", "", CONV_NUMERIC_LITERAL, API_BLAS}},
{"cublasPointerMode_t", {"hipblasPointerMode_t", "rocblas_pointer_mode", CONV_TYPE, API_BLAS}},
{"CUBLAS_POINTER_MODE_HOST", {"HIPBLAS_POINTER_MODE_HOST", "rocblas_pointer_mode_host", CONV_NUMERIC_LITERAL, API_BLAS}},
{"CUBLAS_POINTER_MODE_DEVICE", {"HIPBLAS_POINTER_MODE_DEVICE", "rocblas_pointer_mode_device", CONV_NUMERIC_LITERAL, API_BLAS}},
// Blas Atomics Modes
{"cublasAtomicsMode_t", {"hipblasAtomicsMode_t", "", CONV_TYPE, API_BLAS, HIP_UNSUPPORTED}},
{"CUBLAS_ATOMICS_NOT_ALLOWED", {"HIPBLAS_ATOMICS_NOT_ALLOWED", "", CONV_NUMERIC_LITERAL, API_BLAS, HIP_UNSUPPORTED}},
{"CUBLAS_ATOMICS_ALLOWED", {"HIPBLAS_ATOMICS_ALLOWED", "", CONV_NUMERIC_LITERAL, API_BLAS, HIP_UNSUPPORTED}},
{"cublasAtomicsMode_t", {"hipblasAtomicsMode_t", "rocblas_atomics_mode", CONV_TYPE, API_BLAS, HIP_UNSUPPORTED}},
{"CUBLAS_ATOMICS_NOT_ALLOWED", {"HIPBLAS_ATOMICS_NOT_ALLOWED", "rocblas_atomics_not_allowed", CONV_NUMERIC_LITERAL, API_BLAS, HIP_UNSUPPORTED}},
{"CUBLAS_ATOMICS_ALLOWED", {"HIPBLAS_ATOMICS_ALLOWED", "rocblas_atomics_allowed", CONV_NUMERIC_LITERAL, API_BLAS, HIP_UNSUPPORTED}},
// Blas Data Type
{"cublasDataType_t", {"hipblasDatatype_t", "", CONV_TYPE, API_BLAS}},
{"cublasDataType_t", {"hipblasDatatype_t", "rocblas_datatype", CONV_TYPE, API_BLAS}},
// Blas Math mode/tensor operation
{"cublasMath_t", {"hipblasMath_t", "", CONV_TYPE, API_BLAS, HIP_UNSUPPORTED}},
{"CUBLAS_DEFAULT_MATH", {"HIPBLAS_DEFAULT_MATH", "", CONV_NUMERIC_LITERAL, API_BLAS, HIP_UNSUPPORTED}},
{"CUBLAS_TENSOR_OP_MATH", {"HIPBLAS_TENSOR_OP_MATH", "", CONV_NUMERIC_LITERAL, API_BLAS, HIP_UNSUPPORTED}},
{"cublasMath_t", {"hipblasMath_t", "", CONV_TYPE, API_BLAS, UNSUPPORTED}},
{"CUBLAS_DEFAULT_MATH", {"HIPBLAS_DEFAULT_MATH", "", CONV_NUMERIC_LITERAL, API_BLAS, UNSUPPORTED}},
{"CUBLAS_TENSOR_OP_MATH", {"HIPBLAS_TENSOR_OP_MATH", "", CONV_NUMERIC_LITERAL, API_BLAS, UNSUPPORTED}},
// Blass different GEMM algorithms
{"cublasGemmAlgo_t", {"hipblasGemmAlgo_t", "", CONV_TYPE, API_BLAS}},
{"CUBLAS_GEMM_DFALT", {"HIPBLAS_GEMM_DEFAULT", "", CONV_NUMERIC_LITERAL, API_BLAS}}, // -1 // 160
{"CUBLAS_GEMM_DEFAULT", {"HIPBLAS_GEMM_DEFAULT", "", CONV_NUMERIC_LITERAL, API_BLAS}}, // -1 // 160
{"CUBLAS_GEMM_ALGO0", {"HIPBLAS_GEMM_ALGO0", "", CONV_NUMERIC_LITERAL, API_BLAS, HIP_UNSUPPORTED}}, // 0
{"CUBLAS_GEMM_ALGO1", {"HIPBLAS_GEMM_ALGO1", "", CONV_NUMERIC_LITERAL, API_BLAS, HIP_UNSUPPORTED}}, // 1
{"CUBLAS_GEMM_ALGO2", {"HIPBLAS_GEMM_ALGO2", "", CONV_NUMERIC_LITERAL, API_BLAS, HIP_UNSUPPORTED}}, // 2
{"CUBLAS_GEMM_ALGO3", {"HIPBLAS_GEMM_ALGO3", "", CONV_NUMERIC_LITERAL, API_BLAS, HIP_UNSUPPORTED}}, // 3
{"CUBLAS_GEMM_ALGO4", {"HIPBLAS_GEMM_ALGO4", "", CONV_NUMERIC_LITERAL, API_BLAS, HIP_UNSUPPORTED}}, // 4
{"CUBLAS_GEMM_ALGO5", {"HIPBLAS_GEMM_ALGO5", "", CONV_NUMERIC_LITERAL, API_BLAS, HIP_UNSUPPORTED}}, // 5
{"CUBLAS_GEMM_ALGO6", {"HIPBLAS_GEMM_ALGO6", "", CONV_NUMERIC_LITERAL, API_BLAS, HIP_UNSUPPORTED}}, // 6
{"CUBLAS_GEMM_ALGO7", {"HIPBLAS_GEMM_ALGO7", "", CONV_NUMERIC_LITERAL, API_BLAS, HIP_UNSUPPORTED}}, // 7
{"CUBLAS_GEMM_ALGO8", {"HIPBLAS_GEMM_ALGO8", "", CONV_NUMERIC_LITERAL, API_BLAS, HIP_UNSUPPORTED}}, // 8
{"CUBLAS_GEMM_ALGO9", {"HIPBLAS_GEMM_ALGO9", "", CONV_NUMERIC_LITERAL, API_BLAS, HIP_UNSUPPORTED}}, // 9
{"CUBLAS_GEMM_ALGO10", {"HIPBLAS_GEMM_ALGO10", "", CONV_NUMERIC_LITERAL, API_BLAS, HIP_UNSUPPORTED}}, // 10
{"CUBLAS_GEMM_ALGO11", {"HIPBLAS_GEMM_ALGO11", "", CONV_NUMERIC_LITERAL, API_BLAS, HIP_UNSUPPORTED}}, // 11
{"CUBLAS_GEMM_ALGO12", {"HIPBLAS_GEMM_ALGO12", "", CONV_NUMERIC_LITERAL, API_BLAS, HIP_UNSUPPORTED}}, // 12
{"CUBLAS_GEMM_ALGO13", {"HIPBLAS_GEMM_ALGO13", "", CONV_NUMERIC_LITERAL, API_BLAS, HIP_UNSUPPORTED}}, // 13
{"CUBLAS_GEMM_ALGO14", {"HIPBLAS_GEMM_ALGO14", "", CONV_NUMERIC_LITERAL, API_BLAS, HIP_UNSUPPORTED}}, // 14
{"CUBLAS_GEMM_ALGO15", {"HIPBLAS_GEMM_ALGO15", "", CONV_NUMERIC_LITERAL, API_BLAS, HIP_UNSUPPORTED}}, // 15
{"CUBLAS_GEMM_ALGO16", {"HIPBLAS_GEMM_ALGO16", "", CONV_NUMERIC_LITERAL, API_BLAS, HIP_UNSUPPORTED}}, // 16
{"CUBLAS_GEMM_ALGO17", {"HIPBLAS_GEMM_ALGO17", "", CONV_NUMERIC_LITERAL, API_BLAS, HIP_UNSUPPORTED}}, // 17
{"CUBLAS_GEMM_ALGO18", {"HIPBLAS_GEMM_ALGO18", "", CONV_NUMERIC_LITERAL, API_BLAS, HIP_UNSUPPORTED}}, // 18
{"CUBLAS_GEMM_ALGO19", {"HIPBLAS_GEMM_ALGO19", "", CONV_NUMERIC_LITERAL, API_BLAS, HIP_UNSUPPORTED}}, // 19
{"CUBLAS_GEMM_ALGO20", {"HIPBLAS_GEMM_ALGO20", "", CONV_NUMERIC_LITERAL, API_BLAS, HIP_UNSUPPORTED}}, // 20
{"CUBLAS_GEMM_ALGO21", {"HIPBLAS_GEMM_ALGO21", "", CONV_NUMERIC_LITERAL, API_BLAS, HIP_UNSUPPORTED}}, // 21
{"CUBLAS_GEMM_ALGO22", {"HIPBLAS_GEMM_ALGO22", "", CONV_NUMERIC_LITERAL, API_BLAS, HIP_UNSUPPORTED}}, // 22
{"CUBLAS_GEMM_ALGO23", {"HIPBLAS_GEMM_ALGO23", "", CONV_NUMERIC_LITERAL, API_BLAS, HIP_UNSUPPORTED}}, // 23
{"CUBLAS_GEMM_DEFAULT_TENSOR_OP", {"HIPBLAS_GEMM_DEFAULT_TENSOR_OP", "", CONV_NUMERIC_LITERAL, API_BLAS, HIP_UNSUPPORTED}}, // 99
{"CUBLAS_GEMM_DFALT_TENSOR_OP", {"HIPBLAS_GEMM_DFALT_TENSOR_OP", "", CONV_NUMERIC_LITERAL, API_BLAS, HIP_UNSUPPORTED}}, // 99
{"CUBLAS_GEMM_ALGO0_TENSOR_OP", {"HIPBLAS_GEMM_ALGO0_TENSOR_OP", "", CONV_NUMERIC_LITERAL, API_BLAS, HIP_UNSUPPORTED}}, // 100
{"CUBLAS_GEMM_ALGO1_TENSOR_OP", {"HIPBLAS_GEMM_ALGO1_TENSOR_OP", "", CONV_NUMERIC_LITERAL, API_BLAS, HIP_UNSUPPORTED}}, // 101
{"CUBLAS_GEMM_ALGO2_TENSOR_OP", {"HIPBLAS_GEMM_ALGO2_TENSOR_OP", "", CONV_NUMERIC_LITERAL, API_BLAS, HIP_UNSUPPORTED}}, // 102
{"CUBLAS_GEMM_ALGO3_TENSOR_OP", {"HIPBLAS_GEMM_ALGO3_TENSOR_OP", "", CONV_NUMERIC_LITERAL, API_BLAS, HIP_UNSUPPORTED}}, // 103
{"CUBLAS_GEMM_ALGO4_TENSOR_OP", {"HIPBLAS_GEMM_ALGO4_TENSOR_OP", "", CONV_NUMERIC_LITERAL, API_BLAS, HIP_UNSUPPORTED}}, // 104
{"CUBLAS_GEMM_ALGO5_TENSOR_OP", {"HIPBLAS_GEMM_ALGO5_TENSOR_OP", "", CONV_NUMERIC_LITERAL, API_BLAS, HIP_UNSUPPORTED}}, // 105
{"CUBLAS_GEMM_ALGO6_TENSOR_OP", {"HIPBLAS_GEMM_ALGO6_TENSOR_OP", "", CONV_NUMERIC_LITERAL, API_BLAS, HIP_UNSUPPORTED}}, // 106
{"CUBLAS_GEMM_ALGO7_TENSOR_OP", {"HIPBLAS_GEMM_ALGO7_TENSOR_OP", "", CONV_NUMERIC_LITERAL, API_BLAS, HIP_UNSUPPORTED}}, // 107
{"CUBLAS_GEMM_ALGO8_TENSOR_OP", {"HIPBLAS_GEMM_ALGO8_TENSOR_OP", "", CONV_NUMERIC_LITERAL, API_BLAS, HIP_UNSUPPORTED}}, // 108
{"CUBLAS_GEMM_ALGO9_TENSOR_OP", {"HIPBLAS_GEMM_ALGO9_TENSOR_OP", "", CONV_NUMERIC_LITERAL, API_BLAS, HIP_UNSUPPORTED}}, // 109
{"CUBLAS_GEMM_ALGO10_TENSOR_OP", {"HIPBLAS_GEMM_ALGO10_TENSOR_OP", "", CONV_NUMERIC_LITERAL, API_BLAS, HIP_UNSUPPORTED}}, // 110
{"CUBLAS_GEMM_ALGO11_TENSOR_OP", {"HIPBLAS_GEMM_ALGO11_TENSOR_OP", "", CONV_NUMERIC_LITERAL, API_BLAS, HIP_UNSUPPORTED}}, // 111
{"CUBLAS_GEMM_ALGO12_TENSOR_OP", {"HIPBLAS_GEMM_ALGO12_TENSOR_OP", "", CONV_NUMERIC_LITERAL, API_BLAS, HIP_UNSUPPORTED}}, // 112
{"CUBLAS_GEMM_ALGO13_TENSOR_OP", {"HIPBLAS_GEMM_ALGO13_TENSOR_OP", "", CONV_NUMERIC_LITERAL, API_BLAS, HIP_UNSUPPORTED}}, // 113
{"CUBLAS_GEMM_ALGO14_TENSOR_OP", {"HIPBLAS_GEMM_ALGO14_TENSOR_OP", "", CONV_NUMERIC_LITERAL, API_BLAS, HIP_UNSUPPORTED}}, // 114
{"CUBLAS_GEMM_ALGO15_TENSOR_OP", {"HIPBLAS_GEMM_ALGO15_TENSOR_OP", "", CONV_NUMERIC_LITERAL, API_BLAS, HIP_UNSUPPORTED}}, // 115
{"cublasGemmAlgo_t", {"hipblasGemmAlgo_t", "rocblas_gemm_algo", CONV_TYPE, API_BLAS}},
{"CUBLAS_GEMM_DFALT", {"HIPBLAS_GEMM_DEFAULT", "rocblas_gemm_algo_standard", CONV_NUMERIC_LITERAL, API_BLAS}}, // -1 // 160 // 0b0000000000
{"CUBLAS_GEMM_DEFAULT", {"HIPBLAS_GEMM_DEFAULT", "rocblas_gemm_algo_standard", CONV_NUMERIC_LITERAL, API_BLAS}}, // -1 // 160 // 0b0000000000
{"CUBLAS_GEMM_ALGO0", {"HIPBLAS_GEMM_ALGO0", "", CONV_NUMERIC_LITERAL, API_BLAS, UNSUPPORTED}}, // 0
{"CUBLAS_GEMM_ALGO1", {"HIPBLAS_GEMM_ALGO1", "", CONV_NUMERIC_LITERAL, API_BLAS, UNSUPPORTED}}, // 1
{"CUBLAS_GEMM_ALGO2", {"HIPBLAS_GEMM_ALGO2", "", CONV_NUMERIC_LITERAL, API_BLAS, UNSUPPORTED}}, // 2
{"CUBLAS_GEMM_ALGO3", {"HIPBLAS_GEMM_ALGO3", "", CONV_NUMERIC_LITERAL, API_BLAS, UNSUPPORTED}}, // 3
{"CUBLAS_GEMM_ALGO4", {"HIPBLAS_GEMM_ALGO4", "", CONV_NUMERIC_LITERAL, API_BLAS, UNSUPPORTED}}, // 4
{"CUBLAS_GEMM_ALGO5", {"HIPBLAS_GEMM_ALGO5", "", CONV_NUMERIC_LITERAL, API_BLAS, UNSUPPORTED}}, // 5
{"CUBLAS_GEMM_ALGO6", {"HIPBLAS_GEMM_ALGO6", "", CONV_NUMERIC_LITERAL, API_BLAS, UNSUPPORTED}}, // 6
{"CUBLAS_GEMM_ALGO7", {"HIPBLAS_GEMM_ALGO7", "", CONV_NUMERIC_LITERAL, API_BLAS, UNSUPPORTED}}, // 7
{"CUBLAS_GEMM_ALGO8", {"HIPBLAS_GEMM_ALGO8", "", CONV_NUMERIC_LITERAL, API_BLAS, UNSUPPORTED}}, // 8
{"CUBLAS_GEMM_ALGO9", {"HIPBLAS_GEMM_ALGO9", "", CONV_NUMERIC_LITERAL, API_BLAS, UNSUPPORTED}}, // 9
{"CUBLAS_GEMM_ALGO10", {"HIPBLAS_GEMM_ALGO10", "", CONV_NUMERIC_LITERAL, API_BLAS, UNSUPPORTED}}, // 10
{"CUBLAS_GEMM_ALGO11", {"HIPBLAS_GEMM_ALGO11", "", CONV_NUMERIC_LITERAL, API_BLAS, UNSUPPORTED}}, // 11
{"CUBLAS_GEMM_ALGO12", {"HIPBLAS_GEMM_ALGO12", "", CONV_NUMERIC_LITERAL, API_BLAS, UNSUPPORTED}}, // 12
{"CUBLAS_GEMM_ALGO13", {"HIPBLAS_GEMM_ALGO13", "", CONV_NUMERIC_LITERAL, API_BLAS, UNSUPPORTED}}, // 13
{"CUBLAS_GEMM_ALGO14", {"HIPBLAS_GEMM_ALGO14", "", CONV_NUMERIC_LITERAL, API_BLAS, UNSUPPORTED}}, // 14
{"CUBLAS_GEMM_ALGO15", {"HIPBLAS_GEMM_ALGO15", "", CONV_NUMERIC_LITERAL, API_BLAS, UNSUPPORTED}}, // 15
{"CUBLAS_GEMM_ALGO16", {"HIPBLAS_GEMM_ALGO16", "", CONV_NUMERIC_LITERAL, API_BLAS, UNSUPPORTED}}, // 16
{"CUBLAS_GEMM_ALGO17", {"HIPBLAS_GEMM_ALGO17", "", CONV_NUMERIC_LITERAL, API_BLAS, UNSUPPORTED}}, // 17
{"CUBLAS_GEMM_ALGO18", {"HIPBLAS_GEMM_ALGO18", "", CONV_NUMERIC_LITERAL, API_BLAS, UNSUPPORTED}}, // 18
{"CUBLAS_GEMM_ALGO19", {"HIPBLAS_GEMM_ALGO19", "", CONV_NUMERIC_LITERAL, API_BLAS, UNSUPPORTED}}, // 19
{"CUBLAS_GEMM_ALGO20", {"HIPBLAS_GEMM_ALGO20", "", CONV_NUMERIC_LITERAL, API_BLAS, UNSUPPORTED}}, // 20
{"CUBLAS_GEMM_ALGO21", {"HIPBLAS_GEMM_ALGO21", "", CONV_NUMERIC_LITERAL, API_BLAS, UNSUPPORTED}}, // 21
{"CUBLAS_GEMM_ALGO22", {"HIPBLAS_GEMM_ALGO22", "", CONV_NUMERIC_LITERAL, API_BLAS, UNSUPPORTED}}, // 22
{"CUBLAS_GEMM_ALGO23", {"HIPBLAS_GEMM_ALGO23", "", CONV_NUMERIC_LITERAL, API_BLAS, UNSUPPORTED}}, // 23
{"CUBLAS_GEMM_DEFAULT_TENSOR_OP", {"HIPBLAS_GEMM_DEFAULT_TENSOR_OP", "", CONV_NUMERIC_LITERAL, API_BLAS, UNSUPPORTED}}, // 99
{"CUBLAS_GEMM_DFALT_TENSOR_OP", {"HIPBLAS_GEMM_DFALT_TENSOR_OP", "", CONV_NUMERIC_LITERAL, API_BLAS, UNSUPPORTED}}, // 99
{"CUBLAS_GEMM_ALGO0_TENSOR_OP", {"HIPBLAS_GEMM_ALGO0_TENSOR_OP", "", CONV_NUMERIC_LITERAL, API_BLAS, UNSUPPORTED}}, // 100
{"CUBLAS_GEMM_ALGO1_TENSOR_OP", {"HIPBLAS_GEMM_ALGO1_TENSOR_OP", "", CONV_NUMERIC_LITERAL, API_BLAS, UNSUPPORTED}}, // 101
{"CUBLAS_GEMM_ALGO2_TENSOR_OP", {"HIPBLAS_GEMM_ALGO2_TENSOR_OP", "", CONV_NUMERIC_LITERAL, API_BLAS, UNSUPPORTED}}, // 102
{"CUBLAS_GEMM_ALGO3_TENSOR_OP", {"HIPBLAS_GEMM_ALGO3_TENSOR_OP", "", CONV_NUMERIC_LITERAL, API_BLAS, UNSUPPORTED}}, // 103
{"CUBLAS_GEMM_ALGO4_TENSOR_OP", {"HIPBLAS_GEMM_ALGO4_TENSOR_OP", "", CONV_NUMERIC_LITERAL, API_BLAS, UNSUPPORTED}}, // 104
{"CUBLAS_GEMM_ALGO5_TENSOR_OP", {"HIPBLAS_GEMM_ALGO5_TENSOR_OP", "", CONV_NUMERIC_LITERAL, API_BLAS, UNSUPPORTED}}, // 105
{"CUBLAS_GEMM_ALGO6_TENSOR_OP", {"HIPBLAS_GEMM_ALGO6_TENSOR_OP", "", CONV_NUMERIC_LITERAL, API_BLAS, UNSUPPORTED}}, // 106
{"CUBLAS_GEMM_ALGO7_TENSOR_OP", {"HIPBLAS_GEMM_ALGO7_TENSOR_OP", "", CONV_NUMERIC_LITERAL, API_BLAS, UNSUPPORTED}}, // 107
{"CUBLAS_GEMM_ALGO8_TENSOR_OP", {"HIPBLAS_GEMM_ALGO8_TENSOR_OP", "", CONV_NUMERIC_LITERAL, API_BLAS, UNSUPPORTED}}, // 108
{"CUBLAS_GEMM_ALGO9_TENSOR_OP", {"HIPBLAS_GEMM_ALGO9_TENSOR_OP", "", CONV_NUMERIC_LITERAL, API_BLAS, UNSUPPORTED}}, // 109
{"CUBLAS_GEMM_ALGO10_TENSOR_OP", {"HIPBLAS_GEMM_ALGO10_TENSOR_OP", "", CONV_NUMERIC_LITERAL, API_BLAS, UNSUPPORTED}}, // 110
{"CUBLAS_GEMM_ALGO11_TENSOR_OP", {"HIPBLAS_GEMM_ALGO11_TENSOR_OP", "", CONV_NUMERIC_LITERAL, API_BLAS, UNSUPPORTED}}, // 111
{"CUBLAS_GEMM_ALGO12_TENSOR_OP", {"HIPBLAS_GEMM_ALGO12_TENSOR_OP", "", CONV_NUMERIC_LITERAL, API_BLAS, UNSUPPORTED}}, // 112
{"CUBLAS_GEMM_ALGO13_TENSOR_OP", {"HIPBLAS_GEMM_ALGO13_TENSOR_OP", "", CONV_NUMERIC_LITERAL, API_BLAS, UNSUPPORTED}}, // 113
{"CUBLAS_GEMM_ALGO14_TENSOR_OP", {"HIPBLAS_GEMM_ALGO14_TENSOR_OP", "", CONV_NUMERIC_LITERAL, API_BLAS, UNSUPPORTED}}, // 114
{"CUBLAS_GEMM_ALGO15_TENSOR_OP", {"HIPBLAS_GEMM_ALGO15_TENSOR_OP", "", CONV_NUMERIC_LITERAL, API_BLAS, UNSUPPORTED}}, // 115
// TODO: rename hipblasDatatype_t to hipDataType_t and move from hipBLAS to HIP
{"cudaDataType_t", {"hipblasDatatype_t", "", CONV_TYPE, API_RUNTIME}},
{"cudaDataType", {"hipblasDatatype_t", "", CONV_TYPE, API_RUNTIME}},
{"CUDA_R_16F", {"HIPBLAS_R_16F", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 2 // 150
{"CUDA_C_16F", {"HIPBLAS_C_16F", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 6 // 153
{"CUDA_R_32F", {"HIPBLAS_R_32F", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 0 // 151
{"CUDA_C_32F", {"HIPBLAS_C_32F", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 4 // 154
{"CUDA_R_64F", {"HIPBLAS_R_64F", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 1 // 152
{"CUDA_C_64F", {"HIPBLAS_C_64F", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 5 // 155
{"CUDA_R_8I", {"HIPBLAS_R_8I", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 3 //
{"CUDA_C_8I", {"HIPBLAS_C_8I", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 7 //
{"CUDA_R_8U", {"HIPBLAS_R_8U", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 8 //
{"CUDA_C_8U", {"HIPBLAS_C_8U", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 9 //
{"CUDA_R_32I", {"HIPBLAS_R_32I", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 10 //
{"CUDA_C_32I", {"HIPBLAS_C_32I", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 11 //
{"CUDA_R_32U", {"HIPBLAS_R_32U", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 12 //
{"CUDA_C_32U", {"HIPBLAS_C_32U", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 13 //
{"cudaDataType_t", {"hipblasDatatype_t", "rocblas_datatype_", CONV_TYPE, API_RUNTIME}},
{"cudaDataType", {"hipblasDatatype_t", "rocblas_datatype", CONV_TYPE, API_RUNTIME}},
{"CUDA_R_16F", {"HIPBLAS_R_16F", "rocblas_datatype_f16_r", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 2 // 150
{"CUDA_C_16F", {"HIPBLAS_C_16F", "rocblas_datatype_f16_c", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 6 // 153
{"CUDA_R_32F", {"HIPBLAS_R_32F", "rocblas_datatype_f32_r", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 0 // 151
{"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
{"cublasHandle_t", {"hipblasHandle_t", "", CONV_TYPE, API_BLAS}},
{"cublasHandle_t", {"hipblasHandle_t", "rocblas_handle", CONV_TYPE, API_BLAS}},
// TODO: dereferencing: typedef struct cublasContext *cublasHandle_t;
// {"cublasContext", {"hipblasHandle_t", "", CONV_TYPE, API_BLAS}},
{"cublasContext", {"hipblasHandle_t", "_rocblas_handle", CONV_TYPE, API_BLAS, HIP_UNSUPPORTED}},
};
projects/hip/hipify-clang/src/HipifyAction.cpp
+10 -22
Bestand weergeven
@@ -41,11 +41,10 @@ void HipifyAction::RewriteString(StringRef s, clang::SourceLocation start) {
StringRef name = s.slice(begin, end);
const auto found = CUDA_RENAMES_MAP().find(name);
if (found != CUDA_RENAMES_MAP().end()) {
StringRef repName = TranslateToRoc ? found->second.rocName : found->second.hipName;
StringRef repName = Statistics::isToRoc(found->second) ? found->second.rocName : found->second.hipName;
hipCounter counter = {"[string literal]", "", ConvTypes::CONV_LITERAL, ApiTypes::API_RUNTIME, found->second.supportDegree};
Statistics::current().incrementCounter(counter, name.str());
if ((!TranslateToRoc && (HIP_UNSUPPORTED != (counter.supportDegree & HIP_UNSUPPORTED))) ||
(TranslateToRoc && (ROC_UNSUPPORTED != (counter.supportDegree & ROC_UNSUPPORTED)))) {
if (!Statistics::isUnsupported(counter)) {
clang::SourceLocation sl = start.getLocWithOffset(begin + 1);
ct::Replacement Rep(SM, sl, name.size(), repName.str());
clang::FullSourceLoc fullSL(sl, SM);
@@ -60,7 +59,7 @@ void HipifyAction::RewriteString(StringRef s, clang::SourceLocation start) {
}
/**
* Look at, and consider altering, a given token.
* Look at, and consider altering, a given token.
*
* If it's not a CUDA identifier, nothing happens.
* If it's an unsupported CUDA identifier, a warning is emitted.
@@ -68,7 +67,7 @@ void HipifyAction::RewriteString(StringRef s, clang::SourceLocation start) {
*/
void HipifyAction::RewriteToken(const clang::Token& t) {
clang::SourceManager& SM = getCompilerInstance().getSourceManager();
// String literals containing CUDA references need fixing...
// String literals containing CUDA references need fixing.
if (t.is(clang::tok::string_literal)) {
StringRef s(t.getLiteralData(), t.getLength());
RewriteString(unquoteStr(s), t.getLocation());
@@ -86,26 +85,16 @@ void HipifyAction::RewriteToken(const clang::Token& t) {
Statistics::current().incrementCounter(found->second, name.str());
clang::SourceLocation sl = t.getLocation();
clang::DiagnosticsEngine& DE = getCompilerInstance().getDiagnostics();
bool bWarn = false;
std::string sWarn = "HIP";
if (TranslateToRoc) {
if ((found->second.supportDegree & ROC_UNSUPPORTED) == ROC_UNSUPPORTED) {
sWarn = "ROCm";
bWarn = true;
}
} else {
if ((found->second.supportDegree & HIP_UNSUPPORTED) == HIP_UNSUPPORTED) {
bWarn = true;
}
}
// Warn the user about unsupported identifier.
if (bWarn) {
if (Statistics::isUnsupported(found->second)) {
std::string sWarn;
Statistics::isToRoc(found->second) ? sWarn = "ROC" : sWarn = "HIP";
sWarn = "" + sWarn;
const auto ID = DE.getCustomDiagID(clang::DiagnosticsEngine::Warning, "CUDA identifier is unsupported in %0.");
DE.Report(sl, ID) << sWarn;
return;
}
StringRef repName = TranslateToRoc ? found->second.rocName : found->second.hipName;
StringRef repName = Statistics::isToRoc(found->second) ? found->second.rocName : found->second.hipName;
ct::Replacement Rep(SM, sl, name.size(), repName.str());
clang::FullSourceLoc fullSL(sl, SM);
insertReplacement(Rep, fullSL);
@@ -238,8 +227,7 @@ void HipifyAction::InclusionDirective(clang::SourceLocation hash_loc,
Statistics::current().incrementCounter(found->second, file_name.str());
clang::SourceLocation sl = filename_range.getBegin();
if ((!TranslateToRoc && (HIP_UNSUPPORTED == (found->second.supportDegree & HIP_UNSUPPORTED))) ||
(TranslateToRoc && (ROC_UNSUPPORTED == (found->second.supportDegree & ROC_UNSUPPORTED)))) {
if (Statistics::isUnsupported(found->second)) {
clang::DiagnosticsEngine& DE = getCompilerInstance().getDiagnostics();
DE.Report(sl, DE.getCustomDiagID(clang::DiagnosticsEngine::Warning, "Unsupported CUDA header."));
return;
@@ -249,7 +237,7 @@ void HipifyAction::InclusionDirective(clang::SourceLocation hash_loc,
// Keep the same include type that the user gave.
if (!exclude) {
clang::SmallString<128> includeBuffer;
llvm::StringRef name = TranslateToRoc ? found->second.rocName : found->second.hipName;
llvm::StringRef name = Statistics::isToRoc(found->second) ? found->second.rocName : found->second.hipName;
if (is_angled) {
newInclude = llvm::Twine("<" + name+ ">").toStringRef(includeBuffer);
} else {
projects/hip/hipify-clang/src/Statistics.cpp
+22 -3
Bestand weergeven
@@ -177,8 +177,7 @@ Statistics::Statistics(const std::string& name): fileName(name) {
///////// Counter update routines //////////
void Statistics::incrementCounter(const hipCounter &counter, const std::string& name) {
if ((!TranslateToRoc && (HIP_UNSUPPORTED == (counter.supportDegree & HIP_UNSUPPORTED))) ||
(TranslateToRoc && (ROC_UNSUPPORTED == (counter.supportDegree & ROC_UNSUPPORTED)))) {
if (Statistics::isUnsupported(counter)) {
unsupported.incrementCounter(counter, name);
} else {
supported.incrementCounter(counter, name);
@@ -202,7 +201,7 @@ void Statistics::add(const Statistics &other) {
void Statistics::lineTouched(int lineNumber) {
touchedLinesSet.insert(lineNumber);
touchedLines = touchedLinesSet.size();
touchedLines = unsigned(touchedLinesSet.size());
}
void Statistics::bytesChanged(int bytes) {
@@ -284,5 +283,25 @@ void Statistics::setActive(const std::string& name) {
Statistics::currentStatistics = &stats.at(name);
}
bool Statistics::isToRoc(const hipCounter &counter) {
return TranslateToRoc && counter.apiType == API_BLAS;
}
bool Statistics::isHipUnsupported(const hipCounter &counter) {
return HIP_UNSUPPORTED == (counter.supportDegree & HIP_UNSUPPORTED);
}
bool Statistics::isRocUnsupported(const hipCounter &counter) {
return ROC_UNSUPPORTED == (counter.supportDegree & ROC_UNSUPPORTED);
}
bool Statistics::isUnsupported(const hipCounter &counter) {
if (Statistics::isToRoc(counter)) {
return Statistics::isRocUnsupported(counter);
} else {
return Statistics::isHipUnsupported(counter);
}
}
std::map<std::string, Statistics> Statistics::stats = {};
Statistics* Statistics::currentStatistics = nullptr;
projects/hip/hipify-clang/src/Statistics.h
+11
Bestand weergeven
@@ -225,6 +225,17 @@ public:
* timestamp into the currently active one.
*/
static void setActive(const std::string& name);
// Check the counter and option TranslateToRoc whether it should be translated to Roc or not.
static bool isToRoc(const hipCounter &counter);
// Check whether the counter is HIP_UNSUPPORTED or not.
static bool isHipUnsupported(const hipCounter &counter);
// Check whether the counter is ROC_UNSUPPORTED or not.
static bool isRocUnsupported(const hipCounter &counter);
/**
* Check whether the counter is ROC_UNSUPPORTED/HIP_UNSUPPORTED/UNSUPPORTED or not
* based on counter's API_TYPE and option TranslateToRoc.
*/
static bool isUnsupported(const hipCounter &counter);
// Set this flag in case of hipification errors
bool hasErrors = false;
};
projects/hip/include/hip/hcc_detail/hip_vector_types.h
+1 -1
Bestand weergeven
@@ -34,7 +34,7 @@ THE SOFTWARE.
#include "hip/hcc_detail/host_defines.h"
#if !defined(_MSC_VER)
#if !defined(_MSC_VER) || __HIP_DEVICE_COMPILE__
#if defined(__clang__)
#define __NATIVE_VECTOR__(n, ...) __attribute__((ext_vector_type(n)))
#elif defined(__GNUC__) // N.B.: GCC does not support .xyzw syntax.
projects/hip/samples/2_Cookbook/8_peer2peer/peer2peer.cpp
+2 -1
Bestand weergeven
@@ -43,7 +43,8 @@ using namespace std;
#define HIPCHECK(error) \
{ \
hipError_t localError = error; \
if ((localError != hipSuccess)||(localError != hipErrorPeerAccessAlreadyEnabled)) { \
if ((localError != hipSuccess)&& (localError != hipErrorPeerAccessAlreadyEnabled)&& \
(localError != hipErrorPeerAccessNotEnabled )) { \
printf("%serror: '%s'(%d) from %s at %s:%d%s\n", KRED, hipGetErrorString(localError), \
localError, #error, __FILE__, __LINE__, KNRM); \
failed("API returned error code."); \
projects/hip/src/hip_clang.cpp
+6 -2
Bestand weergeven
@@ -132,8 +132,12 @@ extern "C" void __hipRegisterFunction(
assert(modules && modules->size() >= g_deviceCnt);
for (int deviceId = 0; deviceId < g_deviceCnt; ++deviceId) {
hipFunction_t function;
if (hipSuccess == hipModuleGetFunction(&function, modules->at(deviceId), deviceName) &&
function != nullptr) {
if ((hipSuccess == hipModuleGetFunction(&function, modules->at(deviceId), deviceName) ||
// With code-object-v3, we need to match the kernel descriptor symbol name
(hipSuccess == hipModuleGetFunction(
&function, modules->at(deviceId),
(std::string(deviceName) + std::string(".kd")).c_str()
))) && function != nullptr) {
functions[deviceId] = function;
}
else {
projects/hip/src/hip_peer.cpp
+1 -1
Bestand weergeven
@@ -129,7 +129,7 @@ hipError_t ihipEnablePeerAccess(hipCtx_t peerCtx, unsigned int flags) {
LockedAccessor_CtxCrit_t peerCrit(peerCtx->criticalData());
// Add thisCtx to peerCtx's access list so that new allocations on peer will be made
// visible to this device:
bool isNewPeer = peerCrit->addPeerWatcher(thisCtx, peerCtx);
bool isNewPeer = peerCrit->addPeerWatcher(peerCtx, thisCtx);
if (isNewPeer) {
tprintf(DB_MEM, "device=%s can now see all memory allocated on peer=%s\n",
thisCtx->toString().c_str(), peerCtx->toString().c_str());
projects/hip/tests/hipify-clang/run_test.bat
+4 -3
Bestand weergeven
@@ -8,12 +8,13 @@ set HIPIFY=%1
set IN_FILE=%2
set TMP_FILE=%3
set CUDA_ROOT=%4
set ROC=%5
set all_args=%*
call set clang_args=%%all_args:*%5=%%
set clang_args=%5%clang_args%
call set clang_args=%%all_args:*%6=%%
set clang_args=%6%clang_args%
%HIPIFY% -o=%TMP_FILE% %IN_FILE% %CUDA_ROOT% -- %clang_args%
%HIPIFY% -o=%TMP_FILE% %IN_FILE% %CUDA_ROOT% %ROC% -- %clang_args%
if errorlevel 1 (echo Error: hipify-clang.exe failed with exit code: %errorlevel% && exit /b %errorlevel%)
findstr /v /r /c:"[ ]*//[ ]*[CHECK*|RUN]" %TMP_FILE% | %FILE_CHECK% %IN_FILE%
projects/hip/tests/hipify-clang/run_test.sh
+3 -2
Bestand weergeven
@@ -10,8 +10,9 @@ HIPIFY=$1
IN_FILE=$2
TMP_FILE=$3
CUDA_ROOT=$4
shift 4
ROC=$5
shift 5
# Remaining args are the ones to forward to clang proper.
$HIPIFY -o=$TMP_FILE $IN_FILE $CUDA_ROOT -- $@ && cat $TMP_FILE | sed -Ee 's|//.+|// |g' | FileCheck $IN_FILE
$HIPIFY -o=$TMP_FILE $IN_FILE $CUDA_ROOT $ROC -- $@ && cat $TMP_FILE | sed -Ee 's|//.+|// |g' | FileCheck $IN_FILE
projects/hip/tests/hipify-clang/unit_tests/libraries/cuBLAS/cublas_0_based_indexing.cu
+2
Bestand weergeven
@@ -1,5 +1,6 @@
// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args
// CHECK: #include <hip/hip_runtime.h>
#include <stdio.h>
#include <stdlib.h>
#include <math.h>
@@ -33,6 +34,7 @@ int main(void) {
}
// cublasInit is not supported yet
cublasInit();
// cublasAlloc is not supported yet
stat = cublasAlloc(M*N, sizeof(*a), (void**)&devPtrA);
// CHECK: if (stat != HIPBLAS_STATUS_SUCCESS) {
if (stat != CUBLAS_STATUS_SUCCESS) {
projects/hip/tests/hipify-clang/unit_tests/libraries/cuBLAS/rocBLAS/cublas_0_based_indexing_rocblas.cu
+81
Bestand weergeven
@@ -0,0 +1,81 @@
// RUN: %run_test hipify "%s" "%t" %hipify_args "-roc" %clang_args
// CHECK: #include <hip/hip_runtime.h>
#include <stdio.h>
#include <stdlib.h>
#include <math.h>
// CHECK: #include "rocblas.h"
#include "cublas.h"
#define M 6
#define N 5
#define IDX2C(i,j,ld) (((j)*(ld))+(i))
static __inline__ void modify(float *m, int ldm, int n, int p, int q, float
alpha, float beta) {
// CHECK: rocblas_sscal(n - p, alpha, &m[IDX2C(p, q, ldm)], ldm);
// CHECK: rocblas_sscal(ldm - p, beta, &m[IDX2C(p, q, ldm)], 1);
cublasSscal(n - p, alpha, &m[IDX2C(p, q, ldm)], ldm);
cublasSscal(ldm - p, beta, &m[IDX2C(p, q, ldm)], 1);
}
int main(void) {
int i, j;
// CHECK: rocblas_status stat;
cublasStatus stat;
float* devPtrA;
float* a = 0;
a = (float *)malloc(M * N * sizeof(*a));
if (!a) {
printf("host memory allocation failed");
return EXIT_FAILURE;
}
for (j = 0; j < N; j++) {
for (i = 0; i < M; i++) {
a[IDX2C(i, j, M)] = (float)(i * M + j + 1);
}
}
// cublasInit is not supported yet
cublasInit();
// cublasAlloc is not supported yet
stat = cublasAlloc(M*N, sizeof(*a), (void**)&devPtrA);
// CHECK: if (stat != rocblas_status_success) {
if (stat != CUBLAS_STATUS_SUCCESS) {
printf("device memory allocation failed");
// cublasShutdown is not supported yet
cublasShutdown();
return EXIT_FAILURE;
}
// CHECK: stat = rocblas_set_matrix(M, N, sizeof(*a), a, M, devPtrA, M);
stat = cublasSetMatrix(M, N, sizeof(*a), a, M, devPtrA, M);
// CHECK: if (stat != rocblas_status_success) {
if (stat != CUBLAS_STATUS_SUCCESS) {
printf("data download failed");
// cublasFree is not supported yet
cublasFree(devPtrA);
// cublasShutdown is not supported yet
cublasShutdown();
return EXIT_FAILURE;
}
modify(devPtrA, M, N, 1, 2, 16.0f, 12.0f);
// CHECK: stat = rocblas_get_matrix(M, N, sizeof(*a), devPtrA, M, a, M);
stat = cublasGetMatrix(M, N, sizeof(*a), devPtrA, M, a, M);
// CHECK: if (stat != rocblas_status_success) {
if (stat != CUBLAS_STATUS_SUCCESS) {
printf("data upload failed");
// cublasFree is not supported yet
cublasFree(devPtrA);
// cublasShutdown is not supported yet
cublasShutdown();
return EXIT_FAILURE;
}
// cublasFree is not supported yet
cublasFree(devPtrA);
// cublasShutdown is not supported yet
cublasShutdown();
for (j = 0; j < N; j++) {
for (i = 0; i < M; i++) {
printf("%7.0f", a[IDX2C(i, j, M)]);
}
printf("\n");
}
free(a);
return EXIT_SUCCESS;
}
projects/hip/tests/hipify-clang/unit_tests/libraries/cuBLAS/rocBLAS/cublas_1_based_indexing_rocblas.cu
+90
Bestand weergeven
@@ -0,0 +1,90 @@
// RUN: %run_test hipify "%s" "%t" %hipify_args "-roc" %clang_args
#include <stdio.h>
#include <stdlib.h>
#include <math.h>
// CHECK: #include <hip/hip_runtime.h>
#include <cuda_runtime.h>
// CHECK: #include "rocblas.h"
#include "cublas_v2.h"
#define M 6
#define N 5
#define IDX2F(i,j,ld) ((((j)-1)*(ld))+((i)-1))
// CHECK: static __inline__ void modify(rocblas_handle handle, float *m, int ldm, int
static __inline__ void modify(cublasHandle_t handle, float *m, int ldm, int
n, int p, int q, float alpha, float beta) {
// CHECK: rocblas_sscal(handle, n - p + 1, &alpha, &m[IDX2F(p, q, ldm)], ldm);
// CHECK: rocblas_sscal(handle, ldm - p + 1, &beta, &m[IDX2F(p, q, ldm)], 1);
cublasSscal(handle, n - p + 1, &alpha, &m[IDX2F(p, q, ldm)], ldm);
cublasSscal(handle, ldm - p + 1, &beta, &m[IDX2F(p, q, ldm)], 1);
}
int main(void) {
// CHECK: hipError_t cudaStat;
// CHECK: rocblas_status stat;
// CHECK: rocblas_handle handle;
cudaError_t cudaStat;
cublasStatus_t stat;
cublasHandle_t handle;
int i, j;
float* devPtrA;
float* a = 0;
a = (float *)malloc(M * N * sizeof(*a));
if (!a) {
printf("host memory allocation failed");
return EXIT_FAILURE;
}
for (j = 1; j <= N; j++) {
for (i = 1; i <= M; i++) {
a[IDX2F(i, j, M)] = (float)((i - 1) * M + j);
}
}
// CHECK: cudaStat = hipMalloc((void**)&devPtrA, M*N * sizeof(*a));
cudaStat = cudaMalloc((void**)&devPtrA, M*N * sizeof(*a));
// CHECK: if (cudaStat != hipSuccess) {
if (cudaStat != cudaSuccess) {
printf("device memory allocation failed");
return EXIT_FAILURE;
}
// CHECK: stat = rocblas_create_handle(&handle);
stat = cublasCreate(&handle);
// CHECK: if (stat != rocblas_status_success) {
if (stat != CUBLAS_STATUS_SUCCESS) {
printf("CUBLAS initialization failed\n");
return EXIT_FAILURE;
}
// CHECK: stat = rocblas_set_matrix(M, N, sizeof(*a), a, M, devPtrA, M);
stat = cublasSetMatrix(M, N, sizeof(*a), a, M, devPtrA, M);
// CHECK: if (stat != rocblas_status_success) {
if (stat != CUBLAS_STATUS_SUCCESS) {
printf("data download failed");
// CHECK: hipFree(devPtrA);
// CHECK: rocblas_destroy_handle(handle);
cudaFree(devPtrA);
cublasDestroy(handle);
return EXIT_FAILURE;
}
modify(handle, devPtrA, M, N, 2, 3, 16.0f, 12.0f);
// CHECK: stat = rocblas_get_matrix(M, N, sizeof(*a), devPtrA, M, a, M);
stat = cublasGetMatrix(M, N, sizeof(*a), devPtrA, M, a, M);
// CHECK: if (stat != rocblas_status_success) {
if (stat != CUBLAS_STATUS_SUCCESS) {
printf("data upload failed");
// CHECK: hipFree(devPtrA);
// CHECK: rocblas_destroy_handle(handle);
cudaFree(devPtrA);
cublasDestroy(handle);
return EXIT_FAILURE;
}
// CHECK: hipFree(devPtrA);
// CHECK: rocblas_destroy_handle(handle);
cudaFree(devPtrA);
cublasDestroy(handle);
for (j = 1; j <= N; j++) {
for (i = 1; i <= M; i++) {
printf("%7.0f", a[IDX2F(i, j, M)]);
}
printf("\n");
}
free(a);
return EXIT_SUCCESS;
}
projects/hip/tests/hipify-clang/unit_tests/libraries/cuBLAS/rocBLAS/cublas_sgemm_matrix_multiplication_rocblas.cu
+108
Bestand weergeven
@@ -0,0 +1,108 @@
// RUN: %run_test hipify "%s" "%t" %hipify_args "-roc" %clang_args
#include <stdio.h>
#include <stdlib.h>
// CHECK: #include <hip/hip_runtime.h>
#include <cuda_runtime.h>
// CHECK: #include "rocblas.h"
#include "cublas_v2.h"
#define IDX2C(i,j,ld) (((j)*(ld))+(i))
#define m 6
#define n 4
#define k 5
int main(void) {
// CHECK: hipError_t cudaStat;
// CHECK: rocblas_status stat;
// CHECK: rocblas_handle handle;
cudaError_t cudaStat;
cublasStatus_t stat;
cublasHandle_t handle;
int i, j;
float * a;
float * b;
float * c;
a = (float *)malloc(m*k * sizeof(float));
b = (float *)malloc(k*n * sizeof(float));
c = (float *)malloc(m*n * sizeof(float));
int ind = 11;
for (j = 0; j<k; j++) {
for (i = 0; i<m; i++) {
a[IDX2C(i, j, m)] = (float)ind++;
}
}
printf("a:\n");
for (i = 0; i<m; i++) {
for (j = 0; j<k; j++) {
printf(" %5.0f", a[IDX2C(i, j, m)]);
}
printf("\n");
}
ind = 11;
for (j = 0; j<n; j++) {
for (i = 0; i<k; i++) {
b[IDX2C(i, j, k)] = (float)ind++;
}
}
printf("b:\n");
for (i = 0; i<k; i++) {
for (j = 0; j<n; j++) {
printf(" %5.0f", b[IDX2C(i, j, k)]);
}
printf("\n");
}
ind = 11;
for (j = 0; j<n; j++) {
for (i = 0; i<m; i++) {
c[IDX2C(i, j, m)] = (float)ind++;
}
}
printf("c:\n");
for (i = 0; i<m; i++) {
for (j = 0; j<n; j++) {
printf(" %5.0f", c[IDX2C(i, j, m)]);
}
printf("\n");
}
float * d_a;
float * d_b;
float * d_c;
// CHECK: cudaStat = hipMalloc((void **)& d_a, m*k * sizeof(*a));
// CHECK: cudaStat = hipMalloc((void **)& d_b, k*n * sizeof(*b));
// CHECK: cudaStat = hipMalloc((void **)& d_c, m*n * sizeof(*c));
cudaStat = cudaMalloc((void **)& d_a, m*k * sizeof(*a));
cudaStat = cudaMalloc((void **)& d_b, k*n * sizeof(*b));
cudaStat = cudaMalloc((void **)& d_c, m*n * sizeof(*c));
// CHECK: stat = rocblas_create_handle(&handle);
stat = cublasCreate(&handle);
// CHECK: stat = rocblas_set_matrix(m, k, sizeof(*a), a, m, d_a, m);
// CHECK: stat = rocblas_set_matrix(k, n, sizeof(*b), b, k, d_b, k);
// CHECK: stat = rocblas_set_matrix(m, n, sizeof(*c), c, m, d_c, m);
stat = cublasSetMatrix(m, k, sizeof(*a), a, m, d_a, m);
stat = cublasSetMatrix(k, n, sizeof(*b), b, k, d_b, k);
stat = cublasSetMatrix(m, n, sizeof(*c), c, m, d_c, m);
float al = 1.0f;
float bet = 1.0f;
// CHECK: stat = rocblas_sgemm(handle, rocblas_operation_none, rocblas_operation_none, m, n, k, &al, d_a, m, d_b, k, &bet, d_c, m);
stat = cublasSgemm(handle, CUBLAS_OP_N, CUBLAS_OP_N, m, n, k, &al, d_a, m, d_b, k, &bet, d_c, m);
// CHECK: stat = rocblas_get_matrix(m, n, sizeof(*c), d_c, m, c, m);
stat = cublasGetMatrix(m, n, sizeof(*c), d_c, m, c, m);
printf("c after Sgemm :\n");
for (i = 0; i<m; i++) {
for (j = 0; j<n; j++) {
printf(" %7.0f", c[IDX2C(i, j, m)]);
}
printf("\n");
}
// CHECK: hipFree(d_a);
// CHECK: hipFree(d_b);
// CHECK: hipFree(d_c);
// CHECK: rocblas_destroy_handle(handle);
cudaFree(d_a);
cudaFree(d_b);
cudaFree(d_c);
cublasDestroy(handle);
free(a);
free(b);
free(c);
return EXIT_SUCCESS;
}