From 758c8fbf8f4d866e42ee9a3541153603c1ccf67d Mon Sep 17 00:00:00 2001 From: Teja Alaghari Date: Sat, 4 Jan 2025 23:53:36 +0800 Subject: [PATCH] Removed unsupported AI migration test cases --- clang/lib/DPCT/RulesAsm/AsmMigration.cpp | 24 ++++------------ clang/lib/DPCT/RulesAsm/Parser/AsmParser.cpp | 3 +- clang/lib/DPCT/SrcAPI/APINames_ASM.inc | 1 - clang/test/dpct/asm/prefetch.cu | 22 --------------- .../dpct/asm/prefetch_eviction_priority.cu | 28 ------------------- clang/test/dpct/asm/prefetch_tensormap.cu | 28 ------------------- clang/test/dpct/asm/prefetchu.cu | 23 --------------- clang/test/dpct/asm/prefetchu_default.cu | 23 --------------- 8 files changed, 6 insertions(+), 146 deletions(-) delete mode 100644 clang/test/dpct/asm/prefetch_eviction_priority.cu delete mode 100644 clang/test/dpct/asm/prefetch_tensormap.cu delete mode 100644 clang/test/dpct/asm/prefetchu.cu delete mode 100644 clang/test/dpct/asm/prefetchu_default.cu diff --git a/clang/lib/DPCT/RulesAsm/AsmMigration.cpp b/clang/lib/DPCT/RulesAsm/AsmMigration.cpp index 8a6be131f079..aec904a0efe6 100644 --- a/clang/lib/DPCT/RulesAsm/AsmMigration.cpp +++ b/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,9 +1307,7 @@ 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 Store(CurrInst); @@ -1320,7 +1315,7 @@ class SYCLGen : public SYCLGenBase { const auto *Src = dyn_cast_or_null(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: diff --git a/clang/lib/DPCT/RulesAsm/Parser/AsmParser.cpp b/clang/lib/DPCT/RulesAsm/Parser/AsmParser.cpp index 7a361649af80..a93115dc1899 100644 --- a/clang/lib/DPCT/RulesAsm/Parser/AsmParser.cpp +++ b/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)); diff --git a/clang/lib/DPCT/SrcAPI/APINames_ASM.inc b/clang/lib/DPCT/SrcAPI/APINames_ASM.inc index 4cf53303a53b..64b28b41015b 100644 --- a/clang/lib/DPCT/SrcAPI/APINames_ASM.inc +++ b/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") diff --git a/clang/test/dpct/asm/prefetch.cu b/clang/test/dpct/asm/prefetch.cu index 1c9afbd10ef8..833dc295e0fe 100644 --- a/clang/test/dpct/asm/prefetch.cu +++ b/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 diff --git a/clang/test/dpct/asm/prefetch_eviction_priority.cu b/clang/test/dpct/asm/prefetch_eviction_priority.cu deleted file mode 100644 index 3a102a336a4e..000000000000 --- a/clang/test/dpct/asm/prefetch_eviction_priority.cu +++ /dev/null @@ -1,28 +0,0 @@ -// UNSUPPORTED: cuda-8.0, cuda-9.0, cuda-9.1, cuda-9.2, cuda-10.0, cuda-10.1, cuda-10.2, cuda-11.0, cuda-11.1, cuda-11.2, cuda-11.3 -// UNSUPPORTED: v8.0, v9.0, v9.1, v9.2, v10.0, v10.1, v10.2, v11.0, v11.1, v11.2, v11.3 -// RUN: dpct --format-range=none -out-root %T/prefetch_eviction_priority %s --use-experimental-features=prefetch --cuda-include-path="%cuda-path/include" -- -std=c++14 -x cuda --cuda-host-only -// RUN: FileCheck %s --match-full-lines --input-file %T/prefetch_eviction_priority/prefetch_eviction_priority.dp.cpp -// RUN: %if BUILD_LIT %{icpx -c -DBUILD_TEST -fsycl %T/prefetch_eviction_priority/prefetch_eviction_priority.dp.cpp -o %T/prefetch_eviction_priority/prefetch_eviction_priority.dp.o %} - -// clang-format off -#include - -// Unsupported syntax: -// prefetch.global.level::eviction_priority [a]; // prefetch to data cache -// .level::eviction_priority = { .L2::evict_last, .L2::evict_normal }; - -__global__ void prefetch(int *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)); -#endif // BUILD_TEST -} - -// clang-format on diff --git a/clang/test/dpct/asm/prefetch_tensormap.cu b/clang/test/dpct/asm/prefetch_tensormap.cu deleted file mode 100644 index c6d37bfc6866..000000000000 --- a/clang/test/dpct/asm/prefetch_tensormap.cu +++ /dev/null @@ -1,28 +0,0 @@ -// UNSUPPORTED: cuda-8.0, cuda-9.0, cuda-9.1, cuda-9.2, cuda-10.0, cuda-10.1, cuda-10.2, cuda-11.0, cuda-11.1, cuda-11.2, cuda-11.3, cuda-11.4, cuda-11.5, cuda-11.6, cuda-11.7 -// UNSUPPORTED: v8.0, v9.0, v9.1, v9.2, v10.0, v10.1, v10.2, v11.0, v11.1, v11.2, v11.3, v11.4, v11.5, v11.6, v11.7 -// RUN: dpct --format-range=none -out-root %T/prefetch_tensormap %s --use-experimental-features=prefetch --cuda-include-path="%cuda-path/include" -- -std=c++14 -x cuda --cuda-host-only -// RUN: FileCheck %s --match-full-lines --input-file %T/prefetch_tensormap/prefetch_tensormap.dp.cpp -// RUN: %if BUILD_LIT %{icpx -c -DBUILD_TEST -fsycl %T/prefetch_tensormap/prefetch_tensormap.dp.cpp -o %T/prefetch_tensormap/prefetch_tensormap.dp.o %} - -// clang-format off -#include - -// Unsupported syntax: -// prefetch{.tensormap_space}.tensormap [a]; // prefetch the tensormap -// .tensormap_space = { .const, .param }; - -__global__ void prefetch_tensormap(int *arr) { -#ifndef BUILD_TEST - /* prefetch of tensormap space */ - // CHECK: /* - // CHECK-NEXT: DPCT1007:{{.*}} Migration of prefetch.const.tensormap [%0]; is not supported. - // CHECK-NEXT: */ - asm volatile ("prefetch.const.tensormap [%0];" : : "l"(arr)); - // CHECK: /* - // CHECK-NEXT: DPCT1007:{{.*}} Migration of prefetch.param.tensormap [%0]; is not supported. - // CHECK-NEXT: */ - asm volatile ("prefetch.param.tensormap [%0];" : : "l"(arr)); -#endif // BUILD_TEST -} - -// clang-format on diff --git a/clang/test/dpct/asm/prefetchu.cu b/clang/test/dpct/asm/prefetchu.cu deleted file mode 100644 index 4952dd75641f..000000000000 --- a/clang/test/dpct/asm/prefetchu.cu +++ /dev/null @@ -1,23 +0,0 @@ -// UNSUPPORTED: cuda-8.0 -// UNSUPPORTED: v8.0 -// RUN: dpct --format-range=none -out-root %T/prefetchu %s --use-experimental-features=prefetch --cuda-include-path="%cuda-path/include" -- -std=c++14 -x cuda --cuda-host-only -// RUN: FileCheck %s --match-full-lines --input-file %T/prefetchu/prefetchu.dp.cpp -// RUN: %if BUILD_LIT %{icpx -c -DBUILD_TEST -fsycl %T/prefetchu/prefetchu.dp.cpp -o %T/prefetchu/prefetchu.dp.o %} - -// clang-format off -#include - -// Unsupported syntax: -// prefetchu.L1 [a]; // prefetch to uniform cache - -__global__ void prefetchu(int *arr) { -#ifndef BUILD_TEST - /* prefetch of uniform address space */ - // CHECK: /* - // CHECK-NEXT: DPCT1007:{{.*}} Migration of prefetchu.L1 [%0]; is not supported. - // CHECK-NEXT: */ - asm volatile ("prefetchu.L1 [%0];" : : "l"(arr)); -#endif // BUILD_TEST -} - -// clang-format on diff --git a/clang/test/dpct/asm/prefetchu_default.cu b/clang/test/dpct/asm/prefetchu_default.cu deleted file mode 100644 index b506a916487d..000000000000 --- a/clang/test/dpct/asm/prefetchu_default.cu +++ /dev/null @@ -1,23 +0,0 @@ -// UNSUPPORTED: cuda-8.0 -// UNSUPPORTED: v8.0 -// RUN: dpct --format-range=none -out-root %T/prefetchu_default %s --cuda-include-path="%cuda-path/include" -- -std=c++14 -x cuda --cuda-host-only -// RUN: FileCheck %s --match-full-lines --input-file %T/prefetchu_default/prefetchu_default.dp.cpp -// RUN: %if BUILD_LIT %{icpx -c -DBUILD_TEST -fsycl %T/prefetchu_default/prefetchu_default.dp.cpp -o %T/prefetchu_default/prefetchu_default.dp.o %} - -// clang-format off -#include - -// Unsupported syntax: -// prefetchu.L1 [a]; // prefetch to uniform cache - -__global__ void prefetchu(int *arr) { -#ifndef BUILD_TEST - /* prefetch of uniform address space */ - // CHECK: /* - // CHECK-NEXT: DPCT1053:{{.*}} Migration of device assembly code is not supported. - // CHECK-NEXT: */ - asm volatile ("prefetchu.L1 [%0];" : : "l"(arr)); -#endif // BUILD_TEST -} - -// clang-format on