From cfbbcc828d33bed336afdd9e13e5cb796f482104 Mon Sep 17 00:00:00 2001
From: koparasy <parasyris1@llnl.gov>
Date: Fri, 7 Feb 2025 12:40:51 -0800
Subject: [PATCH] [HIP] Compile device code

---
 clang/lib/CIR/CodeGen/CIRGenCall.cpp         |  3 ++-
 clang/lib/CIR/CodeGen/CIRGenModule.cpp       |  2 +-
 clang/lib/CIR/CodeGen/CIRGenTypes.cpp        |  6 +++++-
 clang/test/CIR/CodeGen/HIP/simple-device.cpp | 14 ++++++++++++++
 4 files changed, 22 insertions(+), 3 deletions(-)
 create mode 100644 clang/test/CIR/CodeGen/HIP/simple-device.cpp

diff --git a/clang/lib/CIR/CodeGen/CIRGenCall.cpp b/clang/lib/CIR/CodeGen/CIRGenCall.cpp
index 9a15e3337f32..035b30b1f077 100644
--- a/clang/lib/CIR/CodeGen/CIRGenCall.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenCall.cpp
@@ -1641,7 +1641,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);
   }
diff --git a/clang/lib/CIR/CodeGen/CIRGenModule.cpp b/clang/lib/CIR/CodeGen/CIRGenModule.cpp
index 412369ed07ef..8a51b5c556f3 100644
--- a/clang/lib/CIR/CodeGen/CIRGenModule.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenModule.cpp
@@ -3114,7 +3114,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");
   }
diff --git a/clang/lib/CIR/CodeGen/CIRGenTypes.cpp b/clang/lib/CIR/CodeGen/CIRGenTypes.cpp
index 66d6a57e242a..70f7b681bc86 100644
--- a/clang/lib/CIR/CodeGen/CIRGenTypes.cpp
+++ b/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");
diff --git a/clang/test/CIR/CodeGen/HIP/simple-device.cpp b/clang/test/CIR/CodeGen/HIP/simple-device.cpp
new file mode 100644
index 000000000000..e627a90dc410
--- /dev/null
+++ b/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