lambda: expected single user

[johnbowen42]

template<typename T>
inline auto jit_variable(T arg)
{
#if defined RAJA_ENABLE_JIT
  return proteus::jit_variable(std::forward<T>(arg));
#else
  return std::forward<T>(arg);
#endif
}

if we define a wrapper method like this inside a portability layer, we get

[proteus][FATAL ERROR] /g/g11/bowen36/RAJA/build_lc_toss4-amdclang-proteus-6.4.2-gfx90a/_deps/proteus-src/src/pass/ProteusPass.cpp:1262 => Expected single user
PLEASE submit a bug report to https://github.com/llvm/llvm-project/issues/ and include the crash backtrace, preprocessed source, and associated run script.
Stack dump:
1.      <eof> parser at end of file
2.      Optimizer
3.      Running pass "(anonymous namespace)::ProteusPass" on module "/g/g11/bowen36/RAJA/examples/forall-jit.cpp"
 #0 0x0000000002365c68 llvm::sys::PrintStackTrace(llvm::raw_ostream&, int) (/opt/rocm-6.4.2/lib/llvm/bin/clang-19+0x2365c68)
 #1 0x00000000023634ec SignalHandler(int) (/opt/rocm-6.4.2/lib/llvm/bin/clang-19+0x23634ec)
 #2 0x0000155555116990 __restore_rt (/lib64/libpthread.so.0+0x12990)
 #3 0x0000155553b855af raise (/lib64/libc.so.6+0x4e5af)
 #4 0x0000155553b58ee5 abort (/lib64/libc.so.6+0x21ee5)
 #5 0x0000155553586863 proteus::reportFatalError(llvm::Twine const&, char const*, unsigned int) (/g/g11/bowen36/RAJA/build_lc_toss4-amdclang-proteus-6.4.2-gfx90a/lib/libProteusPass.so+0x62863)
 #6 0x0000155553586962 (/g/g11/bowen36/RAJA/build_lc_toss4-amdclang-proteus-6.4.2-gfx90a/lib/libProteusPass.so+0x62962)
 #7 0x000015555357213a (anonymous namespace)::ProteusPassImpl::run(llvm::Module&, bool) ProteusPass.cpp:0:0
 #8 0x000015555356dc6e llvm::detail::PassModel<llvm::Module, (anonymous namespace)::ProteusPass, llvm::AnalysisManager<llvm::Module>>::run(llvm::Module&, llvm::AnalysisManager<llvm::Module>&) ProteusPass.cpp:0:0
 #9 0x0000000001def754 llvm::PassManager<llvm::Module, llvm::AnalysisManager<llvm::Module>>::run(llvm::Module&, llvm::AnalysisManager<llvm::Module>&) (/opt/rocm-6.4.2/lib/llvm/bin/clang-19+0x1def754)
