Merge pull request #1736 from aakanksha555/amd-staging

[HIPIFY][Tensor][feature] Introduce `cuTensor` support - Part 1 - w/o hipify-perl and documentation

CMakeLists.txt

@@ -278,6 +278,7 @@ if(HIPIFY_CLANG_TESTS OR HIPIFY_CLANG_TESTS_ONLY)
   set(CUDA_TOOLKIT_ROOT_DIR "" CACHE PATH "Path to CUDA Toolkit to use in hipify-clang unit testing")
   set(CUDA_SDK_ROOT_DIR     "" CACHE PATH "Path to CUDA Toolkit Samples")
   set(CUDA_DNN_ROOT_DIR     "" CACHE PATH "Path to cuDNN")
+  set(CUDA_TENSOR_ROOT_DIR  "" CACHE PATH "Path to cuTENSOR")
   set(CUDA_CUB_ROOT_DIR     "" CACHE PATH "Path to CUB")
 
   if(NOT CUDA_TOOLKIT_ROOT_DIR STREQUAL "" AND NOT CUDA_TOOLKIT_ROOT_DIR MATCHES "OFF")
@@ -294,6 +295,7 @@ if(HIPIFY_CLANG_TESTS OR HIPIFY_CLANG_TESTS_ONLY)
   message(STATUS "   - CUDA Toolkit path     : ${CUDA_TOOLKIT_ROOT_DIR}")
   message(STATUS "   - CUDA Samples path     : ${CUDA_SDK_ROOT_DIR}")
   message(STATUS "   - cuDNN path            : ${CUDA_DNN_ROOT_DIR}")
+  message(STATUS "   - cuTENSOR path         : ${CUDA_TENSOR_ROOT_DIR}")
   message(STATUS "   - CUB path              : ${CUDA_CUB_ROOT_DIR}")
 
   if(${CMAKE_VERSION} VERSION_LESS "3.27.0")
@@ -332,6 +334,10 @@ if(HIPIFY_CLANG_TESTS OR HIPIFY_CLANG_TESTS_ONLY)
     set(CUDA_DNN_ROOT_DIR OFF)
   endif()
 
+  if(CUDA_TENSOR_ROOT_DIR STREQUAL "")
+      set(CUDA_TENSOR_ROOT_DIR OFF)
+  endif()
+
   if(CUDA_CUB_ROOT_DIR STREQUAL "")
     if(CUDA_VERSION VERSION_GREATER_EQUAL "11.0")
       if(WIN32 OR (UNIX AND CUDA_VERSION VERSION_LESS "11.4" OR CUDA_VERSION VERSION_GREATER_EQUAL "11.6"))
@@ -348,6 +354,7 @@ if(HIPIFY_CLANG_TESTS OR HIPIFY_CLANG_TESTS_ONLY)
   message(STATUS "   - CUDA Toolkit path     : ${CUDA_TOOLKIT_ROOT_DIR}")
   message(STATUS "   - CUDA Samples path     : ${CUDA_SDK_ROOT_DIR}")
   message(STATUS "   - cuDNN path            : ${CUDA_DNN_ROOT_DIR}")
