Intel® oneAPI Data Parallel C++
Support for Intel® oneAPI DPC++ Compiler, Intel® oneAPI DPC++ Library, Intel ICX Compiler , Intel® DPC++ Compatibility Tool, and GDB*
560 Discussions

How can I debug kernel code of DPC++ with ESIMD extension

tanzl_ustc
Beginner
2,137 Views

Hello, I'm a senior student. I'm studying DPC++ recently.

I found some problems in debugging the kernel code of DPC++ with ESIMD extension.

As  shown in the following image, the guide book told me to use "stream" to debug.

tanzl_ustc_0-1642474771469.png

However, this method will fail on DPC + + code with esimd extension.

tanzl_ustc_1-1642475100649.png

As shown in the figure above, when I add the stream method in lines 49 and 52 of the kernel code, the official sample code that can be compiled normally will report errors in compilation (the compilation command I use is: clang + + - fsycl vadd_usm.cpp - O vaddout). When I remove the "sycl_esimd_kernel" in line 51, the compilation can pass normally, but I can't continue to use esimd extension.

If "sycl_esimd_kernel" and all code related to ESIMD extension are removed from the code, the stream method can be used normally

I also tried to debug kernel code with std:: cout and GDB, but they all failed.

Could you tell me an effective method to debug kernel code and track the value of internal variables? Thank you!

 

0 Kudos
11 Replies
HemanthCH_Intel
Moderator
2,094 Views

Hi,


Thanks for reaching out to us.


There are some restrictions that apply when using the ESIMD extension. For more information please refer to the below link:

https://www.intel.com/content/www/us/en/develop/documentation/oneapi-dpcpp-cpp-compiler-dev-guide-and-reference/top/optimization-and-programming-guide/vectorization/explicit-vector-programming/explicit-simd-sycl-extension.html


Could you please let us know the OS version, IDE version(Visual Studio/Eclipse), and OneAPI toolkit version you are using?


Thanks & Regards,

Hemanth.


0 Kudos
tanzl_ustc
Beginner
2,075 Views

Hi,

 

Thanks for your early reply! I have read the restrictions that apply when using the ESIMD extension.

The OS version:Linux user-X299-UD4-Pro 5.4.48-xe-max #1 SMP Sun Dec 12 17:04:51 CST 2021 x86_64 x86_64 x86_64 GNU/Linux

tanzl_ustc_0-1642559671435.png

 

IDE version: I am using Visual Studio Code version 1.63.2

 

Oneapi toolkit version:l_BaseKit_p_2022.1.1.119

 

Besides, initially I just simply followed the instructions in 

https://registrationcenter-download.intel.com/akdlm/irc_nas/18445/l_BaseKit_p_2022.1.1.119_offline.sh

After downloading the toolkit I did nothing else.

 

But the ESIMD programs can run normally after I set the environment with the instructions offered by

 

https://github.com/intel/llvm/blob/sycl/sycl/doc/GetStartedGuide.md#build-dpc-toolchain

 

even without the download of oneapi toolkit.

 

So it seems to me that oneapi toolkit may not be a critical factor to run the program

 

THANKS FOR READING

0 Kudos
HemanthCH_Intel
Moderator
2,064 Views

Hi,

 

We can debug the kernel using the method "int cl::sycl::ext::oneapi::experimental::printf(const FormatT * __format,/Args... args )".

For using the sycl::ext::oneapi::experimental::printf method, we need to add the below code:

 

 

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


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

 

 

I have attached a zip folder which has 3 files: 

1)vadd_usm.cpp (without using sycl::ext::oneapi::experimental::printf method)

2)vadd_usm_esimd.cpp (using sycl::ext::oneapi::experimental::printf method)

3)output screenshot.

 

Could you please refer to the vadd_usm_esimd code for using the sycl::ext::oneapi::experimental::printf method and the screenshot attached.

 

Thanks & Regards,

Hemanth.

 

0 Kudos
tanzl_ustc
Beginner
2,045 Views

Gratitude! Now I can use printf in kernel function.

 

However, the output seemed to be strange.I tried vadd_usm_esimd.cpp without any changes, but the output is like:

tanzl_ustc_0-1642667474357.png

 

This is a little bit weird to me,because normally the message "Passed" should occur after "inside kernel:......".

 

It looks like there is a problem with the execution order of the program.

 

I was wondering if it is due to the problems of my environment?

 

Thank you for reading.

 

 

 

0 Kudos
tanzl_ustc
Beginner
2,016 Views

Besides, I think I have been trapped in a really confusing problem:

Also in vadd_usm.cpp, when I run vadd_usm_ver1, the output is :

tanzl_ustc_1-1642852880225.png

when run vadd_usm_ver2.cpp( line 51 and line 52 were changed compared to vadd_usm_ver1.cpp) ,the output is:

tanzl_ustc_2-1642852988236.png

I am really confused why such little changes could easily lead to a huge difference on output. And in my own program gemm_genx_dpcpp.cpp, output generated by line 155 also shows that there might be a problem when I try to print out the inner value of simd vector.Even if I reduce the number of threads to 1, this problem still exists.

 

I have attached a zip with three files:

1.vadd_usm_ver1.cpp

2.vadd_usm_ver2.cpp

3.gemm_genx_dpcpp.cpp

 

Could you please check the files and see where the problem is.

 

Thank you very much.

 

 

0 Kudos
v-klochkov
Employee
1,511 Views

Hello and thank you for reporting the issue here.

 

I tried to reproduce the problem with inconsistent result of printf.
Accordingly to our internal tracker the original problem was on Gen9 ("Intel(R) UHD Graphics 630").
With the fresh clang++ compiler (https://github.com/intel/llvm/releases/tag/sycl-nightly%2F20220726) and GPU RT (1.3.23368), the test runs properly on Gen9.

 

I changed only 1 line in the original source to work-around a known problem/error with simd subscript operator used in RHS/read-context:

    sycl::ext::oneapi::experimental::printf(FMT, i, va[i]);
    //==>
    sycl::ext::oneapi::experimental::printf(FMT, i, (int)va[i]);

A test case based on the one reported above is added to the test suite: https://github.com/intel/llvm-test-suite/pull/1111

0 Kudos
HemanthCH_Intel
Moderator
1,938 Views

Hi,


Thanks for reporting us.


We have reported the issue to the concerned development team. They are looking into your issue and will get back to you soon after fixing your issue.


Thanks & Regards,

Hemanth.


0 Kudos
tanzl_ustc
Beginner
1,840 Views

Hi,

sorry for bothering.

I was wondering if my issues were fixed?

Thanks & Regards

Tanzl

0 Kudos
HemanthCH_Intel
Moderator
1,814 Views

Hi,


>>"I was wondering if my issues were fixed?

As mentioned earlier, we have reported this issue to the concerned development team. They are looking into your issue.


Thanks & Regards,

Hemanth.


0 Kudos
HemanthCH_Intel
Moderator
1,414 Views

Hi,


Thank you for your patience. The issue raised by you has been fixed in the 2022.2 version. Please download and let us know if this resolves your issue.


Thanks & Regards,

Hemanth


0 Kudos
HemanthCH_Intel
Moderator
1,396 Views

Hi,


We assume that your issue is resolved. If you need any additional information, please post a new question as this thread will no longer be monitored by Intel.


Thanks & Regards,

Hemanth


0 Kudos
Reply