llvm / llvm/llvm-project

`-Wunused-template` many false positives for GPU code

Open
#219,968 8 comments 0 reactions 0 assignees View on GitHub
clang:diagnostics clang:frontend false-positive
Dominant language
LLVM
Stars
40.5k
Forks
18.7k
PR merge metrics
PR metrics pending

Description

`-Wunused-template` (added by default to `-Wall` in https://github.com/llvm/llvm-project/pull/208001), is essentially unusable for templated GPU libraries such as CCCL (https://github.com/NVIDIA/cccl).

Generic GPU libraries are often header-only and expose all of the GPU kernels in those headers. `nvcc`, however, when preprocessing kernels will generate `static` stub functions for those kernels resulting in many warnings such as:
```
cub/cub/device/dispatch/kernels/kernel_transform.cuh:1150:53: error: unused function template '__wrapper__device_stub_transform_kernel' [-Werror,-Wunused-template]
1150 | __attribute((__visibility__("hidden"))) static void __wrapper__device_stub_transform_kernel(const Offset &
| ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
cub/cub/device/dispatch/kernels/kernel_transform.cuh:1150:53: error: unused function template 'transform_kernel' [-Werror,-Wunused-template]
1150 | __attribute((__visibility__("hidden"))) static void transform_kernel(const Offset
| ^~~~~~~~~~~~~~~~
cub/cub/device/dispatch/kernels/kernel_for_each.cuh:94:53: error: unused function template '__wrapper__device_stub_dynamic_kernel' [-Werror,-Wunused-template]
94 | __attribute((__visibility__("hidden"))) static void __wrapper__device_stub_dynamic_kernel(const OffsetT &num_items, OpT &op) { ::cudaLaunchKernel(0, 0, 0, 0, 0, 0);}
| ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
```
Note the `static` does not exist in the original declarations.

```sh
$ clang --version
Homebrew clang version 23.1.0
Target: x86_64-unknown-linux-gnu
Thread model: posix
InstalledDir: /home/linuxbrew/.linuxbrew/Cellar/llvm/23.1.0/bin
```

Contributor guide

Open the contributing guide

Research direction

Start with the -Wunused-template diagnostic and reproduce the reported warnings using the CUDA-style declarations shown in the issue. Compare the original non-static declarations with the generated static stubs; done means the false positives are no longer emitted without hiding legitimate unused-template warnings.

Written by the indexing model from the issue text.

Assessment

Tech stack
cpp
Domain
compilers
Issue type
Bug
Difficulty
4/5
Estimated time
3-5 days
Activity status
Active
Clarity
Mostly clear
Newbie friendliness
48/100

Get new issues in your inbox

A short digest of beginner-friendly GitHub issues.