Jason-VanBeusekom wrote:
[16c1a68](https://github.com/llvm/llvm-project/pull/164326/commits/16c1a6888b87644f24c07a75c067d9e25eb1ac3e)
Addresses the above feedback (minus documentation).
I refactored the code to be more readable about the states, while, also,
reducing the nested conditionals and added comments to describe the states.
Below I wrote a modified version of `shouldSkipAliasEmission` with debug prints
better helps show the logic and show the results for various cases, of note for
some of the incorrect hip cases the code does compiler under certain cases:
```c++
static bool shouldSkipAliasEmission(const CodeGenModule &CGM,
const ValueDecl *Global) {
const LangOptions &LangOpts = CGM.getLangOpts();
if (!(LangOpts.OpenMPIsTargetDevice || LangOpts.CUDA)) {
printf("Not cuda / openmp target returning false\n");
return false;
}
const auto *AA = Global->getAttr<AliasAttr>();
GlobalDecl AliaseeGD;
// Check if the aliasee exists.
if (!CGM.lookupRepresentativeDecl(AA->getAliasee(), AliaseeGD)) {
if (LangOpts.CUDA) {
printf("HIP / CUDA aliasse not found returning true\n");
// In CUDA/HIP, if the aliasee is not found, skip the alias emission.
// This is not a hard error as this branch is executed for both the host
// and device, with no respect to where the aliasee is defined.
return true;
}
printf("OpenMP aliasee not found error\n");
// For OpenMP, lookupRepresentativeDecl should always find the aliasee, this
// is an error
CGM.getDiags().Report(AA->getLocation(), diag::err_alias_to_undefined)
<< false << true;
return false;
}
const auto *AliaseeDecl = dyn_cast<ValueDecl>(AliaseeGD.getDecl());
if (LangOpts.OpenMPIsTargetDevice) {
if (!AliaseeDecl ||
!OMPDeclareTargetDeclAttr::isDeclareTargetDeclaration(AliaseeDecl)) {
printf("OpenMP aliasse not target decl returning true\n");
return true;
}
printf("OpenMP aliasse is target decl returning false\n");
return false;
}
// CUDA / HIP
const bool HasDeviceAttr = Global->hasAttr<CUDADeviceAttr>();
const bool AliaseeHasDeviceAttr =
AliaseeDecl && AliaseeDecl->hasAttr<CUDADeviceAttr>();
if (LangOpts.CUDAIsDevice) {
if (!HasDeviceAttr || !AliaseeHasDeviceAttr) {
// On device, skip alias emission if either the alias or the aliasee
// is not marked with __device__.
printf(
"Skipping HIP / CUDA alias emission on device for host only alias\n");
return true;
}
printf("Emitting HIP / CUDA alias on device\n");
return false;
}
printf("Emitting HIP / CUDA alias on host\n");
// CUDA / HIP Host
// we know that the aliasee exists from above, so we know to emit
return false;
}
```
And prints out the following for each case:
### OpenMP:
```c++
#pragma omp declare target
int __One(void) { return 1; }
#pragma omp end declare target
int One(void) __attribute__ ((weak, alias("__One")));
```
#### Outputs:
```
Not cuda / openmp target returning false
OpenMP aliasse is target decl returning false
OpenMP aliasse is target decl returning false
```
---
``` c++
#pragma omp declare target
int __One(void) { return 1; }
int One(void) __attribute__ ((weak, alias("__One")));
#pragma omp end declare target
```
#### Outputs:
```
Not cuda / openmp target returning false
OpenMP aliasse is target decl returning false
```
---
```c++
int __One(void) { return 1; }
int One(void) __attribute__ ((weak, alias("__One")));
```
#### Outputs:
```
Not cuda / openmp target returning false
OpenMP aliasse not target decl returning true
```
---
```c++
int __One(void) { return 1; }
#pragma omp declare target
int One(void) __attribute__ ((weak, alias("__One")));
#pragma omp end declare target
```
#### Outputs:
```
Not cuda / openmp target returning false
OpenMP aliasse not target decl returning true
```
##### Results in runtime error (expected):
> omptarget error: Failed to load image "a provided binary image is malformed"
> failed to load binary 0x55b530291d60
> omptarget fatal error 0: "the plugin backend is in an invalid or unsupported
> state" failed to load images on device '0'
---
### HIP:
```c+
__device__ int __DevTwo(void) { return 2; }
__device__ int DevTwo(void) __attribute__ ((weak, alias("__DevTwo")));
```
#### Outputs:
```
Emitting HIP / CUDA alias on device
HIP / CUDA aliasse not found returning true
```
---
```c++
int __DevTwo(void) { return 2; }
int DevTwo(void) __attribute__ ((weak, alias("__DevTwo")));
```
#### Outputs:
```
HIP / CUDA aliasse not found returning true
Emitting HIP / CUDA alias on host
```
---
```c++
int __DevTwo(void) { return 2; }
__device__ int DevTwo(void) __attribute__ ((weak, alias("__DevTwo")));
```
#### Outputs:
```
HIP / CUDA aliasse not found returning true
Emitting HIP / CUDA alias on host
```
(compiles with no error, calling aliasee on host works, calling alias on
device results in compiler error calling alaisee on host results in compiler
error, calling on device results in runtime error [memory access fault])
we could add check to have error here, but not sure if we want to
---
```c++
__device__ int __DevTwo(void) { return 2; }
int DevTwo(void) __attribute__ ((weak, alias("__DevTwo")));
```
#### Outputs:
```
Skipping HIP / CUDA alias emission on device for host only alias
HIP / CUDA aliasse not found returning true
```
-get compiler error when trying to call alias on device
-able to call aliasee on device
-call alias on host results in segfault at runtime
-call aliasee on host results in compiler error
could add check to have error here, but not sure if we want to
https://github.com/llvm/llvm-project/pull/164326
_______________________________________________
cfe-commits mailing list
[email protected]
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
