[CIR][HIP] Compile HIP device code (#1322)

Depends on #1319

Author: koparasy

clang/lib/CIR/CodeGen/CIRGenCall.cpp

@@ -1651,7 +1651,8 @@ static void getTrivialDefaultFunctionAttributes(
   // AFAIK, neither of them support exceptions in device code.
   if (langOpts.SYCLIsDevice)
     llvm_unreachable("NYI");
-  if (langOpts.OpenCL || (langOpts.CUDA && langOpts.CUDAIsDevice)) {
+  if (langOpts.OpenCL ||
+      ((langOpts.CUDA || langOpts.HIP) && langOpts.CUDAIsDevice)) {
     auto noThrow = cir::NoThrowAttr::get(CGM.getBuilder().getContext());
     funcAttrs.set(noThrow.getMnemonic(), noThrow);
   }

clang/lib/CIR/CodeGen/CIRGenModule.cpp

@@ -3112,7 +3112,7 @@ void CIRGenModule::emitDeferred(unsigned recursionLimit) {
   // Emit CUDA/HIP static device variables referenced by host code only. Note we
   // should not clear CUDADeviceVarODRUsedByHost since it is still needed for
   // further handling.
-  if (getLangOpts().CUDA && getLangOpts().CUDAIsDevice &&
+  if ((getLangOpts().CUDA || getLangOpts().HIP) && getLangOpts().CUDAIsDevice &&
       !getASTContext().CUDADeviceVarODRUsedByHost.empty()) {
     llvm_unreachable("NYI");
   }

clang/lib/CIR/CodeGen/CIRGenTypes.cpp

@@ -348,7 +348,11 @@ mlir::Type CIRGenTypes::convertType(QualType T) {
 
   // For the device-side compilation, CUDA device builtin surface/texture types
   // may be represented in different types.
-  if (astContext.getLangOpts().CUDAIsDevice) {
+  // NOTE: CUDAIsDevice is true when building also HIP code.
+  //  1. There is no SurfaceType on HIP,
+  //  2. There is Texture memory on HIP but accessing the memory goes through
+  //  calls to the runtime. e.g. for a 2D: `tex2D<float>(tex, x, y);`
+  if (astContext.getLangOpts().CUDA && astContext.getLangOpts().CUDAIsDevice) {
     if (Ty->isCUDADeviceBuiltinSurfaceType() ||
         Ty->isCUDADeviceBuiltinTextureType())
       llvm_unreachable("NYI");

clang/test/CIR/CodeGen/HIP/simple-device.cpp

@@ -0,0 +1,14 @@
+#include "../Inputs/cuda.h"
+
+// RUN: %clang_cc1 -triple=amdgcn-amd-amdhsa -x hip -fcuda-is-device \
+// RUN:            -fclangir -emit-cir -o - %s | FileCheck %s
+
+// This shouldn't emit.
+__host__ void host_fn(int *a, int *b, int *c) {}
+
+// CHECK-NOT: cir.func @_Z7host_fnPiS_S_
+
+// This should emit as a normal C++ function.
+__device__ void device_fn(int* a, double b, float c) {}
+
+// CIR: cir.func @_Z9device_fnPidf