search

Quuxplusone / LLVMBugzillaTest

0 stars 0 forks source link

CUDA support for __global__ lambdas #26481

Open Quuxplusone opened 8 years ago

Quuxplusone commented 8 years ago
Bugzilla Link PR26482
Status NEW
Importance P normal
Reported by Justin Lebar (:jlebar) (justin.lebar@gmail.com)
Reported on 2016-02-04 15:46:18 -0800
Last modified on 2016-05-16 17:40:30 -0700
Version unspecified
Hardware PC Linux
CC hfinkel@anl.gov, llvm-bugs@lists.llvm.org, richard-llvm@metafoo.co.uk, tra@google.com
Fixed by commit(s)
Attachments
Blocks
Blocked by
See also 
The following code should compile (and generate a kernel lambda we can invoke).

void foo() {
  auto x = []()__global__ { };
}

Instead it gives the following error:

test.cu:2:15: error: kernel function type '<dependent type> () const' must have
void return type
  auto x = [] () __global__ { };
Quuxplusone commented 8 years ago 
You get the same error for

  __global__ auto foo() {}

Do you also consider that case to be a bug?
Quuxplusone commented 8 years ago 
(In reply to comment #1)
> You get the same error for
>
>   __global__ auto foo() {}
>
> Do you also consider that case to be a bug?

Aha.  Yes, I think that's also a bug.  So the return-type issue is one thing
that needs to be fixed to support __global__ lambdas (with implicit return
types), although I suspect it's not the only thing.
Quuxplusone commented 8 years ago

Another problem is that global function must be a free function or static method. Observed lambda's operator() instances currently appear to be neither.