[CIR] Add support for AliasOp

clang/include/clang/CIR/Dialect/Builder/CIRBaseBuilder.h

@@ -113,7 +113,7 @@ class CIRBaseBuilderTy : public mlir::OpBuilder {
     return getPointerTo(cir::VoidType::get(getContext()), as);
   }
 
-  cir::MethodAttr getMethodAttr(cir::MethodType ty, cir::FuncOp methodFuncOp) {
+  cir::MethodAttr getMethodAttr(cir::MethodType ty, cir::CIRCallableOpInterface methodFuncOp) {
     auto methodFuncSymbolRef = mlir::FlatSymbolRefAttr::get(methodFuncOp);
     return cir::MethodAttr::get(ty, methodFuncSymbolRef);
   }
@@ -681,7 +681,7 @@ class CIRBaseBuilderTy : public mlir::OpBuilder {
     return callOp;
   }
 
-  cir::CallOp createCallOp(mlir::Location loc, cir::FuncOp callee,
+  cir::CallOp createCallOp(mlir::Location loc, cir::CIRCallableOpInterface callee,
                            mlir::ValueRange operands = mlir::ValueRange(),
                            cir::CallingConv callingConv = cir::CallingConv::C,
                            cir::SideEffect sideEffect = cir::SideEffect::All,
@@ -737,7 +737,7 @@ class CIRBaseBuilderTy : public mlir::OpBuilder {
   }
 
   cir::CallOp
-  createTryCallOp(mlir::Location loc, cir::FuncOp callee,
+  createTryCallOp(mlir::Location loc, cir::CIRCallableOpInterface callee,
                   mlir::ValueRange operands,
                   cir::CallingConv callingConv = cir::CallingConv::C,
                   cir::SideEffect sideEffect = cir::SideEffect::All,

clang/include/clang/CIR/Dialect/IR/CIROps.td

@@ -3762,8 +3762,95 @@ def CIR_OptionalPriorityAttr : OptionalAttr<
   >
 >;
 
+def CIR_AliasOp : CIR_Op<"alias", [
+  AutomaticAllocationScope, 
+  DeclareOpInterfaceMethods<CIRGlobalValueInterface>,
+  CIRCallableOpInterface,
+  CallableOpInterface, FunctionOpInterface, IsolatedFromAbove]> {
+  let summary = "Declare or define an alias.";
+
+  let description = [{
+    Declare or define an alias that can be used for either 
+    a global variable or a function (only support function for now).
+
+    This operation models the LLVMIR AliasOp to ensure easy lowering:
+    >`llvm.mlir.alias` is a top level operation that defines a global alias for
+    >global variables and functions. The operation is always initialized by
+    >using a initializer region which could be a direct map to another global
+    >value or contain some address computation on top of it.
+
+    It also mirrors CIR_FuncOp to allow for generation of functions and calling of said
+    functions. 
+
+    Examples:
+    ```
+       // Defining an alias. 
+       "cir.alias"() <{aliasee = @_ZN1AC2Ev, calling_conv = 1 : i32, function_type = !cir.func<(!cir.ptr<!rec_A>)>, global_visibility = #cir<visibility default>, linkage = 0 : i32, sym_name = "_ZN1AC1Ev"}> ({
+        }) {sym_visibility = "private"} : () -> () 
+
+
+       // Call an alias op like how you would call a FuncOp
+       cir.call @_ZN1AC1Ev(%0) : (!cir.ptr<!rec_A>) -> () 
+
+    ```
+
+
+
+  }];
+
+  let arguments = (ins 
+    SymbolNameAttr:$sym_name,
+    TypeAttrOf<CIR_FuncType>:$function_type, FlatSymbolRefAttr:$aliasee,
+    CIR_VisibilityAttr:$global_visibility,
+    DefaultValuedAttr<CIR_GlobalLinkageKind,
+                      "GlobalLinkageKind::ExternalLinkage">:$linkage,
+    OptionalAttr<DictArrayAttr>:$arg_attrs,
+    OptionalAttr<DictArrayAttr>:$res_attrs,
+    UnitAttr:$comdat,
+    UnitAttr:$dso_local,
+    UnitAttr:$no_proto,
+    UnitAttr:$builtin,
+    CIR_OptionalPriorityAttr:$global_ctor_priority,
+    CIR_OptionalPriorityAttr:$global_dtor_priority,
+    CIR_ExtraFuncAttr:$extra_attrs,
+    OptionalAttr<CIR_CXXSpecialMemberAttr>:$cxx_special_member,
+    DefaultValuedAttr<CIR_CallingConv, "CallingConv::C">:$calling_conv
+   );
+
+  let regions = (region AnyRegion:$body);
+  let skipDefaultBuilders = 1;
+
+  let builders = [OpBuilder<(ins
+    "llvm::StringRef":$name, "FuncType":$type,
+    "llvm::StringRef":$aliasee,
+    CArg<"GlobalLinkageKind", "GlobalLinkageKind::ExternalLinkage">:$linkage,
+    CArg<"CallingConv", "CallingConv::C">:$callingConv,
+    CArg<"llvm::ArrayRef<mlir::NamedAttribute>", "{}">:$attrs,
+    CArg<"llvm::ArrayRef<mlir::DictionaryAttr>", "{}">:$argAttrs)
+  >];
+
+  let extraClassDeclaration = [{
+    /// Returns the region on the current operation that is callable.
+    /// Since this is just an alias, it should return nullptr.
+
+    /// Actualy definition will be generated once lowered to LLVMIR
+    ::mlir::Region *getCallableRegion() {
+      return nullptr;
+    }
+
+    //===------------------------------------------------------------------===//
+    // SymbolOpInterface Methods
+    //===------------------------------------------------------------------===//
+
+    bool isDeclaration() {
+      return true;
+    }
+  }];
+}
+
 def FuncOp : CIR_Op<"func", [
   AutomaticAllocationScope, CallableOpInterface, FunctionOpInterface,
+  CIRCallableOpInterface,
   DeclareOpInterfaceMethods<CIRGlobalValueInterface>,
   HasAtMostOneOfAttrs<["global_ctor_priority", "global_dtor_priority"]>,
   IsolatedFromAbove
@@ -3889,35 +3976,6 @@ def FuncOp : CIR_Op<"func", [
     /// function.
     ::mlir::Region *getCallableRegion();
 
-    /// Returns the results types that the callable region produces when
-    /// executed.
-    llvm::ArrayRef<mlir::Type> getCallableResults() {
-      return getFunctionType().getReturnTypes();
-    }
-
-    /// Returns the argument attributes for all callable region arguments or
-    /// null if there are none.
-    ::mlir::ArrayAttr getCallableArgAttrs() {
-      return getArgAttrs().value_or(nullptr);
-    }
-
-    /// Returns the result attributes for all callable region results or null if
-    /// there are none.
-    ::mlir::ArrayAttr getCallableResAttrs() {
-      return getResAttrs().value_or(nullptr);
-    }
-
-    /// Returns the argument types of this function.
-    llvm::ArrayRef<mlir::Type> getArgumentTypes() {
-       return getFunctionType().getInputs();
-    }
-
-    /// Returns 0 or 1 result type of this function (0 in the case of a function
-    /// returing void)
-    llvm::ArrayRef<mlir::Type> getResultTypes() {
-       return getFunctionType().getReturnTypes();
-    }
-
     //===------------------------------------------------------------------===//
     // SymbolOpInterface Methods
     //===------------------------------------------------------------------===//

clang/include/clang/CIR/Interfaces/CIROpInterfaces.h

@@ -18,6 +18,8 @@
 #include "clang/AST/DeclTemplate.h"
 #include "clang/AST/Mangle.h"
 #include "clang/CIR/Dialect/IR/CIROpsEnums.h"
+#include "clang/CIR/Dialect/IR/CIRTypes.h"
+#include "clang/CIR/Dialect/IR/CIRAttrs.h"
 
 namespace cir {} // namespace cir
 

clang/include/clang/CIR/Interfaces/CIROpInterfaces.td

@@ -43,6 +43,123 @@ let cppNamespace = "::cir" in {
     ];
   }
 
+  def CIRCallableOpInterface
+      : OpInterface<"CIRCallableOpInterface"> {
+    let description = [{
+       The CIRCallableOpInterface is created due to the flexibility that 
+       either AliasOp or FuncOp is returned from the same function.
+
+       It also handles all the default and similar methods that AliasOp and FuncOp might have so
+       we don't have to define them twice.
+
+       Notice that even though we are declaring methods from CallableOpInterface and
+       FunctionOpInterface, we cannot add them to the inheritance in OpInterface<...>, as
+       it will produce clashing of definitions.
+    }];
+    let methods = [
+      InterfaceMethod<
+          "Returns the region on the current operation that is callable. This may "
+          "return null in the case of an external callable object, e.g. an external "
+          "function.",
+          "::mlir::Region*", "getCallableRegion", (ins)>,
+      InterfaceMethod<
+          "Returns the type of this function.",
+          "cir::FuncType", "getFunctionType",
+          (ins)>,
+      InterfaceMethod<
+          "Returns the results types that the callable region produces when "
+          "executed.",
+          "llvm::ArrayRef<mlir::Type>", "getCallableResults",
+          (ins), [{}],
+      /*defaultImplementation=*/[{
+          return $_op.getFunctionType().getReturnTypes();
+      }]>,
+      InterfaceMethod<
+          "Returns the result attributes for all callable region results or null if "
+          "there are none.",
+          "::mlir::ArrayAttr", "getCallableResAttrs",
+          (ins), [{}],
+      /*defaultImplementation=*/[{
+          return $_op.getResAttrs().value_or(nullptr);
+      }]>,
+      InterfaceMethod<
+          "Returns the argument types of this function.",
+          "llvm::ArrayRef<mlir::Type>", "getArgumentTypes",
+          (ins), [{}],
+      /*defaultImplementation=*/[{
+       return $_op.getFunctionType().getInputs();
+      }]>,
+      InterfaceMethod<
+          "Returns 0 or 1 result type of this function (0 in the case of a function "
+          "returing void).",
+          "llvm::ArrayRef<mlir::Type>", "getResultTypes",
+          (ins), [{}],
+      /*defaultImplementation=*/[{
+       return $_op.getFunctionType().getReturnTypes();
+      }]>,
+      InterfaceMethod<
+          "Return the calling convention of the called operation",
+          "cir::CallingConv", "getCallingConv", (ins)>,
+      InterfaceMethod<
+          "Return if the callable is prototype defined or not",
+          "bool", "getNoProto", (ins)>,
+      InterfaceMethod<
+          "Set the attribute indicating if the callable is prototype defined or not",
+          "void", "setNoProtoAttr", (ins "::mlir::UnitAttr":$noProto)>,
+      InterfaceMethod<
+          "Set the attribute indicating if the callable is prototype defined or not",
+          "::mlir::UnitAttr", "getNoProtoAttr", (ins)>,
+      InterfaceMethod<
+          "Return if the callable is builtin or not",
+          "::mlir::UnitAttr", "getBuiltinAttr", (ins)>,
+      InterfaceMethod<
+          "Set the attribute indicating if the callable is builtin or not",
+          "void", "setBuiltinAttr", (ins "::mlir::UnitAttr":$builtin)>,
+      InterfaceMethod<
+          "Set the function's calling convention",
+          "void", "setCallingConv", (ins "::cir::CallingConv":$attr)>,
+      InterfaceMethod<
+          "Set the function's calling convention",
+          "::cir::ExtraFuncAttributesAttr", "getExtraAttrs", (ins)>,
+      InterfaceMethod<
+          "Set the function's calling convention",
+          "void", "setExtraAttrsAttr", (ins "::cir::ExtraFuncAttributesAttr":$attr)>,
+      InterfaceMethod<
+          "Return if the callable is a declaration or not",
+          "bool", "isDeclaration", (ins)>,
+      InterfaceMethod<
+          "Return function symbol name",
+          "::mlir::StringRef", "getSymName", (ins)>,
+      InterfaceMethod<
+          "Return function symbol name",
+          "::mlir::StringAttr", "getSymNameAttr", (ins)>,
+      InterfaceMethod<
+          "Return function linkage kind",
+          "::cir::GlobalLinkageKind", "getLinkage", (ins)>,
+      InterfaceMethod<
+          "Set function linkage kind",
+          "void", "setLinkage", (ins "::cir::GlobalLinkageKind":$kind)>,
+      InterfaceMethod<
+          "Set function linkage kind attr",
+          "void", "setLinkageAttr", (ins "::cir::GlobalLinkageKindAttr":$kind)>,
+      InterfaceMethod<
+          "Set function linkage kind",
+          "void", "setGlobalVisibilityAttr", (ins "::cir::VisibilityAttr":$attr)>,
+      InterfaceMethod<
+          "Set function custructor priotity",
+          "void", "setGlobalCtorPriority", (ins "::std::optional<uint32_t>":$prio)>,
+      InterfaceMethod<
+          "Set function custructor priotity",
+          "void", "setGlobalDtorPriority", (ins "::std::optional<uint32_t>":$prio)>,
+      InterfaceMethod<
+          "Set special member attribute",
+          "void", "setCxxSpecialMemberAttr", (ins "mlir::Attribute":$attr)>,
+
+
+
+    ];
+  }
+
   def CIRGlobalValueInterface
       : OpInterface<"CIRGlobalValueInterface", [Symbol]> {
 

clang/lib/CIR/CodeGen/CIRGenBuiltin.cpp

@@ -2887,7 +2887,7 @@ mlir::Value CIRGenFunction::evaluateOrEmitBuiltinObjectSize(
 
 /// Given a builtin id for a function like "__builtin_fabsf", return a Function*
 /// for "fabsf".
-cir::FuncOp CIRGenModule::getBuiltinLibFunction(const FunctionDecl *FD,
+cir::CIRCallableOpInterface CIRGenModule::getBuiltinLibFunction(const FunctionDecl *FD,
                                                 unsigned BuiltinID) {
   assert(astContext.BuiltinInfo.isLibFunction(BuiltinID));
 

clang/lib/CIR/CodeGen/CIRGenCUDANV.cpp

@@ -59,9 +59,9 @@ class CIRGenNVCUDARuntime : public CIRGenCUDARuntime {
   std::unique_ptr<MangleContext> deviceMC;
 
 private:
-  void emitDeviceStubBodyLegacy(CIRGenFunction &cgf, cir::FuncOp fn,
+  void emitDeviceStubBodyLegacy(CIRGenFunction &cgf, cir::CIRCallableOpInterface fn,
                                 FunctionArgList &args);
-  void emitDeviceStubBodyNew(CIRGenFunction &cgf, cir::FuncOp fn,
+  void emitDeviceStubBodyNew(CIRGenFunction &cgf, cir::CIRCallableOpInterface fn,
                              FunctionArgList &args);
   std::string addPrefixToName(StringRef FuncName) const;
   std::string addUnderscoredPrefixToName(StringRef FuncName) const;
@@ -70,10 +70,10 @@ class CIRGenNVCUDARuntime : public CIRGenCUDARuntime {
   CIRGenNVCUDARuntime(CIRGenModule &cgm);
   ~CIRGenNVCUDARuntime();
 
-  void emitDeviceStub(CIRGenFunction &cgf, cir::FuncOp fn,
+  void emitDeviceStub(CIRGenFunction &cgf, cir::CIRCallableOpInterface fn,
                       FunctionArgList &args) override;
 
-  mlir::Operation *getKernelHandle(cir::FuncOp fn, GlobalDecl GD) override;
+  mlir::Operation *getKernelHandle(cir::CIRCallableOpInterface fn, GlobalDecl GD) override;
 
   void internalizeDeviceSideVar(const VarDecl *d,
                                 cir::GlobalLinkageKind &linkage) override;
@@ -109,13 +109,13 @@ CIRGenNVCUDARuntime::addUnderscoredPrefixToName(StringRef FuncName) const {
 }
 
 void CIRGenNVCUDARuntime::emitDeviceStubBodyLegacy(CIRGenFunction &cgf,
-                                                   cir::FuncOp fn,
+                                                   cir::CIRCallableOpInterface fn,
                                                    FunctionArgList &args) {
   llvm_unreachable("NYI");
 }
 
 void CIRGenNVCUDARuntime::emitDeviceStubBodyNew(CIRGenFunction &cgf,
-                                                cir::FuncOp fn,
+                                                cir::CIRCallableOpInterface fn,
                                                 FunctionArgList &args) {
 
   // This requires arguments to be sent to kernels in a different way.
@@ -198,7 +198,7 @@ void CIRGenNVCUDARuntime::emitDeviceStubBodyNew(CIRGenFunction &cgf,
       builder.createAlloca(loc, cir::PointerType::get(streamTy), streamTy,
                            "stream", cgm.getPointerAlign());
 
-  cir::FuncOp popConfig = cgm.createRuntimeFunction(
+  cir::CIRCallableOpInterface popConfig = cgm.createRuntimeFunction(
       cir::FuncType::get({gridDim.getType(), blockDim.getType(),
                           sharedMem.getType(), stream.getType()},
                          cgm.SInt32Ty),
@@ -257,7 +257,7 @@ void CIRGenNVCUDARuntime::emitDeviceStubBodyNew(CIRGenFunction &cgf,
                launchArgs);
 }
 
-void CIRGenNVCUDARuntime::emitDeviceStub(CIRGenFunction &cgf, cir::FuncOp fn,
+void CIRGenNVCUDARuntime::emitDeviceStub(CIRGenFunction &cgf, cir::CIRCallableOpInterface fn,
                                          FunctionArgList &args) {
   if (auto globalOp =
           llvm::dyn_cast<cir::GlobalOp>(KernelHandles[fn.getSymName()])) {
@@ -275,7 +275,7 @@ void CIRGenNVCUDARuntime::emitDeviceStub(CIRGenFunction &cgf, cir::FuncOp fn,
     emitDeviceStubBodyLegacy(cgf, fn, args);
 }
 
-mlir::Operation *CIRGenNVCUDARuntime::getKernelHandle(cir::FuncOp fn,
+mlir::Operation *CIRGenNVCUDARuntime::getKernelHandle(cir::CIRCallableOpInterface fn,
                                                       GlobalDecl GD) {
 
   // Check if we already have a kernel handle for this function

clang/lib/CIR/CodeGen/CIRGenCUDARuntime.h

@@ -38,13 +38,13 @@ class CIRGenCUDARuntime {
   CIRGenCUDARuntime(CIRGenModule &cgm) : cgm(cgm) {}
   virtual ~CIRGenCUDARuntime();
 
-  virtual void emitDeviceStub(CIRGenFunction &cgf, cir::FuncOp fn,
+  virtual void emitDeviceStub(CIRGenFunction &cgf, cir::CIRCallableOpInterface fn,
                               FunctionArgList &args) = 0;
 
   virtual RValue emitCUDAKernelCallExpr(CIRGenFunction &cgf,
                                         const CUDAKernelCallExpr *expr,
                                         ReturnValueSlot retValue);
-  virtual mlir::Operation *getKernelHandle(cir::FuncOp fn, GlobalDecl GD) = 0;
+  virtual mlir::Operation *getKernelHandle(cir::CIRCallableOpInterface fn, GlobalDecl GD) = 0;
   virtual void internalizeDeviceSideVar(const VarDecl *d,
                                         cir::GlobalLinkageKind &linkage) = 0;
   /// Returns function or variable name on device side even if the current