18c247cc0b
Per my reading of the OpenACC specification (and as supported by secondary documentation, such as code examples, or presentations), it's valid to call "acc_get_cuda_stream"/"acc_set_cuda_stream" also with "acc_async_sync", "acc_async_noval" arguments, not just with the nonnegative values as currently implemented. libgomp/ PR libgomp/88370 * libgomp.texi (acc_get_current_cuda_context, acc_get_cuda_stream) (acc_set_cuda_stream): Clarify. * oacc-cuda.c (acc_get_cuda_stream, acc_set_cuda_stream): Use "async_valid_p". * plugin/plugin-nvptx.c (nvptx_set_cuda_stream): Refuse "async == acc_async_sync". * testsuite/libgomp.oacc-c-c++-common/acc_set_cuda_stream-1.c: New file. * testsuite/libgomp.oacc-c-c++-common/async_queue-1.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/lib-84.c: Update. * testsuite/libgomp.oacc-c-c++-common/lib-85.c: Likewise. From-SVN: r267147
43 lines
1.1 KiB
C
43 lines
1.1 KiB
C
/* Verify expected nvptx plugin behavior for "acc_set_cuda_stream" for
|
|
"acc_async_sync". */
|
|
|
|
/* { dg-do run { target openacc_nvidia_accel_selected } } */
|
|
/* { dg-set-target-env-var GOMP_DEBUG "1" } */
|
|
|
|
#undef NDEBUG
|
|
#include <assert.h>
|
|
#include <openacc.h>
|
|
|
|
int main(void)
|
|
{
|
|
int async = 42;
|
|
|
|
/* Initialize. */
|
|
#pragma acc parallel async(acc_async_sync)
|
|
;
|
|
#pragma acc parallel async(async)
|
|
;
|
|
#pragma acc wait
|
|
|
|
void *cuda_stream_sync = acc_get_cuda_stream (acc_async_sync);
|
|
assert (cuda_stream_sync == NULL);
|
|
void *cuda_stream_async = acc_get_cuda_stream (async);
|
|
assert (cuda_stream_async != NULL);
|
|
int ret = acc_set_cuda_stream (acc_async_sync, cuda_stream_async);
|
|
assert (ret == 0);
|
|
void *cuda_stream_sync_ = acc_get_cuda_stream (acc_async_sync);
|
|
assert (cuda_stream_sync_ == cuda_stream_sync);
|
|
void *cuda_stream_async_ = acc_get_cuda_stream (async);
|
|
assert (cuda_stream_async_ == cuda_stream_async);
|
|
|
|
#pragma acc parallel async(acc_async_sync)
|
|
;
|
|
#pragma acc parallel async(async)
|
|
;
|
|
#pragma acc wait
|
|
|
|
return 0;
|
|
}
|
|
|
|
/* { dg-output "Refusing request to set CUDA stream associated with \"acc_async_sync\"" } */
|