[HIP] Add hip input kind and codegen for kernel launching

[yaxunl]

Summary

HIP is a language similar to CUDA (https://github.com/ROCm-Developer-Tools/HIP/blob/master/docs/markdown/hip_kernel_language.md ).
The language syntax is very similar, which allows a hip program to be compiled as a CUDA program by Clang. The main difference
is the host API. HIP has a set of vendor neutral host API which can be implemented on different platforms. Currently there is open source
implementation of HIP runtime on amdgpu target (https://github.com/ROCm-Developer-Tools/HIP).

This patch adds support of input kind and language standard hip.

When hip file is compiled, both LangOpts.CUDA and LangOpts.HIP is turned on. This allows compilation of hip program as CUDA
in most cases and only special handling of hip program is needed LangOpts.HIP is checked.

This patch also adds support of kernel launching of HIP program using HIP host API.

When -x hip is not specified, there is no behaviour change for CUDA.

Patch by Greg Rodgers.
Revised and lit test added by Yaxun Liu.

[tra]

The changes appear to cover only some of the functionality needed to enable HIP support. Do you have more patches in queue? Having complete picture would help to make sense of the overall plan.
I did ask for it in D42800, but I don't think I've got the answers. It would help a lot if you or @gregrodgers could write a doc somewhere outlining overall plan for HIP support in clang, what are the main issues that need to be dealt with, and at least a general idea on how to handle them.

As far as "add -x hip, and tweak runtime glue codegen" goes, the change looks OK, but it's not very useful all by itself. It leaves a lot of other issues unsolved and no clear plan on whether/when/how you are planning to deal with them.

As things stand right now, with this patch clang will still attempt to include CUDA headers, which, among other things will provide threadIdx/blockIdx and other CUDA-specific features.
Perhaps it would make sense to disable pre-inclusion of CUDA headers and, probably, disable use of CUDA's libdevice bitcode library if we're compiling with -x hip (i.e. -nocudainc -nocudalib).
If you do depend on CUDA headers, then, I suspect, you may need to adjust some wrapper headers we use for CUDA and that change should probably come before this one.

test/CodeGenCUDA/device-stub.cu

2–9

Please wrap the long lines.

[yaxunl]

In D44984#1050526, @tra wrote:

The changes appear to cover only some of the functionality needed to enable HIP support. Do you have more patches in queue? Having complete picture would help to make sense of the overall plan.
I did ask for it in D42800, but I don't think I've got the answers. It would help a lot if you or @gregrodgers could write a doc somewhere outlining overall plan for HIP support in clang, what are the main issues that need to be dealt with, and at least a general idea on how to handle them.

As far as "add -x hip, and tweak runtime glue codegen" goes, the change looks OK, but it's not very useful all by itself. It leaves a lot of other issues unsolved and no clear plan on whether/when/how you are planning to deal with them.

As things stand right now, with this patch clang will still attempt to include CUDA headers, which, among other things will provide threadIdx/blockIdx and other CUDA-specific features.
Perhaps it would make sense to disable pre-inclusion of CUDA headers and, probably, disable use of CUDA's libdevice bitcode library if we're compiling with -x hip (i.e. -nocudainc -nocudalib).
If you do depend on CUDA headers, then, I suspect, you may need to adjust some wrapper headers we use for CUDA and that change should probably come before this one.

Hi Artem, I am responsible for upstreaming Greg's work and addressing reviewers' comments.

Yes we already have a basic working implementation of HIP compiler due to Greg's work. I will either update D42800 or create a new review about the toolchain changes for compiling and linking HIP programs. Essentially HIP has its own header files and device libraries which are taken care of by the toolchain patch.

Since the header file and library seem not to affect this patch, is it OK to defer their changes to be part of the toolchain patch?

Thanks.

[tra]

In D44984#1050557, @yaxunl wrote:

Yes we already have a basic working implementation of HIP compiler due to Greg's work.

That is great, but it's not necessarily true that all these changes will make it into clang/llvm as is. LLVM/Clang is a community effort and it helps a lot to get the changes in when the community understands what is it you're planning to do. I personally am very glad to see AMD moving towards making clang a viable compiler for AMD GPUs, but there's only so much I'll be able to do to help you with reviews if all I have is either piecemeal patches with little idea how they all fit together or one humongous patch I would have no time to dive in and really understand. Considering that compilation for GPU is a fairly niche market my bet is that your progress will be bottlenecked by the code reviews. Whatever you can do to make reviewers jobs easier by giving more context will help a lot with upstreaming the patches.

I will either update D42800 or create a new review about the toolchain changes for compiling and linking HIP programs. Essentially HIP has its own header files and device libraries which are taken care of by the toolchain patch.

Fair enough. I'll wait for the rest of the patches. If you have multiple pending patches, it helps if you could arrange them as dependent patches in phabricator. It makes it easier to see the big picture.

Since the header file and library seem not to affect this patch, is it OK to defer their changes to be part of the toolchain patch?

I'm not sure I understand. Could you elaborate?

[rjmccall]

You should send an RFC to cfe-dev about adding this new language mode. I understand that it's very similar to an existing language mode that we already support, and that's definitely we'll consider, but we shouldn't just agree to add new language modes in patch review.

[yaxunl]

In D44984#1050672, @rjmccall wrote:

You should send an RFC to cfe-dev about adding this new language mode. I understand that it's very similar to an existing language mode that we already support, and that's definitely we'll consider, but we shouldn't just agree to add new language modes in patch review.

RFC sent http://lists.llvm.org/pipermail/cfe-dev/2018-March/057426.html

Thanks.

[yaxunl]

Since the header file and library seem not to affect this patch, is it OK to defer their changes to be part of the toolchain patch?

I'm not sure I understand. Could you elaborate?

clang -cc1 does not include __clang_cuda_runtime_wrapper.h by default when clang -cc1 is called directly to compile CUDA programs. CUDA toolchain adds -include __clang_cuda_runtime_wrapper.h when compiling CUDA program as kernel code. Therefore if clang -cc1 is used to compile HIP program in lit test, there is no need to use -fnocudainc.

This patch mainly changes kernel launching API function names. The implement and testing of this change does not depend on the CUDA/HIP header files. A minimum header like test/CodeGenCUDA/Input/cuda.h is sufficient for testing this patch.

Basically this patch is only concerns about -cc1 and therefore is independent of the toolchain changes about header and library files.

[yaxunl]

ping. Any further changes need to be done for this patch? Thanks.

[rjmccall]

lib/CodeGen/CGCUDANV.cpp

98

Can you take these as StringRefs or Twines?

104

I think "addUnderscoredPrefixToName" would be better.

134

Please move this comment down into the else clause (and terminate it with a semicolon) and add your own declaration comment in your clause.

lib/Frontend/CompilerInvocation.cpp

2109

Why is this done here? We infer the language mode from the input kind somewhere else.

[yaxunl]

lib/Frontend/CompilerInvocation.cpp

2109

It is usually done through CompilerInvocation::setLangDefaults. However, HIP does not have its own input kind nor is it defined as a language standard. Therefore it cannot use CompilerInvocation::setLangDefaults to set Opts.HIP.

[yaxunl]

Revised by John's comments.

[rjmccall]

lib/Frontend/CompilerInvocation.cpp

2109

What are the values of -x if not input kinds or language standards?

[yaxunl]

lib/Frontend/CompilerInvocation.cpp

2109

I will add hip as input kind and language standard since it really is both.

[yaxunl]

Add hip as input kind and language standard.

[tra]

lib/CodeGen/CGCUDANV.cpp

51–52

const CodeGenModule &CGM

lib/Frontend/InitPreprocessor.cpp

466–469

Is __CUDA__ supposed to be set during HIP compilation? My guess is that __HIP__ and __CUDA__ should be mutually exclusive.
You do set LangOpts.CUDA during HIP compilation, so this should be changed to if (CUDA && ! HIP)

lib/Sema/SemaDecl.cpp

9054–9055

This would be somewhat easier to read:

if (II && II->isStr(getLangOpts().HIP ? "hipConfigureCall" : "cudaConfigureCall") && ...

test/CodeGenCUDA/device-stub.cu

3–10

The changes in this file do not seem to have anything related to the code changes in this patch.
Did you intend to add some HIP tests here?

[yaxunl]

lib/Frontend/InitPreprocessor.cpp

466–469

HIP documentation does not require __CUDA__ to be defined. Will make changes as you suggested.

[yaxunl]

Revised by Artem's comments.

[yaxunl]

ping

[rjmccall]

Otherwise LGTM.

lib/CodeGen/CGCUDANV.cpp

51–52

Why doesn't the CGNVCUDARuntime just hold on to a reference to the CGM? That's what we do with all the other separated singletons (like the CGCXXABI), and it would let you avoid some of the redundant fields like Context and TheModule.

[tra]

lib/CodeGen/CGCUDANV.cpp

51–52

Actually, CGCUDARuntime already has CGM field, so the CGM argument can be just dropped.

[yaxunl]

Remove CodeGenModule argument from addPrefix* functions.

[tra]

test/CodeGenCUDA/device-stub.cu

3–10

Do you need these changes?

[yaxunl]

test/CodeGenCUDA/device-stub.cu

3–10

Sorry, some changes about HIP were lost during revision. I will get back those changes.

[yaxunl]

Add back HIP related changes to the tests.

[rjmccall]

Thank you.

include/clang/Basic/LangOptions.def

 LANGOPT(ShortEnums        , 1, 0, "short enum types")
 
 LANGOPT(OpenCL            , 1, 0, "OpenCL")
 LANGOPT(OpenCLVersion     , 32, 0, "OpenCL version")
 LANGOPT(NativeHalfType    , 1, 0, "Native half type support")
 LANGOPT(NativeHalfArgsAndReturns, 1, 0, "Native half args and returns")
 LANGOPT(HalfArgsAndReturns, 1, 0, "half args and returns")
 LANGOPT(CUDA              , 1, 0, "CUDA")
+LANGOPT(HIP               , 1, 0, "HIP")
 LANGOPT(OpenMP            , 32, 0, "OpenMP support and version of OpenMP (31, 40 or 45)")
 LANGOPT(OpenMPSimd        , 1, 0, "Use SIMD only OpenMP support.")
 LANGOPT(OpenMPUseTLS      , 1, 0, "Use TLS for threadprivates or runtime calls")
 LANGOPT(OpenMPIsDevice    , 1, 0, "Generate code only for OpenMP target device")
 LANGOPT(OpenMPCUDAMode    , 1, 0, "Generate code for OpenMP pragmas in SIMT/SPMD mode")
 LANGOPT(RenderScript      , 1, 0, "RenderScript")
 
 LANGOPT(CUDAIsDevice      , 1, 0, "compiling for CUDA device")

lib/CodeGen/CGCUDANV.cpp

private:
   llvm::SmallVector<std::pair<llvm::GlobalVariable *, unsigned>, 16> DeviceVars;
   /// Keeps track of variable containing handle of GPU binary. Populated by
   /// ModuleCtorFunction() and used to create corresponding cleanup calls in
   /// ModuleDtorFunction()
   llvm::GlobalVariable *GpuBinaryHandle = nullptr;
 
   llvm::Constant *getSetupArgumentFn() const;
   llvm::Constant *getLaunchFn() const;
+  std::string addPrefixToName(CodeGenModule &CGM, std::string FuncName) const;
+  std::string addPrefixToNameBar(CodeGenModule &CGM,
+                                 std::string FuncName) const;
 
   /// Creates a function to register all kernel stubs generated in this module.
   llvm::Function *makeRegisterGlobalsFn();
 
   /// Helper function that generates a constant string and returns a pointer to
   /// the start of the string.  The result of this function can be used anywhere
   /// where the C code specifies const char*.
   llvm::Constant *makeConstantString(const std::string &Str,
public:
   /// Creates module constructor function
   llvm::Function *makeModuleCtorFunction() override;
   /// Creates module destructor function
   llvm::Function *makeModuleDtorFunction() override;
 };
 
 }
 
+std::string CGNVCUDARuntime::addPrefixToName(CodeGenModule &CGM,
+                                             std::string FuncName) const {
+  if (CGM.getLangOpts().HIP)
+    return ((Twine("hip") + Twine(FuncName)).str());
+  return ((Twine("cuda") + Twine(FuncName)).str());
+}
+std::string CGNVCUDARuntime::addPrefixToNameBar(CodeGenModule &CGM,
+                                                std::string FuncName) const {
+  if (CGM.getLangOpts().HIP)
+    return ((Twine("__hip") + Twine(FuncName)).str());
+  return ((Twine("__cuda") + Twine(FuncName)).str());
+}
+
 CGNVCUDARuntime::CGNVCUDARuntime(CodeGenModule &CGM)
     : CGCUDARuntime(CGM), Context(CGM.getLLVMContext()),
       TheModule(CGM.getModule()) {
   CodeGen::CodeGenTypes &Types = CGM.getTypes();
   ASTContext &Ctx = CGM.getContext();
 
   IntTy = CGM.IntTy;
   SizeTy = CGM.SizeTy;
   VoidTy = CGM.VoidTy;
 
   CharPtrTy = llvm::PointerType::getUnqual(Types.ConvertType(Ctx.CharTy));
   VoidPtrTy = cast<llvm::PointerType>(Types.ConvertType(Ctx.VoidPtrTy));
   VoidPtrPtrTy = VoidPtrTy->getPointerTo();
 }
 
 llvm::Constant *CGNVCUDARuntime::getSetupArgumentFn() const {
   // cudaError_t cudaSetupArgument(void *, size_t, size_t)
   llvm::Type *Params[] = {VoidPtrTy, SizeTy, SizeTy};
-  return CGM.CreateRuntimeFunction(llvm::FunctionType::get(IntTy,
-                                                           Params, false),
-                                   "cudaSetupArgument");
+  return CGM.CreateRuntimeFunction(
+      llvm::FunctionType::get(IntTy, Params, false),
+      addPrefixToName(CGM, "SetupArgument"));
 }
 
 llvm::Constant *CGNVCUDARuntime::getLaunchFn() const {
   // cudaError_t cudaLaunch(char *)
+  if (CGM.getLangOpts().HIP)
+    return CGM.CreateRuntimeFunction(
+        llvm::FunctionType::get(IntTy, CharPtrTy, false), "hipLaunchByPtr");
+  else
     return CGM.CreateRuntimeFunction(
         llvm::FunctionType::get(IntTy, CharPtrTy, false), "cudaLaunch");
 }
 
 void CGNVCUDARuntime::emitDeviceStub(CodeGenFunction &CGF,
                                      FunctionArgList &Args) {
   EmittedKernels.push_back(CGF.CurFn);
   emitDeviceStubBody(CGF, Args);
 }
 

 /// \endcode
 llvm::Function *CGNVCUDARuntime::makeRegisterGlobalsFn() {
   // No need to register anything
   if (EmittedKernels.empty() && DeviceVars.empty())
     return nullptr;
 
   llvm::Function *RegisterKernelsFunc = llvm::Function::Create(
       llvm::FunctionType::get(VoidTy, VoidPtrPtrTy, false),
-      llvm::GlobalValue::InternalLinkage, "__cuda_register_globals", &TheModule);
+      llvm::GlobalValue::InternalLinkage,
+      addPrefixToNameBar(CGM, "_register_globals"), &TheModule);
   llvm::BasicBlock *EntryBB =
       llvm::BasicBlock::Create(Context, "entry", RegisterKernelsFunc);
   CGBuilderTy Builder(CGM, Context);
   Builder.SetInsertPoint(EntryBB);
 
   // void __cudaRegisterFunction(void **, const char *, char *, const char *,
   //                             int, uint3*, uint3*, dim3*, dim3*, int*)
   llvm::Type *RegisterFuncParams[] = {
       VoidPtrPtrTy, CharPtrTy, CharPtrTy, CharPtrTy, IntTy,
       VoidPtrTy,    VoidPtrTy, VoidPtrTy, VoidPtrTy, IntTy->getPointerTo()};
   llvm::Constant *RegisterFunc = CGM.CreateRuntimeFunction(
       llvm::FunctionType::get(IntTy, RegisterFuncParams, false),
-      "__cudaRegisterFunction");
+      addPrefixToNameBar(CGM, "RegisterFunction"));
 
   // Extract GpuBinaryHandle passed as the first argument passed to
   // __cuda_register_globals() and generate __cudaRegisterFunction() call for
   // each emitted kernel.
   llvm::Argument &GpuBinaryHandlePtr = *RegisterKernelsFunc->arg_begin();
   for (llvm::Function *Kernel : EmittedKernels) {
     llvm::Constant *KernelName = makeConstantString(Kernel->getName());
     llvm::Constant *NullPtr = llvm::ConstantPointerNull::get(VoidPtrTy);
     llvm::Value *Args[] = {
         &GpuBinaryHandlePtr, Builder.CreateBitCast(Kernel, VoidPtrTy),
         KernelName, KernelName, llvm::ConstantInt::get(IntTy, -1), NullPtr,
         NullPtr, NullPtr, NullPtr,
         llvm::ConstantPointerNull::get(IntTy->getPointerTo())};
     Builder.CreateCall(RegisterFunc, Args);
   }
 
   // void __cudaRegisterVar(void **, char *, char *, const char *,
   //                        int, int, int, int)
   llvm::Type *RegisterVarParams[] = {VoidPtrPtrTy, CharPtrTy, CharPtrTy,
                                      CharPtrTy,    IntTy,     IntTy,
                                      IntTy,        IntTy};
   llvm::Constant *RegisterVar = CGM.CreateRuntimeFunction(
       llvm::FunctionType::get(IntTy, RegisterVarParams, false),
-      "__cudaRegisterVar");
+      addPrefixToNameBar(CGM, "RegisterVar"));
   for (auto &Pair : DeviceVars) {
     llvm::GlobalVariable *Var = Pair.first;
     unsigned Flags = Pair.second;
     llvm::Constant *VarName = makeConstantString(Var->getName());
     uint64_t VarSize =
         CGM.getDataLayout().getTypeAllocSize(Var->getValueType());
     llvm::Value *Args[] = {
         &GpuBinaryHandlePtr,
llvm::Function *CGNVCUDARuntime::makeModuleCtorFunction() {
   if (GpuBinaryFileName.empty())
     return nullptr;
 
   // void __cuda_register_globals(void* handle);
   llvm::Function *RegisterGlobalsFunc = makeRegisterGlobalsFn();
   // void ** __cudaRegisterFatBinary(void *);
   llvm::Constant *RegisterFatbinFunc = CGM.CreateRuntimeFunction(
       llvm::FunctionType::get(VoidPtrPtrTy, VoidPtrTy, false),
-      "__cudaRegisterFatBinary");
+      addPrefixToNameBar(CGM, "RegisterFatBinary"));
   // struct { int magic, int version, void * gpu_binary, void * dont_care };
   llvm::StructType *FatbinWrapperTy =
       llvm::StructType::get(IntTy, IntTy, VoidPtrTy, VoidPtrTy);
 
   // Register GPU binary with the CUDA runtime, store returned handle in a
   // global variable and save a reference in GpuBinaryHandle to be cleaned up
   // in destructor on exit. Then associate all known kernels with the GPU binary
   // handle so CUDA runtime can figure out what to call on the GPU side.
   llvm::ErrorOr<std::unique_ptr<llvm::MemoryBuffer>> GpuBinaryOrErr =
       llvm::MemoryBuffer::getFileOrSTDIN(GpuBinaryFileName);
   if (std::error_code EC = GpuBinaryOrErr.getError()) {
     CGM.getDiags().Report(diag::err_cannot_open_file)
         << GpuBinaryFileName << EC.message();
     return nullptr;
   }
 
   llvm::Function *ModuleCtorFunc = llvm::Function::Create(
       llvm::FunctionType::get(VoidTy, VoidPtrTy, false),
-      llvm::GlobalValue::InternalLinkage, "__cuda_module_ctor", &TheModule);
+      llvm::GlobalValue::InternalLinkage,
+      addPrefixToNameBar(CGM, "_module_ctor"), &TheModule);
   llvm::BasicBlock *CtorEntryBB =
       llvm::BasicBlock::Create(Context, "entry", ModuleCtorFunc);
   CGBuilderTy CtorBuilder(CGM, Context);
 
   CtorBuilder.SetInsertPoint(CtorEntryBB);
 
   const char *FatbinConstantName =
       CGM.getTriple().isMacOSX() ? "__NV_CUDA,__nv_fatbin" : ".nv_fatbin";
llvm::Function *CGNVCUDARuntime::makeModuleCtorFunction() {
   // Fatbin version.
   Values.addInt(IntTy, 1);
   // Data.
   Values.add(makeConstantString(GpuBinaryOrErr.get()->getBuffer(), "",
                                 FatbinConstantName, 8));
   // Unused in fatbin v1.
   Values.add(llvm::ConstantPointerNull::get(VoidPtrTy));
   llvm::GlobalVariable *FatbinWrapper = Values.finishAndCreateGlobal(
-      "__cuda_fatbin_wrapper", CGM.getPointerAlign(),
+      addPrefixToNameBar(CGM, "_fatbin_wrapper"), CGM.getPointerAlign(),
       /*constant*/ true);
   FatbinWrapper->setSection(FatbinSectionName);
 
   // GpuBinaryHandle = __cudaRegisterFatBinary(&FatbinWrapper);
   llvm::CallInst *RegisterFatbinCall = CtorBuilder.CreateCall(
       RegisterFatbinFunc, CtorBuilder.CreateBitCast(FatbinWrapper, VoidPtrTy));
   GpuBinaryHandle = new llvm::GlobalVariable(
       TheModule, VoidPtrPtrTy, false, llvm::GlobalValue::InternalLinkage,
-      llvm::ConstantPointerNull::get(VoidPtrPtrTy), "__cuda_gpubin_handle");
+      llvm::ConstantPointerNull::get(VoidPtrPtrTy),
+      addPrefixToNameBar(CGM, "_gpubin_handle"));
+
   CtorBuilder.CreateAlignedStore(RegisterFatbinCall, GpuBinaryHandle,
                                  CGM.getPointerAlign());
 
   // Call __cuda_register_globals(GpuBinaryHandle);
   if (RegisterGlobalsFunc)
     CtorBuilder.CreateCall(RegisterGlobalsFunc, RegisterFatbinCall);
 
   CtorBuilder.CreateRetVoid();

 llvm::Function *CGNVCUDARuntime::makeModuleDtorFunction() {
   // No need for destructor if we don't have a handle to unregister.
   if (!GpuBinaryHandle)
     return nullptr;
 
   // void __cudaUnregisterFatBinary(void ** handle);
   llvm::Constant *UnregisterFatbinFunc = CGM.CreateRuntimeFunction(
       llvm::FunctionType::get(VoidTy, VoidPtrPtrTy, false),
-      "__cudaUnregisterFatBinary");
+      addPrefixToNameBar(CGM, "UnregisterFatBinary"));
 
   llvm::Function *ModuleDtorFunc = llvm::Function::Create(
       llvm::FunctionType::get(VoidTy, VoidPtrTy, false),
-      llvm::GlobalValue::InternalLinkage, "__cuda_module_dtor", &TheModule);
+      llvm::GlobalValue::InternalLinkage,
+      addPrefixToNameBar(CGM, "_module_dtor"), &TheModule);
+
   llvm::BasicBlock *DtorEntryBB =
       llvm::BasicBlock::Create(Context, "entry", ModuleDtorFunc);
   CGBuilderTy DtorBuilder(CGM, Context);
   DtorBuilder.SetInsertPoint(DtorEntryBB);
 
   auto HandleValue =
       DtorBuilder.CreateAlignedLoad(GpuBinaryHandle, CGM.getPointerAlign());
   DtorBuilder.CreateCall(UnregisterFatbinFunc, HandleValue);
 
   DtorBuilder.CreateRetVoid();
   return ModuleDtorFunc;
 }
 
 CGCUDARuntime *CodeGen::CreateNVCUDARuntime(CodeGenModule &CGM) {
   return new CGNVCUDARuntime(CGM);
 }

