[BUG] given attribute ‘noinline’

[Sarna555]

While compiling raft 23.06 I get the error message, moving to 23.04 solved the problem

../cmake-build-debug/_deps/raft-src/cpp/include/raft/util/cache_util.cuh:331:27: error: inline function ‘void raft::cache::get_cache_idx(int*, int, int*, int, int, int*, int*, bool*, int)’ given attribute ‘noinline’ [-Werror]
  331 | __global__ inline void get_cache_idx(int* keys,
      |                           ^~~~~~~~~~~~~

env:
compiling with g++ 12 and nvcc 12.1

$ c++ --version
c++ (Ubuntu 12.1.0-2ubuntu1~22.04) 12.1.0
Copyright (C) 2022 Free Software Foundation, Inc.
This is free software; see the source for copying conditions.  There is NO
warranty; not even for MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.

$ nvcc --version
nvcc: NVIDIA (R) Cuda compiler driver
Copyright (c) 2005-2023 NVIDIA Corporation
Built on Mon_Apr__3_17:16:06_PDT_2023
Cuda compilation tools, release 12.1, V12.1.105
Build cuda_12.1.r12.1/compiler.32688072_0

[achirkin]

Just tried the debug mode and got the same issue.
@ahendriksen , @cjnolet , shall we maybe revert #1492 ?

[ahendriksen]

I see. Indeed nvcc 12.1 issues a warning, which we treat as an error.

PR #1492 did solve an issue. So reverting it will bring the issue back.

There are two other possible solutions to add weak linkage to the kernel:

Option 1: explicity mark as weak:

// Make kernel have weak linkage:
#define RAFT_WEAK_FUNCTION __attribute__((weak))

__global__ RAFT_WEAK_FUNCTION void kernel_weak() {
}

Option 2: add a dummy template parameter (has the same effect, but adds an unused template parameter)

template <int dummy = 0>
__global__ void kernel_dummy_template() {
}

Godbolt: https://godbolt.org/z/q6E3ffaY1

Can you try if either of these two options works?

[achirkin]

I just did something similar to the dummy option, it works
template <typename = void>

[ahendriksen]

This was a very easy issue to reproduce on godbolt, but very difficult in our repository. Get_cache_idx is used in test/matrix.cu.

However, version 1.7.2 of CUB eats the "inline qualifier ignored for "global" function" warning. That is why this issue was not caught in our CI.

NVIDIA/cccl#319