Skip to content

Commit

Permalink
Added OpenACC C main example
Browse files Browse the repository at this point in the history
  • Loading branch information
jefflarkin committed Jul 16, 2013
1 parent 4be2a4c commit 497826b
Show file tree
Hide file tree
Showing 3 changed files with 55 additions and 1 deletion.
5 changes: 4 additions & 1 deletion Makefile
Original file line number Diff line number Diff line change
Expand Up @@ -8,10 +8,13 @@ FC=pgfortran
FFLAGS=-fast -acc -ta=nvidia -Minfo=accel
LDFLAGS=-Mcuda

EXES=cuda_main cuf_main cuf_openacc_main thrust
EXES=cuda_main cuf_main cuf_openacc_main openacc_c_main thrust

all: $(EXES)

openacc_c_main: saxpy_cuda.o openacc_c_main.o
$(CC) -o $@ $(CFLAGS) $^ $(LDFLAGS)

cuda_main: saxpy_openacc_c.o cuda_main.o
$(CC) -o $@ $(CFLAGS) $^ $(LDFLAGS)

Expand Down
34 changes: 34 additions & 0 deletions openacc_c_main.c
Original file line number Diff line number Diff line change
@@ -0,0 +1,34 @@
#include <stdio.h>
#include <stdlib.h>
#include <unistd.h>

extern void saxpy(int,float,float*,float*);

int main(int argc, char **argv)
{
float *x, *y, tmp;
int n = 1<<20, i;

x = (float*)malloc(n*sizeof(float));
y = (float*)malloc(n*sizeof(float));

#pragma acc data create(x[0:n]) copyout(y[0:n])
{
#pragma acc kernels
{
for( i = 0; i < n; i++)
{
x[i] = 1.0f;
y[i] = 0.0f;
}
}

#pragma acc host_data use_device(x,y)
{
saxpy(n, 2.0, x, y);
}
}

fprintf(stdout, "y[0] = %f\n",y[0]);
return 0;
}
17 changes: 17 additions & 0 deletions saxpy_cuda.cu
Original file line number Diff line number Diff line change
@@ -0,0 +1,17 @@
__global__
void saxpy_kernel(int n, float a, float *x, float *y)
{
int i = blockDim.x * blockIdx.x + threadIdx.x;

if ( i < n )
y[i] += a * x[i];
}
extern "C" void saxpy(int n ,float a, float *x, float *y)
{
dim3 griddim, blockdim;

blockdim = dim3(128,1,1);
griddim = dim3(n/blockdim.x,1,1);

saxpy_kernel<<<griddim,blockdim>>>(n,a,x,y);
}

0 comments on commit 497826b

Please sign in to comment.