[AMDGPU] Add gfx1030 target

[arsenm]

clang/lib/Basic/Targets/AMDGPU.cpp

174

Unrelated, but I think we should stop emitting the features in the target-features string if it's present on the base device

llvm/lib/Target/AMDGPU/AMDGPU.td

396

We don't use underscores in any other subtarget feature name. change to -?

500

Does this need its own feature apart from the instruction set?

512

mad-f32-insts may be a sufficient name?

llvm/lib/Target/AMDGPU/AMDGPUCodeGenPrepare.cpp

926–930

I prefer to use positive conditions and swap the select operand order

llvm/lib/Target/AMDGPU/SIISelLowering.cpp

1139

Volatile should not be necessary (at least not unconditionally)

llvm/lib/Target/AMDGPU/SIShrinkInstructions.cpp

223

It seems like a bug this would be missing?

llvm/lib/Target/AMDGPU/SOPInstructions.td

799–800

This is a problem. ReadCycleCounter should probably be marked IntrNoMem. I've also been thinking we need to introduce memory free versions of setreg pseudos that only touch the mode register

llvm/test/CodeGen/AMDGPU/GlobalISel/urem.i64.ll

3119

This shouldn't have changed

[arsenm]

llvm/lib/Target/AMDGPU/SOPInstructions.td

799–800

s_getreg should definitely not have side effects. it's already a problem that it is mayLoad. We don't get CSE of addrspacecast apertures because of this. Is marking readcyclecounter IntrInaccessibleMemOnly enough to fix this?

[rampitec]

llvm/lib/Target/AMDGPU/AMDGPU.td

396

OK, in a follow-up patch.

500

Yes, it has been dropped.

512

OK, in a follow-up patch.

llvm/lib/Target/AMDGPU/AMDGPUCodeGenPrepare.cpp

926–930

OK, in a follow-up patch.

llvm/lib/Target/AMDGPU/SIISelLowering.cpp

1139

We add it for other atomics here.

llvm/lib/Target/AMDGPU/SOPInstructions.td

799–800

Well, since it now reads timer... We can try to figure it out in subsequent patches.

llvm/test/CodeGen/AMDGPU/GlobalISel/urem.i64.ll

3119

This is changed because there is no -mcpu in the test.

[kzhuravl]

LGTM. This was reviewed downstream. Additional review feedback should be addressed in the follow up patch to simplify the merge downstream. Thank you.

[foad]

llvm/docs/AMDGPUUsage.rst

266–267

Seems odd to list xnack as a "supported fetaure" when the hardware doesn't support it.

[rampitec]

llvm/docs/AMDGPUUsage.rst

266–267

Yeah, it is off, but we may want to review the whole table.

[foad]

llvm/lib/Target/AMDGPU/AMDGPU.td

1245

The ! is obviously wrong in this definition, but if I remove it, all the tests still pass. So does this predicate actually control anything?

[rampitec]

llvm/lib/Target/AMDGPU/AMDGPU.td

1245

We don't select these, so only AssemblerPredicate is actually used. This is an obvious typo to fix.

clang/include/clang/Basic/Cuda.h

