File: ompx_bare_ballot_sync.c

package info (click to toggle)
swiftlang 6.2.3-1
  • links: PTS, VCS
  • area: main
  • in suites: forky, sid
  • size: 2,856,264 kB
  • sloc: cpp: 9,995,718; ansic: 2,234,019; asm: 1,092,167; python: 313,940; objc: 82,726; f90: 80,126; lisp: 38,373; pascal: 25,580; sh: 20,378; ml: 5,058; perl: 4,751; makefile: 4,725; awk: 3,535; javascript: 3,018; xml: 918; fortran: 664; cs: 573; ruby: 396
file content (38 lines) | stat: -rw-r--r-- 840 bytes parent folder | download | duplicates (4) 
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
 
// RUN: %libomptarget-compilexx-run-and-check-generic
//
// REQUIRES: gpu

#include <assert.h>
#include <ompx.h>
#include <stdint.h>
#include <stdio.h>
#include <stdlib.h>

int main(int argc, char *argv[]) {
  const int num_blocks = 1;
  const int block_size = 256;
  const int N = num_blocks * block_size;
  int *res = (int *)malloc(N * sizeof(int));

#pragma omp target teams ompx_bare num_teams(num_blocks) thread_limit(block_size) \
        map(from: res[0:N])
  {
    int tid = ompx_thread_id_x();
    uint64_t mask = ompx_ballot_sync(~0LU, tid & 0x1);
#if defined __AMDGCN_WAVEFRONT_SIZE && __AMDGCN_WAVEFRONT_SIZE == 64
    res[tid] = mask == 0xaaaaaaaaaaaaaaaa;
#else
    res[tid] = mask == 0xaaaaaaaa;
#endif
  }

  for (int i = 0; i < N; ++i)
    assert(res[i]);

  // CHECK: PASS
  printf("PASS\n");

  free(res);

  return 0;
}