26b777444b
We finally got our buildbot added (to staging, at least) so we want to start running L0 tests in CI. We need `check-offload` to pass though, so XFAIL everything failing. There's a couple `UNSUPPORTED` as well, those are for sporadic fails. Also make set the `gpu` and `intelgpu` LIT variables when testing the `spirv64-intel` triple. We have no DeviceRTL yet so basically everything fails, but we manage to get ``` Total Discovered Tests: 432 Unsupported : 169 (39.12%) Passed : 67 (15.51%) Expectedly Failed: 196 (45.37%) ``` We still don't build the level zero plugin by default and these tests don't run unless the plugin was built, so this has no effect on most builds. --------- Signed-off-by: Nick Sarnie <nick.sarnie@intel.com>
60 lines
1.6 KiB
C++
60 lines
1.6 KiB
C++
// clang-format off
|
|
// RUN: %libomptarget-compilexx-generic
|
|
// RUN: env LIBOMPTARGET_INFO=63 %libomptarget-run-generic 2>&1 | %fcheck-generic
|
|
// REQUIRES: gpu
|
|
// XFAIL: intelgpu
|
|
// clang-format on
|
|
|
|
#include <ompx.h>
|
|
|
|
#include <cassert>
|
|
#include <vector>
|
|
|
|
// CHECK: "PluginInterface" device 0 info: Launching kernel __omp_offloading_{{.*}} with [2,4,6] blocks and [32,4,2] threads in BARE mode
|
|
|
|
int main(int argc, char *argv[]) {
|
|
int bs[3] = {32u, 4u, 2u};
|
|
int gs[3] = {2u, 4u, 6u};
|
|
int n = bs[0] * bs[1] * bs[2] * gs[0] * gs[1] * gs[2];
|
|
std::vector<int> x_buf(n);
|
|
std::vector<int> y_buf(n);
|
|
std::vector<int> z_buf(n);
|
|
|
|
auto x = x_buf.data();
|
|
auto y = y_buf.data();
|
|
auto z = z_buf.data();
|
|
for (int i = 0; i < n; ++i) {
|
|
x[i] = i;
|
|
y[i] = i + 1;
|
|
}
|
|
|
|
#pragma omp target teams ompx_bare num_teams(gs[0], gs[1], gs[2]) \
|
|
thread_limit(bs[0], bs[1], bs[2]) map(to : x[ : n], y[ : n]) \
|
|
map(from : z[ : n])
|
|
{
|
|
int tid_x = ompx_thread_id_x();
|
|
int tid_y = ompx_thread_id_y();
|
|
int tid_z = ompx_thread_id_z();
|
|
int gid_x = ompx_block_id_x();
|
|
int gid_y = ompx_block_id_y();
|
|
int gid_z = ompx_block_id_z();
|
|
int bs_x = ompx_block_dim_x();
|
|
int bs_y = ompx_block_dim_y();
|
|
int bs_z = ompx_block_dim_z();
|
|
int bs = bs_x * bs_y * bs_z;
|
|
int gs_x = ompx_grid_dim_x();
|
|
int gs_y = ompx_grid_dim_y();
|
|
int gid = gid_z * gs_y * gs_x + gid_y * gs_x + gid_x;
|
|
int tid = tid_z * bs_x * bs_y + tid_y * bs_x + tid_x;
|
|
int i = gid * bs + tid;
|
|
z[i] = x[i] + y[i];
|
|
}
|
|
|
|
for (int i = 0; i < n; ++i) {
|
|
if (z[i] != (2 * i + 1))
|
|
return 1;
|
|
}
|
|
|
|
return 0;
|
|
}
|