[HIP] Emit target-id module flag

[yaxunl]

Summary

Separate this patch from https://reviews.llvm.org/D60620 since it depends on https://reviews.llvm.org/D80750

[tra]

clang/lib/CodeGen/CodeGenModule.cpp

603–605

Make it a TargetIdStr initializer and just pass TargetIdStr instead?

clang/test/CodeGenCXX/conditional-temporaries.cpp

67–68

This looks like an unrelated change.

clang/lib/CodeGen/CodeGenModule.cpp

if (LangOpts.CUDAIsDevice && getTriple().isNVPTX()) {
     // Indicate whether __nvvm_reflect should be configured to flush denormal
     // floating point values to 0.  (This corresponds to its "__CUDA_FTZ"
     // property.)
     getModule().addModuleFlag(llvm::Module::Override, "nvvm-reflect-ftz",
                               CodeGenOpts.FP32DenormalMode.Output !=
                                   llvm::DenormalMode::IEEE);
   }
 
+  if (auto TargetID = getTarget().getTargetID()) {
+    auto TargetIDStr = TargetID.getValue();
+    // Empty target ID is emitted as empty string in module flag.
+    getModule().addModuleFlag(
+        llvm::Module::MergeTargetID, "target-id",
+        llvm::MDString::get(
+            getModule().getContext(),
+            TargetIDStr == ""
+                ? TargetIDStr
+                : (Twine(getTriple().str()) + "-" + TargetIDStr).str()));
+  }
+
   // Emit OpenCL specific module metadata: OpenCL/SPIR version.
   if (LangOpts.OpenCL) {
     EmitOpenCLMetadata();
     // Emit SPIR version.
     if (getTriple().isSPIR()) {
       // SPIR v2.0 s2.12 - The SPIR version used by the module is stored in the
       // opencl.spir.version named metadata.
       // C++ is backwards compatible with OpenCL v2.0.

clang/test/CodeGenCUDA/target-id.hip

+// REQUIRES: clang-driver
+// REQUIRES: x86-registered-target
+// REQUIRES: amdgpu-registered-target
+
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device \
+// RUN:   -target-cpu gfx908 -target-feature +xnack \
+// RUN:   -target-feature -sram-ecc \
+// RUN:   -emit-llvm -o - %s | FileCheck %s
+
+#include "Inputs/cuda.h"
+
+// CHECK: !{i32 8, !"target-id", !"amdgcn-amd-amdhsa-gfx908:xnack+:sram-ecc-"}
+__global__ void foo() {}

clang/test/CodeGenCXX/conditional-temporaries.cpp

int getDtorCalls() {
   // AMDGCN-NEWPM-OPT: ret i32 [[RET]]
   return dtorcalls;
 }
 
 // CHECK-OPT-LABEL: define zeroext i1 @_Z7successv()
 bool success() {
   // CHECK-LEGACY-OPT: ret i1 true
   // X64-NEWPM-OPT: ret i1 true
-  // AMDGCN-NEWPM-OPT: [[CTORS:%.*]] = load i32, i32* addrspacecast (i32 addrspace(1)* @_ZN12_GLOBAL__N_19ctorcallsE to i32*), align 4, !tbaa !2
-  // AMDGCN-NEWPM-OPT: [[DTORS:%.*]] = load i32, i32* addrspacecast (i32 addrspace(1)* @_ZN12_GLOBAL__N_19dtorcallsE to i32*), align 4, !tbaa !2
+  // AMDGCN-NEWPM-OPT: [[CTORS:%.*]] = load i32, i32* addrspacecast (i32 addrspace(1)* @_ZN12_GLOBAL__N_19ctorcallsE to i32*), align 4
+  // AMDGCN-NEWPM-OPT: [[DTORS:%.*]] = load i32, i32* addrspacecast (i32 addrspace(1)* @_ZN12_GLOBAL__N_19dtorcallsE to i32*), align 4
   // AMDGCN-NEWPM-OPT: %cmp = icmp eq i32 [[CTORS]], [[DTORS]]
   // AMDGCN-NEWPM-OPT: ret i1 %cmp
   return ctorcalls == dtorcalls;
 }
 
 struct X { ~X(); int f(); };
 int g(int, int, int);
 // CHECK-LABEL: @_Z16lifetime_nontriv

clang/test/CodeGenOpenCL/target-id.cl

+// REQUIRES: clang-driver
+// REQUIRES: x86-registered-target
+// REQUIRES: amdgpu-registered-target
+
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa \
+// RUN:   -target-cpu gfx908 -target-feature +xnack \
+// RUN:   -target-feature -sram-ecc \
+// RUN:   -emit-llvm -o - %s | FileCheck -check-prefix=ID1 %s
+
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa \
+// RUN:   -target-cpu fiji \
+// RUN:   -emit-llvm -o - %s | FileCheck -check-prefix=ID2 %s
+
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa \
+// RUN:   -emit-llvm -o - %s | FileCheck -check-prefix=NONE %s
+
+// ID1: !{i32 8, !"target-id", !"amdgcn-amd-amdhsa-gfx908:sram-ecc-:xnack+"}
+// ID2: !{i32 8, !"target-id", !"amdgcn-amd-amdhsa-gfx803"}
+// NONE: !{i32 8, !"target-id", !""}
+
+kernel void foo() {}