Merge 'master' into 'amd-master'

Change-Id: Ic79e86ea3cd8cd7f85d7f9f2a24e6acfad02d942

hipamd/docs/markdown/hip_terms.md

@@ -33,7 +33,7 @@
 |Fast Math|`__cos(f)`|`__cos(f)`|`hc::`<br>`fast_math::cos(f)`|`concurrency::`<br>`fast_math::cos(f)`|`native_cos(f)`
 |Vector|`float4`|`float4`|`hc::`<br>`short_vector::float4`|`concurrency::`<br>`graphics::float_4`|`float4`
 
-###Notes
+### Notes
 1. For HC and C++AMP, assume a captured _tiled_ext_ named "t_ext" and captured _extent_ named "ext".  These languages use captured variables to pass information to the kernel rather than using special built-in functions so the exact variable name may vary.
 2. The indexing functions (starting with `thread-index`) show the terminology for a 1D grid.  Some APIs use reverse order of xyz / 012 indexing for 3D grids.
 3. HC allows tile dimensions to be specified at runtime while C++AMP requires that tile dimensions be specified at compile-time.  Thus hc syntax for tile dims is `t_ext.tile_dim[0]` while C++AMP is t_ext.tile_dim0.

hipamd/hipify-clang/src/CUDA2HipMap.cpp