lib/Frontend/CompilerInvocation.cpp

     IsHeaderFile =
         !Preprocessed && !ModuleMap && XValue.consume_back("-header");
 
     // Principal languages.
     DashX = llvm::StringSwitch<InputKind>(XValue)
                 .Case("c", InputKind::C)
                 .Case("cl", InputKind::OpenCL)
                 .Case("cuda", InputKind::CUDA)
+                .Case("hip", InputKind::CUDA)
                 .Case("c++", InputKind::CXX)
                 .Case("objective-c", InputKind::ObjC)
                 .Case("objective-c++", InputKind::ObjCXX)
                 .Case("renderscript", InputKind::RenderScript)
                 .Default(InputKind::Unknown);
 
     // "objc[++]-cpp-output" is an acceptable synonym for
     // "objective-c[++]-cpp-output".

         << A->getAsString(Args) << A->getValue();
     }
     else
       LangStd = OpenCLLangStd;
   }
 
   Opts.IncludeDefaultHeader = Args.hasArg(OPT_finclude_default_header);
 
+  if (const Arg *A = Args.getLastArg(OPT_x)) {
+    if (!strcmp(A->getValue(), "hip"))
+      Opts.HIP = true;
+  }
+
   llvm::Triple T(TargetOpts.Triple);
   CompilerInvocation::setLangDefaults(Opts, IK, T, PPOpts, LangStd);
 
   // -cl-strict-aliasing needs to emit diagnostic in the case where CL > 1.0.
   // This option should be deprecated for CL > 1.0 because
   // this option was added for compatibility with OpenCL 1.0.
   if (Args.getLastArg(OPT_cl_strict_aliasing)
        && Opts.OpenCLVersion > 100) {

lib/Frontend/InitPreprocessor.cpp

if (TI.isLittleEndian())
       Builder.defineMacro("__ENDIAN_LITTLE__");
 
     if (LangOpts.FastRelaxedMath)
       Builder.defineMacro("__FAST_RELAXED_MATH__");
   }
   // Not "standard" per se, but available even with the -undef flag.
   if (LangOpts.AsmPreprocessor)
     Builder.defineMacro("__ASSEMBLER__");
   if (LangOpts.CUDA)
     Builder.defineMacro("__CUDA__");
