[SYCLomatic] Refine migration of CUDA_ARRAY_DESCRIPTOR and 2 related APIs.

clang/examples/DPCT/Driver/cuTexRefSetAddress2D.cu

@@ -1,7 +1,7 @@
 void test(const CUDA_ARRAY_DESCRIPTOR *pa, CUdeviceptr d, size_t s) {
   // Start
   CUtexref t;
-  cuTexRefSetAddress2D(t /*CUtexref*/, pa /*size_t **/, d /*CUdeviceptr*/,
-                       s /*size_t*/);
+  cuTexRefSetAddress2D(t /*CUtexref*/, pa /*const CUDA_ARRAY_DESCRIPTOR **/,
+                       d /*CUdeviceptr*/, s /*size_t*/);
   // End
 }

clang/lib/DPCT/APINames.inc

@@ -1625,7 +1625,7 @@ ENTRY(cuGraphicsUnregisterResource, cuGraphicsUnregisterResource, false, NO_FLAG
 ENTRY(cuInit, cuInit, true, NO_FLAG, P4, "DPCT1026/DPCT1027")
 ENTRY(cuArray3DCreate, cuArray3DCreate_v2, false, NO_FLAG, P4, "comment")
 ENTRY(cuArray3DGetDescriptor, cuArray3DGetDescriptor_v2, false, NO_FLAG, P4, "comment")
-ENTRY(cuArrayCreate, cuArrayCreate_v2, true, NO_FLAG, P4, "DPCT1073")
+ENTRY(cuArrayCreate, cuArrayCreate_v2, true, NO_FLAG, P4, "Successful")
 ENTRY(cuArrayDestroy, cuArrayDestroy, true, NO_FLAG, P4, "Successful")
 ENTRY(cuArrayGetDescriptor, cuArrayGetDescriptor_v2, false, NO_FLAG, P4, "comment")
 ENTRY(cuDeviceGetByPCIBusId, cuDeviceGetByPCIBusId, false, NO_FLAG, P4, "comment")

clang/lib/DPCT/APINamesTexture.inc

@@ -181,17 +181,11 @@ FEATURE_REQUEST_FACTORY(HelperFeatureEnum::device_ext,
 ENTRY_UNSUPPORTED("cudaGetTextureObjectResourceViewDesc",
                   Diagnostics::API_NOT_MIGRATED)
 
-CONDITIONAL_FACTORY_ENTRY(
-    CheckWarning1073(1),
-    FEATURE_REQUEST_FACTORY(
-        HelperFeatureEnum::device_ext,
-        ASSIGNABLE_FACTORY(ASSIGN_FACTORY_ENTRY(
-            "cuArrayCreate_v2", DEREF(0),
-            NEW(MapNames::getDpctNamespace() + "image_matrix",
-                STRUCT_DISMANTLE(1, "channel_type_ct1", "channel_num_ct1",
-                                 "x_ct1", "y_ct1"))))),
-    UNSUPPORT_FACTORY_ENTRY("cuArrayCreate_v2",
-                            Diagnostics::CANNOT_CAPUTURE_AGUMENTS, ARG(1)))
+FEATURE_REQUEST_FACTORY(HelperFeatureEnum::device_ext,
+                        ASSIGNABLE_FACTORY(ASSIGN_FACTORY_ENTRY(
+                            "cuArrayCreate_v2", DEREF(0),
+                            NEW(MapNames::getDpctNamespace() + "image_matrix",
+                                ARG(1)))))
 ASSIGNABLE_FACTORY(DELETER_FACTORY_ENTRY("cuArrayDestroy", ARG(0)))
 ENTRY_UNSUPPORTED("cuTexObjectGetResourceViewDesc", Diagnostics::API_NOT_MIGRATED)
 FEATURE_REQUEST_FACTORY(
@@ -236,8 +230,7 @@ FEATURE_REQUEST_FACTORY(HelperFeatureEnum::device_ext,
 FEATURE_REQUEST_FACTORY(HelperFeatureEnum::device_ext,
                         ASSIGNABLE_FACTORY(MEMBER_CALL_FACTORY_ENTRY(
                             "cuTexRefSetAddress2D_v3", ARG(0), true, "attach",
-                            ARG(2), STRUCT_DISMANTLE(1, "x_ct1", "y_ct1"),
-                            ARG(3))))
+                            ARG(1), ARG(2), ARG(3))))
 
 FEATURE_REQUEST_FACTORY(
     HelperFeatureEnum::device_ext,

clang/lib/DPCT/ASTTraversal.cpp

@@ -1753,7 +1753,7 @@ void TypeInDeclRule::registerMatcher(MatchFinder &MF) {
               "cudaLaunchAttributeValue", "cusparseSpSMDescr_t",
               "cusparseConstSpMatDescr_t", "cusparseSpSMAlg_t",
               "cusparseConstDnMatDescr_t", "cudaMemcpy3DParms", "CUDA_MEMCPY3D",
-              "CUDA_MEMCPY2D"))))))
+              "CUDA_MEMCPY2D", "CUDA_ARRAY_DESCRIPTOR"))))))
           .bind("cudaTypeDef"),
       this);
 
