unitrace reports work-group size {0; 0; 0} for OpenMP offload kernels
Nadie ha tomado este issue todavía.
Evaluación
- Dificultad
- 3/5
- Tiempo estimado
- 1-2 días
- Aptitud para principiantes
- 70/100
- Tipo de issue
- Error
- Claridad
- Bien especificado
- Estado de actividad
- Activo
- Stack tecnológico
- cpp
- Área
- devtools, performance
Línea de trabajo
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.
Escrito por el modelo de indexación a partir del texto del issue.
Descripción
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))) {
- Lenguaje dominante
- C++
- Estrellas
- 278
- Forks
- 75
- Métricas de merge de PR
- Sin PR fusionados en 30 d
Preparar el entorno
- Sin Dockerfile ni archivo de Docker Compose
- Tiene una plantilla de pull request
- Leer la guía de contribución
Primeros pasos
- Lee el issue completo y luego la guía de contribución del proyecto.
- Comenta en el issue que vas a ocuparte — evita que dos personas hagan lo mismo.
- Haz un fork del repositorio y trabaja en una rama.
- Abre un pull request que haga referencia al número del issue.
Más de intel/pti-gpu
-
Dificultad 2/5 1-3 horas Aptitud para principiantes 72/100
-
Dificultad 1/5 Menos de una hora Aptitud para principiantes 92/100
-
Dificultad 4/5 3-5 días Aptitud para principiantes 48/100
-
Dificultad 4/5 3-5 días Aptitud para principiantes 20/100
-
Dificultad 3/5 1-2 días Aptitud para principiantes 50/100
Todos los issues de intel/pti-gpu
Issues similares
-
Dificultad 2/5 1-3 horas Aptitud para principiantes 85/100
Los mantenedores suelen responder en 1 día
-
parse_selu lowering loses catastrophic cancellation → silent 0.0 for moderate positive inputsAbierto
Dificultad 2/5 1-3 horas Aptitud para principiantes 74/100
ROCm/AMDMIGraphX#5490 ·
Los mantenedores suelen responder en 3 días
-
符号半角全角问题Abierto
Dificultad 1/5 Menos de una hora Aptitud para principiantes 82/100
-
Dificultad 2/5 1-3 horas Aptitud para principiantes 74/100
arvidn/libtorrent#9126 ·
Los mantenedores suelen responder en 1 día
-
bug
Dificultad 2/5 1-3 horas Aptitud para principiantes 72/100
Sunrisepeak/mcpp-language-server#80 ·
Los mantenedores suelen responder en 1 día