Handle jumps into controlled sequences.

Author: AlexVlx

clang/include/clang/Basic/DiagnosticSemaKinds.td

@@ -13312,4 +13312,6 @@ def err_amdgcn_predicate_type_needs_explicit_bool_cast
     : Error<"%0 must be explicitly cast to %1; however, please note that this "
             "is almost always an error and that it prevents the effective "
             "guarding of target dependent code, and thus should be avoided">;
+def note_amdgcn_protected_by_predicate
+    : Note<"jump enters statement controlled by AMDGPU feature predicate">;
 } // end of sema component.

clang/include/clang/Sema/SemaAMDGPU.h

@@ -15,12 +15,15 @@
 
 #include "clang/AST/ASTFwd.h"
 #include "clang/Sema/SemaBase.h"
+#include "llvm/ADT/SmallPtrSet.h"
 
 namespace clang {
 class AttributeCommonInfo;
+class Expr;
 class ParsedAttr;
 
 class SemaAMDGPU : public SemaBase {
+  llvm::SmallPtrSet<Expr *, 32> ExpandedPredicates;
 public:
   SemaAMDGPU(Sema &S);
 
@@ -68,6 +71,7 @@ class SemaAMDGPU : public SemaBase {
   /// Expand a valid use of the feature identification builtins into its
   /// corresponding sequence of instructions.
   Expr *ExpandAMDGPUPredicateBI(CallExpr *CE);
+  bool IsPredicate(Expr *E) const;
 };
 } // namespace clang
 

clang/lib/Sema/JumpDiagnostics.cpp

@@ -19,6 +19,7 @@
 #include "clang/AST/StmtOpenACC.h"
 #include "clang/AST/StmtOpenMP.h"
 #include "clang/Basic/SourceLocation.h"
+#include "clang/Sema/SemaAMDGPU.h"
 #include "clang/Sema/SemaInternal.h"
 #include "llvm/ADT/BitVector.h"
 using namespace clang;
@@ -367,15 +368,19 @@ void JumpScopeChecker::BuildScopeInformation(Stmt *S,
 
   case Stmt::IfStmtClass: {
     IfStmt *IS = cast<IfStmt>(S);
+    bool AMDGPUPredicate = false;
     if (!(IS->isConstexpr() || IS->isConsteval() ||
-          IS->isObjCAvailabilityCheck()))
+          IS->isObjCAvailabilityCheck() ||
+          (AMDGPUPredicate = this->S.AMDGPU().IsPredicate(IS->getCond()))))
       break;
 
     unsigned Diag = diag::note_protected_by_if_available;
     if (IS->isConstexpr())
       Diag = diag::note_protected_by_constexpr_if;
     else if (IS->isConsteval())
       Diag = diag::note_protected_by_consteval_if;
+    else if (AMDGPUPredicate)
+      Diag = diag::note_amdgcn_protected_by_predicate;
 
     if (VarDecl *Var = IS->getConditionVariable())
       BuildScopeInformation(Var, ParentScope);

clang/lib/Sema/SemaAMDGPU.cpp

@@ -375,7 +375,8 @@ Expr *SemaAMDGPU::ExpandAMDGPUPredicateBI(CallExpr *CE) {
   auto Loc = CE->getExprLoc();
 
   if (!CE->getBuiltinCallee())
-    return IntegerLiteral::Create(Ctx, False, BoolTy, Loc);
+    return *ExpandedPredicates.insert(
+        IntegerLiteral::Create(Ctx, False, BoolTy, Loc)).first;
 
   auto P = false;
   auto BI = CE->getBuiltinCallee();
@@ -398,7 +399,7 @@ Expr *SemaAMDGPU::ExpandAMDGPUPredicateBI(CallExpr *CE) {
     }
     if (Ctx.getTargetInfo().getTriple().isSPIRV()) {
       CE->setType(BoolTy);
-      return CE;
+      return *ExpandedPredicates.insert(CE).first;
     }
 
     if (auto TID = Ctx.getTargetInfo().getTargetID())
@@ -412,7 +413,7 @@ Expr *SemaAMDGPU::ExpandAMDGPUPredicateBI(CallExpr *CE) {
 
     if (Ctx.getTargetInfo().getTriple().isSPIRV()) {
       CE->setType(BoolTy);
-      return CE;
+      return *ExpandedPredicates.insert(CE).first;
     }
 
     auto *FD = cast<FunctionDecl>(Arg->getReferencedDeclOfCallee());
@@ -424,6 +425,11 @@ Expr *SemaAMDGPU::ExpandAMDGPUPredicateBI(CallExpr *CE) {
     P = Builtin::evaluateRequiredTargetFeatures(RF, CF);
   }
 
-  return IntegerLiteral::Create(Ctx, P ? True : False, BoolTy, Loc);
+  return *ExpandedPredicates.insert(
+      IntegerLiteral::Create(Ctx, P ? True : False, BoolTy, Loc)).first;
+}
+
+bool SemaAMDGPU::IsPredicate(Expr *E) const {
+  return ExpandedPredicates.contains(E);
 }
 } // namespace clang

clang/test/SemaHIP/amdgpu-feature-builtins-cant-jump.hip

@@ -0,0 +1,62 @@
+// REQUIRES: amdgpu-registered-target
+// REQUIRES: spirv-registered-target
+// RUN: %clang_cc1 -fsyntax-only -verify -triple amdgcn -target-cpu gfx900 -Wno-unused-value %s
+// RUN: %clang_cc1 -fsyntax-only -verify -triple amdgcn -target-cpu gfx1201 -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
+
+#define __device__ __attribute__((device))
+#define __global__ __attribute__((global))
+
+__device__ void f(int *ptr, int size, bool f) {
+    int i = 0;
+    if (f)
+        goto label; // expected-error {{cannot jump from this goto statement to its label}}
+
+    if (__builtin_amdgcn_processor_is("gfx900")) { // expected-note {{jump enters statement controlled by AMDGPU feature predicate}}
+        for (i = 0; i < size; ++i) {
+            label:
+            ptr[i] = i;
+        }
+    }
+}
+
+__device__ void g(int *ptr, int size, bool f) {
+    int i = 0;
+    if (f)
+        goto label; // expected-error {{cannot jump from this goto statement to its label}}
+
+    if (__builtin_amdgcn_is_invocable(__builtin_amdgcn_s_sleep_var)) { // expected-note {{jump enters statement controlled by AMDGPU feature predicate}}
+        for (i = 0; i < size; ++i) {
+            label:
+            ptr[i] = i;
+        }
+    }
+}
+
+__global__ void h(int *ptr, int size, bool f) {
+    int i = 0;
+    if (f)
+        goto label; // expected-error {{cannot jump from this goto statement to its label}}
+
+    if (__builtin_amdgcn_processor_is("gfx900")) { // expected-note {{jump enters statement controlled by AMDGPU feature predicate}}
+        for (i = 0; i < size; ++i) {
+            label:
+            ptr[i] = i;
+        }
+    }
+}
+
+__global__ void i(int *ptr, int size, bool f) {
+    int i = 0;
+    if (f)
+        goto label; // expected-error {{cannot jump from this goto statement to its label}}
+
+    if (__builtin_amdgcn_is_invocable(__builtin_amdgcn_s_sleep_var)) { // expected-note {{jump enters statement controlled by AMDGPU feature predicate}}
+        for (i = 0; i < size; ++i) {
+            label:
+            ptr[i] = i;
+        }
+    }
+}