1330d8983SJohannes Doerfert // RUN: %libomptarget-compile-generic 2330d8983SJohannes Doerfert // RUN: env LIBOMPTARGET_INFO=63 %libomptarget-run-generic 2>&1 | \ 3330d8983SJohannes Doerfert // RUN: %fcheck-generic 4330d8983SJohannes Doerfert // 588234488SEthan Luis McDonough // REQUIRES: gpu 6330d8983SJohannes Doerfert 7330d8983SJohannes Doerfert #include <assert.h> 8330d8983SJohannes Doerfert #include <ompx.h> 9330d8983SJohannes Doerfert #include <stdio.h> 10330d8983SJohannes Doerfert #include <stdlib.h> 11330d8983SJohannes Doerfert 12330d8983SJohannes Doerfert int main(int argc, char *argv[]) { 13330d8983SJohannes Doerfert const int num_blocks = 64; 14330d8983SJohannes Doerfert const int block_size = 64; 15330d8983SJohannes Doerfert const int N = num_blocks * block_size; 16330d8983SJohannes Doerfert int *data = (int *)malloc(N * sizeof(int)); 17330d8983SJohannes Doerfert 18*92376c3fSShilei Tian // CHECK: "PluginInterface" device 0 info: Launching kernel __omp_offloading_{{.*}} with [64,1,1] blocks and [64,1,1] threads in SPMD mode 19330d8983SJohannes Doerfert 20330d8983SJohannes Doerfert #pragma omp target teams ompx_bare num_teams(num_blocks) thread_limit(block_size) map(from: data[0:N]) 21330d8983SJohannes Doerfert { 22330d8983SJohannes Doerfert int bid = ompx_block_id_x(); 23330d8983SJohannes Doerfert int bdim = ompx_block_dim_x(); 24330d8983SJohannes Doerfert int tid = ompx_thread_id_x(); 25330d8983SJohannes Doerfert int idx = bid * bdim + tid; 26330d8983SJohannes Doerfert data[idx] = idx; 27330d8983SJohannes Doerfert } 28330d8983SJohannes Doerfert 29330d8983SJohannes Doerfert for (int i = 0; i < N; ++i) 30330d8983SJohannes Doerfert assert(data[i] == i); 31330d8983SJohannes Doerfert 32330d8983SJohannes Doerfert // CHECK: PASS 33330d8983SJohannes Doerfert printf("PASS\n"); 34330d8983SJohannes Doerfert 35330d8983SJohannes Doerfert return 0; 36330d8983SJohannes Doerfert } 37