From 22a520dbd4d3975e8730efefe8bad14cb699db99 Mon Sep 17 00:00:00 2001
From: "Jiang, Zhiwei" <zhiwei.jiang@intel.com>
Date: Fri, 1 Mar 2024 15:53:02 +0800
Subject: [PATCH] Update

Signed-off-by: Jiang, Zhiwei <zhiwei.jiang@intel.com>
---
 clang/lib/DPCT/APINamesCUBLAS.inc             | 16 +--
 .../dpct-rt/include/dpct/blas_utils.hpp       | 98 +++++++++++--------
 clang/test/dpct/cublas-usm.cu                 | 16 +--
 clang/test/dpct/cublasIsamax_etc.cu           | 16 +--
 clang/test/dpct/cublasRegularCZ.cu            | 16 +--
 clang/test/dpct/cublas_curandInMacro.cu       |  4 +-
 clang/test/dpct/macro_test.cu                 |  2 +-
 clang/test/dpct/nestedqueue.cu                |  2 +-
 .../query_api_mapping/cuBLAS/blas_part2.cu    |  2 +-
 .../query_api_mapping/cuBLAS/blas_part3.cu    |  4 +-
 .../query_api_mapping/cuBLAS/blas_part4.cu    |  2 +-
 .../query_api_mapping/cuBLAS/blas_part5.cu    |  2 +-
 .../query_api_mapping/cuBLAS/blas_part6.cu    |  6 +-
 clang/test/dpct/test_libaray_index/test3.cu   |  2 +-
 14 files changed, 100 insertions(+), 88 deletions(-)

