Revert "[CIR][HIP] Use CUDA attributes for HIP global functions (llvm#1333)"

Broke CI jobs

This reverts commit db307ce.

Author: bcardosolopes

clang/lib/CIR/CodeGen/CIRGenModule.cpp

@@ -2356,9 +2356,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");
   }
 

clang/test/CIR/CodeGen/CUDA/simple.cu

@@ -10,8 +10,6 @@
 // RUN:            %s -o %t.cir
 // RUN: FileCheck --check-prefix=CIR-DEVICE --input-file=%t.cir %s
 
-// XFAIL: *
-
 // Attribute for global_fn
 // CIR-HOST: [[Kernel:#[a-zA-Z_0-9]+]] = {{.*}}#cir.cuda_kernel_name<_Z9global_fni>{{.*}}
 

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

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