[CIR][Dialect] Add convergent attribute to functions for SIMT languages (#840)

seven-mile · lanza · commit 73d2403648cc · 2025-03-13T19:09:51.000-07:00
Fix #805. This PR includes end-to-end implementation. The `convergent` attribute is set depending on languages, which is wrapped as `langOpts.assumeFunctionsAreConvergent()`. Therefore, in ClangIR, every `cir.func` under `#cir.lang<opencl_c>` is set to be convergent. After lowering to LLVM IR, `PostOrderFunctionAttrs` pass will remove unnecessary `convergent` then. In other words, we will still see `convergent` on every function with `-O0`, but not with default optimization level. The test taken from `CodeGenOpenCL/convergent.cl` is a bit complicated. However, the core of it is that `convergent` is set properly for `convfun()` `non_convfun()` `f()` and `g()`. Merge of two `if` is more or less a result of generating the same LLVM IR as OG.
diff --git a/clang/include/clang/CIR/Dialect/IR/CIRAttrs.td b/clang/include/clang/CIR/Dialect/IR/CIRAttrs.td
@@ -1042,6 +1042,10 @@ def NoThrowAttr : CIRUnitAttr<"NoThrow", "nothrow"> {
   let storageType = [{ NoThrowAttr }];
 }
 
+def ConvergentAttr : CIRUnitAttr<"Convergent", "convergent"> {
+  let storageType = [{ ConvergentAttr }];
+}
+
 class CIR_GlobalCtorDtor<string name, string attrMnemonic,
                          string sum, string desc>
     : CIR_Attr<"Global" # name, "global_" # attrMnemonic> {
diff --git a/clang/lib/CIR/CodeGen/CIRGenCall.cpp b/clang/lib/CIR/CodeGen/CIRGenCall.cpp
@@ -440,6 +440,9 @@ void CIRGenModule::constructAttributeList(StringRef Name,
     if (TargetDecl->hasAttr<ArmLocallyStreamingAttr>())
       ;
   }
+
+  getDefaultFunctionAttributes(Name, HasOptnone, AttrOnCallSite, funcAttrs);
+
 }
 
 static mlir::cir::CIRCallOpInterface
@@ -1559,3 +1562,40 @@ mlir::Value CIRGenFunction::buildVAArg(VAArgExpr *VE, Address &VAListAddr) {
   auto vaList = buildVAListRef(VE->getSubExpr()).getPointer();
   return builder.create<mlir::cir::VAArgOp>(loc, type, vaList);
 }
+
+static void getTrivialDefaultFunctionAttributes(
+    StringRef name, bool hasOptnone, const CodeGenOptions &codeGenOpts,
+    const LangOptions &langOpts, bool attrOnCallSite, CIRGenModule &CGM,
+    mlir::NamedAttrList &funcAttrs) {
+
+  if (langOpts.assumeFunctionsAreConvergent()) {
+    // Conservatively, mark all functions and calls in CUDA and OpenCL as
+    // convergent (meaning, they may call an intrinsically convergent op, such
+    // as __syncthreads() / barrier(), and so can't have certain optimizations
+    // applied around them).  LLVM will remove this attribute where it safely
+    // can.
+
+    auto convgt = mlir::cir::ConvergentAttr::get(CGM.getBuilder().getContext());
+    funcAttrs.set(convgt.getMnemonic(), convgt);
+  }
+}
+
+void CIRGenModule::getTrivialDefaultFunctionAttributes(
+    StringRef name, bool hasOptnone, bool attrOnCallSite,
+    mlir::NamedAttrList &funcAttrs) {
+  ::getTrivialDefaultFunctionAttributes(name, hasOptnone, getCodeGenOpts(),
+                                        getLangOpts(), attrOnCallSite, *this,
+                                        funcAttrs);
+}
+
+void CIRGenModule::getDefaultFunctionAttributes(StringRef name, bool hasOptnone,
+                                                bool attrOnCallSite,
+                                                mlir::NamedAttrList &funcAttrs) {
+  getTrivialDefaultFunctionAttributes(name, hasOptnone, attrOnCallSite,
+                                      funcAttrs);
+  // If we're just getting the default, get the default values for mergeable
+  // attributes.
+  if (!attrOnCallSite) {
+    // TODO(cir): addMergableDefaultFunctionAttributes(codeGenOpts, funcAttrs);
+  }
+}
diff --git a/clang/lib/CIR/CodeGen/CIRGenModule.h b/clang/lib/CIR/CodeGen/CIRGenModule.h
@@ -296,6 +296,19 @@ class CIRGenModule : public CIRGenTypeCache {
                               mlir::cir::CallingConv &callingConv,
                               bool AttrOnCallSite, bool IsThunk);
 
