diff --git a/clang/lib/Sema/SemaExpr.cpp b/clang/lib/Sema/SemaExpr.cpp index aad43b7560c150..ad6e7183cb3a42 100644 --- a/clang/lib/Sema/SemaExpr.cpp +++ b/clang/lib/Sema/SemaExpr.cpp @@ -13874,20 +13874,6 @@ inline QualType Sema::CheckBitwiseOperands(ExprResult &LHS, ExprResult &RHS, return ResultTy; } -static inline bool IsAMDGPUPredicateBI(Expr *E) { - if (!E->getType()->isVoidType()) - return false; - - if (auto *CE = dyn_cast(E)) { - if (auto *BI = CE->getDirectCallee()) - if (BI->getName() == "__builtin_amdgcn_processor_is" || - BI->getName() == "__builtin_amdgcn_is_invocable") - return true; - } - - return false; -} - // C99 6.5.[13,14] inline QualType Sema::CheckLogicalOperands(ExprResult &LHS, ExprResult &RHS, SourceLocation Loc, @@ -13992,9 +13978,6 @@ inline QualType Sema::CheckLogicalOperands(ExprResult &LHS, ExprResult &RHS, // The following is safe because we only use this method for // non-overloadable operands. - if (IsAMDGPUPredicateBI(LHS.get()) && IsAMDGPUPredicateBI(RHS.get())) - return Context.VoidTy; - // C++ [expr.log.and]p1 // C++ [expr.log.or]p1 // The operands are both contextually converted to type bool. diff --git a/clang/test/CodeGenHipStdPar/select-accelerator-code-pass-ordering.cpp b/clang/test/CodeGenHipStdPar/select-accelerator-code-pass-ordering.cpp index f03b783199a413..44557284fc5811 100644 --- a/clang/test/CodeGenHipStdPar/select-accelerator-code-pass-ordering.cpp +++ b/clang/test/CodeGenHipStdPar/select-accelerator-code-pass-ordering.cpp @@ -7,7 +7,6 @@ // HIPSTDPAR-PRE: Running pass: EntryExitInstrumenterPass // HIPSTDPAR-PRE-NEXT: Running pass: EntryExitInstrumenterPass // HIPSTDPAR-PRE-NOT: Running pass: HipStdParAcceleratorCodeSelectionPass -// HIPSTDPAR-PRE-NEXT: Running pass: AMDGPUExpandFeaturePredicatesPass // HIPSTDPAR-PRE-NEXT: Running pass: AlwaysInlinerPass // Ensure Pass HipStdParAcceleratorCodeSelectionPass is invoked in PostLink. diff --git a/clang/test/SemaHIP/amdgpu-is-invocable.hip b/clang/test/SemaHIP/amdgpu-is-invocable.hip deleted file mode 100644 index 214d7769a595f9..00000000000000 --- a/clang/test/SemaHIP/amdgpu-is-invocable.hip +++ /dev/null @@ -1,21 +0,0 @@ -// REQUIRES: amdgpu-registered-target -// REQUIRES: spirv-registered-target -// RUN: %clang_cc1 -fsyntax-only -verify -triple amdgcn -Wno-unused-value %s -// RUN: %clang_cc1 -fsyntax-only -verify -triple spirv64-amd-amdhsa -Wno-unused-value %s -// RUN: %clang_cc1 -fsyntax-only -verify -triple x86_64 -aux-triple amdgcn -Wno-unused-value %s -// RUN: %clang_cc1 -fsyntax-only -verify -triple x86_64 -aux-triple spirv64-amd-amdhsa -Wno-unused-value %s - -// expected-no-diagnostics - -#define __device__ __attribute__((device)) -#define __global__ __attribute__((global)) - -__device__ void foo() { - if (__builtin_amdgcn_is_invocable(__builtin_amdgcn_permlanex16)) - return __builtin_trap(); -} - -__global__ void bar() { - if (__builtin_amdgcn_is_invocable(__builtin_amdgcn_permlanex16)) - return __builtin_trap(); -} diff --git a/clang/test/SemaHIP/amdgpu-processor-is.hip b/clang/test/SemaHIP/amdgpu-processor-is.hip deleted file mode 100644 index 0f7211fd75d900..00000000000000 --- a/clang/test/SemaHIP/amdgpu-processor-is.hip +++ /dev/null @@ -1,21 +0,0 @@ -// REQUIRES: amdgpu-registered-target -// REQUIRES: spirv-registered-target -// RUN: %clang_cc1 -fsyntax-only -verify -triple amdgcn -Wno-unused-value %s -// RUN: %clang_cc1 -fsyntax-only -verify -triple spirv64-amd-amdhsa -Wno-unused-value %s -// RUN: %clang_cc1 -fsyntax-only -verify -triple x86_64 -aux-triple amdgcn -Wno-unused-value %s -// RUN: %clang_cc1 -fsyntax-only -verify -triple x86_64 -aux-triple spirv64-amd-amdhsa -Wno-unused-value %s - -// expected-no-diagnostics - -#define __device__ __attribute__((device)) -#define __global__ __attribute__((global)) - -__device__ void foo() { - if (__builtin_amdgcn_processor_is("gfx900")) - return __builtin_trap(); -} - -__global__ void bar() { - if (__builtin_amdgcn_processor_is("gfx900")) - return __builtin_trap(); -} diff --git a/llvm/lib/Target/AMDGPU/AMDGPU.h b/llvm/lib/Target/AMDGPU/AMDGPU.h index 0ee44a24efd115..c6dd1dbb624499 100644 --- a/llvm/lib/Target/AMDGPU/AMDGPU.h +++ b/llvm/lib/Target/AMDGPU/AMDGPU.h @@ -483,15 +483,6 @@ extern char &AMDGPUPrintfRuntimeBindingID; void initializeAMDGPUResourceUsageAnalysisWrapperPassPass(PassRegistry &); extern char &AMDGPUResourceUsageAnalysisID; -struct AMDGPUExpandFeaturePredicatesPass - : PassInfoMixin { - const AMDGPUTargetMachine &TM; - AMDGPUExpandFeaturePredicatesPass(const AMDGPUTargetMachine &ATM) : TM(ATM) {} - PreservedAnalyses run(Module &M, ModuleAnalysisManager &AM); - - static bool isRequired() { return true; } -}; - struct AMDGPUPrintfRuntimeBindingPass : OptionalPassInfoMixin { PreservedAnalyses run(Module &M, ModuleAnalysisManager &AM); diff --git a/llvm/lib/Target/AMDGPU/AMDGPUExpandFeaturePredicates.cpp b/llvm/lib/Target/AMDGPU/AMDGPUExpandFeaturePredicates.cpp deleted file mode 100644 index a0b2a08c76ae24..00000000000000 --- a/llvm/lib/Target/AMDGPU/AMDGPUExpandFeaturePredicates.cpp +++ /dev/null @@ -1,170 +0,0 @@ -//===- AMDGPUExpandFeaturePredicates.cpp - Feature Predicate Expander Pass ===// -// -// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. -// See https://llvm.org/LICENSE.txt for license information. -// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception -// -//===----------------------------------------------------------------------===// -// This file implements a pass that deals with expanding AMDGCN generic feature -// predicates into target specific quantities / sequences. In this context, a -// generic feature predicate is an implementation detail global variable that -// is inserted by the FE as a consequence of using either the __builtin_cpu_is -// or the __builtin_amdgcn_is_invocable special builtins on an abstract target -// (AMDGCNSPIRV). These placeholder globals are used to guide target specific -// lowering, once the concrete target is known, by way of constant folding their -// value all the way into a terminator (i.e. a controlled block) or into a no -// live use scenario. We hard fail if the folding fails, to avoid obtuse BE -// errors or opaque run time errors. This pass should run as early as possible / -// immediately after Clang CodeGen, so that the optimisation pipeline and the BE -// operate with concrete target data. -//===----------------------------------------------------------------------===// - -#include "AMDGPU.h" -#include "AMDGPUTargetMachine.h" -#include "GCNSubtarget.h" - -#include "llvm/ADT/SmallPtrSet.h" -#include "llvm/ADT/SmallVector.h" -#include "llvm/ADT/StringRef.h" -#include "llvm/Analysis/ConstantFolding.h" -#include "llvm/IR/Constants.h" -#include "llvm/IR/Function.h" -#include "llvm/IR/Module.h" -#include "llvm/Pass.h" -#include "llvm/Passes/CodeGenPassBuilder.h" -#include "llvm/Transforms/Utils/Local.h" - -#include -#include - -using namespace llvm; - -namespace { -template void collectUsers(Value *V, C &Container) { - assert(V && "Must pass an existing Value!"); - - for (auto &&U : V->users()) - if (auto *I = dyn_cast(U)) - Container.insert(Container.end(), I); -} - -inline void setPredicate(const GCNSubtarget &ST, GlobalVariable *P) { - const bool IsFeature = P->getName().starts_with("llvm.amdgcn.has"); - const size_t Offset = - IsFeature ? sizeof("llvm.amdgcn.has") : sizeof("llvm.amdgcn.is"); - - std::string PV = P->getName().substr(Offset).str(); - const bool HasFeatureList = IsFeature && !PV.empty(); - if (HasFeatureList) { - size_t Dx = PV.find(','); - while (Dx != std::string::npos) { - PV.insert(++Dx, {'+'}); - - Dx = PV.find(',', Dx); - } - PV.insert(PV.cbegin(), '+'); - } - - Type *PTy = P->getValueType(); - P->setLinkage(GlobalValue::PrivateLinkage); - P->setExternallyInitialized(false); - - if (IsFeature) - P->setInitializer( - ConstantInt::getBool(PTy, !HasFeatureList || ST.checkFeatures(PV))); - else - P->setInitializer(ConstantInt::getBool(PTy, PV == ST.getCPU())); -} - -std::pair -unfoldableFound(Function *Caller, GlobalVariable *P, Instruction *NoFold) { - std::string W; - raw_string_ostream OS(W); - - OS << "Impossible to constant fold feature predicate: " << *P << " used by " - << *NoFold << ", please simplify.\n"; - - Caller->getContext().diagnose( - DiagnosticInfoUnsupported(*Caller, W, NoFold->getDebugLoc(), DS_Error)); - - return {PreservedAnalyses::none(), false}; -} - -std::pair -handlePredicate(const GCNSubtarget &ST, FunctionAnalysisManager &FAM, - SmallPtrSet &Predicated, GlobalVariable *P) { - setPredicate(ST, P); - - SmallPtrSet ToFold; - collectUsers(P, ToFold); - - if (ToFold.empty()) - return {PreservedAnalyses::all(), true}; - - do { - Instruction *I = *ToFold.begin(); - ToFold.erase(I); - - I->dropDroppableUses(); - - Function *F = I->getParent()->getParent(); - auto &DT = FAM.getResult(*F); - DomTreeUpdater DTU(DT, DomTreeUpdater::UpdateStrategy::Lazy); - - if (auto *C = ConstantFoldInstruction(I, P->getDataLayout())) { - collectUsers(I, ToFold); - I->replaceAllUsesWith(C); - I->eraseFromParent(); - continue; - } else if (I->isTerminator() && - ConstantFoldTerminator(I->getParent(), true, nullptr, &DTU)) { - Predicated.insert(F); - - continue; - } - - return unfoldableFound(I->getParent()->getParent(), P, I); - } while (!ToFold.empty()); - - return {PreservedAnalyses::none(), true}; -} -} // Unnamed namespace. - -PreservedAnalyses -AMDGPUExpandFeaturePredicatesPass::run(Module &M, ModuleAnalysisManager &MAM) { - if (M.empty()) - return PreservedAnalyses::all(); - - SmallVector Predicates; - for (auto &&G : M.globals()) { - if (!G.isDeclaration() || !G.hasName()) - continue; - if (G.getName().starts_with("llvm.amdgcn.")) - Predicates.push_back(&G); - } - - if (Predicates.empty()) - return PreservedAnalyses::all(); - - const auto &ST = TM.getSubtarget( - *find_if(M, [](auto &&F) { return !F.isIntrinsic(); })); - - auto &FAM = MAM.getResult(M).getManager(); - SmallPtrSet Predicated; - auto Ret = PreservedAnalyses::all(); - for (auto &&P : Predicates) { - auto R = handlePredicate(ST, FAM, Predicated, P); - - if (!R.second) - break; - - Ret.intersect(R.first); - } - - for (auto &&P : Predicates) - P->eraseFromParent(); - for (auto &&F : Predicated) - removeUnreachableBlocks(*F); - - return Ret; -} diff --git a/llvm/lib/Target/AMDGPU/AMDGPUPassRegistry.def b/llvm/lib/Target/AMDGPU/AMDGPUPassRegistry.def index e3ec6bbaae37ce..2a6560b309e622 100644 --- a/llvm/lib/Target/AMDGPU/AMDGPUPassRegistry.def +++ b/llvm/lib/Target/AMDGPU/AMDGPUPassRegistry.def @@ -16,8 +16,6 @@ #ifndef MODULE_PASS #define MODULE_PASS(NAME, CREATE_PASS) #endif -MODULE_PASS("amdgpu-expand-feature-predicates", - AMDGPUExpandFeaturePredicatesPass(*this)) MODULE_PASS("amdgpu-always-inline", AMDGPUAlwaysInlinePass()) MODULE_PASS("amdgpu-export-kernel-runtime-handles", AMDGPUExportKernelRuntimeHandlesPass()) MODULE_PASS("amdgpu-lower-buffer-fat-pointers", @@ -33,8 +31,6 @@ MODULE_PASS("amdgpu-printf-runtime-binding", AMDGPUPrintfRuntimeBindingPass()) MODULE_PASS("amdgpu-remove-incompatible-functions", AMDGPURemoveIncompatibleFunctionsPass(*this)) MODULE_PASS("amdgpu-lower-exec-sync", AMDGPULowerExecSyncPass()) MODULE_PASS("amdgpu-sw-lower-lds", AMDGPUSwLowerLDSPass(*this)) -MODULE_PASS("amdgpu-expand-feature-predicates", - AMDGPUExpandFeaturePredicatesPass(*this)) #undef MODULE_PASS #ifndef MODULE_PASS_WITH_PARAMS diff --git a/llvm/lib/Target/AMDGPU/AMDGPUTargetMachine.cpp b/llvm/lib/Target/AMDGPU/AMDGPUTargetMachine.cpp index 57efcf0ed2e467..aa76b346db7838 100644 --- a/llvm/lib/Target/AMDGPU/AMDGPUTargetMachine.cpp +++ b/llvm/lib/Target/AMDGPU/AMDGPUTargetMachine.cpp @@ -969,16 +969,6 @@ void AMDGPUTargetMachine::registerPassBuilderCallbacks(PassBuilder &PB) { #define GET_PASS_REGISTRY "AMDGPUPassRegistry.def" #include "llvm/Passes/TargetPassRegistry.inc" - PB.registerPipelineStartEPCallback( - [this](ModulePassManager &PM, OptimizationLevel Level) { - PM.addPass(AMDGPUExpandFeaturePredicatesPass(*this)); - }); - - PB.registerFullLinkTimeOptimizationEarlyEPCallback( - [this](ModulePassManager &PM, OptimizationLevel Level) { - PM.addPass(AMDGPUExpandFeaturePredicatesPass(*this)); - }); - PB.registerPipelineParsingCallback( [this](StringRef Name, CGSCCPassManager &PM, ArrayRef Pipeline) { diff --git a/llvm/lib/Target/AMDGPU/CMakeLists.txt b/llvm/lib/Target/AMDGPU/CMakeLists.txt index 392d561fe3237f..ae8f1c0fad5bae 100644 --- a/llvm/lib/Target/AMDGPU/CMakeLists.txt +++ b/llvm/lib/Target/AMDGPU/CMakeLists.txt @@ -54,7 +54,6 @@ add_llvm_target(AMDGPUCodeGen AMDGPUCodeGenPrepare.cpp AMDGPUCombinerHelper.cpp AMDGPUCtorDtorLowering.cpp - AMDGPUExpandFeaturePredicates.cpp AMDGPUExportClustering.cpp AMDGPUExportKernelRuntimeHandles.cpp AMDGPUFrameLowering.cpp diff --git a/llvm/test/CodeGen/AMDGPU/amdgpu-expand-feature-predicates-empty-feature.ll b/llvm/test/CodeGen/AMDGPU/amdgpu-expand-feature-predicates-empty-feature.ll deleted file mode 100644 index 87b452b5df5344..00000000000000 --- a/llvm/test/CodeGen/AMDGPU/amdgpu-expand-feature-predicates-empty-feature.ll +++ /dev/null @@ -1,99 +0,0 @@ -; NOTE: Assertions have been autogenerated by utils/update_test_checks.py UTC_ARGS: --version 5 -; REQUIRES: amdgpu-registered-target - -; Predicate guarding a builtin with no required target features -; (e.g. `__builtin_amdgcn_is_invocable(__builtin_amdgcn_fence)`) is -; encoded as a global named `llvm.amdgcn.has.` with an empty feature suffix. - -; RUN: opt -S -mtriple=amdgcn-amd-amdhsa -mcpu=gfx906 -passes='amdgpu-expand-feature-predicates' %s -o - | FileCheck --check-prefix=GFX906 %s -; RUN: opt -S -mtriple=amdgcn-amd-amdhsa -mcpu=gfx950 -passes='amdgpu-expand-feature-predicates' %s -o - | FileCheck --check-prefix=GFX950 %s -; RUN: opt -S -mtriple=amdgcn-amd-amdhsa -mcpu=gfx1250 -passes='amdgpu-expand-feature-predicates' %s -o - | FileCheck --check-prefix=GFX1250 %s - -@llvm.amdgcn.has. = external addrspace(1) externally_initialized constant i1 -@llvm.amdgcn.has.gfx1250-insts = external addrspace(1) externally_initialized constant i1 - -declare void @llvm.amdgcn.s.wait.asynccnt(i16 immarg) - -define amdgpu_kernel void @always_invocable(ptr addrspace(1) %p) { -; GFX906-LABEL: define amdgpu_kernel void @always_invocable( -; GFX906-SAME: ptr addrspace(1) [[P:%.*]]) #[[ATTR1:[0-9]+]] { -; GFX906-NEXT: [[ENTRY:.*:]] -; GFX906-NEXT: br label %[[THEN:.*]] -; GFX906: [[THEN]]: -; GFX906-NEXT: store i32 1, ptr addrspace(1) [[P]], align 4 -; GFX906-NEXT: br label %[[SKIP:.*]] -; GFX906: [[SKIP]]: -; GFX906-NEXT: ret void -; -; GFX950-LABEL: define amdgpu_kernel void @always_invocable( -; GFX950-SAME: ptr addrspace(1) [[P:%.*]]) #[[ATTR1:[0-9]+]] { -; GFX950-NEXT: [[ENTRY:.*:]] -; GFX950-NEXT: br label %[[THEN:.*]] -; GFX950: [[THEN]]: -; GFX950-NEXT: store i32 1, ptr addrspace(1) [[P]], align 4 -; GFX950-NEXT: br label %[[SKIP:.*]] -; GFX950: [[SKIP]]: -; GFX950-NEXT: ret void -; -; GFX1250-LABEL: define amdgpu_kernel void @always_invocable( -; GFX1250-SAME: ptr addrspace(1) [[P:%.*]]) #[[ATTR1:[0-9]+]] { -; GFX1250-NEXT: [[ENTRY:.*:]] -; GFX1250-NEXT: br label %[[THEN:.*]] -; GFX1250: [[THEN]]: -; GFX1250-NEXT: store i32 1, ptr addrspace(1) [[P]], align 4 -; GFX1250-NEXT: br label %[[SKIP:.*]] -; GFX1250: [[SKIP]]: -; GFX1250-NEXT: ret void -; -entry: - %has.empty = load i1, ptr addrspace(1) @llvm.amdgcn.has., align 1 - br i1 %has.empty, label %then, label %skip - -then: - store i32 1, ptr addrspace(1) %p, align 4 - br label %skip - -skip: - ret void -} - -define amdgpu_kernel void @gated_asynccnt(ptr addrspace(1) %p) { -; GFX906-LABEL: define amdgpu_kernel void @gated_asynccnt( -; GFX906-SAME: ptr addrspace(1) [[P:%.*]]) #[[ATTR1]] { -; GFX906-NEXT: [[ENTRY:.*:]] -; GFX906-NEXT: br label %[[SKIP:.*]] -; GFX906: [[SKIP]]: -; GFX906-NEXT: store i32 2, ptr addrspace(1) [[P]], align 4 -; GFX906-NEXT: ret void -; -; GFX950-LABEL: define amdgpu_kernel void @gated_asynccnt( -; GFX950-SAME: ptr addrspace(1) [[P:%.*]]) #[[ATTR1]] { -; GFX950-NEXT: [[ENTRY:.*:]] -; GFX950-NEXT: br label %[[SKIP:.*]] -; GFX950: [[SKIP]]: -; GFX950-NEXT: store i32 2, ptr addrspace(1) [[P]], align 4 -; GFX950-NEXT: ret void -; -; GFX1250-LABEL: define amdgpu_kernel void @gated_asynccnt( -; GFX1250-SAME: ptr addrspace(1) [[P:%.*]]) #[[ATTR1]] { -; GFX1250-NEXT: [[ENTRY:.*:]] -; GFX1250-NEXT: br label %[[WAIT:.*]] -; GFX1250: [[WAIT]]: -; GFX1250-NEXT: call void @llvm.amdgcn.s.wait.asynccnt(i16 0) -; GFX1250-NEXT: br label %[[SKIP:.*]] -; GFX1250: [[SKIP]]: -; GFX1250-NEXT: store i32 2, ptr addrspace(1) [[P]], align 4 -; GFX1250-NEXT: ret void -; -entry: - %has.gfx1250 = load i1, ptr addrspace(1) @llvm.amdgcn.has.gfx1250-insts, align 1 - br i1 %has.gfx1250, label %wait, label %skip - -wait: - call void @llvm.amdgcn.s.wait.asynccnt(i16 0) - br label %skip - -skip: - store i32 2, ptr addrspace(1) %p, align 4 - ret void -} diff --git a/llvm/test/CodeGen/AMDGPU/amdgpu-expand-feature-predicates-unfoldable.ll b/llvm/test/CodeGen/AMDGPU/amdgpu-expand-feature-predicates-unfoldable.ll deleted file mode 100644 index bfc35d8c76e372..00000000000000 --- a/llvm/test/CodeGen/AMDGPU/amdgpu-expand-feature-predicates-unfoldable.ll +++ /dev/null @@ -1,28 +0,0 @@ -; REQUIRES: amdgpu-registered-target - -; RUN: not opt -mtriple=amdgcn-amd-amdhsa -mcpu=gfx906 -passes='amdgpu-expand-feature-predicates' < %s 2>&1 | FileCheck %s - -; CHECK: error:{{.*}}in function kernel void (ptr addrspace(1), i32, ptr addrspace(1)): Impossible to constant fold feature predicate: @llvm.amdgcn.is.gfx803 = private addrspace(1) constant i1 false used by %call = call i1 %1(i1 zeroext false), please simplify. - -@llvm.amdgcn.is.gfx803 = external addrspace(1) externally_initialized constant i1 - -declare void @llvm.amdgcn.s.sleep(i32 immarg) #1 - -define amdgpu_kernel void @kernel(ptr addrspace(1) readnone captures(none) %p.coerce, i32 %x, ptr addrspace(1) %pfn.coerce) { -entry: - %0 = ptrtoint ptr addrspace(1) %pfn.coerce to i64 - %1 = inttoptr i64 %0 to ptr - %2 = ptrtoint ptr addrspace(1) %pfn.coerce to i64 - %3 = load i1, ptr addrspace(1) @llvm.amdgcn.is.gfx803, align 1 - %call = call i1 %1(i1 zeroext %3) - br i1 %call, label %if.gfx803, label %if.end - -if.gfx803: - call void @llvm.amdgcn.s.sleep(i32 0) - br label %if.end - -if.end: - ret void -} - -attributes #1 = { nocallback nofree nosync nounwind willreturn } diff --git a/llvm/test/CodeGen/AMDGPU/amdgpu-expand-feature-predicates.ll b/llvm/test/CodeGen/AMDGPU/amdgpu-expand-feature-predicates.ll deleted file mode 100644 index a16a7fc31da222..00000000000000 --- a/llvm/test/CodeGen/AMDGPU/amdgpu-expand-feature-predicates.ll +++ /dev/null @@ -1,284 +0,0 @@ -; NOTE: Assertions have been autogenerated by utils/update_test_checks.py UTC_ARGS: --version 5 -; REQUIRES: amdgpu-registered-target - -; RUN: opt -S -mtriple=amdgcn-amd-amdhsa -mcpu=gfx906 -passes='amdgpu-expand-feature-predicates' %s -o - | FileCheck --check-prefix=GFX906 %s -; RUN: opt -S -mtriple=amdgcn-amd-amdhsa -mcpu=gfx1010 -passes='amdgpu-expand-feature-predicates' %s -o - | FileCheck --check-prefix=GFX1010 %s -; RUN: opt -S -mtriple=amdgcn-amd-amdhsa -mcpu=gfx1101 -passes='amdgpu-expand-feature-predicates' %s -o - | FileCheck --check-prefix=GFX1101 %s -; RUN: opt -S -mtriple=amdgcn-amd-amdhsa -mcpu=gfx1201 -passes='amdgpu-expand-feature-predicates' %s -o - | FileCheck --check-prefix=GFX1201 %s -; RUN: opt -S -mtriple=amdgcn-amd-amdhsa -mcpu=gfx1201 -mattr=+wavefrontsize64 -passes='amdgpu-expand-feature-predicates' %s -o - | FileCheck --check-prefix=GFX1201-W64 %s - -;; The IR was derived from the following source: -;; extern "C" __global__ void kernel(int* p, int x) -;; { -;; if (__builtin_amdgcn_processor_is("gfx1201") || -;; __builtin_amdgcn_is_invocable(__builtin_amdgcn_s_sleep_var)) -;; __builtin_amdgcn_s_sleep_var(x); -;; if (!__builtin_amdgcn_processor_is("gfx906")) -;; __builtin_amdgcn_s_wait_event_export_ready(); -;; else if (__builtin_amdgcn_processor_is("gfx1010") || -;; __builtin_amdgcn_processor_is("gfx1101")) -;; __builtin_amdgcn_s_ttracedata_imm(1); -;; while (__builtin_amdgcn_processor_is("gfx1101")) *p += x; -;; do { -;; *p -= x; -;; } while (__builtin_amdgcn_processor_is("gfx1010")); -;; for (; __builtin_amdgcn_processor_is("gfx1201"); ++*p) break; -;; -;; if (__builtin_amdgcn_is_invocable(__builtin_amdgcn_s_wait_event_export_ready)) -;; __builtin_amdgcn_s_wait_event_export_ready(); -;; else if (__builtin_amdgcn_is_invocable(__builtin_amdgcn_s_ttracedata_imm)) -;; __builtin_amdgcn_s_ttracedata_imm(1); -;; -;; do { -;; *p -= x; -;; } while (__builtin_amdgcn_is_invocable(__builtin_amdgcn_global_load_tr_b64_i32)); -;; for (; __builtin_amdgcn_is_invocable(__builtin_amdgcn_permlane64); ++*p) break; -;; } - -@llvm.amdgcn.is.gfx1201 = external addrspace(1) externally_initialized constant i1 -@llvm.amdgcn.has.gfx12-insts = external addrspace(1) externally_initialized constant i1 -@llvm.amdgcn.is.gfx906 = external addrspace(1) externally_initialized constant i1 -@llvm.amdgcn.is.gfx1010 = external addrspace(1) externally_initialized constant i1 -@llvm.amdgcn.is.gfx1101 = external addrspace(1) externally_initialized constant i1 -@llvm.amdgcn.has.gfx11-insts = external addrspace(1) externally_initialized constant i1 -@llvm.amdgcn.has.gfx10-insts = external addrspace(1) externally_initialized constant i1 -@"llvm.amdgcn.has.gfx12-insts,wavefrontsize64" = external addrspace(1) externally_initialized constant i1 - -declare void @llvm.amdgcn.s.sleep.var(i32) -declare void @llvm.amdgcn.s.wait.event.export.ready() -declare void @llvm.amdgcn.s.ttracedata.imm(i16 immarg) - -define amdgpu_kernel void @kernel(ptr addrspace(1) %p.coerce, i32 %x) { -; GFX906-LABEL: define amdgpu_kernel void @kernel( -; GFX906-SAME: ptr addrspace(1) [[P_COERCE:%.*]], i32 [[X:%.*]]) #[[ATTR2:[0-9]+]] { -; GFX906-NEXT: [[ENTRY:.*:]] -; GFX906-NEXT: [[TMP0:%.*]] = ptrtoint ptr addrspace(1) [[P_COERCE]] to i64 -; GFX906-NEXT: [[TMP1:%.*]] = inttoptr i64 [[TMP0]] to ptr -; GFX906-NEXT: br label %[[IF_GFX1201_OR_GFX12_INSTS:.*]] -; GFX906: [[IF_GFX1201_OR_GFX12_INSTS]]: -; GFX906-NEXT: br label %[[IF_NOT_GFX907:.*]] -; GFX906: [[IF_NOT_GFX907]]: -; GFX906-NEXT: br label %[[IF_GFX1010_OR_GFX1101:.*]] -; GFX906: [[IF_GFX1010_OR_GFX1101]]: -; GFX906-NEXT: br label %[[LOR_NOT_GFX1010:.*]] -; GFX906: [[LOR_NOT_GFX1010]]: -; GFX906-NEXT: br label %[[FOR_COND:.*]] -; GFX906: [[FOR_COND]]: -; GFX906-NEXT: [[DOTPROMOTED:%.*]] = load i32, ptr [[TMP1]], align 4 -; GFX906-NEXT: [[SUB_PEEL:%.*]] = sub nsw i32 [[DOTPROMOTED]], [[X]] -; GFX906-NEXT: store i32 [[SUB_PEEL]], ptr [[TMP1]], align 4 -; GFX906-NEXT: br label %[[IF_GFX11_INSTS:.*]] -; GFX906: [[IF_GFX11_INSTS]]: -; GFX906-NEXT: br label %[[IF_GFX10_INSTS:.*]] -; GFX906: [[IF_GFX10_INSTS]]: -; GFX906-NEXT: call void @llvm.assume(i1 true) -; GFX906-NEXT: [[DOTPROMOTED9:%.*]] = load i32, ptr [[TMP1]], align 4 -; GFX906-NEXT: [[SUB13_PEEL:%.*]] = sub nsw i32 [[DOTPROMOTED9]], [[X]] -; GFX906-NEXT: store i32 [[SUB13_PEEL]], ptr [[TMP1]], align 4 -; GFX906-NEXT: ret void -; -; GFX1010-LABEL: define amdgpu_kernel void @kernel( -; GFX1010-SAME: ptr addrspace(1) [[P_COERCE:%.*]], i32 [[X:%.*]]) #[[ATTR2:[0-9]+]] { -; GFX1010-NEXT: [[ENTRY:.*:]] -; GFX1010-NEXT: [[TMP0:%.*]] = ptrtoint ptr addrspace(1) [[P_COERCE]] to i64 -; GFX1010-NEXT: [[TMP1:%.*]] = inttoptr i64 [[TMP0]] to ptr -; GFX1010-NEXT: br label %[[IF_GFX1201_OR_GFX12_INSTS:.*]] -; GFX1010: [[IF_GFX1201_OR_GFX12_INSTS]]: -; GFX1010-NEXT: br label %[[IF_NOT_GFX906:.*]] -; GFX1010: [[IF_NOT_GFX906]]: -; GFX1010-NEXT: br label %[[LOR_NOT_GFX1010:.*]] -; GFX1010: [[LOR_NOT_GFX1010]]: -; GFX1010-NEXT: call void @llvm.amdgcn.s.wait.event.export.ready() -; GFX1010-NEXT: br label %[[IF_END6:.*]] -; GFX1010: [[IF_END6]]: -; GFX1010-NEXT: call void @llvm.assume(i1 true) -; GFX1010-NEXT: call void @llvm.assume(i1 true) -; GFX1010-NEXT: br label %[[FOR_COND:.*]] -; GFX1010: [[FOR_COND]]: -; GFX1010-NEXT: [[DOTPROMOTED:%.*]] = load i32, ptr [[TMP1]], align 4 -; GFX1010-NEXT: [[SUB_PEEL:%.*]] = sub nsw i32 [[DOTPROMOTED]], [[X]] -; GFX1010-NEXT: store i32 [[SUB_PEEL]], ptr [[TMP1]], align 4 -; GFX1010-NEXT: br label %[[IF_GFX11_INSTS:.*]] -; GFX1010: [[IF_GFX11_INSTS]]: -; GFX1010-NEXT: br label %[[IF_GFX10_INSTS:.*]] -; GFX1010: [[IF_GFX10_INSTS]]: -; GFX1010-NEXT: call void @llvm.amdgcn.s.ttracedata.imm(i16 1) -; GFX1010-NEXT: br label %[[IF_END11:.*]] -; GFX1010: [[IF_END11]]: -; GFX1010-NEXT: call void @llvm.assume(i1 true) -; GFX1010-NEXT: [[DOTPROMOTED9:%.*]] = load i32, ptr [[TMP1]], align 4 -; GFX1010-NEXT: [[SUB13_PEEL:%.*]] = sub nsw i32 [[DOTPROMOTED9]], [[X]] -; GFX1010-NEXT: store i32 [[SUB13_PEEL]], ptr [[TMP1]], align 4 -; GFX1010-NEXT: ret void -; -; GFX1101-LABEL: define amdgpu_kernel void @kernel( -; GFX1101-SAME: ptr addrspace(1) [[P_COERCE:%.*]], i32 [[X:%.*]]) #[[ATTR2:[0-9]+]] { -; GFX1101-NEXT: [[ENTRY:.*:]] -; GFX1101-NEXT: [[TMP0:%.*]] = ptrtoint ptr addrspace(1) [[P_COERCE]] to i64 -; GFX1101-NEXT: [[TMP1:%.*]] = inttoptr i64 [[TMP0]] to ptr -; GFX1101-NEXT: br label %[[IF_GFX1201_OR_GFX12_INSTS:.*]] -; GFX1101: [[IF_GFX1201_OR_GFX12_INSTS]]: -; GFX1101-NEXT: br label %[[IF_END:.*]] -; GFX1101: [[IF_END]]: -; GFX1101-NEXT: br label %[[IF_NOT_GFX907:.*]] -; GFX1101: [[IF_NOT_GFX907]]: -; GFX1101-NEXT: call void @llvm.amdgcn.s.wait.event.export.ready() -; GFX1101-NEXT: br label %[[IF_NOT_GFX906:.*]] -; GFX1101: [[IF_NOT_GFX906]]: -; GFX1101-NEXT: call void @llvm.assume(i1 true) -; GFX1101-NEXT: call void @llvm.assume(i1 true) -; GFX1101-NEXT: br label %[[FOR_COND:.*]] -; GFX1101: [[FOR_COND]]: -; GFX1101-NEXT: [[DOTPROMOTED:%.*]] = load i32, ptr [[TMP1]], align 4 -; GFX1101-NEXT: [[SUB_PEEL:%.*]] = sub nsw i32 [[DOTPROMOTED]], [[X]] -; GFX1101-NEXT: store i32 [[SUB_PEEL]], ptr [[TMP1]], align 4 -; GFX1101-NEXT: br label %[[IF_GFX11_INSTS:.*]] -; GFX1101: [[IF_GFX11_INSTS]]: -; GFX1101-NEXT: call void @llvm.amdgcn.s.wait.event.export.ready() -; GFX1101-NEXT: br label %[[IF_ELSE8:.*]] -; GFX1101: [[IF_ELSE8]]: -; GFX1101-NEXT: call void @llvm.assume(i1 true) -; GFX1101-NEXT: [[DOTPROMOTED9:%.*]] = load i32, ptr [[TMP1]], align 4 -; GFX1101-NEXT: [[SUB13_PEEL:%.*]] = sub nsw i32 [[DOTPROMOTED9]], [[X]] -; GFX1101-NEXT: store i32 [[SUB13_PEEL]], ptr [[TMP1]], align 4 -; GFX1101-NEXT: ret void -; -; GFX1201-LABEL: define amdgpu_kernel void @kernel( -; GFX1201-SAME: ptr addrspace(1) [[P_COERCE:%.*]], i32 [[X:%.*]]) #[[ATTR2:[0-9]+]] { -; GFX1201-NEXT: [[ENTRY:.*:]] -; GFX1201-NEXT: [[TMP0:%.*]] = ptrtoint ptr addrspace(1) [[P_COERCE]] to i64 -; GFX1201-NEXT: [[TMP1:%.*]] = inttoptr i64 [[TMP0]] to ptr -; GFX1201-NEXT: br label %[[LOR_NOT_GFX1201:.*]] -; GFX1201: [[LOR_NOT_GFX1201]]: -; GFX1201-NEXT: call void @llvm.amdgcn.s.sleep.var(i32 [[X]]) -; GFX1201-NEXT: br label %[[IF_NOT_GFX906:.*]] -; GFX1201: [[IF_NOT_GFX906]]: -; GFX1201-NEXT: br label %[[IF_GFX1010_OR_GFX1101:.*]] -; GFX1201: [[IF_GFX1010_OR_GFX1101]]: -; GFX1201-NEXT: call void @llvm.amdgcn.s.wait.event.export.ready() -; GFX1201-NEXT: br label %[[IF_END6:.*]] -; GFX1201: [[IF_END6]]: -; GFX1201-NEXT: call void @llvm.assume(i1 true) -; GFX1201-NEXT: call void @llvm.assume(i1 true) -; GFX1201-NEXT: br label %[[FOR_COND:.*]] -; GFX1201: [[FOR_COND]]: -; GFX1201-NEXT: [[DOTPROMOTED:%.*]] = load i32, ptr [[TMP1]], align 4 -; GFX1201-NEXT: [[SUB_PEEL:%.*]] = sub nsw i32 [[DOTPROMOTED]], [[X]] -; GFX1201-NEXT: store i32 [[SUB_PEEL]], ptr [[TMP1]], align 4 -; GFX1201-NEXT: br label %[[IF_GFX11_INSTS:.*]] -; GFX1201: [[IF_GFX11_INSTS]]: -; GFX1201-NEXT: call void @llvm.amdgcn.s.wait.event.export.ready() -; GFX1201-NEXT: br label %[[IF_ELSE8:.*]] -; GFX1201: [[IF_ELSE8]]: -; GFX1201-NEXT: call void @llvm.assume(i1 true) -; GFX1201-NEXT: [[DOTPROMOTED9:%.*]] = load i32, ptr [[TMP1]], align 4 -; GFX1201-NEXT: [[SUB13_PEEL:%.*]] = sub nsw i32 [[DOTPROMOTED9]], [[X]] -; GFX1201-NEXT: store i32 [[SUB13_PEEL]], ptr [[TMP1]], align 4 -; GFX1201-NEXT: ret void -; -; GFX1201-W64-LABEL: define amdgpu_kernel void @kernel( -; GFX1201-W64-SAME: ptr addrspace(1) [[P_COERCE:%.*]], i32 [[X:%.*]]) #[[ATTR2:[0-9]+]] { -; GFX1201-W64-NEXT: [[ENTRY:.*:]] -; GFX1201-W64-NEXT: [[TMP0:%.*]] = ptrtoint ptr addrspace(1) [[P_COERCE]] to i64 -; GFX1201-W64-NEXT: [[TMP1:%.*]] = inttoptr i64 [[TMP0]] to ptr -; GFX1201-W64-NEXT: br label %[[LOR_NOT_GFX1201:.*]] -; GFX1201-W64: [[LOR_NOT_GFX1201]]: -; GFX1201-W64-NEXT: call void @llvm.amdgcn.s.sleep.var(i32 [[X]]) -; GFX1201-W64-NEXT: br label %[[IF_NOT_GFX906:.*]] -; GFX1201-W64: [[IF_NOT_GFX906]]: -; GFX1201-W64-NEXT: br label %[[IF_GFX1010_OR_GFX1101:.*]] -; GFX1201-W64: [[IF_GFX1010_OR_GFX1101]]: -; GFX1201-W64-NEXT: call void @llvm.amdgcn.s.wait.event.export.ready() -; GFX1201-W64-NEXT: br label %[[IF_END6:.*]] -; GFX1201-W64: [[IF_END6]]: -; GFX1201-W64-NEXT: call void @llvm.assume(i1 true) -; GFX1201-W64-NEXT: call void @llvm.assume(i1 true) -; GFX1201-W64-NEXT: br label %[[FOR_COND:.*]] -; GFX1201-W64: [[FOR_COND]]: -; GFX1201-W64-NEXT: [[DOTPROMOTED:%.*]] = load i32, ptr [[TMP1]], align 4 -; GFX1201-W64-NEXT: [[SUB_PEEL:%.*]] = sub nsw i32 [[DOTPROMOTED]], [[X]] -; GFX1201-W64-NEXT: store i32 [[SUB_PEEL]], ptr [[TMP1]], align 4 -; GFX1201-W64-NEXT: br label %[[IF_GFX11_INSTS:.*]] -; GFX1201-W64: [[IF_GFX11_INSTS]]: -; GFX1201-W64-NEXT: call void @llvm.amdgcn.s.wait.event.export.ready() -; GFX1201-W64-NEXT: br label %[[IF_ELSE8:.*]] -; GFX1201-W64: [[IF_ELSE8]]: -; GFX1201-W64-NEXT: call void @llvm.assume(i1 true) -; GFX1201-W64-NEXT: [[DOTPROMOTED9:%.*]] = load i32, ptr [[TMP1]], align 4 -; GFX1201-W64-NEXT: [[SUB13_PEEL:%.*]] = sub nsw i32 [[DOTPROMOTED9]], [[X]] -; GFX1201-W64-NEXT: store i32 [[SUB13_PEEL]], ptr [[TMP1]], align 4 -; GFX1201-W64-NEXT: ret void -; -entry: - %0 = ptrtoint ptr addrspace(1) %p.coerce to i64 - %1 = inttoptr i64 %0 to ptr - %2 = load i1, ptr addrspace(1) @llvm.amdgcn.is.gfx1201, align 1 - br i1 %2, label %if.gfx1201.or.gfx12-insts, label %lor.not.gfx1201 - -lor.not.gfx1201: - %3 = load i1, ptr addrspace(1) @llvm.amdgcn.has.gfx12-insts, align 1 - br i1 %3, label %if.gfx1201.or.gfx12-insts, label %if.end - -if.gfx1201.or.gfx12-insts: - call void @llvm.amdgcn.s.sleep.var(i32 %x) - br label %if.end - -if.end: - %4 = load i1, ptr addrspace(1) @llvm.amdgcn.is.gfx906, align 1 - br i1 %4, label %if.gfx906, label %if.not.gfx906 - -if.not.gfx906: - call void @llvm.amdgcn.s.wait.event.export.ready() - br label %if.end6 - -if.gfx906: - %5 = load i1, ptr addrspace(1) @llvm.amdgcn.is.gfx1010, align 1 - br i1 %5, label %if.gfx1010.or.gfx1101, label %lor.not.gfx1010 - -lor.not.gfx1010: - %6 = load i1, ptr addrspace(1) @llvm.amdgcn.is.gfx1101, align 1 - br i1 %6, label %if.gfx1010.or.gfx1101, label %for.cond - -if.gfx1010.or.gfx1101: - call void @llvm.amdgcn.s.ttracedata.imm(i16 1) - br label %if.end6 - -if.end6: - %.pr.pr = load i1, ptr addrspace(1) @llvm.amdgcn.is.gfx1101, align 1 - %7 = icmp ne i1 %.pr.pr, true - call void @llvm.assume(i1 %7) - %.pr6.pr = load i1, ptr addrspace(1) @llvm.amdgcn.is.gfx1010, align 1 - %8 = icmp ne i1 %.pr6.pr, true - call void @llvm.assume(i1 %8) - br label %for.cond - -for.cond: - %.promoted = load i32, ptr %1, align 4 - %sub.peel = sub nsw i32 %.promoted, %x - store i32 %sub.peel, ptr %1, align 4 - %9 = load i1, ptr addrspace(1) @llvm.amdgcn.has.gfx11-insts, align 1 - br i1 %9, label %if.gfx11-insts, label %if.else8 - -if.gfx11-insts: - call void @llvm.amdgcn.s.wait.event.export.ready() - br label %if.end11 - -if.else8: - %10 = load i1, ptr addrspace(1) @llvm.amdgcn.has.gfx10-insts, align 1 - br i1 %10, label %if.gfx10-insts, label %if.end11 - -if.gfx10-insts: - call void @llvm.amdgcn.s.ttracedata.imm(i16 1) - br label %if.end11 - -if.end11: - %.pr7 = load i1, ptr addrspace(1) @"llvm.amdgcn.has.gfx12-insts,wavefrontsize64", align 1 - %11 = icmp ne i1 %.pr7, true - call void @llvm.assume(i1 %11) - %.promoted9 = load i32, ptr %1, align 4 - %sub13.peel = sub nsw i32 %.promoted9, %x - store i32 %sub13.peel, ptr %1, align 4 - ret void -} - -declare void @llvm.assume(i1 noundef) diff --git a/llvm/test/CodeGen/AMDGPU/print-pipeline-passes.ll b/llvm/test/CodeGen/AMDGPU/print-pipeline-passes.ll index 93f43b274e28d0..144d6a7a544055 100644 --- a/llvm/test/CodeGen/AMDGPU/print-pipeline-passes.ll +++ b/llvm/test/CodeGen/AMDGPU/print-pipeline-passes.ll @@ -12,12 +12,9 @@ ; RUN: opt -mtriple=amdgcn--amdhsa -S -passes="lto-pre-link" -print-pipeline-passes -amdgpu-internalize-symbols %s -o - | FileCheck --check-prefix=PRE %s ; RUN: opt -mtriple=amdgcn--amdhsa -S -passes="lto-pre-link" -print-pipeline-passes -amdgpu-internalize-symbols %s -o - | FileCheck --check-prefix=PRE %s -; CHECK: amdgpu-expand-feature-predicates ; CHECK: amdgpu-attributor -; O0: amdgpu-expand-feature-predicates ; O0-NOT: amdgpu-attributor -; PRE: amdgpu-expand-feature-predicates ; PRE-NOT: internalize ; PRE-NOT: amdgpu-attributor ; PRE-NOT: printfToRuntime