Community
cancel
Showing results for 
Search instead for 
Did you mean: 
leggett__charles
Beginner
643 Views

calling printf from within kernel

 

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?

Tags (1)
0 Kudos
12 Replies
leggett__charles
Beginner
643 Views

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
Moderator
643 Views

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
Beginner
643 Views

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
Moderator
643 Views

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
Moderator
643 Views

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
Moderator
643 Views

Hi Charles,

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

leggett__charles
Beginner
643 Views

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
Moderator
643 Views

Hi Charles,

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

-Abhishek

AbhishekD_Intel
Moderator
643 Views

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.

127 Views

Hi Charles,

Can you confirm whether the issue is resolved with Gold?


Regards,

Subarna


111 Views

Closing this thread as the bug is already fixed.


111 Views

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


Reply