[llvm-bugs] [Bug 48866] New: Kernel launch failure with lambdas in CUDA on Windows

[via llvm-bugs]

https://bugs.llvm.org/show_bug.cgi?id=48866

            Bug ID: 48866
           Summary: Kernel launch failure with lambdas in CUDA on Windows
           Product: clang
           Version: unspecified
          Hardware: PC
                OS: Windows NT
            Status: NEW
          Severity: normal
          Priority: P
         Component: CUDA
          Assignee: unassignedclangbugs at nondot.org
          Reporter: joachim at joameyer.de
                CC: llvm-bugs at lists.llvm.org

Given multiple (nested) lambdas as CUDA kernel parameters, the mangling/naming
seems to be inconsistent between host and device and thus `cudaLaunchKernel`
fails with error code 98. Might be related to the lambda numbering issues,
solved in Clang 10 (https://reviews.llvm.org/D68818)?

A rather minimal reproducer is:
```
#include <exception>
#include <iostream>

template <class T>
__global__ void kernel(int *t, int n, T lambda)
{
    if(threadIdx.x < n)
        t[threadIdx.x] = lambda(t[threadIdx.x]);
}

int main(){
    int* t=nullptr;
    cudaMalloc(&t,12);
    auto lambda = [](int t){
        return t++;
    };
    kernel<<<1, 32>>>(t, 3, lambda);

    [&]()
    {
        auto lambda2 = [](int t){
            return t++;
        };
        kernel<<<1, 32>>>(t, 3, lambda2);
    }();
    if(cudaGetLastError())
    {
        std::cout << "err";
        cudaFree(t);
        std::terminate();
    }
    cudaFree(t);
    return 0;
}
```

-- 
You are receiving this mail because:
You are on the CC list for the bug.
-------------- next part --------------
An HTML attachment was scrubbed...
URL: <http://lists.llvm.org/pipermail/llvm-bugs/attachments/20210124/55cac50f/attachment.html>