Skip to content

Commit

Permalink
[CIR][Dialect] Add convergent attribute to functions for SIMT langu…
Browse files Browse the repository at this point in the history
…ages (llvm#840)

Fix llvm#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.
  • Loading branch information
seven-mile authored and lanza committed Oct 14, 2024
1 parent 143e2a5 commit bff487f
Show file tree
Hide file tree
Showing 5 changed files with 166 additions and 0 deletions.
4 changes: 4 additions & 0 deletions clang/include/clang/CIR/Dialect/IR/CIRAttrs.td
Original file line number Diff line number Diff line change
Expand Up @@ -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> {
Expand Down
40 changes: 40 additions & 0 deletions clang/lib/CIR/CodeGen/CIRGenCall.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -440,6 +440,9 @@ void CIRGenModule::constructAttributeList(StringRef Name,
if (TargetDecl->hasAttr<ArmLocallyStreamingAttr>())
;
}

getDefaultFunctionAttributes(Name, HasOptnone, AttrOnCallSite, funcAttrs);

}

static mlir::cir::CIRCallOpInterface
Expand Down Expand Up @@ -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);
}
}
13 changes: 13 additions & 0 deletions clang/lib/CIR/CodeGen/CIRGenModule.h
Original file line number Diff line number Diff line change
Expand Up @@ -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
Expand Down
2 changes: 2 additions & 0 deletions clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVMIR.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -108,6 +108,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())) {
Expand Down
107 changes: 107 additions & 0 deletions clang/test/CIR/CodeGen/OpenCL/convergent.cl
Original file line number Diff line number Diff line change
@@ -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{{.*}} }

0 comments on commit bff487f

Please sign in to comment.