[MLIR][EmitC] Emit private functions as static

[AlexDenisov]

No description provided.

[llvmbot]

@llvm/pr-subscribers-mlir

@llvm/pr-subscribers-mlir-emitc

Author: AlexDenisov (AlexDenisov)

Changes

---

Full diff: https://github.com/llvm/llvm-project/pull/78336.diff

2 Files Affected:

• (modified) mlir/lib/Target/Cpp/TranslateToCpp.cpp (+3)
• (modified) mlir/test/Target/Cpp/call.mlir (+6-3)

diff --git a/mlir/lib/Target/Cpp/TranslateToCpp.cpp b/mlir/lib/Target/Cpp/TranslateToCpp.cpp
index c32cb03caf9db69..e0baa24f143d8b5 100644
--- a/mlir/lib/Target/Cpp/TranslateToCpp.cpp
+++ b/mlir/lib/Target/Cpp/TranslateToCpp.cpp
@@ -745,6 +745,9 @@ static LogicalResult printOperation(CppEmitter &emitter,
 
   CppEmitter::Scope scope(emitter);
   raw_indented_ostream &os = emitter.ostream();
+  if (functionOp.isPrivate()) {
+    os << "static ";
+  }
   if (failed(emitter.emitTypes(functionOp.getLoc(),
                                functionOp.getFunctionType().getResults())))
     return failure();
diff --git a/mlir/test/Target/Cpp/call.mlir b/mlir/test/Target/Cpp/call.mlir
index 2bcdc87205184fb..09dd7642bb680aa 100644
--- a/mlir/test/Target/Cpp/call.mlir
+++ b/mlir/test/Target/Cpp/call.mlir
@@ -1,16 +1,16 @@
 // RUN: mlir-translate -mlir-to-cpp %s | FileCheck %s -check-prefix=CPP-DEFAULT
 // RUN: mlir-translate -mlir-to-cpp -declare-variables-at-top %s | FileCheck %s -check-prefix=CPP-DECLTOP
 
-func.func @emitc_call_opaque() {
+func.func private @emitc_call_opaque() {
   %0 = emitc.call_opaque "func_a" () : () -> i32
   %1 = emitc.call_opaque "func_b" () : () -> i32
   return
 }
-// CPP-DEFAULT: void emitc_call_opaque() {
+// CPP-DEFAULT: static void emitc_call_opaque() {
 // CPP-DEFAULT-NEXT: int32_t [[V0:[^ ]*]] = func_a();
 // CPP-DEFAULT-NEXT: int32_t [[V1:[^ ]*]] = func_b();
 
-// CPP-DECLTOP: void emitc_call_opaque() {
+// CPP-DECLTOP: static void emitc_call_opaque() {
 // CPP-DECLTOP-NEXT: int32_t [[V0:[^ ]*]];
 // CPP-DECLTOP-NEXT: int32_t [[V1:[^ ]*]];
 // CPP-DECLTOP-NEXT: [[V0:]] = func_a();
@@ -22,12 +22,15 @@ func.func @emitc_call_opaque_two_results() {
   %1:2 = emitc.call_opaque "two_results" () : () -> (i32, i32)
   return
 }
+
+// CPP-DEFAULT-NOT: static
 // CPP-DEFAULT: void emitc_call_opaque_two_results() {
 // CPP-DEFAULT-NEXT: size_t [[V1:[^ ]*]] = 0;
 // CPP-DEFAULT-NEXT: int32_t [[V2:[^ ]*]];
 // CPP-DEFAULT-NEXT: int32_t [[V3:[^ ]*]];
 // CPP-DEFAULT-NEXT: std::tie([[V2]], [[V3]]) = two_results();
 
+// CPP-DECLTOP-NOT: static
 // CPP-DECLTOP: void emitc_call_opaque_two_results() {
 // CPP-DECLTOP-NEXT: size_t [[V1:[^ ]*]];
 // CPP-DECLTOP-NEXT: int32_t [[V2:[^ ]*]];

[marbre]

Thanks for the pull-request @AlexDenisov. I am not sure if the concept of an private symbol in MLIR or rather a func.func private directly matches with the C/C++ static keyword. However, I fully agree and see the need to emit static functions. I already have a emitc.func somewhere on my stack which should allow attributes like static and inline and could prioritize this if there is a need. This could also allow to implement a conversion pass that converts a func.func and emitc.func with private mapped to static. WDYT?

[AlexDenisov]

I agree that it doesn't map directly, OTOH I'm seeing CPP target as an approximation.
Re: different implementation: that approach sounds good to me and totally makes sense.
I don't need this functionality urgently (I can work with a local patch in the meantime), but it would be nice to get static functions in eventually.

Would it also make sense to get this functionality in and then migrate to emitc.func? Just trying to avoid this slipping through the cracks 😅

[joker-eph]

Why isn't private mapping to static for functions?

[AlexDenisov]

Why isn't private mapping to static for functions?

I guess the mapping is not one-to-one: static from C doesn't necessarily map directly to private, although I think private maps to static perfectly fine in this case.

[marbre]

Why isn't private mapping to static for functions?

I guess the mapping is not one-to-one: static from C doesn't necessarily map directly to private, although I think private maps to static perfectly fine in this case.

Yeah, that is what I had in mind. I think adding an emitc.call op makes sense anyway as it would allow to emit further keywords such as inline, the combination static inline as well as further keywords like __device__, __host__ and __global__ with view to emit CUDA code.

I agree that it doesn't map directly, OTOH I'm seeing CPP target as an approximation. Re: different implementation: that approach sounds good to me and totally makes sense. I don't need this functionality urgently (I can work with a local patch in the meantime), but it would be nice to get static functions in eventually.

I already restarted my work on an emitc.func but this also requires me to implement an emitc.return (func.return is restricted to have func.func as parent) and also a new emitc.call op (the emitc.call op we previously had was renamed to emitc.call_opaque back in November with c4fd1fd).

Would it also make sense to get this functionality in and then migrate to emitc.func? Just trying to avoid this slipping through the cracks 😅

After several discussions there seemed to be kind of conses that one should rather extend the dialect and do what is necessary within conversions instead of extending support in the emitter for non-EmitC functions. Extending func.func support isn't a big thing and the patch is quite tiny, but I still would prefer to keep the way we go for other options as well. For example scf.for support in the emitter was replaced with an emitc.for op implementation and an SCF to EmitC conversion with 2633d94.
Would it be okay if we only accept your patch if I am unable to complete my work within a reasonable time?

[AlexDenisov]

Would it be okay if we only accept your patch if I am unable to complete my work within a reasonable time?

Totally fine with me, no pressure 🙌

After several discussions

Are these happening in the public space somewhere? I have some questions/suggestions about function prototypes, but would like to discuss them before submitting any patches as the implementation is a bit more involved.

[marbre]

Are these happening in the public space somewhere? I have some questions/suggestions about function prototypes, but would like to discuss them before submitting any patches as the implementation is a bit more involved.

Some of these were in person at last year's C4ML and EuroLLVM, others via email, discord or discourse and even in pull requests. However, I'd love to move the entire discussion to discourse! So please feel free to open new topics with your questions and suggestions at https://discourse.llvm.org/c/mlir/31 and give me and my colleague @simon-camp a ping. If the questions/suggestions are not directly related to upstream, https://github.com/iml130/mlir-emitc/issues can be an alternative place.

[AlexDenisov]

Implemented here #78336 https://github.com/llvm/llvm-project/pull//79612

UPD: fixed link

[marbre]

Implemented here #78336

If somebody tries to follow, the functionally was actually implemented with #79612, which just got merged.