diff --git a/clang/lib/DPCT/APINamesCUBLAS.inc b/clang/lib/DPCT/APINamesCUBLAS.inc
index 12d84aecdcdd..84daaec269c9 100644
--- a/clang/lib/DPCT/APINamesCUBLAS.inc
+++ b/clang/lib/DPCT/APINamesCUBLAS.inc
@@ -1529,7 +1529,7 @@ WARNING_FACTORY_ENTRY(
     "cublasIsamax_v2",
     LAMBDA_FACTORY_ENTRY(
         "cublasIsamax_v2",
-        DECL(MapNames::getDpctNamespace() + "blas::out_mem_int_t", "res",
+        DECL(MapNames::getDpctNamespace() + "blas::out_mem_int64_int_t", "res",
              MEMBER_CALL(ARG(0), true, "get_queue"), ARG(4)),
         CALL("oneapi::mkl::blas::column_major::iamax",
              MEMBER_CALL(ARG(0), true, "get_queue"), ARG(1),
@@ -1542,7 +1542,7 @@ WARNING_FACTORY_ENTRY(
     "cublasIdamax_v2",
     LAMBDA_FACTORY_ENTRY(
         "cublasIdamax_v2",
-        DECL(MapNames::getDpctNamespace() + "blas::out_mem_int_t", "res",
+        DECL(MapNames::getDpctNamespace() + "blas::out_mem_int64_int_t", "res",
              MEMBER_CALL(ARG(0), true, "get_queue"), ARG(4)),
         CALL("oneapi::mkl::blas::column_major::iamax",
              MEMBER_CALL(ARG(0), true, "get_queue"), ARG(1),
@@ -1555,7 +1555,7 @@ WARNING_FACTORY_ENTRY(
     "cublasIcamax_v2",
     LAMBDA_FACTORY_ENTRY(
         "cublasIcamax_v2",
-        DECL(MapNames::getDpctNamespace() + "blas::out_mem_int_t", "res",
+        DECL(MapNames::getDpctNamespace() + "blas::out_mem_int64_int_t", "res",
              MEMBER_CALL(ARG(0), true, "get_queue"), ARG(4)),
         CALL("oneapi::mkl::blas::column_major::iamax",
              MEMBER_CALL(ARG(0), true, "get_queue"), ARG(1),
@@ -1568,7 +1568,7 @@ WARNING_FACTORY_ENTRY(
     "cublasIzamax_v2",
     LAMBDA_FACTORY_ENTRY(
         "cublasIzamax_v2",
-        DECL(MapNames::getDpctNamespace() + "blas::out_mem_int_t", "res",
+        DECL(MapNames::getDpctNamespace() + "blas::out_mem_int64_int_t", "res",
              MEMBER_CALL(ARG(0), true, "get_queue"), ARG(4)),
         CALL("oneapi::mkl::blas::column_major::iamax",
              MEMBER_CALL(ARG(0), true, "get_queue"), ARG(1),
@@ -1582,7 +1582,7 @@ WARNING_FACTORY_ENTRY(
     "cublasIsamin_v2",
     LAMBDA_FACTORY_ENTRY(
         "cublasIsamin_v2",
-        DECL(MapNames::getDpctNamespace() + "blas::out_mem_int_t", "res",
+        DECL(MapNames::getDpctNamespace() + "blas::out_mem_int64_int_t", "res",
              MEMBER_CALL(ARG(0), true, "get_queue"), ARG(4)),
         CALL("oneapi::mkl::blas::column_major::iamin",
              MEMBER_CALL(ARG(0), true, "get_queue"), ARG(1),
@@ -1595,7 +1595,7 @@ WARNING_FACTORY_ENTRY(
     "cublasIdamin_v2",
     LAMBDA_FACTORY_ENTRY(
         "cublasIdamin_v2",
-        DECL(MapNames::getDpctNamespace() + "blas::out_mem_int_t", "res",
+        DECL(MapNames::getDpctNamespace() + "blas::out_mem_int64_int_t", "res",
              MEMBER_CALL(ARG(0), true, "get_queue"), ARG(4)),
         CALL("oneapi::mkl::blas::column_major::iamin",
              MEMBER_CALL(ARG(0), true, "get_queue"), ARG(1),
@@ -1608,7 +1608,7 @@ WARNING_FACTORY_ENTRY(
     "cublasIcamin_v2",
     LAMBDA_FACTORY_ENTRY(
         "cublasIcamin_v2",
-        DECL(MapNames::getDpctNamespace() + "blas::out_mem_int_t", "res",
+        DECL(MapNames::getDpctNamespace() + "blas::out_mem_int64_int_t", "res",
              MEMBER_CALL(ARG(0), true, "get_queue"), ARG(4)),
         CALL("oneapi::mkl::blas::column_major::iamin",
              MEMBER_CALL(ARG(0), true, "get_queue"), ARG(1),
@@ -1621,7 +1621,7 @@ WARNING_FACTORY_ENTRY(
     "cublasIzamin_v2",
     LAMBDA_FACTORY_ENTRY(
         "cublasIzamin_v2",
-        DECL(MapNames::getDpctNamespace() + "blas::out_mem_int_t", "res",
+        DECL(MapNames::getDpctNamespace() + "blas::out_mem_int64_int_t", "res",
              MEMBER_CALL(ARG(0), true, "get_queue"), ARG(4)),
         CALL("oneapi::mkl::blas::column_major::iamin",
              MEMBER_CALL(ARG(0), true, "get_queue"), ARG(1),
diff --git a/clang/runtime/dpct-rt/include/dpct/blas_utils.hpp b/clang/runtime/dpct-rt/include/dpct/blas_utils.hpp
index ba690f1db084..d395eac272b9 100644
--- a/clang/runtime/dpct-rt/include/dpct/blas_utils.hpp
+++ b/clang/runtime/dpct-rt/include/dpct/blas_utils.hpp
@@ -20,17 +20,16 @@
 
 namespace dpct {
 namespace blas {
-namespace detail {
-template <typename target_t, typename source_t, bool has_input>
-class scalar_memory_base_t {
+enum class mem_inout { in, out, inout };
+template <typename target_t, typename source_t, mem_inout io> class mem_base_t {
 public:
-  scalar_memory_base_t(sycl::queue q, source_t *source,
+  mem_base_t(sycl::queue q, source_t *source,
 #ifdef DPCT_USM_LEVEL_NONE
-                       sycl::buffer<target_t> target
+             sycl::buffer<target_t> target
 #else
-                       target_t *target
+             target_t *target
 #endif
-                       )
+             )
       : _q(q), _source(source), _target(target),
         _source_attribute(dpct::detail::get_pointer_attribute(_q, _source)) {
   }
@@ -52,20 +51,18 @@ class scalar_memory_base_t {
   dpct::detail::pointer_access_attribute _source_attribute;
 };
 
-template <typename target_t, typename source_t, bool has_input>
-class scalar_memory_t
-    : public scalar_memory_base_t<target_t, source_t, has_input> {
-  static_assert(
-      !has_input &&
-      "input is not supported if target_t and source_t are not same.");
-  using base_t = scalar_memory_base_t<target_t, source_t, has_input>;
+template <typename target_t, typename source_t, mem_inout io>
+class mem_t : public mem_base_t<target_t, source_t, io> {
+  static_assert(io == mem_inout::out && "Only mem_inout::out is supported if "
+                                        "target_t and source_t are not same.");
+  using base_t = mem_base_t<target_t, source_t, io>;
   using base_t::_q;
   using base_t::_source;
   using base_t::_source_attribute;
   using base_t::_target;
 
 public:
-  scalar_memory_t(sycl::queue q, source_t *source)
+  mem_t(sycl::queue q, source_t *source)
       : base_t(q, source,
 #ifdef DPCT_USM_LEVEL_NONE
                sycl::buffer<target_t>(sycl::range<1>(1))
@@ -75,7 +72,7 @@ class scalar_memory_t
         ) {
   }
 
-  ~scalar_memory_t() {
+  ~mem_t() {
 #ifdef DPCT_USM_LEVEL_NONE
     source_t temp = static_cast<source_t>(_target.get_host_access()[0]);
     if (_source_attribute ==
@@ -98,36 +95,48 @@ class scalar_memory_t
   }
 };
 
-template <typename target_t, bool has_input>
-class scalar_memory_t<target_t, target_t, has_input>
-    : public scalar_memory_base_t<target_t, target_t, has_input> {
-  using base_t = scalar_memory_base_t<target_t, target_t, has_input>;
+template <typename target_t, mem_inout io>
+class mem_t<target_t, target_t, io>
+    : public mem_base_t<target_t, target_t, io> {
+  using base_t = mem_base_t<target_t, target_t, io>;
   using base_t::_q;
   using base_t::_source;
   using base_t::_source_attribute;
   using base_t::_target;
+  size_t _ele_num;
 #ifndef DPCT_USM_LEVEL_NONE
   bool _need_free = true;
 #endif
 
-public:
-  scalar_memory_t(sycl::queue q, target_t *source)
-      : base_t(q, source,
-               dpct::detail::get_pointer_attribute(q, source) !=
-                       dpct::detail::pointer_access_attribute::host_only
-                   ?
 #ifdef DPCT_USM_LEVEL_NONE
-                   dpct::get_buffer<target_t>(source)
+  sycl::buffer<target_t>
 #else
-                   source
+  target_t *
 #endif
-                   :
+  construct_member_variable_target(sycl::queue q, target_t *source,
+                                   size_t ele_num) {
 #ifdef DPCT_USM_LEVEL_NONE
-                   sycl::buffer<target_t>(source, sycl::range<1>(1))
+    if (dpct::detail::get_pointer_attribute(q, source) !=
+        dpct::detail::pointer_access_attribute::host_only) {
+      target_t *host_ptr = dpct::get_host_ptr<target_t>(source);
+      return sycl::buffer<target_t>(host_ptr, sycl::range<1>(ele_num));
+    } else {
+      return sycl::buffer<target_t>(source, sycl::range<1>(ele_num));
+    }
 #else
-                   sycl::malloc_shared<target_t>(1, q)
+    if (dpct::detail::get_pointer_attribute(q, source) !=
+        dpct::detail::pointer_access_attribute::host_only) {
+      return source;
+    } else {
+      return sycl::malloc_shared<target_t>(ele_num, q);
+    }
 #endif
-        ) {
+  }
+
+public:
+  mem_t(sycl::queue q, target_t *source, size_t ele_num = 1)
+      : base_t(q, source, construct_member_variable_target(q, source, ele_num)),
+        _ele_num(ele_num) {
     if (_source_attribute !=
         dpct::detail::pointer_access_attribute::host_only) {
 #ifndef DPCT_USM_LEVEL_NONE
@@ -137,33 +146,36 @@ class scalar_memory_t<target_t, target_t, has_input>
     }
 
 #ifndef DPCT_USM_LEVEL_NONE
-    if constexpr (has_input) {
+    if constexpr (io != mem_inout::out) {
       if (_source_attribute ==
           dpct::detail::pointer_access_attribute::host_only) {
-        *_target = *_source;
+        _q.memcpy(_target, _source, sizeof(target_t) * _ele_num).wait();
       }
     }
 #endif
   }
 
-  ~scalar_memory_t() {
+  ~mem_t() {
 #ifndef DPCT_USM_LEVEL_NONE
     if (!_need_free) {
       return;
     }
-    if (_source_attribute ==
-        dpct::detail::pointer_access_attribute::host_only) {
-      _q.wait();
-      *_source = *_target;
+    if constexpr (io != mem_inout::in) {
+      if (_source_attribute ==
+          dpct::detail::pointer_access_attribute::host_only) {
+        _q.wait();
+        _q.memcpy(_source, _target, sizeof(target_t) * _ele_num).wait();
+      }
     }
     sycl::free(_target, _q);
 #endif
   }
 };
-} // namespace detail
-using out_mem_int_t = detail::scalar_memory_t<std::int64_t, int, false>;
-using out_mem_int64_t =
-    detail::scalar_memory_t<std::int64_t, std::int64_t, false>;
+using out_mem_int64_int_t = mem_t<std::int64_t, int, mem_inout::out>;
+using out_mem_int64_t = mem_t<std::int64_t, std::int64_t, mem_inout::out>;
+using out_mem_float_t = mem_t<float, float, mem_inout::out>;
+using inout_mem_float_t = mem_t<float, float, mem_inout::inout>;
+using in_mem_float_t = mem_t<float, float, mem_inout::in>;
 
 class descriptor {
 public:
diff --git a/clang/test/dpct/cublas-usm.cu b/clang/test/dpct/cublas-usm.cu
index 2fc3b17143e1..2a16ce07007e 100644
--- a/clang/test/dpct/cublas-usm.cu
+++ b/clang/test/dpct/cublas-usm.cu
@@ -83,50 +83,50 @@ int main() {
   //level 1
 
   //CHECK:a = [&]() {
-  //CHECK-NEXT:dpct::blas::out_mem_int_t res(handle->get_queue(), result);
+  //CHECK-NEXT:dpct::blas::out_mem_int64_int_t res(handle->get_queue(), result);
   //CHECK-NEXT:oneapi::mkl::blas::column_major::iamax(handle->get_queue(), N, x_S, N, res.get_memory(), oneapi::mkl::index_base::one);
   //CHECK-NEXT:return 0;
   //CHECK-NEXT:}();
   a = cublasIsamax(handle, N, x_S, N, result);
   //CHECK:[&]() {
-  //CHECK-NEXT:dpct::blas::out_mem_int_t res(handle->get_queue(), result);
+  //CHECK-NEXT:dpct::blas::out_mem_int64_int_t res(handle->get_queue(), result);
   //CHECK-NEXT:oneapi::mkl::blas::column_major::iamax(handle->get_queue(), N, x_D, N, res.get_memory(), oneapi::mkl::index_base::one);
   //CHECK-NEXT:return 0;
   //CHECK-NEXT:}();
   cublasIdamax(handle, N, x_D, N, result);
   //CHECK:a = [&]() {
-  //CHECK-NEXT:dpct::blas::out_mem_int_t res(handle->get_queue(), result);
+  //CHECK-NEXT:dpct::blas::out_mem_int64_int_t res(handle->get_queue(), result);
   //CHECK-NEXT:oneapi::mkl::blas::column_major::iamax(handle->get_queue(), N, (std::complex<float>*)x_C, N, res.get_memory(), oneapi::mkl::index_base::one);
   //CHECK-NEXT:return 0;
   //CHECK-NEXT:}();
   a = cublasIcamax(handle, N, x_C, N, result);
   //CHECK:[&]() {
-  //CHECK-NEXT:dpct::blas::out_mem_int_t res(handle->get_queue(), result);
+  //CHECK-NEXT:dpct::blas::out_mem_int64_int_t res(handle->get_queue(), result);
   //CHECK-NEXT:oneapi::mkl::blas::column_major::iamax(handle->get_queue(), N, (std::complex<double>*)x_Z, N, res.get_memory(), oneapi::mkl::index_base::one);
   //CHECK-NEXT:return 0;
   //CHECK-NEXT:}();
   cublasIzamax(handle, N, x_Z, N, result);
 
   //CHECK:a = [&]() {
-  //CHECK-NEXT:dpct::blas::out_mem_int_t res(handle->get_queue(), result);
+  //CHECK-NEXT:dpct::blas::out_mem_int64_int_t res(handle->get_queue(), result);
   //CHECK-NEXT:oneapi::mkl::blas::column_major::iamin(handle->get_queue(), N, x_S, N, res.get_memory(), oneapi::mkl::index_base::one);
   //CHECK-NEXT:return 0;
   //CHECK-NEXT:}();
   a = cublasIsamin(handle, N, x_S, N, result);
   //CHECK:[&]() {
-  //CHECK-NEXT:dpct::blas::out_mem_int_t res(handle->get_queue(), result);
+  //CHECK-NEXT:dpct::blas::out_mem_int64_int_t res(handle->get_queue(), result);
   //CHECK-NEXT:oneapi::mkl::blas::column_major::iamin(handle->get_queue(), N, x_D, N, res.get_memory(), oneapi::mkl::index_base::one);
   //CHECK-NEXT:return 0;
   //CHECK-NEXT:}();
   cublasIdamin(handle, N, x_D, N, result);
   //CHECK:a = [&]() {
-  //CHECK-NEXT:dpct::blas::out_mem_int_t res(handle->get_queue(), result);
+  //CHECK-NEXT:dpct::blas::out_mem_int64_int_t res(handle->get_queue(), result);
   //CHECK-NEXT:oneapi::mkl::blas::column_major::iamin(handle->get_queue(), N, (std::complex<float>*)x_C, N, res.get_memory(), oneapi::mkl::index_base::one);
   //CHECK-NEXT:return 0;
   //CHECK-NEXT:}();
   a = cublasIcamin(handle, N, x_C, N, result);
   //CHECK:[&]() {
-  //CHECK-NEXT:dpct::blas::out_mem_int_t res(handle->get_queue(), result);
+  //CHECK-NEXT:dpct::blas::out_mem_int64_int_t res(handle->get_queue(), result);
   //CHECK-NEXT:oneapi::mkl::blas::column_major::iamin(handle->get_queue(), N, (std::complex<double>*)x_Z, N, res.get_memory(), oneapi::mkl::index_base::one);
   //CHECK-NEXT:return 0;
   //CHECK-NEXT:}();
diff --git a/clang/test/dpct/cublasIsamax_etc.cu b/clang/test/dpct/cublasIsamax_etc.cu
index e60bd4eab976..8768115c40a7 100644
--- a/clang/test/dpct/cublasIsamax_etc.cu
+++ b/clang/test/dpct/cublasIsamax_etc.cu
@@ -38,12 +38,12 @@ int main() {
   //level1
   //cublasI<t>amax
   // CHECK: status = [&]() {
-  // CHECK-NEXT: dpct::blas::out_mem_int_t res(handle->get_queue(), result);
+  // CHECK-NEXT: dpct::blas::out_mem_int64_int_t res(handle->get_queue(), result);
   // CHECK-NEXT: oneapi::mkl::blas::column_major::iamax(handle->get_queue(), n, dpct::rvalue_ref_to_lvalue_ref(dpct::get_buffer<float>(x_S)), incx, res.get_memory(), oneapi::mkl::index_base::one);
   // CHECK-NEXT: return 0;
   // CHECK-NEXT: }();
   // CHECK: [&]() {
-  // CHECK-NEXT: dpct::blas::out_mem_int_t res(handle->get_queue(), result);
+  // CHECK-NEXT: dpct::blas::out_mem_int64_int_t res(handle->get_queue(), result);
   // CHECK-NEXT: oneapi::mkl::blas::column_major::iamax(handle->get_queue(), n, dpct::rvalue_ref_to_lvalue_ref(dpct::get_buffer<float>(x_S)), incx, res.get_memory(), oneapi::mkl::index_base::one);
   // CHECK-NEXT: return 0;
   // CHECK-NEXT: }();
@@ -51,12 +51,12 @@ int main() {
   cublasIsamax(handle, n, x_S, incx, result);
 
   // CHECK: status = [&]() {
-  // CHECK-NEXT: dpct::blas::out_mem_int_t res(handle->get_queue(), result);
+  // CHECK-NEXT: dpct::blas::out_mem_int64_int_t res(handle->get_queue(), result);
   // CHECK-NEXT: oneapi::mkl::blas::column_major::iamax(handle->get_queue(), n, dpct::rvalue_ref_to_lvalue_ref(dpct::get_buffer<double>(x_D)), incx, res.get_memory(), oneapi::mkl::index_base::one);
   // CHECK-NEXT: return 0;
   // CHECK-NEXT: }();
   // CHECK: [&]() {
-  // CHECK-NEXT: dpct::blas::out_mem_int_t res(handle->get_queue(), result);
+  // CHECK-NEXT: dpct::blas::out_mem_int64_int_t res(handle->get_queue(), result);
   // CHECK-NEXT: oneapi::mkl::blas::column_major::iamax(handle->get_queue(), n, dpct::rvalue_ref_to_lvalue_ref(dpct::get_buffer<double>(x_D)), incx, res.get_memory(), oneapi::mkl::index_base::one);
   // CHECK-NEXT: return 0;
   // CHECK-NEXT: }();
@@ -65,12 +65,12 @@ int main() {
 
   //cublasI<t>amin
   // CHECK: status = [&]() {
-  // CHECK-NEXT: dpct::blas::out_mem_int_t res(handle->get_queue(), result);
+  // CHECK-NEXT: dpct::blas::out_mem_int64_int_t res(handle->get_queue(), result);
   // CHECK-NEXT: oneapi::mkl::blas::column_major::iamin(handle->get_queue(), n, dpct::rvalue_ref_to_lvalue_ref(dpct::get_buffer<float>(x_S)), incx, res.get_memory(), oneapi::mkl::index_base::one);
   // CHECK-NEXT: return 0;
   // CHECK-NEXT: }();
   // CHECK: [&]() {
-  // CHECK-NEXT: dpct::blas::out_mem_int_t res(handle->get_queue(), result);
+  // CHECK-NEXT: dpct::blas::out_mem_int64_int_t res(handle->get_queue(), result);
   // CHECK-NEXT: oneapi::mkl::blas::column_major::iamin(handle->get_queue(), n, dpct::rvalue_ref_to_lvalue_ref(dpct::get_buffer<float>(x_S)), incx, res.get_memory(), oneapi::mkl::index_base::one);
   // CHECK-NEXT: return 0;
   // CHECK-NEXT: }();
@@ -78,12 +78,12 @@ int main() {
   cublasIsamin(handle, n, x_S, incx, result);
 
   // CHECK: status = [&]() {
-  // CHECK-NEXT: dpct::blas::out_mem_int_t res(handle->get_queue(), result);
+  // CHECK-NEXT: dpct::blas::out_mem_int64_int_t res(handle->get_queue(), result);
   // CHECK-NEXT: oneapi::mkl::blas::column_major::iamin(handle->get_queue(), n, dpct::rvalue_ref_to_lvalue_ref(dpct::get_buffer<double>(x_D)), incx, res.get_memory(), oneapi::mkl::index_base::one);
   // CHECK-NEXT: return 0;
   // CHECK-NEXT: }();
   // CHECK: [&]() {
-  // CHECK-NEXT: dpct::blas::out_mem_int_t res(handle->get_queue(), result);
+  // CHECK-NEXT: dpct::blas::out_mem_int64_int_t res(handle->get_queue(), result);
   // CHECK-NEXT: oneapi::mkl::blas::column_major::iamin(handle->get_queue(), n, dpct::rvalue_ref_to_lvalue_ref(dpct::get_buffer<double>(x_D)), incx, res.get_memory(), oneapi::mkl::index_base::one);
   // CHECK-NEXT: return 0;
   // CHECK-NEXT: }();
diff --git a/clang/test/dpct/cublasRegularCZ.cu b/clang/test/dpct/cublasRegularCZ.cu
index 9d25d164e567..382b2177fa51 100644
--- a/clang/test/dpct/cublasRegularCZ.cu
+++ b/clang/test/dpct/cublasRegularCZ.cu
@@ -22,12 +22,12 @@ int main(){
 
   //level 1
   // CHECK: status = [&]() {
-  // CHECK-NEXT: dpct::blas::out_mem_int_t res(handle->get_queue(), result);
+  // CHECK-NEXT: dpct::blas::out_mem_int64_int_t res(handle->get_queue(), result);
   // CHECK-NEXT: oneapi::mkl::blas::column_major::iamax(handle->get_queue(), n, dpct::rvalue_ref_to_lvalue_ref(dpct::get_buffer<std::complex<float>>(x_c)), incx, res.get_memory(), oneapi::mkl::index_base::one);
   // CHECK-NEXT: return 0;
   // CHECK-NEXT: }();
   // CHECK: [&]() {
-  // CHECK-NEXT: dpct::blas::out_mem_int_t res(handle->get_queue(), result);
+  // CHECK-NEXT: dpct::blas::out_mem_int64_int_t res(handle->get_queue(), result);
   // CHECK-NEXT: oneapi::mkl::blas::column_major::iamax(handle->get_queue(), n, dpct::rvalue_ref_to_lvalue_ref(dpct::get_buffer<std::complex<float>>(x_c)), incx, res.get_memory(), oneapi::mkl::index_base::one);
   // CHECK-NEXT: return 0;
   // CHECK-NEXT: }();
@@ -35,12 +35,12 @@ int main(){
   cublasIcamax(handle, n, x_c, incx, result);
 
   // CHECK: status = [&]() {
-  // CHECK-NEXT: dpct::blas::out_mem_int_t res(handle->get_queue(), result);
+  // CHECK-NEXT: dpct::blas::out_mem_int64_int_t res(handle->get_queue(), result);
   // CHECK-NEXT: oneapi::mkl::blas::column_major::iamax(handle->get_queue(), n, dpct::rvalue_ref_to_lvalue_ref(dpct::get_buffer<std::complex<double>>(x_z)), incx, res.get_memory(), oneapi::mkl::index_base::one);
   // CHECK-NEXT: return 0;
   // CHECK-NEXT: }();
   // CHECK: [&]() {
-  // CHECK-NEXT: dpct::blas::out_mem_int_t res(handle->get_queue(), result);
+  // CHECK-NEXT: dpct::blas::out_mem_int64_int_t res(handle->get_queue(), result);
   // CHECK-NEXT: oneapi::mkl::blas::column_major::iamax(handle->get_queue(), n, dpct::rvalue_ref_to_lvalue_ref(dpct::get_buffer<std::complex<double>>(x_z)), incx, res.get_memory(), oneapi::mkl::index_base::one);
   // CHECK-NEXT: return 0;
   // CHECK-NEXT: }();
@@ -48,12 +48,12 @@ int main(){
   cublasIzamax(handle, n, x_z, incx, result);
 
   // CHECK: status = [&]() {
-  // CHECK-NEXT: dpct::blas::out_mem_int_t res(handle->get_queue(), result);
+  // CHECK-NEXT: dpct::blas::out_mem_int64_int_t res(handle->get_queue(), result);
   // CHECK-NEXT: oneapi::mkl::blas::column_major::iamin(handle->get_queue(), n, dpct::rvalue_ref_to_lvalue_ref(dpct::get_buffer<std::complex<float>>(x_c)), incx, res.get_memory(), oneapi::mkl::index_base::one);
   // CHECK-NEXT: return 0;
   // CHECK-NEXT: }();
   // CHECK: [&]() {
-  // CHECK-NEXT: dpct::blas::out_mem_int_t res(handle->get_queue(), result);
+  // CHECK-NEXT: dpct::blas::out_mem_int64_int_t res(handle->get_queue(), result);
   // CHECK-NEXT: oneapi::mkl::blas::column_major::iamin(handle->get_queue(), n, dpct::rvalue_ref_to_lvalue_ref(dpct::get_buffer<std::complex<float>>(x_c)), incx, res.get_memory(), oneapi::mkl::index_base::one);
   // CHECK-NEXT: return 0;
   // CHECK-NEXT: }();
@@ -61,12 +61,12 @@ int main(){
   cublasIcamin(handle, n, x_c, incx, result);
 
   // CHECK: status = [&]() {
-  // CHECK-NEXT: dpct::blas::out_mem_int_t res(handle->get_queue(), result);
+  // CHECK-NEXT: dpct::blas::out_mem_int64_int_t res(handle->get_queue(), result);
   // CHECK-NEXT: oneapi::mkl::blas::column_major::iamin(handle->get_queue(), n, dpct::rvalue_ref_to_lvalue_ref(dpct::get_buffer<std::complex<double>>(x_z)), incx, res.get_memory(), oneapi::mkl::index_base::one);
   // CHECK-NEXT: return 0;
   // CHECK-NEXT: }();
   // CHECK: [&]() {
-  // CHECK-NEXT: dpct::blas::out_mem_int_t res(handle->get_queue(), result);
+  // CHECK-NEXT: dpct::blas::out_mem_int64_int_t res(handle->get_queue(), result);
   // CHECK-NEXT: oneapi::mkl::blas::column_major::iamin(handle->get_queue(), n, dpct::rvalue_ref_to_lvalue_ref(dpct::get_buffer<std::complex<double>>(x_z)), incx, res.get_memory(), oneapi::mkl::index_base::one);
   // CHECK-NEXT: return 0;
   // CHECK-NEXT: }();
diff --git a/clang/test/dpct/cublas_curandInMacro.cu b/clang/test/dpct/cublas_curandInMacro.cu
index 097ab1a62bc9..73124ae805d0 100644
--- a/clang/test/dpct/cublas_curandInMacro.cu
+++ b/clang/test/dpct/cublas_curandInMacro.cu
@@ -51,7 +51,7 @@ int main() {
     cublasErrCheck(cublasSgemm(handle, (cublasOperation_t)trans0, (cublasOperation_t)trans1, N, N, N, &alpha_S, d_A_S, N, d_B_S, N, &beta_S, d_C_S, N));
 
     // CHECK: cublasErrCheck([&]() {
-    // CHECK-NEXT: dpct::blas::out_mem_int_t res(handle->get_queue(), result);
+    // CHECK-NEXT: dpct::blas::out_mem_int64_int_t res(handle->get_queue(), result);
     // CHECK-NEXT: oneapi::mkl::blas::column_major::iamax(handle->get_queue(), N, dpct::rvalue_ref_to_lvalue_ref(dpct::get_buffer<float>(x_S)), N, res.get_memory(), oneapi::mkl::index_base::one);
     // CHECK-NEXT: return 0;
     // CHECK-NEXT: }());
@@ -82,7 +82,7 @@ int main() {
     cublasErrCheck(cublasCgemm(handle, (cublasOperation_t)trans0, (cublasOperation_t)trans1, N, N, N, &alpha_C, d_A_C, N, d_B_C, N, &beta_C, d_C_C, N));
 
     // CHECK: cublasErrCheck([&]() {
-    // CHECK-NEXT: dpct::blas::out_mem_int_t res(handle->get_queue(), result);
+    // CHECK-NEXT: dpct::blas::out_mem_int64_int_t res(handle->get_queue(), result);
     // CHECK-NEXT: oneapi::mkl::blas::column_major::iamax(handle->get_queue(), N, dpct::rvalue_ref_to_lvalue_ref(dpct::get_buffer<std::complex<float>>(x_C)), N, res.get_memory(), oneapi::mkl::index_base::one);
     // CHECK-NEXT: return 0;
     // CHECK-NEXT: }());
diff --git a/clang/test/dpct/macro_test.cu b/clang/test/dpct/macro_test.cu
index 53344e9a662c..c59454338833 100644
--- a/clang/test/dpct/macro_test.cu
+++ b/clang/test/dpct/macro_test.cu
@@ -1234,7 +1234,7 @@ void foo34() {
 //CHECK-NEXT:   int amax(dpct::blas::descriptor_ptr handle, const int n, const float *X,     \
 //CHECK-NEXT:            const int incX, int &result) try {                                  \
 //CHECK-NEXT:     return [&]() {                                                             \
-//CHECK-NEXT:       dpct::blas::out_mem_int_t res(handle->get_queue(), &result);             \
+//CHECK-NEXT:       dpct::blas::out_mem_int64_int_t res(handle->get_queue(), &result);       \
 //CHECK-NEXT:       oneapi::mkl::blas::column_major::iamax(handle->get_queue(), n, X, incX,  \
 //CHECK-NEXT:                                              res.get_memory(),                 \
 //CHECK-NEXT:                                              oneapi::mkl::index_base::one);    \
diff --git a/clang/test/dpct/nestedqueue.cu b/clang/test/dpct/nestedqueue.cu
index a2b947dbd34b..e72f82f0d23f 100644
--- a/clang/test/dpct/nestedqueue.cu
+++ b/clang/test/dpct/nestedqueue.cu
@@ -67,7 +67,7 @@ __host__ void foo4(){
   int incx=1;
   int* result =0;
   // CHECK: [&]() {
-  // CHECK-NEXT: dpct::blas::out_mem_int_t res(handle->get_queue(), result);
+  // CHECK-NEXT: dpct::blas::out_mem_int64_int_t res(handle->get_queue(), result);
   // CHECK-NEXT: oneapi::mkl::blas::column_major::iamax(handle->get_queue(), n, dpct::rvalue_ref_to_lvalue_ref(dpct::get_buffer<float>(x_S)), incx, res.get_memory(), oneapi::mkl::index_base::one);
   // CHECK-NEXT: return 0;
   // CHECK-NEXT: }();
diff --git a/clang/test/dpct/query_api_mapping/cuBLAS/blas_part2.cu b/clang/test/dpct/query_api_mapping/cuBLAS/blas_part2.cu
index 5ae54c4b1c61..c84b4c15ccf2 100644
--- a/clang/test/dpct/query_api_mapping/cuBLAS/blas_part2.cu
+++ b/clang/test/dpct/query_api_mapping/cuBLAS/blas_part2.cu
@@ -219,7 +219,7 @@
 // cublasIzamax-NEXT:                x /*const cuDoubleComplex **/, incx /*int*/, res /*int **/);
 // cublasIzamax-NEXT: Is migrated to (with the option --no-dry-pattern):
 // cublasIzamax-NEXT:   [&]() {
-// cublasIzamax-NEXT:   dpct::blas::out_mem_int_t res(handle->get_queue(), res);
+// cublasIzamax-NEXT:   dpct::blas::out_mem_int64_int_t res(handle->get_queue(), res);
 // cublasIzamax-NEXT:   oneapi::mkl::blas::column_major::iamax(handle->get_queue(), n, (std::complex<double>*)x, incx, res.get_memory(), oneapi::mkl::index_base::one);
 // cublasIzamax-NEXT:   return 0;
 // cublasIzamax-NEXT:   }();
diff --git a/clang/test/dpct/query_api_mapping/cuBLAS/blas_part3.cu b/clang/test/dpct/query_api_mapping/cuBLAS/blas_part3.cu
index b3f035c44509..5a5c53d75456 100644
--- a/clang/test/dpct/query_api_mapping/cuBLAS/blas_part3.cu
+++ b/clang/test/dpct/query_api_mapping/cuBLAS/blas_part3.cu
@@ -124,7 +124,7 @@
 // cublasIcamin-NEXT:                incx /*int*/, res /*int **/);
 // cublasIcamin-NEXT: Is migrated to (with the option --no-dry-pattern):
 // cublasIcamin-NEXT:   [&]() {
-// cublasIcamin-NEXT:   dpct::blas::out_mem_int_t res(handle->get_queue(), res);
+// cublasIcamin-NEXT:   dpct::blas::out_mem_int64_int_t res(handle->get_queue(), res);
 // cublasIcamin-NEXT:   oneapi::mkl::blas::column_major::iamin(handle->get_queue(), n, (std::complex<float>*)x, incx, res.get_memory(), oneapi::mkl::index_base::one);
 // cublasIcamin-NEXT:   return 0;
 // cublasIcamin-NEXT:   }();
@@ -135,7 +135,7 @@
 // cublasIcamax-NEXT:                incx /*int*/, res /*int **/);
 // cublasIcamax-NEXT: Is migrated to (with the option --no-dry-pattern):
 // cublasIcamax-NEXT:   [&]() {
-// cublasIcamax-NEXT:   dpct::blas::out_mem_int_t res(handle->get_queue(), res);
+// cublasIcamax-NEXT:   dpct::blas::out_mem_int64_int_t res(handle->get_queue(), res);
 // cublasIcamax-NEXT:   oneapi::mkl::blas::column_major::iamax(handle->get_queue(), n, (std::complex<float>*)x, incx, res.get_memory(), oneapi::mkl::index_base::one);
 // cublasIcamax-NEXT:   return 0;
 // cublasIcamax-NEXT:   }();
diff --git a/clang/test/dpct/query_api_mapping/cuBLAS/blas_part4.cu b/clang/test/dpct/query_api_mapping/cuBLAS/blas_part4.cu
index 16ddc13ccb56..af18fc4afbb9 100644
--- a/clang/test/dpct/query_api_mapping/cuBLAS/blas_part4.cu
+++ b/clang/test/dpct/query_api_mapping/cuBLAS/blas_part4.cu
@@ -247,7 +247,7 @@
 // cublasIzamin-NEXT:                x /*const cuDoubleComplex **/, incx /*int*/, res /*int **/);
 // cublasIzamin-NEXT: Is migrated to (with the option --no-dry-pattern):
 // cublasIzamin-NEXT:   [&]() {
-// cublasIzamin-NEXT:   dpct::blas::out_mem_int_t res(handle->get_queue(), res);
+// cublasIzamin-NEXT:   dpct::blas::out_mem_int64_int_t res(handle->get_queue(), res);
 // cublasIzamin-NEXT:   oneapi::mkl::blas::column_major::iamin(handle->get_queue(), n, (std::complex<double>*)x, incx, res.get_memory(), oneapi::mkl::index_base::one);
 // cublasIzamin-NEXT:   return 0;
 // cublasIzamin-NEXT:   }();
diff --git a/clang/test/dpct/query_api_mapping/cuBLAS/blas_part5.cu b/clang/test/dpct/query_api_mapping/cuBLAS/blas_part5.cu
index 6fc33d8d06e6..cef3c4c0faa8 100644
--- a/clang/test/dpct/query_api_mapping/cuBLAS/blas_part5.cu
+++ b/clang/test/dpct/query_api_mapping/cuBLAS/blas_part5.cu
@@ -177,7 +177,7 @@
 // cublasIsamax-NEXT:                incx /*int*/, res /*int **/);
 // cublasIsamax-NEXT: Is migrated to (with the option --no-dry-pattern):
 // cublasIsamax-NEXT:   [&]() {
-// cublasIsamax-NEXT:   dpct::blas::out_mem_int_t res(handle->get_queue(), res);
+// cublasIsamax-NEXT:   dpct::blas::out_mem_int64_int_t res(handle->get_queue(), res);
 // cublasIsamax-NEXT:   oneapi::mkl::blas::column_major::iamax(handle->get_queue(), n, x, incx, res.get_memory(), oneapi::mkl::index_base::one);
 // cublasIsamax-NEXT:   return 0;
 // cublasIsamax-NEXT:   }();
diff --git a/clang/test/dpct/query_api_mapping/cuBLAS/blas_part6.cu b/clang/test/dpct/query_api_mapping/cuBLAS/blas_part6.cu
index eaa54a0801ee..0cc2e53d5364 100644
--- a/clang/test/dpct/query_api_mapping/cuBLAS/blas_part6.cu
+++ b/clang/test/dpct/query_api_mapping/cuBLAS/blas_part6.cu
@@ -13,7 +13,7 @@
 // cublasIdamin-NEXT:                incx /*int*/, res /*int **/);
 // cublasIdamin-NEXT: Is migrated to (with the option --no-dry-pattern):
 // cublasIdamin-NEXT:   [&]() {
-// cublasIdamin-NEXT:   dpct::blas::out_mem_int_t res(handle->get_queue(), res);
+// cublasIdamin-NEXT:   dpct::blas::out_mem_int64_int_t res(handle->get_queue(), res);
 // cublasIdamin-NEXT:   oneapi::mkl::blas::column_major::iamin(handle->get_queue(), n, x, incx, res.get_memory(), oneapi::mkl::index_base::one);
 // cublasIdamin-NEXT:   return 0;
 // cublasIdamin-NEXT:   }();
@@ -140,7 +140,7 @@
 // cublasIdamax-NEXT:                incx /*int*/, res /*int **/);
 // cublasIdamax-NEXT: Is migrated to (with the option --no-dry-pattern):
 // cublasIdamax-NEXT:   [&]() {
-// cublasIdamax-NEXT:   dpct::blas::out_mem_int_t res(handle->get_queue(), res);
+// cublasIdamax-NEXT:   dpct::blas::out_mem_int64_int_t res(handle->get_queue(), res);
 // cublasIdamax-NEXT:   oneapi::mkl::blas::column_major::iamax(handle->get_queue(), n, x, incx, res.get_memory(), oneapi::mkl::index_base::one);
 // cublasIdamax-NEXT:   return 0;
 // cublasIdamax-NEXT:   }();
@@ -162,7 +162,7 @@
 // cublasIsamin-NEXT:                incx /*int*/, res /*int **/);
 // cublasIsamin-NEXT: Is migrated to (with the option --no-dry-pattern):
 // cublasIsamin-NEXT:   [&]() {
-// cublasIsamin-NEXT:   dpct::blas::out_mem_int_t res(handle->get_queue(), res);
+// cublasIsamin-NEXT:   dpct::blas::out_mem_int64_int_t res(handle->get_queue(), res);
 // cublasIsamin-NEXT:   oneapi::mkl::blas::column_major::iamin(handle->get_queue(), n, x, incx, res.get_memory(), oneapi::mkl::index_base::one);
 // cublasIsamin-NEXT:   return 0;
 // cublasIsamin-NEXT:   }();
diff --git a/clang/test/dpct/test_libaray_index/test3.cu b/clang/test/dpct/test_libaray_index/test3.cu
index 0aa04cb565c5..c87ddc6d1e71 100644
--- a/clang/test/dpct/test_libaray_index/test3.cu
+++ b/clang/test/dpct/test_libaray_index/test3.cu
@@ -12,7 +12,7 @@ int foo () {
   int *result;
 
   //CHECK:[&]() {
-  //CHECK-NEXT:dpct::blas::out_mem_int_t res(handle->get_queue(), result);
+  //CHECK-NEXT:dpct::blas::out_mem_int64_int_t res(handle->get_queue(), result);
   //CHECK-NEXT:oneapi::mkl::blas::column_major::iamax(handle->get_queue(), N, x1, N, res.get_memory(), oneapi::mkl::index_base::one);
   //CHECK-NEXT:return 0;
   //CHECK-NEXT:}();