[mlir][emitc][mlir-translate] Add support for function attributes in emitc

[jfurtek]

Summary

This diff adds support for generic function attributes in the EmitC dialect,
along with mlir-translate support for emitting these attributes.

The motivating use case for the changes in this diff is the ability to leverage
the existing EmitC infrastructure for the generation of C code for GPU
targets. A canonical example would be the generation of C code that contains
the __global__ specifier for kernel functions, as used by C/C++ toolchains in
CUDA and HIP programming environments. However, the approach is general enough
to accommodate both standard C/C++ attributes (e.g. static, [[nodiscard]],
...) and platform-specific specifiers (e.g. __declspec(dllexport)), as well as
strings that refer to preprocessor macros in an included header file.

The approach is as follows:

• Add a (discardable) dense string elements attribute to the EmitC dialect to hold attributes as generic strings
• Modify the C++ translation code used by mlir-translate to write function attribute strings before the function

[jfurtek]

• Add newline to EOF

[jpienaar]

mlir/lib/Dialect/EmitC/IR/EmitC.cpp

49

Do we expect arbitrary emitc ops to have this attribute?

[jfurtek]

mlir/lib/Dialect/EmitC/IR/EmitC.cpp

49

I don't expect arbitrary emitc ops to have this attribute. At the moment, no emitc ops would have this attribute. It is used exclusively for operations that implement FunctionOpInterface.

IIUC, this function is called for a discardable attribute, which must have a name of the form emitc.some_attribute_name. As func_attr is currently the only discardable attribute in the emitc dialect, any other name would be an error.

Would it be preferable to return success() for an unknown discardable attribute? (It seems like the Linalg implementation returns an error for an unexpected attribute name, but some other implementations of verifyOperationAttribute() (e.g. the gpu dialect) return success().)

[christopherbate]

I think GPU function support for EmitC may be better off as a GPUToEmitC conversion pass where gpu.func gets converted to func.func. This would subsume the pass you added as well as the changes to the emitter.

There are other aspects of gpu.func (e.g. the attached workgroup memory attributions) that need to get lowered into emitc calls as well which are currently not handled below in the emitter changes. It's easy to imagine needing to pass user-options to this conversion as well, and you probably don't want to try to wire up more flags into the translation step than absolutely necessary.

Other gpu dialect ops would need to get lowered to calls (e.g. gpu.thread_id and associated operations). It makes sense to add a dedicated conversion pass rather than adding more extensions to the emitter.

mlir/include/mlir/Dialect/EmitC/IR/EmitCBase.td

39

I don't think this one and line 41 really belong in EmitC.

mlir/include/mlir/Dialect/EmitC/IR/EmitCTypes.td

75 ↗

