4394aa685c
This PR adds checks for when emitting weak aliases in: `void
CodeGenModule::EmitGlobal(GlobalDecl GD)`, before for device compilation
for OpenMP, HIP and Cuda, clang would look for the aliasee even if it
was never marked for device compilation.
For OpenMP the following case now works:
> Failed before when compiling with device, ie: `clang -fopenmp
-fopenmp-targets=amdgcn-amd-amdhsa`
> ```
> int __Two(void) { return 2; }
> int Two(void) __attribute__ ((weak, alias("__Two")));
> ```
For HIP / Cuda:
>
> ```
> int __HostFunc(void) { return 42; }
> int HostFunc(void) __attribute__ ((weak, alias("__HostFunc")));
> ```
For HIP:
>Failed before on HIP, Cuda fails due to: `NVPTX aliasee must not be
'.weak'` error
> ```
> __device__ int __One(void) { return 2; }
> __device__ int One(void) __attribute__ ((weak, alias("__One")));
> ```
Included are Codegen LIT tests for the above cases, and also cases for
weak alias cases that currently work in clang.
Fixes https://github.com/llvm/llvm-project/issues/117369
126 lines
5.7 KiB
C++
126 lines
5.7 KiB
C++
// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --check-globals all --version 6
|
|
// REQUIRES: amdgpu-registered-target
|
|
// RUN: %clang_cc1 -x hip -triple x86_64-unknown-linux-gnu -aux-triple amdgcn-amd-amdhsa -emit-llvm-bc %s -o %t-host.bc
|
|
// RUN: %clang_cc1 -x hip -triple x86_64-unknown-linux-gnu -aux-triple amdgcn-amd-amdhsa -emit-llvm %s -o - | FileCheck %s --check-prefix=HOST
|
|
// RUN: %clang_cc1 -x hip -triple amdgcn-amd-amdhsa -aux-triple x86_64-unknown-linux-gnu -emit-llvm %s -fcuda-is-device -o - | FileCheck %s --check-prefix=DEVICE
|
|
|
|
#define __device__ __attribute__((device))
|
|
#define __host__ __attribute__((host))
|
|
|
|
extern "C" {
|
|
//.
|
|
// HOST: @HostFunc = weak alias i32 (), ptr @__HostFunc
|
|
// HOST: @HostFunc_ = alias i32 (), ptr @__HostFunc
|
|
// HOST: @HostVar = weak alias i32, ptr @__HostVar
|
|
// HOST: @HostVar_ = alias i32, ptr @__HostVar
|
|
// HOST: @Two = weak alias i32 (), ptr @__Two
|
|
// HOST: @Two_ = alias i32 (), ptr @__Two
|
|
// HOST: @_Z5Threev = weak alias i32 (), ptr @_Z7__Threev
|
|
// HOST: @_Z6Three_v = alias i32 (), ptr @_Z7__Threev
|
|
// HOST: @_Z4Fourv = weak alias i32 (), ptr @_Z6__Fourv
|
|
// HOST: @_Z4Fourf = weak alias float (float), ptr @_Z6__Fourf
|
|
//.
|
|
// DEVICE: @One = weak alias i32 (), ptr @__One
|
|
// DEVICE: @One_ = alias i32 (), ptr @__One
|
|
// DEVICE: @Two = weak alias i32 (), ptr @__Two
|
|
// DEVICE: @Two_ = alias i32 (), ptr @__Two
|
|
// DEVICE: @_Z5Threev = weak alias i32 (), ptr @_Z7__Threev
|
|
// DEVICE: @_Z6Three_v = alias i32 (), ptr @_Z7__Threev
|
|
// DEVICE: @_Z4Fourv = weak alias i32 (), ptr @_Z6__Fourv
|
|
// DEVICE: @_Z4Fourf = weak alias float (float), ptr @_Z6__Fourf
|
|
//.
|
|
// HOST-LABEL: define dso_local i32 @__HostFunc(
|
|
// HOST-SAME: ) #[[ATTR0:[0-9]+]] {
|
|
// HOST-NEXT: [[ENTRY:.*:]]
|
|
// HOST-NEXT: ret i32 42
|
|
//
|
|
int __HostFunc(void) { return 42; }
|
|
int __HostVar = 1;
|
|
int HostFunc(void) __attribute__((weak, alias("__HostFunc")));
|
|
int HostFunc_(void) __attribute__((alias("__HostFunc")));
|
|
extern int __attribute__((weak, alias("__HostVar"))) HostVar;
|
|
extern int __attribute__((alias("__HostVar"))) HostVar_;
|
|
|
|
// DEVICE-LABEL: define dso_local i32 @__One(
|
|
// DEVICE-SAME: ) #[[ATTR0:[0-9]+]] {
|
|
// DEVICE-NEXT: [[ENTRY:.*:]]
|
|
// DEVICE-NEXT: [[RETVAL:%.*]] = alloca i32, align 4, addrspace(5)
|
|
// DEVICE-NEXT: [[RETVAL_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[RETVAL]] to ptr
|
|
// DEVICE-NEXT: ret i32 1
|
|
//
|
|
__device__ int __One(void) { return 1; }
|
|
__device__ int One(void) __attribute__((weak, alias("__One")));
|
|
__device__ int One_(void) __attribute__((alias("__One")));
|
|
|
|
// HOST-LABEL: define dso_local i32 @__Two(
|
|
// HOST-SAME: ) #[[ATTR0]] {
|
|
// HOST-NEXT: [[ENTRY:.*:]]
|
|
// HOST-NEXT: ret i32 2
|
|
//
|
|
// DEVICE-LABEL: define dso_local i32 @__Two(
|
|
// DEVICE-SAME: ) #[[ATTR0]] {
|
|
// DEVICE-NEXT: [[ENTRY:.*:]]
|
|
// DEVICE-NEXT: [[RETVAL:%.*]] = alloca i32, align 4, addrspace(5)
|
|
// DEVICE-NEXT: [[RETVAL_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[RETVAL]] to ptr
|
|
// DEVICE-NEXT: ret i32 2
|
|
//
|
|
__host__ __device__ int __Two(void) { return 2; }
|
|
__host__ __device__ int Two(void) __attribute__((weak, alias("__Two")));
|
|
__host__ __device__ int Two_(void) __attribute__((alias("__Two")));
|
|
}
|
|
|
|
// HOST-LABEL: define linkonce_odr noundef i32 @_Z7__Threev(
|
|
// HOST-SAME: ) #[[ATTR0]] comdat {
|
|
// HOST-NEXT: [[ENTRY:.*:]]
|
|
// HOST-NEXT: ret i32 5
|
|
//
|
|
// DEVICE-LABEL: define linkonce_odr noundef i32 @_Z7__Threev(
|
|
// DEVICE-SAME: ) #[[ATTR0]] comdat {
|
|
// DEVICE-NEXT: [[ENTRY:.*:]]
|
|
// DEVICE-NEXT: [[RETVAL:%.*]] = alloca i32, align 4, addrspace(5)
|
|
// DEVICE-NEXT: [[RETVAL_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[RETVAL]] to ptr
|
|
// DEVICE-NEXT: ret i32 5
|
|
//
|
|
__host__ __device__ constexpr int __Three(void) { return 5; }
|
|
__host__ __device__ int Three(void) __attribute__((weak, alias("_Z7__Threev")));
|
|
__host__ __device__ int Three_(void) __attribute__((alias("_Z7__Threev")));
|
|
|
|
|
|
// HOST-LABEL: define dso_local noundef i32 @_Z6__Fourv(
|
|
// HOST-SAME: ) #[[ATTR0]] {
|
|
// HOST-NEXT: [[ENTRY:.*:]]
|
|
// HOST-NEXT: ret i32 2
|
|
//
|
|
// DEVICE-LABEL: define dso_local noundef i32 @_Z6__Fourv(
|
|
// DEVICE-SAME: ) #[[ATTR0]] {
|
|
// DEVICE-NEXT: [[ENTRY:.*:]]
|
|
// DEVICE-NEXT: [[RETVAL:%.*]] = alloca i32, align 4, addrspace(5)
|
|
// DEVICE-NEXT: [[RETVAL_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[RETVAL]] to ptr
|
|
// DEVICE-NEXT: ret i32 2
|
|
//
|
|
__host__ __device__ int __Four(void) { return 2; }
|
|
// HOST-LABEL: define dso_local noundef float @_Z6__Fourf(
|
|
// HOST-SAME: float noundef [[F:%.*]]) #[[ATTR0]] {
|
|
// HOST-NEXT: [[ENTRY:.*:]]
|
|
// HOST-NEXT: [[F_ADDR:%.*]] = alloca float, align 4
|
|
// HOST-NEXT: store float [[F]], ptr [[F_ADDR]], align 4
|
|
// HOST-NEXT: [[TMP0:%.*]] = load float, ptr [[F_ADDR]], align 4
|
|
// HOST-NEXT: [[MUL:%.*]] = fmul contract float 2.000000e+00, [[TMP0]]
|
|
// HOST-NEXT: ret float [[MUL]]
|
|
//
|
|
// DEVICE-LABEL: define dso_local noundef float @_Z6__Fourf(
|
|
// DEVICE-SAME: float noundef [[F:%.*]]) #[[ATTR0]] {
|
|
// DEVICE-NEXT: [[ENTRY:.*:]]
|
|
// DEVICE-NEXT: [[RETVAL:%.*]] = alloca float, align 4, addrspace(5)
|
|
// DEVICE-NEXT: [[F_ADDR:%.*]] = alloca float, align 4, addrspace(5)
|
|
// DEVICE-NEXT: [[RETVAL_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[RETVAL]] to ptr
|
|
// DEVICE-NEXT: [[F_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[F_ADDR]] to ptr
|
|
// DEVICE-NEXT: store float [[F]], ptr [[F_ADDR_ASCAST]], align 4
|
|
// DEVICE-NEXT: [[TMP0:%.*]] = load float, ptr [[F_ADDR_ASCAST]], align 4
|
|
// DEVICE-NEXT: [[MUL:%.*]] = fmul contract float 2.000000e+00, [[TMP0]]
|
|
// DEVICE-NEXT: ret float [[MUL]]
|
|
//
|
|
__host__ __device__ float __Four(float f) { return 2.0f * f; }
|
|
__host__ __device__ int Four(void) __attribute__((weak, alias("_Z6__Fourv")));
|
|
__host__ __device__ float Four(float f) __attribute__((weak, alias("_Z6__Fourf")));
|