diff --git a/clang/lib/CIR/CodeGen/CIRGenModule.cpp b/clang/lib/CIR/CodeGen/CIRGenModule.cpp
index c58d260e166a..412369ed07ef 100644
--- a/clang/lib/CIR/CodeGen/CIRGenModule.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenModule.cpp
@@ -515,7 +515,8 @@ void CIRGenModule::emitGlobal(GlobalDecl GD) {
   assert(!Global->hasAttr<IFuncAttr>() && "NYI");
   assert(!Global->hasAttr<CPUDispatchAttr>() && "NYI");
 
-  if (langOpts.CUDA) {
+  if (langOpts.CUDA || langOpts.HIP) {
+    // clang uses the same flag when building HIP code
     if (langOpts.CUDAIsDevice) {
       // This will implicitly mark templates and their
       // specializations as __host__ __device__.
@@ -3217,8 +3218,7 @@ void CIRGenModule::Release() {
   if (astContext.getTargetInfo().getTriple().isWasm())
     llvm_unreachable("NYI");
 
-  if (getTriple().isAMDGPU() ||
-      (getTriple().isSPIRV() && getTriple().getVendor() == llvm::Triple::AMD)) {
+  if (getTriple().isSPIRV() && getTriple().getVendor() == llvm::Triple::AMD) {
     llvm_unreachable("NYI");
   }
 
@@ -3229,9 +3229,7 @@ void CIRGenModule::Release() {
   if (!astContext.CUDAExternalDeviceDeclODRUsedByHost.empty()) {
     llvm_unreachable("NYI");
   }
-  if (langOpts.HIP && !getLangOpts().OffloadingNewDriver) {
-    llvm_unreachable("NYI");
-  }
+
   assert(!MissingFeatures::emitLLVMUsed());
   assert(!MissingFeatures::sanStats());
 
diff --git a/clang/lib/CIR/CodeGen/TargetInfo.cpp b/clang/lib/CIR/CodeGen/TargetInfo.cpp
index 7669dad59eb8..07dca811985e 100644
--- a/clang/lib/CIR/CodeGen/TargetInfo.cpp
+++ b/clang/lib/CIR/CodeGen/TargetInfo.cpp
@@ -329,6 +329,30 @@ class NVPTXTargetCIRGenInfo : public TargetCIRGenInfo {
 
 } // namespace
 
+//===----------------------------------------------------------------------===//
+// AMDGPU ABI Implementation
+//===----------------------------------------------------------------------===//
+
+namespace {
+
+class AMDGPUABIInfo : public ABIInfo {
+public:
+  AMDGPUABIInfo(CIRGenTypes &cgt) : ABIInfo(cgt) {}
+
+  cir::ABIArgInfo classifyReturnType(QualType retTy) const;
+  cir::ABIArgInfo classifyArgumentType(QualType ty) const;
+
+  void computeInfo(CIRGenFunctionInfo &fnInfo) const override;
+};
+
+class AMDGPUTargetCIRGenInfo : public TargetCIRGenInfo {
+public:
+  AMDGPUTargetCIRGenInfo(CIRGenTypes &cgt)
+      : TargetCIRGenInfo(std::make_unique<AMDGPUABIInfo>(cgt)) {}
+};
+
+} // namespace
+
 // TODO(cir): remove the attribute once this gets used.
 LLVM_ATTRIBUTE_UNUSED
 static bool classifyReturnType(const CIRGenCXXABI &CXXABI,
@@ -495,6 +519,34 @@ void NVPTXABIInfo::computeInfo(CIRGenFunctionInfo &fnInfo) const {
     fnInfo.getReturnInfo() = cir::ABIArgInfo::getDirect(CGT.convertType(retTy));
 }
 
+// Skeleton only. Implement when used in TargetLower stage.
+cir::ABIArgInfo AMDGPUABIInfo::classifyReturnType(QualType retTy) const {
+  llvm_unreachable("not yet implemented");
+}
+
+cir::ABIArgInfo AMDGPUABIInfo::classifyArgumentType(QualType ty) const {
+  llvm_unreachable("not yet implemented");
+}
+
+void AMDGPUABIInfo::computeInfo(CIRGenFunctionInfo &fnInfo) const {
+  // Top level CIR has unlimited arguments and return types. Lowering for ABI
+  // specific concerns should happen during a lowering phase. Assume everything
+  // is direct for now.
+  for (CIRGenFunctionInfo::arg_iterator it = fnInfo.arg_begin(),
+                                        ie = fnInfo.arg_end();
+       it != ie; ++it) {
+    if (testIfIsVoidTy(it->type))
+      it->info = cir::ABIArgInfo::getIgnore();
+    else
+      it->info = cir::ABIArgInfo::getDirect(CGT.convertType(it->type));
+  }
+  auto retTy = fnInfo.getReturnType();
+  if (testIfIsVoidTy(retTy))
+    fnInfo.getReturnInfo() = cir::ABIArgInfo::getIgnore();
+  else
+    fnInfo.getReturnInfo() = cir::ABIArgInfo::getDirect(CGT.convertType(retTy));
+}
+
 ABIInfo::~ABIInfo() {}
 
 bool ABIInfo::isPromotableIntegerTypeForABI(QualType Ty) const {
@@ -690,5 +742,9 @@ const TargetCIRGenInfo &CIRGenModule::getTargetCIRGenInfo() {
   case llvm::Triple::nvptx64: {
     return SetCIRGenInfo(new NVPTXTargetCIRGenInfo(genTypes));
   }
+
+  case llvm::Triple::amdgcn: {
+    return SetCIRGenInfo(new AMDGPUTargetCIRGenInfo(genTypes));
+  }
   }
 }
diff --git a/clang/test/CIR/CodeGen/HIP/simple.cpp b/clang/test/CIR/CodeGen/HIP/simple.cpp
new file mode 100644
index 000000000000..ec4110da10d7
--- /dev/null
+++ b/clang/test/CIR/CodeGen/HIP/simple.cpp
@@ -0,0 +1,16 @@
+#include "../Inputs/cuda.h"
+
+// RUN: %clang_cc1 -triple=amdgcn-amd-amdhsa -x hip -fclangir \
+// RUN:            -emit-cir %s -o %t.cir
+// RUN: FileCheck --check-prefix=CIR --input-file=%t.cir %s
+
+
+// This should emit as a normal C++ function.
+__host__ void host_fn(int *a, int *b, int *c) {}
+
+// CIR: cir.func @_Z7host_fnPiS_S_
+
+// This shouldn't emit.
+__device__ void device_fn(int* a, double b, float c) {}
+
+// CHECK-NOT: cir.func @_Z9device_fnPidf