From f7ac3f8a4fd8092c6dafa009180c9738eeb66de3 Mon Sep 17 00:00:00 2001 From: seven-mile Date: Sat, 3 Aug 2024 21:30:02 +0800 Subject: [PATCH] [CIR][CodeGen] Set CIR function calling conventions --- clang/lib/CIR/CodeGen/CIRGenFunctionInfo.h | 12 +++++-- clang/lib/CIR/CodeGen/CIRGenModule.cpp | 32 +++++++++++++++++-- clang/lib/CIR/CodeGen/CIRGenModule.h | 4 +++ .../CIR/CodeGen/OpenCL/kernel-arg-metadata.cl | 4 +-- .../CIR/CodeGen/OpenCL/kernel-attributes.cl | 10 +++--- .../CIR/CodeGen/OpenCL/spir-calling-conv.cl | 20 ++++++++++++ clang/test/CIR/CodeGen/OpenCL/spirv-target.cl | 2 +- 7 files changed, 71 insertions(+), 13 deletions(-) create mode 100644 clang/test/CIR/CodeGen/OpenCL/spir-calling-conv.cl diff --git a/clang/lib/CIR/CodeGen/CIRGenFunctionInfo.h b/clang/lib/CIR/CodeGen/CIRGenFunctionInfo.h index 4c9df914ee7c..a07f62fe28d7 100644 --- a/clang/lib/CIR/CodeGen/CIRGenFunctionInfo.h +++ b/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; } diff --git a/clang/lib/CIR/CodeGen/CIRGenModule.cpp b/clang/lib/CIR/CodeGen/CIRGenModule.cpp index 7b8815414e09..2cdeab0ab83b 100644 --- a/clang/lib/CIR/CodeGen/CIRGenModule.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenModule.cpp @@ -2360,9 +2360,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()); } diff --git a/clang/lib/CIR/CodeGen/CIRGenModule.h b/clang/lib/CIR/CodeGen/CIRGenModule.h index 895f6a54d403..b33e80d478c5 100644 --- a/clang/lib/CIR/CodeGen/CIRGenModule.h +++ b/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); diff --git a/clang/test/CIR/CodeGen/OpenCL/kernel-arg-metadata.cl b/clang/test/CIR/CodeGen/OpenCL/kernel-arg-metadata.cl index 7961e0e26244..ccc8ce967e50 100644 --- a/clang/test/CIR/CodeGen/OpenCL/kernel-arg-metadata.cl +++ b/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{{.+}} -// 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]] = !{} diff --git a/clang/test/CIR/CodeGen/OpenCL/kernel-attributes.cl b/clang/test/CIR/CodeGen/OpenCL/kernel-attributes.cl index 6badc7ce47ba..8a32f1d8088d 100644 --- a/clang/test/CIR/CodeGen/OpenCL/kernel-attributes.cl +++ b/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{{.+}} // 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, 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{{.+}} // 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} diff --git a/clang/test/CIR/CodeGen/OpenCL/spir-calling-conv.cl b/clang/test/CIR/CodeGen/OpenCL/spir-calling-conv.cl new file mode 100644 index 000000000000..96550f721bf5 --- /dev/null +++ b/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); +} diff --git a/clang/test/CIR/CodeGen/OpenCL/spirv-target.cl b/clang/test/CIR/CodeGen/OpenCL/spirv-target.cl index 523ffaf405e9..dadf4e6022b5 100644 --- a/clang/test/CIR/CodeGen/OpenCL/spirv-target.cl +++ b/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