topic Hi Charles, in Intel® oneAPI DPC++/C++ Compiler

calling printf from within kernel

leggett__charles — Mon, 10 Feb 2020 13:31:03 GMT

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?

Figured it out after digging

leggett__charles — Tue, 11 Feb 2020 08:37:57 GMT

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.

Hi Charles,

AbhishekD_Intel — Tue, 11 Feb 2020 14:14:48 GMT

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.

yes - please note that the

leggett__charles — Tue, 11 Feb 2020 17:07:00 GMT

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?

Hi Charles,

AbhishekD_Intel — Wed, 12 Feb 2020 07:20:00 GMT

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

Hi Charles,

AbhishekD_Intel — Wed, 12 Feb 2020 13:31:00 GMT

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

Hi Charles,

AbhishekD_Intel — Tue, 18 Feb 2020 04:40:16 GMT

Hi Charles,

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

Hi Abhishek -

leggett__charles — Tue, 18 Feb 2020 06:07:23 GMT

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.

Hi Charles,

AbhishekD_Intel — Tue, 18 Feb 2020 06:53:25 GMT

Hi Charles,

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

-Abhishek

Hi Charles,

AbhishekD_Intel — Fri, 21 Feb 2020 13:36:39 GMT

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.

Re:calling printf from within kernel

Subarnarek_G_Intel — Mon, 31 May 2021 06:33:07 GMT

Hi Charles,

Can you confirm whether the issue is resolved with Gold?

Regards,

Subarna

Re:calling printf from within kernel

Subarnarek_G_Intel — Wed, 02 Jun 2021 15:01:48 GMT

Closing this thread as the bug is already fixed.

Re:calling printf from within kernel

Subarnarek_G_Intel — Wed, 02 Jun 2021 15:02:04 GMT

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