Index: clang/lib/CodeGen/CodeGenModule.cpp
===================================================================
--- clang/lib/CodeGen/CodeGenModule.cpp
+++ clang/lib/CodeGen/CodeGenModule.cpp
@@ -593,6 +593,18 @@
                                   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();
Index: clang/test/CodeGenCUDA/target-id.hip
===================================================================
--- /dev/null
+++ clang/test/CodeGenCUDA/target-id.hip
@@ -0,0 +1,13 @@
+// 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() {}
Index: clang/test/CodeGenCXX/conditional-temporaries.cpp
===================================================================
--- clang/test/CodeGenCXX/conditional-temporaries.cpp
+++ clang/test/CodeGenCXX/conditional-temporaries.cpp
@@ -64,8 +64,8 @@
 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;
Index: clang/test/CodeGenOpenCL/target-id.cl
===================================================================
--- /dev/null
+++ clang/test/CodeGenOpenCL/target-id.cl
@@ -0,0 +1,21 @@
+// 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() {}