[CIR][CodeGen] Set CIR function calling conventions (#780)

This PR adds the counterparts of methods `SetFunctionAttributes`
`SetLLVMFunctionAttributes` from OG CodeGen, in order to set proper
calling conv for `cir.func` ops.

`spir-calling-conv.cl` is the dedicated test, while other OpenCL-related
tests are also updated. It removes previous workaround of incorrect
calling conv and better synchronizes with the original tests. (These
changes are not musts).

Author: seven-mile

clang/lib/CIR/CodeGen/CIRGenFunctionInfo.h

@@ -250,9 +250,17 @@ class CIRGenFunctionInfo final
     return getExtParameterInfos()[argIndex];
   }
 
-  /// getCallingConvention - REturn the user specified calling convention, which
+  /// getCallingConvention - Return the user specified calling convention, which
   /// has been translated into a CIR CC.
-  mlir::cir::CallingConv getCallingConvention() const { return CallingConvention; }
+  mlir::cir::CallingConv getCallingConvention() const {
+    return CallingConvention;
+  }
+
+  /// getEffectiveCallingConvention - Return the actual calling convention to
+  /// use, which may depend on the ABI.
+  mlir::cir::CallingConv getEffectiveCallingConvention() const {
+    return EffectiveCallingConvention;
+  }
 
   clang::CanQualType getReturnType() const { return getArgsBuffer()[0].type; }
 

clang/lib/CIR/CodeGen/CIRGenModule.cpp

@@ -2361,9 +2361,35 @@ void CIRGenModule::setExtraAttributesForFunc(FuncOp f,
       builder.getContext(), attrs.getDictionary(builder.getContext())));
 }
 