@@ -11494,21 +11494,7 @@ TextModification *ReplaceMemberAssignAsSetMethod(const Expr *E,
                                         ME, MethodName, ReplacedArg, ExtraArg);
 }
 
-void MemoryDataTypeRule::emplaceCuArrayDescDeclarations(const VarDecl *VD) {
-  if (DpctGlobalInfo::isCommentsEnabled()) {
-    emplaceTransformation(ReplaceVarDecl::getVarDeclReplacement(
-        VD, "// These variables are defined for info of image_matrix."));
-  }
-  emplaceParamDecl(VD, "size_t", false, "0", "x", "y");
-  emplaceParamDecl(VD, "unsigned", false, "0", "channel_num");
-  emplaceParamDecl(VD, MapNames::getClNamespace() + "image_channel_type", false,
-                   "0", "channel_type");
-}
-
 void MemoryDataTypeRule::registerMatcher(MatchFinder &MF) {
-  MF.addMatcher(varDecl(hasType(namedDecl(hasAnyName("CUDA_ARRAY_DESCRIPTOR"))))
-                    .bind("decl"),
-                this);
   MF.addMatcher(memberExpr(hasObjectExpression(declRefExpr(hasType(namedDecl(
                                hasAnyName("CUDA_ARRAY_DESCRIPTOR"))))))
                     .bind("arrayMember"),
@@ -11530,19 +11516,12 @@ void MemoryDataTypeRule::registerMatcher(MatchFinder &MF) {
 }
 
 void MemoryDataTypeRule::runRule(const MatchFinder::MatchResult &Result) {
-  if (auto VD = getNodeAsType<VarDecl>(Result, "decl")) {
-    if (isa<ParmVarDecl>(VD))
-      return;
-    auto TypeName = DpctGlobalInfo::getUnqualifiedTypeName(VD->getType());
-    if (TypeName == "CUDA_ARRAY_DESCRIPTOR")
-      emplaceCuArrayDescDeclarations(VD);
-  } else if (auto ME = getNodeAsType<MemberExpr>(Result, "arrayMember")) {
-    if (auto DRE =
-            dyn_cast<DeclRefExpr>(ME->getBase()->IgnoreImplicitAsWritten())) {
-      emplaceTransformation(new ReplaceStmt(
-          ME, getArrayDescMemberName(DRE->getDecl()->getName(),
-                                     ME->getMemberDecl()->getName().str())));
-    }
+  if (auto ME = getNodeAsType<MemberExpr>(Result, "arrayMember")) {
+    const auto &Replace = MapNames::findReplacedName(
+        ArrayDescMemberNames, ME->getMemberDecl()->getName().str());
+    if (!Replace.empty())
+      emplaceTransformation(new ReplaceToken(
+          ME->getMemberLoc(), ME->getEndLoc(), std::string(Replace)));
   } else if (auto CE = getNodeAsType<CallExpr>(Result, "makeData")) {
     if (auto FD = CE->getDirectCallee()) {
       auto Name = FD->getName();

clang/lib/DPCT/ASTTraversal.h

@@ -1527,8 +1527,6 @@ class MemoryDataTypeRule : public NamedMigrationRule<MemoryDataTypeRule> {
   const static std::vector<std::string> RemoveMember;
 
 public:
-  void emplaceCuArrayDescDeclarations(const VarDecl *VD);
-
   static std::string getArrayDescMemberName(StringRef BaseName,
                                             const std::string &Member) {
     auto Itr = ArrayDescMemberNames.find(Member);

clang/lib/DPCT/CallExprRewriterCommon.h

@@ -324,23 +324,6 @@ makeCombinedArg(std::function<T1(const CallExpr *)> Part1,
   };
 }
 
-inline std::function<std::vector<RenameWithSuffix>(const CallExpr *)>
-makeStructDismantler(unsigned Idx, const std::vector<std::string> &Suffixes) {
-  return [=](const CallExpr *C) -> std::vector<RenameWithSuffix> {
-    std::vector<RenameWithSuffix> Ret;
-    if (auto DRE = dyn_cast_or_null<DeclRefExpr>(
-            getDereferencedExpr(C->getArg(Idx)))) {
-      Ret.reserve(Suffixes.size());
-      auto Origin = DRE->getDecl()->getName();
-      std::transform(Suffixes.begin(), Suffixes.end(), std::back_inserter(Ret),
-                     [&](StringRef Suffix) -> RenameWithSuffix {
-                       return RenameWithSuffix(Origin, Suffix);
-                     });
-    }
-    return Ret;
-  };
-}
-
 inline std::function<std::string(const CallExpr *)>
 makeExtendStr(unsigned Idx, const std::string Suffix) {
   return [=](const CallExpr *C) -> std::string {
@@ -1561,17 +1544,6 @@ createDerefExprRewriterFactory(
       std::forward<std::function<ArgT(const CallExpr *)>>(ArgCreator));
 }
 
-class CheckWarning1073 {
-  unsigned Idx;
-
-public:
-  CheckWarning1073(unsigned I) : Idx(I) {}
-  bool operator()(const CallExpr *C) {
-    auto DerefE = getDereferencedExpr(C->getArg(Idx));
-    return DerefE && isa<DeclRefExpr>(DerefE);
-  }
-};
-
 // sycl has 2 overloading of malloc_device
 // 1. sycl::malloc_device(Addr, Size)
 // 2. sycl::malloc_device<type>(Addr, Size)
@@ -2050,7 +2022,6 @@ const std::string MipmapNeedBindlessImage =
 #define ADDROF(x) makeAddrOfExprCreator(x)
 #define DEREF(x) makeDerefExprCreator(x)
 #define DEREF_CAST_IF_NEED(T, S) makeDerefCastIfNeedExprCreator(T, S)
-#define STRUCT_DISMANTLE(idx, ...) makeStructDismantler(idx, {__VA_ARGS__})
 #define ARG(x) makeCallArgCreator(x)
 #define ARG_WC(x) makeDerefArgCreatorWithCall(x)
 #define TEMPLATE_ARG(x) makeCallArgCreatorFromTemplateArg(x)

clang/lib/DPCT/Diagnostics.inc

@@ -176,8 +176,8 @@ DEF_WARNING(OUT_OF_PLACE_FFT_EXEC, 1071, LOW_LEVEL, "deprecated")
 DEF_COMMENT(OUT_OF_PLACE_FFT_EXEC, 1071, LOW_LEVEL, "deprecated")
 DEF_WARNING(UNSUPPORT_FREE_MEMORY_SIZE, 1072, LOW_LEVEL, "SYCL currently does not support getting the available memory on the current device. You may need to adjust the code.")
 DEF_COMMENT(UNSUPPORT_FREE_MEMORY_SIZE, 1072, LOW_LEVEL, "SYCL currently does not support getting the available memory on the current device. You may need to adjust the code.")
-DEF_WARNING(CANNOT_CAPUTURE_AGUMENTS, 1073, MEDIUM_LEVEL, "The field values of parameter '%0' could not be deduced, so the call was not migrated. You need to update this code manually.")
-DEF_COMMENT(CANNOT_CAPUTURE_AGUMENTS, 1073, MEDIUM_LEVEL, "The field values of parameter '{0}' could not be deduced, so the call was not migrated. You need to update this code manually.")
+DEF_WARNING(CANNOT_CAPUTURE_AGUMENTS, 1073, MEDIUM_LEVEL, "deprecated")
+DEF_COMMENT(CANNOT_CAPUTURE_AGUMENTS, 1073, MEDIUM_LEVEL, "deprecated")
 DEF_WARNING(TEX_FLAG_UNSUPPORT, 1074, MEDIUM_LEVEL, "The SYCL Image class does not support some of the flags used in the original code. Unsupported flags were ignored. Data read from SYCL Image could not be normalized as specified in the original code.")
 DEF_COMMENT(TEX_FLAG_UNSUPPORT, 1074, MEDIUM_LEVEL, "The SYCL Image class does not support some of the flags used in the original code. Unsupported flags were ignored. Data read from SYCL Image could not be normalized as specified in the original code.")
 DEF_WARNING(CHECK_RELATED_QUEUE, 1075, LOW_LEVEL, "deprecated")

clang/lib/DPCT/MapNames.cpp

@@ -373,6 +373,8 @@ void MapNames::setExplicitNamespaceMap() {
                                       HelperFeatureEnum::device_ext)},
       {"cudaMemcpyKind",
        std::make_shared<TypeNameRule>(getDpctNamespace() + "memcpy_direction")},
+      {"CUDA_ARRAY_DESCRIPTOR", std::make_shared<TypeNameRule>(
+                                    getDpctNamespace() + "image_matrix_desc")},
       {"cudaMemcpy3DParms",
        std::make_shared<TypeNameRule>(getDpctNamespace() + "memcpy_parameter")},
       {"CUDA_MEMCPY3D",
@@ -4123,8 +4125,8 @@ const MapNames::MapTy MemoryDataTypeRule::ExtentMemberNames{
     {"width", "[0]"}, {"height", "[1]"}, {"depth", "[2]"}};
 
 const MapNames::MapTy MemoryDataTypeRule::ArrayDescMemberNames{
-    {"Width", "x"},
-    {"Height", "y"},
+    {"Width", "width"},
+    {"Height", "height"},
     {"Format", "channel_type"},
     {"NumChannels", "channel_num"}};
 

clang/lib/DPCT/TypeNames.inc

@@ -41,7 +41,6 @@ ENTRY_TYPE(CUgraphNode, false, NO_FLAG, P4, "comment")
 ENTRY_TYPE(CUgraphicsResource, false, NO_FLAG, P4, "comment")
 
 // CUDA Runtime Library
-ENTRY_TYPE(CUDA_ARRAY_DESCRIPTOR*, false, NO_FLAG, P4, "comment")
 ENTRY_TYPE(cudaKernelNodeParams, false, NO_FLAG, P4, "comment")
 
 // cuDNN Library

clang/runtime/dpct-rt/include/dpct/image.hpp

@@ -139,6 +139,14 @@ static image_wrapper_base *create_image_wrapper(image_channel channel, int dims)
 
 } // namespace detail
 
+struct image_matrix_desc {
+  size_t width = 0;
+  size_t height = 0;
+  sycl::image_channel_type channel_type =
+      sycl::image_channel_type::signed_int32;
+  unsigned channel_num = 0;
+};
+
 /// Image channel info, include channel number, order, data width and type
 class image_channel {
   image_channel_data_type _type = image_channel_data_type::signed_int;
@@ -331,6 +339,9 @@ class image_matrix {
     }
     _host_data = std::malloc(_range[0] * _range[1] * _channel.get_total_size());
   }
+  image_matrix(const image_matrix_desc *desc)
+      : image_matrix(desc->channel_type, desc->channel_num, desc->width,
+                     desc->height) {}
 
   /// Construct a new image class with the matrix data.
   template <int dimensions> sycl::image<dimensions> *create_image() {
@@ -626,6 +637,15 @@ class image_wrapper_base {
     image_wrapper_base::set_data(
         image_data(const_cast<void *>(data), x, y, pitch, channel));
   }
+  /// Attach device_ptr data to this class.
+  void attach(const image_matrix_desc *desc, device_ptr ptr, size_t pitch) {
+    detach();
+    image_channel channel;
+    channel.set_channel_num(desc->channel_num);
+    channel.set_channel_type(desc->channel_type);
+    image_wrapper_base::set_data(
+        image_data(ptr, desc->width, desc->height, pitch, channel));
+  }
   /// Detach data.
   virtual void detach() {}
 

clang/test/dpct/know_unsupported_type.cu

@@ -8,10 +8,7 @@
 #include "nvml.h"
 #include <vector>
 int main(int argc, char **argv) {
-    // CHECK: /*
-    // CHECK-NEXT: DPCT1082:{{[0-9]+}}: Migration of CUDA_ARRAY_DESCRIPTOR * type is not supported.
-    // CHECK-NEXT: */
-    // CHECK-NEXT: CUDA_ARRAY_DESCRIPTOR *pcad;
+    // CHECK: dpct::image_matrix_desc *pcad;
     CUDA_ARRAY_DESCRIPTOR *pcad;
     // CHECK: dpct::memcpy_parameter *p1c3d;
     cudaMemcpy3DParms *p1c3d;

clang/test/dpct/query_api_mapping/Driver/test.cu

@@ -313,11 +313,11 @@
 // RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=cuTexRefSetAddress2D | FileCheck %s -check-prefix=CUTEXREFSETADDRESS2D
 // CUTEXREFSETADDRESS2D: CUDA API:
 // CUTEXREFSETADDRESS2D-NEXT:   CUtexref t;
-// CUTEXREFSETADDRESS2D-NEXT:   cuTexRefSetAddress2D(t /*CUtexref*/, pa /*size_t **/, d /*CUdeviceptr*/,
-// CUTEXREFSETADDRESS2D-NEXT:                        s /*size_t*/);
+// CUTEXREFSETADDRESS2D-NEXT:   cuTexRefSetAddress2D(t /*CUtexref*/, pa /*const CUDA_ARRAY_DESCRIPTOR **/,
+// CUTEXREFSETADDRESS2D-NEXT:                        d /*CUdeviceptr*/, s /*size_t*/);
 // CUTEXREFSETADDRESS2D-NEXT: Is migrated to:
 // CUTEXREFSETADDRESS2D-NEXT:   dpct::image_wrapper_base_p t;
-// CUTEXREFSETADDRESS2D-NEXT:   t->attach(d, s);
+// CUTEXREFSETADDRESS2D-NEXT:   t->attach(pa, d, s);
 
 // RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=cuTexRefSetAddressMode | FileCheck %s -check-prefix=CUTEXREFSETADDRESSMODE
 // CUTEXREFSETADDRESSMODE: CUDA API:

clang/test/dpct/texture_driver.cu

@@ -1,5 +1,6 @@
 // RUN: dpct --format-range=none --usm-level=none -out-root %T/texture_driver %s --cuda-include-path="%cuda-path/include" --sycl-named-lambda -- -x cuda --cuda-host-only -std=c++14
 // RUN: FileCheck --input-file %T/texture_driver/texture_driver.dp.cpp --match-full-lines %s
+// RUN: %if build_lit %{icpx -c -fsycl %T/texture_driver/texture_driver.dp.cpp -o %T//texture_driver/texture_driver.dp.o %}
 
 #include "cuda.h"
 #include <stdio.h>
@@ -19,26 +20,22 @@ void funcT(T t) {}
 
 int main() {
 
-  // CHECK: size_t halfDesc_x_ct1, halfDesc_y_ct1;
-  // CHECK-NEXT: unsigned halfDesc_channel_num_ct1;
-  // CHECK-NEXT: sycl::image_channel_type halfDesc_channel_type_ct1;
-  // CHECK-NEXT: halfDesc_y_ct1 = 32;
-  // CHECK-NEXT: halfDesc_x_ct1 = 64;
-  // CHECK-NEXT: halfDesc_channel_type_ct1 = sycl::image_channel_type::fp16;
-  // CHECK-NEXT: halfDesc_channel_num_ct1 = 1;
+  // CHECK: dpct::image_matrix_desc halfDesc;
+  // CHECK-NEXT: halfDesc.height = 32;
+  // CHECK-NEXT: halfDesc.width = 64;
+  // CHECK-NEXT: halfDesc.channel_type = sycl::image_channel_type::fp16;
+  // CHECK-NEXT: halfDesc.channel_num = 1;
   CUDA_ARRAY_DESCRIPTOR halfDesc;
   halfDesc.Height = 32;
   halfDesc.Width = 64;
   halfDesc.Format = CU_AD_FORMAT_HALF;
   halfDesc.NumChannels = 1;
 
-  // CHECK: size_t float4Desc_x_ct1, float4Desc_y_ct1;
-  // CHECK-NEXT: unsigned float4Desc_channel_num_ct1;
-  // CHECK-NEXT: sycl::image_channel_type float4Desc_channel_type_ct1;
-  // CHECK-NEXT: float4Desc_x_ct1 = 64;
-  // CHECK-NEXT: float4Desc_channel_type_ct1 = sycl::image_channel_type::fp32;
-  // CHECK-NEXT: float4Desc_channel_num_ct1 = 4;
-  // CHECK-NEXT: float4Desc_y_ct1 = 32;
+  // CHECK: dpct::image_matrix_desc float4Desc;
+  // CHECK-NEXT: float4Desc.width = 64;
+  // CHECK-NEXT: float4Desc.channel_type = sycl::image_channel_type::fp32;
+  // CHECK-NEXT: float4Desc.channel_num = 4;
+  // CHECK-NEXT: float4Desc.height = 32;
   CUDA_ARRAY_DESCRIPTOR float4Desc;
   float4Desc.Width = 64;
   float4Desc.Format = CU_AD_FORMAT_FLOAT;
@@ -47,8 +44,8 @@ int main() {
 
   // CHECK: dpct::image_matrix **a_ptr = new dpct::image_matrix_p;
   // CHECK-NEXT: dpct::image_matrix_p a42;
-  // CHECK-NEXT: *a_ptr = new dpct::image_matrix(halfDesc_channel_type_ct1, halfDesc_channel_num_ct1, halfDesc_x_ct1, halfDesc_y_ct1);
-  // CHECK-NEXT: a42 = new dpct::image_matrix(float4Desc_channel_type_ct1, float4Desc_channel_num_ct1, float4Desc_x_ct1, float4Desc_y_ct1);
+  // CHECK-NEXT: *a_ptr = new dpct::image_matrix(&halfDesc);
+  // CHECK-NEXT: a42 = new dpct::image_matrix(&float4Desc);
   // CHECK-NEXT: delete (*a_ptr);
   // CHECK-NEXT: delete a42;
   // CHECK-NEXT: delete a_ptr;
@@ -64,25 +61,25 @@ int main() {
   {
     int errorCode;
 
-    // CHECK: errorCode = DPCT_CHECK_ERROR(a42 = new dpct::image_matrix(float4Desc_channel_type_ct1, float4Desc_channel_num_ct1, float4Desc_x_ct1, float4Desc_y_ct1));
+    // CHECK: errorCode = DPCT_CHECK_ERROR(a42 = new dpct::image_matrix(&float4Desc));
     errorCode = cuArrayCreate(&a42, &float4Desc);
     // CHECK: errorCode = DPCT_CHECK_ERROR(delete a42);
     errorCode = cuArrayDestroy(a42);
 
 
-    // CHECK: cudaCheck(DPCT_CHECK_ERROR(a42 = new dpct::image_matrix(float4Desc_channel_type_ct1, float4Desc_channel_num_ct1, float4Desc_x_ct1, float4Desc_y_ct1)));
+    // CHECK: cudaCheck(DPCT_CHECK_ERROR(a42 = new dpct::image_matrix(&float4Desc)));
     cudaCheck(cuArrayCreate(&a42, &float4Desc));
     // CHECK: cudaCheck(DPCT_CHECK_ERROR(delete a42));
     cudaCheck(cuArrayDestroy(a42));
 
 
-    // CHECK: func(DPCT_CHECK_ERROR(a42 = new dpct::image_matrix(float4Desc_channel_type_ct1, float4Desc_channel_num_ct1, float4Desc_x_ct1, float4Desc_y_ct1)));
+    // CHECK: func(DPCT_CHECK_ERROR(a42 = new dpct::image_matrix(&float4Desc)));
     func(cuArrayCreate(&a42, &float4Desc));
     // CHECK: func(DPCT_CHECK_ERROR(delete a42));
     func(cuArrayDestroy(a42));
 
 
-    // CHECK: funcT(DPCT_CHECK_ERROR(a42 = new dpct::image_matrix(float4Desc_channel_type_ct1, float4Desc_channel_num_ct1, float4Desc_x_ct1, float4Desc_y_ct1)));
+    // CHECK: funcT(DPCT_CHECK_ERROR(a42 = new dpct::image_matrix(&float4Desc)));
     funcT(cuArrayCreate(&a42, &float4Desc));
     // CHECK: funcT(DPCT_CHECK_ERROR(delete a42));
     funcT(cuArrayDestroy(a42));
@@ -92,26 +89,14 @@ int main() {
 void create_array_fail() {
   CUarray a;
   unsigned i;
-  // CHECK: CUDA_ARRAY_DESCRIPTOR d[20], *p;
+  // CHECK: dpct::image_matrix_desc d[20], *p;
   CUDA_ARRAY_DESCRIPTOR d[20], *p;
   p = &d[5];
 
-  // CHECK: /*
-  // CHECK-NEXT:  DPCT1073:{{[0-9]+}}: The field values of parameter 'd' could not be deduced, so the call was not migrated. You need to update this code manually.
-  // CHECK-NEXT: */
-  // CHECK-NEXT: cuArrayCreate(&a, d);
-  // CHECK-NEXT: /*
-  // CHECK-NEXT:  DPCT1073:{{[0-9]+}}: The field values of parameter 'p' could not be deduced, so the call was not migrated. You need to update this code manually.
-  // CHECK-NEXT: */
-  // CHECK-NEXT: cuArrayCreate(&a, p);
-  // CHECK-NEXT: /*
-  // CHECK-NEXT:  DPCT1073:{{[0-9]+}}: The field values of parameter 'p + i' could not be deduced, so the call was not migrated. You need to update this code manually.
-  // CHECK-NEXT: */
-  // CHECK-NEXT: cuArrayCreate(&a, p + i);
-  // CHECK-NEXT: /*
-  // CHECK-NEXT:  DPCT1073:{{[0-9]+}}: The field values of parameter '&d[i]' could not be deduced, so the call was not migrated. You need to update this code manually.
-  // CHECK-NEXT: */
-  // CHECK-NEXT: cuArrayCreate(&a, &d[i]);
+  // CHECK: a = new dpct::image_matrix(d);
+  // CHECK-NEXT: a = new dpct::image_matrix(p);
+  // CHECK-NEXT: a = new dpct::image_matrix(p + i);
+  // CHECK-NEXT: a = new dpct::image_matrix(&d[i]);
   cuArrayCreate(&a, d);
   cuArrayCreate(&a, p);
   cuArrayCreate(&a, p + i);