[llvm-bugs] [Bug 26341] New: Cuda __device__ lambda does not work

[via llvm-bugs]

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

            Bug ID: 26341
           Summary: Cuda __device__ lambda does not work
           Product: clang
           Version: trunk
          Hardware: PC
                OS: Linux
            Status: NEW
          Severity: normal
          Priority: P
         Component: -New Bugs
          Assignee: unassignedclangbugs at nondot.org
          Reporter: crtrott at sandia.gov
                CC: llvm-bugs at lists.llvm.org
    Classification: Unclassified

Created attachment 15731
  --> https://llvm.org/bugs/attachment.cgi?id=15731&action=edit
Reproduction code

This is an experimental feature in CUDA 7.5 and expected to be non experimental
in CUDA 8. The attached simple test code fails to compile. 

Code:
  run( [=] __device__ (int i) {
    d_a[i] = d_c;
  });


Error:
main.cpp:23:12: error: lambda requires '()' before attribute specifier
  run( [=] __device__ (int i) {
           ^
           () 
/usr/local/cuda/include/host_defines.h:189:9: note: expanded from macro
'__device__'
        __location__(device)
        ^
/usr/local/cuda/include/host_defines.h:88:9: note: expanded from macro
'__location__'
        __annotate__(a)
        ^
/usr/local/cuda/include/host_defines.h:86:9: note: expanded from macro
'__annotate__'
        __attribute__((a))
        ^
main.cpp:23:23: error: expected body of lambda expression
  run( [=] __device__ (int i) {
                      ^
main.cpp:23:28: error: expected '(' for function-style cast or type
construction
  run( [=] __device__ (int i) {
                       ~~~ ^
3 errors generated.


I actually like to see to be able to do __host__ __device__ (which does not
work with NVCC right now) so that you can dispatch the lambda to the host or to
the device. Effectively I just want the operator of the auto generated struct
to be marked __host__ __device__. I am responsible to only capture stuff which
is accesible on the device. Furthermore one might want to restrict a __device__
lambda to only capture by value. i.e. [=] __host__ __device__ (...) {...}

-- 
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/20160127/46518f7d/attachment.html>