-void CIRGenModule::setFunctionAttributes(GlobalDecl GD, mlir::cir::FuncOp F,
-                                         bool IsIncompleteFunction,
-                                         bool IsThunk) {
+void CIRGenModule::setCIRFunctionAttributes(GlobalDecl GD,
+                                            const CIRGenFunctionInfo &info,
+                                            mlir::cir::FuncOp func,
+                                            bool isThunk) {
+  // TODO(cir): More logic of constructAttributeList is needed.
+  // NOTE(cir): Here we only need CallConv, so a call to constructAttributeList
+  // is omitted for simplicity.
+  mlir::cir::CallingConv callingConv = info.getEffectiveCallingConvention();
+
+  // TODO(cir): Check X86_VectorCall incompatibility with WinARM64EC
+
+  func.setCallingConv(callingConv);
+}
+
+void CIRGenModule::setFunctionAttributes(GlobalDecl globalDecl,
+                                         mlir::cir::FuncOp func,
+                                         bool isIncompleteFunction,
+                                         bool isThunk) {
+  // NOTE(cir): Original CodeGen checks if this is an intrinsic. In CIR we
+  // represent them in dedicated ops. The correct attributes are ensured during
+  // translation to LLVM. Thus, we don't need to check for them here.
+
+  if (!isIncompleteFunction) {
+    setCIRFunctionAttributes(globalDecl,
+                             getTypes().arrangeGlobalDeclaration(globalDecl),
+                             func, isThunk);
+  }
+
+  // TODO(cir): Complete the remaining part of the function.
   assert(!MissingFeatures::setFunctionAttributes());
 }
 

clang/lib/CIR/CodeGen/CIRGenModule.h

@@ -551,6 +551,10 @@ class CIRGenModule : public CIRGenTypeCache {
   void setFunctionAttributes(GlobalDecl GD, mlir::cir::FuncOp F,
                              bool IsIncompleteFunction, bool IsThunk);
 
+  /// Set the CIR function attributes (sext, zext, etc).
+  void setCIRFunctionAttributes(GlobalDecl GD, const CIRGenFunctionInfo &info,
+                                 mlir::cir::FuncOp func, bool isThunk);
+
   void buildGlobalDefinition(clang::GlobalDecl D,
                              mlir::Operation *Op = nullptr);
   void buildGlobalFunctionDefinition(clang::GlobalDecl D, mlir::Operation *Op);

clang/test/CIR/CodeGen/OpenCL/kernel-arg-metadata.cl

@@ -6,7 +6,7 @@
 __kernel void kernel_function() {}
 
 // CIR: #fn_attr[[ATTR:[0-9]*]] = {{.+}}cl.kernel_arg_metadata = #cir.cl.kernel_arg_metadata<addr_space = [], access_qual = [], type = [], base_type = [], type_qual = []>{{.+}}
-// CIR: cir.func @kernel_function() extra(#fn_attr[[ATTR]])
+// CIR: cir.func @kernel_function() cc(spir_kernel) extra(#fn_attr[[ATTR]])
 
-// LLVM: define {{.*}}void @kernel_function() {{[^{]+}} !kernel_arg_addr_space ![[MD:[0-9]+]] !kernel_arg_access_qual ![[MD]] !kernel_arg_type ![[MD]] !kernel_arg_base_type ![[MD]] !kernel_arg_type_qual ![[MD]] {
+// LLVM: define {{.*}}spir_kernel void @kernel_function() {{[^{]+}} !kernel_arg_addr_space ![[MD:[0-9]+]] !kernel_arg_access_qual ![[MD]] !kernel_arg_type ![[MD]] !kernel_arg_base_type ![[MD]] !kernel_arg_type_qual ![[MD]] {
 // LLVM: ![[MD]] = !{}

clang/test/CIR/CodeGen/OpenCL/kernel-attributes.cl

@@ -1,6 +1,6 @@
-// RUN: %clang_cc1 -fclangir -emit-cir -triple x86_64-unknown-linux-gnu %s -o %t.cir
+// RUN: %clang_cc1 -fclangir -emit-cir -triple spirv64-unknown-unknown %s -o %t.cir
 // RUN: FileCheck %s --input-file=%t.cir --check-prefix=CIR
-// RUN: %clang_cc1 -fclangir -emit-llvm -triple x86_64-unknown-linux-gnu %s -o %t.ll
+// RUN: %clang_cc1 -fclangir -emit-llvm -triple spirv64-unknown-unknown %s -o %t.ll
 // RUN: FileCheck %s --input-file=%t.ll --check-prefix=LLVM
 
 typedef unsigned int uint4 __attribute__((ext_vector_type(4)));
@@ -11,7 +11,7 @@ kernel  __attribute__((vec_type_hint(int))) __attribute__((reqd_work_group_size(
 // CIR-DAG: #fn_attr[[KERNEL1:[0-9]*]] = {{.+}}cl.kernel_metadata = #cir.cl.kernel_metadata<reqd_work_group_size = [1 : i32, 2 : i32, 4 : i32], vec_type_hint = !s32i, vec_type_hint_signedness = 1>{{.+}}
 // CIR-DAG: cir.func @kernel1{{.+}} extra(#fn_attr[[KERNEL1]])
 
-// LLVM-DAG: define{{.*}}@kernel1(i32 {{[^%]*}}%0) {{[^{]+}} !reqd_work_group_size ![[MD1_REQD_WG:[0-9]+]] !vec_type_hint ![[MD1_VEC_TYPE:[0-9]+]]
+// LLVM-DAG: define {{(dso_local )?}}spir_kernel void @kernel1(i32 {{[^%]*}}%0) {{[^{]+}} !reqd_work_group_size ![[MD1_REQD_WG:[0-9]+]] !vec_type_hint ![[MD1_VEC_TYPE:[0-9]+]]
 // LLVM-DAG: [[MD1_VEC_TYPE]] = !{i32 undef, i32 1}
 // LLVM-DAG: [[MD1_REQD_WG]] = !{i32 1, i32 2, i32 4}
 
@@ -21,7 +21,7 @@ kernel __attribute__((vec_type_hint(uint4))) __attribute__((work_group_size_hint
 // CIR-DAG: #fn_attr[[KERNEL2:[0-9]*]] = {{.+}}cl.kernel_metadata = #cir.cl.kernel_metadata<work_group_size_hint = [8 : i32, 16 : i32, 32 : i32], vec_type_hint = !cir.vector<!u32i x 4>, vec_type_hint_signedness = 0>{{.+}}
 // CIR-DAG: cir.func @kernel2{{.+}} extra(#fn_attr[[KERNEL2]])
 
-// LLVM-DAG: define{{.*}}@kernel2(i32 {{[^%]*}}%0) {{[^{]+}} !vec_type_hint ![[MD2_VEC_TYPE:[0-9]+]] !work_group_size_hint ![[MD2_WG_SIZE:[0-9]+]]
+// LLVM-DAG: define {{(dso_local )?}}spir_kernel void @kernel2(i32 {{[^%]*}}%0) {{[^{]+}} !vec_type_hint ![[MD2_VEC_TYPE:[0-9]+]] !work_group_size_hint ![[MD2_WG_SIZE:[0-9]+]]
 // LLVM-DAG: [[MD2_VEC_TYPE]] = !{<4 x i32> undef, i32 0}
 // LLVM-DAG: [[MD2_WG_SIZE]] = !{i32 8, i32 16, i32 32}
 
@@ -31,5 +31,5 @@ kernel __attribute__((intel_reqd_sub_group_size(8))) void kernel3(int a) {}
 // CIR-DAG: #fn_attr[[KERNEL3:[0-9]*]] = {{.+}}cl.kernel_metadata = #cir.cl.kernel_metadata<intel_reqd_sub_group_size = 8 : i32>{{.+}}
 // CIR-DAG: cir.func @kernel3{{.+}} extra(#fn_attr[[KERNEL3]])
 
-// LLVM-DAG: define{{.*}}@kernel3(i32 {{[^%]*}}%0) {{[^{]+}} !intel_reqd_sub_group_size ![[MD3_INTEL:[0-9]+]]
+// LLVM-DAG: define {{(dso_local )?}}spir_kernel void @kernel3(i32 {{[^%]*}}%0) {{[^{]+}} !intel_reqd_sub_group_size ![[MD3_INTEL:[0-9]+]]
 // LLVM-DAG: [[MD3_INTEL]] = !{i32 8}

clang/test/CIR/CodeGen/OpenCL/spir-calling-conv.cl

@@ -0,0 +1,20 @@
+// RUN: %clang_cc1 -fclangir %s -O0 -triple "spirv64-unknown-unknown" -emit-cir -o %t.cir
+// RUN: FileCheck %s --input-file=%t.cir --check-prefix=CIR
+// RUN: %clang_cc1 -fclangir %s -O0 -triple "spirv64-unknown-unknown" -emit-llvm -o %t.ll
+// RUN: FileCheck %s --input-file=%t.ll --check-prefix=LLVM
+
+// CIR: cir.func {{.*}}@get_dummy_id{{.*}} cc(spir_function)
+// LLVM-DAG: declare{{.*}} spir_func i32 @get_dummy_id(
+int get_dummy_id(int D);
+
+// CIR: cir.func {{.*}}@bar{{.*}} cc(spir_kernel)
+// LLVM-DAG: declare{{.*}} spir_kernel void @bar(
+kernel void bar(global int *A);
+
+// CIR: cir.func {{.*}}@foo{{.*}} cc(spir_kernel)
+// LLVM-DAG: define{{.*}} spir_kernel void @foo(
+kernel void foo(global int *A) {
+  int id = get_dummy_id(0);
+  A[id] = id;
+  bar(A);
+}

clang/test/CIR/CodeGen/OpenCL/spirv-target.cl

@@ -15,7 +15,7 @@ typedef struct {
 } my_st;
 
 // CIR-SPIRV64: cir.func @func(
-// LLVM-SPIRV64: @func(
+// LLVM-SPIRV64: define spir_kernel void @func(
 kernel void func(global long *arg) {
   int res1[sizeof(my_st)  == 24 ? 1 : -1]; // expected-no-diagnostics
   int res2[sizeof(void *) ==  8 ? 1 : -1]; // expected-no-diagnostics