void fdst_gpu(float *data, float *data2, float *data3, int Nx, int Ny, int Lx)
{
float s;
s = sqrt(2.0/(Nx+1));
#pragma acc data copy(data[0:Nx*Ny]), create(data2[0:Lx*Ny], data3[0:2*Lx*Ny])
{
expand_data(data, data2, Nx, Ny, Lx);
expand_idata(data2, data3, Nx, Ny, Lx);
// Copy data to device at start of region and back to host and end of region
// Inside this region the device data pointer will be used
#pragma acc host_data use_device(data3)
{
void *stream = acc_get_cuda_stream(acc_async_sync);
cuda_fft(data3, Lx, Ny, stream);
}
#pragma acc parallel loop independent
for (int i=0;i<Ny;i++)
{
#pragma acc loop independent
for (int j=0;j<Nx;j++) data[Nx*i+j] = -1.0*s*data3[2*Lx*i+2*j+3]/2;
}
}
}
int
main (int argc, char **argv)
{
float atime;
CUstream stream;
CUresult r;
acc_init (acc_device_nvidia);
(void) acc_get_device_num (acc_device_nvidia);
init_timers (1);
stream = (CUstream) acc_get_cuda_stream (0);
if (stream != NULL)
abort ();
r = cuStreamCreate (&stream, CU_STREAM_DEFAULT);
if (r != CUDA_SUCCESS)
{
fprintf (stderr, "cuStreamCreate failed: %d\n", r);
abort ();
}
if (!acc_set_cuda_stream (0, stream))
abort ();
start_timer (0);
acc_wait_all_async (0);
acc_wait (0);
atime = stop_timer (0);
if (0.010 < atime)
{
fprintf (stderr, "actual time too long\n");
abort ();
}
fini_timers ();
acc_shutdown (acc_device_nvidia);
exit (0);
}
int
main (int argc, char **argv)
{
CUstream stream;
CUresult r;
struct timeval tv1, tv2;
time_t t1;
acc_init (acc_device_nvidia);
stream = (CUstream) acc_get_cuda_stream (0);
if (stream != NULL)
abort ();
r = cuStreamCreate (&stream, CU_STREAM_DEFAULT);
if (r != CUDA_SUCCESS)
{
fprintf (stderr, "cuStreamCreate failed: %d\n", r);
abort ();
}
if (!acc_set_cuda_stream (0, stream))
abort ();
gettimeofday (&tv1, NULL);
acc_wait_all_async (0);
acc_wait (0);
gettimeofday (&tv2, NULL);
t1 = ((tv2.tv_sec - tv1.tv_sec) * 1000000) + (tv2.tv_usec - tv1.tv_usec);
if (t1 > 1000)
{
fprintf (stderr, "too long\n");
abort ();
}
acc_shutdown (acc_device_nvidia);
exit (0);
}
int main(int argc, char *argv[])
{
int n = 256;
float *data = malloc(2*n*sizeof(float));
int i;
// Initialize interleaved input data on host
float w = 7.0;
float x;
for(i=0; i<2*n; i+=2) {
x = (float)i/2.0/(n-1);
data[i] = cos(2*M_PI*w*x);
data[i+1] = 0.0;
}
// Copy data to device at start of region and back to host and end of region
#pragma acc data copy(data[0:2*n])
{
// Inside this region the device data pointer will be used
#pragma acc host_data use_device(data)
{
void *stream = acc_get_cuda_stream(acc_async_sync);
launchCUFFT(data, n, stream);
}
}
// Find the frequency
int max_id = 0;
for(i=0; i<n; i+=2) {
if( data[i] > data[max_id] )
max_id = i;
}
printf("frequency = %d\n", max_id/2);
return 0;
}
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)
{
fprintf (stderr, "cuModuleLoad failed: %d\n", r);
abort ();
}
r = cuModuleGetFunction (&delay, module, "delay");
if (r != CUDA_SUCCESS)
{
fprintf (stderr, "cuModuleGetFunction 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 ();
r = cuStreamCreate (&stream, CU_STREAM_DEFAULT);
if (r != CUDA_SUCCESS)
{
fprintf (stderr, "cuStreamCreate failed: %d\n", r);
abort ();
}
if (!acc_set_cuda_stream (0, stream))
abort ();
r = cuLaunchKernel (delay, 1, 1, 1, 1, 1, 1, 0, stream, kargs, 0);
if (r != CUDA_SUCCESS)
{
fprintf (stderr, "cuLaunchKernel failed: %d\n", r);
abort ();
}
if (acc_async_test (0) != 0)
{
fprintf (stderr, "asynchronous operation not running\n");
abort ();
}
sleep (1);
if (acc_async_test (0) != 1)
{
fprintf (stderr, "found asynchronous operation still running\n");
abort ();
}
acc_unmap_data (a);
free (a);
acc_free (d_a);
acc_shutdown (acc_device_nvidia);
exit (0);
}
int
main (int argc, char **argv)
{
CUdevice dev;
CUfunction delay2;
CUmodule module;
CUresult r;
int N;
int i;
CUstream *streams;
unsigned long **a, **d_a, *tid, ticks;
int nbytes;
void *kargs[3];
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)
{
fprintf (stderr, "cuModuleLoad failed: %d\n", r);
abort ();
}
r = cuModuleGetFunction (&delay2, module, "delay2");
if (r != CUDA_SUCCESS)
{
fprintf (stderr, "cuModuleGetFunction failed: %d\n", r);
abort ();
}
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 *));
d_a = (unsigned long **) malloc (N * sizeof (unsigned long *));
tid = (unsigned long *) malloc (N * sizeof (unsigned long));
for (i = 0; i < N; i++)
{
a[i] = (unsigned long *) malloc (sizeof (unsigned long));
*a[i] = N;
d_a[i] = (unsigned long *) acc_malloc (nbytes);
tid[i] = i;
acc_map_data (a[i], d_a[i], nbytes);
streams[i] = (CUstream) acc_get_cuda_stream (i);
if (streams[i] != NULL)
abort ();
r = cuStreamCreate (&streams[i], CU_STREAM_DEFAULT);
if (r != CUDA_SUCCESS)
{
fprintf (stderr, "cuStreamCreate failed: %d\n", r);
abort ();
}
if (!acc_set_cuda_stream (i, streams[i]))
abort ();
}
for (i = 0; i < N; i++)
{
kargs[0] = (void *) &d_a[i];
kargs[1] = (void *) &ticks;
kargs[2] = (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);
for (i = 0; i < N; i++)
{
acc_copyout (a[i], nbytes);
if (*a[i] != i)
abort ();
}
free (streams);
for (i = 0; i < N; i++)
{
free (a[i]);
}
free (a);
free (d_a);
free (tid);
acc_shutdown (acc_device_nvidia);
exit (0);
}
