topic Re:How can I debug kernel code of DPC++ with ESIMD extension in Intel® oneAPI DPC++/C++ Compiler

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

tanzl_ustc — Tue, 18 Jan 2022 03:17:12 GMT

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.

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

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!

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

HemanthCH_Intel — Tue, 18 Jan 2022 13:12:11 GMT

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.

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

tanzl_ustc — Wed, 19 Jan 2022 03:11:46 GMT

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

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

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

HemanthCH_Intel — Fri, 27 May 2022 09:43:17 GMT

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.

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

tanzl_ustc — Thu, 20 Jan 2022 08:39:55 GMT

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:

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.

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

tanzl_ustc — Sat, 22 Jan 2022 12:15:03 GMT

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 :

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

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.

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

HemanthCH_Intel — Thu, 27 Jan 2022 10:14:44 GMT

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.

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

tanzl_ustc — Tue, 15 Feb 2022 06:50:24 GMT

Hi,

sorry for bothering.

I was wondering if my issues were fixed?

Thanks & Regards

Tanzl

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

HemanthCH_Intel — Wed, 23 Feb 2022 08:04:01 GMT

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.

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

v-klochkov — Tue, 26 Jul 2022 22:38:47 GMT

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

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

HemanthCH_Intel — Mon, 29 Aug 2022 07:35:25 GMT

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

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

HemanthCH_Intel — Mon, 05 Sep 2022 07:18:47 GMT

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