diff --git a/clang/examples/clang-interpreter/main.cpp b/clang/examples/clang-interpreter/main.cpp
index 134d70774f875a..6ac142bffdffca 100644
--- a/clang/examples/clang-interpreter/main.cpp
+++ b/clang/examples/clang-interpreter/main.cpp
@@ -173,11 +173,7 @@ int main(int argc, const char **argv) {
   // Initialize a compiler invocation object from the clang (-cc1) arguments.
   const llvm::opt::ArgStringList &CCArgs = Cmd.getArguments();
   std::unique_ptr<CompilerInvocation> CI(new CompilerInvocation);
-  CompilerInvocation::CreateFromArgs(*CI,
-                                     const_cast<const char **>(CCArgs.data()),
-                                     const_cast<const char **>(CCArgs.data()) +
-                                       CCArgs.size(),
-                                     Diags);
+  CompilerInvocation::CreateFromArgs(*CI, CCArgs, Diags);
 
   // Show the invocation, with -v.
   if (CI->getHeaderSearchOpts().Verbose) {
diff --git a/compiler-rt/CMakeLists.txt b/compiler-rt/CMakeLists.txt
index d28fd20d581362..50d0eb0ed9034b 100644
--- a/compiler-rt/CMakeLists.txt
+++ b/compiler-rt/CMakeLists.txt
@@ -463,24 +463,30 @@ add_subdirectory(include)
 option(COMPILER_RT_USE_LIBCXX
   "Enable compiler-rt to use libc++ from the source tree" ON)
 if(COMPILER_RT_USE_LIBCXX)
-  foreach(path IN ITEMS ${LLVM_MAIN_SRC_DIR}/projects/libcxx
-                        ${LLVM_MAIN_SRC_DIR}/runtimes/libcxx
-                        ${LLVM_MAIN_SRC_DIR}/../libcxx
-                        ${LLVM_EXTERNAL_LIBCXX_SOURCE_DIR})
-    if(IS_DIRECTORY ${path})
-      set(COMPILER_RT_LIBCXX_PATH ${path})
-      break()
-    endif()
-  endforeach()
-  foreach(path IN ITEMS ${LLVM_MAIN_SRC_DIR}/projects/libcxxabi
-                        ${LLVM_MAIN_SRC_DIR}/runtimes/libcxxabi
-                        ${LLVM_MAIN_SRC_DIR}/../libcxxabi
-                        ${LLVM_EXTERNAL_LIBCXXABI_SOURCE_DIR})
-    if(IS_DIRECTORY ${path})
-      set(COMPILER_RT_LIBCXXABI_PATH ${path})
-      break()
-    endif()
-  endforeach()
+  if(LLVM_ENABLE_PROJECTS_USED)
+    # Don't use libcxx if LLVM_ENABLE_PROJECTS does not enable it.
+    set(COMPILER_RT_LIBCXX_PATH ${LLVM_EXTERNAL_LIBCXX_SOURCE_DIR})
+    set(COMPILER_RT_LIBCXXABI_PATH ${LLVM_EXTERNAL_LIBCXXABI_SOURCE_DIR})
+  else()
+    foreach(path IN ITEMS ${LLVM_MAIN_SRC_DIR}/projects/libcxx
+                          ${LLVM_MAIN_SRC_DIR}/runtimes/libcxx
+                          ${LLVM_MAIN_SRC_DIR}/../libcxx
+                          ${LLVM_EXTERNAL_LIBCXX_SOURCE_DIR})
+      if(IS_DIRECTORY ${path})
+        set(COMPILER_RT_LIBCXX_PATH ${path})
+        break()
+      endif()
+    endforeach()
+    foreach(path IN ITEMS ${LLVM_MAIN_SRC_DIR}/projects/libcxxabi
+                          ${LLVM_MAIN_SRC_DIR}/runtimes/libcxxabi
+                          ${LLVM_MAIN_SRC_DIR}/../libcxxabi
+                          ${LLVM_EXTERNAL_LIBCXXABI_SOURCE_DIR})
+      if(IS_DIRECTORY ${path})
+        set(COMPILER_RT_LIBCXXABI_PATH ${path})
+        break()
+      endif()
+    endforeach()
+  endif()
 endif()
 
 set(COMPILER_RT_LLD_PATH ${LLVM_MAIN_SRC_DIR}/tools/lld)
diff --git a/lldb/source/Plugins/LanguageRuntime/ObjC/AppleObjCRuntime/AppleObjCRuntimeV2.cpp b/lldb/source/Plugins/LanguageRuntime/ObjC/AppleObjCRuntime/AppleObjCRuntimeV2.cpp
index 1ff3fc42dd0a5b..5284775343b8ba 100644
--- a/lldb/source/Plugins/LanguageRuntime/ObjC/AppleObjCRuntime/AppleObjCRuntimeV2.cpp
+++ b/lldb/source/Plugins/LanguageRuntime/ObjC/AppleObjCRuntime/AppleObjCRuntimeV2.cpp
@@ -1601,7 +1601,7 @@ AppleObjCRuntimeV2::UpdateISAToDescriptorMapSharedCache() {
     // use that in our jitted expression.  Else fall back to the old
     // class_getName.
     static ConstString g_class_getName_symbol_name("class_getName");
-    static ConstString g_class_getNameRaw_symbol_name("class_getNameRaw");
+    static ConstString g_class_getNameRaw_symbol_name("objc_debug_class_getNameRaw");
     ConstString class_name_getter_function_name = g_class_getName_symbol_name;
 
     ObjCLanguageRuntime *objc_runtime = ObjCLanguageRuntime::Get(*process);
diff --git a/llvm/lib/Target/AMDGPU/AMDGPUInstructionSelector.cpp b/llvm/lib/Target/AMDGPU/AMDGPUInstructionSelector.cpp
index 2421f294b121e2..2fe463c190b58c 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPUInstructionSelector.cpp
+++ b/llvm/lib/Target/AMDGPU/AMDGPUInstructionSelector.cpp
@@ -295,14 +295,7 @@ bool AMDGPUInstructionSelector::selectG_AND_OR_XOR(MachineInstr &I) const {
   if (DstRB->getID() == AMDGPU::SGPRRegBankID) {
     unsigned InstOpc = getLogicalBitOpcode(I.getOpcode(), Size > 32);
     I.setDesc(TII.get(InstOpc));
-
-    const TargetRegisterClass *RC
-      = TRI.getConstrainedRegClassForOperand(Dst, MRI);
-    if (!RC)
-      return false;
-    return RBI.constrainGenericRegister(DstReg, *RC, MRI) &&
-           RBI.constrainGenericRegister(Src0.getReg(), *RC, MRI) &&
-           RBI.constrainGenericRegister(Src1.getReg(), *RC, MRI);
+    return constrainSelectedInstRegOperands(I, TII, TRI, RBI);
   }
 
   return false;
diff --git a/llvm/test/CodeGen/AMDGPU/GlobalISel/inst-select-and.mir b/llvm/test/CodeGen/AMDGPU/GlobalISel/inst-select-and.mir
index d3b877d72e28c3..799b7cd61204f9 100644
--- a/llvm/test/CodeGen/AMDGPU/GlobalISel/inst-select-and.mir
+++ b/llvm/test/CodeGen/AMDGPU/GlobalISel/inst-select-and.mir
@@ -54,14 +54,14 @@ body: |
     ; WAVE64: liveins: $sgpr0, $sgpr1
     ; WAVE64: [[COPY:%[0-9]+]]:sreg_32_xm0 = COPY $sgpr0
     ; WAVE64: [[COPY1:%[0-9]+]]:sreg_32_xm0 = COPY $sgpr1
-    ; WAVE64: [[S_AND_B32_:%[0-9]+]]:sreg_32_xm0 = S_AND_B32 [[COPY]], [[COPY1]]
+    ; WAVE64: [[S_AND_B32_:%[0-9]+]]:sreg_32 = S_AND_B32 [[COPY]], [[COPY1]]
     ; WAVE64: S_ENDPGM 0, implicit [[S_AND_B32_]]
     ; WAVE32-LABEL: name: and_s1_sgpr_sgpr_sgpr
     ; WAVE32: liveins: $sgpr0, $sgpr1
     ; WAVE32: $vcc_hi = IMPLICIT_DEF
     ; WAVE32: [[COPY:%[0-9]+]]:sreg_32_xm0 = COPY $sgpr0
     ; WAVE32: [[COPY1:%[0-9]+]]:sreg_32_xm0 = COPY $sgpr1
-    ; WAVE32: [[S_AND_B32_:%[0-9]+]]:sreg_32_xm0 = S_AND_B32 [[COPY]], [[COPY1]]
+    ; WAVE32: [[S_AND_B32_:%[0-9]+]]:sreg_32 = S_AND_B32 [[COPY]], [[COPY1]]
     ; WAVE32: S_ENDPGM 0, implicit [[S_AND_B32_]]
     %0:sgpr(s32) = COPY $sgpr0
     %1:sgpr(s32) = COPY $sgpr1
@@ -119,14 +119,14 @@ body: |
     ; WAVE64: liveins: $sgpr0, $sgpr1
     ; WAVE64: [[COPY:%[0-9]+]]:sreg_32_xm0 = COPY $sgpr0
     ; WAVE64: [[COPY1:%[0-9]+]]:sreg_32_xm0 = COPY $sgpr1
-    ; WAVE64: [[S_AND_B32_:%[0-9]+]]:sreg_32_xm0 = S_AND_B32 [[COPY]], [[COPY1]]
+    ; WAVE64: [[S_AND_B32_:%[0-9]+]]:sreg_32 = S_AND_B32 [[COPY]], [[COPY1]]
     ; WAVE64: S_ENDPGM 0, implicit [[S_AND_B32_]]
     ; WAVE32-LABEL: name: and_s16_sgpr_sgpr_sgpr
     ; WAVE32: liveins: $sgpr0, $sgpr1
     ; WAVE32: $vcc_hi = IMPLICIT_DEF
     ; WAVE32: [[COPY:%[0-9]+]]:sreg_32_xm0 = COPY $sgpr0
     ; WAVE32: [[COPY1:%[0-9]+]]:sreg_32_xm0 = COPY $sgpr1
-    ; WAVE32: [[S_AND_B32_:%[0-9]+]]:sreg_32_xm0 = S_AND_B32 [[COPY]], [[COPY1]]
+    ; WAVE32: [[S_AND_B32_:%[0-9]+]]:sreg_32 = S_AND_B32 [[COPY]], [[COPY1]]
     ; WAVE32: S_ENDPGM 0, implicit [[S_AND_B32_]]
     %0:sgpr(s32) = COPY $sgpr0
     %1:sgpr(s32) = COPY $sgpr1
@@ -182,16 +182,16 @@ body: |
     liveins: $sgpr0, $sgpr1
     ; WAVE64-LABEL: name: and_s32_sgpr_sgpr_sgpr
     ; WAVE64: liveins: $sgpr0, $sgpr1
-    ; WAVE64: [[COPY:%[0-9]+]]:sreg_32_xm0 = COPY $sgpr0
-    ; WAVE64: [[COPY1:%[0-9]+]]:sreg_32_xm0 = COPY $sgpr1
-    ; WAVE64: [[S_AND_B32_:%[0-9]+]]:sreg_32_xm0 = S_AND_B32 [[COPY]], [[COPY1]]
+    ; WAVE64: [[COPY:%[0-9]+]]:sreg_32 = COPY $sgpr0
+    ; WAVE64: [[COPY1:%[0-9]+]]:sreg_32 = COPY $sgpr1
+    ; WAVE64: [[S_AND_B32_:%[0-9]+]]:sreg_32 = S_AND_B32 [[COPY]], [[COPY1]]
     ; WAVE64: S_ENDPGM 0, implicit [[S_AND_B32_]]
     ; WAVE32-LABEL: name: and_s32_sgpr_sgpr_sgpr
     ; WAVE32: liveins: $sgpr0, $sgpr1
     ; WAVE32: $vcc_hi = IMPLICIT_DEF
-    ; WAVE32: [[COPY:%[0-9]+]]:sreg_32_xm0 = COPY $sgpr0
-    ; WAVE32: [[COPY1:%[0-9]+]]:sreg_32_xm0 = COPY $sgpr1
-    ; WAVE32: [[S_AND_B32_:%[0-9]+]]:sreg_32_xm0 = S_AND_B32 [[COPY]], [[COPY1]]
+    ; WAVE32: [[COPY:%[0-9]+]]:sreg_32 = COPY $sgpr0
+    ; WAVE32: [[COPY1:%[0-9]+]]:sreg_32 = COPY $sgpr1
+    ; WAVE32: [[S_AND_B32_:%[0-9]+]]:sreg_32 = S_AND_B32 [[COPY]], [[COPY1]]
     ; WAVE32: S_ENDPGM 0, implicit [[S_AND_B32_]]
     %0:sgpr(s32) = COPY $sgpr0
     %1:sgpr(s32) = COPY $sgpr1
@@ -211,16 +211,16 @@ body: |
     liveins: $sgpr0_sgpr1, $sgpr2_sgpr3
     ; WAVE64-LABEL: name: and_s64_sgpr_sgpr_sgpr
     ; WAVE64: liveins: $sgpr0_sgpr1, $sgpr2_sgpr3
-    ; WAVE64: [[COPY:%[0-9]+]]:sreg_64_xexec = COPY $sgpr0_sgpr1
-    ; WAVE64: [[COPY1:%[0-9]+]]:sreg_64_xexec = COPY $sgpr2_sgpr3
-    ; WAVE64: [[S_AND_B64_:%[0-9]+]]:sreg_64_xexec = S_AND_B64 [[COPY]], [[COPY1]]
+    ; WAVE64: [[COPY:%[0-9]+]]:sreg_64 = COPY $sgpr0_sgpr1
+    ; WAVE64: [[COPY1:%[0-9]+]]:sreg_64 = COPY $sgpr2_sgpr3
+    ; WAVE64: [[S_AND_B64_:%[0-9]+]]:sreg_64 = S_AND_B64 [[COPY]], [[COPY1]]
     ; WAVE64: S_ENDPGM 0, implicit [[S_AND_B64_]]
     ; WAVE32-LABEL: name: and_s64_sgpr_sgpr_sgpr
     ; WAVE32: liveins: $sgpr0_sgpr1, $sgpr2_sgpr3
     ; WAVE32: $vcc_hi = IMPLICIT_DEF
-    ; WAVE32: [[COPY:%[0-9]+]]:sreg_64_xexec = COPY $sgpr0_sgpr1
-    ; WAVE32: [[COPY1:%[0-9]+]]:sreg_64_xexec = COPY $sgpr2_sgpr3
-    ; WAVE32: [[S_AND_B64_:%[0-9]+]]:sreg_64_xexec = S_AND_B64 [[COPY]], [[COPY1]]
+    ; WAVE32: [[COPY:%[0-9]+]]:sreg_64 = COPY $sgpr0_sgpr1
+    ; WAVE32: [[COPY1:%[0-9]+]]:sreg_64 = COPY $sgpr2_sgpr3
+    ; WAVE32: [[S_AND_B64_:%[0-9]+]]:sreg_64 = S_AND_B64 [[COPY]], [[COPY1]]
     ; WAVE32: S_ENDPGM 0, implicit [[S_AND_B64_]]
     %0:sgpr(s64) = COPY $sgpr0_sgpr1
     %1:sgpr(s64) = COPY $sgpr2_sgpr3
@@ -240,16 +240,16 @@ body: |
     liveins: $sgpr0, $sgpr1
     ; WAVE64-LABEL: name: and_v2s16_sgpr_sgpr_sgpr
     ; WAVE64: liveins: $sgpr0, $sgpr1
-    ; WAVE64: [[COPY:%[0-9]+]]:sreg_32_xm0 = COPY $sgpr0
-    ; WAVE64: [[COPY1:%[0-9]+]]:sreg_32_xm0 = COPY $sgpr1
-    ; WAVE64: [[S_AND_B32_:%[0-9]+]]:sreg_32_xm0 = S_AND_B32 [[COPY]], [[COPY1]]
+    ; WAVE64: [[COPY:%[0-9]+]]:sreg_32 = COPY $sgpr0
+    ; WAVE64: [[COPY1:%[0-9]+]]:sreg_32 = COPY $sgpr1
+    ; WAVE64: [[S_AND_B32_:%[0-9]+]]:sreg_32 = S_AND_B32 [[COPY]], [[COPY1]]
     ; WAVE64: S_ENDPGM 0, implicit [[S_AND_B32_]]
     ; WAVE32-LABEL: name: and_v2s16_sgpr_sgpr_sgpr
     ; WAVE32: liveins: $sgpr0, $sgpr1
     ; WAVE32: $vcc_hi = IMPLICIT_DEF
-    ; WAVE32: [[COPY:%[0-9]+]]:sreg_32_xm0 = COPY $sgpr0
-    ; WAVE32: [[COPY1:%[0-9]+]]:sreg_32_xm0 = COPY $sgpr1
-    ; WAVE32: [[S_AND_B32_:%[0-9]+]]:sreg_32_xm0 = S_AND_B32 [[COPY]], [[COPY1]]
+    ; WAVE32: [[COPY:%[0-9]+]]:sreg_32 = COPY $sgpr0
+    ; WAVE32: [[COPY1:%[0-9]+]]:sreg_32 = COPY $sgpr1
+    ; WAVE32: [[S_AND_B32_:%[0-9]+]]:sreg_32 = S_AND_B32 [[COPY]], [[COPY1]]
     ; WAVE32: S_ENDPGM 0, implicit [[S_AND_B32_]]
     %0:sgpr(<2 x s16>) = COPY $sgpr0
     %1:sgpr(<2 x s16>) = COPY $sgpr1
@@ -269,16 +269,16 @@ body: |
     liveins: $sgpr0_sgpr1, $sgpr2_sgpr3
     ; WAVE64-LABEL: name: and_v2s32_sgpr_sgpr_sgpr
     ; WAVE64: liveins: $sgpr0_sgpr1, $sgpr2_sgpr3
-    ; WAVE64: [[COPY:%[0-9]+]]:sreg_64_xexec = COPY $sgpr0_sgpr1
-    ; WAVE64: [[COPY1:%[0-9]+]]:sreg_64_xexec = COPY $sgpr2_sgpr3
-    ; WAVE64: [[S_AND_B64_:%[0-9]+]]:sreg_64_xexec = S_AND_B64 [[COPY]], [[COPY1]]
+    ; WAVE64: [[COPY:%[0-9]+]]:sreg_64 = COPY $sgpr0_sgpr1
+    ; WAVE64: [[COPY1:%[0-9]+]]:sreg_64 = COPY $sgpr2_sgpr3
+    ; WAVE64: [[S_AND_B64_:%[0-9]+]]:sreg_64 = S_AND_B64 [[COPY]], [[COPY1]]
     ; WAVE64: S_ENDPGM 0, implicit [[S_AND_B64_]]
     ; WAVE32-LABEL: name: and_v2s32_sgpr_sgpr_sgpr
     ; WAVE32: liveins: $sgpr0_sgpr1, $sgpr2_sgpr3
     ; WAVE32: $vcc_hi = IMPLICIT_DEF
-    ; WAVE32: [[COPY:%[0-9]+]]:sreg_64_xexec = COPY $sgpr0_sgpr1
-    ; WAVE32: [[COPY1:%[0-9]+]]:sreg_64_xexec = COPY $sgpr2_sgpr3
-    ; WAVE32: [[S_AND_B64_:%[0-9]+]]:sreg_64_xexec = S_AND_B64 [[COPY]], [[COPY1]]
+    ; WAVE32: [[COPY:%[0-9]+]]:sreg_64 = COPY $sgpr0_sgpr1
+    ; WAVE32: [[COPY1:%[0-9]+]]:sreg_64 = COPY $sgpr2_sgpr3
+    ; WAVE32: [[S_AND_B64_:%[0-9]+]]:sreg_64 = S_AND_B64 [[COPY]], [[COPY1]]
     ; WAVE32: S_ENDPGM 0, implicit [[S_AND_B64_]]
     %0:sgpr(<2 x s32>) = COPY $sgpr0_sgpr1
     %1:sgpr(<2 x s32>) = COPY $sgpr2_sgpr3
@@ -298,16 +298,16 @@ body: |
     liveins: $sgpr0_sgpr1, $sgpr2_sgpr3
     ; WAVE64-LABEL: name: and_v4s16_sgpr_sgpr_sgpr
     ; WAVE64: liveins: $sgpr0_sgpr1, $sgpr2_sgpr3
-    ; WAVE64: [[COPY:%[0-9]+]]:sreg_64_xexec = COPY $sgpr0_sgpr1
-    ; WAVE64: [[COPY1:%[0-9]+]]:sreg_64_xexec = COPY $sgpr2_sgpr3
-    ; WAVE64: [[S_AND_B64_:%[0-9]+]]:sreg_64_xexec = S_AND_B64 [[COPY]], [[COPY1]]
+    ; WAVE64: [[COPY:%[0-9]+]]:sreg_64 = COPY $sgpr0_sgpr1
+    ; WAVE64: [[COPY1:%[0-9]+]]:sreg_64 = COPY $sgpr2_sgpr3
+    ; WAVE64: [[S_AND_B64_:%[0-9]+]]:sreg_64 = S_AND_B64 [[COPY]], [[COPY1]]
     ; WAVE64: S_ENDPGM 0, implicit [[S_AND_B64_]]
     ; WAVE32-LABEL: name: and_v4s16_sgpr_sgpr_sgpr
     ; WAVE32: liveins: $sgpr0_sgpr1, $sgpr2_sgpr3
     ; WAVE32: $vcc_hi = IMPLICIT_DEF
-    ; WAVE32: [[COPY:%[0-9]+]]:sreg_64_xexec = COPY $sgpr0_sgpr1
-    ; WAVE32: [[COPY1:%[0-9]+]]:sreg_64_xexec = COPY $sgpr2_sgpr3
-    ; WAVE32: [[S_AND_B64_:%[0-9]+]]:sreg_64_xexec = S_AND_B64 [[COPY]], [[COPY1]]
+    ; WAVE32: [[COPY:%[0-9]+]]:sreg_64 = COPY $sgpr0_sgpr1
+    ; WAVE32: [[COPY1:%[0-9]+]]:sreg_64 = COPY $sgpr2_sgpr3
+    ; WAVE32: [[S_AND_B64_:%[0-9]+]]:sreg_64 = S_AND_B64 [[COPY]], [[COPY1]]
     ; WAVE32: S_ENDPGM 0, implicit [[S_AND_B64_]]
     %0:sgpr(<4 x s16>) = COPY $sgpr0_sgpr1
     %1:sgpr(<4 x s16>) = COPY $sgpr2_sgpr3
@@ -432,11 +432,11 @@ tracksRegLiveness: true
 body: |
   bb.0:
     ; WAVE64-LABEL: name: and_s1_sgpr_undef_sgpr_undef_sgpr
-    ; WAVE64: [[S_AND_B32_:%[0-9]+]]:sreg_32_xm0 = S_AND_B32 undef %1:sreg_32_xm0, undef %2:sreg_32_xm0
+    ; WAVE64: [[S_AND_B32_:%[0-9]+]]:sreg_32 = S_AND_B32 undef %1:sreg_32, undef %2:sreg_32
     ; WAVE64: S_ENDPGM 0, implicit [[S_AND_B32_]]
     ; WAVE32-LABEL: name: and_s1_sgpr_undef_sgpr_undef_sgpr
     ; WAVE32: $vcc_hi = IMPLICIT_DEF
-    ; WAVE32: [[S_AND_B32_:%[0-9]+]]:sreg_32_xm0 = S_AND_B32 undef %1:sreg_32_xm0, undef %2:sreg_32_xm0
+    ; WAVE32: [[S_AND_B32_:%[0-9]+]]:sreg_32 = S_AND_B32 undef %1:sreg_32, undef %2:sreg_32
     ; WAVE32: S_ENDPGM 0, implicit [[S_AND_B32_]]
     %2:sgpr(s1) = G_AND undef %0:sgpr(s1), undef %1:sgpr(s1)
     S_ENDPGM 0, implicit %2
@@ -513,16 +513,14 @@ body:             |
 
     ; WAVE64-LABEL: name: copy_select_constrain_vcc_result_reg_wave32
     ; WAVE64: liveins: $vgpr0
-    ; WAVE64: [[COPY:%[0-9]+]]:vgpr(s32) = COPY $vgpr0
-    ; WAVE64: [[TRUNC:%[0-9]+]]:vgpr(s1) = G_TRUNC [[COPY]](s32)
-    ; WAVE64: [[C:%[0-9]+]]:sgpr(s1) = G_CONSTANT i1 true
-    ; WAVE64: [[C1:%[0-9]+]]:sgpr(s32) = G_CONSTANT i32 0
-    ; WAVE64: [[DEF:%[0-9]+]]:sgpr(p1) = G_IMPLICIT_DEF
-    ; WAVE64: [[COPY1:%[0-9]+]]:vcc(s1) = COPY [[TRUNC]](s1)
-    ; WAVE64: [[COPY2:%[0-9]+]]:vcc(s1) = COPY [[C]](s1)
-    ; WAVE64: [[S_AND_B32_:%[0-9]+]]:sreg_64_xexec(s1) = S_AND_B32 [[COPY1]](s1), [[COPY2]](s1)
-    ; WAVE64: [[COPY3:%[0-9]+]]:sreg_32_xm0(s1) = COPY [[S_AND_B32_]](s1)
-    ; WAVE64: S_ENDPGM 0, implicit [[COPY3]](s1)
+    ; WAVE64: [[COPY:%[0-9]+]]:vgpr_32 = COPY $vgpr0
+    ; WAVE64: [[S_MOV_B32_:%[0-9]+]]:sreg_32_xm0 = S_MOV_B32 1
+    ; WAVE64: [[COPY1:%[0-9]+]]:sreg_32 = COPY [[COPY]]
+    ; WAVE64: [[COPY2:%[0-9]+]]:sreg_32 = COPY [[S_MOV_B32_]]
+    ; WAVE64: [[S_AND_B32_:%[0-9]+]]:sreg_32 = S_AND_B32 [[COPY1]], [[COPY2]]
+    ; WAVE64: [[COPY3:%[0-9]+]]:sreg_64_xexec = COPY [[S_AND_B32_]]
+    ; WAVE64: [[COPY4:%[0-9]+]]:sreg_32_xm0 = COPY [[COPY3]]
+    ; WAVE64: S_ENDPGM 0, implicit [[COPY4]]
     ; WAVE32-LABEL: name: copy_select_constrain_vcc_result_reg_wave32
     ; WAVE32: liveins: $vgpr0
     ; WAVE32: $vcc_hi = IMPLICIT_DEF
@@ -566,16 +564,14 @@ body:             |
     ; WAVE64: S_ENDPGM 0, implicit [[COPY1]]
     ; WAVE32-LABEL: name: copy_select_constrain_vcc_result_reg_wave64
     ; WAVE32: liveins: $vgpr0
-    ; WAVE32: [[COPY:%[0-9]+]]:vgpr(s32) = COPY $vgpr0
-    ; WAVE32: [[TRUNC:%[0-9]+]]:vgpr(s1) = G_TRUNC [[COPY]](s32)
-    ; WAVE32: [[C:%[0-9]+]]:sgpr(s1) = G_CONSTANT i1 true
-    ; WAVE32: [[C1:%[0-9]+]]:sgpr(s32) = G_CONSTANT i32 0
-    ; WAVE32: [[DEF:%[0-9]+]]:sgpr(p1) = G_IMPLICIT_DEF
-    ; WAVE32: [[COPY1:%[0-9]+]]:vcc(s1) = COPY [[TRUNC]](s1)
-    ; WAVE32: [[COPY2:%[0-9]+]]:vcc(s1) = COPY [[C]](s1)
-    ; WAVE32: [[S_AND_B32_:%[0-9]+]]:sreg_32_xm0_xexec(s1) = S_AND_B32 [[COPY1]](s1), [[COPY2]](s1)
-    ; WAVE32: [[COPY3:%[0-9]+]]:sreg_64_xexec(s1) = COPY [[S_AND_B32_]](s1)
-    ; WAVE32: S_ENDPGM 0, implicit [[COPY3]](s1)
+    ; WAVE32: $vcc_hi = IMPLICIT_DEF
+    ; WAVE32: [[COPY:%[0-9]+]]:vgpr_32 = COPY $vgpr0
+    ; WAVE32: [[S_MOV_B32_:%[0-9]+]]:sreg_32_xm0 = S_MOV_B32 1
+    ; WAVE32: [[COPY1:%[0-9]+]]:sreg_32 = COPY [[COPY]]
+    ; WAVE32: [[COPY2:%[0-9]+]]:sreg_32 = COPY [[S_MOV_B32_]]
+    ; WAVE32: [[S_AND_B32_:%[0-9]+]]:sreg_32_xm0_xexec = S_AND_B32 [[COPY1]], [[COPY2]]
+    ; WAVE32: [[COPY3:%[0-9]+]]:sreg_64_xexec = COPY [[S_AND_B32_]]
+    ; WAVE32: S_ENDPGM 0, implicit [[COPY3]]
     %1:vgpr(s32) = COPY $vgpr0
     %0:vgpr(s1) = G_TRUNC %1(s32)
     %2:sgpr(s1) = G_CONSTANT i1 true
@@ -588,3 +584,32 @@ body:             |
     S_ENDPGM 0, implicit %3
 
 ...
+
+---
+
+name:            and_s32_sgpr_sgpr_sgpr_result_reg_class
+legalized:       true
+regBankSelected: true
+tracksRegLiveness: true
+
+body: |
+  bb.0:
+    liveins: $sgpr0, $sgpr1
+    ; WAVE64-LABEL: name: and_s32_sgpr_sgpr_sgpr_result_reg_class
+    ; WAVE64: liveins: $sgpr0, $sgpr1
+    ; WAVE64: [[COPY:%[0-9]+]]:sreg_32 = COPY $sgpr0
+    ; WAVE64: [[COPY1:%[0-9]+]]:sreg_32 = COPY $sgpr1
+    ; WAVE64: [[S_AND_B32_:%[0-9]+]]:sreg_32 = S_AND_B32 [[COPY]], [[COPY1]]
+    ; WAVE64: S_ENDPGM 0, implicit [[S_AND_B32_]]
+    ; WAVE32-LABEL: name: and_s32_sgpr_sgpr_sgpr_result_reg_class
+    ; WAVE32: liveins: $sgpr0, $sgpr1
+    ; WAVE32: $vcc_hi = IMPLICIT_DEF
+    ; WAVE32: [[COPY:%[0-9]+]]:sreg_32 = COPY $sgpr0
+    ; WAVE32: [[COPY1:%[0-9]+]]:sreg_32 = COPY $sgpr1
+    ; WAVE32: [[S_AND_B32_:%[0-9]+]]:sreg_32 = S_AND_B32 [[COPY]], [[COPY1]]
+    ; WAVE32: S_ENDPGM 0, implicit [[S_AND_B32_]]
+    %0:sgpr(s32) = COPY $sgpr0
+    %1:sgpr(s32) = COPY $sgpr1
+    %2:sreg_32(s32) = G_AND %0, %1
+    S_ENDPGM 0, implicit %2
+...
diff --git a/llvm/test/CodeGen/AMDGPU/GlobalISel/inst-select-or.mir b/llvm/test/CodeGen/AMDGPU/GlobalISel/inst-select-or.mir
index d102761158159c..6e65cda1df93db 100644
--- a/llvm/test/CodeGen/AMDGPU/GlobalISel/inst-select-or.mir
+++ b/llvm/test/CodeGen/AMDGPU/GlobalISel/inst-select-or.mir
@@ -54,14 +54,14 @@ body: |
     ; WAVE64: liveins: $sgpr0, $sgpr1
     ; WAVE64: [[COPY:%[0-9]+]]:sreg_32_xm0 = COPY $sgpr0
     ; WAVE64: [[COPY1:%[0-9]+]]:sreg_32_xm0 = COPY $sgpr1
-    ; WAVE64: [[S_OR_B32_:%[0-9]+]]:sreg_32_xm0 = S_OR_B32 [[COPY]], [[COPY1]]
+    ; WAVE64: [[S_OR_B32_:%[0-9]+]]:sreg_32 = S_OR_B32 [[COPY]], [[COPY1]]
     ; WAVE64: S_ENDPGM 0, implicit [[S_OR_B32_]]
     ; WAVE32-LABEL: name: or_s1_sgpr_sgpr_sgpr
     ; WAVE32: liveins: $sgpr0, $sgpr1
     ; WAVE32: $vcc_hi = IMPLICIT_DEF
     ; WAVE32: [[COPY:%[0-9]+]]:sreg_32_xm0 = COPY $sgpr0
     ; WAVE32: [[COPY1:%[0-9]+]]:sreg_32_xm0 = COPY $sgpr1
-    ; WAVE32: [[S_OR_B32_:%[0-9]+]]:sreg_32_xm0 = S_OR_B32 [[COPY]], [[COPY1]]
+    ; WAVE32: [[S_OR_B32_:%[0-9]+]]:sreg_32 = S_OR_B32 [[COPY]], [[COPY1]]
     ; WAVE32: S_ENDPGM 0, implicit [[S_OR_B32_]]
     %0:sgpr(s32) = COPY $sgpr0
     %1:sgpr(s32) = COPY $sgpr1
@@ -119,14 +119,14 @@ body: |
     ; WAVE64: liveins: $sgpr0, $sgpr1
     ; WAVE64: [[COPY:%[0-9]+]]:sreg_32_xm0 = COPY $sgpr0
     ; WAVE64: [[COPY1:%[0-9]+]]:sreg_32_xm0 = COPY $sgpr1
-    ; WAVE64: [[S_OR_B32_:%[0-9]+]]:sreg_32_xm0 = S_OR_B32 [[COPY]], [[COPY1]]
+    ; WAVE64: [[S_OR_B32_:%[0-9]+]]:sreg_32 = S_OR_B32 [[COPY]], [[COPY1]]
     ; WAVE64: S_ENDPGM 0, implicit [[S_OR_B32_]]
     ; WAVE32-LABEL: name: or_s16_sgpr_sgpr_sgpr
     ; WAVE32: liveins: $sgpr0, $sgpr1
     ; WAVE32: $vcc_hi = IMPLICIT_DEF
     ; WAVE32: [[COPY:%[0-9]+]]:sreg_32_xm0 = COPY $sgpr0
     ; WAVE32: [[COPY1:%[0-9]+]]:sreg_32_xm0 = COPY $sgpr1
-    ; WAVE32: [[S_OR_B32_:%[0-9]+]]:sreg_32_xm0 = S_OR_B32 [[COPY]], [[COPY1]]
+    ; WAVE32: [[S_OR_B32_:%[0-9]+]]:sreg_32 = S_OR_B32 [[COPY]], [[COPY1]]
     ; WAVE32: S_ENDPGM 0, implicit [[S_OR_B32_]]
     %0:sgpr(s32) = COPY $sgpr0
     %1:sgpr(s32) = COPY $sgpr1
@@ -182,16 +182,16 @@ body: |
     liveins: $sgpr0, $sgpr1
     ; WAVE64-LABEL: name: or_s32_sgpr_sgpr_sgpr
     ; WAVE64: liveins: $sgpr0, $sgpr1
-    ; WAVE64: [[COPY:%[0-9]+]]:sreg_32_xm0 = COPY $sgpr0
-    ; WAVE64: [[COPY1:%[0-9]+]]:sreg_32_xm0 = COPY $sgpr1
-    ; WAVE64: [[S_OR_B32_:%[0-9]+]]:sreg_32_xm0 = S_OR_B32 [[COPY]], [[COPY1]]
+    ; WAVE64: [[COPY:%[0-9]+]]:sreg_32 = COPY $sgpr0
+    ; WAVE64: [[COPY1:%[0-9]+]]:sreg_32 = COPY $sgpr1
+    ; WAVE64: [[S_OR_B32_:%[0-9]+]]:sreg_32 = S_OR_B32 [[COPY]], [[COPY1]]
     ; WAVE64: S_ENDPGM 0, implicit [[S_OR_B32_]]
     ; WAVE32-LABEL: name: or_s32_sgpr_sgpr_sgpr
     ; WAVE32: liveins: $sgpr0, $sgpr1
     ; WAVE32: $vcc_hi = IMPLICIT_DEF
-    ; WAVE32: [[COPY:%[0-9]+]]:sreg_32_xm0 = COPY $sgpr0
-    ; WAVE32: [[COPY1:%[0-9]+]]:sreg_32_xm0 = COPY $sgpr1
-    ; WAVE32: [[S_OR_B32_:%[0-9]+]]:sreg_32_xm0 = S_OR_B32 [[COPY]], [[COPY1]]
+    ; WAVE32: [[COPY:%[0-9]+]]:sreg_32 = COPY $sgpr0
+    ; WAVE32: [[COPY1:%[0-9]+]]:sreg_32 = COPY $sgpr1
+    ; WAVE32: [[S_OR_B32_:%[0-9]+]]:sreg_32 = S_OR_B32 [[COPY]], [[COPY1]]
     ; WAVE32: S_ENDPGM 0, implicit [[S_OR_B32_]]
     %0:sgpr(s32) = COPY $sgpr0
     %1:sgpr(s32) = COPY $sgpr1
@@ -211,16 +211,16 @@ body: |
     liveins: $sgpr0_sgpr1, $sgpr2_sgpr3
     ; WAVE64-LABEL: name: or_s64_sgpr_sgpr_sgpr
     ; WAVE64: liveins: $sgpr0_sgpr1, $sgpr2_sgpr3
-    ; WAVE64: [[COPY:%[0-9]+]]:sreg_64_xexec = COPY $sgpr0_sgpr1
-    ; WAVE64: [[COPY1:%[0-9]+]]:sreg_64_xexec = COPY $sgpr2_sgpr3
-    ; WAVE64: [[S_OR_B64_:%[0-9]+]]:sreg_64_xexec = S_OR_B64 [[COPY]], [[COPY1]]
+    ; WAVE64: [[COPY:%[0-9]+]]:sreg_64 = COPY $sgpr0_sgpr1
+    ; WAVE64: [[COPY1:%[0-9]+]]:sreg_64 = COPY $sgpr2_sgpr3
+    ; WAVE64: [[S_OR_B64_:%[0-9]+]]:sreg_64 = S_OR_B64 [[COPY]], [[COPY1]]
     ; WAVE64: S_ENDPGM 0, implicit [[S_OR_B64_]]
     ; WAVE32-LABEL: name: or_s64_sgpr_sgpr_sgpr
     ; WAVE32: liveins: $sgpr0_sgpr1, $sgpr2_sgpr3
     ; WAVE32: $vcc_hi = IMPLICIT_DEF
-    ; WAVE32: [[COPY:%[0-9]+]]:sreg_64_xexec = COPY $sgpr0_sgpr1
-    ; WAVE32: [[COPY1:%[0-9]+]]:sreg_64_xexec = COPY $sgpr2_sgpr3
-    ; WAVE32: [[S_OR_B64_:%[0-9]+]]:sreg_64_xexec = S_OR_B64 [[COPY]], [[COPY1]]
+    ; WAVE32: [[COPY:%[0-9]+]]:sreg_64 = COPY $sgpr0_sgpr1
+    ; WAVE32: [[COPY1:%[0-9]+]]:sreg_64 = COPY $sgpr2_sgpr3
+    ; WAVE32: [[S_OR_B64_:%[0-9]+]]:sreg_64 = S_OR_B64 [[COPY]], [[COPY1]]
     ; WAVE32: S_ENDPGM 0, implicit [[S_OR_B64_]]
     %0:sgpr(s64) = COPY $sgpr0_sgpr1
     %1:sgpr(s64) = COPY $sgpr2_sgpr3
@@ -240,16 +240,16 @@ body: |
     liveins: $sgpr0, $sgpr1
     ; WAVE64-LABEL: name: or_v2s16_sgpr_sgpr_sgpr
     ; WAVE64: liveins: $sgpr0, $sgpr1
-    ; WAVE64: [[COPY:%[0-9]+]]:sreg_32_xm0 = COPY $sgpr0
-    ; WAVE64: [[COPY1:%[0-9]+]]:sreg_32_xm0 = COPY $sgpr1
-    ; WAVE64: [[S_OR_B32_:%[0-9]+]]:sreg_32_xm0 = S_OR_B32 [[COPY]], [[COPY1]]
+    ; WAVE64: [[COPY:%[0-9]+]]:sreg_32 = COPY $sgpr0
+    ; WAVE64: [[COPY1:%[0-9]+]]:sreg_32 = COPY $sgpr1
+    ; WAVE64: [[S_OR_B32_:%[0-9]+]]:sreg_32 = S_OR_B32 [[COPY]], [[COPY1]]
     ; WAVE64: S_ENDPGM 0, implicit [[S_OR_B32_]]
     ; WAVE32-LABEL: name: or_v2s16_sgpr_sgpr_sgpr
     ; WAVE32: liveins: $sgpr0, $sgpr1
     ; WAVE32: $vcc_hi = IMPLICIT_DEF
-    ; WAVE32: [[COPY:%[0-9]+]]:sreg_32_xm0 = COPY $sgpr0
-    ; WAVE32: [[COPY1:%[0-9]+]]:sreg_32_xm0 = COPY $sgpr1
-    ; WAVE32: [[S_OR_B32_:%[0-9]+]]:sreg_32_xm0 = S_OR_B32 [[COPY]], [[COPY1]]
+    ; WAVE32: [[COPY:%[0-9]+]]:sreg_32 = COPY $sgpr0
+    ; WAVE32: [[COPY1:%[0-9]+]]:sreg_32 = COPY $sgpr1
+    ; WAVE32: [[S_OR_B32_:%[0-9]+]]:sreg_32 = S_OR_B32 [[COPY]], [[COPY1]]
     ; WAVE32: S_ENDPGM 0, implicit [[S_OR_B32_]]
     %0:sgpr(<2 x s16>) = COPY $sgpr0
     %1:sgpr(<2 x s16>) = COPY $sgpr1
@@ -269,16 +269,16 @@ body: |
     liveins: $sgpr0_sgpr1, $sgpr2_sgpr3
     ; WAVE64-LABEL: name: or_v2s32_sgpr_sgpr_sgpr
     ; WAVE64: liveins: $sgpr0_sgpr1, $sgpr2_sgpr3
-    ; WAVE64: [[COPY:%[0-9]+]]:sreg_64_xexec = COPY $sgpr0_sgpr1
-    ; WAVE64: [[COPY1:%[0-9]+]]:sreg_64_xexec = COPY $sgpr2_sgpr3
-    ; WAVE64: [[S_OR_B64_:%[0-9]+]]:sreg_64_xexec = S_OR_B64 [[COPY]], [[COPY1]]
+    ; WAVE64: [[COPY:%[0-9]+]]:sreg_64 = COPY $sgpr0_sgpr1
+    ; WAVE64: [[COPY1:%[0-9]+]]:sreg_64 = COPY $sgpr2_sgpr3
+    ; WAVE64: [[S_OR_B64_:%[0-9]+]]:sreg_64 = S_OR_B64 [[COPY]], [[COPY1]]
     ; WAVE64: S_ENDPGM 0, implicit [[S_OR_B64_]]
     ; WAVE32-LABEL: name: or_v2s32_sgpr_sgpr_sgpr
     ; WAVE32: liveins: $sgpr0_sgpr1, $sgpr2_sgpr3
     ; WAVE32: $vcc_hi = IMPLICIT_DEF
-    ; WAVE32: [[COPY:%[0-9]+]]:sreg_64_xexec = COPY $sgpr0_sgpr1
-    ; WAVE32: [[COPY1:%[0-9]+]]:sreg_64_xexec = COPY $sgpr2_sgpr3
-    ; WAVE32: [[S_OR_B64_:%[0-9]+]]:sreg_64_xexec = S_OR_B64 [[COPY]], [[COPY1]]
+    ; WAVE32: [[COPY:%[0-9]+]]:sreg_64 = COPY $sgpr0_sgpr1
+    ; WAVE32: [[COPY1:%[0-9]+]]:sreg_64 = COPY $sgpr2_sgpr3
+    ; WAVE32: [[S_OR_B64_:%[0-9]+]]:sreg_64 = S_OR_B64 [[COPY]], [[COPY1]]
     ; WAVE32: S_ENDPGM 0, implicit [[S_OR_B64_]]
     %0:sgpr(<2 x s32>) = COPY $sgpr0_sgpr1
     %1:sgpr(<2 x s32>) = COPY $sgpr2_sgpr3
@@ -298,16 +298,16 @@ body: |
     liveins: $sgpr0_sgpr1, $sgpr2_sgpr3
     ; WAVE64-LABEL: name: or_v4s16_sgpr_sgpr_sgpr
     ; WAVE64: liveins: $sgpr0_sgpr1, $sgpr2_sgpr3
-    ; WAVE64: [[COPY:%[0-9]+]]:sreg_64_xexec = COPY $sgpr0_sgpr1
-    ; WAVE64: [[COPY1:%[0-9]+]]:sreg_64_xexec = COPY $sgpr2_sgpr3
-    ; WAVE64: [[S_OR_B64_:%[0-9]+]]:sreg_64_xexec = S_OR_B64 [[COPY]], [[COPY1]]
+    ; WAVE64: [[COPY:%[0-9]+]]:sreg_64 = COPY $sgpr0_sgpr1
+    ; WAVE64: [[COPY1:%[0-9]+]]:sreg_64 = COPY $sgpr2_sgpr3
+    ; WAVE64: [[S_OR_B64_:%[0-9]+]]:sreg_64 = S_OR_B64 [[COPY]], [[COPY1]]
     ; WAVE64: S_ENDPGM 0, implicit [[S_OR_B64_]]
     ; WAVE32-LABEL: name: or_v4s16_sgpr_sgpr_sgpr
     ; WAVE32: liveins: $sgpr0_sgpr1, $sgpr2_sgpr3
     ; WAVE32: $vcc_hi = IMPLICIT_DEF
-    ; WAVE32: [[COPY:%[0-9]+]]:sreg_64_xexec = COPY $sgpr0_sgpr1
-    ; WAVE32: [[COPY1:%[0-9]+]]:sreg_64_xexec = COPY $sgpr2_sgpr3
-    ; WAVE32: [[S_OR_B64_:%[0-9]+]]:sreg_64_xexec = S_OR_B64 [[COPY]], [[COPY1]]
+    ; WAVE32: [[COPY:%[0-9]+]]:sreg_64 = COPY $sgpr0_sgpr1
+    ; WAVE32: [[COPY1:%[0-9]+]]:sreg_64 = COPY $sgpr2_sgpr3
+    ; WAVE32: [[S_OR_B64_:%[0-9]+]]:sreg_64 = S_OR_B64 [[COPY]], [[COPY1]]
     ; WAVE32: S_ENDPGM 0, implicit [[S_OR_B64_]]
     %0:sgpr(<4 x s16>) = COPY $sgpr0_sgpr1
     %1:sgpr(<4 x s16>) = COPY $sgpr2_sgpr3
@@ -432,11 +432,11 @@ tracksRegLiveness: true
 body: |
   bb.0:
     ; WAVE64-LABEL: name: or_s1_sgpr_undef_sgpr_undef_sgpr
-    ; WAVE64: [[S_OR_B32_:%[0-9]+]]:sreg_32_xm0 = S_OR_B32 undef %1:sreg_32_xm0, undef %2:sreg_32_xm0
+    ; WAVE64: [[S_OR_B32_:%[0-9]+]]:sreg_32 = S_OR_B32 undef %1:sreg_32, undef %2:sreg_32
     ; WAVE64: S_ENDPGM 0, implicit [[S_OR_B32_]]
     ; WAVE32-LABEL: name: or_s1_sgpr_undef_sgpr_undef_sgpr
     ; WAVE32: $vcc_hi = IMPLICIT_DEF
-    ; WAVE32: [[S_OR_B32_:%[0-9]+]]:sreg_32_xm0 = S_OR_B32 undef %1:sreg_32_xm0, undef %2:sreg_32_xm0
+    ; WAVE32: [[S_OR_B32_:%[0-9]+]]:sreg_32 = S_OR_B32 undef %1:sreg_32, undef %2:sreg_32
     ; WAVE32: S_ENDPGM 0, implicit [[S_OR_B32_]]
     %2:sgpr(s1) = G_OR undef %0:sgpr(s1), undef %1:sgpr(s1)
     S_ENDPGM 0, implicit %2
@@ -513,16 +513,14 @@ body:             |
 
     ; WAVE64-LABEL: name: copy_select_constrain_vcc_result_reg_wave32
     ; WAVE64: liveins: $vgpr0
-    ; WAVE64: [[COPY:%[0-9]+]]:vgpr(s32) = COPY $vgpr0
-    ; WAVE64: [[TRUNC:%[0-9]+]]:vgpr(s1) = G_TRUNC [[COPY]](s32)
-    ; WAVE64: [[C:%[0-9]+]]:sgpr(s1) = G_CONSTANT i1 true
-    ; WAVE64: [[C1:%[0-9]+]]:sgpr(s32) = G_CONSTANT i32 0
-    ; WAVE64: [[DEF:%[0-9]+]]:sgpr(p1) = G_IMPLICIT_DEF
-    ; WAVE64: [[COPY1:%[0-9]+]]:vcc(s1) = COPY [[TRUNC]](s1)
-    ; WAVE64: [[COPY2:%[0-9]+]]:vcc(s1) = COPY [[C]](s1)
-    ; WAVE64: [[S_OR_B32_:%[0-9]+]]:sreg_64_xexec(s1) = S_OR_B32 [[COPY1]](s1), [[COPY2]](s1)
-    ; WAVE64: [[COPY3:%[0-9]+]]:sreg_32_xm0(s1) = COPY [[S_OR_B32_]](s1)
-    ; WAVE64: S_ENDPGM 0, implicit [[COPY3]](s1)
+    ; WAVE64: [[COPY:%[0-9]+]]:vgpr_32 = COPY $vgpr0
+    ; WAVE64: [[S_MOV_B32_:%[0-9]+]]:sreg_32_xm0 = S_MOV_B32 1
+    ; WAVE64: [[COPY1:%[0-9]+]]:sreg_32 = COPY [[COPY]]
+    ; WAVE64: [[COPY2:%[0-9]+]]:sreg_32 = COPY [[S_MOV_B32_]]
+    ; WAVE64: [[S_OR_B32_:%[0-9]+]]:sreg_32 = S_OR_B32 [[COPY1]], [[COPY2]]
+    ; WAVE64: [[COPY3:%[0-9]+]]:sreg_64_xexec = COPY [[S_OR_B32_]]
+    ; WAVE64: [[COPY4:%[0-9]+]]:sreg_32_xm0 = COPY [[COPY3]]
+    ; WAVE64: S_ENDPGM 0, implicit [[COPY4]]
     ; WAVE32-LABEL: name: copy_select_constrain_vcc_result_reg_wave32
     ; WAVE32: liveins: $vgpr0
     ; WAVE32: $vcc_hi = IMPLICIT_DEF
@@ -566,16 +564,14 @@ body:             |
     ; WAVE64: S_ENDPGM 0, implicit [[COPY1]]
     ; WAVE32-LABEL: name: copy_select_constrain_vcc_result_reg_wave64
     ; WAVE32: liveins: $vgpr0
-    ; WAVE32: [[COPY:%[0-9]+]]:vgpr(s32) = COPY $vgpr0
-    ; WAVE32: [[TRUNC:%[0-9]+]]:vgpr(s1) = G_TRUNC [[COPY]](s32)
-    ; WAVE32: [[C:%[0-9]+]]:sgpr(s1) = G_CONSTANT i1 true
-    ; WAVE32: [[C1:%[0-9]+]]:sgpr(s32) = G_CONSTANT i32 0
-    ; WAVE32: [[DEF:%[0-9]+]]:sgpr(p1) = G_IMPLICIT_DEF
-    ; WAVE32: [[COPY1:%[0-9]+]]:vcc(s1) = COPY [[TRUNC]](s1)
-    ; WAVE32: [[COPY2:%[0-9]+]]:vcc(s1) = COPY [[C]](s1)
-    ; WAVE32: [[S_OR_B32_:%[0-9]+]]:sreg_32_xm0_xexec(s1) = S_OR_B32 [[COPY1]](s1), [[COPY2]](s1)
-    ; WAVE32: [[COPY3:%[0-9]+]]:sreg_64_xexec(s1) = COPY [[S_OR_B32_]](s1)
-    ; WAVE32: S_ENDPGM 0, implicit [[COPY3]](s1)
+    ; WAVE32: $vcc_hi = IMPLICIT_DEF
+    ; WAVE32: [[COPY:%[0-9]+]]:vgpr_32 = COPY $vgpr0
+    ; WAVE32: [[S_MOV_B32_:%[0-9]+]]:sreg_32_xm0 = S_MOV_B32 1
+    ; WAVE32: [[COPY1:%[0-9]+]]:sreg_32 = COPY [[COPY]]
+    ; WAVE32: [[COPY2:%[0-9]+]]:sreg_32 = COPY [[S_MOV_B32_]]
+    ; WAVE32: [[S_OR_B32_:%[0-9]+]]:sreg_32_xm0_xexec = S_OR_B32 [[COPY1]], [[COPY2]]
+    ; WAVE32: [[COPY3:%[0-9]+]]:sreg_64_xexec = COPY [[S_OR_B32_]]
+    ; WAVE32: S_ENDPGM 0, implicit [[COPY3]]
     %1:vgpr(s32) = COPY $vgpr0
     %0:vgpr(s1) = G_TRUNC %1(s32)
     %2:sgpr(s1) = G_CONSTANT i1 true
diff --git a/llvm/test/CodeGen/AMDGPU/GlobalISel/inst-select-xor.mir b/llvm/test/CodeGen/AMDGPU/GlobalISel/inst-select-xor.mir
index 74555ab9940eaa..b6d88e1a5f6e79 100644
--- a/llvm/test/CodeGen/AMDGPU/GlobalISel/inst-select-xor.mir
+++ b/llvm/test/CodeGen/AMDGPU/GlobalISel/inst-select-xor.mir
@@ -54,14 +54,14 @@ body: |
     ; WAVE64: liveins: $sgpr0, $sgpr1
     ; WAVE64: [[COPY:%[0-9]+]]:sreg_32_xm0 = COPY $sgpr0
     ; WAVE64: [[COPY1:%[0-9]+]]:sreg_32_xm0 = COPY $sgpr1
-    ; WAVE64: [[S_XOR_B32_:%[0-9]+]]:sreg_32_xm0 = S_XOR_B32 [[COPY]], [[COPY1]]
+    ; WAVE64: [[S_XOR_B32_:%[0-9]+]]:sreg_32 = S_XOR_B32 [[COPY]], [[COPY1]]
     ; WAVE64: S_ENDPGM 0, implicit [[S_XOR_B32_]]
     ; WAVE32-LABEL: name: xor_s1_sgpr_sgpr_sgpr
     ; WAVE32: liveins: $sgpr0, $sgpr1
     ; WAVE32: $vcc_hi = IMPLICIT_DEF
     ; WAVE32: [[COPY:%[0-9]+]]:sreg_32_xm0 = COPY $sgpr0
     ; WAVE32: [[COPY1:%[0-9]+]]:sreg_32_xm0 = COPY $sgpr1
-    ; WAVE32: [[S_XOR_B32_:%[0-9]+]]:sreg_32_xm0 = S_XOR_B32 [[COPY]], [[COPY1]]
+    ; WAVE32: [[S_XOR_B32_:%[0-9]+]]:sreg_32 = S_XOR_B32 [[COPY]], [[COPY1]]
     ; WAVE32: S_ENDPGM 0, implicit [[S_XOR_B32_]]
     %0:sgpr(s32) = COPY $sgpr0
     %1:sgpr(s32) = COPY $sgpr1
@@ -119,14 +119,14 @@ body: |
     ; WAVE64: liveins: $sgpr0, $sgpr1
     ; WAVE64: [[COPY:%[0-9]+]]:sreg_32_xm0 = COPY $sgpr0
     ; WAVE64: [[COPY1:%[0-9]+]]:sreg_32_xm0 = COPY $sgpr1
-    ; WAVE64: [[S_XOR_B32_:%[0-9]+]]:sreg_32_xm0 = S_XOR_B32 [[COPY]], [[COPY1]]
+    ; WAVE64: [[S_XOR_B32_:%[0-9]+]]:sreg_32 = S_XOR_B32 [[COPY]], [[COPY1]]
     ; WAVE64: S_ENDPGM 0, implicit [[S_XOR_B32_]]
     ; WAVE32-LABEL: name: xor_s16_sgpr_sgpr_sgpr
     ; WAVE32: liveins: $sgpr0, $sgpr1
     ; WAVE32: $vcc_hi = IMPLICIT_DEF
     ; WAVE32: [[COPY:%[0-9]+]]:sreg_32_xm0 = COPY $sgpr0
     ; WAVE32: [[COPY1:%[0-9]+]]:sreg_32_xm0 = COPY $sgpr1
-    ; WAVE32: [[S_XOR_B32_:%[0-9]+]]:sreg_32_xm0 = S_XOR_B32 [[COPY]], [[COPY1]]
+    ; WAVE32: [[S_XOR_B32_:%[0-9]+]]:sreg_32 = S_XOR_B32 [[COPY]], [[COPY1]]
     ; WAVE32: S_ENDPGM 0, implicit [[S_XOR_B32_]]
     %0:sgpr(s32) = COPY $sgpr0
     %1:sgpr(s32) = COPY $sgpr1
@@ -182,16 +182,16 @@ body: |
     liveins: $sgpr0, $sgpr1
     ; WAVE64-LABEL: name: xor_s32_sgpr_sgpr_sgpr
     ; WAVE64: liveins: $sgpr0, $sgpr1
-    ; WAVE64: [[COPY:%[0-9]+]]:sreg_32_xm0 = COPY $sgpr0
-    ; WAVE64: [[COPY1:%[0-9]+]]:sreg_32_xm0 = COPY $sgpr1
-    ; WAVE64: [[S_XOR_B32_:%[0-9]+]]:sreg_32_xm0 = S_XOR_B32 [[COPY]], [[COPY1]]
+    ; WAVE64: [[COPY:%[0-9]+]]:sreg_32 = COPY $sgpr0
+    ; WAVE64: [[COPY1:%[0-9]+]]:sreg_32 = COPY $sgpr1
+    ; WAVE64: [[S_XOR_B32_:%[0-9]+]]:sreg_32 = S_XOR_B32 [[COPY]], [[COPY1]]
     ; WAVE64: S_ENDPGM 0, implicit [[S_XOR_B32_]]
     ; WAVE32-LABEL: name: xor_s32_sgpr_sgpr_sgpr
     ; WAVE32: liveins: $sgpr0, $sgpr1
     ; WAVE32: $vcc_hi = IMPLICIT_DEF
-    ; WAVE32: [[COPY:%[0-9]+]]:sreg_32_xm0 = COPY $sgpr0
-    ; WAVE32: [[COPY1:%[0-9]+]]:sreg_32_xm0 = COPY $sgpr1
-    ; WAVE32: [[S_XOR_B32_:%[0-9]+]]:sreg_32_xm0 = S_XOR_B32 [[COPY]], [[COPY1]]
+    ; WAVE32: [[COPY:%[0-9]+]]:sreg_32 = COPY $sgpr0
+    ; WAVE32: [[COPY1:%[0-9]+]]:sreg_32 = COPY $sgpr1
+    ; WAVE32: [[S_XOR_B32_:%[0-9]+]]:sreg_32 = S_XOR_B32 [[COPY]], [[COPY1]]
     ; WAVE32: S_ENDPGM 0, implicit [[S_XOR_B32_]]
     %0:sgpr(s32) = COPY $sgpr0
     %1:sgpr(s32) = COPY $sgpr1
@@ -211,16 +211,16 @@ body: |
     liveins: $sgpr0_sgpr1, $sgpr2_sgpr3
     ; WAVE64-LABEL: name: xor_s64_sgpr_sgpr_sgpr
     ; WAVE64: liveins: $sgpr0_sgpr1, $sgpr2_sgpr3
-    ; WAVE64: [[COPY:%[0-9]+]]:sreg_64_xexec = COPY $sgpr0_sgpr1
-    ; WAVE64: [[COPY1:%[0-9]+]]:sreg_64_xexec = COPY $sgpr2_sgpr3
-    ; WAVE64: [[S_XOR_B64_:%[0-9]+]]:sreg_64_xexec = S_XOR_B64 [[COPY]], [[COPY1]]
+    ; WAVE64: [[COPY:%[0-9]+]]:sreg_64 = COPY $sgpr0_sgpr1
+    ; WAVE64: [[COPY1:%[0-9]+]]:sreg_64 = COPY $sgpr2_sgpr3
+    ; WAVE64: [[S_XOR_B64_:%[0-9]+]]:sreg_64 = S_XOR_B64 [[COPY]], [[COPY1]]
     ; WAVE64: S_ENDPGM 0, implicit [[S_XOR_B64_]]
     ; WAVE32-LABEL: name: xor_s64_sgpr_sgpr_sgpr
     ; WAVE32: liveins: $sgpr0_sgpr1, $sgpr2_sgpr3
     ; WAVE32: $vcc_hi = IMPLICIT_DEF
-    ; WAVE32: [[COPY:%[0-9]+]]:sreg_64_xexec = COPY $sgpr0_sgpr1
-    ; WAVE32: [[COPY1:%[0-9]+]]:sreg_64_xexec = COPY $sgpr2_sgpr3
-    ; WAVE32: [[S_XOR_B64_:%[0-9]+]]:sreg_64_xexec = S_XOR_B64 [[COPY]], [[COPY1]]
+    ; WAVE32: [[COPY:%[0-9]+]]:sreg_64 = COPY $sgpr0_sgpr1
+    ; WAVE32: [[COPY1:%[0-9]+]]:sreg_64 = COPY $sgpr2_sgpr3
+    ; WAVE32: [[S_XOR_B64_:%[0-9]+]]:sreg_64 = S_XOR_B64 [[COPY]], [[COPY1]]
     ; WAVE32: S_ENDPGM 0, implicit [[S_XOR_B64_]]
     %0:sgpr(s64) = COPY $sgpr0_sgpr1
     %1:sgpr(s64) = COPY $sgpr2_sgpr3
@@ -240,16 +240,16 @@ body: |
     liveins: $sgpr0, $sgpr1
     ; WAVE64-LABEL: name: xor_v2s16_sgpr_sgpr_sgpr
     ; WAVE64: liveins: $sgpr0, $sgpr1
-    ; WAVE64: [[COPY:%[0-9]+]]:sreg_32_xm0 = COPY $sgpr0
-    ; WAVE64: [[COPY1:%[0-9]+]]:sreg_32_xm0 = COPY $sgpr1
-    ; WAVE64: [[S_XOR_B32_:%[0-9]+]]:sreg_32_xm0 = S_XOR_B32 [[COPY]], [[COPY1]]
+    ; WAVE64: [[COPY:%[0-9]+]]:sreg_32 = COPY $sgpr0
+    ; WAVE64: [[COPY1:%[0-9]+]]:sreg_32 = COPY $sgpr1
+    ; WAVE64: [[S_XOR_B32_:%[0-9]+]]:sreg_32 = S_XOR_B32 [[COPY]], [[COPY1]]
     ; WAVE64: S_ENDPGM 0, implicit [[S_XOR_B32_]]
     ; WAVE32-LABEL: name: xor_v2s16_sgpr_sgpr_sgpr
     ; WAVE32: liveins: $sgpr0, $sgpr1
     ; WAVE32: $vcc_hi = IMPLICIT_DEF
-    ; WAVE32: [[COPY:%[0-9]+]]:sreg_32_xm0 = COPY $sgpr0
-    ; WAVE32: [[COPY1:%[0-9]+]]:sreg_32_xm0 = COPY $sgpr1
-    ; WAVE32: [[S_XOR_B32_:%[0-9]+]]:sreg_32_xm0 = S_XOR_B32 [[COPY]], [[COPY1]]
+    ; WAVE32: [[COPY:%[0-9]+]]:sreg_32 = COPY $sgpr0
+    ; WAVE32: [[COPY1:%[0-9]+]]:sreg_32 = COPY $sgpr1
+    ; WAVE32: [[S_XOR_B32_:%[0-9]+]]:sreg_32 = S_XOR_B32 [[COPY]], [[COPY1]]
     ; WAVE32: S_ENDPGM 0, implicit [[S_XOR_B32_]]
     %0:sgpr(<2 x s16>) = COPY $sgpr0
     %1:sgpr(<2 x s16>) = COPY $sgpr1
@@ -269,16 +269,16 @@ body: |
     liveins: $sgpr0_sgpr1, $sgpr2_sgpr3
     ; WAVE64-LABEL: name: xor_v2s32_sgpr_sgpr_sgpr
     ; WAVE64: liveins: $sgpr0_sgpr1, $sgpr2_sgpr3
-    ; WAVE64: [[COPY:%[0-9]+]]:sreg_64_xexec = COPY $sgpr0_sgpr1
-    ; WAVE64: [[COPY1:%[0-9]+]]:sreg_64_xexec = COPY $sgpr2_sgpr3
-    ; WAVE64: [[S_XOR_B64_:%[0-9]+]]:sreg_64_xexec = S_XOR_B64 [[COPY]], [[COPY1]]
+    ; WAVE64: [[COPY:%[0-9]+]]:sreg_64 = COPY $sgpr0_sgpr1
+    ; WAVE64: [[COPY1:%[0-9]+]]:sreg_64 = COPY $sgpr2_sgpr3
+    ; WAVE64: [[S_XOR_B64_:%[0-9]+]]:sreg_64 = S_XOR_B64 [[COPY]], [[COPY1]]
     ; WAVE64: S_ENDPGM 0, implicit [[S_XOR_B64_]]
     ; WAVE32-LABEL: name: xor_v2s32_sgpr_sgpr_sgpr
     ; WAVE32: liveins: $sgpr0_sgpr1, $sgpr2_sgpr3
     ; WAVE32: $vcc_hi = IMPLICIT_DEF
-    ; WAVE32: [[COPY:%[0-9]+]]:sreg_64_xexec = COPY $sgpr0_sgpr1
-    ; WAVE32: [[COPY1:%[0-9]+]]:sreg_64_xexec = COPY $sgpr2_sgpr3
-    ; WAVE32: [[S_XOR_B64_:%[0-9]+]]:sreg_64_xexec = S_XOR_B64 [[COPY]], [[COPY1]]
+    ; WAVE32: [[COPY:%[0-9]+]]:sreg_64 = COPY $sgpr0_sgpr1
+    ; WAVE32: [[COPY1:%[0-9]+]]:sreg_64 = COPY $sgpr2_sgpr3
+    ; WAVE32: [[S_XOR_B64_:%[0-9]+]]:sreg_64 = S_XOR_B64 [[COPY]], [[COPY1]]
     ; WAVE32: S_ENDPGM 0, implicit [[S_XOR_B64_]]
     %0:sgpr(<2 x s32>) = COPY $sgpr0_sgpr1
     %1:sgpr(<2 x s32>) = COPY $sgpr2_sgpr3
@@ -298,16 +298,16 @@ body: |
     liveins: $sgpr0_sgpr1, $sgpr2_sgpr3
     ; WAVE64-LABEL: name: xor_v4s16_sgpr_sgpr_sgpr
     ; WAVE64: liveins: $sgpr0_sgpr1, $sgpr2_sgpr3
-    ; WAVE64: [[COPY:%[0-9]+]]:sreg_64_xexec = COPY $sgpr0_sgpr1
-    ; WAVE64: [[COPY1:%[0-9]+]]:sreg_64_xexec = COPY $sgpr2_sgpr3
-    ; WAVE64: [[S_XOR_B64_:%[0-9]+]]:sreg_64_xexec = S_XOR_B64 [[COPY]], [[COPY1]]
+    ; WAVE64: [[COPY:%[0-9]+]]:sreg_64 = COPY $sgpr0_sgpr1
+    ; WAVE64: [[COPY1:%[0-9]+]]:sreg_64 = COPY $sgpr2_sgpr3
+    ; WAVE64: [[S_XOR_B64_:%[0-9]+]]:sreg_64 = S_XOR_B64 [[COPY]], [[COPY1]]
     ; WAVE64: S_ENDPGM 0, implicit [[S_XOR_B64_]]
     ; WAVE32-LABEL: name: xor_v4s16_sgpr_sgpr_sgpr
     ; WAVE32: liveins: $sgpr0_sgpr1, $sgpr2_sgpr3
     ; WAVE32: $vcc_hi = IMPLICIT_DEF
-    ; WAVE32: [[COPY:%[0-9]+]]:sreg_64_xexec = COPY $sgpr0_sgpr1
-    ; WAVE32: [[COPY1:%[0-9]+]]:sreg_64_xexec = COPY $sgpr2_sgpr3
-    ; WAVE32: [[S_XOR_B64_:%[0-9]+]]:sreg_64_xexec = S_XOR_B64 [[COPY]], [[COPY1]]
+    ; WAVE32: [[COPY:%[0-9]+]]:sreg_64 = COPY $sgpr0_sgpr1
+    ; WAVE32: [[COPY1:%[0-9]+]]:sreg_64 = COPY $sgpr2_sgpr3
+    ; WAVE32: [[S_XOR_B64_:%[0-9]+]]:sreg_64 = S_XOR_B64 [[COPY]], [[COPY1]]
     ; WAVE32: S_ENDPGM 0, implicit [[S_XOR_B64_]]
     %0:sgpr(<4 x s16>) = COPY $sgpr0_sgpr1
     %1:sgpr(<4 x s16>) = COPY $sgpr2_sgpr3
@@ -432,11 +432,11 @@ tracksRegLiveness: true
 body: |
   bb.0:
     ; WAVE64-LABEL: name: xor_s1_sgpr_undef_sgpr_undef_sgpr
-    ; WAVE64: [[S_XOR_B32_:%[0-9]+]]:sreg_32_xm0 = S_XOR_B32 undef %1:sreg_32_xm0, undef %2:sreg_32_xm0
+    ; WAVE64: [[S_XOR_B32_:%[0-9]+]]:sreg_32 = S_XOR_B32 undef %1:sreg_32, undef %2:sreg_32
     ; WAVE64: S_ENDPGM 0, implicit [[S_XOR_B32_]]
     ; WAVE32-LABEL: name: xor_s1_sgpr_undef_sgpr_undef_sgpr
     ; WAVE32: $vcc_hi = IMPLICIT_DEF
-    ; WAVE32: [[S_XOR_B32_:%[0-9]+]]:sreg_32_xm0 = S_XOR_B32 undef %1:sreg_32_xm0, undef %2:sreg_32_xm0
+    ; WAVE32: [[S_XOR_B32_:%[0-9]+]]:sreg_32 = S_XOR_B32 undef %1:sreg_32, undef %2:sreg_32
     ; WAVE32: S_ENDPGM 0, implicit [[S_XOR_B32_]]
     %2:sgpr(s1) = G_XOR undef %0:sgpr(s1), undef %1:sgpr(s1)
     S_ENDPGM 0, implicit %2
@@ -513,16 +513,14 @@ body:             |
 
     ; WAVE64-LABEL: name: copy_select_constrain_vcc_result_reg_wave32
     ; WAVE64: liveins: $vgpr0
-    ; WAVE64: [[COPY:%[0-9]+]]:vgpr(s32) = COPY $vgpr0
-    ; WAVE64: [[TRUNC:%[0-9]+]]:vgpr(s1) = G_TRUNC [[COPY]](s32)
-    ; WAVE64: [[C:%[0-9]+]]:sgpr(s1) = G_CONSTANT i1 true
-    ; WAVE64: [[C1:%[0-9]+]]:sgpr(s32) = G_CONSTANT i32 0
-    ; WAVE64: [[DEF:%[0-9]+]]:sgpr(p1) = G_IMPLICIT_DEF
-    ; WAVE64: [[COPY1:%[0-9]+]]:vcc(s1) = COPY [[TRUNC]](s1)
-    ; WAVE64: [[COPY2:%[0-9]+]]:vcc(s1) = COPY [[C]](s1)
-    ; WAVE64: [[S_XOR_B32_:%[0-9]+]]:sreg_64_xexec(s1) = S_XOR_B32 [[COPY1]](s1), [[COPY2]](s1)
-    ; WAVE64: [[COPY3:%[0-9]+]]:sreg_32_xm0(s1) = COPY [[S_XOR_B32_]](s1)
-    ; WAVE64: S_ENDPGM 0, implicit [[COPY3]](s1)
+    ; WAVE64: [[COPY:%[0-9]+]]:vgpr_32 = COPY $vgpr0
+    ; WAVE64: [[S_MOV_B32_:%[0-9]+]]:sreg_32_xm0 = S_MOV_B32 1
+    ; WAVE64: [[COPY1:%[0-9]+]]:sreg_32 = COPY [[COPY]]
+    ; WAVE64: [[COPY2:%[0-9]+]]:sreg_32 = COPY [[S_MOV_B32_]]
+    ; WAVE64: [[S_XOR_B32_:%[0-9]+]]:sreg_32 = S_XOR_B32 [[COPY1]], [[COPY2]]
+    ; WAVE64: [[COPY3:%[0-9]+]]:sreg_64_xexec = COPY [[S_XOR_B32_]]
+    ; WAVE64: [[COPY4:%[0-9]+]]:sreg_32_xm0 = COPY [[COPY3]]
+    ; WAVE64: S_ENDPGM 0, implicit [[COPY4]]
     ; WAVE32-LABEL: name: copy_select_constrain_vcc_result_reg_wave32
     ; WAVE32: liveins: $vgpr0
     ; WAVE32: $vcc_hi = IMPLICIT_DEF
@@ -566,16 +564,14 @@ body:             |
     ; WAVE64: S_ENDPGM 0, implicit [[COPY1]]
     ; WAVE32-LABEL: name: copy_select_constrain_vcc_result_reg_wave64
     ; WAVE32: liveins: $vgpr0
-    ; WAVE32: [[COPY:%[0-9]+]]:vgpr(s32) = COPY $vgpr0
-    ; WAVE32: [[TRUNC:%[0-9]+]]:vgpr(s1) = G_TRUNC [[COPY]](s32)
-    ; WAVE32: [[C:%[0-9]+]]:sgpr(s1) = G_CONSTANT i1 true
-    ; WAVE32: [[C1:%[0-9]+]]:sgpr(s32) = G_CONSTANT i32 0
-    ; WAVE32: [[DEF:%[0-9]+]]:sgpr(p1) = G_IMPLICIT_DEF
-    ; WAVE32: [[COPY1:%[0-9]+]]:vcc(s1) = COPY [[TRUNC]](s1)
-    ; WAVE32: [[COPY2:%[0-9]+]]:vcc(s1) = COPY [[C]](s1)
-    ; WAVE32: [[S_XOR_B32_:%[0-9]+]]:sreg_32_xm0_xexec(s1) = S_XOR_B32 [[COPY1]](s1), [[COPY2]](s1)
-    ; WAVE32: [[COPY3:%[0-9]+]]:sreg_64_xexec(s1) = COPY [[S_XOR_B32_]](s1)
-    ; WAVE32: S_ENDPGM 0, implicit [[COPY3]](s1)
+    ; WAVE32: $vcc_hi = IMPLICIT_DEF
+    ; WAVE32: [[COPY:%[0-9]+]]:vgpr_32 = COPY $vgpr0
+    ; WAVE32: [[S_MOV_B32_:%[0-9]+]]:sreg_32_xm0 = S_MOV_B32 1
+    ; WAVE32: [[COPY1:%[0-9]+]]:sreg_32 = COPY [[COPY]]
+    ; WAVE32: [[COPY2:%[0-9]+]]:sreg_32 = COPY [[S_MOV_B32_]]
+    ; WAVE32: [[S_XOR_B32_:%[0-9]+]]:sreg_32_xm0_xexec = S_XOR_B32 [[COPY1]], [[COPY2]]
+    ; WAVE32: [[COPY3:%[0-9]+]]:sreg_64_xexec = COPY [[S_XOR_B32_]]
+    ; WAVE32: S_ENDPGM 0, implicit [[COPY3]]
     %1:vgpr(s32) = COPY $vgpr0
     %0:vgpr(s1) = G_TRUNC %1(s32)
     %2:sgpr(s1) = G_CONSTANT i1 true
diff --git a/llvm/test/LTO/Resolution/X86/not-prevailing-weak-aliasee.ll b/llvm/test/LTO/Resolution/X86/not-prevailing-weak-aliasee.ll
index 5f484bf3bfb080..0112a3871b218c 100644
--- a/llvm/test/LTO/Resolution/X86/not-prevailing-weak-aliasee.ll
+++ b/llvm/test/LTO/Resolution/X86/not-prevailing-weak-aliasee.ll
@@ -15,7 +15,7 @@
 ; CHECK: declare hidden void @__b
 ; CHECK: declare void @b
 
-target datalayout = "e-m:e-p270:32:32-p271:32:32-p272:64:64-i64:64-f80:128-n8:16:32:64-S128"
+target datalayout = "e-m:e-i64:64-f80:128-n8:16:32:64-S128"
 target triple = "x86_64-unknown-linux-gnu"
 
 @a = hidden alias void (), void ()* @__a
diff --git a/llvm/test/ThinLTO/X86/printer.ll b/llvm/test/ThinLTO/X86/printer.ll
index ad5489658250d4..79f4f5645cd60f 100644
--- a/llvm/test/ThinLTO/X86/printer.ll
+++ b/llvm/test/ThinLTO/X86/printer.ll
@@ -4,7 +4,7 @@
 ; CHECK-BEFORE: *** IR Dump Before GlobalDCEPass ***
 ; CHECK-AFTER: *** IR Dump After GlobalDCEPass ***
 
-target datalayout = "e-m:e-p270:32:32-p271:32:32-p272:64:64-i64:64-f80:128-n8:16:32:64-S128"
+target datalayout = "e-m:e-i64:64-f80:128-n8:16:32:64-S128"
 target triple = "x86_64-unknown-linux-gnu"
 
 define i32 @foo() {
diff --git a/openmp/libomptarget/deviceRTLs/nvptx/src/omptarget-nvptx.h b/openmp/libomptarget/deviceRTLs/nvptx/src/omptarget-nvptx.h
index 9fcd1a925834e4..5519f6476556fb 100644
--- a/openmp/libomptarget/deviceRTLs/nvptx/src/omptarget-nvptx.h
+++ b/openmp/libomptarget/deviceRTLs/nvptx/src/omptarget-nvptx.h
@@ -51,16 +51,9 @@
 #ifndef CUDA_VERSION
 #error CUDA_VERSION macro is undefined, something wrong with cuda.
 #elif CUDA_VERSION >= 9000
-#define __SHFL_DOWN_SYNC(mask, var, delta, width)                              \
-  __shfl_down_sync((mask), (var), (delta), (width))
 #define __ACTIVEMASK() __activemask()
-#define __SYNCWARP(Mask) __syncwarp(Mask)
 #else
-#define __SHFL_DOWN_SYNC(mask, var, delta, width)                              \
-  __shfl_down((var), (delta), (width))
 #define __ACTIVEMASK() __ballot(1)
-// In Cuda < 9.0 no need to sync threads in warps.
-#define __SYNCWARP(Mask)
 #endif // CUDA_VERSION
 
 #define __SYNCTHREADS_N(n) asm volatile("bar.sync %0;" : : "r"(n) : "memory");
diff --git a/openmp/libomptarget/deviceRTLs/nvptx/src/reduction.cu b/openmp/libomptarget/deviceRTLs/nvptx/src/reduction.cu
index c9256383487854..e5e76d553117e0 100644
--- a/openmp/libomptarget/deviceRTLs/nvptx/src/reduction.cu
+++ b/openmp/libomptarget/deviceRTLs/nvptx/src/reduction.cu
@@ -15,6 +15,7 @@
 #include <stdio.h>
 
 #include "omptarget-nvptx.h"
+#include "target_impl.h"
 
 EXTERN
 void __kmpc_nvptx_end_reduce(int32_t global_tid) {}
@@ -23,14 +24,14 @@ EXTERN
 void __kmpc_nvptx_end_reduce_nowait(int32_t global_tid) {}
 
 EXTERN int32_t __kmpc_shuffle_int32(int32_t val, int16_t delta, int16_t size) {
-  return __SHFL_DOWN_SYNC(0xFFFFFFFF, val, delta, size);
+  return __kmpc_impl_shfl_down_sync(0xFFFFFFFF, val, delta, size);
 }
 
 EXTERN int64_t __kmpc_shuffle_int64(int64_t val, int16_t delta, int16_t size) {
    int lo, hi;
    asm volatile("mov.b64 {%0,%1}, %2;" : "=r"(lo), "=r"(hi) : "l"(val));
-   hi = __SHFL_DOWN_SYNC(0xFFFFFFFF, hi, delta, size);
-   lo = __SHFL_DOWN_SYNC(0xFFFFFFFF, lo, delta, size);
+   hi = __kmpc_impl_shfl_down_sync(0xFFFFFFFF, hi, delta, size);
+   lo = __kmpc_impl_shfl_down_sync(0xFFFFFFFF, lo, delta, size);
    asm volatile("mov.b64 %0, {%1,%2};" : "=l"(val) : "r"(lo), "r"(hi));
    return val;
 }
diff --git a/openmp/libomptarget/deviceRTLs/nvptx/src/supporti.h b/openmp/libomptarget/deviceRTLs/nvptx/src/supporti.h
index ceed7d3f7c81e7..c1a84679649405 100644
--- a/openmp/libomptarget/deviceRTLs/nvptx/src/supporti.h
+++ b/openmp/libomptarget/deviceRTLs/nvptx/src/supporti.h
@@ -14,6 +14,8 @@
 // Execution Parameters
 ////////////////////////////////////////////////////////////////////////////////
 
+#include "target_impl.h"
+
 INLINE void setExecutionParameters(ExecutionMode EMode, RuntimeMode RMode) {
   execution_param = EMode;
   execution_param |= RMode;
@@ -203,7 +205,7 @@ INLINE int IsTeamMaster(int ompThreadId) { return (ompThreadId == 0); }
 
 INLINE void IncParallelLevel(bool ActiveParallel) {
   unsigned Active = __ACTIVEMASK();
-  __SYNCWARP(Active);
+  __kmpc_impl_syncwarp(Active);
   unsigned LaneMaskLt;
   asm("mov.u32 %0, %%lanemask_lt;" : "=r"(LaneMaskLt));
   unsigned Rank = __popc(Active & LaneMaskLt);
@@ -212,12 +214,12 @@ INLINE void IncParallelLevel(bool ActiveParallel) {
         (1 + (ActiveParallel ? OMP_ACTIVE_PARALLEL_LEVEL : 0));
     __threadfence();
   }
-  __SYNCWARP(Active);
+  __kmpc_impl_syncwarp(Active);
 }
 
 INLINE void DecParallelLevel(bool ActiveParallel) {
   unsigned Active = __ACTIVEMASK();
-  __SYNCWARP(Active);
+  __kmpc_impl_syncwarp(Active);
   unsigned LaneMaskLt;
   asm("mov.u32 %0, %%lanemask_lt;" : "=r"(LaneMaskLt));
   unsigned Rank = __popc(Active & LaneMaskLt);
@@ -226,7 +228,7 @@ INLINE void DecParallelLevel(bool ActiveParallel) {
         (1 + (ActiveParallel ? OMP_ACTIVE_PARALLEL_LEVEL : 0));
     __threadfence();
   }
-  __SYNCWARP(Active);
+  __kmpc_impl_syncwarp(Active);
 }
 
 ////////////////////////////////////////////////////////////////////////////////
diff --git a/openmp/libomptarget/deviceRTLs/nvptx/src/target_impl.h b/openmp/libomptarget/deviceRTLs/nvptx/src/target_impl.h
index caa9feafe037c2..91883eaea54de1 100644
--- a/openmp/libomptarget/deviceRTLs/nvptx/src/target_impl.h
+++ b/openmp/libomptarget/deviceRTLs/nvptx/src/target_impl.h
@@ -43,6 +43,7 @@ INLINE int __kmpc_impl_popc(uint32_t x) { return __popc(x); }
 #endif
 
 // In Cuda 9.0, the *_sync() version takes an extra argument 'mask'.
+
 INLINE int32_t __kmpc_impl_shfl_sync(__kmpc_impl_lanemask_t Mask, int32_t Var,
                                      int32_t SrcLane) {
 #if CUDA_VERSION >= 9000
@@ -52,6 +53,22 @@ INLINE int32_t __kmpc_impl_shfl_sync(__kmpc_impl_lanemask_t Mask, int32_t Var,
 #endif // CUDA_VERSION
 }
 
-INLINE void __kmpc_impl_syncwarp(int32_t Mask) { __SYNCWARP(Mask); }
+INLINE int32_t __kmpc_impl_shfl_down_sync(__kmpc_impl_lanemask_t Mask,
+                                          int32_t Var, uint32_t Delta,
+                                          int32_t Width) {
+#if CUDA_VERSION >= 9000
+  return __shfl_down_sync(Mask, Var, Delta, Width);
+#else
+  return __shfl_down(Var, Delta, Width);
+#endif // CUDA_VERSION
+}
+
+INLINE void __kmpc_impl_syncwarp(__kmpc_impl_lanemask_t Mask) {
+#if CUDA_VERSION >= 9000
+  __syncwarp(Mask);
+#else
+  // In Cuda < 9.0 no need to sync threads in warps.
+#endif // CUDA_VERSION
+}
 
 #endif