[all-commits] [llvm/llvm-project] 03e42c: Reland [CUDA][HIP] Fix CTAD for host/device constr...

[Yaxun (Sam) Liu via All-commits]

  Branch: refs/heads/main
  Home:   https://github.com/llvm/llvm-project
  Commit: 03e42c97c7c0f6df51e3298cafa60b1f1e59c0b9
      https://github.com/llvm/llvm-project/commit/03e42c97c7c0f6df51e3298cafa60b1f1e59c0b9
  Author: Yaxun (Sam) Liu <yaxun.liu at amd.com>
  Date:   2025-12-05 (Fri, 05 Dec 2025)

  Changed paths:
    M clang/docs/HIPSupport.rst
    M clang/docs/ReleaseNotes.rst
    M clang/include/clang/Basic/DiagnosticSemaKinds.td
    M clang/lib/Sema/SemaCUDA.cpp
    M clang/lib/Sema/SemaDeclAttr.cpp
    M clang/lib/Sema/SemaTemplateDeductionGuide.cpp
    A clang/test/SemaCUDA/deduction-guide-attrs.cu
    A clang/test/SemaCUDA/deduction-guide-overload.cu
    A clang/test/SemaCUDA/deduction-guide.cu

  Log Message:
  -----------
  Reland [CUDA][HIP] Fix CTAD for host/device constructors (#168711) (#170481)

Clang currently does not allow using CTAD in CUDA/HIP device functions
since deduction guides are treated as host-only. This patch fixes that
by treating deduction guides as host+device. The rationale is that
deduction guides do not actually generate code in IR, and there is an
existing check for device/host correctness for constructors.

The patch also suppresses duplicate implicit deduction guides from
host/device constructors with identical signatures and constraints to
prevent ambiguity.

For CUDA/HIP, deduction guides are now always implicitly enabled for
both host and device, which matches nvcc's effective behavior. Unlike
nvcc, which silently ignores explicit CUDA/HIP target attributes on
deduction guides, Clang diagnoses device- and host-only attributes as
errors to keep the syntax clean and avoid confusion.

This ensures CTAD works correctly in CUDA/HIP for constructors with
different target attributes and provides clearer diagnostics when users
attempt to annotate deduction guides with CUDA/HIP target attributes.

Example:

```
  #include <tuple>

  __host__ __device__ void func()
  {
    std::tuple<int, int> t = std::tuple(1, 1);
  }
```

This compiles with nvcc but fails with clang for CUDA/HIP without this
fix.

Reference: https://godbolt.org/z/WhT1GrhWE

Fixes: https://github.com/ROCm/ROCm/issues/5646

Fixes: https://github.com/llvm/llvm-project/issues/146646



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