如何调试下面的DPC++程序，把内核卸载到GPU上？

Question

我想在英特尔GNU项目调试器中运行DPC++程序。我已从英特尔OneAPI Basekit下载了英特尔GDB。它预装了OneAPI Basekit。

要下载的链接为：

https://software.intel.com/content/www/us/en/develop/tools/oneapi/base-toolkit.html#gs.ynm6aj

如何调试下面的DPC++程序，把内核卸载到GPU上？如何在下级和线程之间切换？请查看下面提供的代码。

#include <CL/sycl.hpp>
#include <iostream>
// Location of file: <oneapi-root>/dev-utilities/<version>/include
#include "dpc_common.hpp"
#include "selector.hpp"

using namespace std;
using namespace sycl;

// A device function, called from inside the kernel.
static size_t GetDim(id<1> wi, int dim) {
  return wi[dim];
}

int main(int argc, char *argv[]) {
  constexpr size_t length = 64;
  int input[length];
  int output[length];

  // Initialize the input
  for (int i = 0; i < length; i++)
    input[i] = i + 100;

  try {
    CustomSelector selector(GetDeviceType(argc, argv));
    queue q(selector, dpc_common::exception_handler);
    cout << "[SYCL] Using device: ["
         << q.get_device().get_info<info::device::name>()
         << "] from ["
         << q.get_device().get_platform().get_info<info::platform::name>()
         << "]\n";

    range data_range{length};
    buffer buffer_in{input, data_range};
    buffer buffer_out{output, data_range};

    q.submit([&](auto &h) {
      accessor in(buffer_in, h, read_only);
      accessor out(buffer_out, h, write_only);

      // kernel-start
      h.parallel_for(data_range, [=](id<1> index) {
        size_t id0 = GetDim(index, 0);
        int element = in[index];  // breakpoint-here
        int result = element + 50;
        if (id0 % 2 == 0) {
          result = result + 50;  // then-branch
        } else {
          result = -1;  // else-branch
        }
        out[index] = result;
      });
      // kernel-end
    });

    q.wait_and_throw();
  } catch (sycl::exception const& e) {
    cout << "fail; synchronous exception occurred: " << e.what() << "\n";
    return -1;
  }

  // Verify the output
  for (int i = 0; i < length; i++) {
    int result = (i % 2 == 0) ? (input[i] + 100) : -1;
    if (output[i] != result) {
      cout << "fail; element " << i << " is " << output[i] << "\n";
      return -1;
    }
  }

  cout << "success; result is correct.\n";
  return 0;
}

Answer 1

可以使用调试器列出应用程序的线程。打印的信息包括线程ids和线程当前停止的位置。对于GPU线程，调试器还会打印活动的SIMD通道。GDB is以以下格式显示线程：

<inferior_number>.<thread_number>:<SIMD Lane/s>

您可以使用" thread“命令切换线程和SIMD通道以更改上下文，例如"thread 3:4”、"thread :6“或"thread 7”。第一个命令切换到线程3和SIMD通道4。第二个命令切换到当前线程内的SIMD通道6。第三个命令切换到线程7。选定的默认通道将是先前选定的通道(如果它处于活动状态)或线程中的第一个活动通道。有关更多详细信息，请参阅下面的链接，该链接解释了如何在GPU设备上调试代码。

https://software.intel.com/content/www/us/en/develop/documentation/debugging-dpcpp-linux/top/debug-a-dpc-application-on-a-gpu/basic-debugging-1.html