Skip to content

Commit

Permalink
Revert "[CIR][HIP] Use CUDA attributes for HIP global functions (#1333)"
Browse files Browse the repository at this point in the history
Broke CI jobs

This reverts commit db307ce.
  • Loading branch information
bcardosolopes committed Feb 13, 2025
1 parent db307ce commit 0bdd896
Show file tree
Hide file tree
Showing 3 changed files with 24 additions and 26 deletions.
5 changes: 2 additions & 3 deletions clang/lib/CIR/CodeGen/CIRGenModule.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -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");
}

Expand Down
14 changes: 14 additions & 0 deletions clang/test/CIR/CodeGen/HIP/simple-device.cpp
Original file line number Diff line number Diff line change
@@ -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
31 changes: 8 additions & 23 deletions clang/test/CIR/CodeGen/HIP/simple.cpp
Original file line number Diff line number Diff line change
@@ -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

0 comments on commit 0bdd896

Please sign in to comment.