+  message(STATUS "   - cuTENSOR path         : ${CUDA_TENSOR_ROOT_DIR}")
   message(STATUS "   - CUB path              : ${CUDA_CUB_ROOT_DIR}")
 
   if((CUDA_VERSION VERSION_LESS "7.0") OR (LLVM_PACKAGE_VERSION VERSION_LESS "3.8") OR

src/CUDA2HIP.cpp

@@ -67,6 +67,8 @@ const std::map <llvm::StringRef, hipCounter> CUDA_INCLUDE_MAP {
   {"curand_uniform.h",                                      {"hiprand/hiprand_kernel.h",                              "rocrand/rocrand_uniform.h",                                      CONV_INCLUDE,                API_RAND, 0}},
   // cuDNN includes
   {"cudnn.h",                                               {"hipDNN.h",                               "miopen/miopen.h", CONV_INCLUDE_CUDA_MAIN_H,    API_DNN, 0}},
+  // cuTensor includes
+  {"cutensor.h",                                            {"hiptensor.h",                                           "", CONV_INCLUDE_CUDA_MAIN_H,    API_TENSOR, 0}},
   // cuFFT includes
   {"cufft.h",                                               {"hipfft/hipfft.h",                                       "", CONV_INCLUDE_CUDA_MAIN_H,    API_FFT, 0}},
   {"cufftXt.h",                                             {"hipfft/hipfftXt.h",                                     "", CONV_INCLUDE,                API_FFT, 0}},
@@ -127,6 +129,8 @@ const std::map<llvm::StringRef, hipCounter> &CUDA_RENAMES_MAP() {
   ret.insert(CUDA_DEVICE_TYPE_NAME_MAP.begin(), CUDA_DEVICE_TYPE_NAME_MAP.end());
   ret.insert(CUDA_SOLVER_TYPE_NAME_MAP.begin(), CUDA_SOLVER_TYPE_NAME_MAP.end());
   ret.insert(CUDA_SOLVER_FUNCTION_MAP.begin(), CUDA_SOLVER_FUNCTION_MAP.end());
+  ret.insert(CUDA_TENSOR_TYPE_NAME_MAP.begin(), CUDA_TENSOR_TYPE_NAME_MAP.end());
+  ret.insert(CUDA_TENSOR_FUNCTION_MAP.begin(), CUDA_TENSOR_FUNCTION_MAP.end());
   return ret;
 };
 

src/CUDA2HIP.h

@@ -85,6 +85,10 @@ extern const std::map<llvm::StringRef, hipCounter> CUDA_RTC_FUNCTION_MAP;
 extern const std::map<llvm::StringRef, hipCounter> CUDA_SOLVER_TYPE_NAME_MAP;
 // Maps the names of CUDA SOLVER API functions to the corresponding HIP functions
 extern const std::map<llvm::StringRef, hipCounter> CUDA_SOLVER_FUNCTION_MAP;
+// Maps the names of CUDA TENSOR API types to the corresponding HIP types
+extern const std::map<llvm::StringRef, hipCounter> CUDA_TENSOR_TYPE_NAME_MAP;
+// Maps the names of CUDA TENSOR API functions to the corresponding HIP functions
+extern const std::map<llvm::StringRef, hipCounter> CUDA_TENSOR_FUNCTION_MAP;
 
 /**
   * The union of all the above maps, except includes.

src/CUDA2HIP_TENSOR_API_functions.cpp

@@ -0,0 +1,26 @@
+/*
+Copyright (c) 2015 - present Advanced Micro Devices, Inc. All rights reserved.
+
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+
+#include "CUDA2HIP.h"
+
+const std::map<llvm::StringRef, hipCounter> CUDA_TENSOR_FUNCTION_MAP {
+};

src/CUDA2HIP_TENSOR_API_types.cpp

@@ -0,0 +1,53 @@
+/*
+Copyright (c) 2015 - present Advanced Micro Devices, Inc. All rights reserved.
+
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+
+#include "CUDA2HIP.h"
+
+// Map of all functions
+const std::map<llvm::StringRef, hipCounter> CUDA_TENSOR_TYPE_NAME_MAP {
+  // cuTENSOR defines
+
+
+  // cuTENSOR enums
+    {"cutensorDataType_t",                               {"",                                                         "",           CONV_TYPE, API_TENSOR, 1, HIP_UNSUPPORTED}},
+    {"cutensorOperator_t",                               {"",                                                         "",           CONV_TYPE, API_TENSOR, 1, HIP_UNSUPPORTED}},
+
+    {"cutensorStatus_t",                                 {"hiptensorStatus_t",                                        "",           CONV_TYPE, API_TENSOR, 1}},
+    {"CUTENSOR_STATUS_SUCCESS",                          {"HIPTENSOR_STATUS_SUCCESS",                                 "",           CONV_NUMERIC_LITERAL, API_TENSOR, 1}},
+    {"CUTENSOR_STATUS_NOT_INITIALIZED",                  {"HIPTENSOR_STATUS_NOT_INITIALIZED",                         "",           CONV_NUMERIC_LITERAL, API_TENSOR, 1}},
+    {"CUTENSOR_STATUS_ALLOC_FAILED",                     {"HIPTENSOR_STATUS_ALLOC_FAILED",                            "",           CONV_NUMERIC_LITERAL, API_TENSOR, 1}},
+    {"CUTENSOR_STATUS_INVALID_VALUE",                    {"HIPTENSOR_STATUS_INVALID_VALUE",                           "",           CONV_NUMERIC_LITERAL, API_TENSOR, 1}},
+    {"CUTENSOR_STATUS_ARCH_MISMATCH",                    {"HIPTENSOR_STATUS_ARCH_MISMATCH",                           "",           CONV_NUMERIC_LITERAL, API_TENSOR, 1}},
+    {"CUTENSOR_STATUS_MAPPING_ERROR",                    {"",                                                         "",           CONV_NUMERIC_LITERAL, API_TENSOR, 1}},
+    {"CUTENSOR_STATUS_EXECUTION_FAILED",                 {"HIPTENSOR_STATUS_EXECUTION_FAILED",                        "",           CONV_NUMERIC_LITERAL, API_TENSOR, 1}},
+    {"CUTENSOR_STATUS_INTERNAL_ERROR",                   {"HIPTENSOR_STATUS_INTERNAL_ERROR",                          "",           CONV_NUMERIC_LITERAL, API_TENSOR, 1}},
+    {"CUTENSOR_STATUS_NOT_SUPPORTED",                    {"HIPTENSOR_STATUS_NOT_SUPPORTED",                           "",           CONV_NUMERIC_LITERAL, API_TENSOR, 1}},
+    {"CUTENSOR_STATUS_LICENSE_ERROR",                    {"",                                                         "",           CONV_NUMERIC_LITERAL, API_TENSOR, 1}},
+    {"CUTENSOR_STATUS_CUBLAS_ERROR",                     {"",                                                         "",           CONV_NUMERIC_LITERAL, API_TENSOR, 1}},
+    {"CUTENSOR_STATUS_CUDA_ERROR",                       {"",                                                         "",           CONV_NUMERIC_LITERAL, API_TENSOR, 1}},
+    {"CUTENSOR_STATUS_INSUFFICIENT_WORKSPACE",           {"HIPTENSOR_STATUS_INSUFFICIENT_WORKSPACE",                  "",           CONV_NUMERIC_LITERAL, API_TENSOR, 1}},
+    {"CUTENSOR_STATUS_INSUFFICIENT_DRIVER",              {"HIPTENSOR_STATUS_INSUFFICIENT_DRIVER",                     "",           CONV_NUMERIC_LITERAL, API_TENSOR, 1}},
+    {"CUTENSOR_STATUS_IO_ERROR",                         {"HIPTENSOR_STATUS_IO_ERROR",                                "",           CONV_NUMERIC_LITERAL, API_TENSOR, 1}},
+};
+
+const std::map<llvm::StringRef, cudaAPIversions> CUDA_TENSOR_TYPE_NAME_VER_MAP {
+};

src/Statistics.cpp

@@ -145,7 +145,8 @@ const char *apiNames[NUM_API_TYPES] = {
   "cuSOLVER API",
   "CUB API",
   "CAFFE2 API",
-  "RTC API"
+  "RTC API",
+  "TENSOR API"
 };
 
 const char *apiTypes[NUM_API_TYPES] = {
@@ -160,7 +161,8 @@ const char *apiTypes[NUM_API_TYPES] = {
   "API_SPARSE",
   "API_SOLVER",
   "API_CAFFE2",
-  "API_RTC"
+  "API_RTC",
+  "API_TENSOR"
 };
 
 namespace {

src/Statistics.h

@@ -178,6 +178,7 @@ enum ApiTypes {
   API_CUB,
   API_CAFFE2,
   API_RTC,
+  API_TENSOR,
   API_LAST
 };
 constexpr int NUM_API_TYPES = (int) ApiTypes::API_LAST;

tests/lit.cfg

@@ -248,15 +248,37 @@ if config.cuda_dnn_root and config.cuda_dnn_root != "OFF":
 if config.cuda_cub_root and config.cuda_cub_root != "OFF":
     clang_arguments += " -I'%s'"
 
+# cuTensor ROOT
+if config.cuda_tensor_root and config.cuda_tensor_root != "OFF":
+    clang_arguments += " -I'%s'/include"
+
 i_subst = 0
 if config.cuda_sdk_root != "OFF":
     i_subst = 1
 if config.cuda_dnn_root != "OFF":
     i_subst += 2
 if config.cuda_cub_root != "OFF":
     i_subst += 4
-
-if i_subst == 7:
+if config.cuda_tensor_root != "OFF":
+    i_subst += 8
+
+if i_subst == 15:
+    config.substitutions.append(("%clang_args", clang_arguments % (config.cuda_sdk_root, config.cuda_dnn_root, config.cuda_cub_root, config.cuda_tensor_root)))
+elif i_subst == 14:
+    config.substitutions.append(("%clang_args", clang_arguments % (config.cuda_dnn_root, config.cuda_cub_root, config.cuda_tensor_root)))
+elif i_subst == 13:
+    config.substitutions.append(("%clang_args", clang_arguments % (config.cuda_sdk_root, config.cuda_cub_root, config.cuda_tensor_root)))
+elif i_subst == 12:
+    config.substitutions.append(("%clang_args", clang_arguments % (config.cuda_cub_root, config.cuda_tensor_root)))
+elif i_subst == 11:
+    config.substitutions.append(("%clang_args", clang_arguments % (config.cuda_sdk_root, config.cuda_dnn_root, config.cuda_tensor_root)))
+elif i_subst == 10:
+    config.substitutions.append(("%clang_args", clang_arguments % (config.cuda_dnn_root, config.cuda_tensor_root)))
+elif i_subst == 9:
+    config.substitutions.append(("%clang_args", clang_arguments % (config.cuda_sdk_root, config.cuda_tensor_root)))
+elif i_subst == 8:
+    config.substitutions.append(("%clang_args", clang_arguments % config.cuda_tensor_root))
+elif i_subst == 7:
     config.substitutions.append(("%clang_args", clang_arguments % (config.cuda_sdk_root, config.cuda_dnn_root, config.cuda_cub_root)))
 elif i_subst == 6:
     config.substitutions.append(("%clang_args", clang_arguments % (config.cuda_dnn_root, config.cuda_cub_root)))

tests/lit.site.cfg.in

@@ -12,6 +12,7 @@ config.cuda_root = "@CUDA_TOOLKIT_ROOT_DIR@"
 config.cuda_sdk_root = "@CUDA_SDK_ROOT_DIR@"
 config.cuda_dnn_root = "@CUDA_DNN_ROOT_DIR@"
 config.cuda_cub_root = "@CUDA_CUB_ROOT_DIR@"
+config.cuda_tensor_root = "@CUDA_TENSOR_ROOT_DIR@"
 config.cuda_version_major = int("@CUDA_VERSION_MAJOR@")
 config.cuda_version_minor = int("@CUDA_VERSION_MINOR@")
 config.cuda_version_full = "@CUDA_VERSION_FULL@"

tests/unit_tests/synthetic/libraries/cutensor2hiptensor.cu

@@ -0,0 +1,37 @@
+// RUN: %run_test hipify "%s" "%t" %hipify_args 3 --amap --skip-excluded-preprocessor-conditional-blocks --experimental %clang_args -D__CUDA_API_VERSION_INTERNAL -ferror-limit=500
+
+// CHECK: #include <hip/hip_runtime.h>
+#include <cuda_runtime.h>
+#include <stdio.h>
+// CHECK: #include "hiptensor.h"
+#include "cutensor.h"
+// CHECK-NOT: #include "hiptensor.h"
+
+int main() {
+
+    // CHECK: hiptensorStatus_t tensorStatus_t;
+    // CHECK-NEXT hiptensorStatus_t tensorStatus;
+    // CHECK-NEXT hiptensorStatus_t TENSOR_STATUS_SUCCESS = HIPTENSOR_STATUS_SUCCESS;
+    // CHECK-NEXT hiptensorStatus_t TENSOR_STATUS_NOT_INITIALIZED = HIPTENSOR_STATUS_NOT_INITIALIZED;
+    // CHECK-NEXT hiptensorStatus_t TENSOR_STATUS_ALLOC_FAILED = HIPTENSOR_STATUS_ALLOC_FAILED;
+    // CHECK-NEXT hiptensorStatus_t TENSOR_STATUS_INVALID_VALUE = HIPTENSOR_STATUS_INVALID_VALUE;
+    // CHECK-NEXT hiptensorStatus_t TENSOR_STATUS_ARCH_MISMATCH = HIPTENSOR_STATUS_ARCH_MISMATCH;
+    // CHECK-NEXT hiptensorStatus_t TENSOR_STATUS_EXECUTION_FAILED = HIPTENSOR_STATUS_EXECUTION_FAILED;
+    // CHECK-NEXT hiptensorStatus_t TENSOR_STATUS_INTERNAL_ERROR = HIPTENSOR_STATUS_INTERNAL_ERROR;
+    // CHECK-NEXT hiptensorStatus_t TENSOR_STATUS_NOT_SUPPORTED = HIPTENSOR_STATUS_NOT_SUPPORTED;
+    // CHECK-NEXT hiptensorStatus_t TENSOR_STATUS_INSUFFICIENT_WORKSPACE = HIPTENSOR_STATUS_INSUFFICIENT_WORKSPACE;
+    // CHECK-NEXT hiptensorStatus_t TENSOR_STATUS_INSUFFICIENT_DRIVER = HIPTENSOR_STATUS_INSUFFICIENT_DRIVER;
+    // CHECK-NEXT hiptensorStatus_t TENSOR_STATUS_IO_ERROR = HIPTENSOR_STATUS_IO_ERROR;
+    cutensorStatus_t tensorStatus_t;
+    cutensorStatus_t TENSOR_STATUS_SUCCESS = CUTENSOR_STATUS_SUCCESS;
+    cutensorStatus_t TENSOR_STATUS_NOT_INITIALIZED = CUTENSOR_STATUS_NOT_INITIALIZED;
+    cutensorStatus_t TENSOR_STATUS_ALLOC_FAILED = CUTENSOR_STATUS_ALLOC_FAILED;
+    cutensorStatus_t TENSOR_STATUS_INVALID_VALUE = CUTENSOR_STATUS_INVALID_VALUE;
+    cutensorStatus_t TENSOR_STATUS_ARCH_MISMATCH = CUTENSOR_STATUS_ARCH_MISMATCH;
+    cutensorStatus_t TENSOR_STATUS_EXECUTION_FAILED = CUTENSOR_STATUS_EXECUTION_FAILED;
+    cutensorStatus_t TENSOR_STATUS_INTERNAL_ERROR = CUTENSOR_STATUS_INTERNAL_ERROR;
+    cutensorStatus_t TENSOR_STATUS_NOT_SUPPORTED = CUTENSOR_STATUS_NOT_SUPPORTED;
+    cutensorStatus_t TENSOR_STATUS_INSUFFICIENT_WORKSPACE = CUTENSOR_STATUS_INSUFFICIENT_WORKSPACE;
+    cutensorStatus_t TENSOR_STATUS_INSUFFICIENT_DRIVER = CUTENSOR_STATUS_INSUFFICIENT_DRIVER;
+    cutensorStatus_t TENSOR_STATUS_IO_ERROR = CUTENSOR_STATUS_IO_ERROR;
+}