unitrace reports work-group size {0; 0; 0} for OpenMP offload kernels

Đang mở
  • Có gì không ổn?
#128 0 bình luận 0 reaction 0 người được giao Xem trên GitHub

Chưa có ai nhận issue này.

Đánh giá

Độ khó
3/5
Thời gian dự kiến
1-2 ngày
Mức phù hợp với người mới
70/100
Loại issue
Lỗi
Độ rõ ràng
Đặc tả rõ ràng
Mức độ hoạt động
Sôi nổi
Công nghệ
cpp
Lĩnh vực
devtools, performance

Hướng nghiên cứu

Start in tools/unitrace/src/levelzero/ze_collector.h at the zeKernelSetGroupSize handler and the launch-call handler shown in the issue. Trace how launch parameters are recorded, then build and run the provided OpenMP offload reproducer with unitrace. Done when the reported group size is {64; 1; 1} rather than {0; 0; 0}, and existing launch handling still works.

Do mô hình lập chỉ mục viết ra từ nội dung của issue.

Mô tả

unitrace reports the work-group size of every OpenMP offload kernel as {0; 0; 0}. The OpenMP runtime (libomptarget, Level Zero plugin) launches kernels with zeCommandListAppendLaunchKernelWithArguments and passes the group size in its pGroupSizes argument. It does not call zeKernelSetGroupSize. unitrace records the group size only in its handler of zeKernelSetGroupSize, so for these kernels the size stays 0.

Reproducer, built with icpx -fiopenmp -fopenmp-targets=spir64 repro.cpp -o repro:

#include <cstdio>

int main() {
  const int N = 1 << 20;
  static float a[N];
#pragma omp target teams distribute parallel for map(from : a[0 : N])
  for (int i = 0; i < N; ++i)
    a[i] = 2.0f * i;
  printf("%f\n", a[N - 1]);
  return 0;
}

unitrace -d -v ./repro reports:

"__omp_offloading_821_3df0002__Z4main_l6[SIMD32 {16384; 1; 1} {0; 0; 0}]", 1, 9760, ...

With the change below the same run reports {16384; 1; 1} {64; 1; 1}. The change passes the group size of the launch call to the existing zeKernelSetGroupSize handler:

--- a/tools/unitrace/src/levelzero/ze_collector.h
+++ b/tools/unitrace/src/levelzero/ze_collector.h
@@ -4908,6 +4908,12 @@ class ZeCollector {
     ze_instance_data.instrument_ = false;
     ze_instance_data.query_ = nullptr;
     ZeCollector* collector = reinterpret_cast<ZeCollector*>(global_data);
+    // This launch call passes the group size instead of zeKernelSetGroupSize.
+    if (params->pgroupSizes != nullptr) {
+      ze_group_size_t gs = *(params->pgroupSizes);
+      ze_kernel_set_group_size_params_t p{params->phKernel, &gs.groupSizeX, &gs.groupSizeY, &gs.groupSizeZ};
+      OnExitKernelSetGroupSize(&p, ZE_RESULT_SUCCESS, global_data, instance_data);
+    }
     if (collector->ShouldCaptureAppend(*(params->phCommandList),
                                        *(params->pnumWaitEvents),
                                        *(params->pphWaitEvents))) {
Ngôn ngữ chính
C++
Star
278
Fork
75
Không có pull request nào được merge trong 30 ngày

Chuẩn bị môi trường

Bắt đầu từ đâu

  1. Đọc hết issue, rồi đọc hướng dẫn đóng góp của dự án.
  2. Bình luận trên issue rằng bạn sẽ nhận — tránh hai người làm cùng một việc.
  3. Fork repository và làm thay đổi trên một nhánh.
  4. Mở pull request có tham chiếu số hiệu của issue.

Issue khác của intel/pti-gpu

Tất cả issue của intel/pti-gpu

Issue tương tự

Thêm issue về C++

Nhận issue mới trong hộp thư của bạn

Bản tóm tắt ngắn những issue GitHub phù hợp với người mới.