[all-commits] [llvm/llvm-project] e6de9e: [CUDA][HIP] Externalize kernels in anonymous name ...

[Yaxun (Sam) Liu via All-commits]

  Branch: refs/heads/release/14.x
  Home:   https://github.com/llvm/llvm-project
  Commit: e6de9ed37308e46560243229dd78e84542f37ead
      https://github.com/llvm/llvm-project/commit/e6de9ed37308e46560243229dd78e84542f37ead
  Author: Yaxun (Sam) Liu <yaxun.liu at amd.com>
  Date:   2022-05-24 (Tue, 24 May 2022)

  Changed paths:
    M clang/include/clang/AST/ASTContext.h
    M clang/lib/AST/ASTContext.cpp
    M clang/lib/CodeGen/CGCUDANV.cpp
    M clang/lib/CodeGen/CodeGenModule.cpp
    M clang/lib/CodeGen/CodeGenModule.h
    A clang/test/CodeGenCUDA/kernel-in-anon-ns.cu

  Log Message:
  -----------
  [CUDA][HIP] Externalize kernels in anonymous name space

kernels in anonymous name space needs to have unique name
to avoid duplicate symbols.

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

Reviewed by: Artem Belevich

Differential Revision: https://reviews.llvm.org/D123353

(cherry picked from commit 4ea1d435099f992cc16127619b0feb64e070630d)


  Commit: 29f1039a7285a5c3a9c353d054140bf2556d4c4d
      https://github.com/llvm/llvm-project/commit/29f1039a7285a5c3a9c353d054140bf2556d4c4d
  Author: Yaxun (Sam) Liu <yaxun.liu at amd.com>
  Date:   2022-05-24 (Tue, 24 May 2022)

  Changed paths:
    M clang/lib/AST/ASTContext.cpp
    M clang/lib/CodeGen/CodeGenModule.cpp
    M clang/test/CodeGenCUDA/device-var-linkage.cu
    M clang/test/CodeGenCUDA/kernel-in-anon-ns.cu
    M clang/test/CodeGenCUDA/managed-var.cu
    M clang/test/CodeGenCUDA/static-device-var-rdc.cu

  Log Message:
  -----------
  [CUDA][HIP] Externalize kernels with internal linkage

This patch is a continuation of https://reviews.llvm.org/D123353.

Not only kernels in anonymous namespace, but also template
kernels with template arguments in anonymous namespace
need to be externalized.

To be more generic, this patch checks the linkage of a kernel
assuming the kernel does not have __global__ attribute. If
the linkage is internal then clang will externalize it.

This patch also fixes the postfix for externalized symbol
since nvptx does not allow '.' in symbol name.

Reviewed by: Artem Belevich

Differential Revision: https://reviews.llvm.org/D124189

Fixes: https://github.com/llvm/llvm-project/issues/54560
(cherry picked from commit 04fb81674ed7981397ffe70fe6a07b7168f6fe2f)


Compare: https://github.com/llvm/llvm-project/compare/fecfc8394484...29f1039a7285