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
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
93
94
95
96
97
98
99
100
101
102
103
104
105
106
107
108
109
110
111
112
113
114
115
116
117
118
119
120
121
122
123
124
125
126
127
128
129
130
131
132
133
134
135
136
137
138
139
140
141
142
143
144
145
146
147
148
149
150
151
152
153
154
155
156
157
158
159
|
/* { dg-require-effective-target openacc_cuda } */
/* { dg-require-effective-target openacc_cudart } */
/* { dg-additional-options "-lcuda -lcudart" } */
/* NOTE: This file is also included by libgomp.c-c++-common/interop-cuda-libonly.c
to test the fallback version, which defines USE_CUDA_FALLBACK_HEADER. */
/* Minimal check whether CUDA works - by checking whether the API routines
seem to work. This includes a fallback if the header is not
available. */
#include <assert.h>
#include <omp.h>
#if __has_include(<cuda.h>) && __has_include(<cudaTypedefs.h>) && __has_include(<cuda_runtime.h>) && !defined(USE_CUDA_FALLBACK_HEADER)
#include <cuda.h>
#include <cudaTypedefs.h>
#include <cuda_runtime.h>
#else
/* Add a poor man's fallback declaration. */
#if USE_CUDA_FALLBACK_HEADER
// Don't warn.
#elif !__has_include(<cuda.h>)
#warning "Using GCC's cuda.h as fallback for cuda.h"
#elif !__has_include(<cudaTypedefs.h>)
#warning "Using GCC's cuda.h as fallback for cudaTypedefs.h"
#else
#warning "Using GCC's cuda.h as fallback for cuda_runtime.h"
#endif
#include "../../../include/cuda/cuda.h"
typedef int cudaError_t;
typedef CUstream cudaStream_t;
enum {
cudaSuccess = 0
};
enum cudaDeviceAttr {
cudaDevAttrClockRate = 13,
cudaDevAttrMaxGridDimX = 5
};
cudaError_t cudaDeviceGetAttribute (int *, enum cudaDeviceAttr, int);
cudaError_t cudaStreamQuery(cudaStream_t);
CUresult cuCtxGetApiVersion(CUcontext, unsigned int *);
CUresult cuStreamGetCtx (CUstream, CUcontext *);
#endif
int
main ()
{
int ivar;
unsigned uvar;
omp_interop_rc_t res;
omp_interop_t obj_cuda = omp_interop_none;
omp_interop_t obj_cuda_driver = omp_interop_none;
cudaError_t cuda_err;
CUresult cu_err;
#pragma omp interop init(target, targetsync, prefer_type("cuda") : obj_cuda) \
init(target, targetsync, prefer_type("cuda_driver") : obj_cuda_driver) \
omp_interop_fr_t fr = (omp_interop_fr_t) omp_get_interop_int (obj_cuda, omp_ipr_fr_id, &res);
assert (res == omp_irc_success);
assert (fr == omp_ifr_cuda);
fr = (omp_interop_fr_t) omp_get_interop_int (obj_cuda_driver, omp_ipr_fr_id, &res);
assert (res == omp_irc_success);
assert (fr == omp_ifr_cuda_driver);
ivar = (int) omp_get_interop_int (obj_cuda, omp_ipr_vendor, &res);
assert (res == omp_irc_success);
assert (ivar == 11);
ivar = (int) omp_get_interop_int (obj_cuda_driver, omp_ipr_vendor, &res);
assert (res == omp_irc_success);
assert (ivar == 11);
/* Check whether the omp_ipr_device -> cudaDevice_t yields a valid device. */
CUdevice cu_dev = (int) omp_get_interop_int (obj_cuda_driver, omp_ipr_device, &res);
assert (res == omp_irc_success);
/* Assume a clock size is available and > 1 GHz; value is in kHz. */
cu_err = cuDeviceGetAttribute (&ivar, cudaDevAttrClockRate, cu_dev);
assert (cu_err == CUDA_SUCCESS);
assert (ivar > 1000000 /* kHz */);
/* Assume that the MaxGridDimX is available and > 1024. */
cu_err = cuDeviceGetAttribute (&ivar, cudaDevAttrMaxGridDimX, cu_dev);
assert (cu_err == CUDA_SUCCESS);
assert (ivar > 1024);
int cuda_dev = (int) omp_get_interop_int (obj_cuda, omp_ipr_device, &res);
assert (res == omp_irc_success);
assert (cuda_dev == (CUdevice) cu_dev); // Assume they are the same ...
/* Assume a clock size is available and > 1 GHz; value is in kHz. */
cuda_err = cudaDeviceGetAttribute (&ivar, cudaDevAttrClockRate, cuda_dev);
assert (cuda_err == cudaSuccess);
assert (ivar > 1000000 /* kHz */);
/* Assume that the MaxGridDimX is available and > 1024. */
cuda_err = cudaDeviceGetAttribute (&ivar, cudaDevAttrMaxGridDimX, cuda_dev);
assert (cuda_err == cudaSuccess);
assert (ivar > 1024);
/* Check whether the omp_ipr_device_context -> CUcontext yields a context. */
CUcontext cu_ctx = (CUcontext) omp_get_interop_ptr (obj_cuda_driver, omp_ipr_device_context, &res);
assert (res == omp_irc_success);
/* Assume API Version > 0 for Nvidia, cudaErrorNotSupported for AMD. */
uvar = 99;
cu_err = cuCtxGetApiVersion (cu_ctx, &uvar);
assert (cu_err == CUDA_SUCCESS);
assert (uvar > 0);
/* Check whether the omp_ipr_targetsync -> cudaStream_t yields a stream. */
cudaStream_t cuda_sm = (cudaStream_t) omp_get_interop_ptr (obj_cuda, omp_ipr_targetsync, &res);
assert (res == omp_irc_success);
CUstream cu_sm = (cudaStream_t) omp_get_interop_ptr (obj_cuda_driver, omp_ipr_targetsync, &res);
assert (res == omp_irc_success);
assert ((void*) cu_sm != (void*) cuda_sm); // Type compatible but should have created two streams
int dev_stream = 99;
#if CUDA_VERSION >= 12080
cuda_err = cudaStreamGetDevice (cuda_sm, &dev_stream);
assert (cuda_err == cudaSuccess);
#else
cu_err = cuStreamGetCtx (cu_sm, &cu_ctx) != CUDA_SUCCESS;
if (cu_err == CUDA_SUCCESS)
cuda_err = cuCtxPushCurrent (cu_ctx) != CUDA_SUCCESS;
if (cu_err == CUDA_SUCCESS)
cuda_err = cuCtxGetDevice (&dev_stream) != CUDA_SUCCESS;
if (cu_err == CUDA_SUCCESS)
cu_err = cuCtxPopCurrent (&cu_ctx) != CUDA_SUCCESS;
assert (cu_err == CUDA_SUCCESS);
#endif
assert (dev_stream == cuda_dev);
/* All jobs should have been completed (as there were none none) */
cuda_err = cudaStreamQuery (cuda_sm);
assert (cuda_err == cudaSuccess);
cu_err = cuStreamQuery (cu_sm);
assert (cu_err == CUDA_SUCCESS);
#pragma omp interop destroy(obj_cuda, obj_cuda_driver)
}
|
