From 0bdd896e5675a1919ad394d0092bab6e43ead9d2 Mon Sep 17 00:00:00 2001
From: Bruno Cardoso Lopes <bruno.cardoso@gmail.com>
Date: Thu, 13 Feb 2025 15:01:46 +0100
Subject: [PATCH] Revert "[CIR][HIP] Use CUDA attributes for HIP global
 functions (#1333)"

Broke CI jobs

This reverts commit db307ce95f657c0e4bc63bc43da8fd8765cbaccb.
---
 clang/lib/CIR/CodeGen/CIRGenModule.cpp       |  5 ++--
 clang/test/CIR/CodeGen/HIP/simple-device.cpp | 14 +++++++++
 clang/test/CIR/CodeGen/HIP/simple.cpp        | 31 +++++---------------
 3 files changed, 24 insertions(+), 26 deletions(-)
 create mode 100644 clang/test/CIR/CodeGen/HIP/simple-device.cpp

diff --git a/clang/lib/CIR/CodeGen/CIRGenModule.cpp b/clang/lib/CIR/CodeGen/CIRGenModule.cpp
index a8376dd8f6c5..10dbd85edc4b 100644
--- a/clang/lib/CIR/CodeGen/CIRGenModule.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenModule.cpp
@@ -2357,9 +2357,8 @@ cir::FuncOp CIRGenModule::GetAddrOfFunction(clang::GlobalDecl GD, mlir::Type Ty,
   // As __global__ functions (kernels) always reside on device,
   // when we access them from host, we must refer to the kernel handle.
   // For CUDA, it's just the device stub. For HIP, it's something different.
-  if ((langOpts.CUDA || langOpts.HIP) && !langOpts.CUDAIsDevice &&
-      cast<FunctionDecl>(GD.getDecl())->hasAttr<CUDAGlobalAttr>() &&
-      cast<FunctionDecl>(GD.getDecl())->isThisDeclarationADefinition()) {
+  if (langOpts.CUDA && !langOpts.CUDAIsDevice && langOpts.HIP &&
+      cast<FunctionDecl>(GD.getDecl())->hasAttr<CUDAGlobalAttr>()) {
     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
diff --git a/clang/test/CIR/CodeGen/HIP/simple.cpp b/clang/test/CIR/CodeGen/HIP/simple.cpp
index 4fa711373d89..ec4110da10d7 100644
--- a/clang/test/CIR/CodeGen/HIP/simple.cpp
+++ b/clang/test/CIR/CodeGen/HIP/simple.cpp
@@ -1,31 +1,16 @@
 #include "../Inputs/cuda.h"
 
-// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fclangir \
-// RUN:            -x hip -emit-cir %s -o %t.cir
-// RUN: FileCheck --check-prefix=CIR-HOST --input-file=%t.cir %s
+// RUN: %clang_cc1 -triple=amdgcn-amd-amdhsa -x hip -fclangir \
+// RUN:            -emit-cir %s -o %t.cir
+// RUN: FileCheck --check-prefix=CIR --input-file=%t.cir %s
 
-// RUN: %clang_cc1 -triple=amdgcn-amd-amdhsa -x hip \
-// RUN:            -fcuda-is-device -emit-cir %s -o %t.cir
-// RUN: FileCheck --check-prefix=CIR-DEVICE --input-file=%t.cir %s
-
-// Attribute for global_fn
-// CIR-HOST: [[Kernel:#[a-zA-Z_0-9]+]] = {{.*}}#cir.cuda_kernel_name<_Z9global_fnv>{{.*}}
 
+// This should emit as a normal C++ function.
 __host__ void host_fn(int *a, int *b, int *c) {}
-// CIR-HOST: cir.func @_Z7host_fnPiS_S_
-// CIR-DEVICE-NOT: cir.func @_Z7host_fnPiS_S_
 
-__device__ void device_fn(int *a, double b, float c) {}
-// CIR-HOST-NOT: cir.func @_Z9device_fnPidf
-// CIR-DEVICE: cir.func @_Z9device_fnPidf
+// CIR: cir.func @_Z7host_fnPiS_S_
 
-#ifdef __AMDGPU__
-__global__ void global_fn() {}
-#else
-__global__ void global_fn();
-#endif
-// CIR-HOST: @_Z24__device_stub__global_fnv(){{.*}}extra([[Kernel]])
-// CIR-DEVICE: @_Z9global_fnv
+// This shouldn't emit.
+__device__ void device_fn(int* a, double b, float c) {}
 
-// Make sure `global_fn` indeed gets emitted
-__host__ void x() { auto v = global_fn; }
+// CHECK-NOT: cir.func @_Z9device_fnPidf