diff --git a/clang/include/clang/Driver/CommonArgs.h b/clang/include/clang/Driver/CommonArgs.h
index 8c861df793311..ad1912247e01f 100644
--- a/clang/include/clang/Driver/CommonArgs.h
+++ b/clang/include/clang/Driver/CommonArgs.h
@@ -144,6 +144,12 @@ void addArchSpecificRPath(const ToolChain &TC, const llvm::opt::ArgList &Args,
 void addOpenMPRuntimeLibraryPath(const ToolChain &TC,
                                  const llvm::opt::ArgList &Args,
                                  llvm::opt::ArgStringList &CmdArgs);
+
+bool addLLVMOffloadingRuntime(const Compilation &C,
+                              llvm::opt::ArgStringList &CmdArgs,
+                              const ToolChain &TC,
+                              const llvm::opt::ArgList &Args);
+
 /// Returns true, if an OpenMP runtime has been added.
 bool addOpenMPRuntime(const Compilation &C, llvm::opt::ArgStringList &CmdArgs,
                       const ToolChain &TC, const llvm::opt::ArgList &Args,
diff --git a/clang/lib/CodeGen/CGCUDANV.cpp b/clang/lib/CodeGen/CGCUDANV.cpp
index 416ed935c1b30..0ea3ed36fae83 100644
--- a/clang/lib/CodeGen/CGCUDANV.cpp
+++ b/clang/lib/CodeGen/CGCUDANV.cpp
@@ -254,9 +254,7 @@ CGNVCUDARuntime::CGNVCUDARuntime(CodeGenModule &CGM)
   VoidTy = CGM.VoidTy;
   PtrTy = CGM.DefaultPtrTy;
 
-  if (CGM.getLangOpts().OffloadViaLLVM)
-    Prefix = "llvm";
-  else if (CGM.getLangOpts().HIP)
+  if (CGM.getLangOpts().HIP)
     Prefix = "hip";
   else
     Prefix = "cuda";
@@ -345,41 +343,52 @@ void CGNVCUDARuntime::emitDeviceStub(CodeGenFunction &CGF,
     emitDeviceStubBodyLegacy(CGF, Args);
 }
 
-/// Build the input as a sized array of pointers so that it can be launched by
-/// the offloading runtime.
+/// CUDA passes the arguments with a level of indirection. For example, a
+/// (void*, short, void*) is passed as {void **, short *, void **} to the launch
+/// function. For the LLVM/Offload launch we include the number of arguments and
+/// their size. Thus, we pass {{void **, short*, void **}, 3, {sizeof(void*),
+/// sizeof(short), sizeof(void*)}}.
 Address CGNVCUDARuntime::prepareKernelArgsLLVMOffload(CodeGenFunction &CGF,
                                                       FunctionArgList &Args) {
-  SmallVector<llvm::Type *> ArgTypes, KernelLaunchParamsTypes;
-  for (auto &Arg : Args)
-    ArgTypes.push_back(CGF.ConvertTypeForMem(Arg->getType()));
-  llvm::StructType *KernelArgsTy = llvm::StructType::create(ArgTypes);
-  llvm::Type *KernelArgsPtrsTy = llvm::ArrayType::get(PtrTy, Args.size());
-
-  auto *Int32Ty = CGF.Builder.getInt32Ty();
-  KernelLaunchParamsTypes.push_back(Int32Ty);
+  SmallVector<llvm::Type *> KernelLaunchParamsTypes;
+
+  auto *Int64Ty = CGF.Builder.getInt64Ty();
+  KernelLaunchParamsTypes.push_back(PtrTy);
+  KernelLaunchParamsTypes.push_back(Int64Ty);
   KernelLaunchParamsTypes.push_back(PtrTy);
 
   llvm::StructType *KernelLaunchParamsTy =
       llvm::StructType::create(KernelLaunchParamsTypes);
-  Address KernelArgs = CGF.CreateTempAllocaWithoutCast(
-      KernelArgsTy, CharUnits::fromQuantity(16), "kernel_args");
-  Address KernelArgsPtrs = CGF.CreateTempAllocaWithoutCast(
-      KernelArgsPtrsTy, CharUnits::fromQuantity(16), "kernel_args_ptrs");
   Address KernelLaunchParams = CGF.CreateTempAllocaWithoutCast(
       KernelLaunchParamsTy, CharUnits::fromQuantity(16),
       "kernel_launch_params");
+  Address KernelArgs = CGF.CreateTempAlloca(
+      PtrTy, LangAS::Default, CharUnits::fromQuantity(16), "kernel_args",
+      llvm::ConstantInt::get(SizeTy, std::max<size_t>(1, Args.size())));
+  Address KernelArgSizes = CGF.CreateTempAlloca(
+      SizeTy, LangAS::Default, CharUnits::fromQuantity(16), "kernel_arg_sizes",
+      llvm::ConstantInt::get(SizeTy, std::max<size_t>(1, Args.size())));
 
-  CGF.Builder.CreateStore(llvm::ConstantInt::get(Int32Ty, Args.size()),
+  CGF.Builder.CreateStore(KernelArgs.emitRawPointer(CGF),
                           CGF.Builder.CreateStructGEP(KernelLaunchParams, 0));
-  CGF.Builder.CreateStore(KernelArgsPtrs.emitRawPointer(CGF),
+  CGF.Builder.CreateStore(llvm::ConstantInt::get(Int64Ty, Args.size()),
                           CGF.Builder.CreateStructGEP(KernelLaunchParams, 1));
+  CGF.Builder.CreateStore(KernelArgSizes.emitRawPointer(CGF),
+                          CGF.Builder.CreateStructGEP(KernelLaunchParams, 2));
 
   for (unsigned i = 0; i < Args.size(); ++i) {
-    auto *ArgVal = CGF.Builder.CreateLoad(CGF.GetAddrOfLocalVar(Args[i]));
-    Address ArgAddr = CGF.Builder.CreateStructGEP(KernelArgs, i);
-    CGF.Builder.CreateStore(ArgVal, ArgAddr);
-    CGF.Builder.CreateStore(ArgAddr.emitRawPointer(CGF),
-                            CGF.Builder.CreateConstArrayGEP(KernelArgsPtrs, i));
+    llvm::Value *VarPtr = CGF.GetAddrOfLocalVar(Args[i]).emitRawPointer(CGF);
+    llvm::Value *VoidVarPtr = CGF.Builder.CreatePointerCast(VarPtr, PtrTy);
+    CGF.Builder.CreateDefaultAlignedStore(
+        VoidVarPtr, CGF.Builder.CreateConstGEP1_32(
+                        PtrTy, KernelArgs.emitRawPointer(CGF), i));
+
+    auto ArgSize = CGM.getDataLayout().getTypeAllocSize(
+        CGM.getTypes().ConvertType(Args[i]->getType()));
+    CGF.Builder.CreateDefaultAlignedStore(
+        llvm::ConstantInt::get(SizeTy, ArgSize),
+        CGF.Builder.CreateConstGEP1_32(SizeTy,
+                                       KernelArgSizes.emitRawPointer(CGF), i));
   }
 
   return KernelLaunchParams;
@@ -408,8 +417,9 @@ Address CGNVCUDARuntime::prepareKernelArgs(CodeGenFunction &CGF,
 // array and kernels are launched using cudaLaunchKernel().
 void CGNVCUDARuntime::emitDeviceStubBodyNew(CodeGenFunction &CGF,
                                             FunctionArgList &Args) {
+  bool UsesLLVMOffloading = CGF.getLangOpts().OffloadViaLLVM;
   // Build the shadow stack entry at the very start of the function.
-  Address KernelArgs = CGF.getLangOpts().OffloadViaLLVM
+  Address KernelArgs = UsesLLVMOffloading
                            ? prepareKernelArgsLLVMOffload(CGF, Args)
                            : prepareKernelArgs(CGF, Args);
 
@@ -435,7 +445,9 @@ void CGNVCUDARuntime::emitDeviceStubBodyNew(CodeGenFunction &CGF,
     else if (CGF.getLangOpts().CUDA)
       KernelLaunchAPI = KernelLaunchAPI + "_ptsz";
   }
-  auto LaunchKernelName = addPrefixToName(KernelLaunchAPI);
+  /// Use __llvmLaunchKernel for LLVMOffload.
+  auto LaunchKernelName = UsesLLVMOffloading ? "__llvm" + KernelLaunchAPI
+                                             : addPrefixToName(KernelLaunchAPI);
   const IdentifierInfo &cudaLaunchKernelII =
       CGM.getContext().Idents.get(LaunchKernelName);
   FunctionDecl *cudaLaunchKernelFD = nullptr;
@@ -1282,9 +1294,6 @@ void CGNVCUDARuntime::createOffloadingEntries() {
   llvm::object::OffloadKind Kind = CGM.getLangOpts().HIP
                                        ? llvm::object::OffloadKind::OFK_HIP
                                        : llvm::object::OffloadKind::OFK_Cuda;
-  // For now, just spoof this as OpenMP because that's the runtime it uses.
-  if (CGM.getLangOpts().OffloadViaLLVM)
-    Kind = llvm::object::OffloadKind::OFK_OpenMP;
 
   llvm::Module &M = CGM.getModule();
   for (KernelInfo &I : EmittedKernels)
diff --git a/clang/lib/Driver/Driver.cpp b/clang/lib/Driver/Driver.cpp
index 38795f7c2ae7a..f27557c09fca3 100644
--- a/clang/lib/Driver/Driver.cpp
+++ b/clang/lib/Driver/Driver.cpp
@@ -905,10 +905,15 @@ getSystemOffloadArchs(Compilation &C, Action::OffloadKind Kind) {
   if (llvm::ErrorOr<std::string> Executable =
           llvm::sys::findProgramByName(Program, {C.getDriver().Dir})) {
     llvm::SmallVector<StringRef> Args{*Executable};
-    if (Kind == Action::OFK_HIP)
-      Args.push_back("--only=amdgpu");
-    else if (Kind == Action::OFK_Cuda)
-      Args.push_back("--only=nvptx");
+    bool UsesLLVMOffloading =
+        C.getArgs().hasFlag(options::OPT_foffload_via_llvm,
+                            options::OPT_fno_offload_via_llvm, false);
+    if (!UsesLLVMOffloading) {
+      if (Kind == Action::OFK_HIP)
+        Args.push_back("--only=amdgpu");
+      else if (Kind == Action::OFK_Cuda)
+        Args.push_back("--only=nvptx");
+    }
     auto StdoutOrErr = C.getDriver().executeProgram(Args);
 
     if (!StdoutOrErr) {
@@ -965,15 +970,20 @@ static TripleSet inferOffloadToolchains(Compilation &C,
       ID = StringToOffloadArch(
           getProcessorFromTargetID(llvm::Triple("amdgcn-amd-amdhsa"), Arch));
 
-    if (Kind == Action::OFK_HIP && !IsAMDOffloadArch(ID)) {
-      C.getDriver().Diag(clang::diag::err_drv_offload_bad_gpu_arch)
-          << "HIP" << Arch;
-      return {};
-    }
-    if (Kind == Action::OFK_Cuda && !IsNVIDIAOffloadArch(ID)) {
-      C.getDriver().Diag(clang::diag::err_drv_offload_bad_gpu_arch)
-          << "CUDA" << Arch;
-      return {};
+    bool UsesLLVMOffloading =
+        C.getArgs().hasFlag(options::OPT_foffload_via_llvm,
+                            options::OPT_fno_offload_via_llvm, false);
+    if (!UsesLLVMOffloading) {
+      if (Kind == Action::OFK_HIP && !IsAMDOffloadArch(ID)) {
+        C.getDriver().Diag(clang::diag::err_drv_offload_bad_gpu_arch)
+            << "HIP" << Arch;
+        return {};
+      }
+      if (Kind == Action::OFK_Cuda && !IsNVIDIAOffloadArch(ID)) {
+        C.getDriver().Diag(clang::diag::err_drv_offload_bad_gpu_arch)
+            << "CUDA" << Arch;
+        return {};
+      }
     }
     if (Kind == Action::OFK_OpenMP &&
         (ID == OffloadArch::Unknown || ID == OffloadArch::Unused)) {
@@ -989,6 +999,8 @@ static TripleSet inferOffloadToolchains(Compilation &C,
 
     llvm::Triple Triple =
         OffloadArchToTriple(C.getDefaultToolChain().getTriple(), ID);
+    if (UsesLLVMOffloading)
+      Triple.setEnvironment(llvm::Triple::LLVM);
 
     // Make a new argument that dispatches this argument to the appropriate
     // toolchain. This is required when we infer it and create potentially
@@ -1032,32 +1044,30 @@ static TripleSet inferOffloadToolchains(Compilation &C,
 
 void Driver::CreateOffloadingDeviceToolChains(Compilation &C,
                                               InputList &Inputs) {
-  bool UseLLVMOffload = C.getInputArgs().hasArg(
-      options::OPT_foffload_via_llvm, options::OPT_fno_offload_via_llvm, false);
   bool IsCuda =
-      llvm::any_of(Inputs,
-                   [](std::pair<types::ID, const llvm::opt::Arg *> &I) {
-                     return types::isCuda(I.first);
-                   }) &&
-      !UseLLVMOffload;
+      llvm::any_of(Inputs, [](std::pair<types::ID, const llvm::opt::Arg *> &I) {
+        return types::isCuda(I.first);
+      });
   bool IsHIP =
       (llvm::any_of(Inputs,
                     [](std::pair<types::ID, const llvm::opt::Arg *> &I) {
                       return types::isHIP(I.first);
                     }) ||
        C.getInputArgs().hasArg(options::OPT_hip_link) ||
-       C.getInputArgs().hasArg(options::OPT_hipstdpar)) &&
-      !UseLLVMOffload;
+       C.getInputArgs().hasArg(options::OPT_hipstdpar));
   bool IsSYCL = C.getInputArgs().hasFlag(options::OPT_fsycl,
                                          options::OPT_fno_sycl, false);
   bool IsOpenMPOffloading =
-      UseLLVMOffload ||
       (C.getInputArgs().hasFlag(options::OPT_fopenmp, options::OPT_fopenmp_EQ,
                                 options::OPT_fno_openmp, false) &&
        (C.getInputArgs().hasArg(options::OPT_offload_targets_EQ) ||
         (C.getInputArgs().hasArg(options::OPT_offload_arch_EQ) &&
          !(IsCuda || IsHIP))));
 
+  // We currently don't support any kind of mixed offloading.
+  if (IsOpenMPOffloading)
+    IsCuda = IsHIP = IsSYCL = false;
+
   llvm::SmallSet<Action::OffloadKind, 4> Kinds;
   const std::pair<bool, Action::OffloadKind> ActiveKinds[] = {
       {IsCuda, Action::OFK_Cuda},
@@ -1143,7 +1153,7 @@ void Driver::CreateOffloadingDeviceToolChains(Compilation &C,
                                      C.getDefaultToolChain().getTriple());
 
       // Emit a warning if the detected CUDA version is too new.
-      if (Kind == Action::OFK_Cuda) {
+      if (Kind == Action::OFK_Cuda && Target.getOS() == llvm::Triple::CUDA) {
         auto &CudaInstallation =
             static_cast<const toolchains::CudaToolChain &>(TC).CudaInstallation;
         if (CudaInstallation.isValid())
@@ -5069,6 +5079,9 @@ Driver::BuildOffloadingActions(Compilation &C, llvm::opt::DerivedArgList &Args,
         getFinalPhase(Args) == phases::Preprocess))
     return HostAction;
 
+  bool UsesLLVMOffloading = Args.hasArg(
+      options::OPT_foffload_via_llvm, options::OPT_fno_offload_via_llvm, false);
+
   ActionList OffloadActions;
   OffloadAction::DeviceDependences DDeps;
 
@@ -5089,7 +5102,6 @@ Driver::BuildOffloadingActions(Compilation &C, llvm::opt::DerivedArgList &Args,
     types::ID InputType = Input.first;
     const Arg *InputArg = Input.second;
 
-    // The toolchain can be active for unsupported file types.
     if ((Kind == Action::OFK_Cuda && !types::isCuda(InputType)) ||
         (Kind == Action::OFK_HIP && !types::isHIP(InputType)))
       continue;
@@ -5184,9 +5196,12 @@ Driver::BuildOffloadingActions(Compilation &C, llvm::opt::DerivedArgList &Args,
       OffloadAction::DeviceDependences DDep;
       DDep.add(*A, *TCAndArch->first, TCAndArch->second, Kind);
 
-      // Compiling CUDA in non-RDC mode uses the PTX output if available.
+      // The legacy CUDA fatbinary path can include PTX alongside the cubin.
+      // The LLVM offload wrapper path feeds these images through a device
+      // linker first, and clang-nvlink-wrapper does not accept PTX as input.
       for (Action *Input : A->getInputs())
-        if (Kind == Action::OFK_Cuda && A->getType() == types::TY_Object &&
+        if (!UsesLLVMOffloading && Kind == Action::OFK_Cuda &&
+            A->getType() == types::TY_Object &&
             !Args.hasFlag(options::OPT_fgpu_rdc, options::OPT_fno_gpu_rdc,
                           false))
           DDep.add(*Input, *TCAndArch->first, TCAndArch->second, Kind);
@@ -5214,7 +5229,7 @@ Driver::BuildOffloadingActions(Compilation &C, llvm::opt::DerivedArgList &Args,
     return HostAction;
 
   OffloadAction::DeviceDependences DDep;
-  if (C.isOffloadingHostKind(Action::OFK_Cuda) &&
+  if (!UsesLLVMOffloading && C.isOffloadingHostKind(Action::OFK_Cuda) &&
       !Args.hasFlag(options::OPT_fgpu_rdc, options::OPT_fno_gpu_rdc, false)) {
     // If we are not in RDC-mode we just emit the final CUDA fatbinary for
     // each translation unit without requiring any linking.
@@ -5222,7 +5237,7 @@ Driver::BuildOffloadingActions(Compilation &C, llvm::opt::DerivedArgList &Args,
         C.MakeAction<LinkJobAction>(OffloadActions, types::TY_CUDA_FATBIN);
     DDep.add(*FatbinAction, *C.getSingleOffloadToolChain<Action::OFK_Cuda>(),
              /*BA=*/{}, Action::OFK_Cuda);
-  } else if (HIPNoRDC && offloadDeviceOnly()) {
+  } else if (!UsesLLVMOffloading && HIPNoRDC && offloadDeviceOnly()) {
     // If we are in device-only non-RDC-mode we just emit the final HIP
     // fatbinary for each translation unit, linking each input individually.
     Action *FatbinAction =
@@ -5230,7 +5245,7 @@ Driver::BuildOffloadingActions(Compilation &C, llvm::opt::DerivedArgList &Args,
     DDep.add(*FatbinAction,
              *C.getOffloadToolChains<Action::OFK_HIP>().first->second,
              /*BA=*/{}, Action::OFK_HIP);
-  } else if (HIPNoRDC) {
+  } else if (!UsesLLVMOffloading && HIPNoRDC) {
     // Host + device assembly: defer to clang-offload-bundler (see
     // BuildActions).
     if (HIPAsmBundleDeviceOut &&
@@ -7091,7 +7106,8 @@ const ToolChain &Driver::getOffloadToolChain(
       // For AMDHSA offloading (HIP, OpenMP), use the unified AMDGPUToolChain
       // This handles both amdgpu-amd-amdhsa and spirv64-amd-amdhsa
       // FIXME: This should not key off language or OS.
-      if (Kind == Action::OFK_HIP || Kind == Action::OFK_OpenMP)
+      if (Kind == Action::OFK_HIP || Kind == Action::OFK_OpenMP ||
+          Kind == Action::OFK_Cuda)
         TC = std::make_unique<toolchains::AMDGPUToolChain>(*this, Target, Args,
                                                            HostTC.get(), Kind);
       break;
diff --git a/clang/lib/Driver/ToolChains/AMDGPU.cpp b/clang/lib/Driver/ToolChains/AMDGPU.cpp
index 7bce060de0596..5c30417365b94 100644
--- a/clang/lib/Driver/ToolChains/AMDGPU.cpp
+++ b/clang/lib/Driver/ToolChains/AMDGPU.cpp
@@ -515,6 +515,24 @@ void RocmInstallationDetector::AddHIPIncludeArgs(const ArgList &DriverArgs,
                             !DriverArgs.hasArg(options::OPT_nohipwrapperinc);
   bool HasHipStdPar = DriverArgs.hasArg(options::OPT_hipstdpar);
 
+  if (DriverArgs.hasFlag(options::OPT_foffload_via_llvm,
+                         options::OPT_fno_offload_via_llvm, false)) {
+    if (DriverArgs.hasFlag(options::OPT_offload_inc,
+                           options::OPT_no_offload_inc, true) &&
+        !DriverArgs.hasArg(options::OPT_nohipwrapperinc) &&
+        !DriverArgs.hasArg(options::OPT_nobuiltininc)) {
+      CC1Args.append({"-include", "__clang_gpu_device_functions.h"});
+
+      SmallString<128> HIPIncludePath(D.ResourceDir);
+      llvm::sys::path::append(HIPIncludePath, "..", "..", "..");
+      llvm::sys::path::append(HIPIncludePath, "include", "offload");
+      CC1Args.push_back("-internal-isystem");
+      CC1Args.push_back(DriverArgs.MakeArgString(HIPIncludePath));
+      CC1Args.append({"-include", "hip/hip_runtime.h"});
+    }
+    return;
+  }
+
   if (!DriverArgs.hasArg(options::OPT_nobuiltininc)) {
     // HIP header includes standard library wrapper headers under clang
     // cuda_wrappers directory. Since these wrapper headers include_next
@@ -699,7 +717,11 @@ AMDGPUToolChain::AMDGPUToolChain(const Driver &D, const llvm::Triple &Triple,
     : Generic_ELF(D, Triple, Args),
       OptionsDefault(
           {{options::OPT_O, "3"}, {options::OPT_cl_std_EQ, "CL1.2"}}),
-      HostTC(HostTC_), UseHIPLinker(Kind == Action::OFK_HIP),
+      HostTC(HostTC_),
+      UseHIPLinker(Kind == Action::OFK_HIP ||
+                   (Kind == Action::OFK_Cuda &&
+                    Args.hasFlag(options::OPT_foffload_via_llvm,
+                                 options::OPT_fno_offload_via_llvm, false))),
       ShouldLinkDeviceLibs(ShouldLinkDeviceLibs) {
   loadMultilibsFromYAML(Args, D);
 
@@ -709,8 +731,10 @@ AMDGPUToolChain::AMDGPUToolChain(const Driver &D, const llvm::Triple &Triple,
   // each tool invocation.
   checkAMDGPUCodeObjectVersion(D, Args);
 
+  bool UsesLLVMOffloading = Args.hasFlag(
+      options::OPT_foffload_via_llvm, options::OPT_fno_offload_via_llvm, false);
   if (Triple.getOS() == llvm::Triple::AMDHSA &&
-      Triple.getEnvironment() != llvm::Triple::LLVM)
+      Triple.getEnvironment() != llvm::Triple::LLVM && !UsesLLVMOffloading)
     RocmInstallation->detectDeviceLibrary();
 
   if (HostTC)
@@ -889,7 +913,10 @@ bool AMDGPUToolChain::isWave64(const llvm::opt::ArgList &DriverArgs,
 void AMDGPUToolChain::addClangTargetOptions(
     const llvm::opt::ArgList &DriverArgs, llvm::opt::ArgStringList &CC1Args,
     BoundArch BA, Action::OffloadKind DeviceOffloadingKind) const {
-  if (DeviceOffloadingKind == Action::OFK_HIP) {
+  bool UsesLLVMOffloading = DriverArgs.hasFlag(
+      options::OPT_foffload_via_llvm, options::OPT_fno_offload_via_llvm, false);
+  if (DeviceOffloadingKind == Action::OFK_HIP ||
+      (DeviceOffloadingKind == Action::OFK_Cuda && UsesLLVMOffloading)) {
     CC1Args.append({"-fcuda-is-device", "-fno-threadsafe-statics"});
 
     if (!DriverArgs.hasFlag(options::OPT_fgpu_rdc, options::OPT_fno_gpu_rdc,
diff --git a/clang/lib/Driver/ToolChains/Clang.cpp b/clang/lib/Driver/ToolChains/Clang.cpp
index 94f9a26aac39f..035bc3f5b4273 100644
--- a/clang/lib/Driver/ToolChains/Clang.cpp
+++ b/clang/lib/Driver/ToolChains/Clang.cpp
@@ -53,6 +53,7 @@
 #include "llvm/Support/Path.h"
 #include "llvm/Support/Process.h"
 #include "llvm/Support/YAMLParser.h"
+#include "llvm/Support/raw_ostream.h"
 #include "llvm/TargetParser/AArch64TargetParser.h"
 #include "llvm/TargetParser/ARMTargetParserCommon.h"
 #include "llvm/TargetParser/Host.h"
@@ -952,9 +953,12 @@ void Clang::AddPreprocessingOptions(Compilation &C, const JobAction &JA,
   // before we -I or -include anything else, because we must pick up the
   // CUDA/HIP/SYCL headers from the particular CUDA/ROCm/SYCL installation,
   // rather than from e.g. /usr/local/include.
-  if (JA.isOffloading(Action::OFK_Cuda))
+  bool UsesLLVMOffloading = Args.hasFlag(
+      options::OPT_foffload_via_llvm, options::OPT_fno_offload_via_llvm, false);
+  if (JA.isOffloading(Action::OFK_Cuda) && !UsesLLVMOffloading) {
     getToolChain().AddCudaIncludeArgs(Args, CmdArgs);
-  if (JA.isOffloading(Action::OFK_HIP))
+  }
+  if (JA.isOffloading(Action::OFK_HIP) && !UsesLLVMOffloading)
     getToolChain().AddHIPIncludeArgs(Args, CmdArgs);
   if (JA.isOffloading(Action::OFK_SYCL))
     getToolChain().addSYCLIncludeArgs(Args, CmdArgs);
@@ -979,17 +983,35 @@ void Clang::AddPreprocessingOptions(Compilation &C, const JobAction &JA,
     CmdArgs.push_back("-include");
     CmdArgs.push_back("__clang_openmp_device_functions.h");
   }
+  bool isCudaInput = llvm::any_of(
+      Inputs, [](const InputInfo &I) { return types::isCuda(I.getType()); });
+  if (UsesLLVMOffloading && JA.isOffloading(Action::OFK_Cuda) &&
+      Args.hasFlag(options::OPT_offload_inc, options::OPT_no_offload_inc,
+                   true) &&
+      !Args.hasArg(options::OPT_nobuiltininc) && isCudaInput) {
+    CmdArgs.append({"-include", "__clang_gpu_device_functions.h"});
+
+    SmallString<128> OffloadCudaInclude(D.Dir);
+    llvm::sys::path::append(OffloadCudaInclude, "..", "include", "offload",
+                            "cuda");
+    CmdArgs.append({"-internal-isystem", Args.MakeArgString(OffloadCudaInclude),
+                    "-include"});
+    CmdArgs.push_back("cuda_runtime.h");
+  }
+  bool isHIPInput = llvm::any_of(
+      Inputs, [](const InputInfo &I) { return types::isHIP(I.getType()); });
+  if (UsesLLVMOffloading && JA.isOffloading(Action::OFK_HIP) &&
+      Args.hasFlag(options::OPT_offload_inc, options::OPT_no_offload_inc,
+                   true) &&
+      !Args.hasArg(options::OPT_nohipwrapperinc) &&
+      !Args.hasArg(options::OPT_nobuiltininc) && isHIPInput) {
+    CmdArgs.append({"-include", "__clang_gpu_device_functions.h"});
 
-  if (Args.hasArg(options::OPT_foffload_via_llvm)) {
-    // Add llvm_wrappers/* to our system include path.  This lets us wrap
-    // standard library headers and other headers.
-    SmallString<128> P(D.ResourceDir);
-    llvm::sys::path::append(P, "include", "llvm_offload_wrappers");
-    CmdArgs.append({"-internal-isystem", Args.MakeArgString(P), "-include"});
-    if (JA.isDeviceOffloading(Action::OFK_OpenMP))
-      CmdArgs.push_back("__llvm_offload_device.h");
-    else
-      CmdArgs.push_back("__llvm_offload_host.h");
+    SmallString<128> OffloadHIPInclude(D.Dir);
+    llvm::sys::path::append(OffloadHIPInclude, "..", "include", "offload");
+    CmdArgs.append({"-internal-isystem", Args.MakeArgString(OffloadHIPInclude),
+                    "-include"});
+    CmdArgs.push_back("hip/hip_runtime.h");
   }
 
   // Add -i* options, and automatically translate to
@@ -1171,7 +1193,7 @@ void Clang::AddPreprocessingOptions(Compilation &C, const JobAction &JA,
       Args.hasFlag(options::OPT_offload_inc, options::OPT_no_offload_inc,
                    true) &&
       !Args.hasArg(options::OPT_nobuiltininc) &&
-      (C.getActiveOffloadKinds() == Action::OFK_OpenMP)) {
+      JA.isDeviceOffloading(Action::OFK_OpenMP)) {
     // TODO: CUDA / HIP include their own headers for some common functions
     // implemented here. We'll need to clean those up so they do not conflict.
     SmallString<128> P(D.ResourceDir);
@@ -5194,6 +5216,8 @@ void Clang::ConstructJob(Compilation &C, const JobAction &JA,
   bool IsSYCLDevice = JA.isDeviceOffloading(Action::OFK_SYCL);
   bool IsOpenMPDevice = JA.isDeviceOffloading(Action::OFK_OpenMP);
   bool IsExtractAPI = isa<ExtractAPIJobAction>(JA);
+  bool UsesLLVMOffloading = Args.hasFlag(
+      options::OPT_foffload_via_llvm, options::OPT_fno_offload_via_llvm, false);
   bool IsDeviceOffloadAction = !(JA.isDeviceOffloading(Action::OFK_None) ||
                                  JA.isDeviceOffloading(Action::OFK_Host));
   bool IsHostOffloadingAction =
@@ -5312,7 +5336,7 @@ void Clang::ConstructJob(Compilation &C, const JobAction &JA,
     }
   }
 
-  if (IsCuda && !IsCudaDevice) {
+  if (IsCuda && !IsCudaDevice && !UsesLLVMOffloading) {
     // We need to figure out which CUDA version we're compiling for, as that
     // determines how we load and launch GPU kernels.
     auto *CTC = static_cast<const toolchains::CudaToolChain *>(
@@ -8313,11 +8337,11 @@ void Clang::ConstructJob(Compilation &C, const JobAction &JA,
   // Host-side offloading compilation receives all device-side outputs. Include
   // them in the host compilation depending on the target. If the host inputs
   // are not empty we use the new-driver scheme, otherwise use the old scheme.
-  if ((IsCuda || IsHIP) && CudaDeviceInput) {
+  if ((IsCuda || IsHIP) && !UsesLLVMOffloading && CudaDeviceInput) {
     CmdArgs.push_back("-fcuda-include-gpubinary");
     CmdArgs.push_back(CudaDeviceInput->getFilename());
   } else if (!HostOffloadingInputs.empty()) {
-    if ((IsCuda || IsHIP) && !IsRDCMode) {
+    if ((IsCuda || IsHIP) && !UsesLLVMOffloading && !IsRDCMode) {
       assert(HostOffloadingInputs.size() == 1 && "Only one input expected");
       CmdArgs.push_back("-fcuda-include-gpubinary");
       CmdArgs.push_back(HostOffloadingInputs.front().getFilename());
diff --git a/clang/lib/Driver/ToolChains/CommonArgs.cpp b/clang/lib/Driver/ToolChains/CommonArgs.cpp
index a2b99dffc383e..37a4392f2941f 100644
--- a/clang/lib/Driver/ToolChains/CommonArgs.cpp
+++ b/clang/lib/Driver/ToolChains/CommonArgs.cpp
@@ -1459,18 +1459,25 @@ void tools::addArchSpecificRPath(const ToolChain &TC, const ArgList &Args,
   }
 }
 
+bool tools::addLLVMOffloadingRuntime(const Compilation &C,
+                                     ArgStringList &CmdArgs,
+                                     const ToolChain &TC, const ArgList &Args) {
+
+  if (!Args.hasFlag(options::OPT_foffload_via_llvm,
+                    options::OPT_fno_offload_via_llvm, false))
+    return false;
+
+  CmdArgs.push_back("-lLLVMOffloadKernel");
+  return true;
+}
+
 bool tools::addOpenMPRuntime(const Compilation &C, ArgStringList &CmdArgs,
                              const ToolChain &TC, const ArgList &Args,
                              bool ForceStaticHostRuntime, bool IsOffloadingHost,
                              bool GompNeedsRT) {
   if (!Args.hasFlag(options::OPT_fopenmp, options::OPT_fopenmp_EQ,
-                    options::OPT_fno_openmp, false)) {
-    // We need libomptarget (liboffload) if it's the choosen offloading runtime.
-    if (Args.hasFlag(options::OPT_foffload_via_llvm,
-                     options::OPT_fno_offload_via_llvm, false))
-      CmdArgs.push_back("-lomptarget");
+                    options::OPT_fno_openmp, false))
     return false;
-  }
 
   Driver::OpenMPRuntimeKind RTKind = TC.getDriver().getOpenMPRuntime(Args);
 
diff --git a/clang/lib/Driver/ToolChains/Cuda.cpp b/clang/lib/Driver/ToolChains/Cuda.cpp
index 5040e451ed9e5..f79109f71efb3 100644
--- a/clang/lib/Driver/ToolChains/Cuda.cpp
+++ b/clang/lib/Driver/ToolChains/Cuda.cpp
@@ -301,6 +301,23 @@ CudaInstallationDetector::CudaInstallationDetector(
 
 void CudaInstallationDetector::AddCudaIncludeArgs(
     const ArgList &DriverArgs, ArgStringList &CC1Args) const {
+  if (DriverArgs.hasFlag(options::OPT_foffload_via_llvm,
+                         options::OPT_fno_offload_via_llvm, false)) {
+    if (DriverArgs.hasFlag(options::OPT_offload_inc,
+                           options::OPT_no_offload_inc, true) &&
+        !DriverArgs.hasArg(options::OPT_nobuiltininc)) {
+      CC1Args.append({"-include", "__clang_gpu_device_functions.h"});
+
+      SmallString<128> CudaIncludePath(D.ResourceDir);
+      llvm::sys::path::append(CudaIncludePath, "..", "..", "..");
+      llvm::sys::path::append(CudaIncludePath, "include", "offload", "cuda");
+      CC1Args.push_back("-internal-isystem");
+      CC1Args.push_back(DriverArgs.MakeArgString(CudaIncludePath));
+      CC1Args.append({"-include", "cuda_runtime.h"});
+    }
+    return;
+  }
+
   if (!DriverArgs.hasArg(options::OPT_nobuiltininc)) {
     // Add cuda_wrappers/* to our system include path.  This lets us wrap
     // standard library headers.
@@ -395,7 +412,9 @@ void NVPTX::Assembler::ConstructJob(Compilation &C, const JobAction &JA,
                                     const char *LinkingOutput) const {
   const auto &TC =
       static_cast<const toolchains::NVPTXToolChain &>(getToolChain());
-  assert(TC.getTriple().isNVPTX() && "Wrong platform");
+
+  bool UsesLLVMOffloading = Args.hasFlag(
+      options::OPT_foffload_via_llvm, options::OPT_fno_offload_via_llvm, false);
 
   BoundArch GPUArch;
   // If this is a CUDA action we need to extract the device architecture
@@ -418,7 +437,7 @@ void NVPTX::Assembler::ConstructJob(Compilation &C, const JobAction &JA,
          "Device action expected to have an architecture.");
 
   // Check that our installation's ptxas supports gpu_arch.
-  if (!Args.hasArg(options::OPT_no_cuda_version_check)) {
+  if (!UsesLLVMOffloading && !Args.hasArg(options::OPT_no_cuda_version_check)) {
     TC.CudaInstallation.CheckCudaVersionSupportsArch(GPUArch.Arch);
   }
 
@@ -491,7 +510,8 @@ void NVPTX::Assembler::ConstructJob(Compilation &C, const JobAction &JA,
                                /*Default=*/true);
   else if (JA.isOffloading(Action::OFK_Cuda))
     // In CUDA we generate relocatable code by default.
-    Relocatable = Args.hasFlag(options::OPT_fgpu_rdc, options::OPT_fno_gpu_rdc,
+    Relocatable = UsesLLVMOffloading ||
+                  Args.hasFlag(options::OPT_fgpu_rdc, options::OPT_fno_gpu_rdc,
                                /*Default=*/false);
   else
     // Otherwise, we are compiling directly and should create linkable output.
@@ -540,7 +560,9 @@ void NVPTX::FatBinary::ConstructJob(Compilation &C, const JobAction &JA,
                                     const char *LinkingOutput) const {
   const auto &TC =
       static_cast<const toolchains::CudaToolChain &>(getToolChain());
-  assert(TC.getTriple().isNVPTX() && "Wrong platform");
+  bool UsesLLVMOffloading = Args.hasFlag(
+      options::OPT_foffload_via_llvm, options::OPT_fno_offload_via_llvm, false);
+  assert((UsesLLVMOffloading || TC.getTriple().isNVPTX()) && "Wrong platform");
 
   ArgStringList CmdArgs;
   if (TC.CudaInstallation.version() <= CudaVersion::CUDA_100)
@@ -588,7 +610,9 @@ void NVPTX::Linker::ConstructJob(Compilation &C, const JobAction &JA,
       static_cast<const toolchains::NVPTXToolChain &>(getToolChain());
   ArgStringList CmdArgs;
 
-  assert(TC.getTriple().isNVPTX() && "Wrong platform");
+  bool UsesLLVMOffloading = Args.hasFlag(
+      options::OPT_foffload_via_llvm, options::OPT_fno_offload_via_llvm, false);
+  assert((UsesLLVMOffloading || TC.getTriple().isNVPTX()) && "Wrong platform");
 
   assert((Output.isFilename() || Output.isNothing()) && "Invalid output.");
   if (Output.isFilename()) {
@@ -886,9 +910,12 @@ void CudaToolChain::addClangTargetOptions(
     BoundArch BA, Action::OffloadKind DeviceOffloadingKind) const {
   HostTC.addClangTargetOptions(DriverArgs, CC1Args, BA, DeviceOffloadingKind);
 
+  bool UsesLLVMOffloading = DriverArgs.hasFlag(
+      options::OPT_foffload_via_llvm, options::OPT_fno_offload_via_llvm, false);
+
   StringRef GpuArch = DriverArgs.getLastArgValue(options::OPT_march_EQ);
   assert((DeviceOffloadingKind == Action::OFK_OpenMP ||
-          DeviceOffloadingKind == Action::OFK_Cuda) &&
+          DeviceOffloadingKind == Action::OFK_Cuda || UsesLLVMOffloading) &&
          "Only OpenMP or CUDA offloading kinds are supported for NVIDIA GPUs.");
 
   CC1Args.append({"-fcuda-is-device", "-mllvm",
@@ -907,6 +934,9 @@ void CudaToolChain::addClangTargetOptions(
       DriverArgs.hasArg(options::OPT_S))
     return;
 
+  if (UsesLLVMOffloading)
+    return;
+
   std::string LibDeviceFile = CudaInstallation.getLibDeviceFile(GpuArch);
   if (LibDeviceFile.empty()) {
     getDriver().Diag(diag::err_drv_no_cuda_libdevice) << GpuArch;
@@ -916,13 +946,6 @@ void CudaToolChain::addClangTargetOptions(
   CC1Args.push_back("-mlink-builtin-bitcode");
   CC1Args.push_back(DriverArgs.MakeArgString(LibDeviceFile));
 
-  // For now, we don't use any Offload/OpenMP device runtime when we offload
-  // CUDA via LLVM/Offload. We should split the Offload/OpenMP device runtime
-  // and include the "generic" (or CUDA-specific) parts.
-  if (DriverArgs.hasFlag(options::OPT_foffload_via_llvm,
-                         options::OPT_fno_offload_via_llvm, false))
-    return;
-
   clang::CudaVersion CudaInstallationVersion = CudaInstallation.version();
 
   if (CudaInstallationVersion >= CudaVersion::UNKNOWN)
@@ -963,6 +986,23 @@ llvm::DenormalMode CudaToolChain::getDefaultDenormalModeForType(
 
 void CudaToolChain::AddCudaIncludeArgs(const ArgList &DriverArgs,
                                        ArgStringList &CC1Args) const {
+  if (DriverArgs.hasFlag(options::OPT_foffload_via_llvm,
+                         options::OPT_fno_offload_via_llvm, false)) {
+    if (DriverArgs.hasFlag(options::OPT_offload_inc,
+                           options::OPT_no_offload_inc, true) &&
+        !DriverArgs.hasArg(options::OPT_nobuiltininc)) {
+      CC1Args.append({"-include", "__clang_gpu_device_functions.h"});
+
+      SmallString<128> CudaIncludePath(getDriver().ResourceDir);
+      llvm::sys::path::append(CudaIncludePath, "..", "..", "..");
+      llvm::sys::path::append(CudaIncludePath, "include", "offload", "cuda");
+      CC1Args.push_back("-internal-isystem");
+      CC1Args.push_back(DriverArgs.MakeArgString(CudaIncludePath));
+      CC1Args.append({"-include", "cuda_runtime.h"});
+    }
+    return;
+  }
+
   // Check our CUDA version if we're going to include the CUDA headers.
   if (DriverArgs.hasFlag(options::OPT_offload_inc, options::OPT_no_offload_inc,
                          true) &&
@@ -1035,6 +1075,10 @@ CudaToolChain::GetCXXStdlibType(const ArgList &Args) const {
 
 void CudaToolChain::AddClangSystemIncludeArgs(const ArgList &DriverArgs,
                                               ArgStringList &CC1Args) const {
+  if (DriverArgs.hasFlag(options::OPT_foffload_via_llvm,
+                         options::OPT_fno_offload_via_llvm, false))
+    return;
+
   HostTC.AddClangSystemIncludeArgs(DriverArgs, CC1Args);
 
   if (DriverArgs.hasFlag(options::OPT_offload_inc, options::OPT_no_offload_inc,
diff --git a/clang/lib/Driver/ToolChains/Gnu.cpp b/clang/lib/Driver/ToolChains/Gnu.cpp
index 24076d8814322..72affac131701 100644
--- a/clang/lib/Driver/ToolChains/Gnu.cpp
+++ b/clang/lib/Driver/ToolChains/Gnu.cpp
@@ -510,6 +510,7 @@ void tools::gnutools::Linker::ConstructJob(Compilation &C, const JobAction &JA,
         // FIXME: Does this really make sense for all GNU toolchains?
         WantPthread = true;
 
+      addLLVMOffloadingRuntime(C, CmdArgs, ToolChain, Args);
       AddRunTimeLibs(ToolChain, D, CmdArgs, Args);
 
       // LLVM support for atomics on 32-bit SPARC V8+ is incomplete, so
diff --git a/clang/lib/Driver/ToolChains/Linux.cpp b/clang/lib/Driver/ToolChains/Linux.cpp
index 1ab385a9ea001..03017f20e391d 100644
--- a/clang/lib/Driver/ToolChains/Linux.cpp
+++ b/clang/lib/Driver/ToolChains/Linux.cpp
@@ -881,7 +881,9 @@ void Linux::addOffloadRTLibs(unsigned ActiveKinds, const ArgList &Args,
   if (!Args.hasFlag(options::OPT_offloadlib, options::OPT_no_offloadlib,
                     true) ||
       Args.hasArg(options::OPT_nostdlib) ||
-      Args.hasArg(options::OPT_no_hip_rt) || Args.hasArg(options::OPT_r))
+      Args.hasArg(options::OPT_no_hip_rt) || Args.hasArg(options::OPT_r) ||
+      Args.hasFlag(options::OPT_foffload_via_llvm,
+                   options::OPT_fno_offload_via_llvm, false))
     return;
 
   llvm::SmallVector<std::pair<StringRef, StringRef>> Libraries;
diff --git a/clang/lib/Headers/__clang_gpu_builtin_vars.h b/clang/lib/Headers/__clang_gpu_builtin_vars.h
index b80248dcd2be3..f3137be0a3181 100644
--- a/clang/lib/Headers/__clang_gpu_builtin_vars.h
+++ b/clang/lib/Headers/__clang_gpu_builtin_vars.h
@@ -6,6 +6,8 @@
 //
 //===-----------------------------------------------------------------------===
 
+#include <stddef.h>
+
 #ifndef __CLANG_GPU_BUILTIN_VARS_H__
 #define __CLANG_GPU_BUILTIN_VARS_H__
 
@@ -20,6 +22,23 @@ static inline __attribute__((device)) const struct {
   }
 } warpSize{};
 
+extern "C" {
+
+typedef struct dim3 {
+  dim3() {}
+  dim3(unsigned x) : x(x) {}
+  unsigned x = 0, y = 0, z = 0;
+} dim3;
+
+// TODO: For some reason the CUDA device compilation requires this declaration
+// to be present on the device while it is only used on the host.
+unsigned __llvmPushCallConfiguration(dim3 gridDim, dim3 blockDim,
+                                     size_t sharedMem = 0, void *stream = 0);
+unsigned __llvmLaunchKernel(const void *func, dim3 gridDim, dim3 blockDim,
+                            void **args, size_t sharedMem = 0,
+                            void *stream = 0);
+}
+
 // Make sure nobody can create instances of the coordinate types, take their
 // address, copy, or assign them.
 #pragma push_macro("__GPU_DISALLOW_BUILTINVAR_ACCESS")
diff --git a/clang/test/CodeGenCUDA/Inputs/cuda.h b/clang/test/CodeGenCUDA/Inputs/cuda.h
index 421fa4dd7dbae..83bb7b7bdbb7f 100644
--- a/clang/test/CodeGenCUDA/Inputs/cuda.h
+++ b/clang/test/CodeGenCUDA/Inputs/cuda.h
@@ -55,7 +55,7 @@ extern "C" hipError_t hipLaunchKernel_spt(const void *func, dim3 gridDim,
 #elif __OFFLOAD_VIA_LLVM__
 extern "C" unsigned __llvmPushCallConfiguration(dim3 gridDim, dim3 blockDim,
                                      size_t sharedMem = 0, void *stream = 0);
-extern "C" unsigned llvmLaunchKernel(const void *func, dim3 gridDim, dim3 blockDim,
+extern "C" unsigned __llvmLaunchKernel(const void *func, dim3 gridDim, dim3 blockDim,
                           void **args, size_t sharedMem = 0, void *stream = 0);
 #else
 typedef struct cudaStream *cudaStream_t;
diff --git a/clang/test/CodeGenCUDA/offload_via_llvm.cu b/clang/test/CodeGenCUDA/offload_via_llvm.cu
index b13a64c81b775..c58e793d11826 100644
--- a/clang/test/CodeGenCUDA/offload_via_llvm.cu
+++ b/clang/test/CodeGenCUDA/offload_via_llvm.cu
@@ -14,9 +14,7 @@
 // HST-NEXT:    [[DOTADDR1:%.*]] = alloca i16, align 2
 // HST-NEXT:    [[DOTADDR2:%.*]] = alloca ptr, align 4
 // HST-NEXT:    [[DOTADDR3:%.*]] = alloca ptr, align 4
-// HST-NEXT:    [[KERNEL_ARGS:%.*]] = alloca [[TMP0]], align 16
-// HST-NEXT:    [[KERNEL_ARGS_PTRS:%.*]] = alloca [4 x ptr], align 16
-// HST-NEXT:    [[KERNEL_LAUNCH_PARAMS:%.*]] = alloca [[TMP1]], align 16
+// HST-NEXT:    [[KERNEL_LAUNCH_PARAMS:%.*]] = alloca [[TMP0]], align 16
 // HST-NEXT:    [[GRID_DIM:%.*]] = alloca [[STRUCT_DIM3:%.*]], align 8
 // HST-NEXT:    [[BLOCK_DIM:%.*]] = alloca [[STRUCT_DIM3]], align 8
 // HST-NEXT:    [[SHMEM_SIZE:%.*]] = alloca i32, align 4
@@ -25,34 +23,34 @@
 // HST-NEXT:    store i16 [[TMP1]], ptr [[DOTADDR1]], align 2
 // HST-NEXT:    store ptr [[TMP2]], ptr [[DOTADDR2]], align 4
 // HST-NEXT:    store ptr [[TMP3]], ptr [[DOTADDR3]], align 4
-// HST-NEXT:    [[TMP4:%.*]] = getelementptr inbounds nuw [[TMP1]], ptr [[KERNEL_LAUNCH_PARAMS]], i32 0, i32 0
-// HST-NEXT:    store i32 4, ptr [[TMP4]], align 16
-// HST-NEXT:    [[TMP5:%.*]] = getelementptr inbounds nuw [[TMP1]], ptr [[KERNEL_LAUNCH_PARAMS]], i32 0, i32 1
-// HST-NEXT:    store ptr [[KERNEL_ARGS_PTRS]], ptr [[TMP5]], align 4
-// HST-NEXT:    [[TMP6:%.*]] = load i32, ptr [[DOTADDR]], align 4
-// HST-NEXT:    [[TMP7:%.*]] = getelementptr inbounds nuw [[TMP0]], ptr [[KERNEL_ARGS]], i32 0, i32 0
-// HST-NEXT:    store i32 [[TMP6]], ptr [[TMP7]], align 16
-// HST-NEXT:    [[TMP8:%.*]] = getelementptr inbounds [4 x ptr], ptr [[KERNEL_ARGS_PTRS]], i32 0, i32 0
-// HST-NEXT:    store ptr [[TMP7]], ptr [[TMP8]], align 16
-// HST-NEXT:    [[TMP9:%.*]] = load i16, ptr [[DOTADDR1]], align 2
-// HST-NEXT:    [[TMP10:%.*]] = getelementptr inbounds nuw [[TMP0]], ptr [[KERNEL_ARGS]], i32 0, i32 1
-// HST-NEXT:    store i16 [[TMP9]], ptr [[TMP10]], align 4
-// HST-NEXT:    [[TMP11:%.*]] = getelementptr inbounds [4 x ptr], ptr [[KERNEL_ARGS_PTRS]], i32 0, i32 1
-// HST-NEXT:    store ptr [[TMP10]], ptr [[TMP11]], align 4
-// HST-NEXT:    [[TMP12:%.*]] = load ptr, ptr [[DOTADDR2]], align 4
-// HST-NEXT:    [[TMP13:%.*]] = getelementptr inbounds nuw [[TMP0]], ptr [[KERNEL_ARGS]], i32 0, i32 2
-// HST-NEXT:    store ptr [[TMP12]], ptr [[TMP13]], align 8
-// HST-NEXT:    [[TMP14:%.*]] = getelementptr inbounds [4 x ptr], ptr [[KERNEL_ARGS_PTRS]], i32 0, i32 2
-// HST-NEXT:    store ptr [[TMP13]], ptr [[TMP14]], align 8
-// HST-NEXT:    [[TMP15:%.*]] = load ptr, ptr [[DOTADDR3]], align 4
-// HST-NEXT:    [[TMP16:%.*]] = getelementptr inbounds nuw [[TMP0]], ptr [[KERNEL_ARGS]], i32 0, i32 3
-// HST-NEXT:    store ptr [[TMP15]], ptr [[TMP16]], align 4
-// HST-NEXT:    [[TMP17:%.*]] = getelementptr inbounds [4 x ptr], ptr [[KERNEL_ARGS_PTRS]], i32 0, i32 3
-// HST-NEXT:    store ptr [[TMP16]], ptr [[TMP17]], align 4
-// HST-NEXT:    [[TMP18:%.*]] = call i32 @__llvmPopCallConfiguration(ptr [[GRID_DIM]], ptr [[BLOCK_DIM]], ptr [[SHMEM_SIZE]], ptr [[STREAM]])
-// HST-NEXT:    [[TMP19:%.*]] = load i32, ptr [[SHMEM_SIZE]], align 4
-// HST-NEXT:    [[TMP20:%.*]] = load ptr, ptr [[STREAM]], align 4
-// HST-NEXT:    [[CALL:%.*]] = call noundef i32 @llvmLaunchKernel(ptr noundef @_Z18__device_stub__fooisPvS_, ptr noundef byval([[STRUCT_DIM3]]) align 4 [[GRID_DIM]], ptr noundef byval([[STRUCT_DIM3]]) align 4 [[BLOCK_DIM]], ptr noundef [[KERNEL_LAUNCH_PARAMS]], i32 noundef [[TMP19]], ptr noundef [[TMP20]]) #[[ATTR3:[0-9]+]]
+// HST-NEXT:    [[KERNEL_ARGS:%.*]] = alloca ptr, i32 4, align 16
+// HST-NEXT:    [[KERNEL_ARG_SIZES:%.*]] = alloca i32, i32 4, align 16
+// HST-NEXT:    [[TMP4:%.*]] = getelementptr inbounds nuw [[TMP0]], ptr [[KERNEL_LAUNCH_PARAMS]], i32 0, i32 0
+// HST-NEXT:    store ptr [[KERNEL_ARGS]], ptr [[TMP4]], align 16
+// HST-NEXT:    [[TMP5:%.*]] = getelementptr inbounds nuw [[TMP0]], ptr [[KERNEL_LAUNCH_PARAMS]], i32 0, i32 1
+// HST-NEXT:    store i64 4, ptr [[TMP5]], align 8
+// HST-NEXT:    [[TMP6:%.*]] = getelementptr inbounds nuw [[TMP0]], ptr [[KERNEL_LAUNCH_PARAMS]], i32 0, i32 2
+// HST-NEXT:    store ptr [[KERNEL_ARG_SIZES]], ptr [[TMP6]], align 16
+// HST-NEXT:    [[TMP7:%.*]] = getelementptr ptr, ptr [[KERNEL_ARGS]], i32 0
+// HST-NEXT:    store ptr [[DOTADDR]], ptr [[TMP7]], align 4
+// HST-NEXT:    [[TMP8:%.*]] = getelementptr i32, ptr [[KERNEL_ARG_SIZES]], i32 0
+// HST-NEXT:    store i32 4, ptr [[TMP8]], align 4
+// HST-NEXT:    [[TMP9:%.*]] = getelementptr ptr, ptr [[KERNEL_ARGS]], i32 1
+// HST-NEXT:    store ptr [[DOTADDR1]], ptr [[TMP9]], align 4
+// HST-NEXT:    [[TMP10:%.*]] = getelementptr i32, ptr [[KERNEL_ARG_SIZES]], i32 1
+// HST-NEXT:    store i32 2, ptr [[TMP10]], align 4
+// HST-NEXT:    [[TMP11:%.*]] = getelementptr ptr, ptr [[KERNEL_ARGS]], i32 2
+// HST-NEXT:    store ptr [[DOTADDR2]], ptr [[TMP11]], align 4
+// HST-NEXT:    [[TMP12:%.*]] = getelementptr i32, ptr [[KERNEL_ARG_SIZES]], i32 2
+// HST-NEXT:    store i32 4, ptr [[TMP12]], align 4
+// HST-NEXT:    [[TMP13:%.*]] = getelementptr ptr, ptr [[KERNEL_ARGS]], i32 3
+// HST-NEXT:    store ptr [[DOTADDR3]], ptr [[TMP13]], align 4
+// HST-NEXT:    [[TMP14:%.*]] = getelementptr i32, ptr [[KERNEL_ARG_SIZES]], i32 3
+// HST-NEXT:    store i32 4, ptr [[TMP14]], align 4
+// HST-NEXT:    [[TMP15:%.*]] = call i32 @__cudaPopCallConfiguration(ptr [[GRID_DIM]], ptr [[BLOCK_DIM]], ptr [[SHMEM_SIZE]], ptr [[STREAM]])
+// HST-NEXT:    [[TMP16:%.*]] = load i32, ptr [[SHMEM_SIZE]], align 4
+// HST-NEXT:    [[TMP17:%.*]] = load ptr, ptr [[STREAM]], align 4
+// HST-NEXT:    [[CALL:%.*]] = call noundef i32 @__llvmLaunchKernel(ptr noundef @_Z18__device_stub__fooisPvS_, ptr noundef byval([[STRUCT_DIM3]]) align 4 [[GRID_DIM]], ptr noundef byval([[STRUCT_DIM3]]) align 4 [[BLOCK_DIM]], ptr noundef [[KERNEL_LAUNCH_PARAMS]], i32 noundef [[TMP16]], ptr noundef [[TMP17]]) #[[ATTR3:[0-9]+]]
 // HST-NEXT:    br label %[[SETUP_END:.*]]
 // HST:       [[SETUP_END]]:
 // HST-NEXT:    ret void
diff --git a/clang/test/Driver/cuda-via-liboffload.cu b/clang/test/Driver/cuda-via-liboffload.cu
index 68dc963e906b2..d30e529f0ce12 100644
--- a/clang/test/Driver/cuda-via-liboffload.cu
+++ b/clang/test/Driver/cuda-via-liboffload.cu
@@ -2,21 +2,20 @@
 // RUN:        --offload-arch=sm_35 --offload-arch=sm_70 %s 2>&1 \
 // RUN: | FileCheck -check-prefix BINDINGS %s
 
-//      BINDINGS: "x86_64-unknown-linux-gnu" - "clang", inputs: ["[[INPUT:.+]]"], output: "[[HOST_BC:.+]]"
-// BINDINGS-NEXT: "nvptx64-nvidia-cuda" - "clang", inputs: ["[[INPUT]]", "[[HOST_BC]]"], output: "[[PTX_SM_35:.+]]"
-// BINDINGS-NEXT: "nvptx64-nvidia-cuda" - "NVPTX::Assembler", inputs: ["[[PTX_SM_35]]"], output: "[[CUBIN_SM_35:.+]]"
-// BINDINGS-NEXT: "nvptx64-nvidia-cuda" - "clang", inputs: ["[[INPUT]]", "[[HOST_BC]]"], output: "[[PTX_SM_70:.+]]"
-// BINDINGS-NEXT: "nvptx64-nvidia-cuda" - "NVPTX::Assembler", inputs: ["[[PTX_SM_70:.+]]"], output: "[[CUBIN_SM_70:.+]]"
+// BINDINGS: "nvptx64-nvidia-cuda-llvm" - "clang", inputs: ["[[INPUT:.+]]"], output: "[[PTX_SM_35:.+]]"
+// BINDINGS-NEXT: "nvptx64-nvidia-cuda-llvm" - "NVPTX::Assembler", inputs: ["[[PTX_SM_35]]"], output: "[[CUBIN_SM_35:.+]]"
+// BINDINGS-NEXT: "nvptx64-nvidia-cuda-llvm" - "clang", inputs: ["[[INPUT]]"], output: "[[PTX_SM_70:.+]]"
+// BINDINGS-NEXT: "nvptx64-nvidia-cuda-llvm" - "NVPTX::Assembler", inputs: ["[[PTX_SM_70:.+]]"], output: "[[CUBIN_SM_70:.+]]"
 // BINDINGS-NEXT: "x86_64-unknown-linux-gnu" - "Offload::Packager", inputs: ["[[CUBIN_SM_35]]", "[[CUBIN_SM_70]]"], output: "[[BINARY:.+]]"
-// BINDINGS-NEXT: "x86_64-unknown-linux-gnu" - "clang", inputs: ["[[HOST_BC]]", "[[BINARY]]"], output: "[[HOST_OBJ:.+]]"
+// BINDINGS-NEXT: "x86_64-unknown-linux-gnu" - "clang", inputs: ["[[INPUT]]", "[[BINARY]]"], output: "[[HOST_OBJ:.+]]"
 // BINDINGS-NEXT: "x86_64-unknown-linux-gnu" - "Offload::Linker", inputs: ["[[HOST_OBJ]]"], output: "a.out"
 
 // RUN: %clang -### -target x86_64-linux-gnu -foffload-via-llvm -ccc-print-bindings \
 // RUN:        --offload-arch=sm_35 --offload-arch=sm_70 %s 2>&1 \
 // RUN: | FileCheck -check-prefix BINDINGS-DEVICE %s
 
-// BINDINGS-DEVICE: # "nvptx64-nvidia-cuda" - "clang", inputs: ["[[INPUT:.+]]"], output: "[[PTX:.+]]"
-// BINDINGS-DEVICE: # "nvptx64-nvidia-cuda" - "NVPTX::Assembler", inputs: ["[[PTX]]"], output: "[[CUBIN:.+]]"
+// BINDINGS-DEVICE: # "nvptx64-nvidia-cuda-llvm" - "clang", inputs: ["[[INPUT:.+]]"], output: "[[PTX:.+]]"
+// BINDINGS-DEVICE: # "nvptx64-nvidia-cuda-llvm" - "NVPTX::Assembler", inputs: ["[[PTX]]"], output: "[[CUBIN:.+]]"
 
 // RUN: %clang -### -target x86_64-linux-gnu -ccc-print-bindings --offload-link -foffload-via-llvm %s 2>&1 | FileCheck -check-prefix DEVICE-LINK %s
 
diff --git a/clang/test/OffloadTools/clang-linker-wrapper/linker-wrapper-image.c b/clang/test/OffloadTools/clang-linker-wrapper/linker-wrapper-image.c
index d05e9d54a108a..083d3340f6f81 100644
--- a/clang/test/OffloadTools/clang-linker-wrapper/linker-wrapper-image.c
+++ b/clang/test/OffloadTools/clang-linker-wrapper/linker-wrapper-image.c
@@ -94,7 +94,7 @@
 // CUDA-NEXT:   br i1 %1, label %while.entry, label %while.end
 //
 //      CUDA: while.entry:
-// CUDA-NEXT:   %entry1 = phi ptr [ @__start_llvm_offload_entries, %entry ], [ %16, %if.end ]
+// CUDA-NEXT:   %entry1 = phi ptr [ @__start_llvm_offload_entries, %entry ], [ %17, %if.end ]
 // CUDA-NEXT:   %2 = getelementptr inbounds %struct.__tgt_offload_entry, ptr %entry1, i32 0, i32 4
 // CUDA-NEXT:   %addr = load ptr, ptr %2, align 8
 // CUDA-NEXT:   %3 = getelementptr inbounds %struct.__tgt_offload_entry, ptr %entry1, i32 0, i32 8
@@ -117,15 +117,16 @@
 // CUDA-NEXT:   %constant = lshr i32 %11, 4
 // CUDA-NEXT:   %12 = and i32 %flags, 32
 // CUDA-NEXT:   %normalized = lshr i32 %12, 5
-// CUDA-NEXT:   %13 = icmp eq i16 %kind, 2
-// CUDA-NEXT:   br i1 %13, label %if.kind, label %if.end
+// CUDA-NEXT:   %13 = and i16 %kind, 2
+// CUDA-NEXT:   %14 = icmp ne i16 %13, 0
+// CUDA-NEXT:   br i1 %14, label %if.kind, label %if.end
 //
 //      CUDA: if.kind:
-// CUDA-NEXT:   %14 = icmp eq i64 %size, 0
-// CUDA-NEXT:   br i1 %14, label %if.then, label %if.else
+// CUDA-NEXT:   %15 = icmp eq i64 %size, 0
+// CUDA-NEXT:   br i1 %15, label %if.then, label %if.else
 //
 //      CUDA: if.then:
-// CUDA-NEXT:   %15 = call i32 @__cudaRegisterFunction(ptr %0, ptr %addr, ptr %name, ptr %name, i32 -1, ptr null, ptr null, ptr null, ptr null, ptr null)
+// CUDA-NEXT:   %16 = call i32 @__cudaRegisterFunction(ptr %0, ptr %addr, ptr %name, ptr %name, i32 -1, ptr null, ptr null, ptr null, ptr null, ptr null)
 // CUDA-NEXT:   br label %if.end
 //
 //      CUDA: if.else:
@@ -151,9 +152,9 @@
 // CUDA-NEXT:   br label %if.end
 //
 //      CUDA: if.end:
-// CUDA-NEXT:   %16 = getelementptr inbounds %struct.__tgt_offload_entry, ptr %entry1, i64 1
-// CUDA-NEXT:   %17 = icmp eq ptr %16, @__stop_llvm_offload_entries
-// CUDA-NEXT:   br i1 %17, label %while.end, label %while.entry
+// CUDA-NEXT:   %17 = getelementptr inbounds %struct.__tgt_offload_entry, ptr %entry1, i64 1
+// CUDA-NEXT:   %18 = icmp eq ptr %17, @__stop_llvm_offload_entries
+// CUDA-NEXT:   br i1 %18, label %while.end, label %while.entry
 //
 //      CUDA: while.end:
 // CUDA-NEXT:   ret void
@@ -236,7 +237,7 @@
 // HIP-NEXT:   br i1 %1, label %while.entry, label %while.end
 //
 //      HIP: while.entry:
-// HIP-NEXT:   %entry1 = phi ptr [ @{{.*offload_entries.*}}, %entry ], [ %16, %if.end ]
+// HIP-NEXT:   %entry1 = phi ptr [ @{{.*offload_entries.*}}, %entry ], [ %17, %if.end ]
 // HIP-NEXT:   %2 = getelementptr inbounds %struct.__tgt_offload_entry, ptr %entry1, i32 0, i32 4
 // HIP-NEXT:   %addr = load ptr, ptr %2, align 8
 // HIP-NEXT:   %3 = getelementptr inbounds %struct.__tgt_offload_entry, ptr %entry1, i32 0, i32 8
@@ -259,15 +260,16 @@
 // HIP-NEXT:   %constant = lshr i32 %11, 4
 // HIP-NEXT:   %12 = and i32 %flags, 32
 // HIP-NEXT:   %normalized = lshr i32 %12, 5
-// HIP-NEXT:   %13 = icmp eq i16 %kind, 4
-// HIP-NEXT:   br i1 %13, label %if.kind, label %if.end
+// HIP-NEXT:   %13 = and i16 %kind, 4
+// HIP-NEXT:   %14 = icmp ne i16 %13, 0
+// HIP-NEXT:   br i1 %14, label %if.kind, label %if.end
 //
 //      HIP: if.kind:
-// HIP-NEXT:   %14 = icmp eq i64 %size, 0
-// HIP-NEXT:   br i1 %14, label %if.then, label %if.else
+// HIP-NEXT:   %15 = icmp eq i64 %size, 0
+// HIP-NEXT:   br i1 %15, label %if.then, label %if.else
 //
 //      HIP: if.then:
-// HIP-NEXT:   %15 = call i32 @__hipRegisterFunction(ptr %0, ptr %addr, ptr %name, ptr %name, i32 -1, ptr null, ptr null, ptr null, ptr null, ptr null)
+// HIP-NEXT:   %16 = call i32 @__hipRegisterFunction(ptr %0, ptr %addr, ptr %name, ptr %name, i32 -1, ptr null, ptr null, ptr null, ptr null, ptr null)
 // HIP-NEXT:   br label %if.end
 //
 //      HIP: if.else:
@@ -295,9 +297,9 @@
 // HIP-NEXT:   br label %if.end
 //
 //      HIP: if.end:
-// HIP-NEXT:   %16 = getelementptr inbounds %struct.__tgt_offload_entry, ptr %entry1, i64 1
-// HIP-NEXT:   %17 = icmp eq ptr %16, @{{.*offload_entries.*}}
-// HIP-NEXT:   br i1 %17, label %while.end, label %while.entry
+// HIP-NEXT:   %17 = getelementptr inbounds %struct.__tgt_offload_entry, ptr %entry1, i64 1
+// HIP-NEXT:   %18 = icmp eq ptr %17, @{{.*offload_entries.*}}
+// HIP-NEXT:   br i1 %18, label %while.end, label %while.entry
 //
 //      HIP: while.end:
 // HIP-NEXT:   ret void
diff --git a/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp b/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp
index c2de6578773c7..b61c4dae85a86 100644
--- a/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp
+++ b/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp
@@ -139,6 +139,13 @@ static bool CanonicalPrefixes = true;
 
 using OffloadingImage = OffloadBinary::OffloadingImage;
 
+static bool usesLLVMOffloadWrapper(ArrayRef<OffloadingImage> Images) {
+  return llvm::any_of(Images, [](const OffloadingImage &Image) {
+    return Triple(Image.StringData.lookup("triple")).getEnvironment() ==
+           Triple::LLVM;
+  });
+}
+
 namespace llvm {
 // Provide DenseMapInfo so that OffloadKind can be used in a DenseMap.
 template <> struct DenseMapInfo<OffloadKind> {
@@ -971,6 +978,9 @@ Expected<SmallVector<std::unique_ptr<MemoryBuffer>>>
 bundleLinkedOutput(ArrayRef<OffloadingImage> Images, const ArgList &Args,
                    OffloadKind Kind) {
   llvm::TimeTraceScope TimeScope("Bundle linked output");
+  if (usesLLVMOffloadWrapper(Images))
+    return bundleOpenMP(Images);
+
   switch (Kind) {
   case OFK_OpenMP:
     return (Verbose && SaveTemps) ? bundleOpenMPVerbose(Images)
@@ -1214,7 +1224,8 @@ linkAndWrapDeviceFiles(ArrayRef<SmallVector<OffloadFile>> LinkerInputFiles,
       continue;
     }
 
-    auto OutputOrErr = wrapDeviceImages(*BundledImagesOrErr, Args, Kind);
+    OffloadKind WrapperKind = usesLLVMOffloadWrapper(Input) ? OFK_OpenMP : Kind;
+    auto OutputOrErr = wrapDeviceImages(*BundledImagesOrErr, Args, WrapperKind);
     if (!OutputOrErr)
       return OutputOrErr.takeError();
     WrappedOutput.push_back(*OutputOrErr);
diff --git a/llvm/lib/Frontend/Offloading/OffloadWrapper.cpp b/llvm/lib/Frontend/Offloading/OffloadWrapper.cpp
index ded603a1e00e3..037b81a7c42fb 100644
--- a/llvm/lib/Frontend/Offloading/OffloadWrapper.cpp
+++ b/llvm/lib/Frontend/Offloading/OffloadWrapper.cpp
@@ -485,10 +485,12 @@ Function *createRegisterGlobalsFunction(Module &M, bool IsHIP,
                               llvm::offloading::OffloadGlobalNormalized));
   auto *Normalized = Builder.CreateLShr(
       NormalizedBit, ConstantInt::get(Type::getInt32Ty(C), 5), "normalized");
-  auto *KindCond = Builder.CreateICmpEQ(
+  auto *KindAnd = Builder.CreateAnd(
       Kind, ConstantInt::get(Type::getInt16Ty(C),
                              IsHIP ? object::OffloadKind::OFK_HIP
                                    : object::OffloadKind::OFK_Cuda));
+  auto *KindCond =
+      Builder.CreateICmpNE(KindAnd, ConstantInt::get(Type::getInt16Ty(C), 0));
   Builder.CreateCondBr(KindCond, IfKindBB, IfEndBB);
   Builder.SetInsertPoint(IfKindBB);
   auto *FnCond = Builder.CreateICmpEQ(
diff --git a/offload/CMakeLists.txt b/offload/CMakeLists.txt
index 2885b20f9c1d8..2dd4446979c05 100644
--- a/offload/CMakeLists.txt
+++ b/offload/CMakeLists.txt
@@ -340,6 +340,7 @@ if(BUILD_LIBOMPTARGET)
 endif()
 
 add_subdirectory(liboffload)
+add_subdirectory(languages)
 
 # Add tests.
 if(OFFLOAD_INCLUDE_TESTS)
diff --git a/offload/languages/CMakeLists.txt b/offload/languages/CMakeLists.txt
new file mode 100644
index 0000000000000..fbd774e645e21
--- /dev/null
+++ b/offload/languages/CMakeLists.txt
@@ -0,0 +1,3 @@
+add_subdirectory(kernel)
+add_subdirectory(cuda)
+add_subdirectory(hip)
diff --git a/offload/languages/cuda/CMakeLists.txt b/offload/languages/cuda/CMakeLists.txt
new file mode 100644
index 0000000000000..f09e9b46ef038
--- /dev/null
+++ b/offload/languages/cuda/CMakeLists.txt
@@ -0,0 +1 @@
+install(FILES ${CMAKE_CURRENT_SOURCE_DIR}/../include/cuda/cuda_runtime.h DESTINATION ${CMAKE_INSTALL_PREFIX}/include/offload/cuda/)
diff --git a/offload/languages/cuda/src/cuda_runtime.cpp b/offload/languages/cuda/src/cuda_runtime.cpp
new file mode 100644
index 0000000000000..a09cf86697b95
--- /dev/null
+++ b/offload/languages/cuda/src/cuda_runtime.cpp
@@ -0,0 +1,15 @@
+//===-- cuda_runtime.cpp - CUDA runtime API implementations ---------------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#include "cuda_runtime.h"
+
+#include "OffloadAPI.h"
+
+#define LANGUAGE cuda
+
+#include "../../kernel/src/LanguageRuntime.cpp"
diff --git a/offload/languages/hip/CMakeLists.txt b/offload/languages/hip/CMakeLists.txt
new file mode 100644
index 0000000000000..ea41dd59adf5c
--- /dev/null
+++ b/offload/languages/hip/CMakeLists.txt
@@ -0,0 +1 @@
+install(FILES ${CMAKE_CURRENT_SOURCE_DIR}/../include/hip/hip_runtime.h DESTINATION ${CMAKE_INSTALL_PREFIX}/include/offload/hip/)
diff --git a/offload/languages/hip/src/hip_runtime.cpp b/offload/languages/hip/src/hip_runtime.cpp
new file mode 100644
index 0000000000000..276fa8334ed77
--- /dev/null
+++ b/offload/languages/hip/src/hip_runtime.cpp
@@ -0,0 +1,16 @@
+//===-- hip_runtime.cpp - HIP runtime API implementations -----------------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#include "hip_runtime.h"
+
+#include "LanguageLaunch.h"
+#include "OffloadAPI.h"
+
+#define LANGUAGE hip
+
+#include "../../kernel/src/LanguageRuntime.cpp"
diff --git a/offload/languages/include/cuda/cuda_runtime.h b/offload/languages/include/cuda/cuda_runtime.h
new file mode 100644
index 0000000000000..a140ce1a1aca0
--- /dev/null
+++ b/offload/languages/include/cuda/cuda_runtime.h
@@ -0,0 +1,24 @@
+//===-- cuda_runtime.h - CUDA runtime API declarations --------------------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef LLVM_OFFLOAD_LANGUAGES_INCLUDE_CUDA_CUDA_RUNTIME_H
+#define LLVM_OFFLOAD_LANGUAGES_INCLUDE_CUDA_CUDA_RUNTIME_H
+
+#define LANGUAGE cuda
+
+#include "../kernel/DefineLanguageNames.inc"
+
+#include "../kernel/LanguageRuntime.h"
+
+#include "../kernel/UndefineLanguageNames.inc"
+
+#undef LANGUAGE
+
+using cudaDeviceProp = cudaDeviceProp_t;
+
+#endif // LLVM_OFFLOAD_LANGUAGES_INCLUDE_CUDA_CUDA_RUNTIME_H
diff --git a/offload/languages/include/hip/hip_runtime.h b/offload/languages/include/hip/hip_runtime.h
new file mode 100644
index 0000000000000..b56e295c904e3
--- /dev/null
+++ b/offload/languages/include/hip/hip_runtime.h
@@ -0,0 +1,61 @@
+//===-- hip_runtime.h - HIP runtime API declarations ----------------------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef LLVM_OFFLOAD_LANGUAGES_INCLUDE_HIP_HIP_RUNTIME_H
+#define LLVM_OFFLOAD_LANGUAGES_INCLUDE_HIP_HIP_RUNTIME_H
+
+#define LANGUAGE hip
+
+#include "../kernel/DefineLanguageNames.inc"
+
+#include "../kernel/LanguageRuntime.h"
+
+#include "../kernel/UndefineLanguageNames.inc"
+
+#undef LANGUAGE
+
+#define hipHostMallocDefault hipHostAllocDefault
+#define hipHostMallocPortable hipHostAllocPortable
+#define hipHostMallocMapped hipHostAllocMapped
+#define hipHostMallocWriteCombined hipHostAllocWriteCombined
+#define hipHostMallocNonCoherent 0x80000000
+
+inline hipError_t hipHostMalloc(void **Ptr, size_t Size, unsigned int Flags) {
+  return hipHostAlloc(Ptr, Size, Flags);
+}
+
+inline hipError_t hipHostFree(void *Ptr) { return ::hipFreeHost(Ptr); }
+
+#ifdef __cplusplus
+template <class T>
+static inline hipError_t hipHostMalloc(T **Ptr, size_t Size,
+                                       unsigned int Flags) {
+  return ::hipHostMalloc((void **)Ptr, Size, Flags);
+}
+
+template <class T> static inline hipError_t hipHostFree(T *Ptr) {
+  return ::hipHostFree((void *)Ptr);
+}
+#endif
+
+#if defined(__AMDGPU__) || defined(__NVPTX__)
+#define HIP_KERNEL_NAME(...) __VA_ARGS__
+
+extern "C" hipError_t hipLaunchKernel(const char *Kernel, dim3 GridDim,
+                                      dim3 BlockDim, void **KernelArgs,
+                                      size_t DynamicSharedMem, void *Stream);
+
+template <typename... AT, typename FT = void (*)(AT...)>
+static inline void hipLaunchKernelGGL(FT Kernel, dim3 GridDim, dim3 BlockDim,
+                                      size_t DynamicSharedMem, void *Stream,
+                                      AT... KernelArgs) {
+  Kernel<<<GridDim, BlockDim, DynamicSharedMem, Stream>>>(KernelArgs...);
+}
+#endif
+
+#endif // LLVM_OFFLOAD_LANGUAGES_INCLUDE_HIP_HIP_RUNTIME_H
diff --git a/offload/languages/include/kernel/DefineLanguageNames.inc b/offload/languages/include/kernel/DefineLanguageNames.inc
new file mode 100644
index 0000000000000..989d119ae6ed7
--- /dev/null
+++ b/offload/languages/include/kernel/DefineLanguageNames.inc
@@ -0,0 +1,44 @@
+//===-- DefineLanguageNames.inc - Kernel language API name definitions ----===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#define COMBINE2(X, Y) X##Y
+#define COMBINE(X, Y) COMBINE2(X, Y)
+
+#define Error_t COMBINE(LANGUAGE, Error_t)
+#define DeviceProp_t COMBINE(LANGUAGE, DeviceProp_t)
+#define Malloc COMBINE(LANGUAGE, Malloc)
+#define Free COMBINE(LANGUAGE, Free)
+#define Memcpy COMBINE(LANGUAGE, Memcpy)
+#define DeviceSynchronize COMBINE(LANGUAGE, DeviceSynchronize)
+#define Success COMBINE(LANGUAGE, Success)
+#define ErrorInvalidValue COMBINE(LANGUAGE, ErrorInvalidValue)
+#define ErrorInvalidDevice COMBINE(LANGUAGE, ErrorInvalidDevice)
+#define ErrorUnknown COMBINE(LANGUAGE, ErrorUnknown)
+#define GetErrorName COMBINE(LANGUAGE, GetErrorName)
+#define GetErrorString COMBINE(LANGUAGE, GetErrorString)
+#define MemcpyKind COMBINE(LANGUAGE, MemcpyKind)
+#define MemcpyHostToHost COMBINE(LANGUAGE, MemcpyHostToHost)
+#define MemcpyHostToDevice COMBINE(LANGUAGE, MemcpyHostToDevice)
+#define MemcpyDeviceToHost COMBINE(LANGUAGE, MemcpyDeviceToHost)
+#define MemcpyDeviceToDevice COMBINE(LANGUAGE, MemcpyDeviceToDevice)
+#define MemcpyDefault COMBINE(LANGUAGE, MemcpyDefault)
+#define GetDevice COMBINE(LANGUAGE, GetDevice)
+#define GetDeviceCount COMBINE(LANGUAGE, GetDeviceCount)
+#define SetDevice COMBINE(LANGUAGE, SetDevice)
+#define HostAlloc COMBINE(LANGUAGE, HostAlloc)
+#define HostAllocDefault COMBINE(LANGUAGE, HostAllocDefault)
+#define HostAllocPortable COMBINE(LANGUAGE, HostAllocPortable)
+#define HostAllocMapped COMBINE(LANGUAGE, HostAllocMapped)
+#define HostAllocWriteCombined COMBINE(LANGUAGE, HostAllocWriteCombined)
+#define MallocHost COMBINE(LANGUAGE, MallocHost)
+#define FreeHost COMBINE(LANGUAGE, FreeHost)
+#define GetDeviceProperties COMBINE(LANGUAGE, GetDeviceProperties)
+#define Stream_t COMBINE(LANGUAGE, Stream_t)
+#define StreamCreate COMBINE(LANGUAGE, StreamCreate)
+#define StreamDestroy COMBINE(LANGUAGE, StreamDestroy)
+#define StreamSynchronize COMBINE(LANGUAGE, StreamSynchronize)
diff --git a/offload/languages/include/kernel/LanguageRuntime.h b/offload/languages/include/kernel/LanguageRuntime.h
new file mode 100644
index 0000000000000..ce1c888edbe94
--- /dev/null
+++ b/offload/languages/include/kernel/LanguageRuntime.h
@@ -0,0 +1,193 @@
+//===-- LanguageRuntime.h - Kernel language runtime API declarations ------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef LLVM_OFFLOAD_LANGUAGES_INCLUDE_KERNEL_LANGUAGE_RUNTIME_H
+#define LLVM_OFFLOAD_LANGUAGES_INCLUDE_KERNEL_LANGUAGE_RUNTIME_H
+
+#include <cstddef>
+#include <cstdint>
+#include <cstdio>
+#include <cstdlib>
+
+enum Error_t : uint32_t {
+  Success = 0,
+  ErrorInvalidValue = 1,
+  ErrorInvalidDevice = 2,
+  ErrorUnknown = 3,
+};
+
+const char *GetErrorName(Error_t Error);
+
+const char *GetErrorString(Error_t Error);
+
+struct DeviceProp_t {
+  char name[256];
+  size_t totalGlobalMem;
+  int warpSize;
+  int multiProcessorCount;
+  int major;
+  int minor;
+  int ECCEnabled;
+  int pciBusID;
+  int pciDeviceID;
+  int pciDomainID;
+  int memoryBusWidth;
+};
+
+enum MemcpyKind {
+  MemcpyHostToHost = 0,
+  MemcpyHostToDevice = 1,
+  MemcpyDeviceToHost = 2,
+  MemcpyDeviceToDevice = 3,
+  MemcpyDefault = 4
+};
+
+enum HostAllocFlags : unsigned int {
+  HostAllocDefault = 0x00,
+  HostAllocPortable = 0x01,
+  HostAllocMapped = 0x02,
+  HostAllocWriteCombined = 0x04,
+};
+
+typedef struct Stream_st *Stream_t;
+
+/// Malloc, with type template overlay.
+///{
+Error_t Malloc(void **Dev_Ptr, size_t Size);
+
+template <class T> static inline Error_t Malloc(T **dev_Ptr, size_t Size) {
+  return ::Malloc((void **)dev_Ptr, Size);
+}
+
+Error_t HostAlloc(void **Ptr, size_t Size, unsigned int Flags);
+
+template <class T>
+static inline Error_t HostAlloc(T **Ptr, size_t Size, unsigned int Flags) {
+  return ::HostAlloc((void **)Ptr, Size, Flags);
+}
+
+Error_t MallocHost(void **Ptr, size_t Size);
+
+template <class T> static inline Error_t MallocHost(T **Ptr, size_t Size) {
+  return ::MallocHost((void **)Ptr, Size);
+}
+///}
+
+/// Free, no type template necessary.
+Error_t Free(void *Dev_Ptr);
+
+/// Memcpy, with type template overlay.
+///{
+Error_t Memcpy(void *Dst, const void *Src, size_t Size, MemcpyKind Kind);
+
+template <class T>
+static inline Error_t Memcpy(T *Dst, const T *Src, size_t Size,
+                             MemcpyKind Kind) {
+  return ::Memcpy((void *)Dst, (const void *)Src, Size, Kind);
+}
+///}
+
+Error_t DeviceSynchronize();
+
+Error_t GetDevice(int *DeviceNo);
+
+Error_t GetDeviceCount(int *Count);
+
+Error_t SetDevice(int DeviceNo);
+
+Error_t FreeHost(void *Ptr);
+
+Error_t GetDeviceProperties(DeviceProp_t *DeviceProp, int DeviceNo);
+
+Error_t StreamCreate(Stream_t *stream);
+
+Error_t StreamDestroy(Stream_t stream);
+
+Error_t StreamSynchronize(Stream_t stream);
+
+#if defined(__AMDGPU__) || defined(__NVPTX__)
+#include <gpuintrin.h>
+
+/// Define \p FIELD as a property backed by component \p OFFSET of Vec.
+#define __LLVM_OFFLOAD_DEVICE_BUILTIN(FIELD, OFFSET)                           \
+  __declspec(property(get = __get_##FIELD,                                     \
+                      put = __put_##FIELD)) unsigned int FIELD;                \
+  __device__ inline __attribute__((always_inline)) T __get_##FIELD(void)       \
+      const {                                                                  \
+    return Vec[OFFSET];                                                        \
+  }                                                                            \
+  __device__ inline __attribute__((always_inline)) T __put_##FIELD(T V) {      \
+    return Vec[OFFSET] = V;                                                    \
+  }
+
+/// Common storage for CUDA/HIP vector aliases such as int4 and float3.
+///
+/// Provides array-style indexing and x/y/z/w component properties over Clang
+/// ext_vector_type storage.
+template <class T, int Size> struct BaseVector {
+  using VT = float __attribute__((ext_vector_type(Size)));
+  VT Vec;
+
+  __device__ __host__ BaseVector() = default;
+
+  /// Construct a vector from component values.
+  template <typename... Args>
+  __device__ __host__ BaseVector(Args... args) : BaseVector({args...}) {}
+
+  /// Return component \p Idx.
+  __device__ __host__ T &operator[](int Idx) { return Vec[Idx]; }
+  __device__ __host__ const T &operator[](int Idx) const { return Vec[Idx]; }
+
+  __LLVM_OFFLOAD_DEVICE_BUILTIN(x, 0);
+  __LLVM_OFFLOAD_DEVICE_BUILTIN(y, 1);
+  __LLVM_OFFLOAD_DEVICE_BUILTIN(z, 2);
+  __LLVM_OFFLOAD_DEVICE_BUILTIN(w, 3);
+};
+
+/// Define the vector alias TY##SIZE and its make_TY##SIZE constructor helper.
+#define __VECTOR_DEF_IMPL(TY, SIZE)                                            \
+  using TY##SIZE = BaseVector<TY, SIZE>;                                       \
+                                                                               \
+  template <typename... Args>                                                  \
+  __device__ __host__ TY##SIZE make_##TY##SIZE(Args... args) {                 \
+    return TY##SIZE(args...);                                                  \
+  }
+
+/// Define the standard 1/2/3/4/8/16 element vector aliases for \p TY.
+#define __VECTOR_DEF(TY)                                                       \
+  __VECTOR_DEF_IMPL(TY, 1)                                                     \
+  __VECTOR_DEF_IMPL(TY, 2)                                                     \
+  __VECTOR_DEF_IMPL(TY, 3)                                                     \
+  __VECTOR_DEF_IMPL(TY, 4)                                                     \
+  __VECTOR_DEF_IMPL(TY, 8)                                                     \
+  __VECTOR_DEF_IMPL(TY, 16)
+
+/// Instantiate CUDA/HIP-style vector types and make_* helpers.
+__VECTOR_DEF(float)
+__VECTOR_DEF(double)
+__VECTOR_DEF(int8_t)
+__VECTOR_DEF(int16_t)
+__VECTOR_DEF(int32_t)
+__VECTOR_DEF(int64_t)
+__VECTOR_DEF(uint8_t)
+__VECTOR_DEF(uint16_t)
+__VECTOR_DEF(uint32_t)
+__VECTOR_DEF(uint64_t)
+__VECTOR_DEF(char)
+__VECTOR_DEF(short)
+__VECTOR_DEF(int)
+__VECTOR_DEF(unsigned)
+__VECTOR_DEF(long)
+
+#undef __VECTOR_DEF_IMPL
+#undef __VECTOR_DEF
+#undef __LLVM_OFFLOAD_DEVICE_BUILTIN
+
+#endif
+
+#endif // LLVM_OFFLOAD_LANGUAGES_INCLUDE_KERNEL_LANGUAGE_RUNTIME_H
diff --git a/offload/languages/include/kernel/UndefineLanguageNames.inc b/offload/languages/include/kernel/UndefineLanguageNames.inc
new file mode 100644
index 0000000000000..b2a9f2fd846ae
--- /dev/null
+++ b/offload/languages/include/kernel/UndefineLanguageNames.inc
@@ -0,0 +1,42 @@
+//===-- UndefineLanguageNames.inc - Kernel language API name undefines ----===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#undef Error_t
+#undef DeviceProp_t
+#undef Malloc
+#undef Free
+#undef Memcpy
+#undef DeviceSynchronize
+#undef Success
+#undef ErrorInvalidValue
+#undef ErrorInvalidDevice
+#undef ErrorUnknown
+#undef GetErrorName
+#undef GetErrorString
+#undef MemcpyKind
+#undef MemcpyHostToHost
+#undef MemcpyHostToDevice
+#undef MemcpyDeviceToHost
+#undef MemcpyDeviceToDevice
+#undef MemcpyDefault
+#undef GetDevice
+#undef GetDeviceCount
+#undef SetDevice
+#undef HostAlloc
+#undef HostAllocFlags
+#undef HostAllocDefault
+#undef HostAllocPortable
+#undef HostAllocMapped
+#undef HostAllocWriteCombined
+#undef MallocHost
+#undef FreeHost
+#undef GetDeviceProperties
+#undef Stream_t
+#undef StreamCreate
+#undef StreamDestroy
+#undef StreamSynchronize
diff --git a/offload/languages/kernel/CMakeLists.txt b/offload/languages/kernel/CMakeLists.txt
new file mode 100644
index 0000000000000..97954a32d26c0
--- /dev/null
+++ b/offload/languages/kernel/CMakeLists.txt
@@ -0,0 +1,53 @@
+add_llvm_library(
+  LLVMOffloadKernel SHARED
+
+  ../cuda/src/cuda_runtime.cpp
+  ../hip/src/hip_runtime.cpp
+  src/LanguageCommon.cpp
+  src/State.cpp
+
+  LINK_COMPONENTS
+  Support
+  Offload
+  )
+
+if(LLVM_HAVE_LINK_VERSION_SCRIPT)
+  target_link_libraries(LLVMOffloadKernel PRIVATE "-Wl,--version-script=${CMAKE_CURRENT_SOURCE_DIR}/exports")
+endif()
+
+target_include_directories(LLVMOffloadKernel PUBLIC
+                            ${CMAKE_CURRENT_BINARY_DIR}/../include
+                            ${CMAKE_CURRENT_SOURCE_DIR}/../include/cuda
+                            ${CMAKE_CURRENT_SOURCE_DIR}/../include/hip
+                            ${CMAKE_CURRENT_SOURCE_DIR}/../include/kernel
+                            ${CMAKE_CURRENT_SOURCE_DIR}/include
+                            ${CMAKE_CURRENT_SOURCE_DIR}/../kernel/include
+                            ${CMAKE_CURRENT_SOURCE_DIR}/../../liboffload/include
+                            ${CMAKE_CURRENT_BINARY_DIR}/../../liboffload/API
+                            ${CMAKE_CURRENT_SOURCE_DIR}/../../include
+                            ${CMAKE_CURRENT_SOURCE_DIR}/../../plugins-nextgen/common/include)
+
+target_compile_options(LLVMOffloadKernel PRIVATE ${offload_compile_flags})
+target_link_options(LLVMOffloadKernel PRIVATE ${offload_link_flags})
+
+target_compile_definitions(LLVMOffloadKernel PRIVATE
+  TARGET_NAME="libLLVMOffloadKernel"
+  DEBUG_PREFIX="LLVMOffloadKernel"
+)
+
+set_target_properties(LLVMOffloadKernel PROPERTIES
+                      RUNTIME_OUTPUT_DIRECTORY "${LLVM_LIBRARY_OUTPUT_INTDIR}/${OFFLOAD_TARGET_SUBDIR}"
+                      POSITION_INDEPENDENT_CODE ON
+                      INSTALL_RPATH "$ORIGIN"
+                      BUILD_RPATH "$ORIGIN:${CMAKE_CURRENT_BINARY_DIR}/..")
+install(TARGETS LLVMOffloadKernel
+        COMPONENT offload
+        RUNTIME DESTINATION "${CMAKE_INSTALL_BINDIR}"
+        LIBRARY DESTINATION "${OFFLOAD_INSTALL_LIBDIR}"
+        ARCHIVE DESTINATION "${OFFLOAD_INSTALL_LIBDIR}")
+
+install(FILES
+        ${CMAKE_CURRENT_SOURCE_DIR}/../include/kernel/DefineLanguageNames.inc
+        ${CMAKE_CURRENT_SOURCE_DIR}/../include/kernel/LanguageRuntime.h
+        ${CMAKE_CURRENT_SOURCE_DIR}/../include/kernel/UndefineLanguageNames.inc
+        DESTINATION ${CMAKE_INSTALL_PREFIX}/include/offload/kernel/)
diff --git a/offload/languages/kernel/exports b/offload/languages/kernel/exports
new file mode 100644
index 0000000000000..3372890ff42c6
--- /dev/null
+++ b/offload/languages/kernel/exports
@@ -0,0 +1,11 @@
+VERS1.0 {
+  global:
+    *cuda*;
+    *hip*;
+    llvmLaunchKernel*;
+    __llvm*;
+    __tgt_register_lib;
+    __tgt_unregister_lib;
+  local:
+    *;
+};
diff --git a/offload/languages/kernel/include/LanguageAliases.inc b/offload/languages/kernel/include/LanguageAliases.inc
new file mode 100644
index 0000000000000..4331e37054dc2
--- /dev/null
+++ b/offload/languages/kernel/include/LanguageAliases.inc
@@ -0,0 +1,48 @@
+//===-- LanguageAliases.inc - Language runtime symbol aliases -------------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+//
+// The LLVM offload kernel runtime implements the shared __llvm* entry points
+// once. These aliases expose the CUDA/HIP ABI names, such as
+// __cudaRegisterFunction and __hipPopCallConfiguration, as aliases to the
+// shared implementations.
+//
+//===----------------------------------------------------------------------===//
+
+// This file intentionally has no include guard. It is included once for each
+// LANGUAGE value to emit each language's alias definitions.
+
+#ifndef LANGUAGE
+#error This file should be included, or used, with a LANGUAGE macro set.
+#endif
+
+#define MA_IMPL2(PREFIX, L, RTY, NAME, ...)                                    \
+  extern "C" [[gnu::alias("__llvm" #NAME)]] RTY PREFIX##L##NAME(__VA_ARGS__);
+
+#define MA_IMPL1(PREFIX, L, RTY, NAME, ...)                                    \
+  MA_IMPL2(PREFIX, L, RTY, NAME, __VA_ARGS__)
+
+#define MAKE_ALIAS(PREFIX, RTY, NAME, ...)                                     \
+  MA_IMPL1(PREFIX, LANGUAGE, RTY, NAME, __VA_ARGS__)
+
+MAKE_ALIAS(__, void, RegisterFunction, const char *, const char *, char *,
+           const char *, int, uint3 *, uint3 *, dim3 *, dim3 *, int *)
+MAKE_ALIAS(__, void, RegisterVar, void **, char *, char *, const char *, int,
+           int, int, int)
+MAKE_ALIAS(__, void, RegisterManagedVar, void **, char *, char *, const char *,
+           size_t, unsigned)
+MAKE_ALIAS(__, void, RegisterSurface, void **, const struct surfaceReference *,
+           const void **, const char *, int, int)
+MAKE_ALIAS(__, void, RegisterTexture, void **, const struct textureReference *,
+           const void **, const char *, int, int, int)
+
+MAKE_ALIAS(__, unsigned, PopCallConfiguration, dim3 *, dim3 *, size_t *,
+           void **)
+
+#undef MAKE_ALIAS
+#undef MA_IMPL1
+#undef MA_IMPL2
diff --git a/offload/languages/kernel/include/LanguageLaunch.h b/offload/languages/kernel/include/LanguageLaunch.h
new file mode 100644
index 0000000000000..161aead85bbef
--- /dev/null
+++ b/offload/languages/kernel/include/LanguageLaunch.h
@@ -0,0 +1,40 @@
+//===-- LanguageLaunch.h - Language launch API declarations ---------------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef LLVM_OFFLOAD_LANGUAGES_KERNEL_INCLUDE_LANGUAGE_LAUNCH_H
+#define LLVM_OFFLOAD_LANGUAGES_KERNEL_INCLUDE_LANGUAGE_LAUNCH_H
+
+#include "OffloadAPI.h"
+#include "Types.h"
+
+#include <algorithm> // for std::max
+#include <cstddef>
+#include <cstdint>
+
+extern "C" {
+
+/// Push call configuration for kernel launch
+unsigned __llvmPushCallConfiguration(dim3 __grid_size, dim3 __block_size,
+                                     size_t __shared_memory, void *__stream);
+
+/// Pop call configuration for kernel launch
+unsigned __llvmPopCallConfiguration(dim3 *__grid_size, dim3 *__block_size,
+                                    size_t *__shared_memory, void **__stream);
+
+/// Internal kernel launch implementation
+ol_result_t __llvmLaunchKernelImpl(const char *KernelID, dim3 GridDim,
+                                   dim3 BlockDim, void *KernelArgsPtr,
+                                   size_t DynamicSharedMem, void *Stream);
+
+/// LLVM-style kernel launch entry point
+unsigned __llvmLaunchKernel(const char *KernelID, dim3 GridDim, dim3 BlockDim,
+                            void *KernelArgsPtr, size_t DynamicSharedMem,
+                            void *Stream);
+}
+
+#endif // LLVM_OFFLOAD_LANGUAGES_KERNEL_INCLUDE_LANGUAGE_LAUNCH_H
diff --git a/offload/languages/kernel/include/LanguageRegistration.h b/offload/languages/kernel/include/LanguageRegistration.h
new file mode 100644
index 0000000000000..7570ce76de546
--- /dev/null
+++ b/offload/languages/kernel/include/LanguageRegistration.h
@@ -0,0 +1,44 @@
+//===-- LanguageRegistration.h - Language registration API declarations ---===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef LLVM_OFFLOAD_LANGUAGES_KERNEL_INCLUDE_LANGUAGE_REGISTRATION_H
+#define LLVM_OFFLOAD_LANGUAGES_KERNEL_INCLUDE_LANGUAGE_REGISTRATION_H
+
+#include "OffloadAPI.h"
+#include "Types.h"
+
+#include <cstdint>
+#include <iterator>
+
+/// Hidden, but exported, Registration API
+///{
+extern "C" {
+
+void __llvmRegisterFunction(const char *Binary, const char *KernelID,
+                            char *KernelName, const char *KernelName1, int,
+                            uint3 *, uint3 *, dim3 *, dim3 *, int *);
+
+void __llvmRegisterVar(void **, char *, char *, const char *, int, int, int,
+                       int);
+
+void __llvmRegisterManagedVar(void **, char *, char *, const char *, size_t,
+                              unsigned);
+
+void __llvmRegisterSurface(void **, const struct surfaceReference *,
+                           const void **, const char *, int, int);
+
+void __llvmRegisterTexture(void **, const struct textureReference *,
+                           const void **, const char *, int, int, int);
+
+struct __tgt_bin_desc;
+void __tgt_register_lib(__tgt_bin_desc *Desc);
+void __tgt_unregister_lib(__tgt_bin_desc *Desc);
+}
+///}
+
+#endif // LLVM_OFFLOAD_LANGUAGES_KERNEL_INCLUDE_LANGUAGE_REGISTRATION_H
diff --git a/offload/languages/kernel/include/LanguageUtils.h b/offload/languages/kernel/include/LanguageUtils.h
new file mode 100644
index 0000000000000..95335f2b66779
--- /dev/null
+++ b/offload/languages/kernel/include/LanguageUtils.h
@@ -0,0 +1,40 @@
+//===-- LanguageUtils.h - Kernel Language utility functions ---------------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#pragma once
+
+#ifndef LANGUAGE
+#error This file should be included, or used, with a LANGUAGE macro set.
+#endif
+
+#include "OffloadAPI.h"
+
+/// Convert an ol_result_t to the active language's Error_t.
+static inline Error_t convertResult(ol_result_t Result) {
+  if (Result == OL_SUCCESS)
+    return Success;
+  switch (Result->Code) {
+  case OL_ERRC_INVALID_VALUE:
+  case OL_ERRC_INVALID_ARGUMENT:
+  case OL_ERRC_INVALID_NULL_POINTER:
+    return ErrorInvalidValue;
+  case OL_ERRC_INVALID_DEVICE:
+    return ErrorInvalidDevice;
+  default:
+    return ErrorUnknown;
+  }
+}
+
+/// Convert a Stream_t to an ol_queue_handle_t.
+static inline Error_t getQueueFromStream(Stream_t Stream,
+                                         ol_queue_handle_t *Queue) {
+  if (!Stream)
+    return ErrorInvalidValue;
+  *Queue = reinterpret_cast<ol_queue_handle_t>(Stream);
+  return Success;
+}
diff --git a/offload/languages/kernel/include/State.h b/offload/languages/kernel/include/State.h
new file mode 100644
index 0000000000000..4c8e9ce101ff0
--- /dev/null
+++ b/offload/languages/kernel/include/State.h
@@ -0,0 +1,149 @@
+//===-- State.h - Kernel language persistent state ------------------------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef LLVM_OFFLOAD_LANGUAGES_KERNEL_INCLUDE_STATE_H
+#define LLVM_OFFLOAD_LANGUAGES_KERNEL_INCLUDE_STATE_H
+
+#include "OffloadAPI.h"
+#include "Types.h"
+
+#include "llvm/ADT/ArrayRef.h"
+#include "llvm/ADT/DenseMap.h"
+#include "llvm/ADT/SmallVector.h"
+#include "llvm/Support/raw_ostream.h"
+
+#define CHECK_FATAL(Result, ...)                                               \
+  if (Result && Result->Code) {                                                \
+    llvm::errs() << __VA_ARGS__ << '\n';                                       \
+    abort();                                                                   \
+  }
+
+namespace llvm {
+namespace offload {
+
+/// Opaque host-side key used to identify a registered kernel.
+///
+/// This is the address emitted in the offload entry table for the kernel
+using KernelIDTy = const void *;
+
+/// Per-thread state used by the language runtime entry points.
+///
+/// Tracks the current thread's default device, optional per-thread queue,
+/// and pending kernel launch configuration.
+struct ThreadStateTy {
+  ~ThreadStateTy();
+
+  /// Return the default queue for the current host thread
+  static ol_queue_handle_t getDefaultQueue();
+
+  /// Return the thread-local default device, or the first discovered device.
+  static ol_device_handle_t getDefaultDevice();
+
+  /// Return the pending kernel launch configuration for this thread.
+  static CallConfigurationTy &getCallConfiguration();
+
+  /// Set the thread-local default device to \p Device and recreate its queue.
+  static void setDefaultDevice(ol_device_handle_t Device);
+
+private:
+  static ThreadStateTy &get();
+
+  void createDefaultQueue(ol_device_handle_t Device);
+
+  ol_device_handle_t DefaultDevice = nullptr;
+  ol_queue_handle_t DefaultQueue = nullptr;
+
+  CallConfigurationTy CC = {};
+
+  ThreadStateTy();
+};
+
+/// Process-wide state shared by CUDA and HIP language entry points.
+///
+/// Owns the discovered devices, host device, process default queue, and maps
+/// from registered binaries and kernels to liboffload handles.
+struct StateTy {
+  ~StateTy();
+
+  friend struct ThreadStateTy;
+
+  /// Return the host device discovered during runtime initialization.
+  static ol_device_handle_t getHostDevice();
+
+  /// Return the number of non-host devices available to kernel languages.
+  static int getDeviceCount();
+
+  /// Return the thread-local default device and write its number to \p
+  /// DeviceNo.
+  static ol_device_handle_t getDevice(int *DeviceNo);
+
+  /// Set the thread-local default device by device number.
+  ///
+  /// \returns the selected device, or nullptr if \p DeviceNo is invalid.
+  static ol_device_handle_t setDefaultDevice(int DeviceNo);
+
+  /// Register \p Kernel for the host-side kernel identifier \p ID.
+  ///
+  /// \p ID is the opaque kernel key emitted by Clang in the offload entry
+  /// table.  It is later passed to the launch entry point to recover the
+  /// corresponding liboffload symbol handle.
+  static void registerKernel(const void *ID, ol_symbol_handle_t Kernel);
+
+  /// Remove any registered kernel handle for the host-side kernel key \p ID.
+  static void unregisterKernel(const void *ID);
+
+  /// Return the registered kernel handle for the host-side kernel key \p ID.
+  static ol_symbol_handle_t getKernel(const void *ID);
+
+  /// Register \p Program for the binary image identifier \p ID.
+  ///
+  /// \p ID is the device image start address from the offload binary
+  /// descriptor.  It keys the loaded program so later function registration
+  /// can look up the program that owns each kernel symbol.
+  static void registerProgram(const void *ID, ol_program_handle_t Program);
+
+  /// Remove and return the loaded program handle for binary image key \p ID.
+  static ol_program_handle_t unregisterProgram(const void *ID);
+
+  /// Return the loaded program handle for binary image key \p ID.
+  static ol_program_handle_t getProgram(const void *ID);
+
+private:
+  static StateTy &get();
+  static StateTy *tryGet();
+  static bool addDevices(ol_device_handle_t Device, void *Payload);
+
+  llvm::ArrayRef<ol_device_handle_t> getDevices() const;
+
+  void addDevice(ol_device_handle_t Device);
+  void setHostDevice(ol_device_handle_t Device);
+
+  void addKernel(KernelIDTy KernelID, ol_symbol_handle_t Kernel);
+  void removeKernel(KernelIDTy KernelID);
+  ol_symbol_handle_t lookupKernel(KernelIDTy KernelID);
+
+  void addProgram(const void *Binary, ol_program_handle_t Program);
+  ol_program_handle_t removeProgram(const void *Binary);
+  ol_program_handle_t lookupProgram(const void *Binary);
+
+  void destroyRegisteredPrograms();
+
+  llvm::DenseMap<const void *, ol_program_handle_t> BinaryRegisterMap;
+  llvm::DenseMap<KernelIDTy, ol_symbol_handle_t> KernelMap;
+  llvm::SmallVector<ol_device_handle_t, 8> Devices;
+
+  ol_queue_handle_t DefaultQueue = nullptr;
+  ol_device_handle_t HostDevice = nullptr;
+
+  StateTy();
+};
+
+} // namespace offload
+} // namespace llvm
+
+#endif // LLVM_OFFLOAD_LANGUAGES_KERNEL_INCLUDE_STATE_H
diff --git a/offload/languages/kernel/include/Types.h b/offload/languages/kernel/include/Types.h
new file mode 100644
index 0000000000000..860b9431a3a24
--- /dev/null
+++ b/offload/languages/kernel/include/Types.h
@@ -0,0 +1,29 @@
+//===-- Types.h - Kernel language API types -------------------------------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef LLVM_OFFLOAD_LANGUAGES_KERNEL_INCLUDE_TYPES_H
+#define LLVM_OFFLOAD_LANGUAGES_KERNEL_INCLUDE_TYPES_H
+
+#include "Types.h"
+#include <cstddef>
+#include <cstdint>
+
+struct uint3 {
+  unsigned x = 0, y = 0, z = 0;
+};
+
+using dim3 = uint3;
+
+struct CallConfigurationTy {
+  dim3 GridSize;
+  dim3 BlockSize;
+  size_t SharedMemory;
+  void *Stream;
+};
+
+#endif // LLVM_OFFLOAD_LANGUAGES_KERNEL_INCLUDE_TYPES_H
diff --git a/offload/languages/kernel/src/LanguageCommon.cpp b/offload/languages/kernel/src/LanguageCommon.cpp
new file mode 100644
index 0000000000000..b2f06a030c246
--- /dev/null
+++ b/offload/languages/kernel/src/LanguageCommon.cpp
@@ -0,0 +1,18 @@
+//===-- LanguageCommon.cpp - Shared CUDA/HIP runtime entry points ---------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#include "LanguageLaunch.cpp"
+#include "LanguageRegistration.cpp"
+
+#define LANGUAGE cuda
+#include "LanguageAliases.inc"
+#undef LANGUAGE
+
+#define LANGUAGE hip
+#include "LanguageAliases.inc"
+#undef LANGUAGE
diff --git a/offload/languages/kernel/src/LanguageLaunch.cpp b/offload/languages/kernel/src/LanguageLaunch.cpp
new file mode 100644
index 0000000000000..71c2b275beb35
--- /dev/null
+++ b/offload/languages/kernel/src/LanguageLaunch.cpp
@@ -0,0 +1,83 @@
+//===-- LanguageLaunch.cpp - Language launch API --------------------------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#include "LanguageLaunch.h"
+#include "State.h"
+
+#include <cstdio>
+
+using RuntimeState = llvm::offload::StateTy;
+using ThreadState = llvm::offload::ThreadStateTy;
+
+extern "C" {
+
+/// Push call configuration for kernel launch
+unsigned __llvmPushCallConfiguration(dim3 GridSize, dim3 BlockSize,
+                                     size_t SharedMemory, void *Stream) {
+  CallConfigurationTy &CC = ThreadState::getCallConfiguration();
+
+  CC.GridSize = GridSize;
+  CC.BlockSize = BlockSize;
+  CC.SharedMemory = SharedMemory;
+  CC.Stream = Stream;
+  return 0;
+}
+
+/// Pop call configuration for kernel launch
+unsigned __llvmPopCallConfiguration(dim3 *GridSize, dim3 *BlockSize,
+                                    size_t *SharedMemory, void **Stream) {
+  CallConfigurationTy &CC = ThreadState::getCallConfiguration();
+  *GridSize = CC.GridSize;
+  *BlockSize = CC.BlockSize;
+  *SharedMemory = CC.SharedMemory;
+  *Stream = CC.Stream;
+  return 0;
+}
+
+/// Internal kernel launch implementation
+ol_result_t __llvmLaunchKernelImpl(const char *KernelID, dim3 GridDim,
+                                   dim3 BlockDim, void *KernelArgsPtr,
+                                   size_t DynamicSharedMem, void *Stream) {
+  ol_device_handle_t Device = ThreadState::getDefaultDevice();
+  ol_symbol_handle_t Kernel = RuntimeState::getKernel(KernelID);
+
+  ol_kernel_launch_size_args_t LaunchSizeArgs;
+  LaunchSizeArgs.Dimensions =
+      1 + (GridDim.y > 1 || BlockDim.y > 1) + (GridDim.z > 1 || BlockDim.z > 1);
+  LaunchSizeArgs.NumGroups.x = GridDim.x;
+  LaunchSizeArgs.NumGroups.y = std::max(GridDim.y, 1u);
+  LaunchSizeArgs.NumGroups.z = std::max(GridDim.z, 1u);
+  LaunchSizeArgs.GroupSize.x = BlockDim.x;
+  LaunchSizeArgs.GroupSize.y = std::max(BlockDim.y, 1u);
+  LaunchSizeArgs.GroupSize.z = std::max(BlockDim.z, 1u);
+  LaunchSizeArgs.DynSharedMemory = DynamicSharedMem;
+
+  ol_queue_handle_t Queue = Stream ? reinterpret_cast<ol_queue_handle_t>(Stream)
+                                   : ThreadState::getDefaultQueue();
+
+  struct OffloadKernelArgs {
+    void **Args;
+    size_t NumArgs;
+    size_t *ArgSizes;
+  };
+  OffloadKernelArgs *OKA = static_cast<OffloadKernelArgs *>(KernelArgsPtr);
+
+  return olLaunchKernel(Queue, Device, Kernel, &LaunchSizeArgs,
+                        /*Properties=*/nullptr, OKA->NumArgs, OKA->Args,
+                        OKA->ArgSizes);
+}
+
+unsigned __llvmLaunchKernel(const char *KernelID, dim3 GridDim, dim3 BlockDim,
+                            void *KernelArgsPtr, size_t DynamicSharedMem,
+                            void *Stream) {
+  ol_result_t Result = __llvmLaunchKernelImpl(
+      KernelID, GridDim, BlockDim, KernelArgsPtr, DynamicSharedMem, Stream);
+  return Result ? Result->Code : 0;
+}
+
+} // extern "C"
diff --git a/offload/languages/kernel/src/LanguageRegistration.cpp b/offload/languages/kernel/src/LanguageRegistration.cpp
new file mode 100644
index 0000000000000..2219e16ea0365
--- /dev/null
+++ b/offload/languages/kernel/src/LanguageRegistration.cpp
@@ -0,0 +1,125 @@
+//===-- LanguageRegistration.cpp - Language registration API --------------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#include "LanguageRegistration.h"
+#include "OffloadAPI.h"
+#include "State.h"
+#include "llvm/ADT/StringRef.h"
+#include "llvm/Frontend/Offloading/Utility.h"
+#include "llvm/Support/Error.h"
+#include "llvm/Support/raw_ostream.h"
+#include <cstdio>
+#include <cstring>
+#include <inttypes.h>
+
+using RuntimeState = llvm::offload::StateTy;
+using ThreadState = llvm::offload::ThreadStateTy;
+
+/// Hidden, but exported, Registration API
+///{
+extern "C" {
+
+void __llvmRegisterFunction(const char *Binary, const char *KernelID,
+                            char *KernelName, const char *KernelName1, int,
+                            uint3 *, uint3 *, dim3 *, dim3 *, int *) {
+  ol_symbol_handle_t Kernel;
+  ol_program_handle_t Program = RuntimeState::getProgram(Binary);
+  ol_result_t Result = olGetSymbol(
+      Program, KernelName, ol_symbol_kind_t::OL_SYMBOL_KIND_KERNEL, &Kernel);
+  CHECK_FATAL(Result, "Failed to get kernel symbol for " << KernelName);
+  RuntimeState::registerKernel(KernelID, Kernel);
+}
+
+void __llvmRegisterVar(void **, char *, char *, const char *, int, int, int,
+                       int) {
+  llvm::errs() << "RegisterVar is not implemented!" << "\n";
+}
+
+void __llvmRegisterManagedVar(void **, char *, char *, const char *, size_t,
+                              unsigned) {
+  llvm::errs() << "RegisterManagedVar is not implemented!" << "\n";
+}
+
+void __llvmRegisterSurface(void **, const struct surfaceReference *,
+                           const void **, const char *, int, int) {
+  llvm::errs() << "RegisterSurface is not implemented!" << "\n";
+}
+
+void __llvmRegisterTexture(void **, const struct textureReference *,
+                           const void **, const char *, int, int, int) {
+  llvm::errs() << "RegisterTexture is not implemented!" << "\n";
+}
+
+/// This struct is a record of the device image information
+struct __tgt_device_image {
+  void *ImageStart; // Pointer to the target code start
+  void *ImageEnd;   // Pointer to the target code end
+  llvm::offloading::EntryTy
+      *EntriesBegin; // Begin of table with all target entries
+  llvm::offloading::EntryTy *EntriesEnd; // End of table (non inclusive)
+};
+
+/// This struct is a record of all the host code that may be offloaded to a
+/// target.
+struct __tgt_bin_desc {
+  int32_t NumDeviceImages;          // Number of device types supported
+  __tgt_device_image *DeviceImages; // Array of device images (1 per dev. type)
+  llvm::offloading::EntryTy
+      *HostEntriesBegin; // Begin of table with all host entries
+  llvm::offloading::EntryTy *HostEntriesEnd; // End of table (non inclusive)
+};
+
+void __tgt_register_lib(__tgt_bin_desc *Desc) {
+  // TODO: For each device, lazily.
+  ol_device_handle_t Device = ThreadState::getDefaultDevice();
+
+  for (int32_t I = 0, E = Desc->NumDeviceImages; I < E; ++I) {
+    ol_program_handle_t Program = nullptr;
+
+    __tgt_device_image &DeviceImage = Desc->DeviceImages[I];
+    void *ProgramData = DeviceImage.ImageStart;
+    size_t ProgramSize =
+        (char *)DeviceImage.ImageEnd - (char *)DeviceImage.ImageStart;
+    ol_result_t Result =
+        olCreateProgram(Device, ProgramData, ProgramSize, &Program);
+
+    if (Result && Result->Code) {
+      fprintf(stderr, "Failed to register device code (%i): %s\n", Result->Code,
+              Result->Details);
+      abort();
+    }
+
+    RuntimeState::registerProgram(DeviceImage.ImageStart, Program);
+
+    for (auto *Entry = DeviceImage.EntriesBegin;
+         Entry != DeviceImage.EntriesEnd; ++Entry) {
+      if (!Entry->Size && !Entry->Flags)
+        __llvmRegisterFunction((const char *)DeviceImage.ImageStart,
+                               (const char *)Entry->Address, Entry->SymbolName,
+                               Entry->SymbolName, 0, nullptr, nullptr, nullptr,
+                               nullptr, nullptr);
+    }
+  }
+}
+
+void __tgt_unregister_lib(__tgt_bin_desc *Desc) {
+  for (int32_t I = 0, E = Desc->NumDeviceImages; I < E; ++I) {
+    __tgt_device_image &DeviceImage = Desc->DeviceImages[I];
+    for (auto *Entry = DeviceImage.EntriesBegin;
+         Entry != DeviceImage.EntriesEnd; ++Entry) {
+      if (!Entry->Size && !Entry->Flags)
+        RuntimeState::unregisterKernel((const char *)Entry->Address);
+    }
+
+    if (ol_program_handle_t Program =
+            RuntimeState::unregisterProgram(DeviceImage.ImageStart))
+      olDestroyProgram(Program);
+  }
+}
+}
+///}
diff --git a/offload/languages/kernel/src/LanguageRuntime.cpp b/offload/languages/kernel/src/LanguageRuntime.cpp
new file mode 100644
index 0000000000000..67700f3b02234
--- /dev/null
+++ b/offload/languages/kernel/src/LanguageRuntime.cpp
@@ -0,0 +1,166 @@
+//===-- LanguageRuntime.cpp - Kernel language runtime API implementation --===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#include "LanguageRuntime.h"
+#include <cassert>
+
+#ifndef LANGUAGE
+#error This file should be included, or used, with a LANGUAGE macro set.
+#endif
+
+#include "LanguageUtils.h"
+#include "OffloadAPI.h"
+#include "State.h"
+#include "Types.h"
+
+#include <cstdio>
+#include <cstdlib>
+#include <cstring>
+
+#define STR(X) #X
+#define LANGUAGE_STR STR(LANGUAGE)
+
+using RuntimeState = llvm::offload::StateTy;
+using ThreadState = llvm::offload::ThreadStateTy;
+
+Error_t Malloc(void **DevPtr, size_t Size) {
+  ol_device_handle_t Device = ThreadState::getDefaultDevice();
+  ol_result_t Result = olMemAlloc(Device, OL_ALLOC_TYPE_DEVICE, Size, DevPtr);
+  return convertResult(Result);
+}
+
+Error_t Free(void *DevPtr) {
+  ol_result_t Result = olMemFree(DevPtr);
+  return convertResult(Result);
+}
+
+Error_t Memcpy(void *Dst, const void *Src, size_t Size, MemcpyKind Kind) {
+  ol_queue_handle_t Queue = ThreadState::getDefaultQueue();
+
+  ol_result_t Result;
+  switch (Kind) {
+  case MemcpyHostToHost: {
+    ol_device_handle_t Host = RuntimeState::getHostDevice();
+    Result = olMemcpy(nullptr, Dst, Host, const_cast<void *>(Src), Host, Size);
+    break;
+  }
+  case MemcpyHostToDevice: {
+    ol_device_handle_t Device = ThreadState::getDefaultDevice();
+    ol_device_handle_t Host = RuntimeState::getHostDevice();
+    Result = olMemcpy(Queue, Dst, Device, const_cast<void *>(Src), Host, Size);
+    break;
+  }
+  case MemcpyDeviceToHost: {
+    ol_device_handle_t Device = ThreadState::getDefaultDevice();
+    ol_device_handle_t Host = RuntimeState::getHostDevice();
+
+    Result = olMemcpy(Queue, Dst, Host, const_cast<void *>(Src), Device, Size);
+    break;
+  }
+  case MemcpyDeviceToDevice: {
+    ol_device_handle_t Device = ThreadState::getDefaultDevice();
+
+    Result =
+        olMemcpy(Queue, Dst, Device, const_cast<void *>(Src), Device, Size);
+    break;
+  }
+  case MemcpyDefault:
+    fprintf(stderr, LANGUAGE_STR "MemcpyDefault is not implemented yet");
+    abort();
+  };
+
+  Result = olSyncQueue(Queue);
+
+  return convertResult(Result);
+}
+
+Error_t DeviceSynchronize() {
+  // TODO: This is not correct. We likely want to pipe this through to the
+  // plugins.
+  ol_queue_handle_t Queue = ThreadState::getDefaultQueue();
+  ol_result_t Result = olSyncQueue(Queue);
+  return convertResult(Result);
+}
+
+Error_t GetDevice(int *DeviceNo) {
+  ol_device_handle_t Device = RuntimeState::getDevice(DeviceNo);
+  if (!Device)
+    return ErrorInvalidDevice;
+  return Success;
+}
+
+Error_t GetDeviceCount(int *Count) {
+  *Count = RuntimeState::getDeviceCount();
+  return Success;
+}
+
+Error_t SetDevice(int DeviceNo) {
+  ol_device_handle_t Device = RuntimeState::setDefaultDevice(DeviceNo);
+  if (!Device)
+    return ErrorInvalidDevice;
+  assert(Device == ThreadState::getDefaultDevice() &&
+         "Set Device is not Default Device");
+  return Success;
+}
+
+Error_t HostAlloc(void **Ptr, size_t Size, unsigned int Flags) {
+  ol_device_handle_t Device = ThreadState::getDefaultDevice();
+  ol_result_t Result = olMemAllocHost(Device, Size, Ptr);
+  return convertResult(Result);
+}
+
+Error_t MallocHost(void **Ptr, size_t Size) {
+  return HostAlloc(Ptr, Size, /* HostAllocDefault */ 0);
+}
+
+Error_t FreeHost(void *Ptr) {
+  ol_result_t Result = olMemFree(Ptr);
+  return convertResult(Result);
+}
+
+Error_t GetDeviceProperties(DeviceProp_t *DeviceProp, int DeviceNo) {
+  ol_device_handle_t Device = ThreadState::getDefaultDevice();
+  size_t NameSize = 0;
+  olGetDeviceInfoSize(Device, OL_DEVICE_INFO_NAME, &NameSize);
+  assert(NameSize <= sizeof(DeviceProp->name) &&
+         "Device name is too long for DeviceProp_t");
+  olGetDeviceInfo(Device, OL_DEVICE_INFO_NAME, NameSize, &DeviceProp->name[0]);
+  olGetDeviceInfo(Device, OL_DEVICE_INFO_GLOBAL_MEM_SIZE, sizeof(size_t),
+                  &DeviceProp->totalGlobalMem);
+  olGetDeviceInfo(Device, OL_DEVICE_INFO_NUM_COMPUTE_UNITS, sizeof(uint32_t),
+                  &DeviceProp->multiProcessorCount);
+  olGetDeviceInfo(Device, OL_DEVICE_INFO_NUM_LANES, sizeof(uint32_t),
+                  &DeviceProp->warpSize);
+  return Success;
+}
+
+Error_t StreamCreate(Stream_t *Stream) {
+  ol_queue_handle_t Queue;
+  ol_result_t Result = olCreateQueue(ThreadState::getDefaultDevice(), &Queue);
+  if (Result == OL_SUCCESS)
+    *Stream = reinterpret_cast<Stream_t>(Queue);
+  return convertResult(Result);
+}
+
+Error_t StreamDestroy(Stream_t Stream) {
+  ol_queue_handle_t Queue;
+  Error_t Err = getQueueFromStream(Stream, &Queue);
+  if (Err != Success)
+    return Err;
+  ol_result_t Result = olDestroyQueue(Queue);
+  return convertResult(Result);
+}
+
+Error_t StreamSynchronize(Stream_t Stream) {
+  ol_queue_handle_t Queue;
+  Error_t Err = getQueueFromStream(Stream, &Queue);
+  if (Err != Success)
+    return Err;
+  ol_result_t Result = olSyncQueue(Queue);
+  return convertResult(Result);
+}
diff --git a/offload/languages/kernel/src/LanguageUtils.cpp b/offload/languages/kernel/src/LanguageUtils.cpp
new file mode 100644
index 0000000000000..e514c6445a0a1
--- /dev/null
+++ b/offload/languages/kernel/src/LanguageUtils.cpp
@@ -0,0 +1,43 @@
+//===-- LanguageUtils.cpp - Kernel language utilities ---------------------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#include "LanguageUtils.h"
+#include "LanguageRuntime.h"
+#include "OffloadAPI.h"
+
+const char *GetErrorName(Error_t Error) {
+  switch (Error) {
+#define LLVM_OFFLOAD_STRINGIFY_IMPL(NAME) #NAME
+#define LLVM_OFFLOAD_STRINGIFY(NAME) LLVM_OFFLOAD_STRINGIFY_IMPL(NAME)
+#define LLVM_OFFLOAD_ERR_STR(NAME)                                             \
+  case NAME:                                                                   \
+    return LLVM_OFFLOAD_STRINGIFY(NAME);
+    LLVM_OFFLOAD_ERR_STR(Success)
+    LLVM_OFFLOAD_ERR_STR(ErrorInvalidValue)
+    LLVM_OFFLOAD_ERR_STR(ErrorInvalidDevice)
+#undef LLVM_OFFLOAD_ERR_STR
+#undef LLVM_OFFLOAD_STRINGIFY
+#undef LLVM_OFFLOAD_STRINGIFY_IMPL
+  default:
+    return "Unrecognized error";
+  };
+}
+
+const char *GetErrorString(Error_t Error) {
+  switch (Error) {
+  case Success:
+    return "No error";
+  case ErrorInvalidValue:
+    return "Invalid argument value";
+  case ErrorInvalidDevice:
+    return "Invalid device number";
+  case ErrorUnknown:
+    return "Unknown error";
+  }
+  return "Unrecognized error";
+}
diff --git a/offload/languages/kernel/src/State.cpp b/offload/languages/kernel/src/State.cpp
new file mode 100644
index 0000000000000..9607220ebf3cc
--- /dev/null
+++ b/offload/languages/kernel/src/State.cpp
@@ -0,0 +1,296 @@
+//===-- State.cpp - Kernel language persistent state ----------------------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#include "State.h"
+#include "Types.h"
+
+#include "OffloadAPI.h"
+#include "llvm/ADT/ArrayRef.h"
+#include "llvm/ADT/DenseMap.h"
+#include "llvm/ADT/SmallPtrSet.h"
+#include "llvm/ADT/SmallVector.h"
+
+#include <atomic>
+#include <cassert>
+#include <cstdio>
+#include <mutex>
+
+using namespace llvm;
+using namespace offload;
+
+// Weak so another runtime object can override the default stream mode.
+__attribute__((weak)) uint32_t PerThreadQueue = 0;
+
+// Process-wide singleton and thread-state registry.
+static std::mutex StateLock;
+static std::atomic<StateTy *> StatePtr = nullptr;
+
+static thread_local ThreadStateTy *ThreadState = nullptr;
+
+static std::mutex ThreadStatesLock;
+using ThreadStatesTy = SmallVector<ThreadStateTy *, 64>;
+static ThreadStatesTy *ThreadStatesPtr = nullptr;
+
+static void deleteThreadState() {
+  // Detach the registry before deletion because deleteThreadState may be called
+  // more than once via atexit and StateTy teardown.
+  std::lock_guard<std::mutex> LG(ThreadStatesLock);
+  ThreadStatesTy *ThreadStates = ThreadStatesPtr;
+  ThreadStatesPtr = nullptr;
+  if (!ThreadStates)
+    return;
+
+  for (auto *TS : *ThreadStates)
+    delete TS;
+  delete ThreadStates;
+  ThreadState = nullptr;
+}
+
+static void deleteState() {
+  StateTy *ST = StatePtr.load();
+  StatePtr.store(nullptr);
+  delete ST;
+}
+
+static void destroyQueue(ol_queue_handle_t &Queue) {
+  if (!Queue)
+    return;
+
+  olSyncQueue(Queue);
+  olDestroyQueue(Queue);
+  Queue = nullptr;
+}
+
+namespace llvm {
+namespace offload {
+
+// ThreadStateTy implementation.
+
+ThreadStateTy::ThreadStateTy() {
+  if (PerThreadQueue) [[unlikely]]
+    createDefaultQueue(getDefaultDevice());
+  atexit(deleteThreadState);
+}
+ThreadStateTy::~ThreadStateTy() { destroyQueue(DefaultQueue); }
+
+ThreadStateTy &ThreadStateTy::get() {
+  auto *&TS = ThreadState;
+  if (!TS) {
+    TS = new ThreadStateTy();
+    std::lock_guard<std::mutex> LG(ThreadStatesLock);
+    if (!ThreadStatesPtr)
+      ThreadStatesPtr = new ThreadStatesTy;
+    ThreadStatesPtr->push_back(TS);
+  }
+  return *TS;
+}
+
+ol_device_handle_t ThreadStateTy::getDefaultDevice() {
+  ol_device_handle_t DD = ThreadStateTy::get().DefaultDevice;
+  if (DD)
+    return DD;
+  for (ol_device_handle_t Device : StateTy::get().getDevices()) {
+    DD = Device;
+    break;
+  }
+  return DD;
+}
+
+ol_queue_handle_t ThreadStateTy::getDefaultQueue() {
+  if (!PerThreadQueue) [[likely]]
+    return StateTy::get().DefaultQueue;
+  return ThreadStateTy::get().DefaultQueue;
+}
+
+CallConfigurationTy &ThreadStateTy::getCallConfiguration() {
+  return ThreadStateTy::get().CC;
+}
+
+void ThreadStateTy::setDefaultDevice(ol_device_handle_t Device) {
+  ThreadStateTy &State = get();
+  State.DefaultDevice = Device;
+  State.createDefaultQueue(Device);
+}
+
+void ThreadStateTy::createDefaultQueue(ol_device_handle_t Device) {
+  if (DefaultQueue)
+    olDestroyQueue(DefaultQueue);
+  CHECK_FATAL(olCreateQueue(Device, &DefaultQueue),
+              "Failed to create per-thread default queue");
+}
+
+// StateTy implementation.
+
+StateTy &StateTy::get() {
+  StateTy *ST = StatePtr.load();
+  if (!ST) [[unlikely]] {
+    std::lock_guard<std::mutex> LG(StateLock);
+    ST = StatePtr.load();
+    if (!ST) {
+      ST = new StateTy();
+      StatePtr.store(ST);
+    }
+  }
+  return *ST;
+}
+
+StateTy *StateTy::tryGet() { return StatePtr.load(); }
+
+ol_device_handle_t StateTy::getHostDevice() { return get().HostDevice; }
+
+int StateTy::getDeviceCount() {
+  int DeviceCount = get().getDevices().size();
+  return DeviceCount;
+}
+
+ol_device_handle_t StateTy::getDevice(int *DeviceNo) {
+  ol_device_handle_t DefaultDevice = ThreadStateTy::getDefaultDevice();
+  int DeviceCount = get().getDevices().size();
+  ArrayRef<ol_device_handle_t> Devices = get().getDevices();
+  for (int i = 0; i < DeviceCount; i++) {
+    if (Devices[i] == DefaultDevice) {
+      *DeviceNo = i;
+      return Devices[i];
+    }
+  }
+  return nullptr;
+}
+
+ol_device_handle_t StateTy::setDefaultDevice(int DeviceNo) {
+  ArrayRef<ol_device_handle_t> Devices = get().getDevices();
+  if (DeviceNo < 0 || DeviceNo >= static_cast<int>(Devices.size()))
+    return nullptr;
+  ol_device_handle_t Device = Devices[DeviceNo];
+  ThreadStateTy::setDefaultDevice(Device);
+  return Device;
+}
+
+ArrayRef<ol_device_handle_t> StateTy::getDevices() const { return Devices; }
+
+void StateTy::addDevice(ol_device_handle_t Device) {
+  Devices.push_back(Device);
+}
+
+void StateTy::setHostDevice(ol_device_handle_t Device) {
+  if (!HostDevice)
+    HostDevice = Device;
+}
+
+void StateTy::addKernel(KernelIDTy KernelID, ol_symbol_handle_t Kernel) {
+  KernelMap[KernelID] = Kernel;
+}
+
+void StateTy::removeKernel(KernelIDTy KernelID) { KernelMap.erase(KernelID); }
+
+ol_symbol_handle_t StateTy::lookupKernel(KernelIDTy KernelID) {
+  return KernelMap[KernelID];
+}
+
+void StateTy::registerKernel(const void *ID, ol_symbol_handle_t Kernel) {
+  get().addKernel(ID, Kernel);
+}
+
+void StateTy::unregisterKernel(const void *ID) {
+  if (StateTy *State = tryGet())
+    State->removeKernel(ID);
+}
+
+ol_symbol_handle_t StateTy::getKernel(const void *ID) {
+  return get().lookupKernel(ID);
+}
+
+void StateTy::addProgram(const void *Binary, ol_program_handle_t Program) {
+  BinaryRegisterMap[Binary] = Program;
+}
+
+ol_program_handle_t StateTy::removeProgram(const void *Binary) {
+  auto It = BinaryRegisterMap.find(Binary);
+  if (It == BinaryRegisterMap.end())
+    return nullptr;
+  ol_program_handle_t Program = It->second;
+  BinaryRegisterMap.erase(It);
+  return Program;
+}
+
+ol_program_handle_t StateTy::lookupProgram(const void *Binary) {
+  assert(BinaryRegisterMap.count(Binary) &&
+         "Program not registered for binary");
+  return BinaryRegisterMap[Binary];
+}
+
+void StateTy::registerProgram(const void *ID, ol_program_handle_t Program) {
+  get().addProgram(ID, Program);
+}
+
+ol_program_handle_t StateTy::unregisterProgram(const void *ID) {
+  if (StateTy *State = tryGet())
+    return State->removeProgram(ID);
+  return nullptr;
+}
+
+ol_program_handle_t StateTy::getProgram(const void *ID) {
+  return get().lookupProgram(ID);
+}
+
+bool StateTy::addDevices(ol_device_handle_t Device, void *Payload) {
+  StateTy &State = *reinterpret_cast<StateTy *>(Payload);
+  ol_platform_handle_t Platform;
+  ol_result_t Result;
+
+  Result = olGetDeviceInfo(Device, OL_DEVICE_INFO_PLATFORM, sizeof(Platform),
+                           &Platform);
+  if (Result && Result->Code)
+    return true;
+
+  ol_platform_backend_t Backend;
+  Result = olGetPlatformInfo(Platform, OL_PLATFORM_INFO_BACKEND,
+                             sizeof(Backend), &Backend);
+  if (Result && Result->Code)
+    return true;
+
+  if (Backend == OL_PLATFORM_BACKEND_HOST)
+    State.setHostDevice(Device);
+  else
+    State.addDevice(Device);
+  return true;
+}
+
+StateTy::StateTy() {
+  CHECK_FATAL(olInit(nullptr), "Failed to initialize the LLVMOffload");
+  CHECK_FATAL(olIterateDevices(StateTy::addDevices, this),
+              "Failed to identify devices");
+
+  if (!PerThreadQueue) [[likely]]
+    if (!Devices.empty()) [[likely]]
+      CHECK_FATAL(olCreateQueue(Devices.front(), &DefaultQueue),
+                  "Failed to create default queue");
+
+  atexit(deleteState);
+}
+
+StateTy::~StateTy() {
+  deleteThreadState();
+  destroyQueue(DefaultQueue);
+  destroyRegisteredPrograms();
+  olShutDown();
+}
+
+void StateTy::destroyRegisteredPrograms() {
+  SmallPtrSet<ol_program_handle_t, 8> Programs;
+  for (auto &It : BinaryRegisterMap)
+    Programs.insert(It.second);
+
+  KernelMap.clear();
+  BinaryRegisterMap.clear();
+
+  for (ol_program_handle_t Program : Programs)
+    olDestroyProgram(Program);
+}
+
+} // namespace offload
+} // namespace llvm
diff --git a/offload/test/lit.cfg b/offload/test/lit.cfg
index ace2b1ea8a749..2981cf5d10fea 100644
--- a/offload/test/lit.cfg
+++ b/offload/test/lit.cfg
@@ -83,7 +83,7 @@ def remove_suffix_if_present(name):
 config.name = 'libomptarget :: ' + config.libomptarget_current_target
 
 # suffixes: A list of file extensions to treat as test files.
-config.suffixes = ['.c', '.cpp', '.cc', '.f90', '.cu', '.td']
+config.suffixes = ['.c', '.cpp', '.cc', '.f90', '.cu', '.hip', '.td']
 
 # excludes: A list of directories to exclude from the testuites.
 config.excludes = ['Inputs', 'unit']
@@ -91,6 +91,11 @@ config.excludes = ['Inputs', 'unit']
 # test_source_root: The root path where tests are located.
 config.test_source_root = os.path.dirname(__file__)
 
+# language includes
+config.test_language_includes = os.path.join(config.test_source_root, "../languages/include")
+config.test_language_cuda_includes = os.path.join(config.test_language_includes, "cuda")
+config.test_language_hip_includes = os.path.join(config.test_language_includes, "hip")
+
 # test_exec_root: The root object directory where output is placed
 config.test_exec_root = config.libomptarget_obj_root
 
@@ -100,6 +105,9 @@ config.test_format = lit.formats.ShTest()
 # compiler flags
 config.test_flags = " -I " + config.test_source_root + \
     " -I " + config.omp_header_directory + \
+    " -I " + config.test_language_includes + \
+    " -I " + config.test_language_cuda_includes + \
+    " -I " + config.test_language_hip_includes + \
     " -L " + config.library_dir + \
     " -L " + config.llvm_library_intdir + \
     " -L " + config.llvm_lib_directory
diff --git a/offload/test/offloading/CUDA/basic_launch.cu b/offload/test/offloading/CUDA/basic_launch.cu
index e017241bb9a74..5ecfc3e9d5601 100644
--- a/offload/test/offloading/CUDA/basic_launch.cu
+++ b/offload/test/offloading/CUDA/basic_launch.cu
@@ -7,25 +7,24 @@
 
 // UNSUPPORTED: aarch64-unknown-linux-gnu
 // UNSUPPORTED: x86_64-unknown-linux-gnu
+// UNSUPPORTED: nvptx64-nvidia-cuda-LTO
+// UNSUPPORTED: amdgcn-amd-amdhsa-LTO
+// UNSUPPORTED: amdgpu-amd-amdhsa-LTO
 // UNSUPPORTED: intelgpu
 
 #include <stdio.h>
 
-extern "C" {
-void *llvm_omp_target_alloc_shared(size_t Size, int DeviceNum);
-void llvm_omp_target_free_shared(void *DevicePtr, int DeviceNum);
-}
-
 __global__ void square(int *A) { *A = 42; }
 
 int main(int argc, char **argv) {
-  int DevNo = 0;
-  int *Ptr = reinterpret_cast<int *>(llvm_omp_target_alloc_shared(4, DevNo));
-  *Ptr = 7;
-  printf("Ptr %p, *Ptr: %i\n", Ptr, *Ptr);
-  // CHECK: Ptr [[Ptr:0x.*]], *Ptr: 7
+  int *Ptr;
+  cudaMalloc(&Ptr, 4);
+  printf("Ptr %p\n", Ptr);
+  // CHECK: Ptr [[Ptr:0x.*]]
   square<<<1, 1>>>(Ptr);
-  printf("Ptr %p, *Ptr: %i\n", Ptr, *Ptr);
-  // CHECK: Ptr [[Ptr]], *Ptr: 42
-  llvm_omp_target_free_shared(Ptr, DevNo);
+  int I = 0;
+  cudaDeviceSynchronize();
+  cudaMemcpy(&I, Ptr, sizeof(int), cudaMemcpyDeviceToHost);
+  printf("I: %i\n", I);
+  // CHECK: I: 42
 }
diff --git a/offload/test/offloading/CUDA/basic_launch_blocks_and_threads.cu b/offload/test/offloading/CUDA/basic_launch_blocks_and_threads.cu
index a428e25d82359..0bfc9e231ffde 100644
--- a/offload/test/offloading/CUDA/basic_launch_blocks_and_threads.cu
+++ b/offload/test/offloading/CUDA/basic_launch_blocks_and_threads.cu
@@ -7,27 +7,25 @@
 
 // UNSUPPORTED: aarch64-unknown-linux-gnu
 // UNSUPPORTED: x86_64-unknown-linux-gnu
+// UNSUPPORTED: nvptx64-nvidia-cuda-LTO
+// UNSUPPORTED: amdgcn-amd-amdhsa-LTO
+// UNSUPPORTED: amdgpu-amd-amdhsa-LTO
 // UNSUPPORTED: intelgpu
 
 #include <stdio.h>
 
-extern "C" {
-void *llvm_omp_target_alloc_shared(size_t Size, int DeviceNum);
-void llvm_omp_target_free_shared(void *DevicePtr, int DeviceNum);
-}
-
 __global__ void square(int *A) {
   __scoped_atomic_fetch_add(A, 1, __ATOMIC_SEQ_CST, __MEMORY_SCOPE_DEVICE);
 }
 
 int main(int argc, char **argv) {
   int DevNo = 0;
-  int *Ptr = reinterpret_cast<int *>(llvm_omp_target_alloc_shared(4, DevNo));
-  *Ptr = 0;
-  printf("Ptr %p, *Ptr: %i\n", Ptr, *Ptr);
-  // CHECK: Ptr [[Ptr:0x.*]], *Ptr: 0
+  int *Ptr, I;
+  cudaMalloc(&Ptr, 4);
+  printf("Ptr %p\n", Ptr);
+  // CHECK: Ptr [[Ptr:0x.*]]
   square<<<7, 6>>>(Ptr);
-  printf("Ptr %p, *Ptr: %i\n", Ptr, *Ptr);
-  // CHECK: Ptr [[Ptr]], *Ptr: 42
-  llvm_omp_target_free_shared(Ptr, DevNo);
+  cudaMemcpy(&I, Ptr, sizeof(int), cudaMemcpyDeviceToHost);
+  printf("I: %i\n", I);
+  // CHECK: I: 42
 }
diff --git a/offload/test/offloading/CUDA/basic_launch_multi_arg.cu b/offload/test/offloading/CUDA/basic_launch_multi_arg.cu
index db2a1e48371b0..7c0c755e71057 100644
--- a/offload/test/offloading/CUDA/basic_launch_multi_arg.cu
+++ b/offload/test/offloading/CUDA/basic_launch_multi_arg.cu
@@ -6,15 +6,13 @@
 // clang-format on
 
 // REQUIRES: gpu
+// UNSUPPORTED: nvptx64-nvidia-cuda-LTO
+// UNSUPPORTED: amdgcn-amd-amdhsa-LTO
+// UNSUPPORTED: amdgpu-amd-amdhsa-LTO
 // UNSUPPORTED: intelgpu
 
 #include <stdio.h>
 
-extern "C" {
-void *llvm_omp_target_alloc_shared(size_t Size, int DeviceNum);
-void llvm_omp_target_free_shared(void *DevicePtr, int DeviceNum);
-}
-
 __global__ void square(int *Dst, short Q, int *Src, short P) {
   *Dst = (Src[0] + Src[1]) * (Q + P);
   Src[0] = Q;
@@ -23,19 +21,19 @@ __global__ void square(int *Dst, short Q, int *Src, short P) {
 
 int main(int argc, char **argv) {
   int DevNo = 0;
-  int *Ptr = reinterpret_cast<int *>(llvm_omp_target_alloc_shared(4, DevNo));
-  int *Src = reinterpret_cast<int *>(llvm_omp_target_alloc_shared(8, DevNo));
-  *Ptr = 7;
-  Src[0] = -2;
-  Src[1] = 8;
-  printf("Ptr %p, *Ptr: %i\n", Ptr, *Ptr);
-  // CHECK: Ptr [[Ptr:0x.*]], *Ptr: 7
-  printf("Src: %i : %i\n", Src[0], Src[1]);
-  // CHECK: Src: -2 : 8
+  int *Src, *Ptr;
+  cudaMalloc(&Ptr, 4);
+  cudaMalloc(&Src, 8);
+
+  int I = 7;
+  int HostSrc[2] = {-2,8};
+  cudaMemcpy(Ptr, &I, sizeof(int), cudaMemcpyHostToDevice);
+  cudaMemcpy(Src, &HostSrc[0], 2*sizeof(int), cudaMemcpyHostToDevice);
   square<<<1, 1>>>(Ptr, 3, Src, 4);
-  printf("Ptr %p, *Ptr: %i\n", Ptr, *Ptr);
-  // CHECK: Ptr [[Ptr]], *Ptr: 42
-  printf("Src: %i : %i\n", Src[0], Src[1]);
-  // CHECK: Src: 3 : 4
-  llvm_omp_target_free_shared(Ptr, DevNo);
+  cudaMemcpy(&I, Ptr, sizeof(int), cudaMemcpyDeviceToHost);
+  cudaMemcpy(&HostSrc[0], Src, 2 * sizeof(int), cudaMemcpyDeviceToHost);
+  printf("I: %i\n", I);
+  // CHECK: I: 42
+  printf("Src: %i, %i\n", HostSrc[0], HostSrc[1]);
+  // CHECK: Src: 3, 4
 }
diff --git a/offload/test/offloading/CUDA/device_api.cu b/offload/test/offloading/CUDA/device_api.cu
new file mode 100644
index 0000000000000..af2b046eee397
--- /dev/null
+++ b/offload/test/offloading/CUDA/device_api.cu
@@ -0,0 +1,45 @@
+// clang-format off
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t
+// RUN: %t | %fcheck-generic
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t -fopenmp
+// RUN: %t | %fcheck-generic
+// clang-format on
+
+// UNSUPPORTED: aarch64-unknown-linux-gnu
+// UNSUPPORTED: x86_64-unknown-linux-gnu
+// UNSUPPORTED: nvptx64-nvidia-cuda-LTO
+// UNSUPPORTED: amdgcn-amd-amdhsa-LTO
+// UNSUPPORTED: amdgpu-amd-amdhsa-LTO
+// UNSUPPORTED: intelgpu
+
+#include <stdio.h>
+
+int main(int argc, char **argv) {
+  int Count = 0;
+  if (cudaGetDeviceCount(&Count) != cudaSuccess)
+    return 1;
+
+  printf("device count: %d\n", Count);
+  // CHECK: device count: {{[1-9][0-9]*}}
+
+  int Device = -1;
+  if (cudaGetDevice(&Device) != cudaSuccess)
+    return 1;
+
+  printf("device: %d\n", Device);
+  // CHECK: device: {{[0-9]+}}
+
+  if (cudaSetDevice(Device) != cudaSuccess)
+    return 1;
+
+  int After = -1;
+  if (cudaGetDevice(&After) != cudaSuccess)
+    return 1;
+
+  printf("device after set: %d\n", After);
+  // CHECK: device after set: {{[0-9]+}}
+
+  cudaError_t Err = cudaSetDevice(-1);
+  printf("set invalid device: %u\n", Err);
+  // CHECK: set invalid device: 2
+}
diff --git a/offload/test/offloading/CUDA/device_properties.cu b/offload/test/offloading/CUDA/device_properties.cu
new file mode 100644
index 0000000000000..8f625f6ccabe1
--- /dev/null
+++ b/offload/test/offloading/CUDA/device_properties.cu
@@ -0,0 +1,40 @@
+// clang-format off
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t
+// RUN: %t | %fcheck-generic
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t -fopenmp
+// RUN: %t | %fcheck-generic
+// clang-format on
+
+// UNSUPPORTED: aarch64-unknown-linux-gnu
+// UNSUPPORTED: x86_64-unknown-linux-gnu
+// UNSUPPORTED: nvptx64-nvidia-cuda-LTO
+// UNSUPPORTED: amdgcn-amd-amdhsa-LTO
+// UNSUPPORTED: amdgpu-amd-amdhsa-LTO
+// UNSUPPORTED: intelgpu
+
+#include <stdio.h>
+
+int main(int argc, char **argv) {
+  cudaDeviceProp Prop = {};
+  cudaError_t Err = cudaGetDeviceProperties(&Prop, 0);
+  if (Err != cudaSuccess) {
+    printf("cudaGetDeviceProperties failed: %u\n", Err);
+    return 1;
+  }
+
+  printf("Device name: %s\n", Prop.name);
+  // CHECK: Device name:
+  printf("Total global memory: %zu\n", Prop.totalGlobalMem);
+  // CHECK: Total global memory:
+  printf("Multiprocessors: %i\n", Prop.multiProcessorCount);
+  // CHECK: Multiprocessors:
+  printf("Warp size: %i\n", Prop.warpSize);
+  // CHECK: Warp size:
+
+  if (!Prop.name[0] || !Prop.totalGlobalMem || !Prop.multiProcessorCount ||
+      !Prop.warpSize)
+    return 1;
+
+  printf("Device properties are populated.\n");
+  // CHECK: Device properties are populated.
+}
diff --git a/offload/test/offloading/CUDA/error_kinds.cu b/offload/test/offloading/CUDA/error_kinds.cu
new file mode 100644
index 0000000000000..5fa3035fe52f0
--- /dev/null
+++ b/offload/test/offloading/CUDA/error_kinds.cu
@@ -0,0 +1,61 @@
+// clang-format off
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t
+// RUN: %t | %fcheck-generic
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t -fopenmp
+// RUN: %t | %fcheck-generic
+// clang-format on
+
+// UNSUPPORTED: aarch64-unknown-linux-gnu
+// UNSUPPORTED: x86_64-unknown-linux-gnu
+// UNSUPPORTED: nvptx64-nvidia-cuda-LTO
+// UNSUPPORTED: amdgcn-amd-amdhsa-LTO
+// UNSUPPORTED: amdgpu-amd-amdhsa-LTO
+// UNSUPPORTED: intelgpu
+
+#include <stdio.h>
+
+static void print_error(const char *Label, cudaError_t Error) {
+  printf("%s value: %u\n", Label, static_cast<unsigned>(Error));
+  printf("%s name: %s\n", Label, cudaGetErrorName(Error));
+  printf("%s string: %s\n", Label, cudaGetErrorString(Error));
+}
+
+int main() {
+  print_error("success", cudaSuccess);
+  // CHECK: success value: 0
+  // CHECK: success name: cudaSuccess
+  // CHECK: success string: No error
+
+  print_error("invalid value", cudaErrorInvalidValue);
+  // CHECK: invalid value value: 1
+  // CHECK: invalid value name: cudaErrorInvalidValue
+  // CHECK: invalid value string: Invalid argument value
+
+  print_error("invalid device", cudaErrorInvalidDevice);
+  // CHECK: invalid device value: 2
+  // CHECK: invalid device name: cudaErrorInvalidDevice
+  // CHECK: invalid device string: Invalid device number
+
+  print_error("unknown", cudaErrorUnknown);
+  // CHECK: unknown value: 3
+  // CHECK: unknown name: Unrecognized error
+  // CHECK: unknown string: Unknown error
+
+  cudaError_t Unrecognized = static_cast<cudaError_t>(999);
+  print_error("unrecognized", Unrecognized);
+  // CHECK: unrecognized value: 999
+  // CHECK: unrecognized name: Unrecognized error
+  // CHECK: unrecognized string: Unrecognized error
+
+  print_error("set invalid device", cudaSetDevice(-1));
+  // CHECK: set invalid device value: 2
+  // CHECK: set invalid device name: cudaErrorInvalidDevice
+  // CHECK: set invalid device string: Invalid device number
+
+  print_error("null stream destroy", cudaStreamDestroy(nullptr));
+  // CHECK: null stream destroy value: 1
+  // CHECK: null stream destroy name: cudaErrorInvalidValue
+  // CHECK: null stream destroy string: Invalid argument value
+
+  return 0;
+}
diff --git a/offload/test/offloading/CUDA/host_alloc.cu b/offload/test/offloading/CUDA/host_alloc.cu
new file mode 100644
index 0000000000000..f23eac7604582
--- /dev/null
+++ b/offload/test/offloading/CUDA/host_alloc.cu
@@ -0,0 +1,40 @@
+// clang-format off
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t
+// RUN: %t | %fcheck-generic
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t -fopenmp
+// RUN: %t | %fcheck-generic
+// clang-format on
+
+// UNSUPPORTED: aarch64-unknown-linux-gnu
+// UNSUPPORTED: x86_64-unknown-linux-gnu
+// UNSUPPORTED: nvptx64-nvidia-cuda-LTO
+// UNSUPPORTED: amdgcn-amd-amdhsa-LTO
+// UNSUPPORTED: amdgpu-amd-amdhsa-LTO
+// UNSUPPORTED: intelgpu
+
+#include <stdio.h>
+
+int main(int argc, char **argv) {
+  int *HostAllocPtr = nullptr;
+  if (cudaHostAlloc(&HostAllocPtr, sizeof(int), cudaHostAllocDefault) !=
+      cudaSuccess)
+    return 1;
+
+  *HostAllocPtr = 17;
+  printf("cudaHostAlloc value: %d\n", *HostAllocPtr);
+  // CHECK: cudaHostAlloc value: 17
+
+  if (cudaFreeHost(HostAllocPtr) != cudaSuccess)
+    return 1;
+
+  int *MallocHostPtr = nullptr;
+  if (cudaMallocHost(&MallocHostPtr, sizeof(int)) != cudaSuccess)
+    return 1;
+
+  *MallocHostPtr = 23;
+  printf("cudaMallocHost value: %d\n", *MallocHostPtr);
+  // CHECK: cudaMallocHost value: 23
+
+  if (cudaFreeHost(MallocHostPtr) != cudaSuccess)
+    return 1;
+}
diff --git a/offload/test/offloading/CUDA/launch_tu.cu b/offload/test/offloading/CUDA/launch_tu.cu
index a46472b514a6c..8b92194ba435e 100644
--- a/offload/test/offloading/CUDA/launch_tu.cu
+++ b/offload/test/offloading/CUDA/launch_tu.cu
@@ -1,31 +1,30 @@
 // clang-format off
 // RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t.launch_tu.o -c
 // RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native -x cuda %S/kernel_tu.cu.inc -o %t.kernel_tu.o -c
-// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %t.launch_tu.o %t.kernel_tu.o -o %t
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native --offload-link %t.launch_tu.o %t.kernel_tu.o -o %t
 // RUN: %t | %fcheck-generic
 // clang-format on
 
 // UNSUPPORTED: aarch64-unknown-linux-gnu
 // UNSUPPORTED: x86_64-unknown-linux-gnu
+// UNSUPPORTED: nvptx64-nvidia-cuda-LTO
+// UNSUPPORTED: amdgcn-amd-amdhsa-LTO
+// UNSUPPORTED: amdgpu-amd-amdhsa-LTO
 // UNSUPPORTED: intelgpu
 
 #include <stdio.h>
 
-extern "C" {
-void *llvm_omp_target_alloc_shared(size_t Size, int DeviceNum);
-void llvm_omp_target_free_shared(void *DevicePtr, int DeviceNum);
-}
-
 extern __global__ void square(int *A);
 
 int main(int argc, char **argv) {
   int DevNo = 0;
-  int *Ptr = reinterpret_cast<int *>(llvm_omp_target_alloc_shared(4, DevNo));
-  *Ptr = 7;
-  printf("Ptr %p, *Ptr: %i\n", Ptr, *Ptr);
-  // CHECK: Ptr [[Ptr:0x.*]], *Ptr: 7
+  int *Ptr;
+  cudaMalloc(&Ptr, 4);
+  printf("Ptr %p\n", Ptr);
+  // CHECK: Ptr [[Ptr:0x.*]]
   square<<<1, 1>>>(Ptr);
-  printf("Ptr %p, *Ptr: %i\n", Ptr, *Ptr);
-  // CHECK: Ptr [[Ptr]], *Ptr: 42
-  llvm_omp_target_free_shared(Ptr, DevNo);
+  int I;
+  cudaMemcpy(&I, Ptr, sizeof(int), cudaMemcpyDeviceToHost);
+  printf("I: %i\n", I);
+  // CHECK: I: 42
 }
diff --git a/offload/test/offloading/CUDA/memcpy_kinds.cu b/offload/test/offloading/CUDA/memcpy_kinds.cu
new file mode 100644
index 0000000000000..a4288ee51ee3b
--- /dev/null
+++ b/offload/test/offloading/CUDA/memcpy_kinds.cu
@@ -0,0 +1,51 @@
+// clang-format off
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t
+// RUN: %t | %fcheck-generic
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t -fopenmp
+// RUN: %t | %fcheck-generic
+// clang-format on
+
+// UNSUPPORTED: aarch64-unknown-linux-gnu
+// UNSUPPORTED: x86_64-unknown-linux-gnu
+// UNSUPPORTED: nvptx64-nvidia-cuda-LTO
+// UNSUPPORTED: amdgcn-amd-amdhsa-LTO
+// UNSUPPORTED: amdgpu-amd-amdhsa-LTO
+// UNSUPPORTED: intelgpu
+
+#include <stdio.h>
+
+int main(int argc, char **argv) {
+  int HostSrc = 11;
+  int HostDst = 0;
+  if (cudaMemcpy(&HostDst, &HostSrc, sizeof(int), cudaMemcpyHostToHost) !=
+      cudaSuccess)
+    return 1;
+
+  printf("host to host: %d\n", HostDst);
+  // CHECK: host to host: 11
+
+  int *DevSrc = nullptr;
+  int *DevDst = nullptr;
+  int Result = 0;
+  if (cudaMalloc(&DevSrc, sizeof(int)) != cudaSuccess)
+    return 1;
+  if (cudaMalloc(&DevDst, sizeof(int)) != cudaSuccess)
+    return 1;
+
+  HostSrc = 42;
+  if (cudaMemcpy(DevSrc, &HostSrc, sizeof(int), cudaMemcpyHostToDevice) !=
+      cudaSuccess)
+    return 1;
+  if (cudaMemcpy(DevDst, DevSrc, sizeof(int), cudaMemcpyDeviceToDevice) !=
+      cudaSuccess)
+    return 1;
+  if (cudaMemcpy(&Result, DevDst, sizeof(int), cudaMemcpyDeviceToHost) !=
+      cudaSuccess)
+    return 1;
+
+  printf("device to device: %d\n", Result);
+  // CHECK: device to device: 42
+
+  cudaFree(DevSrc);
+  cudaFree(DevDst);
+}
diff --git a/offload/test/offloading/CUDA/stream_api.cu b/offload/test/offloading/CUDA/stream_api.cu
new file mode 100644
index 0000000000000..7202751f8207e
--- /dev/null
+++ b/offload/test/offloading/CUDA/stream_api.cu
@@ -0,0 +1,46 @@
+// clang-format off
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t
+// RUN: %t | %fcheck-generic
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t -fopenmp
+// RUN: %t | %fcheck-generic
+// clang-format on
+
+// UNSUPPORTED: aarch64-unknown-linux-gnu
+// UNSUPPORTED: x86_64-unknown-linux-gnu
+// UNSUPPORTED: nvptx64-nvidia-cuda-LTO
+// UNSUPPORTED: amdgcn-amd-amdhsa-LTO
+// UNSUPPORTED: amdgpu-amd-amdhsa-LTO
+// UNSUPPORTED: intelgpu
+
+#include <stdio.h>
+
+__global__ void setValue(int *Out) { *Out = 42; }
+
+int main(int argc, char **argv) {
+  cudaStream_t Stream = nullptr;
+  if (cudaStreamCreate(&Stream) != cudaSuccess)
+    return 1;
+
+  printf("stream created: %d\n", Stream != nullptr);
+  // CHECK: stream created: 1
+
+  int *DevPtr = nullptr;
+  int Result = 0;
+  if (cudaMalloc(&DevPtr, sizeof(int)) != cudaSuccess)
+    return 1;
+
+  setValue<<<1, 1, 0, Stream>>>(DevPtr);
+
+  if (cudaStreamSynchronize(Stream) != cudaSuccess)
+    return 1;
+  if (cudaMemcpy(&Result, DevPtr, sizeof(int), cudaMemcpyDeviceToHost) !=
+      cudaSuccess)
+    return 1;
+
+  printf("stream result: %d\n", Result);
+  // CHECK: stream result: 42
+
+  if (cudaStreamDestroy(Stream) != cudaSuccess)
+    return 1;
+  cudaFree(DevPtr);
+}
diff --git a/offload/test/offloading/CUDA/syncthreads.cu b/offload/test/offloading/CUDA/syncthreads.cu
new file mode 100644
index 0000000000000..0c6048c32f824
--- /dev/null
+++ b/offload/test/offloading/CUDA/syncthreads.cu
@@ -0,0 +1,40 @@
+// clang-format off
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t
+// RUN: %t | %fcheck-generic
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t -fopenmp
+// RUN: %t | %fcheck-generic
+// clang-format on
+
+// UNSUPPORTED: aarch64-unknown-linux-gnu
+// UNSUPPORTED: x86_64-unknown-linux-gnu
+// UNSUPPORTED: nvptx64-nvidia-cuda-LTO
+// UNSUPPORTED: amdgcn-amd-amdhsa-LTO
+// UNSUPPORTED: amdgpu-amd-amdhsa-LTO
+// UNSUPPORTED: intelgpu
+
+#include <stdio.h>
+
+__global__ void reduceBlock(int *Out) {
+  __shared__ int Scratch[64];
+  int Tid = threadIdx.x;
+  Scratch[Tid] = Tid;
+  __syncthreads();
+
+  if (Tid == 0) {
+    int Sum = 0;
+    for (int I = 0; I < 64; ++I)
+      Sum += Scratch[I];
+    Out[0] = Sum;
+  }
+}
+
+int main(int argc, char **argv) {
+  int *DevPtr;
+  int Result = 0;
+  cudaMalloc(&DevPtr, sizeof(int));
+  reduceBlock<<<1, 64>>>(DevPtr);
+  cudaMemcpy(&Result, DevPtr, sizeof(int), cudaMemcpyDeviceToHost);
+
+  printf("sum: %i\n", Result);
+  // CHECK: sum: 2016
+}
diff --git a/offload/test/offloading/CUDA/thread_and_block_id.cu b/offload/test/offloading/CUDA/thread_and_block_id.cu
new file mode 100644
index 0000000000000..074c7e6557d06
--- /dev/null
+++ b/offload/test/offloading/CUDA/thread_and_block_id.cu
@@ -0,0 +1,44 @@
+// clang-format off
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t
+// RUN: %t | %fcheck-generic
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t -fopenmp 
+// RUN: %t | %fcheck-generic
+// clang-format on
+
+// UNSUPPORTED: aarch64-unknown-linux-gnu
+// UNSUPPORTED: aarch64-unknown-linux-gnu-LTO
+// UNSUPPORTED: x86_64-unknown-linux-gnu
+// UNSUPPORTED: x86_64-unknown-linux-gnu-LTO
+// UNSUPPORTED: nvptx64-nvidia-cuda-LTO
+// UNSUPPORTED: amdgcn-amd-amdhsa-LTO
+// UNSUPPORTED: amdgpu-amd-amdhsa-LTO
+
+#include <stdio.h>
+#include <stdlib.h>
+
+__global__ void fill(int *A) {
+  int tid = threadIdx.x + blockDim.x * blockIdx.x;
+  A[tid] = 42;
+}
+
+int main(int argc, char **argv) {
+  int NThreads = 128;
+  int NBlocks = 512;
+  int Size = sizeof(int) * NThreads * NBlocks;
+  int *Ptr = (int*)calloc(1, Size);
+  int *DevPtr;
+  cudaMalloc(&DevPtr, Size);
+  cudaMemcpy(DevPtr, Ptr, Size, cudaMemcpyHostToDevice);
+  printf("DevPtr %p\n", DevPtr);
+  // CHECK: DevPtr [[DevPtr:0x.*]]
+  fill<<<NBlocks, NThreads>>>(DevPtr);
+  cudaMemcpy(Ptr, DevPtr, Size, cudaMemcpyDeviceToHost);
+
+  for (int I = 0; I < NBlocks * NThreads; ++I) {
+    if (Ptr[I] == 42)
+      continue;
+    printf("Error at %i: %i vs %i\n", I, Ptr[I], 42);
+    return 1;
+  }
+  return 0;
+}
diff --git a/offload/test/offloading/HIP/basic_launch.hip b/offload/test/offloading/HIP/basic_launch.hip
new file mode 100644
index 0000000000000..bd2f2a6078671
--- /dev/null
+++ b/offload/test/offloading/HIP/basic_launch.hip
@@ -0,0 +1,30 @@
+// clang-format off
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t
+// RUN: %t | %fcheck-generic
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t -fopenmp 
+// RUN: %t | %fcheck-generic
+// clang-format on
+
+// UNSUPPORTED: aarch64-unknown-linux-gnu
+// UNSUPPORTED: x86_64-unknown-linux-gnu
+// UNSUPPORTED: nvptx64-nvidia-cuda-LTO
+// UNSUPPORTED: amdgcn-amd-amdhsa-LTO
+// UNSUPPORTED: amdgpu-amd-amdhsa-LTO
+// UNSUPPORTED: intelgpu
+
+#include <stdio.h>
+
+__global__ void square(int *A) { *A = 42; }
+
+int main(int argc, char **argv) {
+  int *Ptr;
+  hipMalloc(&Ptr, 4);
+  printf("Ptr %p\n", Ptr);
+  // CHECK: Ptr [[Ptr:0x.*]]
+  square<<<1, 1>>>(Ptr);
+  int I = 0;
+  hipDeviceSynchronize();
+  hipMemcpy(&I, Ptr, sizeof(int), hipMemcpyDeviceToHost);
+  printf("I: %i\n", I);
+  // CHECK: I: 42
+}
diff --git a/offload/test/offloading/HIP/basic_launch_blocks_and_threads.hip b/offload/test/offloading/HIP/basic_launch_blocks_and_threads.hip
new file mode 100644
index 0000000000000..344b98b1636f1
--- /dev/null
+++ b/offload/test/offloading/HIP/basic_launch_blocks_and_threads.hip
@@ -0,0 +1,31 @@
+// clang-format off
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t
+// RUN: %t | %fcheck-generic
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t -fopenmp 
+// RUN: %t | %fcheck-generic
+// clang-format on
+
+// UNSUPPORTED: aarch64-unknown-linux-gnu
+// UNSUPPORTED: x86_64-unknown-linux-gnu
+// UNSUPPORTED: nvptx64-nvidia-cuda-LTO
+// UNSUPPORTED: amdgcn-amd-amdhsa-LTO
+// UNSUPPORTED: amdgpu-amd-amdhsa-LTO
+// UNSUPPORTED: intelgpu
+
+#include <stdio.h>
+
+__global__ void square(int *A) {
+  __scoped_atomic_fetch_add(A, 1, __ATOMIC_SEQ_CST, __MEMORY_SCOPE_DEVICE);
+}
+
+int main(int argc, char **argv) {
+  int DevNo = 0;
+  int *Ptr, I;
+  hipMalloc(&Ptr, 4);
+  printf("Ptr %p\n", Ptr);
+  // CHECK: Ptr [[Ptr:0x.*]]
+  square<<<7, 6>>>(Ptr);
+  hipMemcpy(&I, Ptr, sizeof(int), hipMemcpyDeviceToHost);
+  printf("I: %i\n", I);
+  // CHECK: I: 42
+}
diff --git a/offload/test/offloading/HIP/basic_launch_multi_arg.hip b/offload/test/offloading/HIP/basic_launch_multi_arg.hip
new file mode 100644
index 0000000000000..6e599d6704598
--- /dev/null
+++ b/offload/test/offloading/HIP/basic_launch_multi_arg.hip
@@ -0,0 +1,39 @@
+// clang-format off
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t
+// RUN: %t | %fcheck-generic
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t -fopenmp 
+// RUN: %t | %fcheck-generic
+// clang-format on
+
+// REQUIRES: gpu
+// UNSUPPORTED: nvptx64-nvidia-cuda-LTO
+// UNSUPPORTED: amdgcn-amd-amdhsa-LTO
+// UNSUPPORTED: amdgpu-amd-amdhsa-LTO
+// UNSUPPORTED: intelgpu
+
+#include <stdio.h>
+
+__global__ void square(int *Dst, short Q, int *Src, short P) {
+  *Dst = (Src[0] + Src[1]) * (Q + P);
+  Src[0] = Q;
+  Src[1] = P;
+}
+
+int main(int argc, char **argv) {
+  int DevNo = 0;
+  int *Src, *Ptr;
+  hipMalloc(&Ptr, 4);
+  hipMalloc(&Src, 8);
+
+  int I = 7;
+  int HostSrc[2] = {-2,8};
+  hipMemcpy(Ptr, &I, sizeof(int), hipMemcpyHostToDevice);
+  hipMemcpy(Src, &HostSrc[0], 2*sizeof(int), hipMemcpyHostToDevice);
+  square<<<1, 1>>>(Ptr, 3, Src, 4);
+  hipMemcpy(&I, Ptr, sizeof(int), hipMemcpyDeviceToHost);
+  hipMemcpy(&HostSrc[0], Src, 2 * sizeof(int), hipMemcpyDeviceToHost);
+  printf("I: %i\n", I);
+  // CHECK: I: 42
+  printf("Src: %i, %i\n", HostSrc[0], HostSrc[1]);
+  // CHECK: Src: 3, 4
+}
diff --git a/offload/test/offloading/HIP/device_api.hip b/offload/test/offloading/HIP/device_api.hip
new file mode 100644
index 0000000000000..5fb66e6e45eeb
--- /dev/null
+++ b/offload/test/offloading/HIP/device_api.hip
@@ -0,0 +1,45 @@
+// clang-format off
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t
+// RUN: %t | %fcheck-generic
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t -fopenmp
+// RUN: %t | %fcheck-generic
+// clang-format on
+
+// UNSUPPORTED: aarch64-unknown-linux-gnu
+// UNSUPPORTED: x86_64-unknown-linux-gnu
+// UNSUPPORTED: nvptx64-nvidia-cuda-LTO
+// UNSUPPORTED: amdgcn-amd-amdhsa-LTO
+// UNSUPPORTED: amdgpu-amd-amdhsa-LTO
+// UNSUPPORTED: intelgpu
+
+#include <stdio.h>
+
+int main(int argc, char **argv) {
+  int Count = 0;
+  if (hipGetDeviceCount(&Count) != hipSuccess)
+    return 1;
+
+  printf("device count: %d\n", Count);
+  // CHECK: device count: {{[1-9][0-9]*}}
+
+  int Device = -1;
+  if (hipGetDevice(&Device) != hipSuccess)
+    return 1;
+
+  printf("device: %d\n", Device);
+  // CHECK: device: {{[0-9]+}}
+
+  if (hipSetDevice(Device) != hipSuccess)
+    return 1;
+
+  int After = -1;
+  if (hipGetDevice(&After) != hipSuccess)
+    return 1;
+
+  printf("device after set: %d\n", After);
+  // CHECK: device after set: {{[0-9]+}}
+
+  hipError_t Err = hipSetDevice(-1);
+  printf("set invalid device: %u\n", Err);
+  // CHECK: set invalid device: 2
+}
diff --git a/offload/test/offloading/HIP/device_properties.hip b/offload/test/offloading/HIP/device_properties.hip
new file mode 100644
index 0000000000000..1a9b9a70f8ea9
--- /dev/null
+++ b/offload/test/offloading/HIP/device_properties.hip
@@ -0,0 +1,40 @@
+// clang-format off
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t
+// RUN: %t | %fcheck-generic
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t -fopenmp
+// RUN: %t | %fcheck-generic
+// clang-format on
+
+// UNSUPPORTED: aarch64-unknown-linux-gnu
+// UNSUPPORTED: x86_64-unknown-linux-gnu
+// UNSUPPORTED: nvptx64-nvidia-cuda-LTO
+// UNSUPPORTED: amdgcn-amd-amdhsa-LTO
+// UNSUPPORTED: amdgpu-amd-amdhsa-LTO
+// UNSUPPORTED: intelgpu
+
+#include <stdio.h>
+
+int main(int argc, char **argv) {
+  hipDeviceProp_t Prop = {};
+  hipError_t Err = hipGetDeviceProperties(&Prop, 0);
+  if (Err != hipSuccess) {
+    printf("hipGetDeviceProperties failed: %u\n", Err);
+    return 1;
+  }
+
+  printf("Device name: %s\n", Prop.name);
+  // CHECK: Device name:
+  printf("Total global memory: %zu\n", Prop.totalGlobalMem);
+  // CHECK: Total global memory:
+  printf("Multiprocessors: %i\n", Prop.multiProcessorCount);
+  // CHECK: Multiprocessors:
+  printf("Warp size: %i\n", Prop.warpSize);
+  // CHECK: Warp size:
+
+  if (!Prop.name[0] || !Prop.totalGlobalMem || !Prop.multiProcessorCount ||
+      !Prop.warpSize)
+    return 1;
+
+  printf("Device properties are populated.\n");
+  // CHECK: Device properties are populated.
+}
diff --git a/offload/test/offloading/HIP/error_kinds.hip b/offload/test/offloading/HIP/error_kinds.hip
new file mode 100644
index 0000000000000..af760a4cd4399
--- /dev/null
+++ b/offload/test/offloading/HIP/error_kinds.hip
@@ -0,0 +1,61 @@
+// clang-format off
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t
+// RUN: %t | %fcheck-generic
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t -fopenmp
+// RUN: %t | %fcheck-generic
+// clang-format on
+
+// UNSUPPORTED: aarch64-unknown-linux-gnu
+// UNSUPPORTED: x86_64-unknown-linux-gnu
+// UNSUPPORTED: nvptx64-nvidia-cuda-LTO
+// UNSUPPORTED: amdgcn-amd-amdhsa-LTO
+// UNSUPPORTED: amdgpu-amd-amdhsa-LTO
+// UNSUPPORTED: intelgpu
+
+#include <stdio.h>
+
+static void print_error(const char *Label, hipError_t Error) {
+  printf("%s value: %u\n", Label, static_cast<unsigned>(Error));
+  printf("%s name: %s\n", Label, hipGetErrorName(Error));
+  printf("%s string: %s\n", Label, hipGetErrorString(Error));
+}
+
+int main() {
+  print_error("success", hipSuccess);
+  // CHECK: success value: 0
+  // CHECK: success name: hipSuccess
+  // CHECK: success string: No error
+
+  print_error("invalid value", hipErrorInvalidValue);
+  // CHECK: invalid value value: 1
+  // CHECK: invalid value name: hipErrorInvalidValue
+  // CHECK: invalid value string: Invalid argument value
+
+  print_error("invalid device", hipErrorInvalidDevice);
+  // CHECK: invalid device value: 2
+  // CHECK: invalid device name: hipErrorInvalidDevice
+  // CHECK: invalid device string: Invalid device number
+
+  print_error("unknown", hipErrorUnknown);
+  // CHECK: unknown value: 3
+  // CHECK: unknown name: Unrecognized error
+  // CHECK: unknown string: Unknown error
+
+  hipError_t Unrecognized = static_cast<hipError_t>(999);
+  print_error("unrecognized", Unrecognized);
+  // CHECK: unrecognized value: 999
+  // CHECK: unrecognized name: Unrecognized error
+  // CHECK: unrecognized string: Unrecognized error
+
+  print_error("set invalid device", hipSetDevice(-1));
+  // CHECK: set invalid device value: 2
+  // CHECK: set invalid device name: hipErrorInvalidDevice
+  // CHECK: set invalid device string: Invalid device number
+
+  print_error("null stream destroy", hipStreamDestroy(nullptr));
+  // CHECK: null stream destroy value: 1
+  // CHECK: null stream destroy name: hipErrorInvalidValue
+  // CHECK: null stream destroy string: Invalid argument value
+
+  return 0;
+}
diff --git a/offload/test/offloading/HIP/host_alloc.hip b/offload/test/offloading/HIP/host_alloc.hip
new file mode 100644
index 0000000000000..8b067b39f2f81
--- /dev/null
+++ b/offload/test/offloading/HIP/host_alloc.hip
@@ -0,0 +1,40 @@
+// clang-format off
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t
+// RUN: %t | %fcheck-generic
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t -fopenmp
+// RUN: %t | %fcheck-generic
+// clang-format on
+
+// UNSUPPORTED: aarch64-unknown-linux-gnu
+// UNSUPPORTED: x86_64-unknown-linux-gnu
+// UNSUPPORTED: nvptx64-nvidia-cuda-LTO
+// UNSUPPORTED: amdgcn-amd-amdhsa-LTO
+// UNSUPPORTED: amdgpu-amd-amdhsa-LTO
+// UNSUPPORTED: intelgpu
+
+#include <stdio.h>
+
+int main(int argc, char **argv) {
+  int *HostAllocPtr = nullptr;
+  if (hipHostAlloc(&HostAllocPtr, sizeof(int), hipHostAllocDefault) !=
+      hipSuccess)
+    return 1;
+
+  *HostAllocPtr = 17;
+  printf("hipHostAlloc value: %d\n", *HostAllocPtr);
+  // CHECK: hipHostAlloc value: 17
+
+  if (hipFreeHost(HostAllocPtr) != hipSuccess)
+    return 1;
+
+  int *MallocHostPtr = nullptr;
+  if (hipMallocHost(&MallocHostPtr, sizeof(int)) != hipSuccess)
+    return 1;
+
+  *MallocHostPtr = 23;
+  printf("hipMallocHost value: %d\n", *MallocHostPtr);
+  // CHECK: hipMallocHost value: 23
+
+  if (hipFreeHost(MallocHostPtr) != hipSuccess)
+    return 1;
+}
diff --git a/offload/test/offloading/HIP/kernel_tu.hip.inc b/offload/test/offloading/HIP/kernel_tu.hip.inc
new file mode 100644
index 0000000000000..d7d28a109dfc5
--- /dev/null
+++ b/offload/test/offloading/HIP/kernel_tu.hip.inc
@@ -0,0 +1 @@
+__global__ void square(int *A) { *A = 42; }
diff --git a/offload/test/offloading/HIP/launch_tu.hip b/offload/test/offloading/HIP/launch_tu.hip
new file mode 100644
index 0000000000000..03073029ca211
--- /dev/null
+++ b/offload/test/offloading/HIP/launch_tu.hip
@@ -0,0 +1,30 @@
+// clang-format off
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t.launch_tu.o -c
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native -x hip %S/kernel_tu.hip.inc -o %t.kernel_tu.o -c
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native --offload-link %t.launch_tu.o %t.kernel_tu.o -o %t
+// RUN: %t | %fcheck-generic
+// clang-format on
+
+// UNSUPPORTED: aarch64-unknown-linux-gnu
+// UNSUPPORTED: x86_64-unknown-linux-gnu
+// UNSUPPORTED: nvptx64-nvidia-cuda-LTO
+// UNSUPPORTED: amdgcn-amd-amdhsa-LTO
+// UNSUPPORTED: amdgpu-amd-amdhsa-LTO
+// UNSUPPORTED: intelgpu
+
+#include <stdio.h>
+
+extern __global__ void square(int *A);
+
+int main(int argc, char **argv) {
+  int DevNo = 0;
+  int *Ptr;
+  hipMalloc(&Ptr, 4);
+  printf("Ptr %p\n", Ptr);
+  // CHECK: Ptr [[Ptr:0x.*]]
+  square<<<1, 1>>>(Ptr);
+  int I;
+  hipMemcpy(&I, Ptr, sizeof(int), hipMemcpyDeviceToHost);
+  printf("I: %i\n", I);
+  // CHECK: I: 42
+}
diff --git a/offload/test/offloading/HIP/memcpy_kinds.hip b/offload/test/offloading/HIP/memcpy_kinds.hip
new file mode 100644
index 0000000000000..6755a55aa0794
--- /dev/null
+++ b/offload/test/offloading/HIP/memcpy_kinds.hip
@@ -0,0 +1,51 @@
+// clang-format off
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t
+// RUN: %t | %fcheck-generic
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t -fopenmp
+// RUN: %t | %fcheck-generic
+// clang-format on
+
+// UNSUPPORTED: aarch64-unknown-linux-gnu
+// UNSUPPORTED: x86_64-unknown-linux-gnu
+// UNSUPPORTED: nvptx64-nvidia-cuda-LTO
+// UNSUPPORTED: amdgcn-amd-amdhsa-LTO
+// UNSUPPORTED: amdgpu-amd-amdhsa-LTO
+// UNSUPPORTED: intelgpu
+
+#include <stdio.h>
+
+int main(int argc, char **argv) {
+  int HostSrc = 11;
+  int HostDst = 0;
+  if (hipMemcpy(&HostDst, &HostSrc, sizeof(int), hipMemcpyHostToHost) !=
+      hipSuccess)
+    return 1;
+
+  printf("host to host: %d\n", HostDst);
+  // CHECK: host to host: 11
+
+  int *DevSrc = nullptr;
+  int *DevDst = nullptr;
+  int Result = 0;
+  if (hipMalloc(&DevSrc, sizeof(int)) != hipSuccess)
+    return 1;
+  if (hipMalloc(&DevDst, sizeof(int)) != hipSuccess)
+    return 1;
+
+  HostSrc = 42;
+  if (hipMemcpy(DevSrc, &HostSrc, sizeof(int), hipMemcpyHostToDevice) !=
+      hipSuccess)
+    return 1;
+  if (hipMemcpy(DevDst, DevSrc, sizeof(int), hipMemcpyDeviceToDevice) !=
+      hipSuccess)
+    return 1;
+  if (hipMemcpy(&Result, DevDst, sizeof(int), hipMemcpyDeviceToHost) !=
+      hipSuccess)
+    return 1;
+
+  printf("device to device: %d\n", Result);
+  // CHECK: device to device: 42
+
+  hipFree(DevSrc);
+  hipFree(DevDst);
+}
diff --git a/offload/test/offloading/HIP/stream_api.hip b/offload/test/offloading/HIP/stream_api.hip
new file mode 100644
index 0000000000000..c0e2699822814
--- /dev/null
+++ b/offload/test/offloading/HIP/stream_api.hip
@@ -0,0 +1,46 @@
+// clang-format off
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t
+// RUN: %t | %fcheck-generic
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t -fopenmp
+// RUN: %t | %fcheck-generic
+// clang-format on
+
+// UNSUPPORTED: aarch64-unknown-linux-gnu
+// UNSUPPORTED: x86_64-unknown-linux-gnu
+// UNSUPPORTED: nvptx64-nvidia-cuda-LTO
+// UNSUPPORTED: amdgcn-amd-amdhsa-LTO
+// UNSUPPORTED: amdgpu-amd-amdhsa-LTO
+// UNSUPPORTED: intelgpu
+
+#include <stdio.h>
+
+__global__ void setValue(int *Out) { *Out = 42; }
+
+int main(int argc, char **argv) {
+  hipStream_t Stream = nullptr;
+  if (hipStreamCreate(&Stream) != hipSuccess)
+    return 1;
+
+  printf("stream created: %d\n", Stream != nullptr);
+  // CHECK: stream created: 1
+
+  int *DevPtr = nullptr;
+  int Result = 0;
+  if (hipMalloc(&DevPtr, sizeof(int)) != hipSuccess)
+    return 1;
+
+  setValue<<<1, 1, 0, Stream>>>(DevPtr);
+
+  if (hipStreamSynchronize(Stream) != hipSuccess)
+    return 1;
+  if (hipMemcpy(&Result, DevPtr, sizeof(int), hipMemcpyDeviceToHost) !=
+      hipSuccess)
+    return 1;
+
+  printf("stream result: %d\n", Result);
+  // CHECK: stream result: 42
+
+  if (hipStreamDestroy(Stream) != hipSuccess)
+    return 1;
+  hipFree(DevPtr);
+}
diff --git a/offload/test/offloading/HIP/syncthreads.hip b/offload/test/offloading/HIP/syncthreads.hip
new file mode 100644
index 0000000000000..5962ab5468b86
--- /dev/null
+++ b/offload/test/offloading/HIP/syncthreads.hip
@@ -0,0 +1,40 @@
+// clang-format off
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t
+// RUN: %t | %fcheck-generic
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t -fopenmp
+// RUN: %t | %fcheck-generic
+// clang-format on
+
+// UNSUPPORTED: aarch64-unknown-linux-gnu
+// UNSUPPORTED: x86_64-unknown-linux-gnu
+// UNSUPPORTED: nvptx64-nvidia-cuda-LTO
+// UNSUPPORTED: amdgcn-amd-amdhsa-LTO
+// UNSUPPORTED: amdgpu-amd-amdhsa-LTO
+// UNSUPPORTED: intelgpu
+
+#include <stdio.h>
+
+__global__ void reduceBlock(int *Out) {
+  __shared__ int Scratch[64];
+  int Tid = threadIdx.x;
+  Scratch[Tid] = Tid;
+  __syncthreads();
+
+  if (Tid == 0) {
+    int Sum = 0;
+    for (int I = 0; I < 64; ++I)
+      Sum += Scratch[I];
+    Out[0] = Sum;
+  }
+}
+
+int main(int argc, char **argv) {
+  int *DevPtr;
+  int Result = 0;
+  hipMalloc(&DevPtr, sizeof(int));
+  reduceBlock<<<1, 64>>>(DevPtr);
+  hipMemcpy(&Result, DevPtr, sizeof(int), hipMemcpyDeviceToHost);
+
+  printf("sum: %i\n", Result);
+  // CHECK: sum: 2016
+}
diff --git a/offload/test/offloading/HIP/thread_and_block_id.hip b/offload/test/offloading/HIP/thread_and_block_id.hip
new file mode 100644
index 0000000000000..c9c33c55d72fb
--- /dev/null
+++ b/offload/test/offloading/HIP/thread_and_block_id.hip
@@ -0,0 +1,44 @@
+// clang-format off
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t
+// RUN: %t | %fcheck-generic
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t -fopenmp 
+// RUN: %t | %fcheck-generic
+// clang-format on
+
+// UNSUPPORTED: aarch64-unknown-linux-gnu
+// UNSUPPORTED: aarch64-unknown-linux-gnu-LTO
+// UNSUPPORTED: x86_64-unknown-linux-gnu
+// UNSUPPORTED: x86_64-unknown-linux-gnu-LTO
+// UNSUPPORTED: nvptx64-nvidia-cuda-LTO
+// UNSUPPORTED: amdgcn-amd-amdhsa-LTO
+// UNSUPPORTED: amdgpu-amd-amdhsa-LTO
+
+#include <stdio.h>
+#include <stdlib.h>
+
+__global__ void fill(int *A) {
+  int tid = threadIdx.x + blockDim.x * blockIdx.x;
+  A[tid] = 42;
+}
+
+int main(int argc, char **argv) {
+  int NThreads = 128;
+  int NBlocks = 512;
+  int Size = sizeof(int) * NThreads * NBlocks;
+  int *Ptr = (int*)calloc(1, Size);
+  int *DevPtr;
+  hipMalloc(&DevPtr, Size);
+  hipMemcpy(DevPtr, Ptr, Size, hipMemcpyHostToDevice);
+  printf("DevPtr %p\n", DevPtr);
+  // CHECK: DevPtr [[DevPtr:0x.*]]
+  fill<<<NBlocks, NThreads>>>(DevPtr);
+  hipMemcpy(Ptr, DevPtr, Size, hipMemcpyDeviceToHost);
+
+  for (int I = 0; I < NBlocks * NThreads; ++I) {
+    if (Ptr[I] == 42)
+      continue;
+    printf("Error at %i: %i vs %i\n", I, Ptr[I], 42);
+    return 1;
+  }
+  return 0;
+}