Initial commit

-------------------------------------------------------------------------------

ENVIRONMENT SETUP

-------------------------------------------------------------------------------

* Module Setting

	- Load the following modules: java, cuda, pgi, and cudampi

	(Run the commands below.)

	$ module load java

	$ module load cuda

	$ module load pgi

	$ module cudampi

* OpenARC Environment Setting

	- Set the following environment variables:

	$ export OPENARC_ARCH=0

* Install OpenARC

	- Run the following command:

	$ /home01/kedu01/shared/installOpenARC.bash

	- Set environment variable, openarc, to the OpenARC install directory:

	$ export openarc=$HOME/local/openarc

include ../../../../make.header

########################

# Set the program name #

########################

BENCHMARK = arrayreduction

########################################

# Set the input C source files (CSRCS) #

########################################

CSRCS = arrayreduction.c

#########################################

# Set macros used for the input program #

#########################################

#_N_ ?= 512

#DEFSET_CPU = -D_N_=$(_N_)

#DEFSET_ACC = -D_N_=$(_N_)

#########################################################

# Makefile options that the user can overwrite          #

# OMP: set to 1 to use OpenMP (default: 0)              # 

# MODE: set to profile to use a built-in profiling tool #

#       (default: normal)                               #

#       If this is set to profile, the runtime system   #

#       will print profiling results according to the   #

#       verbosity level set by OPENARCRT_VERBOSITY      #

#       environment variable.                           # 

#########################################################

OMP ?= 0

MODE ?= normal

#########################################################

# Use the following macros to give program-specific     #

# compiler flags and libraries                          #

# - CFLAGS1 and CLIBS1 to compile the input C program   #

# - CFLAGS2 and CLIBS2 to compile the OpenARC-generated #

#   output C++ program                                  # 

#########################################################

#CFLAGS1 =  

#CFLAGS2 =  

#CLIBS1 = 

#CLIBS2 = 

################################################

# TARGET is where the output binary is stored. #

################################################

#TARGET ?= ./bin

include ../../../../make.template

#! /bin/bash

if [ $# -ge 2 ]; then

	inputSize=$1

	verLevel=$2

elif [ $# -eq 1 ]; then

	inputSize=$1

	verLevel=0

else

	inputSize=512

	verLevel=0

fi

openarcrt="../../../../openarcrt"

openarclib="../../../../lib"

if [ ! -f "openarcConf.txt" ]; then

    cp "openarcConf_NORMAL.txt" "openarcConf.txt"

fi

if [ "$inputSize" != "" ]; then

    mv "openarcConf.txt" "openarcConf.txt_tmp"

    cat "openarcConf.txt_tmp" | sed "s|__inputSize__|${inputSize}|g" > "openarcConf.txt"

    rm "openarcConf.txt_tmp"

fi

if [ "$openarcrt" != "" ]; then

    mv "openarcConf.txt" "openarcConf.txt_tmp"

    cat "openarcConf.txt_tmp" | sed "s|__openarcrt__|${openarcrt}|g" > "openarcConf.txt"

    rm "openarcConf.txt_tmp"

fi

java -classpath $openarclib/cetus.jar:$openarclib/antlr.jar openacc.exec.ACC2GPUDriver -verbosity=${verLevel} -gpuConfFile=openarcConf.txt *.c

echo ""

echo "====> To compile the translated output file:"

echo "\$ make"

echo ""

echo "====> To run the compiled binary:"

echo "\$ cd bin; arrayreduction_ACC"

echo ""

#include <stdio.h>

#include <stdlib.h>

#ifndef WORKERS 

#define WORKERS 256

#pragma openarc #define WORKERS 256

#endif

#ifndef GANGS 

#define GANGS 16

#pragma openarc #define GANGS 16

#endif

#define MSIZE 4096

#pragma openarc #define MSIZE 4096

#ifndef TESTMODE

#define TESTMODE 1

#endif

int main(int argc, char** argv) {

    int size;

    float *A, *B, *C;

	float D = 0.0F;

	float D_ref = 0.0F;

    int i;

    int error = 0;

    size = MSIZE;

    A = (float*) malloc(size * sizeof(float));

    B = (float*) malloc(size * sizeof(float));

    C = (float*) malloc(size * sizeof(float));

    for (i = 0; i < size; i++) {

        A[i] = (float) i;

        B[i] = (float) i * 100;

        C[i] = 1.0F;

		D_ref += (float)i;

    }

#pragma acc data copyin(A[0:size], B[0:size])

    {

#if TESTMODE == 1

#pragma acc kernels loop independent gang(GANGS) worker(WORKERS) reduction(+:C[0:MSIZE], D)

        for (i = 0; i < size; i++) {

            C[i] += A[i] + B[i];

			D += (float)i;

        }

#else

#pragma acc kernels loop independent gang(GANGS) worker(WORKERS) reduction(+:C[0:MSIZE])

        for (i = 0; i < size; i++) {

            C[i] += A[i] + B[i];

        }

#pragma acc kernels loop independent gang(GANGS) worker(WORKERS) reduction(+:D)

        for (i = 0; i < size; i++) {

			D += (float)i;

        }

#endif

    }

    for (i = 0; i < size; i++) {

        if (C[i] != (float) i + (float) i * 100 + 1.0F) error++;

    }

    printf("workers:%d, gangs:%d, size:%d, array reduction error:%d\n", WORKERS, GANGS, size, error);

	if( D != D_ref ) {

		printf("scalar reduction failed!\n");

	} else {

		printf("scalar reduction succeeded!\n");

	}

    return 0;

}