calling printf from within kernel

[leggett__charles]

The release notes for beta04 state the following is supported:

   - "printf" support in device code

However, when I try this I get an error message:

error: SYCL kernel cannot call a variadic function

Is there some special incantation?

[leggett__charles]

Figured it out after digging through the compiler source:

#ifdef __SYCL_DEVICE_ONLY__
  #define CONSTANT __attribute__((opencl_constant))
#else
  #define CONSTANT
#endif

static const CONSTANT char FMT[] = "n: %d\n";
sycl::intel::experimental::printf(FMT, 7);

I would suggest fixing the release notes

However, after a number of iterations in a parallel_for, it will segfault.

[AbhishekD_Intel]

Hi Charles,

It is specified in the release notes the complete implementation of "printf" with the example.

Can you be more specific regarding your iterations in parallel_for? It will good if you can share your code snippet.

[leggett__charles]

hi Abhishek -

yes - please note that the release notes are incorrect - they're missing the "__" at the beginning and end of the SYCL_DEVICE_ONLY macro.

In my simple example, I'm printing out some integer counters such as cl::sycl::item<1>.get_linear_id()

 auto k2 = [=](cl::sycl::item<1> item) {
 ...
   unsigned int i = item.get_linear_id(); 
   static const CONSTANT char FMT[] = "n: %u\n";
   sycl::intel::experimental::printf(FMT, i);
...

when the work size is between 2^16 and 2^17, first it stops printing out the integers (though continues to print out the rest of the string part of the statement, then as the work size increases further, it segfaults.

If I try to print out more values in the same statement (eg several integers), it will start failing for smaller work sizes. Maybe a buffer is getting filled?

[AbhishekD_Intel]

Hi Charles,

Thanks for the information we will take a note of this thing and will concern to our engineering team.

We will try to increase the work size as specified and will give you more clear information regarding your issue.

Thanks

-Abhishek

[AbhishekD_Intel]

Hi Charles,

I tried the way you are solving the problem and its working for me, even for the workgroup size more than 2^17 it's executing the code and also printing the whole string with numbers without any segFault, you can find the code that I have used:

#ifdef __SYCL_DEVICE_ONLY__
          #define CONSTANT __attribute__((opencl_constant))
#else
          #define CONSTANT
#endif


auto exception_handler = [] (cl::sycl::exception_list exceptions) {
      for (std::exception_ptr const& e : exceptions) {
         try {
            std::rethrow_exception(e);
         } catch(cl::sycl::exception const& e) {
         std::cout << "Caught asynchronous SYCL exception:\n"
                   << e.what() << std::endl;
         }
      }
   };


        cl::sycl::queue queue(sycl::default_selector{}, exception_handler);

        std::cout << "Running on "
                << queue.get_device().get_info<cl::sycl::info::device::name>()
                << "\n";
        {

                queue.submit([&] (cl::sycl::handler& cgh) {
                                
                                cgh.parallel_for<class vector_addition>(cl::sycl::range<1>(N), [=](cl::sycl::item<1> idx) {
                                                
                                                unsigned int i = idx.get_linear_id();
                                                static const CONSTANT char FMT[] = "n: %u\n";
                                                sycl::intel::experimental::printf(FMT, i);

                                                });
                                });
        }

As the workgroup size limit is between 0 to 4,294,967,295 it will run without any segFault, I will suggest you to debug your code and if possible send your source code so that we can debug the problem.

-Abhishek

[AbhishekD_Intel]

Hi Charles,

I hope your issue got resolved, Can we close this thread?

[leggett__charles]

Hi Abhishek - I don't think it's resolved - there's definitely something odd going on. If you can provide me with an email address, I'll send you my code. Otherwise I will be at the Aurora workshop next week, and can discuss it with the dpcpp developers there.

[AbhishekD_Intel]

Hi Charles,

You can send me your code privately so that we can debug it.

-Abhishek

[AbhishekD_Intel]

Hi Charles,

We tested the code that you shared privately, and on our platform with Ubuntu 18.04 and beta-04 Toolkit and its working perfectly.

We are working on the logs shared by you privately and will get back to you soon.

[Subarnarek_G_Intel]

Hi Charles,

Can you confirm whether the issue is resolved with Gold?

Regards,

Subarna

[Subarnarek_G_Intel]

Closing this thread as the bug is already fixed.

[Subarnarek_G_Intel]

This issue has been resolved and we will no longer respond to this thread. If you require additional assistance from Intel, please start a new thread. Any further interaction in this thread will be considered community only