1 // RUN: %libomptarget-compilexx-generic 2 // RUN: env LIBOMPTARGET_INFO=63 %libomptarget-run-generic 2>&1 | %fcheck-generic 3 // REQUIRES: gpu 4 5 #include <ompx.h> 6 7 #include <cassert> 8 #include <vector> 9 10 // CHECK: "PluginInterface" device 0 info: Launching kernel __omp_offloading_{{.*}} with [2,4,6] blocks and [32,4,2] threads in SPMD mode 11 12 int main(int argc, char *argv[]) { 13 int bs[3] = {32u, 4u, 2u}; 14 int gs[3] = {2u, 4u, 6u}; 15 int n = bs[0] * bs[1] * bs[2] * gs[0] * gs[1] * gs[2]; 16 std::vector<int> x_buf(n); 17 std::vector<int> y_buf(n); 18 std::vector<int> z_buf(n); 19 20 auto x = x_buf.data(); 21 auto y = y_buf.data(); 22 auto z = z_buf.data(); 23 for (int i = 0; i < n; ++i) { 24 x[i] = i; 25 y[i] = i + 1; 26 } 27 28 #pragma omp target teams ompx_bare num_teams(gs[0], gs[1], gs[2]) \ 29 thread_limit(bs[0], bs[1], bs[2]) map(to : x[ : n], y[ : n]) \ 30 map(from : z[ : n]) 31 { 32 int tid_x = ompx_thread_id_x(); 33 int tid_y = ompx_thread_id_y(); 34 int tid_z = ompx_thread_id_z(); 35 int gid_x = ompx_block_id_x(); 36 int gid_y = ompx_block_id_y(); 37 int gid_z = ompx_block_id_z(); 38 int bs_x = ompx_block_dim_x(); 39 int bs_y = ompx_block_dim_y(); 40 int bs_z = ompx_block_dim_z(); 41 int bs = bs_x * bs_y * bs_z; 42 int gs_x = ompx_grid_dim_x(); 43 int gs_y = ompx_grid_dim_y(); 44 int gid = gid_z * gs_y * gs_x + gid_y * gs_x + gid_x; 45 int tid = tid_z * bs_x * bs_y + tid_y * bs_x + tid_x; 46 int i = gid * bs + tid; 47 z[i] = x[i] + y[i]; 48 } 49 50 for (int i = 0; i < n; ++i) { 51 if (z[i] != (2 * i + 1)) 52 return 1; 53 } 54 55 return 0; 56 } 57