[clang] [CIR] Emit target-cpu, target-features, and tune-cpu attrs on cir.func (PR #193458)

[Bruno Cardoso Lopes via cfe-commits]

================
@@ -0,0 +1,32 @@
+#include "../CodeGenCUDA/Inputs/cuda.h"
+
+// REQUIRES: amdgpu-registered-target
+
+// RUN: %clang_cc1 -triple=amdgcn-amd-amdhsa -x hip -fclangir \
+// RUN:            -fcuda-is-device -target-cpu gfx900 -emit-cir %s -o %t.cir
+// RUN: FileCheck --check-prefix=CIR %s --input-file=%t.cir
+
+// RUN: %clang_cc1 -triple=amdgcn-amd-amdhsa -x hip -fclangir \
+// RUN:            -fcuda-is-device -target-cpu gfx900 -emit-llvm %s -o %t.cir.ll
+// RUN: FileCheck --check-prefix=LLVM %s --input-file=%t.cir.ll
+
+// RUN: %clang_cc1 -triple=amdgcn-amd-amdhsa -x hip \
+// RUN:            -fcuda-is-device -target-cpu gfx900 -emit-llvm %s -o %t.ll
+// RUN: FileCheck --check-prefix=LLVM %s --input-file=%t.ll
+
+// AMDGPU with gfx900: target-cpu is set, target-features is absent.
+
+// CIR: cir.func{{.*}} @_Z6kernelv()
+// CIR-SAME: "cir.target-cpu" = "gfx900"
+// CIR-NOT: cir.target-features
+__global__ void kernel() {}
+
+// CIR: cir.func{{.*}} @_Z9device_fnv()
+// CIR-SAME: "cir.target-cpu" = "gfx900"
+// CIR-NOT: cir.target-features
+__device__ void device_fn() {}
----------------
bcardosolopes wrote:

This test only verifies that gfx900 with no extra features doesn't emit `cir.target-features`. Worth adding a case where features differ from the CPU baseline, e.g. `-target-cpu gfx900 -target-feature +sram-ecc`, to verify the delta actually emits only the differing feature.

https://github.com/llvm/llvm-project/pull/193458