[all-commits] [llvm/llvm-project] 04fb81: [CUDA][HIP] Externalize kernels with internal linkage

Yaxun (Sam) Liu via All-commits all-commits at lists.llvm.org
Fri Apr 22 14:07:12 PDT 2022

  Branch: refs/heads/main
  Home:   https://github.com/llvm/llvm-project
  Commit: 04fb81674ed7981397ffe70fe6a07b7168f6fe2f
  Author: Yaxun (Sam) Liu <yaxun.liu at amd.com>
  Date:   2022-04-22 (Fri, 22 Apr 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

More information about the All-commits mailing list