Removed unsupported AI migration test cases

clang/lib/DPCT/RulesAsm/AsmMigration.cpp

@@ -590,7 +590,7 @@ bool SYCLGenBase::emitVariableDeclaration(const InlineAsmVarDecl *D) {
 bool SYCLGenBase::emitAddressExpr(const InlineAsmAddressExpr *Dst) {
   // Address expression only support ld/st & atom instructions.
   if (!CurrInst || !CurrInst->is(asmtok::op_st, asmtok::op_ld, asmtok::op_atom,
-                                 asmtok::op_prefetch, asmtok::op_prefetchu))
+                                 asmtok::op_prefetch))
     return SYCLGenError();
   std::string Type;
   if (tryEmitType(Type, CurrInst->getType(0)))
@@ -618,8 +618,7 @@ bool SYCLGenBase::emitAddressExpr(const InlineAsmAddressExpr *Dst) {
     std::string Reg;
     if (tryEmitStmt(Reg, Dst->getSymbol()))
       return SYCLGenSuccess();
-    if (CurrInst->is(asmtok::op_prefetch, asmtok::op_prefetchu) ||
-        CanSuppressCast(Dst->getSymbol()))
+    if (CurrInst->is(asmtok::op_prefetch) || CanSuppressCast(Dst->getSymbol()))
       OS() << llvm::formatv("{0}", Reg);
     else
       OS() << llvm::formatv("(({0} *)(uintptr_t){1})", Type, Reg);
@@ -1289,9 +1288,7 @@ class SYCLGen : public SYCLGenBase {
 
     AsmStateSpace SS = Inst->getStateSpace();
     if (SS != AsmStateSpace::S_global && SS != AsmStateSpace::none) {
-      report(Diagnostics::API_NOT_MIGRATED, /*UseTextBegin=*/true,
-             GAS->getAsmString()->getString());
-      return SYCLGenSuccess();
+      return SYCLGenError();
     }
 
     if (!(Inst->hasAttr(InstAttr::L1) || Inst->hasAttr(InstAttr::L2)))
@@ -1310,17 +1307,15 @@ class SYCLGen : public SYCLGenBase {
       evictionHint = "evict_normal";
 
     if (!evictionHint.empty()) {
-      report(Diagnostics::API_NOT_MIGRATED, /*UseTextBegin=*/true,
-             GAS->getAsmString()->getString());
-      return SYCLGenSuccess();
+      return SYCLGenError();
     }
 
     llvm::SaveAndRestore<const InlineAsmInstruction *> Store(CurrInst);
     CurrInst = Inst;
     const auto *Src =
         dyn_cast_or_null<InlineAsmAddressExpr>(Inst->getInputOperand(0));
     if (!Src)
-      return false;
+      return SYCLGenError();
 
     OS() << MapNames::getExpNamespace() << "prefetch(";
     if (emitStmt(Src))
@@ -1333,15 +1328,6 @@ class SYCLGen : public SYCLGenBase {
     return SYCLGenSuccess();
   }
 
-  bool handle_prefetchu(const InlineAsmInstruction *Inst) override {
-    if (!DpctGlobalInfo::useExtPrefetch())
-      return SYCLGenError();
-
-    report(Diagnostics::API_NOT_MIGRATED, /*UseTextBegin=*/true,
-           GAS->getAsmString()->getString());
-    return SYCLGenSuccess();
-  }
-
   StringRef GetWiderTypeAsString(const InlineAsmBuiltinType *Type) const {
     switch (Type->getKind()) {
     case InlineAsmBuiltinType::s16:

clang/lib/DPCT/RulesAsm/Parser/AsmParser.cpp

@@ -383,8 +383,7 @@ InlineAsmStmtResult InlineAsmParser::ParseInstruction() {
     Out = nullptr;
   }
   // prefetch{.state}.{level} [%0] has only one input operand and no type.
-  if (Opcode->getTokenID() == asmtok::op_prefetch ||
-      Opcode->getTokenID() == asmtok::op_prefetchu) {
+  if (Opcode->getTokenID() == asmtok::op_prefetch) {
     Ops.push_back(Out.get());
     Out = nullptr;
     Types.push_back(Context.getBuiltinType(InlineAsmBuiltinType::byte));

clang/lib/DPCT/SrcAPI/APINames_ASM.inc

@@ -101,7 +101,6 @@ ENTRY("or", "or", true, NO_FLAG, P1, "Successful")
 ENTRY("pmevent", "pmevent", false, NO_FLAG, P1, "Comment")
 ENTRY("popc", "popc", true, NO_FLAG, P1, "Successful")
 ENTRY("prefetch", "prefetch", true, NO_FLAG, P1, "Partial")
-ENTRY("prefetchu", "prefetchu", true, NO_FLAG, P1, "Partial")
 ENTRY("prmt", "prmt", false, NO_FLAG, P1, "Comment")
 ENTRY("rcp", "rcp", true, NO_FLAG, P1, "Successful")
 ENTRY("red", "red", false, NO_FLAG, P1, "Comment")

clang/test/dpct/asm/prefetch.cu

@@ -40,28 +40,6 @@ __global__ void prefetch(int *arr) {
   /* using Register-Immediate (Displacement) address mode */
   // CHECK: sycl::ext::oneapi::experimental::prefetch(((uint8_t *)((uintptr_t)arr + 2)), sycl::ext::oneapi::experimental::properties{sycl::ext::oneapi::experimental::prefetch_hint_L1});
   asm volatile("prefetch.global.L1 [%0 + 2];" :: "l"(arr));
-
-#ifndef BUILD_TEST
-  /* prefetch of global address space with eviction priority */
-  // CHECK: /*
-  // CHECK-NEXT: DPCT1007:{{.*}} Migration of prefetch.global.L2::evict_last [%0]; is not supported.
-  // CHECK-NEXT: */
-  asm volatile ("prefetch.global.L2::evict_last [%0];" : : "l"(arr));
-  // CHECK: /*
-  // CHECK-NEXT: DPCT1007:{{.*}} Migration of prefetch.global.L2::evict_normal [%0]; is not supported.
-  // CHECK-NEXT: */
-  asm volatile ("prefetch.global.L2::evict_normal [%0];" : : "l"(arr));
-
-  /* prefetch of local address space */
-  // CHECK: /*
-  // CHECK-NEXT: DPCT1007:{{.*}} Migration of prefetch.local.L1 [%0]; is not supported.
-  // CHECK-NEXT: */
-  asm volatile ("prefetch.local.L1 [%0];" : : "l"(arr));
-  // CHECK: /*
-  // CHECK-NEXT: DPCT1007:{{.*}} Migration of prefetch.local.L2 [%0]; is not supported.
-  // CHECK-NEXT: */
-  asm volatile ("prefetch.local.L2 [%0];" : : "l"(arr));
-#endif // BUILD_TEST
 }
 
 // clang-format on