enum class CudaArch {
   GFX902,
   GFX904,
   GFX906,
   GFX908,
   GFX909,
   GFX1010,
   GFX1011,
   GFX1012,
+  GFX1030,
   LAST,
 };
 
 static inline bool IsNVIDIAGpuArch(CudaArch A) {
   return A >= CudaArch::SM_20 && A < CudaArch::GFX600;
 }
 
 static inline bool IsAMDGpuArch(CudaArch A) {

clang/lib/Basic/Targets/AMDGPU.cpp

bool AMDGPUTargetInfo::initFeatureMap(
     llvm::StringMap<bool> &Features, DiagnosticsEngine &Diags, StringRef CPU,
     const std::vector<std::string> &FeatureVec) const {
 
   using namespace llvm::AMDGPU;
 
   // XXX - What does the member GPU mean if device name string passed here?
   if (isAMDGCN(getTriple())) {
     switch (llvm::AMDGPU::parseArchAMDGCN(CPU)) {
+    case GK_GFX1030:
+      Features["ci-insts"] = true;
+      Features["dot1-insts"] = true;
+      Features["dot2-insts"] = true;
+      Features["dot5-insts"] = true;
+      Features["dot6-insts"] = true;
+      Features["dl-insts"] = true;
+      Features["flat-address-space"] = true;
+      Features["16-bit-insts"] = true;
+      Features["dpp"] = true;
+      Features["gfx8-insts"] = true;
+      Features["gfx9-insts"] = true;
+      Features["gfx10-insts"] = true;
+      Features["gfx10-3-insts"] = true;
+      Features["s-memrealtime"] = true;
+      break;
     case GK_GFX1012:
     case GK_GFX1011:
       Features["dot1-insts"] = true;
       Features["dot2-insts"] = true;
       Features["dot5-insts"] = true;
       Features["dot6-insts"] = true;
       LLVM_FALLTHROUGH;
     case GK_GFX1010:

clang/lib/Basic/Targets/NVPTX.cpp

std::string CUDAArchCode = [this] {
       case CudaArch::GFX902:
       case CudaArch::GFX904:
       case CudaArch::GFX906:
       case CudaArch::GFX908:
       case CudaArch::GFX909:
       case CudaArch::GFX1010:
       case CudaArch::GFX1011:
       case CudaArch::GFX1012:
+      case CudaArch::GFX1030:
       case CudaArch::LAST:
         break;
       case CudaArch::UNKNOWN:
         assert(false && "No GPU arch when compiling CUDA device code.");
         return "";
       case CudaArch::SM_20:
         return "200";
       case CudaArch::SM_21:

clang/lib/CodeGen/CGOpenMPRuntimeNVPTX.cpp

if (Clause->getClauseKind() == OMPC_unified_shared_memory) {
       case CudaArch::GFX902:
       case CudaArch::GFX904:
       case CudaArch::GFX906:
       case CudaArch::GFX908:
       case CudaArch::GFX909:
       case CudaArch::GFX1010:
       case CudaArch::GFX1011:
       case CudaArch::GFX1012:
+      case CudaArch::GFX1030:
       case CudaArch::UNKNOWN:
         break;
       case CudaArch::LAST:
         llvm_unreachable("Unexpected Cuda arch.");
       }
     }
   }
   CGOpenMPRuntime::processRequiresDirective(D);
static std::pair<unsigned, unsigned> getSMsBlocksPerSM(CodeGenModule &CGM) {
   case CudaArch::GFX902:
   case CudaArch::GFX904:
   case CudaArch::GFX906:
   case CudaArch::GFX908:
   case CudaArch::GFX909:
   case CudaArch::GFX1010:
   case CudaArch::GFX1011:
   case CudaArch::GFX1012:
+  case CudaArch::GFX1030:
   case CudaArch::UNKNOWN:
     break;
   case CudaArch::LAST:
     llvm_unreachable("Unexpected Cuda arch.");
   }
   llvm_unreachable("Unexpected NVPTX target without ptx feature.");
 }
 

clang/test/CodeGenOpenCL/amdgpu-features.cl

 // REQUIRES: amdgpu-registered-target
 
 // Check that appropriate features are defined for every supported AMDGPU
 // "-target" and "-mcpu" options.
 
 // RUN: %clang_cc1 -triple amdgcn -target-cpu gfx600 -S -emit-llvm -o - %s | FileCheck --check-prefix=GFX600 %s
 // RUN: %clang_cc1 -triple amdgcn -target-cpu gfx601 -S -emit-llvm -o - %s | FileCheck --check-prefix=GFX601 %s
 // RUN: %clang_cc1 -triple amdgcn -target-cpu gfx700 -S -emit-llvm -o - %s | FileCheck --check-prefix=GFX700 %s
 // RUN: %clang_cc1 -triple amdgcn -target-cpu gfx801 -S -emit-llvm -o - %s | FileCheck --check-prefix=GFX801 %s
 // RUN: %clang_cc1 -triple amdgcn -target-cpu gfx904 -S -emit-llvm -o - %s | FileCheck --check-prefix=GFX904 %s
 // RUN: %clang_cc1 -triple amdgcn -target-cpu gfx906 -S -emit-llvm -o - %s | FileCheck --check-prefix=GFX906 %s
 // RUN: %clang_cc1 -triple amdgcn -target-cpu gfx908 -S -emit-llvm -o - %s | FileCheck --check-prefix=GFX908 %s
 // RUN: %clang_cc1 -triple amdgcn -target-cpu gfx1010 -S -emit-llvm -o - %s | FileCheck --check-prefix=GFX1010 %s
 // RUN: %clang_cc1 -triple amdgcn -target-cpu gfx1011 -S -emit-llvm -o - %s | FileCheck --check-prefix=GFX1011 %s
 // RUN: %clang_cc1 -triple amdgcn -target-cpu gfx1012 -S -emit-llvm -o - %s | FileCheck --check-prefix=GFX1012 %s
+// RUN: %clang_cc1 -triple amdgcn -target-cpu gfx1030 -S -emit-llvm -o - %s | FileCheck --check-prefix=GFX1030 %s
 
 // GFX600-NOT: "target-features"
 // GFX601-NOT: "target-features"
 // GFX700: "target-features"="+ci-insts,+flat-address-space"
 // GFX801: "target-features"="+16-bit-insts,+ci-insts,+dpp,+flat-address-space,+gfx8-insts,+s-memrealtime"
 // GFX904: "target-features"="+16-bit-insts,+ci-insts,+dpp,+flat-address-space,+gfx8-insts,+gfx9-insts,+s-memrealtime"
 // GFX906: "target-features"="+16-bit-insts,+ci-insts,+dl-insts,+dot1-insts,+dot2-insts,+dpp,+flat-address-space,+gfx8-insts,+gfx9-insts,+s-memrealtime"
 // GFX908: "target-features"="+16-bit-insts,+ci-insts,+dl-insts,+dot1-insts,+dot2-insts,+dot3-insts,+dot4-insts,+dot5-insts,+dot6-insts,+dpp,+flat-address-space,+gfx8-insts,+gfx9-insts,+mai-insts,+s-memrealtime"
 // GFX1010: "target-features"="+16-bit-insts,+ci-insts,+dl-insts,+dpp,+flat-address-space,+gfx10-insts,+gfx8-insts,+gfx9-insts,+s-memrealtime"
 // GFX1011: "target-features"="+16-bit-insts,+ci-insts,+dl-insts,+dot1-insts,+dot2-insts,+dot5-insts,+dot6-insts,+dpp,+flat-address-space,+gfx10-insts,+gfx8-insts,+gfx9-insts,+s-memrealtime"
 // GFX1012: "target-features"="+16-bit-insts,+ci-insts,+dl-insts,+dot1-insts,+dot2-insts,+dot5-insts,+dot6-insts,+dpp,+flat-address-space,+gfx10-insts,+gfx8-insts,+gfx9-insts,+s-memrealtime"
+// GFX1030: "target-features"="+16-bit-insts,+ci-insts,+dl-insts,+dot1-insts,+dot2-insts,+dot5-insts,+dot6-insts,+dpp,+flat-address-space,+gfx10-3-insts,+gfx10-insts,+gfx8-insts,+gfx9-insts,+s-memrealtime"
 
 kernel void test() {}

clang/test/Driver/amdgpu-macros.cl

 // RUN: %clang -E -dM -target amdgcn -mcpu=gfx902 %s 2>&1 | FileCheck --check-prefixes=ARCH-GCN,GFX902 %s
 // RUN: %clang -E -dM -target amdgcn -mcpu=gfx904 %s 2>&1 | FileCheck --check-prefixes=ARCH-GCN,GFX904 %s
 // RUN: %clang -E -dM -target amdgcn -mcpu=gfx906 %s 2>&1 | FileCheck --check-prefixes=ARCH-GCN,GFX906 %s
 // RUN: %clang -E -dM -target amdgcn -mcpu=gfx908 %s 2>&1 | FileCheck --check-prefixes=ARCH-GCN,GFX908 %s
 // RUN: %clang -E -dM -target amdgcn -mcpu=gfx909 %s 2>&1 | FileCheck --check-prefixes=ARCH-GCN,GFX909 %s
 // RUN: %clang -E -dM -target amdgcn -mcpu=gfx1010 %s 2>&1 | FileCheck --check-prefixes=ARCH-GCN,GFX1010 %s
 // RUN: %clang -E -dM -target amdgcn -mcpu=gfx1011 %s 2>&1 | FileCheck --check-prefixes=ARCH-GCN,GFX1011 %s
 // RUN: %clang -E -dM -target amdgcn -mcpu=gfx1012 %s 2>&1 | FileCheck --check-prefixes=ARCH-GCN,GFX1012 %s
+// RUN: %clang -E -dM -target amdgcn -mcpu=gfx1030 %s 2>&1 | FileCheck --check-prefixes=ARCH-GCN,GFX1030 %s
 
 // GFX600-DAG: #define FP_FAST_FMA 1
 // GFX601-DAG: #define FP_FAST_FMA 1
 // GFX700-DAG: #define FP_FAST_FMA 1
 // GFX701-DAG: #define FP_FAST_FMA 1
 // GFX702-DAG: #define FP_FAST_FMA 1
 // GFX703-DAG: #define FP_FAST_FMA 1
 // GFX704-DAG: #define FP_FAST_FMA 1
 // GFX801-DAG: #define FP_FAST_FMA 1
 // GFX802-DAG: #define FP_FAST_FMA 1
 // GFX803-DAG: #define FP_FAST_FMA 1
 // GFX810-DAG: #define FP_FAST_FMA 1
 // GFX900-DAG: #define FP_FAST_FMA 1
 // GFX902-DAG: #define FP_FAST_FMA 1
 // GFX904-DAG: #define FP_FAST_FMA 1
 // GFX906-DAG: #define FP_FAST_FMA 1
 // GFX908-DAG: #define FP_FAST_FMA 1
 // GFX909-DAG: #define FP_FAST_FMA 1
 // GFX1010-DAG: #define FP_FAST_FMA 1
 // GFX1011-DAG: #define FP_FAST_FMA 1
 // GFX1012-DAG: #define FP_FAST_FMA 1
+// GFX1030-DAG: #define FP_FAST_FMA 1
 
 // GFX600-DAG: #define FP_FAST_FMAF 1
 // GFX601-NOT: #define FP_FAST_FMAF 1
 // GFX700-NOT: #define FP_FAST_FMAF 1
 // GFX701-DAG: #define FP_FAST_FMAF 1
 // GFX702-DAG: #define FP_FAST_FMAF 1
 // GFX703-NOT: #define FP_FAST_FMAF 1
 // GFX704-NOT: #define FP_FAST_FMAF 1
 // GFX801-DAG: #define FP_FAST_FMAF 1
 // GFX802-NOT: #define FP_FAST_FMAF 1
 // GFX803-NOT: #define FP_FAST_FMAF 1
 // GFX810-NOT: #define FP_FAST_FMAF 1
 // GFX900-DAG: #define FP_FAST_FMAF 1
 // GFX902-DAG: #define FP_FAST_FMAF 1
 // GFX904-DAG: #define FP_FAST_FMAF 1
 // GFX906-DAG: #define FP_FAST_FMAF 1
 // GFX908-DAG: #define FP_FAST_FMAF 1
 // GFX909-DAG: #define FP_FAST_FMAF 1
 // GFX1010-DAG: #define FP_FAST_FMAF 1
 // GFX1011-DAG: #define FP_FAST_FMAF 1
 // GFX1012-DAG: #define FP_FAST_FMAF 1
+// GFX1030-DAG: #define FP_FAST_FMAF 1
 
 // ARCH-GCN-DAG: #define __AMDGCN__ 1
 // ARCH-GCN-DAG: #define __AMDGPU__ 1
 // ARCH-GCN-DAG: #define __AMD__ 1
 
 // GFX600-DAG: #define __HAS_FMAF__ 1
 // GFX601-DAG: #define __HAS_FMAF__ 1
 // GFX700-DAG: #define __HAS_FMAF__ 1

 // GFX902-DAG: #define __HAS_FMAF__ 1
 // GFX904-DAG: #define __HAS_FMAF__ 1
 // GFX906-DAG: #define __HAS_FMAF__ 1
 // GFX908-DAG: #define __HAS_FMAF__ 1
 // GFX909-DAG: #define __HAS_FMAF__ 1
 // GFX1010-DAG: #define __HAS_FMAF__ 1
 // GFX1011-DAG: #define __HAS_FMAF__ 1
 // GFX1012-DAG: #define __HAS_FMAF__ 1
+// GFX1030-DAG: #define __HAS_FMAF__ 1
 
 // GFX600-DAG: #define __HAS_FP64__ 1
 // GFX601-DAG: #define __HAS_FP64__ 1
 // GFX700-DAG: #define __HAS_FP64__ 1
 // GFX701-DAG: #define __HAS_FP64__ 1
 // GFX702-DAG: #define __HAS_FP64__ 1
 // GFX703-DAG: #define __HAS_FP64__ 1
 // GFX704-DAG: #define __HAS_FP64__ 1
 // GFX801-DAG: #define __HAS_FP64__ 1
 // GFX802-DAG: #define __HAS_FP64__ 1
 // GFX803-DAG: #define __HAS_FP64__ 1
 // GFX810-DAG: #define __HAS_FP64__ 1
 // GFX900-DAG: #define __HAS_FP64__ 1
 // GFX902-DAG: #define __HAS_FP64__ 1
 // GFX904-DAG: #define __HAS_FP64__ 1
 // GFX906-DAG: #define __HAS_FP64__ 1
 // GFX908-DAG: #define __HAS_FP64__ 1
 // GFX909-DAG: #define __HAS_FP64__ 1
 // GFX1010-DAG: #define __HAS_FP64__ 1
 // GFX1011-DAG: #define __HAS_FP64__ 1
 // GFX1012-DAG: #define __HAS_FP64__ 1
+// GFX1030-DAG: #define __HAS_FP64__ 1
 
 // GFX600-DAG: #define __HAS_LDEXPF__ 1
 // GFX601-DAG: #define __HAS_LDEXPF__ 1
 // GFX700-DAG: #define __HAS_LDEXPF__ 1
 // GFX701-DAG: #define __HAS_LDEXPF__ 1
 // GFX702-DAG: #define __HAS_LDEXPF__ 1
 // GFX703-DAG: #define __HAS_LDEXPF__ 1
 // GFX704-DAG: #define __HAS_LDEXPF__ 1
 // GFX801-DAG: #define __HAS_LDEXPF__ 1
 // GFX802-DAG: #define __HAS_LDEXPF__ 1
 // GFX803-DAG: #define __HAS_LDEXPF__ 1
 // GFX810-DAG: #define __HAS_LDEXPF__ 1
 // GFX900-DAG: #define __HAS_LDEXPF__ 1
 // GFX902-DAG: #define __HAS_LDEXPF__ 1
 // GFX904-DAG: #define __HAS_LDEXPF__ 1
 // GFX906-DAG: #define __HAS_LDEXPF__ 1
 // GFX908-DAG: #define __HAS_LDEXPF__ 1
 // GFX909-DAG: #define __HAS_LDEXPF__ 1
 // GFX1010-DAG: #define __HAS_LDEXPF__ 1
 // GFX1011-DAG: #define __HAS_LDEXPF__ 1
 // GFX1012-DAG: #define __HAS_LDEXPF__ 1
+// GFX1030-DAG: #define __HAS_LDEXPF__ 1
 
 // GFX600-DAG: #define __gfx600__ 1
 // GFX601-DAG: #define __gfx601__ 1
 // GFX700-DAG: #define __gfx700__ 1
 // GFX701-DAG: #define __gfx701__ 1
 // GFX702-DAG: #define __gfx702__ 1
 // GFX703-DAG: #define __gfx703__ 1
 // GFX704-DAG: #define __gfx704__ 1
 // GFX801-DAG: #define __gfx801__ 1
 // GFX802-DAG: #define __gfx802__ 1
 // GFX803-DAG: #define __gfx803__ 1
 // GFX810-DAG: #define __gfx810__ 1
 // GFX900-DAG: #define __gfx900__ 1
 // GFX902-DAG: #define __gfx902__ 1
 // GFX904-DAG: #define __gfx904__ 1
 // GFX906-DAG: #define __gfx906__ 1
 // GFX908-DAG: #define __gfx908__ 1
 // GFX909-DAG: #define __gfx909__ 1
 // GFX1010-DAG: #define __gfx1010__ 1
 // GFX1011-DAG: #define __gfx1011__ 1
 // GFX1012-DAG: #define __gfx1012__ 1
+// GFX1030-DAG: #define __gfx1030__ 1

clang/test/Driver/amdgpu-mcpu.cl

 // RUN: %clang -### -target amdgcn -mcpu=gfx902 %s 2>&1 | FileCheck --check-prefix=GFX902 %s
 // RUN: %clang -### -target amdgcn -mcpu=gfx904 %s 2>&1 | FileCheck --check-prefix=GFX904 %s
 // RUN: %clang -### -target amdgcn -mcpu=gfx906 %s 2>&1 | FileCheck --check-prefix=GFX906 %s
 // RUN: %clang -### -target amdgcn -mcpu=gfx908 %s 2>&1 | FileCheck --check-prefix=GFX908 %s
 // RUN: %clang -### -target amdgcn -mcpu=gfx909 %s 2>&1 | FileCheck --check-prefix=GFX909 %s
 // RUN: %clang -### -target amdgcn -mcpu=gfx1010 %s 2>&1 | FileCheck --check-prefix=GFX1010 %s
 // RUN: %clang -### -target amdgcn -mcpu=gfx1011 %s 2>&1 | FileCheck --check-prefix=GFX1011 %s
 // RUN: %clang -### -target amdgcn -mcpu=gfx1012 %s 2>&1 | FileCheck --check-prefix=GFX1012 %s
+// RUN: %clang -### -target amdgcn -mcpu=gfx1030 %s 2>&1 | FileCheck --check-prefix=GFX1030 %s
 
 // GCNDEFAULT-NOT: -target-cpu
 // GFX600:    "-target-cpu" "gfx600"
 // TAHITI:    "-target-cpu" "tahiti"
 // GFX601:    "-target-cpu" "gfx601"
 // HAINAN:    "-target-cpu" "hainan"
 // OLAND:     "-target-cpu" "oland"
 // PITCAIRN:  "-target-cpu" "pitcairn"

 // GFX902:    "-target-cpu" "gfx902"
 // GFX904:    "-target-cpu" "gfx904"
 // GFX906:    "-target-cpu" "gfx906"
 // GFX908:    "-target-cpu" "gfx908"
 // GFX909:    "-target-cpu" "gfx909"
 // GFX1010:   "-target-cpu" "gfx1010"
 // GFX1011:   "-target-cpu" "gfx1011"
 // GFX1012:   "-target-cpu" "gfx1012"
+// GFX1030:   "-target-cpu" "gfx1030"

llvm/docs/AMDGPUUsage.rst

``gfx1012``                 ``amdgcn``   dGPU  - xnack                   *TBA*
                                                       [off]
                                                     - wavefrontsize64
                                                       [off]
                                                     - cumode
                                                       [off]
                                                                               .. TODO::
                                                                                  Add product
                                                                                  names.
+     ``gfx1030``                 ``amdgcn``   dGPU  - xnack                   *TBA*
+                                                      [off]
+                                                    - wavefrontsize64
+                                                      [off]
+                                                    - cumode
+                                                      [off]
+                                                                              .. TODO
+                                                                                 Add product
+                                                                                 names.
      =========== =============== ============ ===== ================= ======= ======================
 
 .. _amdgpu-target-features:
 
 Target Features
 ---------------
 
 Target features control how code is generated to support certain
.. table:: AMDGPU ``EF_AMDGPU_MACH`` Values
      ``EF_AMDGPU_MACH_AMDGCN_GFX904``  0x02e      ``gfx904``
      ``EF_AMDGPU_MACH_AMDGCN_GFX906``  0x02f      ``gfx906``
      ``EF_AMDGPU_MACH_AMDGCN_GFX908``  0x030      ``gfx908``
      ``EF_AMDGPU_MACH_AMDGCN_GFX909``  0x031      ``gfx909``
      *reserved*                        0x032      Reserved.
      ``EF_AMDGPU_MACH_AMDGCN_GFX1010`` 0x033      ``gfx1010``
      ``EF_AMDGPU_MACH_AMDGCN_GFX1011`` 0x034      ``gfx1011``
      ``EF_AMDGPU_MACH_AMDGCN_GFX1012`` 0x035      ``gfx1012``
+     ``EF_AMDGPU_MACH_AMDGCN_GFX1030`` 0x036      ``gfx1030``
      ================================= ========== =============================
 
 Sections
 --------
 
 An AMDGPU target ELF code object has the standard ELF sections which include:
 
   .. table:: AMDGPU ELF Sections

llvm/include/llvm/BinaryFormat/ELF.h

enum : unsigned {
   EF_AMDGPU_MACH_AMDGCN_GFX904 = 0x02e,
   EF_AMDGPU_MACH_AMDGCN_GFX906 = 0x02f,
   EF_AMDGPU_MACH_AMDGCN_GFX908 = 0x030,
   EF_AMDGPU_MACH_AMDGCN_GFX909 = 0x031,
   // AMDGCN GFX10.
   EF_AMDGPU_MACH_AMDGCN_GFX1010 = 0x033,
   EF_AMDGPU_MACH_AMDGCN_GFX1011 = 0x034,
   EF_AMDGPU_MACH_AMDGCN_GFX1012 = 0x035,
+  EF_AMDGPU_MACH_AMDGCN_GFX1030 = 0x036,
 
   // Reserved for AMDGCN-based processors.
   EF_AMDGPU_MACH_AMDGCN_RESERVED0 = 0x027,
   EF_AMDGPU_MACH_AMDGCN_RESERVED1 = 0x032,
 
   // First/last AMDGCN-based processors.
   EF_AMDGPU_MACH_AMDGCN_FIRST = EF_AMDGPU_MACH_AMDGCN_GFX600,
-  EF_AMDGPU_MACH_AMDGCN_LAST = EF_AMDGPU_MACH_AMDGCN_GFX1012,
+  EF_AMDGPU_MACH_AMDGCN_LAST = EF_AMDGPU_MACH_AMDGCN_GFX1030,
 
   // Indicates if the "xnack" target feature is enabled for all code contained
   // in the object.
   EF_AMDGPU_XNACK = 0x100,
   // Indicates if the "sram-ecc" target feature is enabled for all code
   // contained in the object.
   EF_AMDGPU_SRAM_ECC = 0x200,
 };

llvm/include/llvm/IR/IntrinsicsAMDGPU.td

defset list<AMDGPUImageDimIntrinsic> AMDGPUImageDimIntrinsics = {
 
   defm int_amdgcn_image_store : AMDGPUImageDimIntrinsicsAll<
               "STORE", [], [AMDGPUArg<llvm_anyfloat_ty, "vdata">],
               [IntrWriteMem], [SDNPMemOperand]>;
   defm int_amdgcn_image_store_mip : AMDGPUImageDimIntrinsicsNoMsaa<
               "STORE_MIP", [], [AMDGPUArg<llvm_anyfloat_ty, "vdata">],
               [IntrWriteMem], [SDNPMemOperand], 1>;
 
+  defm int_amdgcn_image_msaa_load
+    : AMDGPUImageDimIntrinsicsAll<"MSAA_LOAD", [llvm_any_ty], [], [IntrReadMem],
+                                  [SDNPMemOperand]>,
+      AMDGPUImageDMaskIntrinsic;
+
   //////////////////////////////////////////////////////////////////////////
   // sample and getlod intrinsics
   //////////////////////////////////////////////////////////////////////////
   multiclass AMDGPUImageDimSampleDims<string opmod,
                                       AMDGPUSampleVariant sample,
                                       bit NoMem = 0> {
     foreach dim = AMDGPUDims.NoMsaa in {
       def !strconcat(NAME, "_", dim.Name) : AMDGPUImageDimIntrinsic<
def int_amdgcn_buffer_atomic_cmpswap : Intrinsic<
    llvm_i32_ty,       // cmp(VGPR)
    llvm_v4i32_ty,     // rsrc(SGPR)
    llvm_i32_ty,       // vindex(VGPR)
    llvm_i32_ty,       // offset(SGPR/VGPR/imm)
    llvm_i1_ty],       // slc(imm)
   [ImmArg<ArgIndex<5>>], "", [SDNPMemOperand]>,
   AMDGPURsrcIntrinsic<2, 0>;
 
+def int_amdgcn_buffer_atomic_csub : AMDGPUBufferAtomic;
 } // defset AMDGPUBufferIntrinsics
 
 // Uses that do not set the done bit should set IntrWriteMem on the
 // call site.
 def int_amdgcn_exp : Intrinsic <[], [
   llvm_i32_ty,       // tgt,
   llvm_i32_ty,       // en
   llvm_any_ty,       // src0 (f32 or i32)
def int_amdgcn_mov_dpp8 :
   Intrinsic<[llvm_anyint_ty],
             [LLVMMatchType<0>, llvm_i32_ty],
             [IntrNoMem, IntrConvergent, ImmArg<ArgIndex<1>>]>;
 
 def int_amdgcn_s_get_waveid_in_workgroup :
   GCCBuiltin<"__builtin_amdgcn_s_get_waveid_in_workgroup">,
   Intrinsic<[llvm_i32_ty], [], [IntrReadMem, IntrInaccessibleMemOnly]>;
 
+class AMDGPUGlobalAtomicRtn<LLVMType vt> : Intrinsic <
+  [vt],
+  [llvm_anyptr_ty,    // vaddr
+   vt],               // vdata(VGPR)
+  [IntrArgMemOnly, NoCapture<ArgIndex<0>>], "", [SDNPMemOperand]>;
+
+def int_amdgcn_global_atomic_csub : AMDGPUGlobalAtomicRtn<llvm_i32_ty>;
+
 //===----------------------------------------------------------------------===//
 // Deep learning intrinsics.
 //===----------------------------------------------------------------------===//
 
 // f32 %r = llvm.amdgcn.fdot2(v2f16 %a, v2f16 %b, f32 %c, i1 %clamp)
 //   %r = %a[0] * %b[0] + %a[1] * %b[1] + %c
 def int_amdgcn_fdot2 :
   GCCBuiltin<"__builtin_amdgcn_fdot2">,

llvm/include/llvm/Support/TargetParser.h

enum GPUKind : uint32_t {
   GK_GFX904 = 62,
   GK_GFX906 = 63,
   GK_GFX908 = 64,
   GK_GFX909 = 65,
 
   GK_GFX1010 = 71,
   GK_GFX1011 = 72,
   GK_GFX1012 = 73,
+  GK_GFX1030 = 75,
 
   GK_AMDGCN_FIRST = GK_GFX600,
-  GK_AMDGCN_LAST = GK_GFX1012,
+  GK_AMDGCN_LAST = GK_GFX1030,
 };
 
 /// Instruction set architecture version.
 struct IsaVersion {
   unsigned Major;
   unsigned Minor;
   unsigned Stepping;
 };

llvm/lib/ObjectYAML/ELFYAML.cpp

case ELF::EM_AMDGPU:
     BCaseMask(EF_AMDGPU_MACH_AMDGCN_GFX902, EF_AMDGPU_MACH);
     BCaseMask(EF_AMDGPU_MACH_AMDGCN_GFX904, EF_AMDGPU_MACH);
     BCaseMask(EF_AMDGPU_MACH_AMDGCN_GFX906, EF_AMDGPU_MACH);
     BCaseMask(EF_AMDGPU_MACH_AMDGCN_GFX908, EF_AMDGPU_MACH);
     BCaseMask(EF_AMDGPU_MACH_AMDGCN_GFX909, EF_AMDGPU_MACH);
     BCaseMask(EF_AMDGPU_MACH_AMDGCN_GFX1010, EF_AMDGPU_MACH);
     BCaseMask(EF_AMDGPU_MACH_AMDGCN_GFX1011, EF_AMDGPU_MACH);
     BCaseMask(EF_AMDGPU_MACH_AMDGCN_GFX1012, EF_AMDGPU_MACH);
+    BCaseMask(EF_AMDGPU_MACH_AMDGCN_GFX1030, EF_AMDGPU_MACH);
     BCase(EF_AMDGPU_XNACK);
     BCase(EF_AMDGPU_SRAM_ECC);
     break;
   case ELF::EM_X86_64:
     break;
   default:
     llvm_unreachable("Unsupported architecture");
   }

llvm/lib/Support/TargetParser.cpp

constexpr GPUInfo R600GPUs[26] = {
   {{"caicos"},  {"caicos"},  GK_CAICOS,  FEATURE_NONE },
   {{"aruba"},   {"cayman"},  GK_CAYMAN,  FEATURE_FMA  },
   {{"cayman"},  {"cayman"},  GK_CAYMAN,  FEATURE_FMA  },
   {{"turks"},   {"turks"},   GK_TURKS,   FEATURE_NONE }
 };
 
 // This table should be sorted by the value of GPUKind
 // Don't bother listing the implicitly true features
-constexpr GPUInfo AMDGCNGPUs[37] = {
+constexpr GPUInfo AMDGCNGPUs[38] = {
   // Name         Canonical    Kind        Features
   //              Name
   {{"gfx600"},    {"gfx600"},  GK_GFX600,  FEATURE_FAST_FMA_F32},
   {{"tahiti"},    {"gfx600"},  GK_GFX600,  FEATURE_FAST_FMA_F32},
   {{"gfx601"},    {"gfx601"},  GK_GFX601,  FEATURE_NONE},
   {{"hainan"},    {"gfx601"},  GK_GFX601,  FEATURE_NONE},
   {{"oland"},     {"gfx601"},  GK_GFX601,  FEATURE_NONE},
   {{"pitcairn"},  {"gfx601"},  GK_GFX601,  FEATURE_NONE},
constexpr GPUInfo AMDGCNGPUs[38] = {
   {{"gfx902"},    {"gfx902"},  GK_GFX902,  FEATURE_FAST_FMA_F32|FEATURE_FAST_DENORMAL_F32},
   {{"gfx904"},    {"gfx904"},  GK_GFX904,  FEATURE_FAST_FMA_F32|FEATURE_FAST_DENORMAL_F32},
   {{"gfx906"},    {"gfx906"},  GK_GFX906,  FEATURE_FAST_FMA_F32|FEATURE_FAST_DENORMAL_F32},
   {{"gfx908"},    {"gfx908"},  GK_GFX908,  FEATURE_FAST_FMA_F32|FEATURE_FAST_DENORMAL_F32},
   {{"gfx909"},    {"gfx909"},  GK_GFX909,  FEATURE_FAST_FMA_F32|FEATURE_FAST_DENORMAL_F32},
   {{"gfx1010"},   {"gfx1010"}, GK_GFX1010, FEATURE_FAST_FMA_F32|FEATURE_FAST_DENORMAL_F32|FEATURE_WAVE32},
   {{"gfx1011"},   {"gfx1011"}, GK_GFX1011, FEATURE_FAST_FMA_F32|FEATURE_FAST_DENORMAL_F32|FEATURE_WAVE32},
   {{"gfx1012"},   {"gfx1012"}, GK_GFX1012, FEATURE_FAST_FMA_F32|FEATURE_FAST_DENORMAL_F32|FEATURE_WAVE32},
+  {{"gfx1030"},   {"gfx1030"}, GK_GFX1030, FEATURE_FAST_FMA_F32|FEATURE_FAST_DENORMAL_F32|FEATURE_WAVE32},
 };
 
 const GPUInfo *getArchEntry(AMDGPU::GPUKind AK, ArrayRef<GPUInfo> Table) {
   GPUInfo Search = { {""}, {""}, AK, AMDGPU::FEATURE_NONE };
 
   auto I = std::lower_bound(Table.begin(), Table.end(), Search,
     [](const GPUInfo &A, const GPUInfo &B) {
       return A.Kind < B.Kind;
AMDGPU::IsaVersion AMDGPU::getIsaVersion(StringRef GPU) {
   case GK_GFX902:  return {9, 0, 2};
   case GK_GFX904:  return {9, 0, 4};
   case GK_GFX906:  return {9, 0, 6};
   case GK_GFX908:  return {9, 0, 8};
   case GK_GFX909:  return {9, 0, 9};
   case GK_GFX1010: return {10, 1, 0};
   case GK_GFX1011: return {10, 1, 1};
   case GK_GFX1012: return {10, 1, 2};
+  case GK_GFX1030: return {10, 3, 0};
   default:         return {0, 0, 0};
   }
 }

llvm/lib/Target/AMDGPU/AMDGPU.td

 >;
 
 def FeatureGFX10Insts : SubtargetFeature<"gfx10-insts",
   "GFX10Insts",
   "true",
   "Additional instructions for GFX10+"
 >;
 
+def FeatureGFX10_3Insts : SubtargetFeature<"gfx10-3-insts",
+  "GFX10_3Insts",
+  "true",
+  "Additional instructions for GFX10.3"
+>;
+
 def FeatureGFX7GFX8GFX9Insts : SubtargetFeature<"gfx7-gfx8-gfx9-insts",
   "GFX7GFX8GFX9Insts",
   "true",
   "Instructions shared in GFX7, GFX8, GFX9"
 >;
 
 def FeatureSMemRealTime : SubtargetFeature<"s-memrealtime",
   "HasSMemRealTime",

 >;
 
 def FeatureNSAEncoding : SubtargetFeature<"nsa-encoding",
   "HasNSAEncoding",
   "true",
   "Support NSA encoding for image instructions"
 >;
 
+def FeatureGFX10_BEncoding : SubtargetFeature<"gfx10_b-encoding",
+  "GFX10_BEncoding",
+  "true",
+  "Encoding format GFX10_B"
+>;
+
 def FeatureIntClamp : SubtargetFeature<"int-clamp-insts",
   "HasIntClamp",
   "true",
   "Support clamp for integer destination"
 >;
 
 def FeatureUnpackedD16VMem : SubtargetFeature<"unpacked-d16-vmem",
   "HasUnpackedD16VMem",

 >;
 
 def FeatureVscnt : SubtargetFeature<"vscnt",
   "HasVscnt",
   "true",
   "Has separate store vscnt counter"
 >;
 
+def FeatureGetWaveIdInst : SubtargetFeature<"get-wave-id-inst",
+  "HasGetWaveIdInst",
+  "true",
+  "Has s_get_waveid_in_workgroup instruction"
+>;
+
+def FeatureSMemTimeInst : SubtargetFeature<"s-memtime-inst",
+  "HasSMemTimeInst",
+  "true",
+  "Has s_memtime instruction"
+>;
+
+def FeatureMadMacF32Insts : SubtargetFeature<"mad-mac-f32-insts",
+  "HasMadMacF32Insts",
+  "true",
+  "Has v_mad_f32/v_mac_f32/v_madak_f32/v_madmk_f32 instructions"
+>;
+
+def FeatureDsSrc2Insts : SubtargetFeature<"ds-src2-insts",
+  "HasDsSrc2Insts",
+  "true",
+  "Has ds_*_src2 instructions"
+>;
+
 def FeatureRegisterBanking : SubtargetFeature<"register-banking",
   "HasRegisterBanking",
   "true",
   "Has register banking"
 >;
 
 def FeatureVOP3Literal : SubtargetFeature<"vop3-literal",
   "HasVOP3Literal",

 class GCNSubtargetFeatureGeneration <string Value,
                                      string FeatureName,
                                      list<SubtargetFeature> Implies> :
         SubtargetFeatureGeneration <Value, FeatureName, "GCNSubtarget", Implies>;
 
 def FeatureSouthernIslands : GCNSubtargetFeatureGeneration<"SOUTHERN_ISLANDS",
     "southern-islands",
   [FeatureFP64, FeatureLocalMemorySize32768, FeatureMIMG_R128,
-  FeatureWavefrontSize64,
-  FeatureLDSBankCount32, FeatureMovrel, FeatureTrigReducedRange,
-  FeatureDoesNotSupportSRAMECC, FeatureDoesNotSupportXNACK]
+  FeatureWavefrontSize64, FeatureSMemTimeInst, FeatureMadMacF32Insts,
+  FeatureDsSrc2Insts, FeatureLDSBankCount32, FeatureMovrel,
+  FeatureTrigReducedRange, FeatureDoesNotSupportSRAMECC,
+  FeatureDoesNotSupportXNACK]
 >;
 
 def FeatureSeaIslands : GCNSubtargetFeatureGeneration<"SEA_ISLANDS",
     "sea-islands",
   [FeatureFP64, FeatureLocalMemorySize65536, FeatureMIMG_R128,
   FeatureWavefrontSize64, FeatureFlatAddressSpace,
   FeatureCIInsts, FeatureMovrel, FeatureTrigReducedRange,
-  FeatureGFX7GFX8GFX9Insts, FeatureDoesNotSupportSRAMECC]
+  FeatureGFX7GFX8GFX9Insts, FeatureSMemTimeInst, FeatureMadMacF32Insts,
+  FeatureDsSrc2Insts, FeatureDoesNotSupportSRAMECC]
 >;
 
 def FeatureVolcanicIslands : GCNSubtargetFeatureGeneration<"VOLCANIC_ISLANDS",
   "volcanic-islands",
   [FeatureFP64, FeatureLocalMemorySize65536, FeatureMIMG_R128,
    FeatureWavefrontSize64, FeatureFlatAddressSpace,
    FeatureGCN3Encoding, FeatureCIInsts, Feature16BitInsts,
    FeatureSMemRealTime, FeatureVGPRIndexMode, FeatureMovrel,
    FeatureScalarStores, FeatureInv2PiInlineImm,
    FeatureSDWA, FeatureSDWAOutModsVOPC, FeatureSDWAMac, FeatureDPP,
-   FeatureIntClamp, FeatureTrigReducedRange, FeatureDoesNotSupportSRAMECC,
-   FeatureGFX8Insts, FeatureGFX7GFX8GFX9Insts, FeatureFastDenormalF32
+   FeatureIntClamp, FeatureTrigReducedRange, FeatureGFX8Insts,
+   FeatureGFX7GFX8GFX9Insts, FeatureSMemTimeInst, FeatureMadMacF32Insts,
+   FeatureDsSrc2Insts, FeatureDoesNotSupportSRAMECC, FeatureFastDenormalF32
   ]
 >;
 
 def FeatureGFX9 : GCNSubtargetFeatureGeneration<"GFX9",
   "gfx9",
   [FeatureFP64, FeatureLocalMemorySize65536,
    FeatureWavefrontSize64, FeatureFlatAddressSpace,
    FeatureGCN3Encoding, FeatureCIInsts, Feature16BitInsts,
    FeatureSMemRealTime, FeatureScalarStores, FeatureInv2PiInlineImm,
    FeatureApertureRegs, FeatureGFX9Insts, FeatureVOP3P, FeatureVGPRIndexMode,
    FeatureFastFMAF32, FeatureDPP, FeatureIntClamp,
    FeatureSDWA, FeatureSDWAOmod, FeatureSDWAScalar, FeatureSDWASdst,
    FeatureFlatInstOffsets, FeatureFlatGlobalInsts, FeatureFlatScratchInsts,
    FeatureAddNoCarryInsts, FeatureGFX8Insts, FeatureGFX7GFX8GFX9Insts,
    FeatureScalarFlatScratchInsts, FeatureScalarAtomics, FeatureR128A16,
-   FeatureFastDenormalF32]
+   FeatureSMemTimeInst, FeatureMadMacF32Insts, FeatureDsSrc2Insts,
+   FeatureFastDenormalF32
+  ]
 >;
 
 def FeatureGFX10 : GCNSubtargetFeatureGeneration<"GFX10",
   "gfx10",
   [FeatureFP64, FeatureLocalMemorySize65536, FeatureMIMG_R128,
    FeatureFlatAddressSpace,
    FeatureCIInsts, Feature16BitInsts,
    FeatureSMemRealTime, FeatureInv2PiInlineImm,
!listconcat(FeatureGroup.GFX10_1_Bugs,
     [FeatureGFX10,
      FeatureLDSBankCount32,
      FeatureDLInsts,
      FeatureNSAEncoding,
      FeatureWavefrontSize32,
      FeatureScalarStores,
      FeatureScalarAtomics,
      FeatureScalarFlatScratchInsts,
+     FeatureGetWaveIdInst,
+     FeatureSMemTimeInst,
+     FeatureMadMacF32Insts,
+     FeatureDsSrc2Insts,
      FeatureLdsMisalignedBug,
      FeatureDoesNotSupportXNACK,
      FeatureCodeObjectV3])>;
 
 def FeatureISAVersion10_1_1 : FeatureSet<
   !listconcat(FeatureGroup.GFX10_1_Bugs,
     [FeatureGFX10,
      FeatureLDSBankCount32,
      FeatureDLInsts,
      FeatureDot1Insts,
      FeatureDot2Insts,
      FeatureDot5Insts,
      FeatureDot6Insts,
      FeatureNSAEncoding,
      FeatureWavefrontSize32,
      FeatureScalarStores,
      FeatureScalarAtomics,
      FeatureScalarFlatScratchInsts,
+     FeatureGetWaveIdInst,
+     FeatureSMemTimeInst,
+     FeatureMadMacF32Insts,
+     FeatureDsSrc2Insts,
      FeatureDoesNotSupportXNACK,
      FeatureCodeObjectV3])>;
 
 def FeatureISAVersion10_1_2 : FeatureSet<
   !listconcat(FeatureGroup.GFX10_1_Bugs,
     [FeatureGFX10,
      FeatureLDSBankCount32,
      FeatureDLInsts,
      FeatureDot1Insts,
      FeatureDot2Insts,
      FeatureDot5Insts,
      FeatureDot6Insts,
      FeatureNSAEncoding,
      FeatureWavefrontSize32,
      FeatureScalarStores,
      FeatureScalarAtomics,
      FeatureScalarFlatScratchInsts,
+     FeatureGetWaveIdInst,
+     FeatureSMemTimeInst,
+     FeatureMadMacF32Insts,
+     FeatureDsSrc2Insts,
      FeatureLdsMisalignedBug,
      FeatureDoesNotSupportXNACK,
      FeatureCodeObjectV3])>;
 
+def FeatureISAVersion10_3_0 : FeatureSet<
+  [FeatureGFX10,
+   FeatureGFX10_BEncoding,
+   FeatureGFX10_3Insts,
+   FeatureLDSBankCount32,
+   FeatureDLInsts,
+   FeatureDot1Insts,
+   FeatureDot2Insts,
+   FeatureDot5Insts,
+   FeatureDot6Insts,
+   FeatureNSAEncoding,
+   FeatureWavefrontSize32,
+   FeatureDoesNotSupportXNACK,
+   FeatureCodeObjectV3]>;
+
 //===----------------------------------------------------------------------===//
 
 def AMDGPUInstrInfo : InstrInfo {
   let guessInstructionProperties = 1;
   let noNamedPositionallyEncodedOperands = 1;
 }
 
 def AMDGPUAsmParser : AsmParser {
def HasFlatGlobalInsts : Predicate<"Subtarget->hasFlatGlobalInsts()">,
   AssemblerPredicate<(all_of FeatureFlatGlobalInsts)>;
 def HasFlatScratchInsts : Predicate<"Subtarget->hasFlatScratchInsts()">,
   AssemblerPredicate<(all_of FeatureFlatScratchInsts)>;
 def HasScalarFlatScratchInsts : Predicate<"Subtarget->hasScalarFlatScratchInsts()">,
   AssemblerPredicate<(all_of FeatureScalarFlatScratchInsts)>;
 def HasD16LoadStore : Predicate<"Subtarget->hasD16LoadStore()">,
   AssemblerPredicate<(all_of FeatureGFX9Insts)>;
 
+def HasGFX10_BEncoding : Predicate<"Subtarget->hasGFX10_BEncoding()">,
+  AssemblerPredicate<(all_of FeatureGFX10_BEncoding)>;
+
 def HasUnpackedD16VMem : Predicate<"Subtarget->hasUnpackedD16VMem()">,
   AssemblerPredicate<(all_of FeatureUnpackedD16VMem)>;
 def HasPackedD16VMem : Predicate<"!Subtarget->hasUnpackedD16VMem()">,
   AssemblerPredicate<(all_of (not FeatureUnpackedD16VMem))>;
 
 def D16PreservesUnusedBits :
   Predicate<"Subtarget->d16PreservesUnusedBits()">,
   AssemblerPredicate<(all_of FeatureGFX9Insts, (not FeatureSRAMECC))>;
def HasDot4Insts : Predicate<"Subtarget->hasDot4Insts()">,
   AssemblerPredicate<(all_of FeatureDot4Insts)>;
 
 def HasDot5Insts : Predicate<"Subtarget->hasDot5Insts()">,
   AssemblerPredicate<(all_of FeatureDot5Insts)>;
 
 def HasDot6Insts : Predicate<"Subtarget->hasDot6Insts()">,
   AssemblerPredicate<(all_of FeatureDot6Insts)>;
 
+def HasGetWaveIdInst : Predicate<"Subtarget->hasGetWaveIdInst()">,
+  AssemblerPredicate<(all_of FeatureGetWaveIdInst)>;
+
 def HasMAIInsts : Predicate<"Subtarget->hasMAIInsts()">,
   AssemblerPredicate<(all_of FeatureMAIInsts)>;
 
+def HasSMemTimeInst : Predicate<"Subtarget->hasSMemTimeInst()">,
+  AssemblerPredicate<(all_of FeatureSMemTimeInst)>;
+
+def HasNoSMemTimeInst : Predicate<"!Subtarget->hasSMemTimeInst()">;
+
 def HasPkFmacF16Inst : Predicate<"Subtarget->hasPkFmacF16Inst()">,
   AssemblerPredicate<(all_of FeaturePkFmacF16Inst)>;
 
+def HasMadMacF32Insts : Predicate<"Subtarget->hasMadMacF32Insts()">,
+  AssemblerPredicate<(all_of FeatureMadMacF32Insts)>;
+
 def HasAtomicFaddInsts : Predicate<"Subtarget->hasAtomicFaddInsts()">,
   AssemblerPredicate<(all_of FeatureAtomicFaddInsts)>;
 
+def HasNoMadMacF32Insts : Predicate<"!Subtarget->hasMadMacF32Insts()">,
+  AssemblerPredicate<(all_of (not FeatureMadMacF32Insts))>;
+
+def HasDsSrc2Insts : Predicate<"!Subtarget->hasDsSrc2Insts()">,
+  AssemblerPredicate<(all_of FeatureDsSrc2Insts)>;
+
 def HasOffset3fBug : Predicate<"!Subtarget->hasOffset3fBug()">,
   AssemblerPredicate<(all_of FeatureOffset3fBug)>;
 
 def EnableLateCFGStructurize : Predicate<
   "EnableLateStructurizeCFG">;
 
 // Include AMDGPU TD files
 include "SISchedule.td"
 include "GCNProcessors.td"
 include "AMDGPUInstrInfo.td"
 include "SIRegisterInfo.td"
 include "AMDGPURegisterBanks.td"
 include "AMDGPUInstructions.td"
 include "SIInstrInfo.td"
 include "AMDGPUCallingConv.td"
 include "AMDGPUSearchableTables.td"

llvm/lib/Target/AMDGPU/AMDGPUCodeGenPrepare.cpp

Value *AMDGPUCodeGenPrepare::expandDivRem24Impl(IRBuilder<> &Builder,
   // fq = trunc(fqm);
   CallInst *FQ = Builder.CreateUnaryIntrinsic(Intrinsic::trunc, FQM);
   FQ->copyFastMathFlags(Builder.getFastMathFlags());
 
   // float fqneg = -fq;
   Value *FQNeg = Builder.CreateFNeg(FQ);
 
   // float fr = mad(fqneg, fb, fa);
-  Value *FR = Builder.CreateIntrinsic(Intrinsic::amdgcn_fmad_ftz,
+  auto FMAD = !ST->hasMadMacF32Insts()
+                  ? Intrinsic::fma
+                  : (Intrinsic::ID)Intrinsic::amdgcn_fmad_ftz;
+  Value *FR = Builder.CreateIntrinsic(FMAD,
                                       {FQNeg->getType()}, {FQNeg, FB, FA}, FQ);
 
   // int iq = (int)fq;
   Value *IQ = IsSigned ? Builder.CreateFPToSI(FQ, I32Ty)
                        : Builder.CreateFPToUI(FQ, I32Ty);
 
   // fr = fabs(fr);
   FR = Builder.CreateUnaryIntrinsic(Intrinsic::fabs, FR, FQ);
 

llvm/lib/Target/AMDGPU/AMDGPUISelDAGToDAG.cpp

void AMDGPUDAGToDAGISel::Select(SDNode *N) {
   }
 
   // isa<MemSDNode> almost works but is slightly too permissive for some DS
   // intrinsics.
   if (Opc == ISD::LOAD || Opc == ISD::STORE || isa<AtomicSDNode>(N) ||
       (Opc == AMDGPUISD::ATOMIC_INC || Opc == AMDGPUISD::ATOMIC_DEC ||
        Opc == ISD::ATOMIC_LOAD_FADD ||
        Opc == AMDGPUISD::ATOMIC_LOAD_FMIN ||
-       Opc == AMDGPUISD::ATOMIC_LOAD_FMAX)) {
+       Opc == AMDGPUISD::ATOMIC_LOAD_FMAX ||
+       Opc == AMDGPUISD::ATOMIC_LOAD_CSUB)) {
     N = glueCopyToM0LDSInit(N);
     SelectCode(N);
     return;
   }
 
   switch (Opc) {
   default:
     break;

llvm/lib/Target/AMDGPU/AMDGPUISelLowering.h

enum NodeType : unsigned {
   TBUFFER_LOAD_FORMAT,
   TBUFFER_LOAD_FORMAT_D16,
   DS_ORDERED_COUNT,
   ATOMIC_CMP_SWAP,
   ATOMIC_INC,
   ATOMIC_DEC,
   ATOMIC_LOAD_FMIN,
   ATOMIC_LOAD_FMAX,
+  ATOMIC_LOAD_CSUB,
   BUFFER_LOAD,
   BUFFER_LOAD_UBYTE,
   BUFFER_LOAD_USHORT,
   BUFFER_LOAD_BYTE,
   BUFFER_LOAD_SHORT,
   BUFFER_LOAD_FORMAT,
   BUFFER_LOAD_FORMAT_D16,
   SBUFFER_LOAD,
enum NodeType : unsigned {
   BUFFER_ATOMIC_SMAX,
   BUFFER_ATOMIC_UMAX,
   BUFFER_ATOMIC_AND,
   BUFFER_ATOMIC_OR,
   BUFFER_ATOMIC_XOR,
   BUFFER_ATOMIC_INC,
   BUFFER_ATOMIC_DEC,
   BUFFER_ATOMIC_CMPSWAP,
+  BUFFER_ATOMIC_CSUB,
   BUFFER_ATOMIC_FADD,
   BUFFER_ATOMIC_PK_FADD,
   ATOMIC_PK_FADD,
 
   LAST_AMDGPU_ISD_NUMBER
 };
 
 
 } // End namespace AMDGPUISD
 
 } // End namespace llvm
 
 #endif

llvm/lib/Target/AMDGPU/AMDGPUISelLowering.cpp

SDValue AMDGPUTargetLowering::LowerDIVREM24(SDValue Op, SelectionDAG &DAG,
 
   // float fqneg = -fq;
   SDValue fqneg = DAG.getNode(ISD::FNEG, DL, FltVT, fq);
 
   MachineFunction &MF = DAG.getMachineFunction();
   const AMDGPUMachineFunction *MFI = MF.getInfo<AMDGPUMachineFunction>();
 
   // float fr = mad(fqneg, fb, fa);
-  unsigned OpCode = !MFI->getMode().allFP32Denormals() ?
+  unsigned OpCode = !Subtarget->hasMadMacF32Insts() ?
+                    (unsigned)ISD::FMA :
+                    !MFI->getMode().allFP32Denormals() ?
                     (unsigned)ISD::FMAD :
                     (unsigned)AMDGPUISD::FMAD_FTZ;
-
   SDValue fr = DAG.getNode(OpCode, DL, FltVT, fqneg, fb, fa);
 
   // int iq = (int)fq;
   SDValue iq = DAG.getNode(ToInt, DL, IntVT, fq);
 
   // fr = fabs(fr);
   fr = DAG.getNode(ISD::FABS, DL, FltVT, fr);
 
if (DAG.MaskedValueIsZero(RHS, APInt::getHighBitsSet(64, 32)) &&
     return;
   }
 
   if (isTypeLegal(MVT::i64)) {
     MachineFunction &MF = DAG.getMachineFunction();
     const SIMachineFunctionInfo *MFI = MF.getInfo<SIMachineFunctionInfo>();
 
     // Compute denominator reciprocal.
-    unsigned FMAD = !MFI->getMode().allFP32Denormals() ?
+    unsigned FMAD = !Subtarget->hasMadMacF32Insts() ?
+                    (unsigned)ISD::FMA :
+                    !MFI->getMode().allFP32Denormals() ?
                     (unsigned)ISD::FMAD :
                     (unsigned)AMDGPUISD::FMAD_FTZ;
 
-
     SDValue Cvt_Lo = DAG.getNode(ISD::UINT_TO_FP, DL, MVT::f32, RHS_Lo);
     SDValue Cvt_Hi = DAG.getNode(ISD::UINT_TO_FP, DL, MVT::f32, RHS_Hi);
     SDValue Mad1 = DAG.getNode(FMAD, DL, MVT::f32, Cvt_Hi,
       DAG.getConstantFP(APInt(32, 0x4f800000).bitsToFloat(), DL, MVT::f32),
       Cvt_Lo);
     SDValue Rcp = DAG.getNode(AMDGPUISD::RCP, DL, MVT::f32, Mad1);
     SDValue Mul1 = DAG.getNode(ISD::FMUL, DL, MVT::f32, Rcp,
       DAG.getConstantFP(APInt(32, 0x5f7ffffc).bitsToFloat(), DL, MVT::f32));
const char* AMDGPUTargetLowering::getTargetNodeName(unsigned Opcode) const {
   NODE_NAME_CASE(TBUFFER_LOAD_FORMAT)
   NODE_NAME_CASE(TBUFFER_LOAD_FORMAT_D16)
   NODE_NAME_CASE(DS_ORDERED_COUNT)
   NODE_NAME_CASE(ATOMIC_CMP_SWAP)
   NODE_NAME_CASE(ATOMIC_INC)
   NODE_NAME_CASE(ATOMIC_DEC)
   NODE_NAME_CASE(ATOMIC_LOAD_FMIN)
   NODE_NAME_CASE(ATOMIC_LOAD_FMAX)
+  NODE_NAME_CASE(ATOMIC_LOAD_CSUB)
   NODE_NAME_CASE(BUFFER_LOAD)
   NODE_NAME_CASE(BUFFER_LOAD_UBYTE)
   NODE_NAME_CASE(BUFFER_LOAD_USHORT)
   NODE_NAME_CASE(BUFFER_LOAD_BYTE)
   NODE_NAME_CASE(BUFFER_LOAD_SHORT)
   NODE_NAME_CASE(BUFFER_LOAD_FORMAT)
   NODE_NAME_CASE(BUFFER_LOAD_FORMAT_D16)
   NODE_NAME_CASE(SBUFFER_LOAD)
const char* AMDGPUTargetLowering::getTargetNodeName(unsigned Opcode) const {
   NODE_NAME_CASE(BUFFER_ATOMIC_SMAX)
   NODE_NAME_CASE(BUFFER_ATOMIC_UMAX)
   NODE_NAME_CASE(BUFFER_ATOMIC_AND)
   NODE_NAME_CASE(BUFFER_ATOMIC_OR)
   NODE_NAME_CASE(BUFFER_ATOMIC_XOR)
   NODE_NAME_CASE(BUFFER_ATOMIC_INC)
   NODE_NAME_CASE(BUFFER_ATOMIC_DEC)
   NODE_NAME_CASE(BUFFER_ATOMIC_CMPSWAP)
+  NODE_NAME_CASE(BUFFER_ATOMIC_CSUB)
   NODE_NAME_CASE(BUFFER_ATOMIC_FADD)
   NODE_NAME_CASE(BUFFER_ATOMIC_PK_FADD)
   NODE_NAME_CASE(ATOMIC_PK_FADD)
 
   case AMDGPUISD::LAST_AMDGPU_ISD_NUMBER: break;
   }
   return nullptr;
 }

llvm/lib/Target/AMDGPU/AMDGPUSearchableTables.td

 def : SourceOfDivergence<int_amdgcn_interp_p2_f16>;
 def : SourceOfDivergence<int_amdgcn_mbcnt_hi>;
 def : SourceOfDivergence<int_amdgcn_mbcnt_lo>;
 def : SourceOfDivergence<int_r600_read_tidig_x>;
 def : SourceOfDivergence<int_r600_read_tidig_y>;
 def : SourceOfDivergence<int_r600_read_tidig_z>;
 def : SourceOfDivergence<int_amdgcn_atomic_inc>;
 def : SourceOfDivergence<int_amdgcn_atomic_dec>;
+def : SourceOfDivergence<int_amdgcn_global_atomic_csub>;
 def : SourceOfDivergence<int_amdgcn_ds_fadd>;
 def : SourceOfDivergence<int_amdgcn_ds_fmin>;
 def : SourceOfDivergence<int_amdgcn_ds_fmax>;
 def : SourceOfDivergence<int_amdgcn_buffer_atomic_swap>;
 def : SourceOfDivergence<int_amdgcn_buffer_atomic_add>;
 def : SourceOfDivergence<int_amdgcn_buffer_atomic_sub>;
 def : SourceOfDivergence<int_amdgcn_buffer_atomic_smin>;
 def : SourceOfDivergence<int_amdgcn_buffer_atomic_umin>;

 def : SourceOfDivergence<int_amdgcn_struct_buffer_atomic_smax>;
 def : SourceOfDivergence<int_amdgcn_struct_buffer_atomic_umax>;
 def : SourceOfDivergence<int_amdgcn_struct_buffer_atomic_and>;
 def : SourceOfDivergence<int_amdgcn_struct_buffer_atomic_or>;
 def : SourceOfDivergence<int_amdgcn_struct_buffer_atomic_xor>;
 def : SourceOfDivergence<int_amdgcn_struct_buffer_atomic_inc>;
 def : SourceOfDivergence<int_amdgcn_struct_buffer_atomic_dec>;
 def : SourceOfDivergence<int_amdgcn_struct_buffer_atomic_cmpswap>;
+def : SourceOfDivergence<int_amdgcn_buffer_atomic_csub>;
 def : SourceOfDivergence<int_amdgcn_ps_live>;
 def : SourceOfDivergence<int_amdgcn_ds_swizzle>;
 def : SourceOfDivergence<int_amdgcn_ds_ordered_add>;
 def : SourceOfDivergence<int_amdgcn_ds_ordered_swap>;
 def : SourceOfDivergence<int_amdgcn_permlane16>;
 def : SourceOfDivergence<int_amdgcn_permlanex16>;
 def : SourceOfDivergence<int_amdgcn_mov_dpp>;
 def : SourceOfDivergence<int_amdgcn_mov_dpp8>;

llvm/lib/Target/AMDGPU/AMDGPUSubtarget.h

public:
   };
 
 private:
   Triple TargetTriple;
 
 protected:
   bool Has16BitInsts;
   bool HasMadMixInsts;
+  bool HasMadMacF32Insts;
+  bool HasDsSrc2Insts;
   bool HasSDWA;
   bool HasVOP3PInsts;
   bool HasMulI24;
   bool HasMulU24;
   bool HasInv2PiInlineImm;
   bool HasFminFmaxLegacy;
   bool EnablePromoteAlloca;
   bool HasTrigReducedRange;
public:
   bool isMesaKernel(const Function &F) const {
     return isMesa3DOS() && !AMDGPU::isShader(F.getCallingConv());
   }
 
   bool isAmdHsaOrMesa(const Function &F) const {
     return isAmdHsaOS() || isMesaKernel(F);
   }
 
+  bool isGCN() const {
+    return TargetTriple.getArch() == Triple::amdgcn;
+  }
+
   bool has16BitInsts() const {
     return Has16BitInsts;
   }
 
   bool hasMadMixInsts() const {
     return HasMadMixInsts;
   }
 
+  bool hasMadMacF32Insts() const {
+    return HasMadMacF32Insts || !isGCN();
+  }
+
+  bool hasDsSrc2Insts() const {
+    return HasDsSrc2Insts;
+  }
+
   bool hasSDWA() const {
     return HasSDWA;
   }
 
   bool hasVOP3PInsts() const {
     return HasVOP3PInsts;
   }
 
protected:
   bool FMA;
   bool MIMG_R128;
   bool IsGCN;
   bool GCN3Encoding;
   bool CIInsts;
   bool GFX8Insts;
   bool GFX9Insts;
   bool GFX10Insts;
+  bool GFX10_3Insts;
   bool GFX7GFX8GFX9Insts;
   bool SGPRInitBug;
   bool HasSMemRealTime;
   bool HasIntClamp;
   bool HasFmaMixInsts;
   bool HasMovrel;
   bool HasVGPRIndexMode;
   bool HasScalarStores;
   bool HasScalarAtomics;
   bool HasSDWAOmod;
   bool HasSDWAScalar;
   bool HasSDWASdst;
   bool HasSDWAMac;
   bool HasSDWAOutModsVOPC;
   bool HasDPP;
   bool HasDPP8;
   bool HasR128A16;
   bool HasGFX10A16;
   bool HasG16;
   bool HasNSAEncoding;
+  bool GFX10_BEncoding;
   bool HasDLInsts;
   bool HasDot1Insts;
   bool HasDot2Insts;
   bool HasDot3Insts;
   bool HasDot4Insts;
   bool HasDot5Insts;
   bool HasDot6Insts;
   bool HasMAIInsts;
   bool HasPkFmacF16Inst;
   bool HasAtomicFaddInsts;
   bool EnableSRAMECC;
   bool DoesNotSupportSRAMECC;
   bool HasNoSdstCMPX;
   bool HasVscnt;
+  bool HasGetWaveIdInst;
+  bool HasSMemTimeInst;
   bool HasRegisterBanking;
   bool HasVOP3Literal;
   bool HasNoDataDepHazard;
   bool FlatAddressSpace;
   bool FlatInstOffsets;
   bool FlatGlobalInsts;
   bool FlatScratchInsts;
   bool ScalarFlatScratchInsts;
public:
   bool hasFlatScratchInsts() const {
     return FlatScratchInsts;
   }
 
   bool hasScalarFlatScratchInsts() const {
     return ScalarFlatScratchInsts;
   }
 
+  bool hasGlobalAddTidInsts() const {
+    return GFX10_BEncoding;
+  }
+
+  bool hasAtomicCSub() const {
+    return GFX10_BEncoding;
+  }
+
   bool hasMultiDwordFlatScratchAddressing() const {
     return getGeneration() >= GFX9;
   }
 
   bool hasFlatSegmentOffsetBug() const {
     return HasFlatSegmentOffsetBug;
   }
 
public:
   bool hasNoSdstCMPX() const {
     return HasNoSdstCMPX;
   }
 
   bool hasVscnt() const {
     return HasVscnt;
   }
 
+  bool hasGetWaveIdInst() const {
+    return HasGetWaveIdInst;
+  }
+
+  bool hasSMemTimeInst() const {
+    return HasSMemTimeInst;
+  }
+
   bool hasRegisterBanking() const {
     return HasRegisterBanking;
   }
 
   bool hasVOP3Literal() const {
     return HasVOP3Literal;
   }
 
public:
   bool hasOffset3fBug() const {
     return HasOffset3fBug;
   }
 
   bool hasNSAEncoding() const {
     return HasNSAEncoding;
   }
 
+  bool hasGFX10_BEncoding() const {
+    return GFX10_BEncoding;
+  }
+
+  bool hasGFX10_3Insts() const {
+    return GFX10_3Insts;
+  }
+
   bool hasMadF16() const;
 
   bool enableSIScheduler() const {
     return EnableSIScheduler;
   }
 
   bool loadStoreOptEnabled() const {
     return EnableLoadStoreOpt;

llvm/lib/Target/AMDGPU/AMDGPUSubtarget.cpp

GCNSubtarget::initializeSubtargetDependencies(const Triple &TT,
 
   return *this;
 }
 
 AMDGPUSubtarget::AMDGPUSubtarget(const Triple &TT) :
   TargetTriple(TT),
   Has16BitInsts(false),
   HasMadMixInsts(false),
+  HasMadMacF32Insts(false),
+  HasDsSrc2Insts(false),
   HasSDWA(false),
   HasVOP3PInsts(false),
   HasMulI24(true),
   HasMulU24(true),
   HasInv2PiInlineImm(false),
   HasFminFmaxLegacy(true),
   EnablePromoteAlloca(false),
   HasTrigReducedRange(false),
GCNSubtarget::GCNSubtarget(const Triple &TT, StringRef GPU, StringRef FS,
     DumpCode(false),
 
     FP64(false),
     GCN3Encoding(false),
     CIInsts(false),
     GFX8Insts(false),
     GFX9Insts(false),
     GFX10Insts(false),
+    GFX10_3Insts(false),
     GFX7GFX8GFX9Insts(false),
     SGPRInitBug(false),
     HasSMemRealTime(false),
     HasIntClamp(false),
     HasFmaMixInsts(false),
     HasMovrel(false),
     HasVGPRIndexMode(false),
     HasScalarStores(false),
     HasScalarAtomics(false),
     HasSDWAOmod(false),
     HasSDWAScalar(false),
     HasSDWASdst(false),
     HasSDWAMac(false),
     HasSDWAOutModsVOPC(false),
     HasDPP(false),
     HasDPP8(false),
     HasR128A16(false),
     HasGFX10A16(false),
     HasG16(false),
     HasNSAEncoding(false),
+    GFX10_BEncoding(false),
     HasDLInsts(false),
     HasDot1Insts(false),
     HasDot2Insts(false),
     HasDot3Insts(false),
     HasDot4Insts(false),
     HasDot5Insts(false),
     HasDot6Insts(false),
     HasMAIInsts(false),
     HasPkFmacF16Inst(false),
     HasAtomicFaddInsts(false),
     EnableSRAMECC(false),
     DoesNotSupportSRAMECC(false),
     HasNoSdstCMPX(false),
     HasVscnt(false),
+    HasGetWaveIdInst(false),
+    HasSMemTimeInst(false),
     HasRegisterBanking(false),
     HasVOP3Literal(false),
     HasNoDataDepHazard(false),
     FlatAddressSpace(false),
     FlatInstOffsets(false),
     FlatGlobalInsts(false),
     FlatScratchInsts(false),
     ScalarFlatScratchInsts(false),

llvm/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp

public:
   bool isGFX9() const {
     return AMDGPU::isGFX9(getSTI());
   }
 
   bool isGFX10() const {
     return AMDGPU::isGFX10(getSTI());
   }
 
+  bool isGFX10_BEncoding() const {
+    return AMDGPU::isGFX10_BEncoding(getSTI());
+  }
+
   bool hasInv2PiInlineImm() const {
     return getFeatureBits()[AMDGPU::FeatureInv2PiInlineImm];
   }
 
   bool hasFlatOffsets() const {
     return getFeatureBits()[AMDGPU::FeatureFlatInstOffsets];
   }
 

llvm/lib/Target/AMDGPU/BUFInstructions.td

 >;
 defm BUFFER_ATOMIC_INC_X2 : MUBUF_Pseudo_Atomics <
   "buffer_atomic_inc_x2", VReg_64, i64, atomic_inc_global_64
 >;
 defm BUFFER_ATOMIC_DEC_X2 : MUBUF_Pseudo_Atomics <
   "buffer_atomic_dec_x2", VReg_64, i64, atomic_dec_global_64
 >;
 
+let SubtargetPredicate = HasGFX10_BEncoding in
+defm BUFFER_ATOMIC_CSUB : MUBUF_Pseudo_Atomics_RTN <
+  "buffer_atomic_csub", VGPR_32, i32, atomic_csub_global_32
+>;
+
 let SubtargetPredicate = isGFX8GFX9 in {
 def BUFFER_STORE_LDS_DWORD : MUBUF_Pseudo_Store_Lds <"buffer_store_lds_dword">;
 }
 
 let SubtargetPredicate = isGFX6 in { // isn't on CI & VI
 /*
 defm BUFFER_ATOMIC_RSUB        : MUBUF_Pseudo_Atomics <"buffer_atomic_rsub">;
 defm BUFFER_ATOMIC_RSUB_X2     : MUBUF_Pseudo_Atomics <"buffer_atomic_rsub_x2">;

 defm : BufferAtomicPatterns<SIbuffer_atomic_umin, i32, "BUFFER_ATOMIC_UMIN">;
 defm : BufferAtomicPatterns<SIbuffer_atomic_smax, i32, "BUFFER_ATOMIC_SMAX">;
 defm : BufferAtomicPatterns<SIbuffer_atomic_umax, i32, "BUFFER_ATOMIC_UMAX">;
 defm : BufferAtomicPatterns<SIbuffer_atomic_and, i32, "BUFFER_ATOMIC_AND">;
 defm : BufferAtomicPatterns<SIbuffer_atomic_or, i32, "BUFFER_ATOMIC_OR">;
 defm : BufferAtomicPatterns<SIbuffer_atomic_xor, i32, "BUFFER_ATOMIC_XOR">;
 defm : BufferAtomicPatterns<SIbuffer_atomic_inc, i32, "BUFFER_ATOMIC_INC">;
 defm : BufferAtomicPatterns<SIbuffer_atomic_dec, i32, "BUFFER_ATOMIC_DEC">;
+defm : BufferAtomicPatterns<SIbuffer_atomic_csub, i32, "BUFFER_ATOMIC_CSUB">;
 defm : BufferAtomicPatterns<SIbuffer_atomic_swap, i64, "BUFFER_ATOMIC_SWAP_X2">;
 defm : BufferAtomicPatterns<SIbuffer_atomic_add, i64,  "BUFFER_ATOMIC_ADD_X2">;
 defm : BufferAtomicPatterns<SIbuffer_atomic_sub, i64, "BUFFER_ATOMIC_SUB_X2">;
 defm : BufferAtomicPatterns<SIbuffer_atomic_smin, i64, "BUFFER_ATOMIC_SMIN_X2">;
 defm : BufferAtomicPatterns<SIbuffer_atomic_umin, i64, "BUFFER_ATOMIC_UMIN_X2">;
 defm : BufferAtomicPatterns<SIbuffer_atomic_smax, i64, "BUFFER_ATOMIC_SMAX_X2">;
 defm : BufferAtomicPatterns<SIbuffer_atomic_umax, i64, "BUFFER_ATOMIC_UMAX_X2">;
 defm : BufferAtomicPatterns<SIbuffer_atomic_and, i64, "BUFFER_ATOMIC_AND_X2">;
def _LDS_OFFSET_gfx10 : MUBUF_Real_gfx10<op, !cast<MUBUF_Pseudo>(NAME#"_LDS_OFFSET")>,
                             MUBUFLdsTable<1, NAME # "_OFFSET_gfx10">;
     def _LDS_OFFEN_gfx10  : MUBUF_Real_gfx10<op, !cast<MUBUF_Pseudo>(NAME#"_LDS_OFFEN")>,
                             MUBUFLdsTable<1, NAME # "_OFFEN_gfx10">;
     def _LDS_IDXEN_gfx10  : MUBUF_Real_gfx10<op, !cast<MUBUF_Pseudo>(NAME#"_LDS_IDXEN")>,
                             MUBUFLdsTable<1, NAME # "_IDXEN_gfx10">;
     def _LDS_BOTHEN_gfx10 : MUBUF_Real_gfx10<op, !cast<MUBUF_Pseudo>(NAME#"_LDS_BOTHEN")>,
                             MUBUFLdsTable<1, NAME # "_BOTHEN_gfx10">;
   }
-  multiclass MUBUF_Real_Atomics_gfx10<bits<8> op> :
-      MUBUF_Real_AllAddr_gfx10<op> {
+  multiclass MUBUF_Real_Atomics_RTN_gfx10<bits<8> op> {
     def _BOTHEN_RTN_gfx10 :
       MUBUF_Real_gfx10<op, !cast<MUBUF_Pseudo>(NAME#"_BOTHEN_RTN")>;
     def _IDXEN_RTN_gfx10 :
       MUBUF_Real_gfx10<op, !cast<MUBUF_Pseudo>(NAME#"_IDXEN_RTN")>;
     def _OFFEN_RTN_gfx10 :
       MUBUF_Real_gfx10<op, !cast<MUBUF_Pseudo>(NAME#"_OFFEN_RTN")>;
     def _OFFSET_RTN_gfx10 :
       MUBUF_Real_gfx10<op, !cast<MUBUF_Pseudo>(NAME#"_OFFSET_RTN")>;
   }
+  multiclass MUBUF_Real_Atomics_gfx10<bits<8> op> :
+      MUBUF_Real_AllAddr_gfx10<op>, MUBUF_Real_Atomics_RTN_gfx10<op>;
 } // End AssemblerPredicate = isGFX10Plus, DecoderNamespace = "GFX10"
 
 defm BUFFER_STORE_BYTE_D16_HI     : MUBUF_Real_AllAddr_gfx10<0x019>;
 defm BUFFER_STORE_SHORT_D16_HI    : MUBUF_Real_AllAddr_gfx10<0x01b>;
 defm BUFFER_LOAD_UBYTE_D16        : MUBUF_Real_AllAddr_gfx10<0x020>;
 defm BUFFER_LOAD_UBYTE_D16_HI     : MUBUF_Real_AllAddr_gfx10<0x021>;
 defm BUFFER_LOAD_SBYTE_D16        : MUBUF_Real_AllAddr_gfx10<0x022>;
 defm BUFFER_LOAD_SBYTE_D16_HI     : MUBUF_Real_AllAddr_gfx10<0x023>;

 defm BUFFER_ATOMIC_XOR_X2      : MUBUF_Real_Atomics_gfx6_gfx7_gfx10<0x05b>;
 defm BUFFER_ATOMIC_INC_X2      : MUBUF_Real_Atomics_gfx6_gfx7_gfx10<0x05c>;
 defm BUFFER_ATOMIC_DEC_X2      : MUBUF_Real_Atomics_gfx6_gfx7_gfx10<0x05d>;
 // FIXME-GFX7: Need to handle hazard for BUFFER_ATOMIC_FCMPSWAP_X2 on GFX7.
 defm BUFFER_ATOMIC_FCMPSWAP_X2 : MUBUF_Real_Atomics_gfx6_gfx7_gfx10<0x05e>;
 defm BUFFER_ATOMIC_FMIN_X2     : MUBUF_Real_Atomics_gfx6_gfx7_gfx10<0x05f>;
 defm BUFFER_ATOMIC_FMAX_X2     : MUBUF_Real_Atomics_gfx6_gfx7_gfx10<0x060>;
 
+defm BUFFER_ATOMIC_CSUB       : MUBUF_Real_Atomics_RTN_gfx10<0x034>;
+
 defm BUFFER_WBINVL1_SC        : MUBUF_Real_gfx6<0x070>;
 defm BUFFER_WBINVL1_VOL       : MUBUF_Real_gfx7<0x070>;
 def  BUFFER_WBINVL1_gfx6_gfx7 : MUBUF_Real_gfx6_gfx7<0x071, BUFFER_WBINVL1>;
 
 //===----------------------------------------------------------------------===//
 // Base ENC_MTBUF for GFX6, GFX7, GFX10.
 //===----------------------------------------------------------------------===//
 

llvm/lib/Target/AMDGPU/DSInstructions.td

def DS_GWS_INIT       : DS_GWS_1D<"ds_gws_init"> {
   let mayLoad = 0;
 }
 def DS_GWS_SEMA_V     : DS_GWS_0D<"ds_gws_sema_v">;
 def DS_GWS_SEMA_BR    : DS_GWS_1D<"ds_gws_sema_br">;
 def DS_GWS_SEMA_P     : DS_GWS_0D<"ds_gws_sema_p">;
 def DS_GWS_BARRIER    : DS_GWS_1D<"ds_gws_barrier">;
 }
 
+let SubtargetPredicate = HasDsSrc2Insts in {
 def DS_ADD_SRC2_U32   : DS_1A<"ds_add_src2_u32">;
 def DS_SUB_SRC2_U32   : DS_1A<"ds_sub_src2_u32">;
 def DS_RSUB_SRC2_U32  : DS_1A<"ds_rsub_src2_u32">;
 def DS_INC_SRC2_U32   : DS_1A<"ds_inc_src2_u32">;
 def DS_DEC_SRC2_U32   : DS_1A<"ds_dec_src2_u32">;
 def DS_MIN_SRC2_I32   : DS_1A<"ds_min_src2_i32">;
 def DS_MAX_SRC2_I32   : DS_1A<"ds_max_src2_i32">;
 def DS_MIN_SRC2_U32   : DS_1A<"ds_min_src2_u32">;

 def DS_AND_SRC2_B64   : DS_1A<"ds_and_src2_b64">;
 def DS_OR_SRC2_B64    : DS_1A<"ds_or_src2_b64">;
 def DS_XOR_SRC2_B64   : DS_1A<"ds_xor_src2_b64">;
 def DS_MIN_SRC2_F64   : DS_1A<"ds_min_src2_f64">;
 def DS_MAX_SRC2_F64   : DS_1A<"ds_max_src2_f64">;
 
 def DS_WRITE_SRC2_B32 : DS_1A<"ds_write_src2_b32">;
 def DS_WRITE_SRC2_B64 : DS_1A<"ds_write_src2_b64">;
+} // End SubtargetPredicate = HasDsSrc2Insts
 
 let Uses = [EXEC], mayLoad = 0, mayStore = 0, isConvergent = 1 in {
 def DS_SWIZZLE_B32 : DS_1A_RET <"ds_swizzle_b32", VGPR_32, 0, SwizzleImm>;
 }
 
 let mayStore = 0 in {
 defm DS_READ_I8      : DS_1A_RET_mc<"ds_read_i8">;
 defm DS_READ_U8      : DS_1A_RET_mc<"ds_read_u8">;

 def DS_PERMUTE_B32  : DS_1A1D_PERMUTE <"ds_permute_b32",
                                        int_amdgcn_ds_permute>;
 def DS_BPERMUTE_B32 : DS_1A1D_PERMUTE <"ds_bpermute_b32",
                                        int_amdgcn_ds_bpermute>;
 }
 
 } // let SubtargetPredicate = isGFX8Plus
 
-let SubtargetPredicate = HasLDSFPAtomics in {
+let SubtargetPredicate = HasLDSFPAtomics, OtherPredicates = [HasDsSrc2Insts] in {
 def DS_ADD_SRC2_F32 : DS_1A<"ds_add_src2_f32">;
 }
 
 //===----------------------------------------------------------------------===//
 // DS Patterns
 //===----------------------------------------------------------------------===//
 
 def : GCNPat <

llvm/lib/Target/AMDGPU/Disassembler/AMDGPUDisassembler.cpp

do {
     // ToDo: better to switch encoding length using some bit predicate
     // but it is unknown yet, so try all we can
 
     // Try to decode DPP and SDWA first to solve conflict with VOP1 and VOP2
     // encodings
     if (Bytes.size() >= 8) {
       const uint64_t QW = eatBytes<uint64_t>(Bytes);
 
+      if (STI.getFeatureBits()[AMDGPU::FeatureGFX10_BEncoding]) {
+        Res = tryDecodeInst(DecoderTableGFX10_B64, MI, QW, Address);
+        if (Res) {
+          if (AMDGPU::getNamedOperandIdx(MI.getOpcode(), AMDGPU::OpName::dpp8)
+              == -1)
+            break;
+          if (convertDPP8Inst(MI) == MCDisassembler::Success)
+            break;
+          MI = MCInst(); // clear
+        }
+      }
+
       Res = tryDecodeInst(DecoderTableDPP864, MI, QW, Address);
       if (Res && convertDPP8Inst(MI) == MCDisassembler::Success)
         break;
 
       MI = MCInst(); // clear
 
       Res = tryDecodeInst(DecoderTableDPP64, MI, QW, Address);
       if (Res) break;
do {
     if (Res) break;
 
     Res = tryDecodeInst(DecoderTableAMDGPU32, MI, DW, Address);
     if (Res) break;
 
     Res = tryDecodeInst(DecoderTableGFX932, MI, DW, Address);
     if (Res) break;
 
+    if (STI.getFeatureBits()[AMDGPU::FeatureGFX10_BEncoding]) {
+      Res = tryDecodeInst(DecoderTableGFX10_B32, MI, DW, Address);
+      if (Res) break;
+    }
+
     Res = tryDecodeInst(DecoderTableGFX1032, MI, DW, Address);
     if (Res) break;
 
     if (Bytes.size() < 4) break;
     const uint64_t QW = ((uint64_t)eatBytes<uint32_t>(Bytes) << 32) | DW;
     Res = tryDecodeInst(DecoderTableGFX864, MI, QW, Address);
     if (Res) break;
 

llvm/lib/Target/AMDGPU/FLATInstructions.td

multiclass FLAT_Global_Load_Pseudo<string opName, RegisterClass regClass, bit HasTiedInput = 0> {
   let is_flat_global = 1, SubtargetPredicate = HasFlatGlobalInsts in {
     def "" : FLAT_Load_Pseudo<opName, regClass, HasTiedInput, 1>,
       GlobalSaddrTable<0, opName>;
     def _SADDR : FLAT_Load_Pseudo<opName, regClass, HasTiedInput, 1, 1>,
       GlobalSaddrTable<1, opName>;
   }
 }
 
+class FLAT_Global_Load_AddTid_Pseudo <string opName, RegisterClass regClass,
+  bit HasTiedOutput = 0, bit HasSignedOffset = 0> : FLAT_Pseudo<
+  opName,
+  (outs regClass:$vdst),
+  !con((ins SReg_64:$saddr, flat_offset:$offset, GLC:$glc, SLC:$slc, DLC:$dlc),
+    !if(HasTiedOutput, (ins regClass:$vdst_in), (ins))),
+  " $vdst, $saddr$offset$glc$slc$dlc"> {
+  let is_flat_global = 1;
+  let has_data = 0;
+  let mayLoad = 1;
+  let has_vaddr = 0;
+  let has_saddr = 1;
+  let enabled_saddr = 1;
+  let maybeAtomic = 1;
+
+  let Constraints = !if(HasTiedOutput, "$vdst = $vdst_in", "");
+  let DisableEncoding = !if(HasTiedOutput, "$vdst_in", "");
+}
+
 multiclass FLAT_Global_Store_Pseudo<string opName, RegisterClass regClass> {
   let is_flat_global = 1, SubtargetPredicate = HasFlatGlobalInsts in {
     def "" : FLAT_Store_Pseudo<opName, regClass, 1>,
       GlobalSaddrTable<0, opName>;
     def _SADDR : FLAT_Store_Pseudo<opName, regClass, 1, 1>,
       GlobalSaddrTable<1, opName>;
   }
 }
 
+class FLAT_Global_Store_AddTid_Pseudo <string opName, RegisterClass vdataClass,
+  bit HasSignedOffset = 0> : FLAT_Pseudo<
+  opName,
+  (outs),
+  !con(
+    (ins vdataClass:$vdata, SReg_64:$saddr),
+      (ins flat_offset:$offset, GLC:$glc, SLC:$slc, DLC:$dlc)),
+  " $vdata, $saddr$offset$glc$slc$dlc"> {
+  let is_flat_global = 1;
+  let mayLoad  = 0;
+  let mayStore = 1;
+  let has_vdst = 0;
+  let has_vaddr = 0;
+  let has_saddr = 1;
+  let enabled_saddr = 1;
+  let maybeAtomic = 1;
+}
+
 class FLAT_Scratch_Load_Pseudo <string opName, RegisterClass regClass,
   bit EnableSaddr = 0>: FLAT_Pseudo<
   opName,
   (outs regClass:$vdst),
   !if(EnableSaddr,
       (ins SReg_32_XEXEC_HI:$saddr, flat_offset:$offset, GLC:$glc, SLC:$slc, DLC:$dlc),
       (ins VGPR_32:$vaddr, flat_offset:$offset, GLC:$glc, SLC:$slc, DLC:$dlc)),
   " $vdst, "#!if(EnableSaddr, "off", "$vaddr")#!if(EnableSaddr, ", $saddr", ", off")#"$offset$glc$slc$dlc"> {

 defm GLOBAL_LOAD_DWORDX4  : FLAT_Global_Load_Pseudo <"global_load_dwordx4", VReg_128>;
 
 defm GLOBAL_LOAD_UBYTE_D16    : FLAT_Global_Load_Pseudo <"global_load_ubyte_d16", VGPR_32, 1>;
 defm GLOBAL_LOAD_UBYTE_D16_HI : FLAT_Global_Load_Pseudo <"global_load_ubyte_d16_hi", VGPR_32, 1>;
 defm GLOBAL_LOAD_SBYTE_D16    : FLAT_Global_Load_Pseudo <"global_load_sbyte_d16", VGPR_32, 1>;
 defm GLOBAL_LOAD_SBYTE_D16_HI : FLAT_Global_Load_Pseudo <"global_load_sbyte_d16_hi", VGPR_32, 1>;
 defm GLOBAL_LOAD_SHORT_D16    : FLAT_Global_Load_Pseudo <"global_load_short_d16", VGPR_32, 1>;
 defm GLOBAL_LOAD_SHORT_D16_HI : FLAT_Global_Load_Pseudo <"global_load_short_d16_hi", VGPR_32, 1>;
+let OtherPredicates = [HasGFX10_BEncoding] in
+def  GLOBAL_LOAD_DWORD_ADDTID : FLAT_Global_Load_AddTid_Pseudo <"global_load_dword_addtid", VGPR_32>;
 
 defm GLOBAL_STORE_BYTE    : FLAT_Global_Store_Pseudo <"global_store_byte", VGPR_32>;
 defm GLOBAL_STORE_SHORT   : FLAT_Global_Store_Pseudo <"global_store_short", VGPR_32>;
 defm GLOBAL_STORE_DWORD   : FLAT_Global_Store_Pseudo <"global_store_dword", VGPR_32>;
 defm GLOBAL_STORE_DWORDX2 : FLAT_Global_Store_Pseudo <"global_store_dwordx2", VReg_64>;
 defm GLOBAL_STORE_DWORDX3 : FLAT_Global_Store_Pseudo <"global_store_dwordx3", VReg_96>;
 defm GLOBAL_STORE_DWORDX4 : FLAT_Global_Store_Pseudo <"global_store_dwordx4", VReg_128>;
+let OtherPredicates = [HasGFX10_BEncoding] in
+def  GLOBAL_STORE_DWORD_ADDTID : FLAT_Global_Store_AddTid_Pseudo <"global_store_dword_addtid", VGPR_32>;
 
 defm GLOBAL_STORE_BYTE_D16_HI  : FLAT_Global_Store_Pseudo <"global_store_byte_d16_hi", VGPR_32>;
 defm GLOBAL_STORE_SHORT_D16_HI : FLAT_Global_Store_Pseudo <"global_store_short_d16_hi", VGPR_32>;
 
 let is_flat_global = 1 in {
 defm GLOBAL_ATOMIC_CMPSWAP : FLAT_Global_Atomic_Pseudo <"global_atomic_cmpswap",
                                VGPR_32, i32, AMDGPUatomic_cmp_swap_global_32, null_frag,
                                v2i32, VReg_64>;

 defm GLOBAL_ATOMIC_XOR_X2 : FLAT_Global_Atomic_Pseudo <"global_atomic_xor_x2",
                               VReg_64, i64, atomic_load_xor_global_64>;
 
 defm GLOBAL_ATOMIC_INC_X2 : FLAT_Global_Atomic_Pseudo <"global_atomic_inc_x2",
                               VReg_64, i64, atomic_inc_global_64>;
 
 defm GLOBAL_ATOMIC_DEC_X2 : FLAT_Global_Atomic_Pseudo <"global_atomic_dec_x2",
                               VReg_64, i64, atomic_dec_global_64>;
+
+let SubtargetPredicate = HasGFX10_BEncoding in
+defm GLOBAL_ATOMIC_CSUB : FLAT_Global_Atomic_Pseudo_RTN <"global_atomic_csub",
+                              VGPR_32, i32, atomic_csub_global_32>;
 } // End is_flat_global = 1
 
 
 
 let SubtargetPredicate = HasFlatScratchInsts in {
 defm SCRATCH_LOAD_UBYTE    : FLAT_Scratch_Load_Pseudo <"scratch_load_ubyte", VGPR_32>;
 defm SCRATCH_LOAD_SBYTE    : FLAT_Scratch_Load_Pseudo <"scratch_load_sbyte", VGPR_32>;
 defm SCRATCH_LOAD_USHORT   : FLAT_Scratch_Load_Pseudo <"scratch_load_ushort", VGPR_32>;

 def : FlatSignedAtomicPat <GLOBAL_ATOMIC_SMAX_RTN, atomic_load_max_global_32, i32>;
 def : FlatSignedAtomicPat <GLOBAL_ATOMIC_UMAX_RTN, atomic_load_umax_global_32, i32>;
 def : FlatSignedAtomicPat <GLOBAL_ATOMIC_SMIN_RTN, atomic_load_min_global_32, i32>;
 def : FlatSignedAtomicPat <GLOBAL_ATOMIC_UMIN_RTN, atomic_load_umin_global_32, i32>;
 def : FlatSignedAtomicPat <GLOBAL_ATOMIC_OR_RTN, atomic_load_or_global_32, i32>;
 def : FlatSignedAtomicPat <GLOBAL_ATOMIC_SWAP_RTN, atomic_swap_global_32, i32>;
 def : FlatSignedAtomicPat <GLOBAL_ATOMIC_CMPSWAP_RTN, AMDGPUatomic_cmp_swap_global_32, i32, v2i32>;
 def : FlatSignedAtomicPat <GLOBAL_ATOMIC_XOR_RTN, atomic_load_xor_global_32, i32>;
+def : FlatSignedAtomicPat <GLOBAL_ATOMIC_CSUB_RTN, atomic_csub_global_32, i32>;
 
 def : FlatSignedAtomicPat <GLOBAL_ATOMIC_ADD_X2_RTN, atomic_load_add_global_64, i64>;
 def : FlatSignedAtomicPat <GLOBAL_ATOMIC_SUB_X2_RTN, atomic_load_sub_global_64, i64>;
 def : FlatSignedAtomicPat <GLOBAL_ATOMIC_INC_X2_RTN, atomic_inc_global_64, i64>;
 def : FlatSignedAtomicPat <GLOBAL_ATOMIC_DEC_X2_RTN, atomic_dec_global_64, i64>;
 def : FlatSignedAtomicPat <GLOBAL_ATOMIC_AND_X2_RTN, atomic_load_and_global_64, i64>;
 def : FlatSignedAtomicPat <GLOBAL_ATOMIC_SMAX_X2_RTN, atomic_load_max_global_64, i64>;
 def : FlatSignedAtomicPat <GLOBAL_ATOMIC_UMAX_X2_RTN, atomic_load_umax_global_64, i64>;
multiclass FLAT_Real_Atomics_gfx10<bits<7> op> :
   FLAT_Real_Base_gfx10<op>,
   FLAT_Real_RTN_gfx10<op>;
 
 multiclass FLAT_Real_GlblAtomics_gfx10<bits<7> op> :
   FLAT_Real_AllAddr_gfx10<op>,
   FLAT_Real_RTN_gfx10<op>,
   FLAT_Real_SADDR_RTN_gfx10<op>;
 
+multiclass FLAT_Real_GlblAtomics_RTN_gfx10<bits<7> op> :
+  FLAT_Real_RTN_gfx10<op>,
+  FLAT_Real_SADDR_RTN_gfx10<op>;
 
 // ENC_FLAT.
 defm FLAT_LOAD_UBYTE            : FLAT_Real_Base_gfx10<0x008>;
 defm FLAT_LOAD_SBYTE            : FLAT_Real_Base_gfx10<0x009>;
 defm FLAT_LOAD_USHORT           : FLAT_Real_Base_gfx10<0x00a>;
 defm FLAT_LOAD_SSHORT           : FLAT_Real_Base_gfx10<0x00b>;
 defm FLAT_LOAD_DWORD            : FLAT_Real_Base_gfx10<0x00c>;
 defm FLAT_LOAD_DWORDX2          : FLAT_Real_Base_gfx10<0x00d>;

 defm GLOBAL_LOAD_SBYTE_D16      : FLAT_Real_AllAddr_gfx10<0x022>;
 defm GLOBAL_LOAD_SBYTE_D16_HI   : FLAT_Real_AllAddr_gfx10<0x023>;
 defm GLOBAL_LOAD_SHORT_D16      : FLAT_Real_AllAddr_gfx10<0x024>;
 defm GLOBAL_LOAD_SHORT_D16_HI   : FLAT_Real_AllAddr_gfx10<0x025>;
 defm GLOBAL_ATOMIC_SWAP         : FLAT_Real_GlblAtomics_gfx10<0x030>;
 defm GLOBAL_ATOMIC_CMPSWAP      : FLAT_Real_GlblAtomics_gfx10<0x031>;
 defm GLOBAL_ATOMIC_ADD          : FLAT_Real_GlblAtomics_gfx10<0x032>;
 defm GLOBAL_ATOMIC_SUB          : FLAT_Real_GlblAtomics_gfx10<0x033>;
+defm GLOBAL_ATOMIC_CSUB         : FLAT_Real_GlblAtomics_RTN_gfx10<0x034>;
 defm GLOBAL_ATOMIC_SMIN         : FLAT_Real_GlblAtomics_gfx10<0x035>;
 defm GLOBAL_ATOMIC_UMIN         : FLAT_Real_GlblAtomics_gfx10<0x036>;
 defm GLOBAL_ATOMIC_SMAX         : FLAT_Real_GlblAtomics_gfx10<0x037>;
 defm GLOBAL_ATOMIC_UMAX         : FLAT_Real_GlblAtomics_gfx10<0x038>;
 defm GLOBAL_ATOMIC_AND          : FLAT_Real_GlblAtomics_gfx10<0x039>;
 defm GLOBAL_ATOMIC_OR           : FLAT_Real_GlblAtomics_gfx10<0x03a>;
 defm GLOBAL_ATOMIC_XOR          : FLAT_Real_GlblAtomics_gfx10<0x03b>;
 defm GLOBAL_ATOMIC_INC          : FLAT_Real_GlblAtomics_gfx10<0x03c>;

 defm GLOBAL_ATOMIC_AND_X2       : FLAT_Real_GlblAtomics_gfx10<0x059>;
 defm GLOBAL_ATOMIC_OR_X2        : FLAT_Real_GlblAtomics_gfx10<0x05a>;
 defm GLOBAL_ATOMIC_XOR_X2       : FLAT_Real_GlblAtomics_gfx10<0x05b>;
 defm GLOBAL_ATOMIC_INC_X2       : FLAT_Real_GlblAtomics_gfx10<0x05c>;
 defm GLOBAL_ATOMIC_DEC_X2       : FLAT_Real_GlblAtomics_gfx10<0x05d>;
 defm GLOBAL_ATOMIC_FCMPSWAP_X2  : FLAT_Real_GlblAtomics_gfx10<0x05e>;
 defm GLOBAL_ATOMIC_FMIN_X2      : FLAT_Real_GlblAtomics_gfx10<0x05f>;
 defm GLOBAL_ATOMIC_FMAX_X2      : FLAT_Real_GlblAtomics_gfx10<0x060>;
-
+defm GLOBAL_LOAD_DWORD_ADDTID   : FLAT_Real_Base_gfx10<0x016>;
+defm GLOBAL_STORE_DWORD_ADDTID  : FLAT_Real_Base_gfx10<0x017>;
 
 // ENC_FLAT_SCRATCH.
 defm SCRATCH_LOAD_UBYTE         : FLAT_Real_AllAddr_gfx10<0x008>;
 defm SCRATCH_LOAD_SBYTE         : FLAT_Real_AllAddr_gfx10<0x009>;
 defm SCRATCH_LOAD_USHORT        : FLAT_Real_AllAddr_gfx10<0x00a>;
 defm SCRATCH_LOAD_SSHORT        : FLAT_Real_AllAddr_gfx10<0x00b>;
 defm SCRATCH_LOAD_DWORD         : FLAT_Real_AllAddr_gfx10<0x00c>;
 defm SCRATCH_LOAD_DWORDX2       : FLAT_Real_AllAddr_gfx10<0x00d>;

llvm/lib/Target/AMDGPU/GCNProcessors.td

 
 def : ProcessorModel<"gfx1011", GFX10SpeedModel,
   FeatureISAVersion10_1_1.Features
 >;
 
 def : ProcessorModel<"gfx1012", GFX10SpeedModel,
   FeatureISAVersion10_1_2.Features
 >;
+
+def : ProcessorModel<"gfx1030", GFX10SpeedModel,
+  FeatureISAVersion10_3_0.Features
+>;

llvm/lib/Target/AMDGPU/MCTargetDesc/AMDGPUTargetStreamer.cpp

StringRef AMDGPUTargetStreamer::getArchNameFromElfMach(unsigned ElfMach) {
   case ELF::EF_AMDGPU_MACH_AMDGCN_GFX902:  AK = GK_GFX902;  break;
   case ELF::EF_AMDGPU_MACH_AMDGCN_GFX904:  AK = GK_GFX904;  break;
   case ELF::EF_AMDGPU_MACH_AMDGCN_GFX906:  AK = GK_GFX906;  break;
   case ELF::EF_AMDGPU_MACH_AMDGCN_GFX908:  AK = GK_GFX908;  break;
   case ELF::EF_AMDGPU_MACH_AMDGCN_GFX909:  AK = GK_GFX909;  break;
   case ELF::EF_AMDGPU_MACH_AMDGCN_GFX1010: AK = GK_GFX1010; break;
   case ELF::EF_AMDGPU_MACH_AMDGCN_GFX1011: AK = GK_GFX1011; break;
   case ELF::EF_AMDGPU_MACH_AMDGCN_GFX1012: AK = GK_GFX1012; break;
+  case ELF::EF_AMDGPU_MACH_AMDGCN_GFX1030: AK = GK_GFX1030; break;
   case ELF::EF_AMDGPU_MACH_NONE:           AK = GK_NONE;    break;
   }
 
   StringRef GPUName = getArchNameAMDGCN(AK);
   if (GPUName != "")
     return GPUName;
   return getArchNameR600(AK);
 }
unsigned AMDGPUTargetStreamer::getElfMach(StringRef GPU) {
   case GK_GFX902:  return ELF::EF_AMDGPU_MACH_AMDGCN_GFX902;
   case GK_GFX904:  return ELF::EF_AMDGPU_MACH_AMDGCN_GFX904;
   case GK_GFX906:  return ELF::EF_AMDGPU_MACH_AMDGCN_GFX906;
   case GK_GFX908:  return ELF::EF_AMDGPU_MACH_AMDGCN_GFX908;
   case GK_GFX909:  return ELF::EF_AMDGPU_MACH_AMDGCN_GFX909;
   case GK_GFX1010: return ELF::EF_AMDGPU_MACH_AMDGCN_GFX1010;
   case GK_GFX1011: return ELF::EF_AMDGPU_MACH_AMDGCN_GFX1011;
   case GK_GFX1012: return ELF::EF_AMDGPU_MACH_AMDGCN_GFX1012;
+  case GK_GFX1030: return ELF::EF_AMDGPU_MACH_AMDGCN_GFX1030;
   case GK_NONE:    return ELF::EF_AMDGPU_MACH_NONE;
   }
 
   llvm_unreachable("unknown GPU");
 }
 
 //===----------------------------------------------------------------------===//
 // AMDGPUTargetAsmStreamer

llvm/lib/Target/AMDGPU/MIMGInstructions.td

 defm IMAGE_SAMPLE_C_CD_CL_G16   : MIMG_Sampler <0x000000eb, AMDGPUSample_c_cd_cl, 0, 1>;
 defm IMAGE_SAMPLE_CD_O_G16      : MIMG_Sampler <0x000000ec, AMDGPUSample_cd_o, 0, 1>;
 defm IMAGE_SAMPLE_CD_CL_O_G16   : MIMG_Sampler <0x000000ed, AMDGPUSample_cd_cl_o, 0, 1>;
 defm IMAGE_SAMPLE_C_CD_O_G16    : MIMG_Sampler <0x000000ee, AMDGPUSample_c_cd_o, 0, 1>;
 defm IMAGE_SAMPLE_C_CD_CL_O_G16 : MIMG_Sampler <0x000000ef, AMDGPUSample_c_cd_cl_o, 0, 1>;
 //def IMAGE_RSRC256 : MIMG_NoPattern_RSRC256 <"image_rsrc256", 0x0000007e>;
 //def IMAGE_SAMPLER : MIMG_NoPattern_ <"image_sampler", 0x0000007f>;
 
+let SubtargetPredicate = HasGFX10_BEncoding in
+defm IMAGE_MSAA_LOAD : MIMG_NoSampler <0x00000080, "image_msaa_load", 1>;
+
 /********** ========================================= **********/
 /********** Table of dimension-aware image intrinsics **********/
 /********** ========================================= **********/
 
 class ImageDimIntrinsicInfo<AMDGPUImageDimIntrinsic I> {
   Intrinsic Intr = I;
   MIMGBaseOpcode BaseOpcode = !cast<MIMGBaseOpcode>(!strconcat("IMAGE_", I.P.OpMod));
   AMDGPUDimProps Dim = I.P.Dim;

llvm/lib/Target/AMDGPU/SIDefines.h

enum Id { // HwRegCode, (6) [5:0]
   ID_SYMBOLIC_FIRST_GFX10_ = ID_TBA_LO,
   ID_TBA_HI = 17,
   ID_TMA_LO = 18,
   ID_TMA_HI = 19,
   ID_FLAT_SCR_LO = 20,
   ID_FLAT_SCR_HI = 21,
   ID_XNACK_MASK = 22,
   ID_POPS_PACKER = 25,
-  ID_SYMBOLIC_LAST_ = 26,
+  ID_SHADER_CYCLES = 29,
+  ID_SYMBOLIC_FIRST_GFX1030_ = ID_SHADER_CYCLES,
+  ID_SYMBOLIC_LAST_ = 30,
   ID_SHIFT_ = 0,
   ID_WIDTH_ = 6,
   ID_MASK_ = (((1 << ID_WIDTH_) - 1) << ID_SHIFT_)
 };
 
 enum Offset : unsigned { // Offset, (5) [10:6]
   OFFSET_DEFAULT_ = 0,
   OFFSET_SHIFT_ = 6,

llvm/lib/Target/AMDGPU/SIISelLowering.cpp

#endif
 
   if (Subtarget->has16BitInsts()) {
     setOperationAction(ISD::FPOW, MVT::f16, Promote);
     setOperationAction(ISD::FLOG, MVT::f16, Custom);
     setOperationAction(ISD::FEXP, MVT::f16, Custom);
     setOperationAction(ISD::FLOG10, MVT::f16, Custom);
   }
 
-  // v_mad_f32 does not support denormals. We report it as unconditionally
-  // legal, and the context where it is formed will disallow it when fp32
-  // denormals are enabled.
+  if (Subtarget->hasMadMacF32Insts())
     setOperationAction(ISD::FMAD, MVT::f32, Legal);
 
   if (!Subtarget->hasBFI()) {
     // fcopysign can be done in a single instruction with BFI.
     setOperationAction(ISD::FCOPYSIGN, MVT::f32, Expand);
     setOperationAction(ISD::FCOPYSIGN, MVT::f64, Expand);
   }
 
   if (!Subtarget->hasBCNT(32))
case Intrinsic::amdgcn_ds_consume: {
     Info.flags = MachineMemOperand::MOLoad | MachineMemOperand::MOStore;
 
     const ConstantInt *Vol = cast<ConstantInt>(CI.getOperand(1));
     if (!Vol->isZero())
       Info.flags |= MachineMemOperand::MOVolatile;
 
     return true;
   }
+  case Intrinsic::amdgcn_global_atomic_csub: {
+    Info.opc = ISD::INTRINSIC_W_CHAIN;
+    Info.memVT = MVT::getVT(CI.getType());
+    Info.ptrVal = CI.getOperand(0);
+    Info.align.reset();
+    Info.flags = MachineMemOperand::MOLoad |
+                 MachineMemOperand::MOStore |
+                 MachineMemOperand::MODereferenceable |
+                 MachineMemOperand::MOVolatile;
+    return true;
+  }
   case Intrinsic::amdgcn_ds_gws_init:
   case Intrinsic::amdgcn_ds_gws_barrier:
   case Intrinsic::amdgcn_ds_gws_sema_v:
   case Intrinsic::amdgcn_ds_gws_sema_br:
   case Intrinsic::amdgcn_ds_gws_sema_p:
   case Intrinsic::amdgcn_ds_gws_sema_release_all: {
     Info.opc = ISD::INTRINSIC_VOID;
 

 }
 
 bool SITargetLowering::isFMADLegal(const SelectionDAG &DAG,
                                    const SDNode *N) const {
   // TODO: Check future ftz flag
   // v_mad_f32/v_mac_f32 do not support denormals.
   EVT VT = N->getValueType(0);
   if (VT == MVT::f32)
-    return !hasFP32Denormals(DAG.getMachineFunction());
+    return Subtarget->hasMadMacF32Insts() &&
+           !hasFP32Denormals(DAG.getMachineFunction());
   if (VT == MVT::f16) {
     return Subtarget->hasMadF16() &&
            !hasFP64FP16Denormals(DAG.getMachineFunction());
   }
 
   return false;
 }
 
if (LoadVT.getScalarType() == MVT::f16)
                                  M, DAG, Ops);
     return getMemIntrinsicNode(AMDGPUISD::TBUFFER_LOAD_FORMAT, DL,
                                Op->getVTList(), Ops, LoadVT, M->getMemOperand(),
                                DAG);
   }
   case Intrinsic::amdgcn_buffer_atomic_swap:
   case Intrinsic::amdgcn_buffer_atomic_add:
   case Intrinsic::amdgcn_buffer_atomic_sub:
+  case Intrinsic::amdgcn_buffer_atomic_csub:
   case Intrinsic::amdgcn_buffer_atomic_smin:
   case Intrinsic::amdgcn_buffer_atomic_umin:
   case Intrinsic::amdgcn_buffer_atomic_smax:
   case Intrinsic::amdgcn_buffer_atomic_umax:
   case Intrinsic::amdgcn_buffer_atomic_and:
   case Intrinsic::amdgcn_buffer_atomic_or:
   case Intrinsic::amdgcn_buffer_atomic_xor: {
     unsigned Slc = cast<ConstantSDNode>(Op.getOperand(6))->getZExtValue();
case Intrinsic::amdgcn_buffer_atomic_swap:
       Opcode = AMDGPUISD::BUFFER_ATOMIC_SWAP;
       break;
     case Intrinsic::amdgcn_buffer_atomic_add:
       Opcode = AMDGPUISD::BUFFER_ATOMIC_ADD;
       break;
     case Intrinsic::amdgcn_buffer_atomic_sub:
       Opcode = AMDGPUISD::BUFFER_ATOMIC_SUB;
       break;
+    case Intrinsic::amdgcn_buffer_atomic_csub:
+      Opcode = AMDGPUISD::BUFFER_ATOMIC_CSUB;
+      break;
     case Intrinsic::amdgcn_buffer_atomic_smin:
       Opcode = AMDGPUISD::BUFFER_ATOMIC_SMIN;
       break;
     case Intrinsic::amdgcn_buffer_atomic_umin:
       Opcode = AMDGPUISD::BUFFER_ATOMIC_UMIN;
       break;
     case Intrinsic::amdgcn_buffer_atomic_smax:
       Opcode = AMDGPUISD::BUFFER_ATOMIC_SMAX;
case Intrinsic::amdgcn_struct_buffer_atomic_cmpswap: {
     EVT VT = Op.getValueType();
     auto *M = cast<MemSDNode>(Op);
     M->getMemOperand()->setOffset(getBufferOffsetForMMO(Ops[5], Ops[6], Ops[7],
                                                         Ops[4]));
 
     return DAG.getMemIntrinsicNode(AMDGPUISD::BUFFER_ATOMIC_CMPSWAP, DL,
                                    Op->getVTList(), Ops, VT, M->getMemOperand());
   }
+  case Intrinsic::amdgcn_global_atomic_csub: {
+    MemSDNode *M = cast<MemSDNode>(Op);
+    SDValue Ops[] = {
+      M->getOperand(0), // Chain
+      M->getOperand(2), // Ptr
+      M->getOperand(3)  // Value
+    };
+
+    return DAG.getMemIntrinsicNode(AMDGPUISD::ATOMIC_LOAD_CSUB, SDLoc(Op),
+                                   M->getVTList(), Ops, M->getMemoryVT(),
+                                   M->getMemOperand());
+  }
 
   default:
     if (const AMDGPU::ImageDimIntrinsicInfo *ImageDimIntr =
             AMDGPU::getImageDimIntrinsicInfo(IntrID))
       return lowerImage(Op, ImageDimIntr, DAG);
 
     return SDValue();
   }

llvm/lib/Target/AMDGPU/SIInstrInfo.td

def SIEncodingFamily {
   int SI = 0;
   int VI = 1;
   int SDWA = 2;
   int SDWA9 = 3;
   int GFX80 = 4;
   int GFX9 = 5;
   int GFX10 = 6;
   int SDWA10 = 7;
+  int GFX10_B = 8;
 }
 
 //===----------------------------------------------------------------------===//
 // SI DAG Nodes
 //===----------------------------------------------------------------------===//
 
 def AMDGPUclamp : SDNode<"AMDGPUISD::CLAMP", SDTFPUnaryOp>;
 

 def SIatomic_inc : SDNode<"AMDGPUISD::ATOMIC_INC", SDTAtomic2,
   [SDNPMayLoad, SDNPMayStore, SDNPMemOperand, SDNPHasChain]
 >;
 
 def SIatomic_dec : SDNode<"AMDGPUISD::ATOMIC_DEC", SDTAtomic2,
   [SDNPMayLoad, SDNPMayStore, SDNPMemOperand, SDNPHasChain]
 >;
 
+def SIatomic_csub : SDNode<"AMDGPUISD::ATOMIC_LOAD_CSUB", SDTAtomic2,
+  [SDNPMayLoad, SDNPMayStore, SDNPMemOperand, SDNPHasChain]
+>;
+
 def SDTAtomic2_f32 : SDTypeProfile<1, 2, [
   SDTCisSameAs<0,2>, SDTCisFP<0>, SDTCisPtrTy<1>
 ]>;
 
 def SIatomic_fmin : SDNode<"AMDGPUISD::ATOMIC_LOAD_FMIN", SDTAtomic2_f32,
   [SDNPMayLoad, SDNPMayStore, SDNPMemOperand, SDNPHasChain]
 >;
 

 def SIbuffer_atomic_umin : SDBufferAtomic <"AMDGPUISD::BUFFER_ATOMIC_UMIN">;
 def SIbuffer_atomic_smax : SDBufferAtomic <"AMDGPUISD::BUFFER_ATOMIC_SMAX">;
 def SIbuffer_atomic_umax : SDBufferAtomic <"AMDGPUISD::BUFFER_ATOMIC_UMAX">;
 def SIbuffer_atomic_and : SDBufferAtomic <"AMDGPUISD::BUFFER_ATOMIC_AND">;
 def SIbuffer_atomic_or : SDBufferAtomic <"AMDGPUISD::BUFFER_ATOMIC_OR">;
 def SIbuffer_atomic_xor : SDBufferAtomic <"AMDGPUISD::BUFFER_ATOMIC_XOR">;
 def SIbuffer_atomic_inc : SDBufferAtomic <"AMDGPUISD::BUFFER_ATOMIC_INC">;
 def SIbuffer_atomic_dec : SDBufferAtomic <"AMDGPUISD::BUFFER_ATOMIC_DEC">;
+def SIbuffer_atomic_csub : SDBufferAtomic <"AMDGPUISD::BUFFER_ATOMIC_CSUB">;
 def SIbuffer_atomic_fadd : SDBufferAtomicNoRtn <"AMDGPUISD::BUFFER_ATOMIC_FADD", f32>;
 def SIbuffer_atomic_pk_fadd : SDBufferAtomicNoRtn <"AMDGPUISD::BUFFER_ATOMIC_PK_FADD", v2f16>;
 
 def SIbuffer_atomic_cmpswap : SDNode <"AMDGPUISD::BUFFER_ATOMIC_CMPSWAP",
   SDTypeProfile<1, 9,
     [SDTCisVT<0, i32>,   // dst
      SDTCisVT<1, i32>,   // src
      SDTCisVT<2, i32>,   // cmp
!if(!eq(SrcVT.Value, v2i16.Value), 1,
         !if(!eq(SrcVT.Value, v4f16.Value), 1, 0)
     ));
 }
 
 //===----------------------------------------------------------------------===//
 // PatFrags for global memory operations
 //===----------------------------------------------------------------------===//
 
+let AddressSpaces = !cast<AddressSpaceList>("LoadAddress_global").AddrSpaces in {
+defm atomic_csub_global : binary_atomic_op<SIatomic_csub>;
+}
+
 foreach as = [ "global", "flat", "constant", "local", "private", "region" ] in {
 let AddressSpaces = !cast<AddressSpaceList>("LoadAddress_"#as).AddrSpaces in {
 
 
 defm atomic_inc_#as : binary_atomic_op<SIatomic_inc>;
 defm atomic_dec_#as : binary_atomic_op<SIatomic_dec>;
 defm atomic_load_fmin_#as : binary_atomic_op<SIatomic_fmin, 0>;
 defm atomic_load_fmax_#as : binary_atomic_op<SIatomic_fmax, 0>;
multiclass SIAtomicM0Glue2 <string op_name, bit is_amdgpu = 0,
 
   let AddressSpaces = StoreAddress_region.AddrSpaces in {
     defm _region_m0 : binary_atomic_op <!cast<SDNode>(NAME#"_glue"), IsInt>;
   }
 }
 
 defm atomic_load_add : SIAtomicM0Glue2 <"LOAD_ADD">;
 defm atomic_load_sub : SIAtomicM0Glue2 <"LOAD_SUB">;
+defm atomic_load_csub : SIAtomicM0Glue2 <"LOAD_CSUB", 1>;
 defm atomic_inc : SIAtomicM0Glue2 <"INC", 1>;
 defm atomic_dec : SIAtomicM0Glue2 <"DEC", 1>;
 defm atomic_load_and : SIAtomicM0Glue2 <"LOAD_AND">;
 defm atomic_load_min : SIAtomicM0Glue2 <"LOAD_MIN">;
 defm atomic_load_max : SIAtomicM0Glue2 <"LOAD_MAX">;
 defm atomic_load_or : SIAtomicM0Glue2 <"LOAD_OR">;
 defm atomic_load_xor : SIAtomicM0Glue2 <"LOAD_XOR">;
 defm atomic_load_umin : SIAtomicM0Glue2 <"LOAD_UMIN">;
def HWREG {
   int TBA_LO = 16;
   int TBA_HI = 17;
   int TMA_LO = 18;
   int TMA_HI = 19;
   int FLAT_SCR_LO = 20;
   int FLAT_SCR_HI = 21;
   int XNACK_MASK = 22;
   int POPS_PACKER = 25;
+  int SHADER_CYCLES = 29;
 }
 
 class getHwRegImm<int Reg, int Offset = 0, int Size = 32> {
   int ret = !or(Reg,
                 !or(!shl(Offset, 6),
                     !shl(!add(Size, -1), 11)));
 }
 

llvm/lib/Target/AMDGPU/SIInstructions.td

class FMADModsPat<ValueType Ty, Instruction inst, SDPatternOperator mad_opr>
   : GCNPat<
   (Ty (mad_opr (Ty (VOP3Mods Ty:$src0, i32:$src0_mod)),
                (Ty (VOP3Mods Ty:$src1, i32:$src1_mod)),
                (Ty (VOP3Mods Ty:$src2, i32:$src2_mod)))),
   (inst $src0_mod, $src0, $src1_mod, $src1,
   $src2_mod, $src2, DSTCLAMP.NONE, DSTOMOD.NONE)
 >;
 
+let SubtargetPredicate = HasMadMacF32Insts in
 def : FMADModsPat<f32, V_MAD_F32, AMDGPUfmad_ftz>;
 def : FMADModsPat<f16, V_MAD_F16, AMDGPUfmad_ftz> {
   let SubtargetPredicate = Has16BitInsts;
 }
 
 class VOPSelectModsPat <ValueType vt> : GCNPat <
   (vt (select i1:$src0, (VOP3Mods vt:$src1, i32:$src1_mods),
                         (VOP3Mods vt:$src2, i32:$src2_mods))),

llvm/lib/Target/AMDGPU/SIShrinkInstructions.cpp

if ((TII->sopkIsZext(SOPKOpc) && isKUImmOperand(TII, Src1)) ||
       (!TII->sopkIsZext(SOPKOpc) && isKImmOperand(TII, Src1))) {
     MI.setDesc(NewDesc);
   }
 }
 
 // Shrink NSA encoded instructions with contiguous VGPRs to non-NSA encoding.
 void SIShrinkInstructions::shrinkMIMG(MachineInstr &MI) {
   const AMDGPU::MIMGInfo *Info = AMDGPU::getMIMGInfo(MI.getOpcode());
-  if (Info->MIMGEncoding != AMDGPU::MIMGEncGfx10NSA)
+  if (!Info || Info->MIMGEncoding != AMDGPU::MIMGEncGfx10NSA)
     return;
 
   MachineFunction *MF = MI.getParent()->getParent();
   const GCNSubtarget &ST = MF->getSubtarget<GCNSubtarget>();
   const SIInstrInfo *TII = ST.getInstrInfo();
   const SIRegisterInfo &TRI = TII->getRegisterInfo();
   int VAddr0Idx =
       AMDGPU::getNamedOperandIdx(MI.getOpcode(), AMDGPU::OpName::vaddr0);

llvm/lib/Target/AMDGPU/SMInstructions.td

 >;
 
 defm S_BUFFER_STORE_DWORDX4 : SM_Pseudo_Stores <
   "s_buffer_store_dwordx4", SReg_128, SReg_128
 >;
 }
 } // End SubtargetPredicate = HasScalarStores
 
+let SubtargetPredicate = HasSMemTimeInst in
 def S_MEMTIME : SM_Time_Pseudo <"s_memtime", int_amdgcn_s_memtime>;
 def S_DCACHE_INV : SM_Inval_Pseudo <"s_dcache_inv", int_amdgcn_s_dcache_inv>;
 
 let SubtargetPredicate = isGFX7GFX8GFX9 in {
 def S_DCACHE_INV_VOL : SM_Inval_Pseudo <"s_dcache_inv_vol", int_amdgcn_s_dcache_inv_vol>;
 } // let SubtargetPredicate = isGFX7GFX8GFX9
 
 let SubtargetPredicate = isGFX8Plus in {
 let OtherPredicates = [HasScalarStores] in {
 def S_DCACHE_WB     : SM_Inval_Pseudo <"s_dcache_wb", int_amdgcn_s_dcache_wb>;
 def S_DCACHE_WB_VOL : SM_Inval_Pseudo <"s_dcache_wb_vol", int_amdgcn_s_dcache_wb_vol>;
 } // End OtherPredicates = [HasScalarStores]
 def S_MEMREALTIME   : SM_Time_Pseudo <"s_memrealtime", int_amdgcn_s_memrealtime>;
 
 defm S_ATC_PROBE        : SM_Pseudo_Probe <"s_atc_probe", SReg_64>;
 let is_buffer = 1 in {
 defm S_ATC_PROBE_BUFFER : SM_Pseudo_Probe <"s_atc_probe_buffer", SReg_128>;
 }
 } // SubtargetPredicate = isGFX8Plus
 
-let SubtargetPredicate = isGFX10Plus in {
+let SubtargetPredicate = isGFX10Plus in
 def S_GL1_INV : SM_Inval_Pseudo<"s_gl1_inv">;
+let SubtargetPredicate = HasGetWaveIdInst in
 def S_GET_WAVEID_IN_WORKGROUP : SM_WaveId_Pseudo <"s_get_waveid_in_workgroup", int_amdgcn_s_get_waveid_in_workgroup>;
-} // End SubtargetPredicate = isGFX10Plus
+
 
 let SubtargetPredicate = HasScalarFlatScratchInsts, Uses = [FLAT_SCR] in {
 defm S_SCRATCH_LOAD_DWORD    : SM_Pseudo_Loads <"s_scratch_load_dword",   SReg_64, SReg_32_XM0_XEXEC>;
 defm S_SCRATCH_LOAD_DWORDX2  : SM_Pseudo_Loads <"s_scratch_load_dwordx2", SReg_64, SReg_64_XEXEC>;
 defm S_SCRATCH_LOAD_DWORDX4  : SM_Pseudo_Loads <"s_scratch_load_dwordx4", SReg_64, SReg_128>;
 
 defm S_SCRATCH_STORE_DWORD   : SM_Pseudo_Stores <"s_scratch_store_dword",   SReg_64, SReg_32_XM0_XEXEC>;
 defm S_SCRATCH_STORE_DWORDX2 : SM_Pseudo_Stores <"s_scratch_store_dwordx2", SReg_64, SReg_64_XEXEC>;

 
 defm : SMLoad_Pattern <"S_BUFFER_LOAD_DWORD",     f32>;
 defm : SMLoad_Pattern <"S_BUFFER_LOAD_DWORDX2",   v2f32>;
 defm : SMLoad_Pattern <"S_BUFFER_LOAD_DWORDX4",   v4f32>;
 defm : SMLoad_Pattern <"S_BUFFER_LOAD_DWORDX8",   v8f32>;
 defm : SMLoad_Pattern <"S_BUFFER_LOAD_DWORDX16",  v16f32>;
 } // End let AddedComplexity = 100
 
+let OtherPredicates = [HasSMemTimeInst] in {
 def : GCNPat <
   (i64 (readcyclecounter)),
   (S_MEMTIME)
 >;
+} // let OtherPredicates = [HasSMemTimeInst]
+
+let OtherPredicates = [HasNoSMemTimeInst] in {
+def : GCNPat <
+  (i64 (readcyclecounter)),
+  (REG_SEQUENCE SReg_64,
+    (S_GETREG_B32 getHwRegImm<HWREG.SHADER_CYCLES, 0, -12>.ret), sub0,
+    (S_MOV_B32 (i32 0)), sub1)
+>;
+} // let OtherPredicates = [HasNoSMemTimeInst]
 
 //===----------------------------------------------------------------------===//
 // GFX10.
 //===----------------------------------------------------------------------===//
 
 class SMEM_Real_gfx10<bits<8> op, SM_Pseudo ps> :
     SM_Real<ps>, SIMCInstr<ps.PseudoInstr, SIEncodingFamily.GFX10>, Enc64 {
   bit glc;

llvm/lib/Target/AMDGPU/SOPInstructions.td

 
 let SubtargetPredicate = isGFX6GFX7GFX8GFX9 in
 def S_CBRANCH_I_FORK : SOPK_Pseudo <
   "s_cbranch_i_fork",
   (outs), (ins SReg_64:$sdst, sopp_brtarget:$simm16),
   "$sdst, $simm16"
 >;
 
+let hasSideEffects = 1 in {
+
 let mayLoad = 1 in {
+// s_getreg_b32 should use hasSideEffects = 1 for tablegen to allow
+// its use in the readcyclecounter selection.
 def S_GETREG_B32 : SOPK_Pseudo <
   "s_getreg_b32",
   (outs SReg_32:$sdst), (ins hwreg:$simm16),
   "$sdst, $simm16"
 >;
 }
 
-let hasSideEffects = 1, mayLoad = 0, mayStore =0 in {
+let mayLoad = 0, mayStore =0 in {
 
 def S_SETREG_B32 : SOPK_Pseudo <
   "s_setreg_b32",
   (outs), (ins SReg_32:$sdst, hwreg:$simm16),
   "$simm16, $sdst",
   [(int_amdgcn_s_setreg (i32 timm:$simm16), i32:$sdst)]> {
 
   // Use custom inserter to optimize some cases to
def S_SETREG_IMM32_B32 : SOPK_Pseudo <
   (outs), (ins i32imm:$imm, hwreg:$simm16),
   "$simm16, $imm"> {
   let Size = 8; // Unlike every other SOPK instruction.
   let has_sdst = 0;
   let Defs = [MODE];
   let Uses = [MODE];
 }
 
+}
 } // End hasSideEffects = 1
 
 class SOPK_WAITCNT<string opName, list<dag> pat=[]> :
     SOPK_Pseudo<
         opName,
         (outs),
         (ins SReg_32:$sdst, s16imm:$simm16),
         "$sdst, $simm16",

llvm/lib/Target/AMDGPU/Utils/AMDGPUAsmUtils.cpp

const char* const IdSymbolic[] = {
   "HW_REG_TBA_HI",
   "HW_REG_TMA_LO",
   "HW_REG_TMA_HI",
   "HW_REG_FLAT_SCR_LO",
   "HW_REG_FLAT_SCR_HI",
   "HW_REG_XNACK_MASK",
   nullptr, // HW_ID1, no predictable values
   nullptr, // HW_ID2, no predictable values
-  "HW_REG_POPS_PACKER"
+  "HW_REG_POPS_PACKER",
+  nullptr,
+  nullptr,
+  nullptr,
+  "HW_REG_SHADER_CYCLES"
 };
 
 } // namespace Hwreg
 
 namespace Swizzle {
 
 // This must be in sync with llvm::AMDGPU::Swizzle::Id enum members, see SIDefines.h.
 const char* const IdSymbolic[] = {

llvm/lib/Target/AMDGPU/Utils/AMDGPUBaseInfo.h

 bool hasG16(const MCSubtargetInfo &STI);
 bool hasPackedD16(const MCSubtargetInfo &STI);
 
 bool isSI(const MCSubtargetInfo &STI);
 bool isCI(const MCSubtargetInfo &STI);
 bool isVI(const MCSubtargetInfo &STI);
 bool isGFX9(const MCSubtargetInfo &STI);
 bool isGFX10(const MCSubtargetInfo &STI);
+bool isGCN3Encoding(const MCSubtargetInfo &STI);
+bool isGFX10_BEncoding(const MCSubtargetInfo &STI);
+bool hasGFX10_3Insts(const MCSubtargetInfo &STI);
 
 /// Is Reg - scalar register
 bool isSGPR(unsigned Reg, const MCRegisterInfo* TRI);
 
 /// Is there any intersection between registers
 bool isRegIntersect(unsigned Reg0, unsigned Reg1, const MCRegisterInfo* TRI);
 
 /// If \p Reg is a pseudo reg, return the correct hardware register given

llvm/lib/Target/AMDGPU/Utils/AMDGPUBaseInfo.cpp

 unsigned getMinWavesPerEU(const MCSubtargetInfo *STI) {
   return 1;
 }
 
 unsigned getMaxWavesPerEU(const MCSubtargetInfo *STI) {
   // FIXME: Need to take scratch memory into account.
   if (!isGFX10(*STI))
     return 10;
-  return 20;
+  return hasGFX10_3Insts(*STI) ? 16 : 20;
 }
 
 unsigned getWavesPerEUForWorkGroup(const MCSubtargetInfo *STI,
                                    unsigned FlatWorkGroupSize) {
   return divideCeil(getWavesPerWorkGroup(STI, FlatWorkGroupSize),
                     getEUsPerCU(STI));
 }
 
unsigned getNumSGPRBlocks(const MCSubtargetInfo *STI, unsigned NumSGPRs) {
   return NumSGPRs / getSGPREncodingGranule(STI) - 1;
 }
 
 unsigned getVGPRAllocGranule(const MCSubtargetInfo *STI,
                              Optional<bool> EnableWavefrontSize32) {
   bool IsWave32 = EnableWavefrontSize32 ?
       *EnableWavefrontSize32 :
       STI->getFeatureBits().test(FeatureWavefrontSize32);
+
+  if (hasGFX10_3Insts(*STI))
+    return IsWave32 ? 16 : 8;
+
   return IsWave32 ? 8 : 4;
 }
 
 unsigned getVGPREncodingGranule(const MCSubtargetInfo *STI,
                                 Optional<bool> EnableWavefrontSize32) {
-  return getVGPRAllocGranule(STI, EnableWavefrontSize32);
+
+  bool IsWave32 = EnableWavefrontSize32 ?
+      *EnableWavefrontSize32 :
+      STI->getFeatureBits().test(FeatureWavefrontSize32);
+
+  return IsWave32 ? 8 : 4;
 }
 
 unsigned getTotalNumVGPRs(const MCSubtargetInfo *STI) {
   if (!isGFX10(*STI))
     return 256;
   return STI->getFeatureBits().test(FeatureWavefrontSize32) ? 1024 : 512;
 }
 
int64_t getHwregId(const StringRef Name) {
   return ID_UNKNOWN_;
 }
 
 static unsigned getLastSymbolicHwreg(const MCSubtargetInfo &STI) {
   if (isSI(STI) || isCI(STI) || isVI(STI))
     return ID_SYMBOLIC_FIRST_GFX9_;
   else if (isGFX9(STI))
     return ID_SYMBOLIC_FIRST_GFX10_;
+  else if (isGFX10(STI) && !isGFX10_BEncoding(STI))
+    return ID_SYMBOLIC_FIRST_GFX1030_;
   else
     return ID_SYMBOLIC_LAST_;
 }
 
 bool isValidHwreg(int64_t Id, const MCSubtargetInfo &STI) {
-  return ID_SYMBOLIC_FIRST_ <= Id && Id < getLastSymbolicHwreg(STI) &&
-         IdSymbolic[Id];
+  return
+    ID_SYMBOLIC_FIRST_ <= Id && Id < getLastSymbolicHwreg(STI) &&
+    IdSymbolic[Id] && (Id != ID_XNACK_MASK || !AMDGPU::isGFX10_BEncoding(STI));
 }
 
 bool isValidHwreg(int64_t Id) {
   return 0 <= Id && isUInt<ID_WIDTH_>(Id);
 }
 
 bool isValidHwregOffset(int64_t Offset) {
   return 0 <= Offset && isUInt<OFFSET_WIDTH_>(Offset);

 bool isGFX10(const MCSubtargetInfo &STI) {
   return STI.getFeatureBits()[AMDGPU::FeatureGFX10];
 }
 
 bool isGCN3Encoding(const MCSubtargetInfo &STI) {
   return STI.getFeatureBits()[AMDGPU::FeatureGCN3Encoding];
 }
 
+bool isGFX10_BEncoding(const MCSubtargetInfo &STI) {
+  return STI.getFeatureBits()[AMDGPU::FeatureGFX10_BEncoding];
+}
+
+bool hasGFX10_3Insts(const MCSubtargetInfo &STI) {
+  return STI.getFeatureBits()[AMDGPU::FeatureGFX10_3Insts];
+}
+
 bool isSGPR(unsigned Reg, const MCRegisterInfo* TRI) {
   const MCRegisterClass SGPRClass = TRI->getRegClass(AMDGPU::SReg_32RegClassID);
   const unsigned FirstSubReg = TRI->getSubReg(Reg, AMDGPU::sub0);
   return SGPRClass.contains(FirstSubReg != 0 ? FirstSubReg : Reg) ||
     Reg == AMDGPU::SCC;
 }
 
 bool isRegIntersect(unsigned Reg0, unsigned Reg1, const MCRegisterInfo* TRI) {

llvm/lib/Target/AMDGPU/VOP2Instructions.td

def VOP_WRITELANE : VOPProfile<[i32, i32, i32, i32]> {
   let HasExtSDWA9 = 0;
 }
 
 //===----------------------------------------------------------------------===//
 // VOP2 Instructions
 //===----------------------------------------------------------------------===//
 
 defm V_CNDMASK_B32 : VOP2eInst <"v_cndmask_b32", VOP2e_I32_I32_I32_I1>;
+let SubtargetPredicate = HasMadMacF32Insts in
 def V_MADMK_F32 : VOP2_Pseudo <"v_madmk_f32", VOP_MADMK_F32, []>;
 
 let isCommutable = 1 in {
 defm V_ADD_F32 : VOP2Inst <"v_add_f32", VOP_F32_F32_F32, any_fadd>;
 defm V_SUB_F32 : VOP2Inst <"v_sub_f32", VOP_F32_F32_F32, fsub>;
 defm V_SUBREV_F32 : VOP2Inst <"v_subrev_f32", VOP_F32_F32_F32, null_frag, "v_sub_f32">;
 defm V_MUL_LEGACY_F32 : VOP2Inst <"v_mul_legacy_f32", VOP_F32_F32_F32, AMDGPUfmul_legacy>;
 defm V_MUL_F32 : VOP2Inst <"v_mul_f32", VOP_F32_F32_F32, any_fmul>;

 defm V_LSHRREV_B32 : VOP2Inst <"v_lshrrev_b32", VOP_I32_I32_I32, lshr_rev, "v_lshr_b32">;
 defm V_ASHRREV_I32 : VOP2Inst <"v_ashrrev_i32", VOP_I32_I32_I32, ashr_rev, "v_ashr_i32">;
 defm V_LSHLREV_B32 : VOP2Inst <"v_lshlrev_b32", VOP_I32_I32_I32, lshl_rev, "v_lshl_b32">;
 defm V_AND_B32 : VOP2Inst <"v_and_b32", VOP_PAT_GEN<VOP_I32_I32_I32>, and>;
 defm V_OR_B32 : VOP2Inst <"v_or_b32", VOP_PAT_GEN<VOP_I32_I32_I32>, or>;
 defm V_XOR_B32 : VOP2Inst <"v_xor_b32", VOP_PAT_GEN<VOP_I32_I32_I32>, xor>;
 
 let mayRaiseFPException = 0 in {
+let SubtargetPredicate = HasMadMacF32Insts in {
 let Constraints = "$vdst = $src2", DisableEncoding="$src2",
     isConvertibleToThreeAddress = 1 in {
 defm V_MAC_F32 : VOP2Inst <"v_mac_f32", VOP_MAC_F32>;
 }
 
 def V_MADAK_F32 : VOP2_Pseudo <"v_madak_f32", VOP_MADAK_F32, []>;
+} // End SubtargetPredicate = HasMadMacF32Insts
 }
 
 // No patterns so that the scalar instructions are always selected.
 // The scalar versions will be replaced with vector when needed later.
 
 // V_ADD_I32, V_SUB_I32, and V_SUBREV_I32 where renamed to *_U32 in VI,
 // but the VI instructions behave the same as the SI versions.
 defm V_ADD_I32 : VOP2bInst <"v_add_i32", VOP2b_I32_I1_I32_I32, null_frag, "v_add_i32", 1>;

 
 let SubtargetPredicate = isGFX6GFX7 in {
 defm V_MIN_LEGACY_F32 : VOP2Inst <"v_min_legacy_f32", VOP_F32_F32_F32, AMDGPUfmin_legacy>;
 defm V_MAX_LEGACY_F32 : VOP2Inst <"v_max_legacy_f32", VOP_F32_F32_F32, AMDGPUfmax_legacy>;
 } // End SubtargetPredicate = isGFX6GFX7
 
 let isCommutable = 1 in {
 let SubtargetPredicate = isGFX6GFX7GFX10 in {
+let OtherPredicates = [HasMadMacF32Insts] in
 defm V_MAC_LEGACY_F32 : VOP2Inst <"v_mac_legacy_f32", VOP_F32_F32_F32>;
 } // End SubtargetPredicate = isGFX6GFX7GFX10
 let SubtargetPredicate = isGFX6GFX7 in {
 defm V_LSHR_B32 : VOP2Inst <"v_lshr_b32", VOP_PAT_GEN<VOP_I32_I32_I32>, srl>;
 defm V_ASHR_I32 : VOP2Inst <"v_ashr_i32", VOP_PAT_GEN<VOP_I32_I32_I32>, sra>;
 defm V_LSHL_B32 : VOP2Inst <"v_lshl_b32", VOP_PAT_GEN<VOP_I32_I32_I32>, shl>;
 } // End SubtargetPredicate = isGFX6GFX7
 } // End isCommutable = 1

 
 let SubtargetPredicate = isGFX6GFX7 in {
   defm : VOP2eInstAliases<V_CNDMASK_B32_e32, V_CNDMASK_B32_e32_gfx6_gfx7>;
 } // End SubtargetPredicate = isGFX6GFX7
 
 defm V_ADD_F32            : VOP2_Real_gfx6_gfx7_gfx10<0x003>;
 defm V_SUB_F32            : VOP2_Real_gfx6_gfx7_gfx10<0x004>;
 defm V_SUBREV_F32         : VOP2_Real_gfx6_gfx7_gfx10<0x005>;
+let OtherPredicates = [HasMadMacF32Insts] in
 defm V_MAC_LEGACY_F32     : VOP2_Real_gfx6_gfx7_gfx10<0x006>;
 defm V_MUL_LEGACY_F32     : VOP2_Real_gfx6_gfx7_gfx10<0x007>;
 defm V_MUL_F32            : VOP2_Real_gfx6_gfx7_gfx10<0x008>;
 defm V_MUL_I32_I24        : VOP2_Real_gfx6_gfx7_gfx10<0x009>;
 defm V_MUL_HI_I32_I24     : VOP2_Real_gfx6_gfx7_gfx10<0x00a>;
 defm V_MUL_U32_U24        : VOP2_Real_gfx6_gfx7_gfx10<0x00b>;
 defm V_MUL_HI_U32_U24     : VOP2_Real_gfx6_gfx7_gfx10<0x00c>;
 defm V_MIN_F32            : VOP2_Real_gfx6_gfx7_gfx10<0x00f>;

 }
 let SubtargetPredicate = HasDot3Insts in {
   defm V_DOT8C_I32_I4  : VOP2_Real_DOT_ACC_gfx9<0x3a>;
 }
 
 let SubtargetPredicate = HasPkFmacF16Inst in {
 defm V_PK_FMAC_F16 : VOP2_Real_e32_vi<0x3c>;
 } // End SubtargetPredicate = HasPkFmacF16Inst
+
+let SubtargetPredicate = HasDot3Insts in {
+  // NB: Opcode conflicts with V_DOT2C_F32_F16
+  let DecoderNamespace = "GFX10_B" in
+  defm V_DOT8C_I32_I4 : VOP2_Real_DOT_ACC_gfx10<0x02>;
+}

llvm/lib/Target/AMDGPU/VOP3Instructions.td

 
 //===----------------------------------------------------------------------===//
 // VOP3 Instructions
 //===----------------------------------------------------------------------===//
 
 let isCommutable = 1 in {
 
 let mayRaiseFPException = 0 in {
+let SubtargetPredicate = HasMadMacF32Insts in {
 def V_MAD_LEGACY_F32 : VOP3Inst <"v_mad_legacy_f32", VOP3_Profile<VOP_F32_F32_F32_F32>>;
 def V_MAD_F32 : VOP3Inst <"v_mad_f32", VOP3_Profile<VOP_F32_F32_F32_F32>, fmad>;
+} // End SubtargetPredicate = HasMadMacInsts
+
+let SubtargetPredicate = HasNoMadMacF32Insts in
+def V_FMA_LEGACY_F32 : VOP3Inst <"v_fma_legacy_f32", VOP3_Profile<VOP_F32_F32_F32_F32>>;
 }
 
 def V_MAD_I32_I24 : VOP3Inst <"v_mad_i32_i24", VOP3_Profile<VOP_I32_I32_I32_I32, VOP3_CLAMP>>;
 def V_MAD_U32_U24 : VOP3Inst <"v_mad_u32_u24", VOP3_Profile<VOP_I32_I32_I32_I32, VOP3_CLAMP>>;
 def V_FMA_F32 : VOP3Inst <"v_fma_f32", VOP3_Profile<VOP_F32_F32_F32_F32>, any_fma>;
 def V_LERP_U8 : VOP3Inst <"v_lerp_u8", VOP3_Profile<VOP_I32_I32_I32_I32>, int_amdgcn_lerp>;
 
 let SchedRW = [WriteDoubleAdd] in {

 defm V_DIV_FMAS_F32    : VOP3_Real_gfx6_gfx7_gfx10<0x16f>;
 defm V_DIV_FMAS_F64    : VOP3_Real_gfx6_gfx7_gfx10<0x170>;
 defm V_MSAD_U8         : VOP3_Real_gfx6_gfx7_gfx10<0x171>;
 defm V_MQSAD_PK_U16_U8 : VOP3_Real_gfx6_gfx7_gfx10<0x173>;
 defm V_TRIG_PREOP_F64  : VOP3_Real_gfx6_gfx7_gfx10<0x174>;
 defm V_DIV_SCALE_F32   : VOP3be_Real_gfx6_gfx7_gfx10<0x16d>;
 defm V_DIV_SCALE_F64   : VOP3be_Real_gfx6_gfx7_gfx10<0x16e>;
 
+// NB: Same opcode as v_mad_legacy_f32
+let DecoderNamespace = "GFX10_B" in
+defm V_FMA_LEGACY_F32  : VOP3_Real_gfx10<0x140>;
+
 //===----------------------------------------------------------------------===//
 // GFX8, GFX9 (VI).
 //===----------------------------------------------------------------------===//
 
 let AssemblerPredicate = isGFX8GFX9, DecoderNamespace = "GFX8" in {
 
 multiclass VOP3_Real_vi<bits<10> op> {
   def _vi : VOP3_Real<!cast<VOP_Pseudo>(NAME), SIEncodingFamily.VI>,

llvm/test/CodeGen/AMDGPU/GlobalISel/udiv.i64.ll

 ; CGP-NEXT:    s_mov_b32 s4, 0xffffff
 ; CGP-NEXT:    v_and_b32_e32 v0, s4, v0
 ; CGP-NEXT:    v_and_b32_e32 v1, s4, v2
 ; CGP-NEXT:    v_cvt_f32_u32_e32 v0, v0
 ; CGP-NEXT:    v_cvt_f32_u32_e32 v1, v1
 ; CGP-NEXT:    v_rcp_f32_e32 v2, v1
 ; CGP-NEXT:    v_mul_f32_e32 v2, v0, v2
 ; CGP-NEXT:    v_trunc_f32_e32 v2, v2
-; CGP-NEXT:    v_mad_f32 v0, -v2, v1, v0
+; CGP-NEXT:    v_fma_f32 v0, -v2, v1, v0
 ; CGP-NEXT:    v_cvt_u32_f32_e32 v2, v2
 ; CGP-NEXT:    v_cmp_ge_f32_e64 s[4:5], |v0|, v1
 ; CGP-NEXT:    v_cndmask_b32_e64 v0, 0, 1, s[4:5]
 ; CGP-NEXT:    v_add_i32_e32 v0, vcc, v2, v0
 ; CGP-NEXT:    v_and_b32_e32 v0, 0xffffff, v0
 ; CGP-NEXT:    v_mov_b32_e32 v1, 0
 ; CGP-NEXT:    s_setpc_b64 s[30:31]
   %num.mask = and i64 %num, 16777215

 ; CGP-NEXT:    v_cvt_f32_u32_e32 v2, v2
 ; CGP-NEXT:    v_cvt_f32_u32_e32 v4, v4
 ; CGP-NEXT:    v_rcp_f32_e32 v5, v3
 ; CGP-NEXT:    v_rcp_f32_e32 v6, v4
 ; CGP-NEXT:    v_mul_f32_e32 v5, v0, v5
 ; CGP-NEXT:    v_mul_f32_e32 v6, v2, v6
 ; CGP-NEXT:    v_trunc_f32_e32 v5, v5
 ; CGP-NEXT:    v_trunc_f32_e32 v6, v6
-; CGP-NEXT:    v_mad_f32 v0, -v5, v3, v0
+; CGP-NEXT:    v_fma_f32 v0, -v5, v3, v0
 ; CGP-NEXT:    v_cvt_u32_f32_e32 v5, v5
-; CGP-NEXT:    v_mad_f32 v2, -v6, v4, v2
+; CGP-NEXT:    v_fma_f32 v2, -v6, v4, v2
 ; CGP-NEXT:    v_cvt_u32_f32_e32 v6, v6
 ; CGP-NEXT:    v_cmp_ge_f32_e64 s[4:5], |v0|, v3
 ; CGP-NEXT:    v_cndmask_b32_e64 v0, 0, 1, s[4:5]
 ; CGP-NEXT:    v_cmp_ge_f32_e64 s[4:5], |v2|, v4
 ; CGP-NEXT:    v_cndmask_b32_e64 v2, 0, 1, s[4:5]
 ; CGP-NEXT:    v_add_i32_e32 v0, vcc, v5, v0
 ; CGP-NEXT:    v_add_i32_e32 v2, vcc, v6, v2
 ; CGP-NEXT:    v_and_b32_e32 v0, s6, v0
 ; CGP-NEXT:    v_and_b32_e32 v2, s6, v2
 ; CGP-NEXT:    v_mov_b32_e32 v3, v1
 ; CGP-NEXT:    s_setpc_b64 s[30:31]
   %num.mask = and <2 x i64> %num, <i64 16777215, i64 16777215>
   %den.mask = and <2 x i64> %den, <i64 16777215, i64 16777215>
   %result = udiv <2 x i64> %num.mask, %den.mask
   ret <2 x i64> %result
 }

llvm/test/CodeGen/AMDGPU/GlobalISel/urem.i64.ll

 ; CGP-NEXT:    s_mov_b32 s4, 0xffffff
 ; CGP-NEXT:    v_and_b32_e32 v0, s4, v0
 ; CGP-NEXT:    v_and_b32_e32 v1, s4, v2
 ; CGP-NEXT:    v_cvt_f32_u32_e32 v2, v0
 ; CGP-NEXT:    v_cvt_f32_u32_e32 v3, v1
 ; CGP-NEXT:    v_rcp_f32_e32 v4, v3
 ; CGP-NEXT:    v_mul_f32_e32 v4, v2, v4
 ; CGP-NEXT:    v_trunc_f32_e32 v4, v4
-; CGP-NEXT:    v_mad_f32 v2, -v4, v3, v2
+; CGP-NEXT:    v_fma_f32 v2, -v4, v3, v2
 ; CGP-NEXT:    v_cvt_u32_f32_e32 v4, v4
 ; CGP-NEXT:    v_cmp_ge_f32_e64 s[4:5], |v2|, v3
 ; CGP-NEXT:    v_cndmask_b32_e64 v2, 0, 1, s[4:5]
 ; CGP-NEXT:    v_add_i32_e32 v2, vcc, v4, v2
 ; CGP-NEXT:    v_mul_lo_u32 v1, v2, v1
 ; CGP-NEXT:    v_sub_i32_e32 v0, vcc, v0, v1
 ; CGP-NEXT:    v_and_b32_e32 v0, 0xffffff, v0
 ; CGP-NEXT:    v_mov_b32_e32 v1, 0

 ; CGP-NEXT:    v_cvt_f32_u32_e32 v7, v2
 ; CGP-NEXT:    v_cvt_f32_u32_e32 v8, v4
 ; CGP-NEXT:    v_rcp_f32_e32 v9, v6
 ; CGP-NEXT:    v_rcp_f32_e32 v10, v8
 ; CGP-NEXT:    v_mul_f32_e32 v9, v5, v9
 ; CGP-NEXT:    v_mul_f32_e32 v10, v7, v10
 ; CGP-NEXT:    v_trunc_f32_e32 v9, v9
 ; CGP-NEXT:    v_trunc_f32_e32 v10, v10
-; CGP-NEXT:    v_mad_f32 v5, -v9, v6, v5
+; CGP-NEXT:    v_fma_f32 v5, -v9, v6, v5
 ; CGP-NEXT:    v_cvt_u32_f32_e32 v9, v9
-; CGP-NEXT:    v_mad_f32 v7, -v10, v8, v7
+; CGP-NEXT:    v_fma_f32 v7, -v10, v8, v7
 ; CGP-NEXT:    v_cvt_u32_f32_e32 v10, v10
 ; CGP-NEXT:    v_cmp_ge_f32_e64 s[4:5], |v5|, v6
 ; CGP-NEXT:    v_cndmask_b32_e64 v5, 0, 1, s[4:5]
 ; CGP-NEXT:    v_cmp_ge_f32_e64 s[4:5], |v7|, v8
 ; CGP-NEXT:    v_cndmask_b32_e64 v6, 0, 1, s[4:5]
 ; CGP-NEXT:    v_add_i32_e32 v5, vcc, v9, v5
 ; CGP-NEXT:    v_add_i32_e32 v6, vcc, v10, v6
 ; CGP-NEXT:    v_mul_lo_u32 v3, v5, v3

llvm/test/CodeGen/AMDGPU/elf-header-flags-mach.ll

 ; RUN: llc -filetype=obj -march=amdgcn -mcpu=gfx902 < %s | llvm-readobj -file-headers - | FileCheck --check-prefixes=ALL,ARCH-GCN,GFX902 %s
 ; RUN: llc -filetype=obj -march=amdgcn -mcpu=gfx904 < %s | llvm-readobj -file-headers - | FileCheck --check-prefixes=ALL,ARCH-GCN,GFX904 %s
 ; RUN: llc -filetype=obj -march=amdgcn -mcpu=gfx906 < %s | llvm-readobj -file-headers - | FileCheck --check-prefixes=ALL,ARCH-GCN,GFX906 %s
 ; RUN: llc -filetype=obj -march=amdgcn -mcpu=gfx908 < %s | llvm-readobj -file-headers - | FileCheck --check-prefixes=ALL,ARCH-GCN,GFX908 %s
 ; RUN: llc -filetype=obj -march=amdgcn -mcpu=gfx909 < %s | llvm-readobj -file-headers - | FileCheck --check-prefixes=ALL,ARCH-GCN,GFX909 %s
 ; RUN: llc -filetype=obj -march=amdgcn -mcpu=gfx1010 < %s | llvm-readobj -file-headers - | FileCheck --check-prefixes=ALL,ARCH-GCN,GFX1010 %s
 ; RUN: llc -filetype=obj -march=amdgcn -mcpu=gfx1011 < %s | llvm-readobj -file-headers - | FileCheck --check-prefixes=ALL,ARCH-GCN,GFX1011 %s
 ; RUN: llc -filetype=obj -march=amdgcn -mcpu=gfx1012 < %s | llvm-readobj -file-headers - | FileCheck --check-prefixes=ALL,ARCH-GCN,GFX1012 %s
+; RUN: llc -filetype=obj -march=amdgcn -mcpu=gfx1030 < %s | llvm-readobj -file-headers - | FileCheck --check-prefixes=ALL,ARCH-GCN,GFX1030 %s
 
 ; ARCH-R600: Arch: r600
 ; ARCH-GCN:  Arch: amdgcn
 
 ; ALL:         Flags [
 ; R600:          EF_AMDGPU_MACH_R600_R600     (0x1)
 ; R630:          EF_AMDGPU_MACH_R600_R630     (0x2)
 ; RS880:         EF_AMDGPU_MACH_R600_RS880    (0x3)

 ; GFX904:        EF_AMDGPU_MACH_AMDGCN_GFX904 (0x2E)
 ; GFX906:        EF_AMDGPU_MACH_AMDGCN_GFX906 (0x2F)
 ; GFX908:        EF_AMDGPU_MACH_AMDGCN_GFX908 (0x30)
 ; GFX908-NEXT:   EF_AMDGPU_SRAM_ECC           (0x200)
 ; GFX909:        EF_AMDGPU_MACH_AMDGCN_GFX909 (0x31)
 ; GFX1010:       EF_AMDGPU_MACH_AMDGCN_GFX1010 (0x33)
 ; GFX1011:       EF_AMDGPU_MACH_AMDGCN_GFX1011 (0x34)
 ; GFX1012:       EF_AMDGPU_MACH_AMDGCN_GFX1012 (0x35)
+; GFX1030:       EF_AMDGPU_MACH_AMDGCN_GFX1030 (0x36)
 ; ALL:         ]
 
 define amdgpu_kernel void @elf_header() {
   ret void
 }

llvm/test/CodeGen/AMDGPU/fadd-fma-fmul-combine.ll

-; RUN: llc -march=amdgcn -mattr=+fast-fmaf -denormal-fp-math-f32=preserve-sign -enable-unsafe-fp-math -verify-machineinstrs < %s | FileCheck -check-prefix=GCN -check-prefix=GCN-FLUSH %s
-; RUN: llc -march=amdgcn -mattr=-fast-fmaf -denormal-fp-math-f32=preserve-sign -enable-unsafe-fp-math -verify-machineinstrs < %s | FileCheck -check-prefix=GCN -check-prefix=GCN-FLUSH %s
+; RUN: llc -march=amdgcn -mattr=+fast-fmaf,+mad-mac-f32-insts -denormal-fp-math-f32=preserve-sign -enable-unsafe-fp-math -verify-machineinstrs < %s | FileCheck -check-prefix=GCN -check-prefix=GCN-FLUSH %s
+; RUN: llc -march=amdgcn -mattr=-fast-fmaf,+mad-mac-f32-insts -denormal-fp-math-f32=preserve-sign -enable-unsafe-fp-math -verify-machineinstrs < %s | FileCheck -check-prefix=GCN -check-prefix=GCN-FLUSH %s
 
-; RUN: llc -march=amdgcn -mattr=+fast-fmaf -denormal-fp-math-f32=ieee  -enable-unsafe-fp-math -verify-machineinstrs < %s | FileCheck -check-prefix=GCN -check-prefix=GCN-FASTFMA %s
-; RUN: llc -march=amdgcn -mattr=-fast-fmaf -denormal-fp-math-f32=ieee -enable-unsafe-fp-math -verify-machineinstrs < %s | FileCheck -check-prefix=GCN -check-prefix=GCN-SLOWFMA %s
+; RUN: llc -march=amdgcn -mattr=+fast-fmaf,+mad-mac-f32-insts -denormal-fp-math-f32=ieee  -enable-unsafe-fp-math -verify-machineinstrs < %s | FileCheck -check-prefix=GCN -check-prefix=GCN-FASTFMA %s
+; RUN: llc -march=amdgcn -mattr=-fast-fmaf,+mad-mac-f32-insts -denormal-fp-math-f32=ieee -enable-unsafe-fp-math -verify-machineinstrs < %s | FileCheck -check-prefix=GCN -check-prefix=GCN-SLOWFMA %s
 
 ; FIXME: This should also fold when fma is actually fast if an FMA
 ; exists in the original program.
 
 ; (fadd (fma x, y, (fmul u, v), z) -> (fma x, y (fma u, v, z))
 
 ; GCN-LABEL: {{^}}fast_add_fmuladd_fmul:
 ; GCN: buffer_load_dword [[X:v[0-9]+]]

llvm/test/CodeGen/AMDGPU/fmuladd.f32.ll

 ; RUN: llc -amdgpu-scalarize-global-loads=false -verify-machineinstrs -mcpu=gfx900 -denormal-fp-math-f32=preserve-sign -fp-contract=on < %s | FileCheck -enable-var-scope -check-prefixes=GCN,GCN-FLUSH,GCN-FLUSH-STRICT,GCN-FLUSH-MAD,GFX9-FLUSH,GCN-FLUSH-FASTFMA,GCN-FLUSH-FASTFMA-STRICT,GFX900 %s
 ; RUN: llc -amdgpu-scalarize-global-loads=false -verify-machineinstrs -mcpu=gfx900 -denormal-fp-math-f32=ieee -fp-contract=on < %s | FileCheck -enable-var-scope -check-prefixes=GCN,GCN-DENORM-STRICT,GCN-DENORM,GFX9-DENORM,GCN-DENORM-FASTFMA,GCN-DENORM-FASTFMA-STRICT,GFX900 %s
 
 ; RUN: llc -amdgpu-scalarize-global-loads=false -verify-machineinstrs -mcpu=gfx906 -denormal-fp-math-f32=preserve-sign -fp-contract=on < %s | FileCheck -enable-var-scope -check-prefixes=GCN,GCN-FLUSH,GCN-FLUSH-STRICT,GCN-FLUSH-FMAC,GFX9-FLUSH,GCN-FLUSH-FASTFMA,GCN-FLUSH-FASTFMA-STRICT,GFX906 %s
 
 ; FIXME: Should probably test this, but sometimes selecting fmac is painful to match.
 ; XUN: llc -amdgpu-scalarize-global-loads=false -verify-machineinstrs -mcpu=gfx906 -denormal-fp-math-f32=ieee -fp-contract=on < %s | FileCheck -enable-var-scope -check-prefixes=GCN,GCN-DENORM-STRICT,GCN-DENORM,GFX9-DENORM,GCN-DENORM-FASTFMA,GCN-DENORM-FASTFMA-STRICT,GFX906 %s
 
+; RUN: llc -amdgpu-scalarize-global-loads=false -verify-machineinstrs -mcpu=gfx1030 -mattr=-fp32-denormals -fp-contract=on < %s | FileCheck -enable-var-scope -check-prefixes=GCN,GCN-FLUSH-STRICT,GCN-FLUSH-FMAC,GCN-FLUSH-FASTFMA,GCN-FLUSH-FASTFMA-STRICT %s
+; RUN: llc -amdgpu-scalarize-global-loads=false -verify-machineinstrs -mcpu=gfx1030 -mattr=+fp32-denormals -fp-contract=on < %s | FileCheck -enable-var-scope -check-prefixes=GCN,GCN-DENORM,GCN-DENORM-FASTFMA-STRICT,GCN-DENORM-STRICT %s
 
 ; Test all permutations of: fp32 denormals, fast fp contract, fp contract enabled for fmuladd, fmaf fast/slow.
 
 target triple = "amdgcn--"
 
 
 declare i32 @llvm.amdgcn.workitem.id.x() #1
 declare float @llvm.fmuladd.f32(float, float, float) #1

llvm/test/CodeGen/AMDGPU/frem.ll

-; RUN:  llc -amdgpu-scalarize-global-loads=false  -march=amdgcn -verify-machineinstrs  < %s | FileCheck -check-prefix=SI -check-prefix=GCN -check-prefix=FUNC %s
+; RUN:  llc -amdgpu-scalarize-global-loads=false  -march=amdgcn -mattr=+mad-mac-f32-insts -verify-machineinstrs  < %s | FileCheck -check-prefix=SI -check-prefix=GCN -check-prefix=FUNC %s
 ; RUN:  llc -amdgpu-scalarize-global-loads=false  -march=amdgcn -mcpu=bonaire -verify-machineinstrs < %s | FileCheck -check-prefix=CI -check-prefix=GCN -check-prefix=FUNC %s
 ; RUN:  llc -amdgpu-scalarize-global-loads=false  -march=amdgcn -mcpu=tonga -mattr=-flat-for-global -verify-machineinstrs < %s | FileCheck -check-prefix=CI -check-prefix=GCN -check-prefix=FUNC %s
 
 ; FUNC-LABEL: {{^}}frem_f32:
 ; GCN-DAG: buffer_load_dword [[X:v[0-9]+]], {{.*$}}
 ; GCN-DAG: buffer_load_dword [[Y:v[0-9]+]], {{.*}} offset:16
 ; GCN: v_div_scale_f32
 

llvm/test/CodeGen/AMDGPU/hsa-note-no-func.ll

 ; RUN: llc < %s -mtriple=amdgcn--amdhsa -mcpu=gfx902 -mattr=-code-object-v3 | FileCheck --check-prefix=HSA --check-prefix=HSA-GFX902 %s
 ; RUN: llc < %s -mtriple=amdgcn--amdhsa -mcpu=gfx904 -mattr=-code-object-v3 | FileCheck --check-prefix=HSA --check-prefix=HSA-GFX904 %s
 ; RUN: llc < %s -mtriple=amdgcn--amdhsa -mcpu=gfx906 -mattr=-code-object-v3 | FileCheck --check-prefix=HSA --check-prefix=HSA-GFX906 %s
 ; RUN: llc < %s -mtriple=amdgcn--amdhsa -mcpu=gfx908 -mattr=-code-object-v3 | FileCheck --check-prefix=HSA --check-prefix=HSA-GFX908 %s
 ; RUN: llc < %s -mtriple=amdgcn--amdhsa -mcpu=gfx909 -mattr=-code-object-v3 | FileCheck --check-prefix=HSA --check-prefix=HSA-GFX909 %s
 ; RUN: llc < %s -mtriple=amdgcn--amdhsa -mcpu=gfx1010 -mattr=-code-object-v3 | FileCheck --check-prefix=HSA --check-prefix=HSA-GFX1010 %s
 ; RUN: llc < %s -mtriple=amdgcn--amdhsa -mcpu=gfx1011 -mattr=-code-object-v3 | FileCheck --check-prefix=HSA --check-prefix=HSA-GFX1011 %s
 ; RUN: llc < %s -mtriple=amdgcn--amdhsa -mcpu=gfx1012 -mattr=-code-object-v3 | FileCheck --check-prefix=HSA --check-prefix=HSA-GFX1012 %s
+; RUN: llc < %s -mtriple=amdgcn--amdhsa -mcpu=gfx1030 -mattr=-code-object-v3 | FileCheck --check-prefix=HSA --check-prefix=HSA-GFX1030 %s
 
 ; HSA: .hsa_code_object_version 2,1
 ; HSA-SI600: .hsa_code_object_isa 6,0,0,"AMD","AMDGPU"
 ; HSA-SI601: .hsa_code_object_isa 6,0,1,"AMD","AMDGPU"
 ; HSA-CI700: .hsa_code_object_isa 7,0,0,"AMD","AMDGPU"
 ; HSA-CI701: .hsa_code_object_isa 7,0,1,"AMD","AMDGPU"
 ; HSA-CI702: .hsa_code_object_isa 7,0,2,"AMD","AMDGPU"
 ; HSA-CI703: .hsa_code_object_isa 7,0,3,"AMD","AMDGPU"
 ; HSA-CI704: .hsa_code_object_isa 7,0,4,"AMD","AMDGPU"
 ; HSA-VI801: .hsa_code_object_isa 8,0,1,"AMD","AMDGPU"
 ; HSA-VI802: .hsa_code_object_isa 8,0,2,"AMD","AMDGPU"
 ; HSA-VI803: .hsa_code_object_isa 8,0,3,"AMD","AMDGPU"
 ; HSA-VI810: .hsa_code_object_isa 8,1,0,"AMD","AMDGPU"
 ; HSA-GFX900: .hsa_code_object_isa 9,0,0,"AMD","AMDGPU"
 ; HSA-GFX902: .hsa_code_object_isa 9,0,2,"AMD","AMDGPU"
 ; HSA-GFX904: .hsa_code_object_isa 9,0,4,"AMD","AMDGPU"
 ; HSA-GFX906: .hsa_code_object_isa 9,0,6,"AMD","AMDGPU"
 ; HSA-GFX908: .hsa_code_object_isa 9,0,8,"AMD","AMDGPU"
 ; HSA-GFX909: .hsa_code_object_isa 9,0,9,"AMD","AMDGPU"
 ; HSA-GFX1010: .hsa_code_object_isa 10,1,0,"AMD","AMDGPU"
 ; HSA-GFX1011: .hsa_code_object_isa 10,1,1,"AMD","AMDGPU"
 ; HSA-GFX1012: .hsa_code_object_isa 10,1,2,"AMD","AMDGPU"
+; HSA-GFX1030: .hsa_code_object_isa 10,3,0,"AMD","AMDGPU"

llvm/test/CodeGen/AMDGPU/idot8s.ll

 ; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py
 ; RUN: llc -mtriple=amdgcn -mcpu=gfx700 -verify-machineinstrs < %s | FileCheck -check-prefixes=GFX7 %s
 ; RUN: llc -mtriple=amdgcn -mcpu=gfx803 -verify-machineinstrs < %s | FileCheck -check-prefixes=GFX8 %s
 ; RUN: llc -mtriple=amdgcn -mcpu=gfx900 -verify-machineinstrs < %s | FileCheck -check-prefixes=GFX9 %s
 ; RUN: llc -mtriple=amdgcn -mcpu=gfx906 -verify-machineinstrs < %s | FileCheck -check-prefixes=GFX9-DL %s
 ; RUN: llc -mtriple=amdgcn -mcpu=gfx1011 -verify-machineinstrs < %s | FileCheck -check-prefixes=GFX10-DL %s
 ; RUN: llc -mtriple=amdgcn -mcpu=gfx1012 -verify-machineinstrs < %s | FileCheck -check-prefixes=GFX10-DL %s
+; RUN: llc -mtriple=amdgcn -mcpu=gfx1030 -verify-machineinstrs < %s | FileCheck -check-prefixes=GFX10-DL %s
 
 define amdgpu_kernel void @idot8_acc32(<8 x i4> addrspace(1)* %src1,
 ; GFX7-LABEL: idot8_acc32:
 ; GFX7:       ; %bb.0: ; %entry
 ; GFX7-NEXT:    s_load_dwordx4 s[4:7], s[0:1], 0x9
 ; GFX7-NEXT:    s_load_dwordx2 s[0:1], s[0:1], 0xd
 ; GFX7-NEXT:    s_mov_b32 s3, 0xf000
 ; GFX7-NEXT:    s_mov_b32 s2, -1

llvm/test/CodeGen/AMDGPU/llvm.amdgcn.atomic.csub.ll

+; RUN: llc < %s -march=amdgcn -mcpu=gfx1030 -verify-machineinstrs | FileCheck %s -check-prefix=GCN
+
+declare i32 @llvm.amdgcn.buffer.atomic.csub(i32, <4 x i32>, i32, i32, i1)
+declare i32 @llvm.amdgcn.global.atomic.csub(i32 addrspace(1)*, i32)
+
+; GCN-LABEL: {{^}}buffer_atomic_csub:
+; GCN: buffer_atomic_csub v0, v1, s[0:3], 0 idxen glc
+define amdgpu_ps void @buffer_atomic_csub(<4 x i32> inreg %rsrc, i32 %data, i32 %vindex) {
+main_body:
+  %ret = call i32 @llvm.amdgcn.buffer.atomic.csub(i32 %data, <4 x i32> %rsrc, i32 %vindex, i32 0, i1 0)
+  ret void
+}
+
+; GCN-LABEL: {{^}}buffer_atomic_csub_off4_slc:
+; GCN: buffer_atomic_csub v0, v1, s[0:3], 0 idxen offset:4 glc slc
+define amdgpu_ps void @buffer_atomic_csub_off4_slc(<4 x i32> inreg %rsrc, i32 %data, i32 %vindex) {
+main_body:
+  %ret = call i32 @llvm.amdgcn.buffer.atomic.csub(i32 %data, <4 x i32> %rsrc, i32 %vindex, i32 4, i1 1)
+  ret void
+}
+
+; GCN-LABEL: {{^}}global_atomic_csub:
+; GCN: global_atomic_csub v{{[0-9]+}}, v[{{[0-9:]+}}], v{{[0-9]+}}, off glc
+define amdgpu_kernel void @global_atomic_csub(i32 addrspace(1)* %ptr, i32 %data) {
+main_body:
+  %ret = call i32 @llvm.amdgcn.global.atomic.csub(i32 addrspace(1)* %ptr, i32 %data)
+  ret void
+}
+
+; GCN-LABEL: {{^}}global_atomic_csub_off4:
+; GCN: global_atomic_csub v{{[0-9]+}}, v[{{[0-9:]+}}], v{{[0-9]+}}, off offset:4 glc
+define amdgpu_kernel void @global_atomic_csub_off4(i32 addrspace(1)* %ptr, i32 %data) {
+main_body:
+  %p = getelementptr i32, i32 addrspace(1)* %ptr, i64 1
+  %ret = call i32 @llvm.amdgcn.global.atomic.csub(i32 addrspace(1)* %p, i32 %data)
+  ret void
+}

llvm/test/CodeGen/AMDGPU/llvm.amdgcn.image.msaa.load.ll

+; RUN: llc -march=amdgcn -mcpu=gfx1030 -verify-machineinstrs -show-mc-encoding < %s | FileCheck -check-prefixes=GCN,GFX10 %s
+
+; GCN-LABEL: {{^}}load_1d:
+; GFX10: image_msaa_load v[0:3], v0, s[0:7] dmask:0xf dim:SQ_RSRC_IMG_1D unorm ;
+define amdgpu_ps <4 x float> @load_1d(<8 x i32> inreg %rsrc, i32 %s) {
+main_body:
+  %v = call <4 x float> @llvm.amdgcn.image.msaa.load.1d.v4f32.i32(i32 15, i32 %s, <8 x i32> %rsrc, i32 0, i32 0)
+  ret <4 x float> %v
+}
+
+; GCN-LABEL: {{^}}load_1d_tfe:
+; GFX10: image_msaa_load v[0:4], v{{[0-9]+}}, s[0:7] dmask:0xf dim:SQ_RSRC_IMG_1D unorm tfe ;
+define amdgpu_ps <4 x float> @load_1d_tfe(<8 x i32> inreg %rsrc, i32 addrspace(1)* inreg %out, i32 %s) {
+main_body:
+  %v = call {<4 x float>,i32} @llvm.amdgcn.image.msaa.load.1d.v4f32i32.i32(i32 15, i32 %s, <8 x i32> %rsrc, i32 1, i32 0)
+  %v.vec = extractvalue {<4 x float>, i32} %v, 0
+  %v.err = extractvalue {<4 x float>, i32} %v, 1
+  store i32 %v.err, i32 addrspace(1)* %out, align 4
+  ret <4 x float> %v.vec
+}
+
+; GCN-LABEL: {{^}}load_1d_lwe:
+; GFX10: image_msaa_load v[0:4], v{{[0-9]+}}, s[0:7] dmask:0xf dim:SQ_RSRC_IMG_1D unorm lwe ;
+define amdgpu_ps <4 x float> @load_1d_lwe(<8 x i32> inreg %rsrc, i32 addrspace(1)* inreg %out, i32 %s) {
+main_body:
+  %v = call {<4 x float>, i32} @llvm.amdgcn.image.msaa.load.1d.v4f32i32.i32(i32 15, i32 %s, <8 x i32> %rsrc, i32 2, i32 0)
+  %v.vec = extractvalue {<4 x float>, i32} %v, 0
+  %v.err = extractvalue {<4 x float>, i32} %v, 1
+  store i32 %v.err, i32 addrspace(1)* %out, align 4
+  ret <4 x float> %v.vec
+}
+
+; GCN-LABEL: {{^}}load_2d:
+; GFX10: image_msaa_load v[0:3], v[0:1], s[0:7] dmask:0xf dim:SQ_RSRC_IMG_2D unorm ;
+define amdgpu_ps <4 x float> @load_2d(<8 x i32> inreg %rsrc, i32 %s, i32 %t) {
+main_body:
+  %v = call <4 x float> @llvm.amdgcn.image.msaa.load.2d.v4f32.i32(i32 15, i32 %s, i32 %t, <8 x i32> %rsrc, i32 0, i32 0)
+  ret <4 x float> %v
+}
+
+; GCN-LABEL: {{^}}load_2d_tfe:
+; GFX10: image_msaa_load v[0:4], v[{{[0-9]+:[0-9]+}}], s[0:7] dmask:0xf dim:SQ_RSRC_IMG_2D unorm tfe ;
+define amdgpu_ps <4 x float> @load_2d_tfe(<8 x i32> inreg %rsrc, i32 addrspace(1)* inreg %out, i32 %s, i32 %t) {
+main_body:
+  %v = call {<4 x float>,i32} @llvm.amdgcn.image.msaa.load.2d.v4f32i32.i32(i32 15, i32 %s, i32 %t, <8 x i32> %rsrc, i32 1, i32 0)
+  %v.vec = extractvalue {<4 x float>, i32} %v, 0
+  %v.err = extractvalue {<4 x float>, i32} %v, 1
+  store i32 %v.err, i32 addrspace(1)* %out, align 4
+  ret <4 x float> %v.vec
+}
+
+; GCN-LABEL: {{^}}load_3d:
+; GFX10: image_msaa_load v[0:3], v[0:2], s[0:7] dmask:0xf dim:SQ_RSRC_IMG_3D unorm ;
+define amdgpu_ps <4 x float> @load_3d(<8 x i32> inreg %rsrc, i32 %s, i32 %t, i32 %r) {
+main_body:
+  %v = call <4 x float> @llvm.amdgcn.image.msaa.load.3d.v4f32.i32(i32 15, i32 %s, i32 %t, i32 %r, <8 x i32> %rsrc, i32 0, i32 0)
+  ret <4 x float> %v
+}
+
+; GCN-LABEL: {{^}}load_3d_tfe_lwe:
+; GFX10: image_msaa_load v[0:4], v[{{[0-9]+:[0-9]+}}], s[0:7] dmask:0xf dim:SQ_RSRC_IMG_3D unorm tfe lwe ;
+define amdgpu_ps <4 x float> @load_3d_tfe_lwe(<8 x i32> inreg %rsrc, i32 addrspace(1)* inreg %out, i32 %s, i32 %t, i32 %r) {
+main_body:
+  %v = call {<4 x float>,i32} @llvm.amdgcn.image.msaa.load.3d.v4f32i32.i32(i32 15, i32 %s, i32 %t, i32 %r, <8 x i32> %rsrc, i32 3, i32 0)
+  %v.vec = extractvalue {<4 x float>, i32} %v, 0
+  %v.err = extractvalue {<4 x float>, i32} %v, 1
+  store i32 %v.err, i32 addrspace(1)* %out, align 4
+  ret <4 x float> %v.vec
+}
+
+; GCN-LABEL: {{^}}load_1darray:
+; GFX10: image_msaa_load v[0:3], v[0:1], s[0:7] dmask:0xf dim:SQ_RSRC_IMG_1D_ARRAY unorm ;
+define amdgpu_ps <4 x float> @load_1darray(<8 x i32> inreg %rsrc, i32 %s, i32 %slice) {
+main_body:
+  %v = call <4 x float> @llvm.amdgcn.image.msaa.load.1darray.v4f32.i32(i32 15, i32 %s, i32 %slice, <8 x i32> %rsrc, i32 0, i32 0)
+  ret <4 x float> %v
+}
+
+; GCN-LABEL: {{^}}load_1darray_tfe:
+; GFX10: image_msaa_load v[0:4], v[{{[0-9]+:[0-9]+}}], s[0:7] dmask:0xf dim:SQ_RSRC_IMG_1D_ARRAY unorm tfe ;
+define amdgpu_ps <4 x float> @load_1darray_tfe(<8 x i32> inreg %rsrc, i32 addrspace(1)* inreg %out, i32 %s, i32 %slice) {
+main_body:
+  %v = call {<4 x float>,i32} @llvm.amdgcn.image.msaa.load.1darray.v4f32i32.i32(i32 15, i32 %s, i32 %slice, <8 x i32> %rsrc, i32 1, i32 0)
+  %v.vec = extractvalue {<4 x float>, i32} %v, 0
+  %v.err = extractvalue {<4 x float>, i32} %v, 1
+  store i32 %v.err, i32 addrspace(1)* %out, align 4
+  ret <4 x float> %v.vec
+}
+
+; GCN-LABEL: {{^}}load_2darray:
+; GFX10: image_msaa_load v[0:3], v[0:2], s[0:7] dmask:0xf dim:SQ_RSRC_IMG_2D_ARRAY unorm ;
+define amdgpu_ps <4 x float> @load_2darray(<8 x i32> inreg %rsrc, i32 %s, i32 %t, i32 %slice) {
+main_body:
+  %v = call <4 x float> @llvm.amdgcn.image.msaa.load.2darray.v4f32.i32(i32 15, i32 %s, i32 %t, i32 %slice, <8 x i32> %rsrc, i32 0, i32 0)
+  ret <4 x float> %v
+}
+
+; GCN-LABEL: {{^}}load_2darray_lwe:
+; GFX10: image_msaa_load v[0:4], v[{{[0-9]+:[0-9]+}}], s[0:7] dmask:0xf dim:SQ_RSRC_IMG_2D_ARRAY unorm lwe ;
+define amdgpu_ps <4 x float> @load_2darray_lwe(<8 x i32> inreg %rsrc, i32 addrspace(1)* inreg %out, i32 %s, i32 %t, i32 %slice) {
+main_body:
+  %v = call {<4 x float>,i32} @llvm.amdgcn.image.msaa.load.2darray.v4f32i32.i32(i32 15, i32 %s, i32 %t, i32 %slice, <8 x i32> %rsrc, i32 2, i32 0)
+  %v.vec = extractvalue {<4 x float>, i32} %v, 0
+  %v.err = extractvalue {<4 x float>, i32} %v, 1
+  store i32 %v.err, i32 addrspace(1)* %out, align 4
+  ret <4 x float> %v.vec
+}
+
+; GCN-LABEL: {{^}}load_2dmsaa:
+; GFX10: image_msaa_load v[0:3], v[0:2], s[0:7] dmask:0xf dim:SQ_RSRC_IMG_2D_MSAA unorm ;
+define amdgpu_ps <4 x float> @load_2dmsaa(<8 x i32> inreg %rsrc, i32 %s, i32 %t, i32 %fragid) {
+main_body:
+  %v = call <4 x float> @llvm.amdgcn.image.msaa.load.2dmsaa.v4f32.i32(i32 15, i32 %s, i32 %t, i32 %fragid, <8 x i32> %rsrc, i32 0, i32 0)
+  ret <4 x float> %v
+}
+
+; GCN-LABEL: {{^}}load_2dmsaa_both:
+; GFX10: image_msaa_load v[0:4], v[{{[0-9]+:[0-9]+}}], s[0:7] dmask:0xf dim:SQ_RSRC_IMG_2D_MSAA unorm tfe lwe ;
+define amdgpu_ps <4 x float> @load_2dmsaa_both(<8 x i32> inreg %rsrc, i32 addrspace(1)* inreg %out, i32 %s, i32 %t, i32 %fragid) {
+main_body:
+  %v = call {<4 x float>,i32} @llvm.amdgcn.image.msaa.load.2dmsaa.v4f32i32.i32(i32 15, i32 %s, i32 %t, i32 %fragid, <8 x i32> %rsrc, i32 3, i32 0)
+  %v.vec = extractvalue {<4 x float>, i32} %v, 0
+  %v.err = extractvalue {<4 x float>, i32} %v, 1
+  store i32 %v.err, i32 addrspace(1)* %out, align 4
+  ret <4 x float> %v.vec
+}
+
+; GCN-LABEL: {{^}}load_2darraymsaa:
+; GFX10: image_msaa_load v[0:3], v[0:3], s[0:7] dmask:0xf dim:SQ_RSRC_IMG_2D_MSAA_ARRAY unorm ;
+define amdgpu_ps <4 x float> @load_2darraymsaa(<8 x i32> inreg %rsrc, i32 %s, i32 %t, i32 %slice, i32 %fragid) {
+main_body:
+  %v = call <4 x float> @llvm.amdgcn.image.msaa.load.2darraymsaa.v4f32.i32(i32 15, i32 %s, i32 %t, i32 %slice, i32 %fragid, <8 x i32> %rsrc, i32 0, i32 0)
+  ret <4 x float> %v
+}
+
+; GCN-LABEL: {{^}}load_2darraymsaa_tfe:
+; GFX10: image_msaa_load v[0:4], v[{{[0-9]+:[0-9]+}}], s[0:7] dmask:0xf dim:SQ_RSRC_IMG_2D_MSAA_ARRAY unorm tfe ;
+define amdgpu_ps <4 x float> @load_2darraymsaa_tfe(<8 x i32> inreg %rsrc, i32 addrspace(1)* inreg %out, i32 %s, i32 %t, i32 %slice, i32 %fragid) {
+main_body:
+  %v = call {<4 x float>,i32} @llvm.amdgcn.image.msaa.load.2darraymsaa.v4f32i32.i32(i32 15, i32 %s, i32 %t, i32 %slice, i32 %fragid, <8 x i32> %rsrc, i32 1, i32 0)
+  %v.vec = extractvalue {<4 x float>, i32} %v, 0
+  %v.err = extractvalue {<4 x float>, i32} %v, 1
+  store i32 %v.err, i32 addrspace(1)* %out, align 4
+  ret <4 x float> %v.vec
+}
+
+; GCN-LABEL: {{^}}load_1d_tfe_V4_dmask3:
+; GFX10: image_msaa_load v[0:3], v{{[0-9]+}}, s[0:7] dmask:0x7 dim:SQ_RSRC_IMG_1D unorm tfe ;
+define amdgpu_ps <4 x float> @load_1d_tfe_V4_dmask3(<8 x i32> inreg %rsrc, i32 addrspace(1)* inreg %out, i32 %s) {
+main_body:
+  %v = call {<4 x float>,i32} @llvm.amdgcn.image.msaa.load.1d.v4f32i32.i32(i32 7, i32 %s, <8 x i32> %rsrc, i32 1, i32 0)
+  %v.vec = extractvalue {<4 x float>, i32} %v, 0
+  %v.err = extractvalue {<4 x float>, i32} %v, 1
+  store i32 %v.err, i32 addrspace(1)* %out, align 4
+  ret <4 x float> %v.vec
+}
+
+; GCN-LABEL: {{^}}load_1d_tfe_V4_dmask2:
+; GFX10: image_msaa_load v[0:2], v{{[0-9]+}}, s[0:7] dmask:0x6 dim:SQ_RSRC_IMG_1D unorm tfe ;
+define amdgpu_ps <4 x float> @load_1d_tfe_V4_dmask2(<8 x i32> inreg %rsrc, i32 addrspace(1)* inreg %out, i32 %s) {
+main_body:
+  %v = call {<4 x float>,i32} @llvm.amdgcn.image.msaa.load.1d.v4f32i32.i32(i32 6, i32 %s, <8 x i32> %rsrc, i32 1, i32 0)
+  %v.vec = extractvalue {<4 x float>, i32} %v, 0
+  %v.err = extractvalue {<4 x float>, i32} %v, 1
+  store i32 %v.err, i32 addrspace(1)* %out, align 4
+  ret <4 x float> %v.vec
+}
+
+; GCN-LABEL: {{^}}load_1d_tfe_V4_dmask1:
+; GFX10: image_msaa_load v[0:1], v{{[0-9]+}}, s[0:7] dmask:0x8 dim:SQ_RSRC_IMG_1D unorm tfe ;
+define amdgpu_ps <4 x float> @load_1d_tfe_V4_dmask1(<8 x i32> inreg %rsrc, i32 addrspace(1)* inreg %out, i32 %s) {
+main_body:
+  %v = call {<4 x float>,i32} @llvm.amdgcn.image.msaa.load.1d.v4f32i32.i32(i32 8, i32 %s, <8 x i32> %rsrc, i32 1, i32 0)
+  %v.vec = extractvalue {<4 x float>, i32} %v, 0
+  %v.err = extractvalue {<4 x float>, i32} %v, 1
+  store i32 %v.err, i32 addrspace(1)* %out, align 4
+  ret <4 x float> %v.vec
+}
+
+; GCN-LABEL: {{^}}load_1d_tfe_V2_dmask1:
+; GFX10: image_msaa_load v[0:1], v{{[0-9]+}}, s[0:7] dmask:0x8 dim:SQ_RSRC_IMG_1D unorm tfe ;
+define amdgpu_ps <2 x float> @load_1d_tfe_V2_dmask1(<8 x i32> inreg %rsrc, i32 addrspace(1)* inreg %out, i32 %s) {
+main_body:
+  %v = call {<2 x float>,i32} @llvm.amdgcn.image.msaa.load.1d.v2f32i32.i32(i32 8, i32 %s, <8 x i32> %rsrc, i32 1, i32 0)
+  %v.vec = extractvalue {<2 x float>, i32} %v, 0
+  %v.err = extractvalue {<2 x float>, i32} %v, 1
+  store i32 %v.err, i32 addrspace(1)* %out, align 4
+  ret <2 x float> %v.vec
+}
+
+; GCN-LABEL: {{^}}load_1d_V1:
+; GFX10: image_msaa_load v0, v0, s[0:7] dmask:0x8 dim:SQ_RSRC_IMG_1D unorm ;
+define amdgpu_ps float @load_1d_V1(<8 x i32> inreg %rsrc, i32 %s) {
+main_body:
+  %v = call float @llvm.amdgcn.image.msaa.load.1d.f32.i32(i32 8, i32 %s, <8 x i32> %rsrc, i32 0, i32 0)
+  ret float %v
+}
+
+; GCN-LABEL: {{^}}load_1d_V2:
+; GFX10: image_msaa_load v[0:1], v0, s[0:7] dmask:0x9 dim:SQ_RSRC_IMG_1D unorm ;
+define amdgpu_ps <2 x float> @load_1d_V2(<8 x i32> inreg %rsrc, i32 %s) {
+main_body:
+  %v = call <2 x float> @llvm.amdgcn.image.msaa.load.1d.v2f32.i32(i32 9, i32 %s, <8 x i32> %rsrc, i32 0, i32 0)
+  ret <2 x float> %v
+}
+
+; GCN-LABEL: {{^}}load_1d_glc:
+; GFX10: image_msaa_load v[0:3], v0, s[0:7] dmask:0xf dim:SQ_RSRC_IMG_1D unorm glc ;
+define amdgpu_ps <4 x float> @load_1d_glc(<8 x i32> inreg %rsrc, i32 %s) {
+main_body:
+  %v = call <4 x float> @llvm.amdgcn.image.msaa.load.1d.v4f32.i32(i32 15, i32 %s, <8 x i32> %rsrc, i32 0, i32 1)
+  ret <4 x float> %v
+}
+
+; GCN-LABEL: {{^}}load_1d_slc:
+; GFX10: image_msaa_load v[0:3], v0, s[0:7] dmask:0xf dim:SQ_RSRC_IMG_1D unorm slc ;
+define amdgpu_ps <4 x float> @load_1d_slc(<8 x i32> inreg %rsrc, i32 %s) {
+main_body:
+  %v = call <4 x float> @llvm.amdgcn.image.msaa.load.1d.v4f32.i32(i32 15, i32 %s, <8 x i32> %rsrc, i32 0, i32 2)
+  ret <4 x float> %v
+}
+
+; GCN-LABEL: {{^}}load_1d_glc_slc:
+; GFX10: image_msaa_load v[0:3], v0, s[0:7] dmask:0xf dim:SQ_RSRC_IMG_1D unorm glc slc ;
+define amdgpu_ps <4 x float> @load_1d_glc_slc(<8 x i32> inreg %rsrc, i32 %s) {
+main_body:
+  %v = call <4 x float> @llvm.amdgcn.image.msaa.load.1d.v4f32.i32(i32 15, i32 %s, <8 x i32> %rsrc, i32 0, i32 3)
+  ret <4 x float> %v
+}
+
+declare <4 x float> @llvm.amdgcn.image.msaa.load.1d.v4f32.i32(i32, i32, <8 x i32>, i32, i32) #1
+declare {float,i32} @llvm.amdgcn.image.msaa.load.1d.f32i32.i32(i32, i32, <8 x i32>, i32, i32) #1
+declare {<2 x float>,i32} @llvm.amdgcn.image.msaa.load.1d.v2f32i32.i32(i32, i32, <8 x i32>, i32, i32) #1
+declare {<4 x float>,i32} @llvm.amdgcn.image.msaa.load.1d.v4f32i32.i32(i32, i32, <8 x i32>, i32, i32) #1
+declare <4 x float> @llvm.amdgcn.image.msaa.load.2d.v4f32.i32(i32, i32, i32, <8 x i32>, i32, i32) #1
+declare {<4 x float>,i32} @llvm.amdgcn.image.msaa.load.2d.v4f32i32.i32(i32, i32, i32, <8 x i32>, i32, i32) #1
+declare <4 x float> @llvm.amdgcn.image.msaa.load.3d.v4f32.i32(i32, i32, i32, i32, <8 x i32>, i32, i32) #1
+declare {<4 x float>,i32} @llvm.amdgcn.image.msaa.load.3d.v4f32i32.i32(i32, i32, i32, i32, <8 x i32>, i32, i32) #1
+declare <4 x float> @llvm.amdgcn.image.msaa.load.1darray.v4f32.i32(i32, i32, i32, <8 x i32>, i32, i32) #1
+declare {<4 x float>,i32} @llvm.amdgcn.image.msaa.load.1darray.v4f32i32.i32(i32, i32, i32, <8 x i32>, i32, i32) #1
+declare <4 x float> @llvm.amdgcn.image.msaa.load.2darray.v4f32.i32(i32, i32, i32, i32, <8 x i32>, i32, i32) #1
+declare {<4 x float>,i32} @llvm.amdgcn.image.msaa.load.2darray.v4f32i32.i32(i32, i32, i32, i32, <8 x i32>, i32, i32) #1
+declare <4 x float> @llvm.amdgcn.image.msaa.load.2dmsaa.v4f32.i32(i32, i32, i32, i32, <8 x i32>, i32, i32) #1
+declare {<4 x float>,i32} @llvm.amdgcn.image.msaa.load.2dmsaa.v4f32i32.i32(i32, i32, i32, i32, <8 x i32>, i32, i32) #1
+declare <4 x float> @llvm.amdgcn.image.msaa.load.2darraymsaa.v4f32.i32(i32, i32, i32, i32, i32, <8 x i32>, i32, i32) #1
+declare {<4 x float>,i32} @llvm.amdgcn.image.msaa.load.2darraymsaa.v4f32i32.i32(i32, i32, i32, i32, i32, <8 x i32>, i32, i32) #1
+
+declare float @llvm.amdgcn.image.msaa.load.1d.f32.i32(i32, i32, <8 x i32>, i32, i32) #1
+declare float @llvm.amdgcn.image.msaa.load.2d.f32.i32(i32, i32, i32, <8 x i32>, i32, i32) #1
+declare <2 x float> @llvm.amdgcn.image.msaa.load.1d.v2f32.i32(i32, i32, <8 x i32>, i32, i32) #1
+
+attributes #0 = { nounwind }
+attributes #1 = { nounwind readonly }

llvm/test/CodeGen/AMDGPU/llvm.amdgcn.s.memtime.ll

 ; RUN: llc -march=amdgcn -mcpu=tahiti -verify-machineinstrs < %s | FileCheck -check-prefix=SI -check-prefix=SIVI -check-prefix=GCN %s
 ; RUN: llc -march=amdgcn -mcpu=tonga -mattr=-flat-for-global -verify-machineinstrs < %s | FileCheck -check-prefix=VI -check-prefix=SIVI -check-prefix=GCN %s
 ; RUN: llc -march=amdgcn -mcpu=gfx1010 -mattr=-flat-for-global -verify-machineinstrs < %s | FileCheck -check-prefix=GCN %s
+; RUN: not --crash llc -march=amdgcn -mcpu=gfx1030 -mattr=-flat-for-global -verify-machineinstrs < %s 2>&1 | FileCheck -check-prefix=GFX1030-ERR %s
 
 declare i64 @llvm.amdgcn.s.memtime() #0
 
 ; GCN-LABEL: {{^}}test_s_memtime:
 ; GCN-DAG: s_memtime s{{\[[0-9]+:[0-9]+\]}}
 ; GCN-DAG: s_load_dwordx2
 ; GCN: lgkmcnt
 ; GCN: {{buffer|global}}_store_dwordx2
 ; SIVI-NOT: lgkmcnt
 ; GCN: s_memtime s{{\[[0-9]+:[0-9]+\]}}
 ; GCN: {{buffer|global}}_store_dwordx2
+; GFX1030-ERR: ERROR
 define amdgpu_kernel void @test_s_memtime(i64 addrspace(1)* %out) #0 {
   %cycle0 = call i64 @llvm.amdgcn.s.memtime()
   store volatile i64 %cycle0, i64 addrspace(1)* %out
 
   %cycle1 = call i64 @llvm.amdgcn.s.memtime()
   store volatile i64 %cycle1, i64 addrspace(1)* %out
   ret void
 }
 
 attributes #0 = { nounwind }

llvm/test/CodeGen/AMDGPU/llvm.amdgcn.sdot4.ll

 ; RUN: llc -march=amdgcn -mcpu=gfx906 -verify-machineinstrs < %s | FileCheck %s --check-prefixes=GCN,GFX906
 ; RUN: llc -march=amdgcn -mcpu=gfx1011 -verify-machineinstrs < %s | FileCheck %s --check-prefixes=GCN,GFX10
 ; RUN: llc -march=amdgcn -mcpu=gfx1012 -verify-machineinstrs < %s | FileCheck %s --check-prefixes=GCN,GFX10
+; RUN: llc -march=amdgcn -mcpu=gfx1030 -verify-machineinstrs < %s | FileCheck %s --check-prefixes=GCN,GFX10
 
 declare i32 @llvm.amdgcn.sdot4(i32 %a, i32 %b, i32 %c, i1 %clamp)
 
 ; GCN-LABEL: {{^}}test_llvm_amdgcn_sdot4_clamp
 ; GFX906: v_dot4_i32_i8 v{{[0-9]+}}, s{{[0-9]+}}, v{{[0-9]+}}, v{{[0-9]+}} clamp{{$}}
 ; GFX10:  v_dot4_i32_i8 v{{[0-9]+}}, s{{[0-9]+}}, s{{[0-9]+}}, v{{[0-9]+}} clamp{{$}}
 define amdgpu_kernel void @test_llvm_amdgcn_sdot4_clamp(
     i32 addrspace(1)* %r,

llvm/test/CodeGen/AMDGPU/llvm.amdgcn.sdot8.ll

 ; RUN: llc -march=amdgcn -mcpu=gfx906 -verify-machineinstrs < %s | FileCheck %s --check-prefixes=GCN,GFX906
 ; RUN: llc -march=amdgcn -mcpu=gfx908 -verify-machineinstrs < %s | FileCheck %s --check-prefixes=GCN,GFX908
-; RUN: llc -march=amdgcn -mcpu=gfx1011 -verify-machineinstrs < %s | FileCheck %s --check-prefixes=GCN,GFX10,GFX1011
-; RUN: llc -march=amdgcn -mcpu=gfx1012 -verify-machineinstrs < %s | FileCheck %s --check-prefixes=GCN,GFX10,GFX1011
+; RUN: llc -march=amdgcn -mcpu=gfx1011 -verify-machineinstrs < %s | FileCheck %s --check-prefixes=GCN,GFX10
+; RUN: llc -march=amdgcn -mcpu=gfx1012 -verify-machineinstrs < %s | FileCheck %s --check-prefixes=GCN,GFX10
+; RUN: llc -march=amdgcn -mcpu=gfx1030 -verify-machineinstrs < %s | FileCheck %s --check-prefixes=GCN,GFX10
 
 declare i32 @llvm.amdgcn.sdot8(i32 %a, i32 %b, i32 %c, i1 %clamp)
 
 ; GCN-LABEL: {{^}}test_llvm_amdgcn_sdot8_clamp
-; GFX906: v_dot8_i32_i4 v{{[0-9]+}}, s{{[0-9]+}}, v{{[0-9]+}}, v{{[0-9]+}} clamp{{$}}
-; GFX908: v_dot8_i32_i4 v{{[0-9]+}}, s{{[0-9]+}}, v{{[0-9]+}}, v{{[0-9]+}} clamp{{$}}
-; GFX10:  v_dot8_i32_i4 v{{[0-9]+}}, s{{[0-9]+}}, s{{[0-9]+}}, v{{[0-9]+}} clamp{{$}}
+; GFX906:  v_dot8_i32_i4 v{{[0-9]+}}, s{{[0-9]+}}, v{{[0-9]+}}, v{{[0-9]+}} clamp{{$}}
+; GFX908:  v_dot8_i32_i4 v{{[0-9]+}}, s{{[0-9]+}}, v{{[0-9]+}}, v{{[0-9]+}} clamp{{$}}
+; GFX10:   v_dot8_i32_i4 v{{[0-9]+}}, s{{[0-9]+}}, s{{[0-9]+}}, v{{[0-9]+}} clamp{{$}}
 define amdgpu_kernel void @test_llvm_amdgcn_sdot8_clamp(
     i32 addrspace(1)* %r,
     <8 x i4> addrspace(1)* %a,
     <8 x i4> addrspace(1)* %b,
     i32 addrspace(1)* %c) {
 entry:
   %a.val = load <8 x i4>, <8 x i4> addrspace(1)* %a
   %b.val = load <8 x i4>, <8 x i4> addrspace(1)* %b
   %a.val.cast = bitcast <8 x i4> %a.val to i32
   %b.val.cast = bitcast <8 x i4> %b.val to i32
   %c.val = load i32, i32 addrspace(1)* %c
   %r.val = call i32 @llvm.amdgcn.sdot8(i32 %a.val.cast, i32 %b.val.cast, i32 %c.val, i1 1)
   store i32 %r.val, i32 addrspace(1)* %r
   ret void
 }
 
 ; GCN-LABEL: {{^}}test_llvm_amdgcn_sdot8_no_clamp
-; GFX906: v_dot8_i32_i4 v{{[0-9]+}}, s{{[0-9]+}}, v{{[0-9]+}}, v{{[0-9]+}}{{$}}
-; GFX908: v_dot8c_i32_i4_e32 v{{[0-9]+}}, s{{[0-9]+}}, v{{[0-9]+}}{{$}}
-; GFX1011:  v_dot8_i32_i4 v{{[0-9]+}}, s{{[0-9]+}}, s{{[0-9]+}}, v{{[0-9]+}}{{$}}
+; GFX906:  v_dot8_i32_i4 v{{[0-9]+}}, s{{[0-9]+}}, v{{[0-9]+}}, v{{[0-9]+}}{{$}}
+; GFX908:  v_dot8c_i32_i4_e32 v{{[0-9]+}}, s{{[0-9]+}}, v{{[0-9]+}}{{$}}
+; GFX10:   v_dot8_i32_i4 v{{[0-9]+}}, s{{[0-9]+}}, s{{[0-9]+}}, v{{[0-9]+}}{{$}}
 define amdgpu_kernel void @test_llvm_amdgcn_sdot8_no_clamp(
     i32 addrspace(1)* %r,
     <8 x i4> addrspace(1)* %a,
     <8 x i4> addrspace(1)* %b,
     i32 addrspace(1)* %c) {
 entry:
   %a.val = load <8 x i4>, <8 x i4> addrspace(1)* %a
   %b.val = load <8 x i4>, <8 x i4> addrspace(1)* %b
   %a.val.cast = bitcast <8 x i4> %a.val to i32
   %b.val.cast = bitcast <8 x i4> %b.val to i32
   %c.val = load i32, i32 addrspace(1)* %c
   %r.val = call i32 @llvm.amdgcn.sdot8(i32 %a.val.cast, i32 %b.val.cast, i32 %c.val, i1 0)
   store i32 %r.val, i32 addrspace(1)* %r
   ret void
 }

llvm/test/CodeGen/AMDGPU/loop-prefetch.ll

+; RUN: llc -march=amdgcn -mcpu=gfx1030 -verify-machineinstrs -asm-verbose=0 < %s | FileCheck --check-prefixes=GCN,GFX10,GFX10-ASM %s
+; RUN: llc -march=amdgcn -mcpu=gfx1030 -verify-machineinstrs < %s -filetype=obj | llvm-objdump -d --arch-name=amdgcn --mcpu=gfx1030 - | FileCheck --check-prefixes=GCN,GFX10,GFX10-DIS %s
+; RUN: llc -march=amdgcn -mcpu=tonga -verify-machineinstrs < %s | FileCheck --check-prefix=GFX8 %s
+
+; GFX8-NOT: s_inst_prefetch
+; GFX8-NOT: .palign 6
+
+; GCN-LABEL: test_loop_64
+; GFX10:          s_movk_i32 s{{[0-9]+}}, 0x400
+; GFX10-DIS-NEXT: {{^$}}
+; GFX10-ASM-NEXT: [[L1:BB[0-9_]+]]:
+; GFX10-DIS-NEXT: <[[L1:BB[0-9_]+]]>:
+; GFX10:          s_sleep 0
+; GFX10:          s_cbranch_scc0 [[L1]]
+; GFX10-NEXT:     s_endpgm
+define amdgpu_kernel void @test_loop_64(i32 addrspace(1)* nocapture %arg) {
+bb:
+  br label %bb2
+
+bb1:                                              ; preds = %bb2
+  ret void
+
+bb2:                                              ; preds = %bb2, %bb
+  %tmp1 = phi i32 [ 0, %bb ], [ %tmp2, %bb2 ]
+  %tmp2 = add nuw nsw i32 %tmp1, 1
+  %tmp3 = icmp eq i32 %tmp2, 1024
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  br i1 %tmp3, label %bb1, label %bb2
+}
+
+; GCN-LABEL: test_loop_128
+; GFX10:          s_movk_i32 s{{[0-9]+}}, 0x400
+; GFX10-ASM-NEXT: .p2align 6
+; GFX10-DIS-NEXT: s_nop 0
+; GFX10-NOT:      s_inst_prefetch
+; GFX10-ASM:      [[L1:BB[0-9_]+]]:
+; GFX10-DIS:      <[[L1:BB[0-9_]+]]>:
+; GFX10:          s_sleep 0
+; GFX10:          s_cbranch_scc0 [[L1]]
+; GFX10-NEXT:     s_endpgm
+define amdgpu_kernel void @test_loop_128(i32 addrspace(1)* nocapture %arg) {
+bb:
+  br label %bb2
+
+bb1:                                              ; preds = %bb2
+  ret void
+
+bb2:                                              ; preds = %bb2, %bb
+  %tmp1 = phi i32 [ 0, %bb ], [ %tmp2, %bb2 ]
+  %tmp2 = add nuw nsw i32 %tmp1, 1
+  %tmp3 = icmp eq i32 %tmp2, 1024
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  br i1 %tmp3, label %bb1, label %bb2
+}
+
+; GCN-LABEL: test_loop_192
+; GFX10:          s_movk_i32 s{{[0-9]+}}, 0x400
+; GFX10-NEXT:     s_inst_prefetch 0x1
+; GFX10-ASM-NEXT: .p2align 6
+; GFX10-DIS-NEXT: s_nop 0
+; GFX10-NOT:      s_inst_prefetch
+; GFX10-ASM:      [[L1:BB[0-9_]+]]:
+; GFX10-DIS:      <[[L1:BB[0-9_]+]]>:
+; GFX10:          s_sleep 0
+; GFX10:          s_cbranch_scc0 [[L1]]
+; GFX10-NEXT:     s_inst_prefetch 0x2
+; GFX10-NEXT:     s_endpgm
+define amdgpu_kernel void @test_loop_192(i32 addrspace(1)* nocapture %arg) {
+bb:
+  br label %bb2
+
+bb1:                                              ; preds = %bb2
+  ret void
+
+bb2:                                              ; preds = %bb2, %bb
+  %tmp1 = phi i32 [ 0, %bb ], [ %tmp2, %bb2 ]
+  %tmp2 = add nuw nsw i32 %tmp1, 1
+  %tmp3 = icmp eq i32 %tmp2, 1024
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  br i1 %tmp3, label %bb1, label %bb2
+}
+
+; GCN-LABEL: test_loop_256
+; GFX10:          s_movk_i32 s{{[0-9]+}}, 0x400
+; GFX10-DIS-NEXT: {{^$}}
+; GFX10-ASM-NEXT: [[L1:BB[0-9_]+]]:
+; GFX10-DIS-NEXT: <[[L1:BB[0-9_]+]]>:
+; GFX10:          s_sleep 0
+; GFX10:          s_cbranch_scc0 [[L1]]
+; GFX10-NEXT:     s_endpgm
+define amdgpu_kernel void @test_loop_256(i32 addrspace(1)* nocapture %arg) {
+bb:
+  br label %bb2
+
+bb1:                                              ; preds = %bb2
+  ret void
+
+bb2:                                              ; preds = %bb2, %bb
+  %tmp1 = phi i32 [ 0, %bb ], [ %tmp2, %bb2 ]
+  %tmp2 = add nuw nsw i32 %tmp1, 1
+  %tmp3 = icmp eq i32 %tmp2, 1024
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  br i1 %tmp3, label %bb1, label %bb2
+}
+
+; GCN-LABEL: test_loop_prefetch_inner_outer
+; GFX10:          s_inst_prefetch 0x1
+; GFX10-ASM-NEXT: .p2align 6
+; GFX10-DIS-NEXT: s_nop 0
+; GFX10-NOT:      s_inst_prefetch
+; GFX10-ASM:      [[L1:BB[0-9_]+]]:
+; GFX10-DIS:      <[[L1:BB[0-9_]+]]>:
+; GFX10-NOT:      s_inst_prefetch
+; GFX10-ASM:      .p2align 6
+; GFX10-DIS:      s_nop 0
+; GFX10-NOT:      s_inst_prefetch
+; GFX10-ASM:      [[L2:BB[0-9_]+]]:
+; GFX10-DIS:      <[[L2:BB[0-9_]+]]>:
+; GFX10-NOT:      s_inst_prefetch
+; GFX10:          s_sleep 0
+; GFX10:          s_cbranch_scc{{[01]}} [[L2]]
+; GFX10-NOT:      s_inst_prefetch
+; GFX10:          s_cbranch_scc{{[01]}} [[L1]]
+; GFX10-NEXT:     s_inst_prefetch 0x2
+; GFX10-NEXT:     s_endpgm
+define amdgpu_kernel void @test_loop_prefetch_inner_outer(i32 addrspace(1)* nocapture %arg) {
+bb:
+  br label %bb2
+
+bb1:
+  ret void
+
+bb2:
+  %tmp1 = phi i32 [ 0, %bb ], [ %tmp2, %bb4 ]
+  %tmp2 = add nuw nsw i32 %tmp1, 1
+  %tmp3 = icmp eq i32 %tmp2, 1024
+  br label %bb3
+
+bb3:
+  %tmp4 = phi i32 [ 0, %bb2 ], [ %tmp5, %bb3 ]
+  %tmp5 = add nuw nsw i32 %tmp4, 1
+  %tmp6 = icmp eq i32 %tmp5, 1024
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  br i1 %tmp6, label %bb4, label %bb3
+
+bb4:
+  br i1 %tmp3, label %bb1, label %bb2
+}
+
+; GCN-LABEL: test_loop_prefetch_inner_outer_noouter
+; GFX10-NOT:      .p2align 6
+; GFX10-NOT:      s_nop
+; GFX10-NOT:      s_inst_prefetch
+; GFX10-ASM:      [[L0:BB[0-9_]+]]:
+; GFX10-DIS:      <[[L0:BB[0-9_]+]]>:
+; GFX10:          s_inst_prefetch 0x1
+; GFX10-ASM-NEXT: .p2align 6
+; GFX10-DIS-NEXT: s_nop 0
+; GFX10-NOT:      s_inst_prefetch
+; GFX10-ASM:      [[L1:BB[0-9_]+]]:
+; GFX10-DIS:      <[[L1:BB[0-9_]+]]>:
+; GFX10-NOT:      s_inst_prefetch
+; GFX10-ASM:      .p2align 6
+; GFX10-DIS:      s_nop 0
+; GFX10-NOT:      s_inst_prefetch
+; GFX10-ASM:      [[L2:BB[0-9_]+]]:
+; GFX10-DIS:      <[[L2:BB[0-9_]+]]>:
+; GFX10-NOT:      s_inst_prefetch
+; GFX10:          s_sleep 0
+; GFX10:          s_cbranch_scc{{[01]}} [[L2]]
+; GFX10-NOT:      s_inst_prefetch
+; GFX10:          s_cbranch_scc{{[01]}} [[L1]]
+; GFX10-NEXT:     s_inst_prefetch 0x2
+; GFX10:          s_cbranch_scc{{[01]}} [[L0]]
+; GFX10-NEXT:     s_endpgm
+define amdgpu_kernel void @test_loop_prefetch_inner_outer_noouter(i32 addrspace(1)* nocapture %arg) {
+bb:
+  br label %bb2
+
+bb1:
+  ret void
+
+bb2:
+  %tmp1 = phi i32 [ 0, %bb ], [ %tmp2, %bb6 ]
+  %tmp2 = add nuw nsw i32 %tmp1, 1
+  %tmp3 = icmp eq i32 %tmp2, 1024
+  br label %bb3
+
+bb3:
+  %tmp4 = phi i32 [ 0, %bb2 ], [ %tmp5, %bb5 ]
+  %tmp5 = add nuw nsw i32 %tmp4, 1
+  %tmp6 = icmp eq i32 %tmp5, 1024
+  br label %bb4
+
+bb4:
+  %tmp7 = phi i32 [ 0, %bb3 ], [ %tmp8, %bb4 ]
+  %tmp8 = add nuw nsw i32 %tmp7, 1
+  %tmp9 = icmp eq i32 %tmp8, 1024
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  br i1 %tmp9, label %bb5, label %bb4
+
+bb5:
+  br i1 %tmp6, label %bb6, label %bb3
+
+bb6:
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  tail call void @llvm.amdgcn.s.sleep(i32 0)
+  br i1 %tmp3, label %bb1, label %bb2
+}
+
+declare void @llvm.amdgcn.s.sleep(i32)

llvm/test/CodeGen/AMDGPU/madmk.ll

-; RUN: llc -march=amdgcn -verify-machineinstrs < %s | FileCheck -check-prefix=GCN -check-prefix=SI %s
+; RUN: llc -march=amdgcn -mattr=+mad-mac-f32-insts -verify-machineinstrs < %s | FileCheck -check-prefix=GCN -check-prefix=SI %s
 ; XUN: llc -march=amdgcn -mcpu=tonga -mattr=-flat-for-global -verify-machineinstrs < %s | FileCheck -check-prefix=GCN -check-prefix=VI %s
 
  ; FIXME: None of these trigger madmk emission anymore. It is still
  ; possible, but requires the correct registers to be used which is
  ; hard to trigger.
 
 declare i32 @llvm.amdgcn.workitem.id.x() nounwind readnone
 declare float @llvm.fabs.f32(float) nounwind readnone

llvm/test/CodeGen/AMDGPU/operand-folding.ll

-; RUN: llc -march=amdgcn -verify-machineinstrs < %s | FileCheck %s
+; RUN: llc -march=amdgcn -mcpu=tahiti -verify-machineinstrs < %s | FileCheck %s
 
 ; CHECK-LABEL: {{^}}fold_sgpr:
 ; CHECK: v_add_i32_e32 v{{[0-9]+}}, vcc, s
 define amdgpu_kernel void @fold_sgpr(i32 addrspace(1)* %out, i32 %fold) #1 {
 entry:
   %tmp0 = icmp ne i32 %fold, 0
   br i1 %tmp0, label %if, label %endif
 

llvm/test/CodeGen/AMDGPU/readcyclecounter.ll

 ; RUN: llc -march=amdgcn -mcpu=tahiti -verify-machineinstrs < %s | FileCheck -enable-var-scope -check-prefix=MEMTIME -check-prefix=SIVI -check-prefix=GCN %s
 ; RUN: llc -march=amdgcn -mcpu=tonga -verify-machineinstrs < %s | FileCheck -enable-var-scope -check-prefix=MEMTIME -check-prefix=SIVI -check-prefix=GCN %s
 ; RUN: llc -march=amdgcn -mcpu=gfx1010 -verify-machineinstrs < %s | FileCheck -enable-var-scope -check-prefix=MEMTIME -check-prefix=GCN %s
+; RUN: llc -march=amdgcn -mcpu=gfx1030 -verify-machineinstrs < %s | FileCheck -enable-var-scope -check-prefix=GETREG -check-prefix=GCN %s
 
 declare i64 @llvm.readcyclecounter() #0
 
 ; GCN-LABEL: {{^}}test_readcyclecounter:
 ; MEMTIME-DAG: s_memtime s{{\[[0-9]+:[0-9]+\]}}
 ; GCN-DAG:     s_load_dwordx2
 ; GCN-DAG:     lgkmcnt
 ; MEMTIME:     store_dwordx2
 ; SIVI-NOT:    lgkmcnt
 ; MEMTIME:     s_memtime s{{\[[0-9]+:[0-9]+\]}}
 ; MEMTIME:     store_dwordx2
 
+; GETREG-DAG:  v_mov_b32_e32 v[[ZERO:[0-9]+]], 0
+; GETREG-DAG:  s_getreg_b32 [[CNT1:s[0-9]+]], hwreg(HW_REG_SHADER_CYCLES, 0, 20)
+; GETREG-DAG:  v_mov_b32_e32 v[[VCNT1:[0-9]+]], [[CNT1]]
+; GETREG:      global_store_dwordx2 v[{{[0-9:]+}}], v{{\[}}[[VCNT1]]:[[ZERO]]], off
+; GETREG:      s_getreg_b32 [[CNT2:s[0-9]+]], hwreg(HW_REG_SHADER_CYCLES, 0, 20)
+; GETREG:      v_mov_b32_e32 v[[VCNT2:[0-9]+]], [[CNT2]]
+; GETREG:      global_store_dwordx2 v[{{[0-9:]+}}], v{{\[}}[[VCNT2]]:[[ZERO]]], off
+
 define amdgpu_kernel void @test_readcyclecounter(i64 addrspace(1)* %out) #0 {
   %cycle0 = call i64 @llvm.readcyclecounter()
   store volatile i64 %cycle0, i64 addrspace(1)* %out
 
   %cycle1 = call i64 @llvm.readcyclecounter()
   store volatile i64 %cycle1, i64 addrspace(1)* %out
   ret void
 }
 
 ; This test used to crash in ScheduleDAG.
 ;
 ; GCN-LABEL: {{^}}test_readcyclecounter_smem:
 ; MEMTIME-DAG: s_memtime
 ; GCN-DAG:     s_load_dword
+; GETREG-DAG:  s_getreg_b32 s1, hwreg(HW_REG_SHADER_CYCLES, 0, 20)
 define amdgpu_cs i32 @test_readcyclecounter_smem(i64 addrspace(4)* inreg %in) #0 {
   %cycle0 = call i64 @llvm.readcyclecounter()
   %in.v = load i64, i64 addrspace(4)* %in
   %r.64 = add i64 %cycle0, %in.v
   %r.32 = trunc i64 %r.64 to i32
   ret i32 %r.32
 }
 
 attributes #0 = { nounwind }

llvm/test/CodeGen/AMDGPU/udiv.ll

 ; RUN:  llc -amdgpu-scalarize-global-loads=false  -march=amdgcn -mcpu=verde -verify-machineinstrs < %s | FileCheck -check-prefix=SI -check-prefix=FUNC %s
 ; RUN:  llc -amdgpu-scalarize-global-loads=false  -march=amdgcn -mcpu=tonga -mattr=-flat-for-global -verify-machineinstrs -denormal-fp-math-f32=preserve-sign < %s | FileCheck -check-prefix=SI -check-prefix=FUNC -check-prefix=VI %s
 
 ; RUN:  llc -amdgpu-scalarize-global-loads=false  -mtriple=amdgcn--amdhsa -mcpu=fiji -denormal-fp-math-f32=ieee < %s | FileCheck -check-prefix=GCN -check-prefix=VI %s
+; RUN:  llc -amdgpu-scalarize-global-loads=false  -mtriple=amdgcn--amdhsa -mcpu=gfx1030 -denormal-fp-math-f32=ieee < %s | FileCheck -check-prefix=GCN -check-prefix=GFX1030 %s
 
 ; RUN:  llc -amdgpu-scalarize-global-loads=false  -march=r600 -mcpu=redwood < %s | FileCheck -check-prefix=EG -check-prefix=FUNC %s
 
 ; FUNC-LABEL: {{^}}udiv_i32:
 ; EG-NOT: SETGE_INT
 ; EG: CF_END
 
 ; SI: v_rcp_iflag_f32_e32

 define amdgpu_kernel void @test_udiv_3_mulhu(i32 %p) {
    %i = udiv i32 %p, 3
    store volatile i32 %i, i32 addrspace(1)* undef
    ret void
 }
 
 ; GCN-LABEL: {{^}}fdiv_test_denormals
 ; VI: v_mad_f32 v{{[0-9]+}}, -v{{[0-9]+}}, v{{[0-9]+}}, v{{[0-9]+}}
+; GFX1030: v_fmac_f32_e64 v{{[0-9]+}}, -v{{[0-9]+}}, v{{[0-9]+}}
 define amdgpu_kernel void @fdiv_test_denormals(i8 addrspace(1)* nocapture readonly %arg) {
 bb:
   %tmp = load i8, i8 addrspace(1)* null, align 1
   %tmp1 = sext i8 %tmp to i32
   %tmp2 = getelementptr inbounds i8, i8 addrspace(1)* %arg, i64 undef
   %tmp3 = load i8, i8 addrspace(1)* %tmp2, align 1
   %tmp4 = sext i8 %tmp3 to i32
   %tmp5 = sdiv i32 %tmp1, %tmp4
   %tmp6 = trunc i32 %tmp5 to i8
   store i8 %tmp6, i8 addrspace(1)* null, align 1
   ret void
 }

llvm/test/CodeGen/AMDGPU/v_mac.ll

-; RUN:  llc -amdgpu-scalarize-global-loads=false -march=amdgcn -denormal-fp-math-f32=preserve-sign -verify-machineinstrs < %s | FileCheck -check-prefix=SI -check-prefix=GCN %s
+; RUN:  llc -amdgpu-scalarize-global-loads=false -march=amdgcn -mattr=+mad-mac-f32-insts -denormal-fp-math-f32=preserve-sign -verify-machineinstrs < %s | FileCheck -check-prefix=SI -check-prefix=GCN %s
 ; RUN:  llc -amdgpu-scalarize-global-loads=false -march=amdgcn -mcpu=tonga -denormal-fp-math=preserve-sign -denormal-fp-math-f32=preserve-sign -mattr=-flat-for-global -verify-machineinstrs < %s | FileCheck -check-prefix=VI -check-prefix=VI-FLUSH -check-prefix=GCN %s
 ; RUN:  llc -amdgpu-scalarize-global-loads=false -march=amdgcn -mcpu=tonga -denormal-fp-math=ieee -denormal-fp-math-f32=preserve-sign -mattr=-flat-for-global -verify-machineinstrs < %s | FileCheck -check-prefix=VI -check-prefix=VI-DENORM -check-prefix=GCN %s
 
 ; GCN-LABEL: {{^}}mac_vvv:
 ; GCN: buffer_load_dword [[A:v[0-9]+]], off, s[{{[0-9]+:[0-9]+}}], 0{{$}}
 ; GCN: buffer_load_dword [[B:v[0-9]+]], off, s[{{[0-9]+:[0-9]+}}], 0 offset:4
 ; GCN: buffer_load_dword [[C:v[0-9]+]], off, s[{{[0-9]+:[0-9]+}}], 0 offset:8
 ; GCN: v_mac_f32_e32 [[C]], [[A]], [[B]]

llvm/test/MC/AMDGPU/gfx1030_err.s

+// RUN: not llvm-mc -arch=amdgcn -mcpu=gfx1030 -show-encoding %s 2>&1 | FileCheck --check-prefix=GFX10 %s
+
+v_dot8c_i32_i4 v5, v1, v2
+// GFX10: error:
+
+v_dot8c_i32_i4 v5, v1, v2 quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
+// GFX10: error:
+
+v_dot8c_i32_i4 v5, v1, v2 quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0 fi:1
+// GFX10: error:
+
+v_dot8c_i32_i4 v5, v1, v2 dpp8:[7,6,5,4,3,2,1,0]
+// GFX10: error:
+
+v_dot8c_i32_i4 v5, v1, v2 dpp8:[7,6,5,4,3,2,1,0] fi:1
+// GFX10: error:
+
+s_get_waveid_in_workgroup s0
+// GFX10: error:
+
+s_memtime s[0:1]
+// GFX10: error:
+
+s_getreg_b32 s2, hwreg(HW_REG_XNACK_MASK)
+// GFX10: error:
+
+v_mac_f32 v0, v1, v2
+// GFX10: error:
+
+v_mad_f32 v0, v1, v2, v3
+// GFX10: error:
+
+v_madak_f32 v0, v1, v2, 1
+// GFX10: error:
+
+v_madmk_f32 v0, v1, 1, v2
+// GFX10: error:
+
+v_mad_legacy_f32 v0, v1, v2, v3
+// GFX10: error:
+
+v_mac_legacy_f32 v0, v1, v2
+// GFX10: error:
+
+ds_add_src2_u32 v1 offset:65535 gds
+// GFX10: error:
+
+ds_add_src2_u32 v1 offset:65535
+// GFX10: error:
+
+ds_add_src2_f32 v1 offset:65535
+// GFX10: error:
+
+ds_sub_src2_u32 v1 offset:65535
+// GFX10: error:
+
+ds_rsub_src2_u32 v1 offset:65535
+// GFX10: error:
+
+ds_inc_src2_u32 v1 offset:65535
+// GFX10: error:
+
+ds_dec_src2_u32 v1 offset:65535
+// GFX10: error:
+
+ds_min_src2_i32 v1 offset:65535
+// GFX10: error:
+
+ds_max_src2_i32 v1 offset:65535
+// GFX10: error:
+
+ds_min_src2_u32 v1 offset:65535
+// GFX10: error:
+
+ds_max_src2_u32 v1 offset:65535
+// GFX10: error:
+
+ds_and_src2_b32 v1 offset:65535
+// GFX10: error:
+
+ds_or_src2_b32 v1 offset:65535
+// GFX10: error:
+
+ds_xor_src2_b32 v1 offset:65535
+// GFX10: error:
+
+ds_min_src2_f32 v1 offset:65535
+// GFX10: error:
+
+ds_max_src2_f32 v1 offset:65535
+// GFX10: error:
+
+ds_add_src2_u64 v1 offset:65535
+// GFX10: error:
+
+ds_sub_src2_u64 v1 offset:65535
+// GFX10: error:
+
+ds_rsub_src2_u64 v1 offset:65535
+// GFX10: error:
+
+ds_inc_src2_u64 v1 offset:65535
+// GFX10: error:
+
+ds_dec_src2_u64 v1 offset:65535
+// GFX10: error:
+
+ds_min_src2_i64 v1 offset:65535
+// GFX10: error:
+
+ds_max_src2_i64 v1 offset:65535
+// GFX10: error:
+
+ds_min_src2_u64 v1 offset:65535
+// GFX10: error:
+
+ds_max_src2_u64 v1 offset:65535
+// GFX10: error:
+
+ds_and_src2_b64 v1 offset:65535
+// GFX10: error:
+
+ds_or_src2_b64 v1 offset:65535
+// GFX10: error:
+
+ds_xor_src2_b64 v1 offset:65535
+// GFX10: error:
+
+ds_min_src2_f64 v1 offset:65535
+// GFX10: error:
+
+ds_max_src2_f64 v1 offset:65535
+// GFX10: error:
+
+ds_write_src2_b32 v1 offset:65535
+// GFX10: error:
+
+ds_write_src2_b64 v1 offset:65535
+// GFX10: error:

llvm/test/MC/AMDGPU/gfx1030_new.s

+// RUN: llvm-mc -arch=amdgcn -mcpu=gfx1030 -show-encoding %s | FileCheck --check-prefix=GFX10 %s
+
+global_load_dword_addtid v1, s[2:3] offset:16
+// GFX10: encoding: [0x10,0x80,0x58,0xdc,0x00,0x00,0x02,0x01]
+
+global_load_dword_addtid v1, s[2:3] offset:16 glc slc dlc
+// GFX10: encoding: [0x10,0x90,0x5b,0xdc,0x00,0x00,0x02,0x01]
+
+global_store_dword_addtid v1, s[2:3] offset:16 glc slc dlc
+// GFX10: encoding: [0x10,0x90,0x5f,0xdc,0x00,0x01,0x02,0x00]
+
+global_store_dword v[254:255], v1, s[2:3] offset:16
+// GFX10: encoding: [0x10,0x80,0x70,0xdc,0xfe,0x01,0x02,0x00]
+
+global_atomic_csub v2, v[0:1], v2, off offset:100 glc slc
+// GFX10: encoding: [0x64,0x80,0xd3,0xdc,0x00,0x02,0x7d,0x02]
+
+global_atomic_csub v2, v[0:1], v2, off
+// GFX10: encoding: [0x00,0x80,0xd1,0xdc,0x00,0x02,0x7d,0x02]
+
+global_atomic_csub v2, v[0:1], v2, s[2:3]
+// GFX10: encoding: [0x00,0x80,0xd1,0xdc,0x00,0x02,0x02,0x02]
+
+global_atomic_csub v2, v[0:1], v2, s[2:3] offset:100 glc slc
+// GFX10: encoding: [0x64,0x80,0xd3,0xdc,0x00,0x02,0x02,0x02]
+
+buffer_atomic_csub v5, off, s[8:11], s3
+// GFX10: encoding: [0x00,0x40,0xd0,0xe0,0x00,0x05,0x02,0x03]
+
+buffer_atomic_csub v5, off, s[8:11], s3 offset:4095 glc
+// GFX10: encoding: [0xff,0x4f,0xd0,0xe0,0x00,0x05,0x02,0x03]
+
+buffer_atomic_csub v5, off, s[8:11], -1 offset:4095 glc
+// GFX10: encoding: [0xff,0x4f,0xd0,0xe0,0x00,0x05,0x02,0xc1]
+
+buffer_atomic_csub v5, v0, s[8:11], s3 offen offset:4095 glc
+// GFX10: encoding: [0xff,0x5f,0xd0,0xe0,0x00,0x05,0x02,0x03]
+
+buffer_atomic_csub v5, v0, s[8:11], s3 idxen offset:4095 glc
+// GFX10: encoding: [0xff,0x6f,0xd0,0xe0,0x00,0x05,0x02,0x03]
+
+buffer_atomic_csub v5, off, s[8:11], s3 glc slc
+// GFX10: encoding: [0x00,0x40,0xd0,0xe0,0x00,0x05,0x42,0x03]
+
+s_getreg_b32 s2, hwreg(HW_REG_SHADER_CYCLES)
+// GFX10: encoding: [0x1d,0xf8,0x02,0xb9]
+
+s_getreg_b32 s2, 29
+// GFX10: s_getreg_b32 s2, hwreg(HW_REG_SHADER_CYCLES, 0, 1) ; encoding: [0x1d,0x00,0x02,0xb9]
+
+s_getreg_b32 s2, hwreg(22)
+// GFX10: s_getreg_b32 s2, hwreg(22) ; encoding: [0x16,0xf8,0x02,0xb9]
+
+v_fma_legacy_f32 v0, v1, v2, v3
+// GFX10: encoding: [0x00,0x00,0x40,0xd5,0x01,0x05,0x0e,0x04]
+
+v_fma_legacy_f32 v0, v1, |v2|, -v3
+// GFX10: encoding: [0x00,0x02,0x40,0xd5,0x01,0x05,0x0e,0x84]
+
+v_fma_legacy_f32 v0, s1, 2.0, -v3
+// GFX10: encoding: [0x00,0x00,0x40,0xd5,0x01,0xe8,0x0d,0x84]
+
+image_msaa_load v[1:4], v5, s[8:15] dmask:0xf dim:SQ_RSRC_IMG_1D
+// GFX10: encoding: [0x01,0x0f,0x00,0xf0,0x05,0x01,0x02,0x00]
+
+image_msaa_load v[1:4], v5, s[8:15] dmask:0xf dim:SQ_RSRC_IMG_1D glc
+// GFX10: encoding: [0x01,0x2f,0x00,0xf0,0x05,0x01,0x02,0x00]
+
+image_msaa_load v5, v[1:2], s[8:15] dmask:0x1 dim:SQ_RSRC_IMG_2D d16
+// GFX10: encoding: [0x09,0x01,0x00,0xf0,0x01,0x05,0x02,0x80]
+
+image_msaa_load v[1:4], v5, s[8:15] dmask:0xf dim:SQ_RSRC_IMG_1D
+// GFX10: encoding: [0x01,0x0f,0x00,0xf0,0x05,0x01,0x02,0x00]
+
+image_msaa_load v14, [v204,v11,v14,v19], s[40:47] dmask:0x1 dim:SQ_RSRC_IMG_2D_MSAA_ARRAY
+// GFX10: encoding: [0x3b,0x01,0x00,0xf0,0xcc,0x0e,0x0a,0x00,0x0b,0x0e,0x13,0x00]

llvm/test/MC/Disassembler/AMDGPU/gfx1030_dasm_new.txt

+# RUN: llvm-mc -arch=amdgcn -mcpu=gfx1030 -disassemble -show-encoding < %s | FileCheck -check-prefix=GFX10 %s
+
+# GFX10: global_load_dword_addtid v1, s[2:3] offset:16
+0x10,0x80,0x58,0xdc,0x00,0x00,0x02,0x01
+
+# GFX10: global_load_dword_addtid v1, s[2:3] offset:16 glc slc dlc
+0x10,0x90,0x5b,0xdc,0x00,0x00,0x02,0x01
+
+# GFX10: global_store_dword_addtid v1, s[2:3] offset:16 glc slc dlc
+0x10,0x90,0x5f,0xdc,0x00,0x01,0x02,0x00
+
+# GFX10: global_store_dword v[254:255], v1, s[2:3] offset:16
+0x10,0x80,0x70,0xdc,0xfe,0x01,0x02,0x00
+
+# GFX10: global_atomic_csub v2, v[0:1], v2, off offset:100 glc slc
+0x64,0x80,0xd3,0xdc,0x00,0x02,0x7d,0x02
+
+# GFX10: global_atomic_csub v2, v[0:1], v2, off glc
+0x00,0x80,0xd1,0xdc,0x00,0x02,0x7d,0x02
+
+# GFX10: global_atomic_csub v2, v[0:1], v2, s[2:3] glc
+0x00,0x80,0xd1,0xdc,0x00,0x02,0x02,0x02
+
+# GFX10: global_atomic_csub v2, v[0:1], v2, s[2:3] offset:100 glc slc
+0x64,0x80,0xd3,0xdc,0x00,0x02,0x02,0x02
+
+# GFX10: buffer_atomic_csub v5, off, s[8:11], s3
+0x00,0x40,0xd0,0xe0,0x00,0x05,0x02,0x03
+
+# GFX10: buffer_atomic_csub v5, off, s[8:11], s3 offset:4095 glc
+0xff,0x4f,0xd0,0xe0,0x00,0x05,0x02,0x03
+
+# GFX10: buffer_atomic_csub v5, off, s[8:11], -1 offset:4095 glc
+0xff,0x4f,0xd0,0xe0,0x00,0x05,0x02,0xc1
+
+# GFX10: buffer_atomic_csub v5, v0, s[8:11], s3 offen offset:4095 glc
+0xff,0x5f,0xd0,0xe0,0x00,0x05,0x02,0x03
+
+# GFX10: buffer_atomic_csub v5, v0, s[8:11], s3 idxen offset:4095 glc
+0xff,0x6f,0xd0,0xe0,0x00,0x05,0x02,0x03
+
+# GFX10: buffer_atomic_csub v5, off, s[8:11], s3 glc slc
+0x00,0x40,0xd0,0xe0,0x00,0x05,0x42,0x03
+
+# GFX10: v_fma_legacy_f32 v0, v1, v2, v3
+0x00,0x00,0x40,0xd5,0x01,0x05,0x0e,0x04
+
+# GFX10: v_fma_legacy_f32 v0, v1, |v2|, -v3
+0x00,0x02,0x40,0xd5,0x01,0x05,0x0e,0x84
+
+# GFX10: v_fma_legacy_f32 v0, s1, 2.0, -v3
+0x00,0x00,0x40,0xd5,0x01,0xe8,0x0d,0x84
+
+# GFX10: image_msaa_load v[1:4], v5, s[8:15] dmask:0xf dim:SQ_RSRC_IMG_1D
+0x01,0x0f,0x00,0xf0,0x05,0x01,0x02,0x00
+
+# GFX10: image_msaa_load v[1:4], v5, s[8:15] dmask:0xf dim:SQ_RSRC_IMG_1D glc
+0x01,0x2f,0x00,0xf0,0x05,0x01,0x02,0x00
+
+# GFX10: image_msaa_load v5, v[1:2], s[8:15] dmask:0x1 dim:SQ_RSRC_IMG_2D d16
+0x09,0x01,0x00,0xf0,0x01,0x05,0x02,0x80
+
+# GFX10: image_msaa_load v[1:4], v5, s[8:15] dmask:0xf dim:SQ_RSRC_IMG_1D
+0x01,0x0f,0x00,0xf0,0x05,0x01,0x02,0x00
+
+# GFX10: image_msaa_load v14, [v204, v11, v14, v19], s[40:47] dmask:0x1 dim:SQ_RSRC_IMG_2D_MSAA_ARRAY
+0x3b,0x01,0x00,0xf0,0xcc,0x0e,0x0a,0x00,0x0b,0x0e,0x13,0x00

llvm/tools/llvm-readobj/ELFDumper.cpp

static const EnumEntry<unsigned> ElfHeaderAMDGPUFlags[] = {
   LLVM_READOBJ_ENUM_ENT(ELF, EF_AMDGPU_MACH_AMDGCN_GFX902),
   LLVM_READOBJ_ENUM_ENT(ELF, EF_AMDGPU_MACH_AMDGCN_GFX904),
   LLVM_READOBJ_ENUM_ENT(ELF, EF_AMDGPU_MACH_AMDGCN_GFX906),
   LLVM_READOBJ_ENUM_ENT(ELF, EF_AMDGPU_MACH_AMDGCN_GFX908),
   LLVM_READOBJ_ENUM_ENT(ELF, EF_AMDGPU_MACH_AMDGCN_GFX909),
   LLVM_READOBJ_ENUM_ENT(ELF, EF_AMDGPU_MACH_AMDGCN_GFX1010),
   LLVM_READOBJ_ENUM_ENT(ELF, EF_AMDGPU_MACH_AMDGCN_GFX1011),
   LLVM_READOBJ_ENUM_ENT(ELF, EF_AMDGPU_MACH_AMDGCN_GFX1012),
+  LLVM_READOBJ_ENUM_ENT(ELF, EF_AMDGPU_MACH_AMDGCN_GFX1030),
   LLVM_READOBJ_ENUM_ENT(ELF, EF_AMDGPU_XNACK),
   LLVM_READOBJ_ENUM_ENT(ELF, EF_AMDGPU_SRAM_ECC)
 };
 
 static const EnumEntry<unsigned> ElfHeaderRISCVFlags[] = {
   ENUM_ENT(EF_RISCV_RVC, "RVC"),
   ENUM_ENT(EF_RISCV_FLOAT_ABI_SINGLE, "single-float ABI"),
   ENUM_ENT(EF_RISCV_FLOAT_ABI_DOUBLE, "double-float ABI"),