@@ -311,19 +311,9 @@ const std::map<llvm::StringRef, hipCounter> CUDA_TYPE_NAME_MAP{
     // EGL Interoperability
     {"cudaEglStreamConnection", {"hipEglStreamConnection", CONV_EGL, API_RUNTIME, HIP_UNSUPPORTED}},    // API_Driver ANALOGUE (CUeglStreamConnection)
 
-    ///////////////////////////// cuBLAS /////////////////////////////
     {"cublasHandle_t", {"hipblasHandle_t", CONV_TYPE, API_BLAS}},
     // TODO: dereferencing: typedef struct cublasContext *cublasHandle_t;
-    //     {"cublasContext", {"hipblasHandle_t", CONV_TYPE, API_BLAS}},
+    // {"cublasContext", {"hipblasHandle_t", CONV_TYPE, API_BLAS}},
-
-    {"cublasOperation_t",   {"hipblasOperation_t",   CONV_TYPE, API_BLAS}},
-    {"cublasStatus_t",      {"hipblasStatus_t",      CONV_TYPE, API_BLAS}},
-    {"cublasFillMode_t",    {"hipblasFillMode_t",    CONV_TYPE, API_BLAS}},
-    {"cublasDiagType_t",    {"hipblasDiagType_t",    CONV_TYPE, API_BLAS}},
-    {"cublasSideMode_t",    {"hipblasSideMode_t",    CONV_TYPE, API_BLAS}},
-    {"cublasPointerMode_t", {"hipblasPointerMode_t", CONV_TYPE, API_BLAS}},
-    {"cublasAtomicsMode_t", {"hipblasAtomicsMode_t", CONV_TYPE, API_BLAS, HIP_UNSUPPORTED}},
-    {"cublasDataType_t",    {"hipblasDataType_t",    CONV_TYPE, API_BLAS, HIP_UNSUPPORTED}},
 
     ///////////////////////////// cuRAND /////////////////////////////
     {"curandStatus",                  {"hiprandStatus_t",                CONV_TYPE, API_RAND}},
@@ -2087,14 +2077,14 @@ const std::map<llvm::StringRef, hipCounter> CUDA_IDENTIFIER_MAP{
     {"cublasGetMathMode",              {"hipblasGetMathMode",              CONV_MATH_FUNC,       API_BLAS, HIP_UNSUPPORTED}},
     {"cublasSetMathMode",              {"hipblasSetMathMode",              CONV_MATH_FUNC,       API_BLAS, HIP_UNSUPPORTED}},
 
-
+    // Blas operations
-
+    {"cublasOperation_t",              {"hipblasOperation_t",              CONV_TYPE, API_BLAS}},
-    // Blas operations (cublasOperation_t)
     {"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}},
 
-    // Blas statuses (cublasStatus_t)
+    // Blas statuses
+    {"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}},
@@ -2102,34 +2092,76 @@ const std::map<llvm::StringRef, hipCounter> CUDA_IDENTIFIER_MAP{
     {"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_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}},
 
-    // Blas Fill Modes (cublasFillMode_t)
+    // 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}},
 
-    // Blas Diag Types (cublasDiagType_t)
+    // 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}},
 
-    // Blas Side Modes (cublasSideMode_t
+    // 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}},
 
-    // Blas Pointer Modes (cublasPointerMode_t)
+    // 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}},
 
-    // Blas Atomics Modes (cublasAtomicsMode_t)
+    // 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}},
 
-    // Blas Data Type (cublasDataType_t)
+    // Blas Data Type
+    {"cublasDataType_t",               {"hipblasDataType_t",               CONV_TYPE, API_BLAS, HIP_UNSUPPORTED}},
     {"CUBLAS_DATA_FLOAT",              {"HIPBLAS_DATA_FLOAT",              CONV_NUMERIC_LITERAL, API_BLAS, HIP_UNSUPPORTED}},
     {"CUBLAS_DATA_DOUBLE",             {"HIPBLAS_DATA_DOUBLE",             CONV_NUMERIC_LITERAL, API_BLAS, HIP_UNSUPPORTED}},
     {"CUBLAS_DATA_HALF",               {"HIPBLAS_DATA_HALF",               CONV_NUMERIC_LITERAL, API_BLAS, HIP_UNSUPPORTED}},
     {"CUBLAS_DATA_INT8",               {"HIPBLAS_DATA_INT8",               CONV_NUMERIC_LITERAL, API_BLAS, HIP_UNSUPPORTED}},
 
+    // 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}},
+
+    // Blass different GEMM algorithms
+    {"cublasGemmAlgo_t",               {"hipblasGemmAlgo_t",               CONV_TYPE, API_BLAS, HIP_UNSUPPORTED}},
+    {"CUBLAS_GEMM_DFALT",              {"HIPBLAS_GEMM_DFALT",              CONV_NUMERIC_LITERAL, API_BLAS, HIP_UNSUPPORTED}},  //  -1
+    {"CUBLAS_GEMM_DEFAULT",            {"HIPBLAS_GEMM_DEFAULT",            CONV_NUMERIC_LITERAL, API_BLAS, HIP_UNSUPPORTED}},  //  -1
+    {"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_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
+
     // Blas1 (v1) Routines
     {"cublasCreate",                   {"hipblasCreate",                   CONV_MATH_FUNC,       API_BLAS}},
     {"cublasDestroy",                  {"hipblasDestroy",                  CONV_MATH_FUNC,       API_BLAS}},
@@ -2354,7 +2386,7 @@ const std::map<llvm::StringRef, hipCounter> CUDA_IDENTIFIER_MAP{
 
     {"cublasCgemm",                    {"hipblasCgemm",                    CONV_MATH_FUNC,       API_BLAS}},
     {"cublasZgemm",                    {"hipblasZgemm",                    CONV_MATH_FUNC,       API_BLAS, HIP_UNSUPPORTED}},
-    {"cublasHgemm",                    {"hipblasHgemm",                    CONV_MATH_FUNC,       API_BLAS, HIP_UNSUPPORTED}},
+    {"cublasHgemm",                    {"hipblasHgemm",                    CONV_MATH_FUNC,       API_BLAS}},
 
     // BATCH GEMM
     {"cublasSgemmBatched",             {"hipblasSgemmBatched",             CONV_MATH_FUNC,       API_BLAS}},
@@ -2414,8 +2446,8 @@ const std::map<llvm::StringRef, hipCounter> CUDA_IDENTIFIER_MAP{
     {"cublasZhemm",                    {"hipblasZhemm",                    CONV_MATH_FUNC,       API_BLAS, HIP_UNSUPPORTED}},
 
     // TRSM
-    {"cublasStrsm",                    {"hipblasStrsm",                    CONV_MATH_FUNC,       API_BLAS, HIP_UNSUPPORTED}},
+    {"cublasStrsm",                    {"hipblasStrsm",                    CONV_MATH_FUNC,       API_BLAS}},
-    {"cublasDtrsm",                    {"hipblasDtrsm",                    CONV_MATH_FUNC,       API_BLAS, HIP_UNSUPPORTED}},
+    {"cublasDtrsm",                    {"hipblasDtrsm",                    CONV_MATH_FUNC,       API_BLAS}},
     {"cublasCtrsm",                    {"hipblasCtrsm",                    CONV_MATH_FUNC,       API_BLAS, HIP_UNSUPPORTED}},
     {"cublasZtrsm",                    {"hipblasZtrsm",                    CONV_MATH_FUNC,       API_BLAS, HIP_UNSUPPORTED}},
 
@@ -2627,6 +2659,8 @@ const std::map<llvm::StringRef, hipCounter> CUDA_IDENTIFIER_MAP{
     //IO in FP16 / FP32, computation in float
     {"cublasSgemmEx",                  {"hipblasSgemmEx",                  CONV_MATH_FUNC,       API_BLAS, HIP_UNSUPPORTED}},
     {"cublasGemmEx",                   {"hipblasGemmEx",                   CONV_MATH_FUNC,       API_BLAS, HIP_UNSUPPORTED}},
+    {"cublasGemmBatchedEx",            {"hipblasGemmBatchedEx",            CONV_MATH_FUNC,       API_BLAS, HIP_UNSUPPORTED}},
+    {"cublasGemmStridedBatchedEx",     {"hipblasGemmStridedBatchedEx",     CONV_MATH_FUNC,       API_BLAS, HIP_UNSUPPORTED}},
     // IO in Int8 complex/cuComplex, computation in cuComplex
     {"cublasCgemmEx",                  {"hipblasCgemmEx",                  CONV_MATH_FUNC,       API_BLAS, HIP_UNSUPPORTED}},