[all-commits] [llvm/llvm-project] b310b1: [flang][cuda] Register device/constant globals as ...

[Matsu via All-commits]

  Branch: refs/heads/main
  Home:   https://github.com/llvm/llvm-project
  Commit: b310b1ab76f24a0e7a88c8ded76f54b816e5d542
      https://github.com/llvm/llvm-project/commit/b310b1ab76f24a0e7a88c8ded76f54b816e5d542
  Author: Matsu <47756807+khaki3 at users.noreply.github.com>
  Date:   2026-07-08 (Wed, 08 Jul 2026)

  Changed paths:
    M flang/include/flang/Optimizer/Dialect/CUF/CUFOps.td
    M flang/lib/Optimizer/Transforms/CUDA/CUFAddConstructor.cpp
    M flang/test/Fir/CUDA/cuda-constructor-2.f90

  Log Message:
  -----------
  [flang][cuda] Register device/constant globals as device-resident under -gpu=mem:unified (#208336)

```fortran
module m
  integer, device   :: aaa;  bind(c, name='aaa_from_c') :: aaa
  integer, constant :: zzz;  bind(c, name='zzz_from_c') :: zzz
end module
```
```c
extern int aaa_from_c;
#pragma acc declare create(aaa_from_c)   // same for zzz_from_c
```

In this code, under `-gpu=mem:unified` a module-scope device/constant
global is registered only as a CUDA variable. Another translation unit
that declares the same symbol on the host can make it be treated as host
memory, so the device symbol is unresolved at run time.

Fix: emit `cuf.register_variable_static` for device/constant globals
under `mem:unified` so they are additionally registered as
device-resident.



To unsubscribe from these emails, change your notification settings at https://github.com/llvm/llvm-project/settings/notifications