llvm-project/offload/test/offloading/ompx_saxpy_mixed.c
Nick Sarnie 78ff5b55fd
[offload][lit] Enable/disable tests on Level Zero when using DeviceRTL (#182128)
Since we can now build the DeviceRTL with SPIR-V, redo the
`XFAIL/UNSUPPORTED` specifications for the tests we see passing/failing
on the Level Zero backend with the DeviceRTL being used.

The tests marked `UNSUPPORTED` hang or sporadically fail and those are
tracked in https://github.com/llvm/llvm-project/issues/182119.

This change will allow us to enable CI testing with the DeviceRTL.

Here are the full test results with this change applied, running only
the `spirv64-intel` `check-offload` tests:

```
Total Discovered Tests: 453
  Unsupported      : 206 (45.47%)
  Passed           : 141 (31.13%)
  Expectedly Failed: 106 (23.40%)
```

31% is not a bad start.

---------

Signed-off-by: Nick Sarnie <nick.sarnie@intel.com>
2026-02-18 21:53:19 +00:00

57 lines
1.2 KiB
C

// RUN: %libomptarget-compileopt-run-and-check-generic
//
// REQUIRES: gpu
// https://github.com/llvm/llvm-project/issues/182119
// UNSUPPORTED: intelgpu
#include <math.h>
#include <omp.h>
#include <ompx.h>
#include <stdio.h>
#include <stdlib.h>
int main(int argc, char **argv) {
int N = 1 << 29;
if (argc > 1)
N = atoi(argv[1]);
float a = 2.f;
float *X = (float *)malloc(sizeof(*X) * N);
float *Y = (float *)malloc(sizeof(*X) * N);
for (int i = 0; i < N; i++) {
X[i] = 1.0f;
Y[i] = 2.0f;
}
int TL = 256;
int NT = (N + TL - 1) / TL;
#pragma omp target data map(to : X [0:N]) map(Y [0:N])
#pragma omp target teams num_teams(NT) thread_limit(TL)
{
#pragma omp parallel
{
int tid = ompx_thread_id_x();
int bid = ompx_block_id_x();
int tdim = ompx_block_dim_x();
int gid = tid + bid * tdim;
if (gid < N)
Y[gid] = a * X[gid] + Y[gid];
}
}
float maxError = 0.0f;
for (int i = 0; i < N; i++) {
maxError = fmax(maxError, fabs(Y[i] - 4.0f));
if (maxError) {
printf("%i %f %f\n", i, maxError, Y[i]);
break;
}
}
// CHECK: Max error: 0.00
printf("Max error: %f\n", maxError);
return 0;
}