+  /// Helper function for getDefaultFunctionAttributes. Builds a set of function
+  /// attributes which can be simply added to a function.
+  void getTrivialDefaultFunctionAttributes(StringRef name, bool hasOptnone,
+                                           bool attrOnCallSite,
+                                           mlir::NamedAttrList &funcAttrs);
+
+  /// Helper function for constructAttributeList and
+  /// addDefaultFunctionDefinitionAttributes.  Builds a set of function
+  /// attributes to add to a function with the given properties.
+  void getDefaultFunctionAttributes(StringRef name, bool hasOptnone,
+                                    bool attrOnCallSite,
+                                    mlir::NamedAttrList &funcAttrs);
+
   /// Will return a global variable of the given type. If a variable with a
   /// different type already exists then a new variable with the right type
   /// will be created and all uses of the old variable will be replaced with a
diff --git a/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVMIR.cpp b/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVMIR.cpp
@@ -109,6 +109,8 @@ class CIRDialectLLVMIRTranslationInterface
           llvmFunc->addFnAttr(llvm::Attribute::OptimizeNone);
         } else if (mlir::dyn_cast<mlir::cir::NoThrowAttr>(attr.getValue())) {
           llvmFunc->addFnAttr(llvm::Attribute::NoUnwind);
+        } else if (mlir::dyn_cast<mlir::cir::ConvergentAttr>(attr.getValue())) {
+          llvmFunc->addFnAttr(llvm::Attribute::Convergent);
         } else if (auto clKernelMetadata =
                        mlir::dyn_cast<mlir::cir::OpenCLKernelMetadataAttr>(
                            attr.getValue())) {
diff --git a/clang/test/CIR/CodeGen/OpenCL/convergent.cl b/clang/test/CIR/CodeGen/OpenCL/convergent.cl
@@ -0,0 +1,107 @@
+// RUN: %clang_cc1 -fclangir -triple spirv64-unknown-unknown -emit-cir %s -o %t.cir
+// RUN: FileCheck %s --input-file=%t.cir --check-prefix=CIR
+// RUN: %clang_cc1 -fclangir -triple spirv64-unknown-unknown -emit-llvm %s -o %t.ll
+// RUN: FileCheck %s --input-file=%t.ll --check-prefix=LLVM
+
+// In ClangIR for OpenCL, all functions should be marked convergent.
+// In LLVM IR, it is initially assumed convergent, but can be deduced to not require it.
+
+// CIR:      #fn_attr[[CONV_NOINLINE_ATTR:[0-9]*]] = #cir<extra({convergent = #cir.convergent, inline = #cir.inline<no>
+// CIR-NEXT: #fn_attr[[CONV_DECL_ATTR:[0-9]*]] = #cir<extra({convergent = #cir.convergent
+// CIR-NEXT: #fn_attr[[CONV_NOTHROW_ATTR:[0-9]*]] = #cir<extra({convergent = #cir.convergent, nothrow = #cir.nothrow
+
+__attribute__((noinline))
+void non_convfun(void) {
+  volatile int* p;
+  *p = 0;
+}
+// CIR: cir.func @non_convfun(){{.*}} extra(#fn_attr[[CONV_NOINLINE_ATTR]])
+// LLVM: define{{.*}} spir_func void @non_convfun() local_unnamed_addr #[[NON_CONV_ATTR:[0-9]+]]
+// LLVM: ret void
+
+// External functions should be assumed convergent.
+void f(void);
+// CIR: cir.func{{.+}} @f(){{.*}} extra(#fn_attr[[CONV_DECL_ATTR]])
+// LLVM: declare {{.+}} spir_func void @f() local_unnamed_addr #[[CONV_ATTR:[0-9]+]]
+void g(void);
+// CIR: cir.func{{.+}} @g(){{.*}} extra(#fn_attr[[CONV_DECL_ATTR]])
+// LLVM: declare {{.+}} spir_func void @g() local_unnamed_addr #[[CONV_ATTR]]
+
+// Test two if's are merged and non_convfun duplicated.
+void test_merge_if(int a) {
+  if (a) {
+    f();
+  }
+  non_convfun();
+  if (a) {
+    g();
+  }
+}
+// CIR: cir.func @test_merge_if{{.*}} extra(#fn_attr[[CONV_NOTHROW_ATTR]])
+
+// The LLVM IR below is equivalent to:
+//    if (a) {
+//      f();
+//      non_convfun();
+//      g();
+//    } else {
+//      non_convfun();
+//    }
+
+// LLVM-LABEL: define{{.*}} spir_func void @test_merge_if
+// LLVM:         %[[tobool:.+]] = icmp eq i32 %[[ARG:.+]], 0
+// LLVM:         br i1 %[[tobool]], label %[[if_end3_critedge:[^,]+]], label %[[if_then:[^,]+]]
+
+// LLVM:       [[if_end3_critedge]]:
+// LLVM:         tail call spir_func void @non_convfun()
+// LLVM:         br label %[[if_end3:[^,]+]]
+
+// LLVM:       [[if_then]]:
+// LLVM:         tail call spir_func void @f()
+// LLVM:         tail call spir_func void @non_convfun()
+// LLVM:         tail call spir_func void @g()
+
+// LLVM:         br label %[[if_end3]]
+
+// LLVM:       [[if_end3]]:
+// LLVM:         ret void
+
+
+void convfun(void) __attribute__((convergent));
+// CIR: cir.func{{.+}} @convfun(){{.*}} extra(#fn_attr[[CONV_DECL_ATTR]])
+// LLVM: declare {{.+}} spir_func void @convfun() local_unnamed_addr #[[CONV_ATTR]]
+
+// Test two if's are not merged.
+void test_no_merge_if(int a) {
+  if (a) {
+    f();
+  }
+  convfun();
+  if(a) {
+    g();
+  }
+}
+// CIR: cir.func @test_no_merge_if{{.*}} extra(#fn_attr[[CONV_NOTHROW_ATTR]])
+
+// LLVM-LABEL: define{{.*}} spir_func void @test_no_merge_if
+// LLVM:         %[[tobool:.+]] = icmp eq i32 %[[ARG:.+]], 0
+// LLVM:         br i1 %[[tobool]], label %[[if_end:[^,]+]], label %[[if_then:[^,]+]]
+// LLVM:       [[if_then]]:
+// LLVM:         tail call spir_func void @f()
+// LLVM-NOT:     call spir_func void @convfun()
+// LLVM-NOT:     call spir_func void @g()
+// LLVM:         br label %[[if_end]]
+// LLVM:       [[if_end]]:
+// LLVM-NOT:     phi i1
+// LLVM:         tail call spir_func void @convfun()
+// LLVM:         br i1 %[[tobool]], label %[[if_end3:[^,]+]], label %[[if_then2:[^,]+]]
+// LLVM:       [[if_then2]]:
+// LLVM:         tail call spir_func void @g()
+// LLVM:         br label %[[if_end3:[^,]+]]
+// LLVM:       [[if_end3]]:
+// LLVM:         ret void
+
+
+// LLVM attribute definitions.
+// LLVM-NOT: attributes #[[NON_CONV_ATTR]] = { {{.*}}convergent{{.*}} }
+// LLVM:     attributes #[[CONV_ATTR]] = { {{.*}}convergent{{.*}} }