(On Diff #486695)

Why does it need it's own string type?

mlir/include/mlir/Dialect/EmitC/Transforms/Passes.h

17 ↗

(On Diff #486695)

This can be moved to line 32/33 and remove line 18.

25 ↗

(On Diff #486695)

You shouldn't need to add this declaration anymore.

mlir/include/mlir/Dialect/EmitC/Transforms/Passes.td

26 ↗

(On Diff #486695)

This line (26) can get omitted and the declarations will be generated.

mlir/lib/Dialect/EmitC/IR/EmitC.cpp

52

Why do you want it to be a DenseStringElementsAttr? Can't it just be a normal ArrayAttr of StringAttr?

[jfurtek]

• Add newline to EOF
• Remove GPU from EmitC dependencies and remove GPU function attribute pass
• Revert small differences from original upstream

[marbre]

Thanks for driving this! We actually had something similar on our long term roadmap and think something function meta attribute like would be desirable. However, for the omp dialect there it is currently discussed on discourse if to go for an omp.function instead of a dialect attribute, see https://discourse.llvm.org/t/rfc-omp-module-and-omp-function-vs-dialect-attributes-to-encode-openmp-properties/67998. I think the same discussion applies here as well. Instead of having emitc.func_attr together with func.func we might want to have emitc.func?

mlir/include/mlir/Dialect/EmitC/IR/EmitCBase.td

def EmitC_Dialect : Dialect {
   let description = [{
     [include "Dialects/emitc.md"]
   }];
 
   let hasConstantMaterializer = 1;
   let useDefaultTypePrinterParser = 1;
   let useDefaultAttributePrinterParser = 1;
   let useFoldAPI = kEmitFoldAdaptorFolder;
+  let hasOperationAttrVerify = 1;
+  let extraClassDeclaration = [{
+    /// Get the name of the attribute used for emitc function attributes.
+    static StringRef getFuncAttrName() { return "emitc.func_attr"; }
+  }];
 }
 
 #endif // MLIR_DIALECT_EMITC_IR_EMITCBASE

mlir/lib/Dialect/EmitC/IR/EmitC.cpp

 /// Materialize a single constant operation from a given attribute value with
 /// the desired resultant type.
 Operation *EmitCDialect::materializeConstant(OpBuilder &builder,
                                              Attribute value, Type type,
                                              Location loc) {
   return builder.create<emitc::ConstantOp>(loc, type, value);
 }
 
+LogicalResult EmitCDialect::verifyOperationAttribute(Operation *op,
+                                                     NamedAttribute attr) {
+  if (attr.getName() == getFuncAttrName()) {
+    auto arrayAttr = dyn_cast<ArrayAttr>(attr.getValue());
+    if (!arrayAttr || !llvm::all_of(arrayAttr, [](Attribute attr) {
+          return attr.isa<StringAttr>();
+        }))
+      return op->emitError() << "attribute '" << attr.getName()
+                             << "' must be an array of strings";
+  }
+  // Pass through other attributes
+  return success();
+}
+
 //===----------------------------------------------------------------------===//
 // ApplyOp
 //===----------------------------------------------------------------------===//
 
 LogicalResult ApplyOp::verify() {
   StringRef applicableOperatorStr = getApplicableOperator();
 
   // Applicable operator must not be empty.

mlir/lib/Target/Cpp/TranslateToCpp.cpp

static LogicalResult printOperation(CppEmitter &emitter,
   if (!emitter.shouldDeclareVariablesAtTop() &&
       functionOp.getBlocks().size() > 1) {
     return functionOp.emitOpError(
         "with multiple blocks needs variables declared at top");
   }
 
   CppEmitter::Scope scope(emitter);
   raw_indented_ostream &os = emitter.ostream();
+  // Emit any function attributes if the (discardable) emitc attribute is
+  // present.
+  if (auto func_attr = functionOp->getAttrOfType<ArrayAttr>(
+          EmitCDialect::getFuncAttrName())) {
+    for (auto &attr : func_attr) {
+      if (auto stringAttr = attr.dyn_cast<StringAttr>())
+        os << stringAttr.getValue() << "\n";
+    }
+  }
+
   if (failed(emitter.emitTypes(functionOp.getLoc(),
                                functionOp.getFunctionType().getResults())))
     return failure();
   os << " " << functionOp.getName();
 
   os << "(";
   if (failed(interleaveCommaWithError(
           functionOp.getArguments(), os,

mlir/test/Target/Cpp/common-cpp.mlir

func.func @opaque_types(%arg0: !emitc.opaque<"bool">, %arg1: !emitc.opaque<"char">) -> !emitc.opaque<"status_t"> {
   %0 = emitc.call "a"(%arg0, %arg1) : (!emitc.opaque<"bool">, !emitc.opaque<"char">) -> (!emitc.opaque<"int">)
   // CHECK: char [[V4:[^ ]*]] = b([[V3]]);
   %1 = emitc.call "b"(%0): (!emitc.opaque<"int">) -> (!emitc.opaque<"char">)
   // CHECK: status_t [[V5:[^ ]*]] = c([[V3]], [[V4]]);
   %2 = emitc.call "c"(%0, %1): (!emitc.opaque<"int">, !emitc.opaque<"char">) -> (!emitc.opaque<"status_t">)
   return %2 : !emitc.opaque<"status_t">
 }
 
-func.func @apply(%arg0: i32) -> !emitc.ptr<i32> {
-  // CHECK: int32_t* [[V2]] = &[[V1]];
+// Note: FileCheck will interpret [[nodiscard]] as a variable, so we use
+// __attribute__ below.
+// CHECK: static
+// CHECK: __attribute__((warn_unused_result))
+// CHECK: int32_t* apply(int32_t [[V1:[^ ]*]]) {
+func.func @apply(%arg0: i32) -> !emitc.ptr<i32> attributes {emitc.func_attr = ["static", "__attribute__((warn_unused_result))"]} {
+  // CHECK: int32_t* [[V2:[^ ]*]] = &[[V1]];
   %0 = emitc.apply "&"(%arg0) : (i32) -> !emitc.ptr<i32>
   // CHECK: int32_t [[V3]] = *[[V2]];
   %1 = emitc.apply "*"(%0) : (!emitc.ptr<i32>) -> (i32)
   return %0 : !emitc.ptr<i32>
 }