#10 0x00000000025b9a5e (anonymous namespace)::EmitAssemblyHelper::RunOptimizationPipeline(clang::BackendAction, std::unique_ptr<llvm::raw_pwrite_stream, std::default_delete<llvm::raw_pwrite_stream>>&, std::unique_ptr<llvm::ToolOutputFile, std::default_delete<llvm::ToolOutputFile>>&, clang::BackendConsumer*) (/opt/rocm-6.4.2/lib/llvm/bin/clang-19+0x25b9a5e)
#11 0x00000000025bcaf3 clang::EmitBackendOutput(clang::DiagnosticsEngine&, clang::HeaderSearchOptions const&, clang::CodeGenOptions const&, clang::TargetOptions const&, clang::LangOptions const&, llvm::StringRef, llvm::Module*, clang::BackendAction, llvm::IntrusiveRefCntPtr<llvm::vfs::FileSystem>, std::unique_ptr<llvm::raw_pwrite_stream, std::default_delete<llvm::raw_pwrite_stream>>, clang::BackendConsumer*) (/opt/rocm-6.4.2/lib/llvm/bin/clang-19+0x25bcaf3)
#12 0x0000000002b3bfed clang::BackendConsumer::HandleTranslationUnit(clang::ASTContext&) (/opt/rocm-6.4.2/lib/llvm/bin/clang-19+0x2b3bfed)
#13 0x0000000004421e7c clang::ParseAST(clang::Sema&, bool, bool) (/opt/rocm-6.4.2/lib/llvm/bin/clang-19+0x4421e7c)
#14 0x0000000002b3ca35 clang::CodeGenAction::ExecuteAction() (/opt/rocm-6.4.2/lib/llvm/bin/clang-19+0x2b3ca35)
#15 0x0000000002db8991 clang::FrontendAction::Execute() (/opt/rocm-6.4.2/lib/llvm/bin/clang-19+0x2db8991)
#16 0x0000000002d465ea clang::CompilerInstance::ExecuteAction(clang::FrontendAction&) (/opt/rocm-6.4.2/lib/llvm/bin/clang-19+0x2d465ea)
#17 0x0000000002e906a3 clang::ExecuteCompilerInvocation(clang::CompilerInstance*) (/opt/rocm-6.4.2/lib/llvm/bin/clang-19+0x2e906a3)
#18 0x0000000000b7d6fc cc1_main(llvm::ArrayRef<char const*>, char const*, void*) (/opt/rocm-6.4.2/lib/llvm/bin/clang-19+0xb7d6fc)
#19 0x0000000000b76bbf ExecuteCC1Tool(llvm::SmallVectorImpl<char const*>&, llvm::ToolContext const&) (/opt/rocm-6.4.2/lib/llvm/bin/clang-19+0xb76bbf)
#20 0x0000000000b79237 clang_main(int, char**, llvm::ToolContext const&) (/opt/rocm-6.4.2/lib/llvm/bin/clang-19+0xb79237)
#21 0x0000000000abd413 main (/opt/rocm-6.4.2/lib/llvm/bin/clang-19+0xabd413)
#22 0x0000155553b71865 __libc_start_main (/lib64/libc.so.6+0x3a865)
#23 0x0000000000b7673e _start (/opt/rocm-6.4.2/lib/llvm/bin/clang-19+0xb7673e)
clang++: error: unable to execute command: Aborted (core dumped)
clang++: error: clang frontend command failed due to signal (use -v to see invocation)

[ggeorgakoudis]

What are the multiple users in this case?

[ZwFink]

This causes a crash with the same error. Maybe it is helpful for your investigation

#include <cstdio>

#include "gpu_common.h"
#include <proteus/JitInterface.h>

// ---------- kernel template (same as existing tests) ----------

template <typename T>
__global__ __attribute__((annotate("jit"))) void kernel(T LB) {
  LB();
}

// ---------- direct launcher: no nesting (works) ----------

template <typename T> void run_direct(T &&LB) {
  proteus::register_lambda(LB);
  kernel<<<1, 1>>>(LB);
  gpuErrCheck(gpuDeviceSynchronize());
}

// ---------- wrapped launcher: one layer of nesting (broken) ----------
// Mirrors: jit_parallel_for_all creates [=](int idx){ body(j,k); }

template <typename Body>
void run_wrapped(Body body) {
  int extra = 1; // simulates j0/k0/k_ext captured by the wrapper
  auto wrapper = [=] __device__() __attribute__((annotate("jit"))) {
    if (extra) // use extra to prevent it from being optimized away
      body();
  };
  proteus::register_lambda(wrapper);
  kernel<<<1, 1>>>(wrapper);
  gpuErrCheck(gpuDeviceSynchronize());
}

int main() {
  int A = 42;

  printf("=== direct ===\n");
  {
    auto lambda = [=, A = proteus::jit_variable(A)] __device__()
        __attribute__((annotate("jit"))) {
      printf("direct A = %d\n", A);
    };
    run_direct(lambda);
  }

  printf("=== wrapped ===\n");
  {
    auto inner = [=, A = proteus::jit_variable(A)] __device__() {
      printf("wrapped A = %d\n", A);
    };
    run_wrapped(inner);
  }

  return 0;
}