如何在内核卸载到 GPU 的情况下调试以下 DPC++ 程序？

Posted 2023-02-22

【中文标题】如何在内核卸载到 GPU 的情况下调试以下 DPC++ 程序？

【英文标题】：How to debug the below DPC++ program with the kernel offloaded to the GPU?

【发布时间】：2021-04-27 05:10:41

【问题描述】：

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

下载链接是：

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

如何在劣质和线程之间切换？请查看下面提供的代码。

#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_rangelength;
    buffer buffer_ininput, data_range;
    buffer buffer_outoutput, 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;

【参考方案1】：

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

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

您可以使用“thread”命令（例如“thread 3:4”、“thread :6”或“thread 7”）切换线程以及 SIMD 通道以更改上下文。第一个命令切换到线程 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-调试-1.html]