[SYCL] Incorrect Kernel Lambda Capture

[agozillon]

#include <CL/sycl.hpp>
using namespace cl::sycl;

int main() {
  queue q;
  q.submit([&](handler &cgh) {
    int w = 512;
    cgh.single_task<class are_you_broken>([=]() {
      printf("%d \n", w);
    });
  });
  q.wait();
  return 0;
}

So it's a pretty simple example, just trying to capture a value and print it, however when using ComputeCPP (1.0) I'll get the correct value 512 inside the single_task, in the Intel implementation I get an arbitrary large integer value (e.g. 946443008). The contrived example should be directly and quickly usable to see if it is indeed a bug in the implementation or if I'm mistaken and its an issue on my end!

I used the following command: clang++ -std=c++11 -fsycl possibly_broken_capture.cpp -o possibly_broken_capture -lsycl -lOpenCL

When I compile just the device side code the header indicates that SemaSYCL does find the parameter and treats it as { kernel_param_kind_t::kind_std_layout, 4, 0 }.

Note: In my test version of the example I have a device selector that forces it to choose Intel OCL, so both SYCL runtimes should find the same OCL implementation.

[Ruyk]

Note that example requires a wait before the end of main: The SYCL queue destructor doesn't have a mandatory wait on it, text below from the SYCL 1.2.1 spec rev 4:

Note that the destructors of other SYCL objects (cl::sycl::queue, cl::sycl::context. . . ) do not block. Only
a cl::sycl::buffer or cl::sycl::image destructor might block. The rationale is that an object without any side
effect on the host does not need to block on destruction as it would impact the performance. So it is up to the programmer to use a method to wait for completion in some cases if this does not fit the goal.

ComputeCpp had a wait on the destructor for historic reasons, but that is soon to be gone (likely in the next release).

[agozillon]

Thank you for clarifying ComputeCpp behavior, I remember a discussion on the behavior you describe recently and it reminds me to be a little more careful in future when I try to compare across implementations!

I have edited the code snippet above to contain a queue wait and ran it on my machine again and unfortunately it still does give an undefined value on execution.

[agozillon]

Spent a little more time looking into this issue I opened and I made the incorrect assumption that it was a capture issue! Made a new test where I write the value to a buffer and then printf on host and the output is correct.

Seems to be a user error on my end as I was working under the incorrect assumption that printf should work on the device, however the SYCL specification defines a stream and not printf as part of the specification so previously an invalid test (use of printf appears to be implementation defined I believe, my apologies).

I suppose this issue can be closed! Unless printf is intended to work on the device in this implementation, in which case perhaps the issue needs to be renamed or refiled?

[keryell]

and then printf on host and the output is correct.

Do not use printf in C++ :-) It is a function from the 1970's...

[keryell]

@agozillon You can close this, if I understand correctly.