7039515e0c
When creating an outline function for device code we're not setting the right calling convention when the target is SPIRV. This results in the calls to the function to be removed by the InstCombine pass as it thinks is not callable.
61 lines
1.2 KiB
C
61 lines
1.2 KiB
C
// RUN: %libomptarget-compileopt-run-and-check-generic
|
|
//
|
|
// REQUIRES: gpu
|
|
|
|
#include <omp.h>
|
|
#include <ompx.h>
|
|
#include <stdio.h>
|
|
#include <stdlib.h>
|
|
#include <string.h>
|
|
|
|
struct info {
|
|
int tid, bid, tdim;
|
|
};
|
|
|
|
int main(int argc, char **argv) {
|
|
int N = 1 << 20;
|
|
if (argc > 1)
|
|
N = atoi(argv[1]);
|
|
|
|
struct info *X = (struct info *)malloc(sizeof(*X) * N);
|
|
memset(X, '0', sizeof(*X) * N);
|
|
|
|
int TL = 256;
|
|
int NT = (N + TL - 1) / TL;
|
|
|
|
#pragma omp target data map(tofrom : X [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) {
|
|
X[gid].tid = tid;
|
|
X[gid].bid = bid;
|
|
X[gid].tdim = tdim;
|
|
};
|
|
}
|
|
}
|
|
|
|
int tid = 0, bid = 0, tdim = 256;
|
|
for (int i = 0; i < N; i++) {
|
|
if (X[i].tid != tid || X[i].bid != bid || X[i].tdim != tdim) {
|
|
printf("%i: %i vs %i, %i vs %i, %i vs %i\n", i, X[i].tid, tid, X[i].bid,
|
|
bid, X[i].tdim, tdim);
|
|
return 1;
|
|
}
|
|
tid++;
|
|
if (tid == tdim) {
|
|
tid = 0;
|
|
bid++;
|
|
}
|
|
}
|
|
|
|
// CHECK: OK
|
|
printf("OK");
|
|
return 0;
|
|
}
|