+  if (LangOpts.HIP)
+    Builder.defineMacro("__HIP__");
 }
 
 /// Initialize the predefined C++ language feature test macros defined in
 /// ISO/IEC JTC1/SC22/WG21 (C++) SD-6: "SG10 Feature Test Recommendations".
 static void InitializeCPlusPlusFeatureTestMacros(const LangOptions &LangOpts,
                                                  MacroBuilder &Builder) {
   // C++98 features.
   if (LangOpts.RTTI)

lib/Sema/SemaCUDA.cpp

bool Sema::PopForceCUDAHostDevice() {
   return true;
 }
 
 ExprResult Sema::ActOnCUDAExecConfigExpr(Scope *S, SourceLocation LLLLoc,
                                          MultiExprArg ExecConfig,
                                          SourceLocation GGGLoc) {
   FunctionDecl *ConfigDecl = Context.getcudaConfigureCallDecl();
   if (!ConfigDecl)
-    return ExprError(Diag(LLLLoc, diag::err_undeclared_var_use)
-                     << "cudaConfigureCall");
+    return ExprError(
+        Diag(LLLLoc, diag::err_undeclared_var_use)
+        << (getLangOpts().HIP ? "hipConfigureCall" : "cudaConfigureCall"));
   QualType ConfigQTy = ConfigDecl->getType();
 
   DeclRefExpr *ConfigDR = new (Context)
       DeclRefExpr(ConfigDecl, false, ConfigQTy, VK_LValue, LLLLoc);
   MarkFunctionReferenced(LLLLoc, ConfigDecl);
 
   return ActOnCallExpr(S, ConfigDR, LLLLoc, ExecConfig, GGGLoc, nullptr,
                        /*IsExecConfig=*/true);

lib/Sema/SemaDecl.cpp

     checkDLLAttributeRedeclaration(*this, Prev, NewFD,
                                    isMemberSpecialization ||
                                        isFunctionTemplateSpecialization,
                                    D.isFunctionDefinition());
   }
 
   if (getLangOpts().CUDA) {
     IdentifierInfo *II = NewFD->getIdentifier();
-    if (II && II->isStr("cudaConfigureCall") && !NewFD->isInvalidDecl() &&
+    if (II &&
+        ((getLangOpts().HIP && II->isStr("hipConfigureCall")) ||
+         (!getLangOpts().HIP && II->isStr("cudaConfigureCall"))) &&
+        !NewFD->isInvalidDecl() &&
         NewFD->getDeclContext()->getRedeclContext()->isTranslationUnit()) {
       if (!R->getAs<FunctionType>()->getReturnType()->isScalarType())
         Diag(NewFD->getLocation(), diag::err_config_scalar_return);
-
       Context.setcudaConfigureCallDecl(NewFD);
     }
 
     // Variadic functions, other than a *declaration* of printf, are not allowed
     // in device-side CUDA code, unless someone passed
     // -fcuda-allow-variadic-functions.
     if (!getLangOpts().CUDAAllowVariadicFunctions && NewFD->isVariadic() &&
         (NewFD->hasAttr<CUDADeviceAttr>() ||

test/CodeGenCUDA/Inputs/cuda.h

 
 struct dim3 {
   unsigned x, y, z;
   __host__ __device__ dim3(unsigned x, unsigned y = 1, unsigned z = 1) : x(x), y(y), z(z) {}
 };
 
 typedef struct cudaStream *cudaStream_t;
 
+#ifdef __HIP__
+int hipConfigureCall(dim3 gridSize, dim3 blockSize, size_t sharedSize = 0,
+                     cudaStream_t stream = 0);
+#else
 int cudaConfigureCall(dim3 gridSize, dim3 blockSize, size_t sharedSize = 0,
                       cudaStream_t stream = 0);
+#endif
 
 extern "C" __device__ int printf(const char*, ...);

test/CodeGenCUDA/device-stub.cu

 // RUN: echo "GPU binary would be here" > %t
-// RUN: %clang_cc1 -triple x86_64-linux-gnu -emit-llvm %s -fcuda-include-gpubinary %t -o - | FileCheck %s
+// RUN: %clang_cc1 -triple x86_64-linux-gnu -emit-llvm %s -fcuda-include-gpubinary %t -o - | FileCheck -check-prefixes=CHECK,CUDA %s
 // RUN: %clang_cc1 -triple x86_64-linux-gnu -emit-llvm %s -fcuda-include-gpubinary %t -o -  -DNOGLOBALS \
 // RUN:   | FileCheck %s -check-prefix=NOGLOBALS
-// RUN: %clang_cc1 -triple x86_64-linux-gnu -emit-llvm %s -o - | FileCheck %s -check-prefix=NOGPUBIN
+// RUN: %clang_cc1 -triple x86_64-linux-gnu -x hip -emit-llvm %s -o - | FileCheck %s -check-prefix=NOGPUBIN
+// RUN: %clang_cc1 -triple x86_64-linux-gnu -x hip -emit-llvm %s -fcuda-include-gpubinary %t -o - | FileCheck -check-prefixes=CHECK,HIP %s
+// RUN: %clang_cc1 -triple x86_64-linux-gnu -x hip -emit-llvm %s -fcuda-include-gpubinary %t -o -  -DNOGLOBALS \
+// RUN:   | FileCheck %s -check-prefix=NOGLOBALS
+// RUN: %clang_cc1 -triple x86_64-linux-gnu -x hip -emit-llvm %s -o - | FileCheck %s -check-prefix=NOGPUBIN
 
 #include "Inputs/cuda.h"
 
 #ifndef NOGLOBALS
 // CHECK-DAG: @device_var = internal global i32
 __device__ int device_var;
 
 // CHECK-DAG: @constant_var = internal global i32
 __constant__ int constant_var;

 
 // Make sure that all parts of GPU code init/cleanup are there:
 // * constant unnamed string with the kernel name
 // CHECK: private unnamed_addr constant{{.*}}kernelfunc{{.*}}\00"
 // * constant unnamed string with GPU binary
 // CHECK: private unnamed_addr constant{{.*GPU binary would be here.*}}\00"
 // CHECK-SAME: section ".nv_fatbin", align 8
 // * constant struct that wraps GPU binary
-// CHECK: @__cuda_fatbin_wrapper = internal constant { i32, i32, i8*, i8* } 
+// CUDA: @__[[PREFIX:cuda]]_fatbin_wrapper = internal constant { i32, i32, i8*, i8* }
+// HIP: @__[[PREFIX:hip]]_fatbin_wrapper = internal constant { i32, i32, i8*, i8* }
 // CHECK-SAME: { i32 1180844977, i32 1, {{.*}}, i8* null }
 // CHECK-SAME: section ".nvFatBinSegment"
 // * variable to save GPU binary handle after initialization
-// CHECK: @__cuda_gpubin_handle = internal global i8** null
+// CHECK: @__[[PREFIX]]_gpubin_handle = internal global i8** null
 // * Make sure our constructor/destructor was added to global ctor/dtor list.
-// CHECK: @llvm.global_ctors = appending global {{.*}}@__cuda_module_ctor
-// CHECK: @llvm.global_dtors = appending global {{.*}}@__cuda_module_dtor
+// CHECK: @llvm.global_ctors = appending global {{.*}}@__[[PREFIX]]_module_ctor
+// CHECK: @llvm.global_dtors = appending global {{.*}}@__[[PREFIX]]_module_dtor
 
 // Test that we build the correct number of calls to cudaSetupArgument followed
 // by a call to cudaLaunch.
 
 // CHECK: define{{.*}}kernelfunc
-// CHECK: call{{.*}}cudaSetupArgument
-// CHECK: call{{.*}}cudaSetupArgument
-// CHECK: call{{.*}}cudaSetupArgument
-// CHECK: call{{.*}}cudaLaunch
+// CHECK: call{{.*}}[[PREFIX]]SetupArgument
+// CHECK: call{{.*}}[[PREFIX]]SetupArgument
+// CHECK: call{{.*}}[[PREFIX]]SetupArgument
+// CHECK: call{{.*}}[[PREFIX]]Launch
 __global__ void kernelfunc(int i, int j, int k) {}
 
 // Test that we've built correct kernel launch sequence.
 // CHECK: define{{.*}}hostfunc
-// CHECK: call{{.*}}cudaConfigureCall
+// CHECK: call{{.*}}[[PREFIX]]ConfigureCall
 // CHECK: call{{.*}}kernelfunc
 void hostfunc(void) { kernelfunc<<<1, 1>>>(1, 1, 1); }
 #endif
 
 // Test that we've built a function to register kernels and global vars.
-// CHECK: define internal void @__cuda_register_globals
-// CHECK: call{{.*}}cudaRegisterFunction(i8** %0, {{.*}}kernelfunc
-// CHECK-DAG: call{{.*}}cudaRegisterVar(i8** %0, {{.*}}device_var{{.*}}i32 0, i32 4, i32 0, i32 0
-// CHECK-DAG: call{{.*}}cudaRegisterVar(i8** %0, {{.*}}constant_var{{.*}}i32 0, i32 4, i32 1, i32 0
-// CHECK-DAG: call{{.*}}cudaRegisterVar(i8** %0, {{.*}}ext_device_var{{.*}}i32 1, i32 4, i32 0, i32 0
-// CHECK-DAG: call{{.*}}cudaRegisterVar(i8** %0, {{.*}}ext_constant_var{{.*}}i32 1, i32 4, i32 1, i32 0
+// CHECK: define internal void @__[[PREFIX]]_register_globals
+// CHECK: call{{.*}}[[PREFIX]]RegisterFunction(i8** %0, {{.*}}kernelfunc
+// CHECK-DAG: call{{.*}}[[PREFIX]]RegisterVar(i8** %0, {{.*}}device_var{{.*}}i32 0, i32 4, i32 0, i32 0
+// CHECK-DAG: call{{.*}}[[PREFIX]]RegisterVar(i8** %0, {{.*}}constant_var{{.*}}i32 0, i32 4, i32 1, i32 0
+// CHECK-DAG: call{{.*}}[[PREFIX]]RegisterVar(i8** %0, {{.*}}ext_device_var{{.*}}i32 1, i32 4, i32 0, i32 0
+// CHECK-DAG: call{{.*}}[[PREFIX]]RegisterVar(i8** %0, {{.*}}ext_constant_var{{.*}}i32 1, i32 4, i32 1, i32 0
 // CHECK: ret void
 
 // Test that we've built contructor..
-// CHECK: define internal void @__cuda_module_ctor
-//   .. that calls __cudaRegisterFatBinary(&__cuda_fatbin_wrapper)
-// CHECK: call{{.*}}cudaRegisterFatBinary{{.*}}__cuda_fatbin_wrapper
-//   .. stores return value in __cuda_gpubin_handle
-// CHECK-NEXT: store{{.*}}__cuda_gpubin_handle
-//   .. and then calls __cuda_register_globals
-// CHECK-NEXT: call void @__cuda_register_globals
+// CHECK: define internal void @__[[PREFIX]]_module_ctor
+//   .. that calls __[[PREFIX]]RegisterFatBinary(&__[[PREFIX]]_fatbin_wrapper)
+// CHECK: call{{.*}}[[PREFIX]]RegisterFatBinary{{.*}}__[[PREFIX]]_fatbin_wrapper
+//   .. stores return value in __[[PREFIX]]_gpubin_handle
+// CHECK-NEXT: store{{.*}}__[[PREFIX]]_gpubin_handle
+//   .. and then calls __[[PREFIX]]_register_globals
+// CHECK-NEXT: call void @__[[PREFIX]]_register_globals
 
 // Test that we've created destructor.
-// CHECK: define internal void @__cuda_module_dtor
-// CHECK: load{{.*}}__cuda_gpubin_handle
-// CHECK-NEXT: call void @__cudaUnregisterFatBinary
+// CHECK: define internal void @__[[PREFIX]]_module_dtor
+// CHECK: load{{.*}}__[[PREFIX]]_gpubin_handle
+// CHECK-NEXT: call void @__[[PREFIX]]UnregisterFatBinary
 
-// There should be no __cuda_register_globals if we have no
+// There should be no __[[PREFIX]]_register_globals if we have no
 // device-side globals, but we still need to register GPU binary.
 // Skip GPU binary string first.
 // NOGLOBALS: @0 = private unnamed_addr constant{{.*}}
-// NOGLOBALS-NOT: define internal void @__cuda_register_globals
-// NOGLOBALS: define internal void @__cuda_module_ctor
-// NOGLOBALS: call{{.*}}cudaRegisterFatBinary{{.*}}__cuda_fatbin_wrapper
-// NOGLOBALS-NOT: call void @__cuda_register_globals
-// NOGLOBALS: define internal void @__cuda_module_dtor
-// NOGLOBALS: call void @__cudaUnregisterFatBinary
+// NOGLOBALS-NOT: define internal void @__{{.*}}_register_globals
+// NOGLOBALS: define internal void @__[[PREFIX:.*]]_module_ctor
+// NOGLOBALS: call{{.*}}[[PREFIX]]RegisterFatBinary{{.*}}__[[PREFIX]]_fatbin_wrapper
+// NOGLOBALS-NOT: call void @__[[PREFIX]]_register_globals
+// NOGLOBALS: define internal void @__[[PREFIX]]_module_dtor
+// NOGLOBALS: call void @__[[PREFIX]]UnregisterFatBinary
 
 // There should be no constructors/destructors if we have no GPU binary.
-// NOGPUBIN-NOT: define internal void @__cuda_register_globals
-// NOGPUBIN-NOT: define internal void @__cuda_module_ctor
-// NOGPUBIN-NOT: define internal void @__cuda_module_dtor
+// NOGPUBIN-NOT: define internal void @__[[PREFIX]]_register_globals
+// NOGPUBIN-NOT: define internal void @__[[PREFIX]]_module_ctor
+// NOGPUBIN-NOT: define internal void @__[[PREFIX]]_module_dtor

test/CodeGenCUDA/kernel-call.cu

-// RUN: %clang_cc1 -emit-llvm %s -o - | FileCheck %s
+// RUN: %clang_cc1 -emit-llvm %s -o - | FileCheck %s --check-prefixes=CUDA,CHECK
+// RUN: %clang_cc1 -x hip -emit-llvm %s -o - | FileCheck %s --check-prefixes=HIP,CHECk
+
 
 #include "Inputs/cuda.h"
 
+// CHECK-LABEL: define void @_Z2g1i(i32 %x)
+// HIP: call{{.*}}hipSetupArgument
+// HIP: call{{.*}}hipLaunchByPtr
+// CUDA: call{{.*}}cudaSetupArgument
+// CUDA: call{{.*}}cudaLaunch
 __global__ void g1(int x) {}
 
+// CHECK-LABEL: define i32 @main
 int main(void) {
-  // CHECK: call{{.*}}cudaConfigureCall
+  // HIP: call{{.*}}hipConfigureCall
+  // CUDA: call{{.*}}cudaConfigureCall
   // CHECK: icmp
   // CHECK: br
   // CHECK: call{{.*}}g1
   g1<<<1, 1>>>(42);
 }