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
18 lines
593 B
Plaintext
18 lines
593 B
Plaintext
// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --check-globals all --version 6
|
|
// RUN: %clang_cc1 -x cuda -triple x86_64-unknown-linux-gnu -aux-triple nvptx64-nvidia-cuda -emit-llvm %s -o - | FileCheck %s --check-prefix=HOST
|
|
|
|
extern "C" {
|
|
|
|
//.
|
|
// HOST: @HostFunc = weak alias i32 (), ptr @__HostFunc
|
|
//.
|
|
// 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 HostFunc(void) __attribute__ ((weak, alias("__HostFunc")));
|
|
|
|
}
|