[all-commits] [llvm/llvm-project] d7e193: [HIP] Fix comdat of template kernel handle (#66283)

Yaxun (Sam) Liu via All-commits all-commits at lists.llvm.org
Thu Sep 14 12:56:15 PDT 2023


  Branch: refs/heads/main
  Home:   https://github.com/llvm/llvm-project
  Commit: d7e1932f857986df039b222619623428dc4ffe30
      https://github.com/llvm/llvm-project/commit/d7e1932f857986df039b222619623428dc4ffe30
  Author: Yaxun (Sam) Liu <yaxun.liu at amd.com>
  Date:   2023-09-14 (Thu, 14 Sep 2023)

  Changed paths:
    M clang/lib/CodeGen/CGCUDANV.cpp
    M clang/test/CodeGenCUDA/kernel-stub-name.cu

  Log Message:
  -----------
  [HIP] Fix comdat of template kernel handle (#66283)

Currently, clang emits LLVM IR that fails verifier for the following
code:

```
template<typename T>
__global__ void foo(T x);

void bar() {
  foo<<<1, 1>>>(0);
}
```
This is due to clang putting the kernel handle for foo into comdat,
which is not allowed, since the kernel handle is a declaration.

The siutation is similar to calling a declaration-only template
function. The callee will be a declaration in LLVM IR and won't be put
into comdat. This is in contrast to calling a template function with
body, which will be put into comdat.

Fixes: SWDEV-419769




More information about the All-commits mailing list