[all-commits] [llvm/llvm-project] 3ab91e: [flang][cuda][openacc] Resolve DEVICE dummies agai...

[Valentin Clement (バレンタイン クレメン) via All-commits]

  Branch: refs/heads/main
  Home:   https://github.com/llvm/llvm-project
  Commit: 3ab91ef29fc59cd48ac88fe0d5226fed6b7fd939
      https://github.com/llvm/llvm-project/commit/3ab91ef29fc59cd48ac88fe0d5226fed6b7fd939
  Author: Valentin Clement (バレンタイン クレメン) <clementval at gmail.com>
  Date:   2026-09-24 (Thu, 24 Sep 2026)

  Changed paths:
    M flang/include/flang/Lower/SymbolMap.h
    M flang/include/flang/Semantics/scope.h
    M flang/include/flang/Semantics/semantics.h
    M flang/include/flang/Semantics/tools.h
    M flang/lib/Lower/ConvertCall.cpp
    M flang/lib/Lower/OpenACC.cpp
    M flang/lib/Lower/SymbolMap.cpp
    M flang/lib/Semantics/check-call.cpp
    M flang/lib/Semantics/expression.cpp
    M flang/lib/Semantics/resolve-directives.cpp
    M flang/lib/Semantics/semantics.cpp
    M flang/lib/Semantics/tools.cpp
    A flang/test/Lower/OpenACC/acc-data-cuda-device.f90

  Log Message:
  -----------
  [flang][cuda][openacc] Resolve DEVICE dummies against !$acc data mapped objects (#225899)

CUDA Fortran generic resolution rejects a call when an actual argument
with
no CUDA data attribute is associated with an ATTRIBUTES(DEVICE) dummy.
That
makes cuBLAS-style generic interfaces unusable inside a structured
OpenACC
data region, even though the mapped objects do have a device copy there:

```
  !$acc data copyin(a, b) copyout(c)
    call cublaszgemm(transa, transb, m, n, k, alpha, a, lda, b, ldb, beta, c, ldc)
  !$acc end data

  error: No specific subroutine of generic 'cublaszgemm' matches the actual
  arguments
    dummy argument 'a=' has ATTRIBUTES(DEVICE) but its associated actual
    argument has no CUDA data attribute
```

Treat an object mapped by an enclosing structured `acc data` construct
as if
the actual had CUDADataAttr::UseDevice, which is the attribute already
used
for `host_data use_device` and which matches either a DEVICE dummy or a
host
dummy. Scalars and unmapped objects keep no attribute, so a specific
taking
host scalars and device arrays is preferred over one requiring every
argument
to be on the device.

Semantics records the mapped objects on the OpenACC construct scope
rather
than on the symbols themselves, so ordinary references inside the region
still denote the host object; only argument association consults the
mapping.
The lookup walks enclosing scopes so nested data regions work, and is
skipped
entirely when the compilation contains no such mapping.

Lowering keeps the host binding in the symbol map and adds a parallel
device-binding stack populated from the `acc.data` operands. A call
argument
is rebuilt from the device binding only when the dummy is DEVICE;
objects
without a whole-object binding, such as common block members and
component
references, are left on the host path.

Only `acc data` is handled; compute constructs are unaffected.



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