Skip to content

[clang][HIP] Don't use the OpenCLKernel CC when targeting AMDGCNSPIRV #110447

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
merged 6 commits into from
Oct 22, 2024
Merged

[clang][HIP] Don't use the OpenCLKernel CC when targeting AMDGCNSPIRV #110447

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
6 commits
Select commit Hold shift + click to select a range
f65d933
Don't use the OpenCLKernel CC when targeting AMDGCNSPIRV.
AlexVlx Sep 30, 2024
100f8bb
Merge branch 'main' of https://github.com/llvm/llvm-project into amdg…
AlexVlx Oct 6, 2024
05ed2fe
Merge branch 'main' of https://github.com/llvm/llvm-project into amdg…
AlexVlx Oct 11, 2024
8d943c3
Merge branch 'main' of https://github.com/llvm/llvm-project into amdg…
AlexVlx Oct 11, 2024
f0f6885
Merge branch 'main' of https://github.com/llvm/llvm-project into amdg…
AlexVlx Oct 15, 2024
f9e5f5f
Merge branch 'main' of https://github.com/llvm/llvm-project into amdg…
AlexVlx Oct 22, 2024
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
10 changes: 8 additions & 2 deletions clang/lib/CodeGen/CGDeclCXX.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -815,7 +815,10 @@ void CodeGenModule::EmitCXXModuleInitFunc(Module *Primary) {
assert(!getLangOpts().CUDA || !getLangOpts().CUDAIsDevice ||
getLangOpts().GPUAllowDeviceInit);
if (getLangOpts().HIP && getLangOpts().CUDAIsDevice) {
Fn->setCallingConv(llvm::CallingConv::AMDGPU_KERNEL);
if (getTriple().isSPIRV())
Fn->setCallingConv(llvm::CallingConv::SPIR_KERNEL);
else
Fn->setCallingConv(llvm::CallingConv::AMDGPU_KERNEL);
Fn->addFnAttr("device-init");
}

Expand Down Expand Up @@ -973,7 +976,10 @@ CodeGenModule::EmitCXXGlobalInitFunc() {
assert(!getLangOpts().CUDA || !getLangOpts().CUDAIsDevice ||
getLangOpts().GPUAllowDeviceInit);
if (getLangOpts().HIP && getLangOpts().CUDAIsDevice) {
Fn->setCallingConv(llvm::CallingConv::AMDGPU_KERNEL);
if (getTriple().isSPIRV())
Fn->setCallingConv(llvm::CallingConv::SPIR_KERNEL);
else
Fn->setCallingConv(llvm::CallingConv::AMDGPU_KERNEL);
Fn->addFnAttr("device-init");
}

Expand Down
8 changes: 4 additions & 4 deletions clang/lib/Sema/SemaType.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -3738,12 +3738,12 @@ static CallingConv getCCForDeclaratorChunk(
}
}
} else if (S.getLangOpts().CUDA) {
// If we're compiling CUDA/HIP code and targeting SPIR-V we need to make
// If we're compiling CUDA/HIP code and targeting HIPSPV we need to make
// sure the kernels will be marked with the right calling convention so that
// they will be visible by the APIs that ingest SPIR-V.
// they will be visible by the APIs that ingest SPIR-V. We do not do this
// when targeting AMDGCNSPIRV, as it does not rely on OpenCL.
llvm::Triple Triple = S.Context.getTargetInfo().getTriple();
if (Triple.getArch() == llvm::Triple::spirv32 ||
Triple.getArch() == llvm::Triple::spirv64) {
if (Triple.isSPIRV() && Triple.getVendor() != llvm::Triple::AMD) {
for (const ParsedAttr &AL : D.getDeclSpec().getAttributes()) {
if (AL.getKind() == ParsedAttr::AT_CUDAGlobal) {
CC = CC_OpenCLKernel;
Expand Down
6 changes: 6 additions & 0 deletions clang/test/CodeGenCUDA/device-init-fun.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -4,11 +4,17 @@
// RUN: -fgpu-allow-device-init -x hip \
// RUN: -fno-threadsafe-statics -emit-llvm -o - %s \
// RUN: | FileCheck %s
// RUN: %clang_cc1 -triple spirv64-amd-amdhsa -fcuda-is-device -std=c++11 \
// RUN: -fgpu-allow-device-init -x hip \
// RUN: -fno-threadsafe-statics -emit-llvm -o - %s \
// RUN: | FileCheck %s --check-prefix=CHECK-SPIRV

#include "Inputs/cuda.h"

// CHECK: define internal amdgpu_kernel void @_GLOBAL__sub_I_device_init_fun.cu() #[[ATTR:[0-9]*]]
// CHECK: attributes #[[ATTR]] = {{.*}}"device-init"
// CHECK-SPIRV: define internal spir_kernel void @_GLOBAL__sub_I_device_init_fun.cu(){{.*}} #[[ATTR:[0-9]*]]
// CHECK-SPIRV: attributes #[[ATTR]] = {{.*}}"device-init"

__device__ void f();

Expand Down
8 changes: 7 additions & 1 deletion clang/test/CodeGenCUDA/kernel-amdgcn.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,31 +1,37 @@
// RUN: %clang_cc1 -triple amdgcn -fcuda-is-device -emit-llvm -x hip %s -o - | FileCheck %s
// RUN: %clang_cc1 -triple spirv64-amd-amdhsa -fcuda-is-device -emit-llvm -x hip %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
#include "Inputs/cuda.h"

// CHECK: define{{.*}} amdgpu_kernel void @_ZN1A6kernelEv
// CHECK-SPIRV: define{{.*}} spir_kernel void @_ZN1A6kernelEv
class A {
public:
static __global__ void kernel(){}
};

// CHECK: define{{.*}} void @_Z10non_kernelv
// CHECK-SPIRV: define{{.*}} void @_Z10non_kernelv
__device__ void non_kernel(){}

// CHECK: define{{.*}} amdgpu_kernel void @_Z6kerneli
// CHECK-SPIRV: define{{.*}} spir_kernel void @_Z6kerneli
__global__ void kernel(int x) {
non_kernel();
}

// CHECK: define{{.*}} amdgpu_kernel void @_Z11EmptyKernelIvEvv
// CHECK-SPIRV: define{{.*}} spir_kernel void @_Z11EmptyKernelIvEvv
template <typename T>
__global__ void EmptyKernel(void) {}

struct Dummy {
/// Type definition of the EmptyKernel kernel entry point
typedef void (*EmptyKernelPtr)();
EmptyKernelPtr Empty() { return EmptyKernel<void>; }
EmptyKernelPtr Empty() { return EmptyKernel<void>; }
};

// CHECK: define{{.*}} amdgpu_kernel void @_Z15template_kernelI1AEvT_{{.*}} #[[ATTR:[0-9][0-9]*]]
// CHECK-SPIRV: define{{.*}} spir_kernel void @_Z15template_kernelI1AEvT_{{.*}} #[[ATTR:[0-9][0-9]*]]
template<class T>
__global__ void template_kernel(T x) {}

Expand Down
Loading