Removed unsupported AI migration test cases

Author: TejaX-Alaghari

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)))
@@ -1303,24 +1300,12 @@ class SYCLGen : public SYCLGenBase {
     else if (Inst->hasAttr(InstAttr::L2))
       PrefetchHint = "L2";
 
-    std::string evictionHint = "";
-    if (Inst->hasAttr(InstAttr::evict_last))
-      evictionHint = "evict_last";
-    else if (Inst->hasAttr(InstAttr::evict_normal))
-      evictionHint = "evict_normal";
-
-    if (!evictionHint.empty()) {
-      report(Diagnostics::API_NOT_MIGRATED, /*UseTextBegin=*/true,
-             GAS->getAsmString()->getString());
-      return SYCLGenSuccess();
-    }
-
     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 +1318,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/AsmLexer.cpp

@@ -265,11 +265,6 @@ bool InlineAsmLexer::lexTokenInternal(InlineAsmToken &Result) {
     Char = getChar(CurPtr);
     if (Char == ':') {
       CurPtr++;
-      Char = getChar(CurPtr);
-      if (isIdentifierStart(Char) || isDigit(Char)) {
-        Result.setFlag(InlineAsmToken::StartOfColonColon);
-        return lexIdentifierContinue(Result, CurPtr);
-      }
       Kind = asmtok::coloncolon;
     } else {
       Kind = asmtok::colon;

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

@@ -335,7 +335,7 @@ InlineAsmStmtResult InlineAsmParser::ParseInstruction() {
   SmallVector<InlineAsmExpr *, 4> Ops;
   std::optional<AsmStateSpace> StateSpace;
 
-  while (Tok.startOfDot() || Tok.startOfColonColon()) {
+  while (Tok.startOfDot()) {
     switch (Tok.getIdentifier()->getFlags()) {
     case InlineAsmIdentifierInfo::BuiltinType:
       Types.push_back(Context.getBuiltinTypeFromTokenKind(Tok.getKind()));
@@ -356,11 +356,6 @@ InlineAsmStmtResult InlineAsmParser::ParseInstruction() {
     ConsumeToken(); // consume instruction attribute
   }
 
-  if (Tok.getKind() == asmtok::coloncolon) {
-    Attrs.push_back(ConvertToInstAttr(Tok.getKind()));
-    ConsumeToken();
-  }
-
   InlineAsmExprResult Pred, Out;
   if ((Out = ParseExpression()).isInvalid())
     return AsmStmtError();
@@ -383,8 +378,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/RulesAsm/Parser/AsmTokenKinds.def

@@ -406,9 +406,6 @@ MODIFIER(sc, ".sc")
 MODIFIER(gl, ".gl")
 MODIFIER(L1, ".L1")
 MODIFIER(L2, ".L2")
-MODIFIER(tensormap, ".tensormap")
-MODIFIER(evict_last, "::evict_last")
-MODIFIER(evict_normal, "::evict_normal")
 
 
 #undef LINKAGE

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