Some additions of redundant "present" clauses dropped.
2018-10-22 Cesar Philippidis <cesar@codesourcery.com>
gcc/testsuite/
* g++.dg/goacc/loop-1.c: New test.
* g++.dg/goacc/loop-2.c: New test.
* g++.dg/goacc/loop-3.c: New test.
2018-10-22 James Norris <jnorris@codesourcery.com>
Cesar Philippidis <cesar@codesourcery.com>
Tom de Vries <tom@codesourcery.com>
libgomp/
* testsuite/libgomp.oacc-fortran/data-3.f90: Update parallel
regions to denote variables copyied in via acc enter data as
present.
* testsuite/libgomp.oacc-c-c++-common/subr.h: Reimplement.
* testsuite/libgomp.oacc-c-c++-common/subr.ptx: Regenerated PTX.
* testsuite/libgomp.oacc-c-c++-common/timer.h: Removed.
* testsuite/libgomp.oacc-c-c++-common/lib-69.c: Change async checks.
* testsuite/libgomp.oacc-c-c++-common/lib-70.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/lib-72.c: Rework kernel i/f and
change async checks.
* testsuite/libgomp.oacc-c-c++-common/lib-73.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/lib-74.c: Rework kernel i/f and
timing checks.
* testsuite/libgomp.oacc-c-c++-common/lib-75.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/lib-76.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/lib-78.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/lib-79.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/lib-81.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/lib-82.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/lib-93.c: New test.
+2018-10-22 Cesar Philippidis <cesar@codesourcery.com>
+
+ * g++.dg/goacc/loop-1.c: New test.
+ * g++.dg/goacc/loop-2.c: New test.
+ * g++.dg/goacc/loop-3.c: New test.
+
2018-12-13 Cesar Philippidis <cesar@codesourcery.com>
Nathan Sidwell <nathan@acm.org>
Julian Brown <julian@codesourcery.com>
--- /dev/null
+void
+f (int i, float j, int k)
+{
+#pragma acc parallel num_gangs (i) num_workers (i) vector_length (i)
+#pragma acc loop gang
+ for (i = 0; i < 20; ++i)
+ ;
+
+#pragma acc parallel num_gangs (j) /* { dg-error "'num_gangs' expression must be integral" } */
+#pragma acc loop gang
+ for (i = 0; i < 20; ++i)
+ ;
+
+#pragma acc parallel num_workers (j) /* { dg-error "'num_workers' expression must be integral" } */
+#pragma acc loop gang
+ for (i = 0; i < 20; ++i)
+ ;
+
+#pragma acc parallel vector_length (j) /* { dg-error "'vector_length' expression must be integral" } */
+#pragma acc loop gang
+ for (i = 0; i < 20; ++i)
+ ;
+}
--- /dev/null
+void
+f (int i, int j, int k)
+{
+#pragma acc kernels
+#pragma acc loop gang
+ for (i = 0; i < 20; ++i)
+ ;
+
+#pragma acc kernels
+#pragma acc loop gang (num: 10)
+ for (i = 0; i < 20; ++i)
+ ;
+
+#pragma acc kernels
+#pragma acc loop gang (static: 10)
+ for (i = 0; i < 20; ++i)
+ ;
+
+#pragma acc kernels
+#pragma acc loop gang (static: 5, num: 10)
+ for (i = 0; i < 20; ++i)
+ ;
+
+
+#pragma acc kernels
+#pragma acc loop gang (static: 5, num: 10, *) /* { dg-error "duplicate operand to clause" } */
+ for (i = 0; i < 20; ++i)
+ ;
+
+#pragma acc kernels
+#pragma acc loop gang (static: 5, num: 10, static: *) /* { dg-error "duplicate 'num' argument" } */
+ for (i = 0; i < 20; ++i)
+ ;
+
+#pragma acc kernels
+#pragma acc loop worker (static: 234) /* { dg-error "expected 'num' before" } */
+ for (i = 0; i < 20; ++i)
+ ;
+
+#pragma acc kernels
+#pragma acc loop worker (num: 234)
+ for (i = 0; i < 20; ++i)
+ ;
+
+#pragma acc kernels
+#pragma acc loop worker (num: 234, num: 12) /* { dg-error "duplicate operand to clause" } */
+ for (i = 0; i < 20; ++i)
+ ;
+
+#pragma acc kernels
+#pragma acc loop vector /* { dg-error "gang, worker and vector must occur in this order in a loop nest" } */
+ for (i = 0; i < 20; ++i)
+#pragma acc loop worker
+ for (j = 0; j < 25; ++j)
+ ;
+
+#pragma acc kernels
+#pragma acc loop worker (length: 20) /* { dg-error "expected 'num' before 'length'" } */
+ for (i = 0; i < 20; ++i)
+#pragma acc loop vector (length: 10)
+ for (j = 0; j < 25; ++j)
+ ;
+
+#pragma acc kernels
+#pragma acc loop worker
+ for (i = 0; i < 20; ++i)
+#pragma acc loop vector
+ for (j = 0; j < 25; ++j)
+ ;
+}
--- /dev/null
+void
+f (int i, int j, int k)
+{
+#pragma acc kernels num_gangs (10) /* { dg-error "'num_gangs' is not valid" } */
+#pragma acc loop gang
+ for (i = 0; i < 20; ++i)
+ ;
+
+#pragma acc kernels num_workers (10) /* { dg-error "'num_workers' is not valid" } */
+#pragma acc loop worker
+ for (i = 0; i < 20; ++i)
+ ;
+
+#pragma acc kernels vector_length (10) /* { dg-error "'vector_length' is not valid" } */
+#pragma acc loop vector
+ for (i = 0; i < 20; ++i)
+ ;
+
+#pragma acc parallel num_gangs (10) num_workers (20) vector_length (32)
+#pragma acc loop gang
+ for (i = 0; i < 20; ++i)
+ ;
+
+#pragma acc parallel num_gangs (i) num_workers (j) vector_length (k)
+#pragma acc loop gang
+ for (i = 0; i < 20; ++i)
+ ;
+
+#pragma acc parallel num_gangs (10, i) /* { dg-error "expected '\\)' before ',' token" } */
+#pragma acc loop gang
+ for (i = 0; i < 20; ++i)
+ ;
+
+#pragma acc parallel num_workers (10, i) /* { dg-error "expected '\\)' before ',' token" } */
+#pragma acc loop gang
+ for (i = 0; i < 20; ++i)
+ ;
+
+#pragma acc parallel vector_length (10, i) /* { dg-error "expected '\\)' before ',' token" } */
+#pragma acc loop gang
+ for (i = 0; i < 20; ++i)
+ ;
+}
+2018-10-22 James Norris <jnorris@codesourcery.com>
+ Cesar Philippidis <cesar@codesourcery.com>
+ Tom de Vries <tom@codesourcery.com>
+
+ * testsuite/libgomp.oacc-fortran/data-3.f90: Update parallel
+ regions to denote variables copyied in via acc enter data as
+ present.
+ * testsuite/libgomp.oacc-c-c++-common/subr.h: Reimplement.
+ * testsuite/libgomp.oacc-c-c++-common/subr.ptx: Regenerated PTX.
+ * testsuite/libgomp.oacc-c-c++-common/timer.h: Removed.
+ * testsuite/libgomp.oacc-c-c++-common/lib-69.c: Change async checks.
+ * testsuite/libgomp.oacc-c-c++-common/lib-70.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/lib-72.c: Rework kernel i/f and
+ change async checks.
+ * testsuite/libgomp.oacc-c-c++-common/lib-73.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/lib-74.c: Rework kernel i/f and
+ timing checks.
+ * testsuite/libgomp.oacc-c-c++-common/lib-75.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/lib-76.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/lib-78.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/lib-79.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/lib-81.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/lib-82.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/lib-93.c: New test.
+
2018-12-13 Cesar Philippidis <cesar@codesourcery.com>
Nathan Sidwell <nathan@acm.org>
Julian Brown <julian@codesourcery.com>
int
main (int argc, char **argv)
{
- CUdevice dev;
CUfunction delay;
CUmodule module;
CUresult r;
CUstream stream;
- unsigned long *a, *d_a, dticks;
- int nbytes;
- float dtime;
- void *kargs[2];
- int clkrate;
- int devnum, nprocs;
acc_init (acc_device_nvidia);
- devnum = acc_get_device_num (acc_device_nvidia);
-
- r = cuDeviceGet (&dev, devnum);
- if (r != CUDA_SUCCESS)
- {
- fprintf (stderr, "cuDeviceGet failed: %d\n", r);
- abort ();
- }
-
- r =
- cuDeviceGetAttribute (&nprocs, CU_DEVICE_ATTRIBUTE_MULTIPROCESSOR_COUNT,
- dev);
- if (r != CUDA_SUCCESS)
- {
- fprintf (stderr, "cuDeviceGetAttribute failed: %d\n", r);
- abort ();
- }
-
- r = cuDeviceGetAttribute (&clkrate, CU_DEVICE_ATTRIBUTE_CLOCK_RATE, dev);
- if (r != CUDA_SUCCESS)
- {
- fprintf (stderr, "cuDeviceGetAttribute failed: %d\n", r);
- abort ();
- }
-
- r = cuModuleLoad (&module, "subr.ptx");
+ r = cuModuleLoad (&module, "./subr.ptx");
if (r != CUDA_SUCCESS)
{
fprintf (stderr, "cuModuleLoad failed: %d\n", r);
abort ();
}
- nbytes = nprocs * sizeof (unsigned long);
-
- dtime = 200.0;
-
- dticks = (unsigned long) (dtime * clkrate);
-
- a = (unsigned long *) malloc (nbytes);
- d_a = (unsigned long *) acc_malloc (nbytes);
-
- acc_map_data (a, d_a, nbytes);
-
- kargs[0] = (void *) &d_a;
- kargs[1] = (void *) &dticks;
-
stream = (CUstream) acc_get_cuda_stream (0);
if (stream != NULL)
abort ();
if (!acc_set_cuda_stream (0, stream))
abort ();
- r = cuLaunchKernel (delay, 1, 1, 1, 1, 1, 1, 0, stream, kargs, 0);
+ r = cuLaunchKernel (delay, 1, 1, 1, 1, 1, 1, 0, stream, NULL, 0);
if (r != CUDA_SUCCESS)
{
fprintf (stderr, "cuLaunchKernel failed: %d\n", r);
abort ();
}
- acc_unmap_data (a);
-
- free (a);
- acc_free (d_a);
-
acc_shutdown (acc_device_nvidia);
exit (0);
/* { dg-additional-options "-lcuda" } */
/* { dg-require-effective-target openacc_cuda } */
+#include <sys/time.h>
#include <stdio.h>
#include <stdlib.h>
#include <unistd.h>
int
main (int argc, char **argv)
{
- CUdevice dev;
CUfunction delay;
CUmodule module;
CUresult r;
- const int N = 10;
+ const int N = 3;
int i;
CUstream streams[N];
- unsigned long *a, *d_a, dticks;
- int nbytes;
- float dtime;
- void *kargs[2];
- int clkrate;
- int devnum, nprocs;
+ struct timeval tv1, tv2;
+ time_t diff;
acc_init (acc_device_nvidia);
- devnum = acc_get_device_num (acc_device_nvidia);
-
- r = cuDeviceGet (&dev, devnum);
- if (r != CUDA_SUCCESS)
- {
- fprintf (stderr, "cuDeviceGet failed: %d\n", r);
- abort ();
- }
-
- r =
- cuDeviceGetAttribute (&nprocs, CU_DEVICE_ATTRIBUTE_MULTIPROCESSOR_COUNT,
- dev);
- if (r != CUDA_SUCCESS)
- {
- fprintf (stderr, "cuDeviceGetAttribute failed: %d\n", r);
- abort ();
- }
-
- r = cuDeviceGetAttribute (&clkrate, CU_DEVICE_ATTRIBUTE_CLOCK_RATE, dev);
- if (r != CUDA_SUCCESS)
- {
- fprintf (stderr, "cuDeviceGetAttribute failed: %d\n", r);
- abort ();
- }
-
r = cuModuleLoad (&module, "subr.ptx");
if (r != CUDA_SUCCESS)
{
abort ();
}
- nbytes = nprocs * sizeof (unsigned long);
-
- dtime = 200.0;
-
- dticks = (unsigned long) (dtime * clkrate);
-
- a = (unsigned long *) malloc (nbytes);
- d_a = (unsigned long *) acc_malloc (nbytes);
-
- acc_map_data (a, d_a, nbytes);
-
- kargs[0] = (void *) &d_a;
- kargs[1] = (void *) &dticks;
-
for (i = 0; i < N; i++)
{
streams[i] = (CUstream) acc_get_cuda_stream (i);
abort ();
}
+ gettimeofday (&tv1, NULL);
+
+ r = cuLaunchKernel (delay, 1, 1, 1, 1, 1, 1, 0, streams[0], NULL, 0);
+ if (r != CUDA_SUCCESS)
+ {
+ fprintf (stderr, "cuLaunchKernel failed: %d\n", r);
+ abort ();
+ }
+
+ r = cuCtxSynchronize ();
+ if (r != CUDA_SUCCESS)
+ {
+ fprintf (stderr, "cuCtxLaunch failed: %d\n", r);
+ abort ();
+ }
+
+ gettimeofday (&tv2, NULL);
+
+ diff = tv2.tv_sec - tv1.tv_sec;
+
for (i = 0; i < N; i++)
{
- r = cuLaunchKernel (delay, 1, 1, 1, 1, 1, 1, 0, streams[i], kargs, 0);
+ r = cuLaunchKernel (delay, 1, 1, 1, 1, 1, 1, 0, streams[i], NULL, 0);
if (r != CUDA_SUCCESS)
{
fprintf (stderr, "cuLaunchKernel failed: %d\n", r);
}
}
- sleep ((int) (dtime / 1000.0f) + 1);
+ sleep ((diff + 1) * N);
for (i = 0; i < N; i++)
{
}
}
- acc_unmap_data (a);
-
- free (a);
- acc_free (d_a);
acc_shutdown (acc_device_nvidia);
int
main (int argc, char **argv)
{
- CUdevice dev;
CUfunction delay;
CUmodule module;
CUresult r;
CUstream stream;
- unsigned long *a, *d_a, dticks;
- int nbytes;
- float dtime;
- void *kargs[2];
- int clkrate;
- int devnum, nprocs;
acc_init (acc_device_nvidia);
- devnum = acc_get_device_num (acc_device_nvidia);
-
- r = cuDeviceGet (&dev, devnum);
- if (r != CUDA_SUCCESS)
- {
- fprintf (stderr, "cuDeviceGet failed: %d\n", r);
- abort ();
- }
-
- r =
- cuDeviceGetAttribute (&nprocs, CU_DEVICE_ATTRIBUTE_MULTIPROCESSOR_COUNT,
- dev);
- if (r != CUDA_SUCCESS)
- {
- fprintf (stderr, "cuDeviceGetAttribute failed: %d\n", r);
- abort ();
- }
-
- r = cuDeviceGetAttribute (&clkrate, CU_DEVICE_ATTRIBUTE_CLOCK_RATE, dev);
- if (r != CUDA_SUCCESS)
- {
- fprintf (stderr, "cuDeviceGetAttribute failed: %d\n", r);
- abort ();
- }
-
r = cuModuleLoad (&module, "subr.ptx");
if (r != CUDA_SUCCESS)
{
abort ();
}
- nbytes = nprocs * sizeof (unsigned long);
-
- dtime = 200.0;
-
- dticks = (unsigned long) (dtime * clkrate);
-
- a = (unsigned long *) malloc (nbytes);
- d_a = (unsigned long *) acc_malloc (nbytes);
-
- acc_map_data (a, d_a, nbytes);
-
- kargs[0] = (void *) &d_a;
- kargs[1] = (void *) &dticks;
-
r = cuStreamCreate (&stream, CU_STREAM_DEFAULT);
if (r != CUDA_SUCCESS)
{
if (!acc_set_cuda_stream (0, stream))
abort ();
- r = cuLaunchKernel (delay, 1, 1, 1, 1, 1, 1, 0, stream, kargs, 0);
+ r = cuLaunchKernel (delay, 1, 1, 1, 1, 1, 1, 0, stream, NULL, 0);
if (r != CUDA_SUCCESS)
{
fprintf (stderr, "cuLaunchKernel failed: %d\n", r);
abort ();
}
- sleep ((int) (dtime / 1000.f) + 1);
+ r = cuCtxSynchronize ();
+ if (r != CUDA_SUCCESS)
+ {
+ fprintf (stderr, "cuCtxSynchronize () failed: %d\n", r);
+ abort ();
+ }
if (acc_async_test_all () != 1)
{
abort ();
}
- acc_unmap_data (a);
-
- free (a);
- acc_free (d_a);
-
acc_shutdown (acc_device_nvidia);
exit (0);
/* { dg-additional-options "-lcuda" } */
/* { dg-require-effective-target openacc_cuda } */
+#include <sys/time.h>
#include <stdio.h>
#include <unistd.h>
#include <stdlib.h>
int
main (int argc, char **argv)
{
- CUdevice dev;
CUfunction delay;
CUmodule module;
CUresult r;
- const int N = 10;
+ const int N = 6;
int i;
CUstream streams[N];
- unsigned long *a, *d_a, dticks;
- int nbytes;
- float dtime;
- void *kargs[2];
- int clkrate;
- int devnum, nprocs;
acc_init (acc_device_nvidia);
- devnum = acc_get_device_num (acc_device_nvidia);
-
- r = cuDeviceGet (&dev, devnum);
- if (r != CUDA_SUCCESS)
- {
- fprintf (stderr, "cuDeviceGet failed: %d\n", r);
- abort ();
- }
-
- r =
- cuDeviceGetAttribute (&nprocs, CU_DEVICE_ATTRIBUTE_MULTIPROCESSOR_COUNT,
- dev);
- if (r != CUDA_SUCCESS)
- {
- fprintf (stderr, "cuDeviceGetAttribute failed: %d\n", r);
- abort ();
- }
-
- r = cuDeviceGetAttribute (&clkrate, CU_DEVICE_ATTRIBUTE_CLOCK_RATE, dev);
- if (r != CUDA_SUCCESS)
- {
- fprintf (stderr, "cuDeviceGetAttribute failed: %d\n", r);
- abort ();
- }
-
r = cuModuleLoad (&module, "subr.ptx");
if (r != CUDA_SUCCESS)
{
abort ();
}
- nbytes = nprocs * sizeof (unsigned long);
-
- dtime = 200.0;
-
- dticks = (unsigned long) (dtime * clkrate);
-
- a = (unsigned long *) malloc (nbytes);
- d_a = (unsigned long *) acc_malloc (nbytes);
-
- acc_map_data (a, d_a, nbytes);
-
- kargs[0] = (void *) &d_a;
- kargs[1] = (void *) &dticks;
-
for (i = 0; i < N; i++)
{
streams[i] = (CUstream) acc_get_cuda_stream (i);
for (i = 0; i < N; i++)
{
- r = cuLaunchKernel (delay, 1, 1, 1, 1, 1, 1, 0, streams[i], kargs, 0);
+ r = cuLaunchKernel (delay, 1, 1, 1, 1, 1, 1, 0, streams[i], NULL, 0);
if (r != CUDA_SUCCESS)
{
fprintf (stderr, "cuLaunchKernel failed: %d\n", r);
abort ();
}
-
}
if (acc_async_test_all () != 0)
abort ();
}
- sleep ((int) (dtime / 1000.0f) + 1);
+ r = cuCtxSynchronize ();
+ if (r != CUDA_SUCCESS)
+ {
+ fprintf (stderr, "cuCtxSynchronize failed: %d\n", r);
+ abort ();
+ }
if (acc_async_test_all () != 1)
{
abort ();
}
- acc_unmap_data (a);
-
- free (a);
- acc_free (d_a);
-
acc_shutdown (acc_device_nvidia);
exit (0);
#include <stdlib.h>
#include <openacc.h>
#include <cuda.h>
-#include "timer.h"
+#include <sys/time.h>
int
main (int argc, char **argv)
{
- CUdevice dev;
CUfunction delay;
CUmodule module;
CUresult r;
CUstream stream;
- unsigned long *a, *d_a, dticks;
- int nbytes;
- float atime, dtime;
- void *kargs[2];
- int clkrate;
- int devnum, nprocs;
+ struct timeval tv1, tv2;
+ time_t t1, t2;
acc_init (acc_device_nvidia);
- devnum = acc_get_device_num (acc_device_nvidia);
-
- r = cuDeviceGet (&dev, devnum);
+ r = cuModuleLoad (&module, "subr.ptx");
if (r != CUDA_SUCCESS)
{
- fprintf (stderr, "cuDeviceGet failed: %d\n", r);
+ fprintf (stderr, "cuModuleLoad failed: %d\n", r);
abort ();
}
- r =
- cuDeviceGetAttribute (&nprocs, CU_DEVICE_ATTRIBUTE_MULTIPROCESSOR_COUNT,
- dev);
+ r = cuModuleGetFunction (&delay, module, "delay");
if (r != CUDA_SUCCESS)
{
- fprintf (stderr, "cuDeviceGetAttribute failed: %d\n", r);
+ fprintf (stderr, "cuModuleGetFunction failed: %d\n", r);
abort ();
}
- r = cuDeviceGetAttribute (&clkrate, CU_DEVICE_ATTRIBUTE_CLOCK_RATE, dev);
- if (r != CUDA_SUCCESS)
- {
- fprintf (stderr, "cuDeviceGetAttribute failed: %d\n", r);
- abort ();
- }
+ gettimeofday (&tv1, NULL);
- r = cuModuleLoad (&module, "subr.ptx");
+ r = cuLaunchKernel (delay, 1, 1, 1, 1, 1, 1, 0, NULL, NULL, 0);
if (r != CUDA_SUCCESS)
{
- fprintf (stderr, "cuModuleLoad failed: %d\n", r);
+ fprintf (stderr, "cuLaunchKernel failed: %d\n", r);
abort ();
}
- r = cuModuleGetFunction (&delay, module, "delay");
+ r = cuCtxSynchronize ();
if (r != CUDA_SUCCESS)
{
- fprintf (stderr, "cuModuleGetFunction failed: %d\n", r);
+ fprintf (stderr, "cuCtxSynchronize failed: %d\n", r);
abort ();
}
- nbytes = nprocs * sizeof (unsigned long);
-
- dtime = 200.0;
-
- dticks = (unsigned long) (dtime * clkrate);
-
- a = (unsigned long *) malloc (nbytes);
- d_a = (unsigned long *) acc_malloc (nbytes);
-
- acc_map_data (a, d_a, nbytes);
+ gettimeofday (&tv2, NULL);
- kargs[0] = (void *) &d_a;
- kargs[1] = (void *) &dticks;
+ t1 = ((tv2.tv_sec - tv1.tv_sec) * 1000000) + (tv2.tv_usec - tv1.tv_usec);
stream = (CUstream) acc_get_cuda_stream (0);
if (stream != NULL)
if (!acc_set_cuda_stream (0, stream))
abort ();
- init_timers (1);
+ gettimeofday (&tv1, NULL);
- start_timer (0);
-
- r = cuLaunchKernel (delay, 1, 1, 1, 1, 1, 1, 0, stream, kargs, 0);
+ r = cuLaunchKernel (delay, 1, 1, 1, 1, 1, 1, 0, stream, NULL, 0);
if (r != CUDA_SUCCESS)
{
fprintf (stderr, "cuLaunchKernel failed: %d\n", r);
}
acc_wait (0);
- /* Test unseen async-argument. */
- acc_wait (1);
- atime = stop_timer (0);
+ gettimeofday (&tv2, NULL);
+
+ t2 = ((tv2.tv_sec - tv1.tv_sec) * 1000000) + (tv2.tv_usec - tv1.tv_usec);
- if (atime < dtime)
+ if (((abs (t2 - t1) / t1) * 100.0) > 1.0)
{
- fprintf (stderr, "actual time < delay time\n");
+ fprintf (stderr, "too long 1\n");
abort ();
}
- start_timer (0);
+ gettimeofday (&tv1, NULL);
acc_wait (0);
- /* Test unseen async-argument. */
- acc_wait (1);
- atime = stop_timer (0);
+ gettimeofday (&tv2, NULL);
+
+ t2 = ((tv2.tv_sec - tv1.tv_sec) * 1000000) + (tv2.tv_usec - tv1.tv_usec);
- if (0.010 < atime)
+ if (t2 > 1000)
{
- fprintf (stderr, "actual time too long\n");
+ fprintf (stderr, "too long 2\n");
abort ();
}
- acc_unmap_data (a);
-
- fini_timers ();
-
- free (a);
- acc_free (d_a);
-
acc_shutdown (acc_device_nvidia);
exit (0);
#include <stdlib.h>
#include <openacc.h>
#include <cuda.h>
-#include "timer.h"
+#include <sys/time.h>
int
main (int argc, char **argv)
{
- CUdevice dev;
CUfunction delay;
CUmodule module;
CUresult r;
- int N;
+ const int N = 2;
int i;
CUstream stream;
- unsigned long *a, *d_a, dticks;
- int nbytes;
- float atime, dtime, hitime, lotime;
- void *kargs[2];
- int clkrate;
- int devnum, nprocs;
+ struct timeval tv1, tv2;
+ time_t t1, t2;
acc_init (acc_device_nvidia);
- devnum = acc_get_device_num (acc_device_nvidia);
-
- r = cuDeviceGet (&dev, devnum);
+ r = cuModuleLoad (&module, "subr.ptx");
if (r != CUDA_SUCCESS)
{
- fprintf (stderr, "cuDeviceGet failed: %d\n", r);
+ fprintf (stderr, "cuModuleLoad failed: %d\n", r);
abort ();
}
- r =
- cuDeviceGetAttribute (&nprocs, CU_DEVICE_ATTRIBUTE_MULTIPROCESSOR_COUNT,
- dev);
+ r = cuModuleGetFunction (&delay, module, "delay");
if (r != CUDA_SUCCESS)
{
- fprintf (stderr, "cuDeviceGetAttribute failed: %d\n", r);
+ fprintf (stderr, "cuModuleGetFunction failed: %d\n", r);
abort ();
}
- r = cuDeviceGetAttribute (&clkrate, CU_DEVICE_ATTRIBUTE_CLOCK_RATE, dev);
- if (r != CUDA_SUCCESS)
- {
- fprintf (stderr, "cuDeviceGetAttribute failed: %d\n", r);
- abort ();
- }
+ gettimeofday (&tv1, NULL);
- r = cuModuleLoad (&module, "subr.ptx");
+ r = cuLaunchKernel (delay, 1, 1, 1, 1, 1, 1, 0, NULL, NULL, 0);
if (r != CUDA_SUCCESS)
{
- fprintf (stderr, "cuModuleLoad failed: %d\n", r);
+ fprintf (stderr, "cuLaunchKernel failed: %d\n", r);
abort ();
}
- r = cuModuleGetFunction (&delay, module, "delay");
+ r = cuCtxSynchronize ();
if (r != CUDA_SUCCESS)
{
- fprintf (stderr, "cuModuleGetFunction failed: %d\n", r);
+ fprintf (stderr, "cuCtxSynchronize failed: %d\n", r);
abort ();
}
- nbytes = nprocs * sizeof (unsigned long);
-
- dtime = 200.0;
-
- dticks = (unsigned long) (dtime * clkrate);
-
- N = nprocs;
-
- a = (unsigned long *) malloc (nbytes);
- d_a = (unsigned long *) acc_malloc (nbytes);
+ gettimeofday (&tv2, NULL);
- acc_map_data (a, d_a, nbytes);
+ t1 = ((tv2.tv_sec - tv1.tv_sec) * 1000000) + (tv2.tv_usec - tv1.tv_usec);
stream = (CUstream) acc_get_cuda_stream (0);
if (stream != NULL)
if (!acc_set_cuda_stream (0, stream))
abort ();
- init_timers (1);
-
- kargs[0] = (void *) &d_a;
- kargs[1] = (void *) &dticks;
-
- start_timer (0);
+ gettimeofday (&tv1, NULL);
for (i = 0; i < N; i++)
{
- r = cuLaunchKernel (delay, 1, 1, 1, 1, 1, 1, 0, stream, kargs, 0);
+ r = cuLaunchKernel (delay, 1, 1, 1, 1, 1, 1, 0, stream, NULL, 0);
if (r != CUDA_SUCCESS)
{
fprintf (stderr, "cuLaunchKernel failed: %d\n", r);
acc_wait (0);
}
- atime = stop_timer (0);
+ gettimeofday (&tv2, NULL);
- hitime = dtime * N;
- hitime += hitime * 0.02;
+ t2 = ((tv2.tv_sec - tv1.tv_sec) * 1000000) + (tv2.tv_usec - tv1.tv_usec);
- lotime = dtime * N;
- lotime -= lotime * 0.02;
+ t1 *= N;
- if (atime > hitime || atime < lotime)
+ if (((abs (t2 - t1) / t1) * 100.0) > 1.0)
{
- fprintf (stderr, "actual time < delay time\n");
+ fprintf (stderr, "too long\n");
abort ();
}
- acc_unmap_data (a);
-
- fini_timers ();
-
- free (a);
- acc_free (d_a);
-
acc_shutdown (acc_device_nvidia);
exit (0);
#include <unistd.h>
#include <openacc.h>
#include <cuda.h>
-#include "timer.h"
+#include <sys/time.h>
int
main (int argc, char **argv)
{
- CUdevice dev;
CUfunction delay;
CUmodule module;
CUresult r;
- int N;
+ const int N = 2;
int i;
CUstream *streams;
- unsigned long *a, *d_a, dticks;
- int nbytes;
- float atime, dtime, hitime, lotime;
- void *kargs[2];
- int clkrate;
- int devnum, nprocs;
+ struct timeval tv1, tv2;
+ time_t t1, t2;
acc_init (acc_device_nvidia);
- devnum = acc_get_device_num (acc_device_nvidia);
-
- r = cuDeviceGet (&dev, devnum);
+ r = cuModuleLoad (&module, "subr.ptx");
if (r != CUDA_SUCCESS)
{
- fprintf (stderr, "cuDeviceGet failed: %d\n", r);
+ fprintf (stderr, "cuModuleLoad failed: %d\n", r);
abort ();
}
- r =
- cuDeviceGetAttribute (&nprocs, CU_DEVICE_ATTRIBUTE_MULTIPROCESSOR_COUNT,
- dev);
+ r = cuModuleGetFunction (&delay, module, "delay");
if (r != CUDA_SUCCESS)
{
- fprintf (stderr, "cuDeviceGetAttribute failed: %d\n", r);
+ fprintf (stderr, "cuModuleGetFunction failed: %d\n", r);
abort ();
}
- r = cuDeviceGetAttribute (&clkrate, CU_DEVICE_ATTRIBUTE_CLOCK_RATE, dev);
- if (r != CUDA_SUCCESS)
- {
- fprintf (stderr, "cuDeviceGetAttribute failed: %d\n", r);
- abort ();
- }
+ gettimeofday (&tv1, NULL);
- r = cuModuleLoad (&module, "subr.ptx");
+ r = cuLaunchKernel (delay, 1, 1, 1, 1, 1, 1, 0, NULL, NULL, 0);
if (r != CUDA_SUCCESS)
{
- fprintf (stderr, "cuModuleLoad failed: %d\n", r);
+ fprintf (stderr, "cuLaunchKernel failed: %d\n", r);
abort ();
}
- r = cuModuleGetFunction (&delay, module, "delay");
+ r = cuCtxSynchronize ();
if (r != CUDA_SUCCESS)
{
- fprintf (stderr, "cuModuleGetFunction failed: %d\n", r);
+ fprintf (stderr, "cuCtxSynchronize failed: %d\n", r);
abort ();
}
- nbytes = nprocs * sizeof (unsigned long);
-
- dtime = 200.0;
-
- dticks = (unsigned long) (dtime * clkrate);
-
- N = nprocs;
+ gettimeofday (&tv2, NULL);
- a = (unsigned long *) malloc (nbytes);
- d_a = (unsigned long *) acc_malloc (nbytes);
-
- acc_map_data (a, d_a, nbytes);
+ t1 = ((tv2.tv_sec - tv1.tv_sec) * 1000000) + (tv2.tv_usec - tv1.tv_usec);
streams = (CUstream *) malloc (N * sizeof (void *));
abort ();
}
- init_timers (1);
-
- kargs[0] = (void *) &d_a;
- kargs[1] = (void *) &dticks;
-
- start_timer (0);
+ gettimeofday (&tv1, NULL);
for (i = 0; i < N; i++)
{
- r = cuLaunchKernel (delay, 1, 1, 1, 1, 1, 1, 0, streams[i], kargs, 0);
+ r = cuLaunchKernel (delay, 1, 1, 1, 1, 1, 1, 0, streams[i], NULL, 0);
if (r != CUDA_SUCCESS)
{
fprintf (stderr, "cuLaunchKernel failed: %d\n", r);
acc_wait (i);
}
- atime = stop_timer (0);
+ gettimeofday (&tv2, NULL);
- hitime = dtime * N;
- hitime += hitime * 0.02;
+ t2 = ((tv2.tv_sec - tv1.tv_sec) * 1000000) + (tv2.tv_usec - tv1.tv_usec);
- lotime = dtime * N;
- lotime -= lotime * 0.02;
+ t1 *= N;
- if (atime > hitime || atime < lotime)
+ if (((abs (t2 - t1) / t1) * 100.0) > 1.0)
{
- fprintf (stderr, "actual time < delay time\n");
+ fprintf (stderr, "too long\n");
abort ();
}
- acc_unmap_data (a);
-
- fini_timers ();
-
free (streams);
- free (a);
- acc_free (d_a);
acc_shutdown (acc_device_nvidia);
#include <unistd.h>
#include <openacc.h>
#include <cuda.h>
-#include "timer.h"
+#include <sys/time.h>
int
main (int argc, char **argv)
{
- CUdevice dev;
CUfunction delay;
CUmodule module;
CUresult r;
CUstream stream;
- unsigned long *a, *d_a, dticks;
- int nbytes;
- float atime, dtime;
- void *kargs[2];
- int clkrate;
- int devnum, nprocs;
+ struct timeval tv1, tv2;
+ time_t t1, t2;
acc_init (acc_device_nvidia);
- devnum = acc_get_device_num (acc_device_nvidia);
-
- r = cuDeviceGet (&dev, devnum);
+ r = cuModuleLoad (&module, "subr.ptx");
if (r != CUDA_SUCCESS)
{
- fprintf (stderr, "cuDeviceGet failed: %d\n", r);
+ fprintf (stderr, "cuModuleLoad failed: %d\n", r);
abort ();
}
- r =
- cuDeviceGetAttribute (&nprocs, CU_DEVICE_ATTRIBUTE_MULTIPROCESSOR_COUNT,
- dev);
+ r = cuModuleGetFunction (&delay, module, "delay");
if (r != CUDA_SUCCESS)
{
- fprintf (stderr, "cuDeviceGetAttribute failed: %d\n", r);
+ fprintf (stderr, "cuModuleGetFunction failed: %d\n", r);
abort ();
}
- r = cuDeviceGetAttribute (&clkrate, CU_DEVICE_ATTRIBUTE_CLOCK_RATE, dev);
- if (r != CUDA_SUCCESS)
- {
- fprintf (stderr, "cuDeviceGetAttribute failed: %d\n", r);
- abort ();
- }
+ gettimeofday (&tv1, NULL);
- r = cuModuleLoad (&module, "subr.ptx");
+ r = cuLaunchKernel (delay, 1, 1, 1, 1, 1, 1, 0, NULL, NULL, 0);
if (r != CUDA_SUCCESS)
{
- fprintf (stderr, "cuModuleLoad failed: %d\n", r);
+ fprintf (stderr, "cuLaunchKernel failed: %d\n", r);
abort ();
}
- r = cuModuleGetFunction (&delay, module, "delay");
+ r = cuCtxSynchronize ();
if (r != CUDA_SUCCESS)
{
- fprintf (stderr, "cuModuleGetFunction failed: %d\n", r);
+ fprintf (stderr, "cuCtxSynchronize failed: %d\n", r);
abort ();
}
- nbytes = nprocs * sizeof (unsigned long);
-
- dtime = 200.0;
-
- dticks = (unsigned long) (dtime * clkrate);
-
- a = (unsigned long *) malloc (nbytes);
- d_a = (unsigned long *) acc_malloc (nbytes);
-
- acc_map_data (a, d_a, nbytes);
+ gettimeofday (&tv2, NULL);
- kargs[0] = (void *) &d_a;
- kargs[1] = (void *) &dticks;
+ t1 = ((tv2.tv_sec - tv1.tv_sec) * 1000000) + (tv2.tv_usec - tv1.tv_usec);
stream = (CUstream) acc_get_cuda_stream (0);
if (stream != NULL)
if (!acc_set_cuda_stream (0, stream))
abort ();
- init_timers (1);
+ gettimeofday (&tv1, NULL);
- start_timer (0);
-
- r = cuLaunchKernel (delay, 1, 1, 1, 1, 1, 1, 0, stream, kargs, 0);
+ r = cuLaunchKernel (delay, 1, 1, 1, 1, 1, 1, 0, stream, NULL, 0);
if (r != CUDA_SUCCESS)
{
fprintf (stderr, "cuLaunchKernel failed: %d\n", r);
acc_wait_all ();
- atime = stop_timer (0);
+ gettimeofday (&tv2, NULL);
+
+ t2 = ((tv2.tv_sec - tv1.tv_sec) * 1000000) + (tv2.tv_usec - tv1.tv_usec);
- if (atime < dtime)
+ if (t2 > (t1 + (t1 * 0.10)))
{
- fprintf (stderr, "actual time < delay time\n");
+ fprintf (stderr, "too long 1\n");
abort ();
}
- start_timer (0);
+ gettimeofday (&tv1, NULL);
acc_wait_all ();
- atime = stop_timer (0);
+ gettimeofday (&tv2, NULL);
+
+ t2 = ((tv2.tv_sec - tv1.tv_sec) * 1000000) + (tv2.tv_usec - tv1.tv_usec);
- if (0.010 < atime)
+ if (t2 > 1000)
{
- fprintf (stderr, "actual time too long\n");
+ fprintf (stderr, "too long 2\n");
abort ();
}
- acc_unmap_data (a);
-
- fini_timers ();
-
- free (a);
- acc_free (d_a);
-
acc_shutdown (acc_device_nvidia);
exit (0);
#include <unistd.h>
#include <openacc.h>
#include <cuda.h>
-#include "timer.h"
+#include <sys/time.h>
int
main (int argc, char **argv)
{
- CUdevice dev;
CUfunction delay;
CUmodule module;
CUresult r;
- int N;
+ const int N = 2;
int i;
CUstream stream;
- unsigned long *a, *d_a, dticks;
- int nbytes;
- float atime, dtime, hitime, lotime;
- void *kargs[2];
- int clkrate;
- int devnum, nprocs;
-
- devnum = 2;
+ struct timeval tv1, tv2;
+ time_t t1, t2;
acc_init (acc_device_nvidia);
- devnum = acc_get_device_num (acc_device_nvidia);
-
- r = cuDeviceGet (&dev, devnum);
+ r = cuModuleLoad (&module, "subr.ptx");
if (r != CUDA_SUCCESS)
{
- fprintf (stderr, "cuDeviceGet failed: %d\n", r);
+ fprintf (stderr, "cuModuleLoad failed: %d\n", r);
abort ();
}
- r =
- cuDeviceGetAttribute (&nprocs, CU_DEVICE_ATTRIBUTE_MULTIPROCESSOR_COUNT,
- dev);
+ r = cuModuleGetFunction (&delay, module, "delay");
if (r != CUDA_SUCCESS)
{
- fprintf (stderr, "cuDeviceGetAttribute failed: %d\n", r);
+ fprintf (stderr, "cuModuleGetFunction failed: %d\n", r);
abort ();
}
- r = cuDeviceGetAttribute (&clkrate, CU_DEVICE_ATTRIBUTE_CLOCK_RATE, dev);
- if (r != CUDA_SUCCESS)
- {
- fprintf (stderr, "cuDeviceGetAttribute failed: %d\n", r);
- abort ();
- }
+ gettimeofday (&tv1, NULL);
- r = cuModuleLoad (&module, "subr.ptx");
+ r = cuLaunchKernel (delay, 1, 1, 1, 1, 1, 1, 0, NULL, NULL, 0);
if (r != CUDA_SUCCESS)
{
- fprintf (stderr, "cuModuleLoad failed: %d\n", r);
+ fprintf (stderr, "cuLaunchKernel failed: %d\n", r);
abort ();
}
- r = cuModuleGetFunction (&delay, module, "delay");
+ r = cuCtxSynchronize ();
if (r != CUDA_SUCCESS)
{
- fprintf (stderr, "cuModuleGetFunction failed: %d\n", r);
+ fprintf (stderr, "cuCtxSynchronize failed: %d\n", r);
abort ();
}
- nbytes = nprocs * sizeof (unsigned long);
-
- dtime = 200.0;
-
- dticks = (unsigned long) (dtime * clkrate);
-
- N = nprocs;
-
- a = (unsigned long *) malloc (nbytes);
- d_a = (unsigned long *) acc_malloc (nbytes);
+ gettimeofday (&tv2, NULL);
- acc_map_data (a, d_a, nbytes);
+ t1 = ((tv2.tv_sec - tv1.tv_sec) * 1000000) + (tv2.tv_usec - tv1.tv_usec);
r = cuStreamCreate (&stream, CU_STREAM_DEFAULT);
if (r != CUDA_SUCCESS)
if (!acc_set_cuda_stream (0, stream))
abort ();
- init_timers (1);
-
- kargs[0] = (void *) &d_a;
- kargs[1] = (void *) &dticks;
-
- start_timer (0);
+ gettimeofday (&tv1, NULL);
for (i = 0; i < N; i++)
{
- r = cuLaunchKernel (delay, 1, 1, 1, 1, 1, 1, 0, stream, kargs, 0);
+ r = cuLaunchKernel (delay, 1, 1, 1, 1, 1, 1, 0, stream, NULL, 0);
if (r != CUDA_SUCCESS)
{
fprintf (stderr, "cuLaunchKernel failed: %d\n", r);
acc_wait (1);
- atime = stop_timer (0);
+ gettimeofday (&tv2, NULL);
if (acc_async_test (0) != 1)
abort ();
if (acc_async_test (1) != 1)
abort ();
- hitime = dtime * N;
- hitime += hitime * 0.02;
+ t2 = ((tv2.tv_sec - tv1.tv_sec) * 1000000) + (tv2.tv_usec - tv1.tv_usec);
- lotime = dtime * N;
- lotime -= lotime * 0.02;
+ t1 *= N;
- if (atime > hitime || atime < lotime)
+ if (((abs (t2 - t1) / t1) * 100.0) > 1.0)
{
- fprintf (stderr, "actual time < delay time\n");
+ fprintf (stderr, "too long\n");
abort ();
}
- acc_unmap_data (a);
-
- fini_timers ();
-
- free (a);
- acc_free (d_a);
-
acc_shutdown (acc_device_nvidia);
exit (0);
#include <unistd.h>
#include <openacc.h>
#include <cuda.h>
-#include "timer.h"
+#include <sys/time.h>
int
main (int argc, char **argv)
{
- CUdevice dev;
CUfunction delay;
CUmodule module;
CUresult r;
- int N;
+ const int N = 2;
int i;
CUstream *streams, stream;
- unsigned long *a, *d_a, dticks;
- int nbytes;
- float atime, dtime;
- void *kargs[2];
- int clkrate;
- int devnum, nprocs;
+ struct timeval tv1, tv2;
+ time_t t1, t2;
acc_init (acc_device_nvidia);
- devnum = acc_get_device_num (acc_device_nvidia);
-
- r = cuDeviceGet (&dev, devnum);
+ r = cuModuleLoad (&module, "subr.ptx");
if (r != CUDA_SUCCESS)
{
- fprintf (stderr, "cuDeviceGet failed: %d\n", r);
+ fprintf (stderr, "cuModuleLoad failed: %d\n", r);
abort ();
}
- r =
- cuDeviceGetAttribute (&nprocs, CU_DEVICE_ATTRIBUTE_MULTIPROCESSOR_COUNT,
- dev);
+ r = cuModuleGetFunction (&delay, module, "delay");
if (r != CUDA_SUCCESS)
{
- fprintf (stderr, "cuDeviceGetAttribute failed: %d\n", r);
+ fprintf (stderr, "cuModuleGetFunction failed: %d\n", r);
abort ();
}
- r = cuDeviceGetAttribute (&clkrate, CU_DEVICE_ATTRIBUTE_CLOCK_RATE, dev);
- if (r != CUDA_SUCCESS)
- {
- fprintf (stderr, "cuDeviceGetAttribute failed: %d\n", r);
- abort ();
- }
+ gettimeofday (&tv1, NULL);
- r = cuModuleLoad (&module, "subr.ptx");
+ r = cuLaunchKernel (delay, 1, 1, 1, 1, 1, 1, 0, NULL, NULL, 0);
if (r != CUDA_SUCCESS)
{
- fprintf (stderr, "cuModuleLoad failed: %d\n", r);
- abort ();
+ fprintf (stderr, "cuLaunchKernel failed: %d\n", r);
+ abort ();
}
- r = cuModuleGetFunction (&delay, module, "delay");
+ r = cuCtxSynchronize ();
if (r != CUDA_SUCCESS)
{
- fprintf (stderr, "cuModuleGetFunction failed: %d\n", r);
- abort ();
+ fprintf (stderr, "cuCtxSynchronize failed: %d\n", r);
+ abort ();
}
- nbytes = nprocs * sizeof (unsigned long);
-
- dtime = 500.0;
-
- dticks = (unsigned long) (dtime * clkrate);
+ gettimeofday (&tv2, NULL);
- N = nprocs;
-
- a = (unsigned long *) malloc (nbytes);
- d_a = (unsigned long *) acc_malloc (nbytes);
-
- acc_map_data (a, d_a, nbytes);
+ t1 = ((tv2.tv_sec - tv1.tv_sec) * 1000000) + (tv2.tv_usec - tv1.tv_usec);
streams = (CUstream *) malloc (N * sizeof (void *));
abort ();
}
- init_timers (1);
-
- kargs[0] = (void *) &d_a;
- kargs[1] = (void *) &dticks;
-
stream = (CUstream) acc_get_cuda_stream (N);
if (stream != NULL)
abort ();
if (!acc_set_cuda_stream (N, stream))
abort ();
- start_timer (0);
+ gettimeofday (&tv1, NULL);
for (i = 0; i < N; i++)
{
- r = cuLaunchKernel (delay, 1, 1, 1, 1, 1, 1, 0, streams[i], kargs, 0);
+ r = cuLaunchKernel (delay, 1, 1, 1, 1, 1, 1, 0, streams[i], NULL, 0);
if (r != CUDA_SUCCESS)
{
fprintf (stderr, "cuLaunchKernel failed: %d\n", r);
}
}
+ gettimeofday (&tv2, NULL);
+
+ t2 = ((tv2.tv_sec - tv1.tv_sec) * 1000000) + (tv2.tv_usec - tv1.tv_usec);
+
acc_wait_all_async (N);
for (i = 0; i <= N; i++)
abort ();
}
- atime = stop_timer (0);
-
- if (atime < dtime)
+ if ((t1 * N) < t2)
{
- fprintf (stderr, "actual time < delay time\n");
+ fprintf (stderr, "too long 1\n");
abort ();
}
- start_timer (0);
+ gettimeofday (&tv1, NULL);
stream = (CUstream) acc_get_cuda_stream (N + 1);
if (stream != NULL)
acc_wait (N + 1);
- atime = stop_timer (0);
+ gettimeofday (&tv2, NULL);
+
+ t1 = ((tv2.tv_sec - tv1.tv_sec) * 1000000) + (tv2.tv_usec - tv1.tv_usec);
- if (0.10 < atime)
+ if (t1 > 1000)
{
- fprintf (stderr, "actual time too long\n");
+ fprintf (stderr, "too long 2\n");
abort ();
}
- start_timer (0);
+ gettimeofday (&tv1, NULL);
acc_wait_all_async (N);
acc_wait (N);
- atime = stop_timer (0);
+ gettimeofday (&tv2, NULL);
- if (0.10 < atime)
+ t1 = ((tv2.tv_sec - tv1.tv_sec) * 1000000) + (tv2.tv_usec - tv1.tv_usec);
+
+ if (t1 > 1000)
{
- fprintf (stderr, "actual time too long\n");
+ fprintf (stderr, "too long 3\n");
abort ();
}
- acc_unmap_data (a);
-
- fini_timers ();
-
free (streams);
- free (a);
- acc_free (d_a);
acc_shutdown (acc_device_nvidia);
int
main (int argc, char **argv)
{
- CUdevice dev;
CUfunction delay2;
CUmodule module;
CUresult r;
- int N;
+ const int N = 32;
int i;
CUstream *streams;
- unsigned long **a, **d_a, *tid, ticks;
+ unsigned long **a, **d_a, *tid;
int nbytes;
- void *kargs[3];
- int clkrate;
- int devnum, nprocs;
+ void *kargs[2];
acc_init (acc_device_nvidia);
- devnum = acc_get_device_num (acc_device_nvidia);
-
- r = cuDeviceGet (&dev, devnum);
- if (r != CUDA_SUCCESS)
- {
- fprintf (stderr, "cuDeviceGet failed: %d\n", r);
- abort ();
- }
-
- r =
- cuDeviceGetAttribute (&nprocs, CU_DEVICE_ATTRIBUTE_MULTIPROCESSOR_COUNT,
- dev);
- if (r != CUDA_SUCCESS)
- {
- fprintf (stderr, "cuDeviceGetAttribute failed: %d\n", r);
- abort ();
- }
-
- r = cuDeviceGetAttribute (&clkrate, CU_DEVICE_ATTRIBUTE_CLOCK_RATE, dev);
- if (r != CUDA_SUCCESS)
- {
- fprintf (stderr, "cuDeviceGetAttribute failed: %d\n", r);
- abort ();
- }
-
r = cuModuleLoad (&module, "subr.ptx");
if (r != CUDA_SUCCESS)
{
nbytes = sizeof (int);
- ticks = (unsigned long) (200.0 * clkrate);
-
- N = nprocs;
-
streams = (CUstream *) malloc (N * sizeof (void *));
a = (unsigned long **) malloc (N * sizeof (unsigned long *));
for (i = 0; i < N; i++)
{
kargs[0] = (void *) &d_a[i];
- kargs[1] = (void *) &ticks;
- kargs[2] = (void *) &tid[i];
+ kargs[1] = (void *) &tid[i];
r = cuLaunchKernel (delay2, 1, 1, 1, 1, 1, 1, 0, streams[i], kargs, 0);
if (r != CUDA_SUCCESS)
fprintf (stderr, "cuLaunchKernel failed: %d\n", r);
abort ();
}
-
- ticks = (unsigned long) (50.0 * clkrate);
}
acc_wait_all_async (0);
--- /dev/null
+/* { dg-do run { target { ! openacc_nvidia_accel_configured } } } */
+
+#include <stdio.h>
+#include <openacc.h>
+
+int
+main (void)
+{
+ fprintf (stderr, "CheCKpOInT\n");
+ acc_init (acc_device_nvidia);
+
+ acc_shutdown (acc_device_nvidia);
+
+ return 0;
+}
+
+/* { dg-output "CheCKpOInT(\n|\r\n|\r).*" } */
+/* { dg-output "device type nvidia not supported" } */
+/* { dg-shouldfail "" } */
-
-#if ACC_DEVICE_TYPE_nvidia
-
#pragma acc routine nohost
-static int clock (void)
-{
- int thetime;
-
- asm __volatile__ ("mov.u32 %0, %%clock;" : "=r"(thetime));
-
- return thetime;
-}
-
-#endif
-
void
-delay (unsigned long *d_o, unsigned long delay)
+delay ()
{
- int start, ticks;
+ int i, sum;
+ const int N = 500000;
- start = clock ();
-
- ticks = 0;
-
- while (ticks < delay)
- ticks = clock () - start;
-
- return;
+ for (i = 0; i < N; i++)
+ sum = sum + 1;
}
+#pragma acc routine nohost
void
-delay2 (unsigned long *d_o, unsigned long delay, unsigned long tid)
+delay2 (unsigned long *d_o, unsigned long tid)
{
- int start, ticks;
+ int i, sum;
+ const int N = 500000;
- start = clock ();
-
- ticks = 0;
-
- while (ticks < delay)
- ticks = clock () - start;
+ for (i = 0; i < N; i++)
+ sum = sum + 1;
d_o[0] = tid;
-
- return;
}
-// BEGIN PREAMBLE
- .version 3.1
- .target sm_30
+ .version 3.1
+ .target sm_30
.address_size 64
-// END PREAMBLE
-// BEGIN FUNCTION DEF: clock
-.func (.param.u32 %out_retval)clock
-{
-.reg.u32 %retval;
- .reg.u64 %hr10;
- .reg.u32 %r22;
- .reg.u32 %r23;
- .reg.u32 %r24;
- .local.align 8 .b8 %frame[8];
- // #APP
-// 7 "subr.c" 1
- mov.u32 %r24, %clock;
-// 0 "" 2
- // #NO_APP
- st.local.u32 [%frame], %r24;
- ld.local.u32 %r22, [%frame];
- mov.u32 %r23, %r22;
- mov.u32 %retval, %r23;
- st.param.u32 [%out_retval], %retval;
- ret;
- }
-// END FUNCTION DEF
-// BEGIN GLOBAL FUNCTION DEF: delay
-.visible .entry delay(.param.u64 %in_ar1, .param.u64 %in_ar2)
-{
- .reg.u64 %ar1;
- .reg.u64 %ar2;
- .reg.u64 %hr10;
- .reg.u64 %r22;
- .reg.u32 %r23;
- .reg.u64 %r24;
- .reg.u64 %r25;
- .reg.u32 %r26;
- .reg.u32 %r27;
- .reg.u32 %r28;
- .reg.u32 %r29;
- .reg.u32 %r30;
- .reg.u64 %r31;
- .reg.pred %r32;
- .local.align 8 .b8 %frame[24];
- ld.param.u64 %ar1, [%in_ar1];
- ld.param.u64 %ar2, [%in_ar2];
- mov.u64 %r24, %ar1;
- st.u64 [%frame+8], %r24;
- mov.u64 %r25, %ar2;
- st.local.u64 [%frame+16], %r25;
+ .visible .entry delay
{
- .param.u32 %retval_in;
- {
- call (%retval_in), clock;
- }
- ld.param.u32 %r26, [%retval_in];
-}
- st.local.u32 [%frame+4], %r26;
- mov.u32 %r27, 0;
- st.local.u32 [%frame], %r27;
- bra $L4;
-$L5:
- {
- .param.u32 %retval_in;
- {
- call (%retval_in), clock;
- }
- ld.param.u32 %r28, [%retval_in];
-}
- mov.u32 %r23, %r28;
- ld.local.u32 %r30, [%frame+4];
- sub.u32 %r29, %r23, %r30;
- st.local.u32 [%frame], %r29;
-$L4:
- ld.local.s32 %r22, [%frame];
- ld.local.u64 %r31, [%frame+16];
- setp.lo.u64 %r32,%r22,%r31;
- @%r32 bra $L5;
+ .reg .u64 %hr10;
+ .reg .u32 %r22;
+ .reg .u32 %r23;
+ .reg .u32 %r24;
+ .reg .u32 %r25;
+ .reg .u32 %r26;
+ .reg .u32 %r27;
+ .reg .u32 %r28;
+ .reg .u32 %r29;
+ .reg .pred %r30;
+ .reg .u64 %frame;
+ .local .align 8 .b8 %farray[16];
+ cvta.local.u64 %frame,%farray;
+ mov.u32 %r22,500000;
+ st.u32 [%frame+8],%r22;
+ mov.u32 %r23,0;
+ st.u32 [%frame],%r23;
+ bra $L2;
+ $L3:
+ ld.u32 %r25,[%frame+4];
+ add.u32 %r24,%r25,1;
+ st.u32 [%frame+4],%r24;
+ ld.u32 %r27,[%frame];
+ add.u32 %r26,%r27,1;
+ st.u32 [%frame],%r26;
+ $L2:
+ ld.u32 %r28,[%frame];
+ ld.u32 %r29,[%frame+8];
+ setp.lt.s32 %r30,%r28,%r29;
+ @%r30
+ bra $L3;
ret;
}
-// END FUNCTION DEF
-// BEGIN GLOBAL FUNCTION DEF: delay2
-.visible .entry delay2(.param.u64 %in_ar1, .param.u64 %in_ar2, .param.u64 %in_ar3)
-{
- .reg.u64 %ar1;
- .reg.u64 %ar2;
- .reg.u64 %ar3;
- .reg.u64 %hr10;
- .reg.u64 %r22;
- .reg.u32 %r23;
- .reg.u64 %r24;
- .reg.u64 %r25;
- .reg.u64 %r26;
- .reg.u32 %r27;
- .reg.u32 %r28;
- .reg.u32 %r29;
- .reg.u32 %r30;
- .reg.u32 %r31;
- .reg.u64 %r32;
- .reg.pred %r33;
- .reg.u64 %r34;
- .reg.u64 %r35;
- .local.align 8 .b8 %frame[32];
- ld.param.u64 %ar1, [%in_ar1];
- ld.param.u64 %ar2, [%in_ar2];
- ld.param.u64 %ar3, [%in_ar3];
- mov.u64 %r24, %ar1;
- st.local.u64 [%frame+8], %r24;
- mov.u64 %r25, %ar2;
- st.local.u64 [%frame+16], %r25;
- mov.u64 %r26, %ar3;
- st.local.u64 [%frame+24], %r26;
- {
- .param.u32 %retval_in;
- {
- call (%retval_in), clock;
- }
- ld.param.u32 %r27, [%retval_in];
-}
- st.local.u32 [%frame+4], %r27;
- mov.u32 %r28, 0;
- st.local.u32 [%frame], %r28;
- bra $L8;
-$L9:
- {
- .param.u32 %retval_in;
+
+ .visible .entry delay2 (.param .u64 %in_ar1, .param .u64 %in_ar2)
{
- call (%retval_in), clock;
- }
- ld.param.u32 %r29, [%retval_in];
-}
- mov.u32 %r23, %r29;
- ld.local.u32 %r31, [%frame+4];
- sub.u32 %r30, %r23, %r31;
- st.local.u32 [%frame], %r30;
-$L8:
- ld.local.s32 %r22, [%frame];
- ld.local.u64 %r32, [%frame+16];
- setp.lo.u64 %r33,%r22,%r32;
- @%r33 bra $L9;
- ld.local.u64 %r34, [%frame+8];
- ld.local.u64 %r35, [%frame+24];
- st.u64 [%r34], %r35;
+ .reg .u64 %ar1;
+ .reg .u64 %ar2;
+ .reg .u64 %hr10;
+ .reg .u64 %r22;
+ .reg .u64 %r23;
+ .reg .u32 %r24;
+ .reg .u32 %r25;
+ .reg .u32 %r26;
+ .reg .u32 %r27;
+ .reg .u32 %r28;
+ .reg .u32 %r29;
+ .reg .u32 %r30;
+ .reg .u32 %r31;
+ .reg .pred %r32;
+ .reg .u64 %r33;
+ .reg .u64 %r34;
+ .reg .u64 %frame;
+ .local .align 8 .b8 %farray[32];
+ cvta.local.u64 %frame,%farray;
+ ld.param.u64 %ar1,[%in_ar1];
+ ld.param.u64 %ar2,[%in_ar2];
+ mov.u64 %r22,%ar1;
+ st.u64 [%frame+16],%r22;
+ mov.u64 %r23,%ar2;
+ st.u64 [%frame+24],%r23;
+ mov.u32 %r24,500000;
+ st.u32 [%frame+8],%r24;
+ mov.u32 %r25,0;
+ st.u32 [%frame],%r25;
+ bra $L5;
+ $L6:
+ ld.u32 %r27,[%frame+4];
+ add.u32 %r26,%r27,1;
+ st.u32 [%frame+4],%r26;
+ ld.u32 %r29,[%frame];
+ add.u32 %r28,%r29,1;
+ st.u32 [%frame],%r28;
+ $L5:
+ ld.u32 %r30,[%frame];
+ ld.u32 %r31,[%frame+8];
+ setp.lt.s32 %r32,%r30,%r31;
+ @%r32
+ bra $L6;
+ ld.u64 %r33,[%frame+16];
+ ld.u64 %r34,[%frame+24];
+ st.u64 [%r33],%r34;
ret;
}
-// END FUNCTION DEF
-
-#include <stdio.h>
-#include <cuda.h>
-
-static int _Tnum_timers;
-static CUevent *_Tstart_events, *_Tstop_events;
-static CUstream _Tstream;
-
-void
-init_timers (int ntimers)
-{
- int i;
- CUresult r;
-
- _Tnum_timers = ntimers;
-
- _Tstart_events = (CUevent *) malloc (_Tnum_timers * sizeof (CUevent));
- _Tstop_events = (CUevent *) malloc (_Tnum_timers * sizeof (CUevent));
-
- r = cuStreamCreate (&_Tstream, CU_STREAM_DEFAULT);
- if (r != CUDA_SUCCESS)
- {
- fprintf (stderr, "cuStreamCreate failed: %d\n", r);
- abort ();
- }
-
- for (i = 0; i < _Tnum_timers; i++)
- {
- r = cuEventCreate (&_Tstart_events[i], CU_EVENT_DEFAULT);
- if (r != CUDA_SUCCESS)
- {
- fprintf (stderr, "cuEventCreate failed: %d\n", r);
- abort ();
- }
-
- r = cuEventCreate (&_Tstop_events[i], CU_EVENT_DEFAULT);
- if (r != CUDA_SUCCESS)
- {
- fprintf (stderr, "cuEventCreate failed: %d\n", r);
- abort ();
- }
- }
-}
-
-void
-fini_timers (void)
-{
- int i;
-
- for (i = 0; i < _Tnum_timers; i++)
- {
- cuEventDestroy (_Tstart_events[i]);
- cuEventDestroy (_Tstop_events[i]);
- }
-
- cuStreamDestroy (_Tstream);
-
- free (_Tstart_events);
- free (_Tstop_events);
-}
-
-void
-start_timer (int timer)
-{
- CUresult r;
-
- r = cuEventRecord (_Tstart_events[timer], _Tstream);
- if (r != CUDA_SUCCESS)
- {
- fprintf (stderr, "cuEventRecord failed: %d\n", r);
- abort ();
- }
-}
-
-float
-stop_timer (int timer)
-{
- CUresult r;
- float etime;
-
- r = cuEventRecord (_Tstop_events[timer], _Tstream);
- if (r != CUDA_SUCCESS)
- {
- fprintf (stderr, "cuEventRecord failed: %d\n", r);
- abort ();
- }
-
- r = cuEventSynchronize (_Tstop_events[timer]);
- if (r != CUDA_SUCCESS)
- {
- fprintf (stderr, "cuEventSynchronize failed: %d\n", r);
- abort ();
- }
-
- r = cuEventElapsedTime (&etime, _Tstart_events[timer], _Tstop_events[timer]);
- if (r != CUDA_SUCCESS)
- {
- fprintf (stderr, "cuEventElapsedTime failed: %d\n", r);
- abort ();
- }
-
- return etime;
-}
c(:) = 0.0
d(:) = 0.0
- !$acc enter data copyin (a(1:N)) create (b(1:N)) create (c(1:N)) create (d(1:N))
+ !$acc enter data copyin (a(1:N)) create (b(1:N)) create (c(1:N)) &
+ !$acc& create (d(1:N))
!$acc parallel async (1)
do i = 1, N
!$acc end parallel
!$acc wait (1)
- !$acc exit data copyout (a(1:N)) copyout (b(1:N)) copyout (c(1:N)) copyout (d(1:N))
+ !$acc exit data copyout (a(1:N)) copyout (b(1:N)) copyout (c(1:N)) &
+ !$acc& copyout (d(1:N))
do i = 1, N
if (a(i) .ne. 3.0) STOP 5
d(:) = 0.0
e(:) = 0.0
- !$acc enter data copyin (a(1:N)) create (b(1:N)) create (c(1:N)) create (d(1:N)) copyin (e(1:N))
+ !$acc enter data copyin (a(1:N)) create (b(1:N)) create (c(1:N)) &
+ !$acc& create (d(1:N)) copyin (e(1:N))
!$acc parallel async (1)
do i = 1, N
!$acc end parallel
!$acc wait (1)
- !$acc exit data copyout (a(1:N)) copyout (b(1:N)) copyout (c(1:N)) copyout (d(1:N)) copyout (e(1:N))
+ !$acc exit data copyout (a(1:N)) copyout (b(1:N)) copyout (c(1:N)) &
+ !$acc& copyout (d(1:N)) copyout (e(1:N))
!$acc exit data delete (N)
do i = 1, N
projects / thirdparty / gcc.git / commitdiff
