diff --git a/README.md b/README.md
index d285c6e1f6852b9f49d02ae91bd193dde07e79fc..5d493035026121da19c8f95ddbfda9da8a77f6a0 100644
--- a/README.md
+++ b/README.md
@@ -39,13 +39,6 @@ Some common options that can be used with CMake are:
* -DLLVM_ENABLE_ASSERTIONS=On --- Compile with assertion checks enabled (default is Yes for Debug builds, No for all other build types).
-## Building hpvm runtime
-HPVM also includes a runtime library which comprises of low-level, target-specific wrappers required by HPVM's code generation.
-```shell
-cd projects/visc-rt
-make
-cd ..
-```
To use hpvm to compile benchmarks set environment variable `LLVM_SRC_ROOT` to llvm directory in your local repository
```shell
export LLVM_SRC_ROOT=<full path to hpvm>/llvm
diff --git a/hpvm/test/README.md b/hpvm/test/README.md
index b391bd2186b696685adb2645c288c01ea7006850..e709ef04195c90e0f91c2a4b4a4b1d2f0b716d1f 100644
--- a/hpvm/test/README.md
+++ b/hpvm/test/README.md
@@ -4,9 +4,7 @@ Tests are provided, along with a template Makefile for user projects.
## Parboil
Several tests from the [parboil suite](http://impact.crhc.illinois.edu/parboil/parboil.aspx) have been ported to HPVM.
To run one of these tests, navigate to its directory under `parboil/benchmarks/`.
-Tests may be built for the cpu or gpu with hpvm, and openCL versions are provided for comparison.
-Check under the `src/` directory in each benchmark to see which versions are available,
-denoted by the names of the subdirectories.
+Tests may be built for the cpu or gpu with hpvm.
```
# sgemm example
cd parboil/benchmarks/sgemm
@@ -16,9 +14,6 @@ make run TARGET=seq VERSION=visc
# HPVM gpu
make TARGET=gpu VERSION=visc
make run TARGET=gpu VERSION=visc
-# openCL
-make VERSION=opencl_base
-make run VERSION=opencl_base
```
## Cava
diff --git a/hpvm/test/parboil/benchmarks/lbm/src/opencl_nvidia/Makefile b/hpvm/test/parboil/benchmarks/lbm/src/opencl_nvidia/Makefile
deleted file mode 100644
index 25131be39c23baa95b34e95444f7e19e1d03b389..0000000000000000000000000000000000000000
--- a/hpvm/test/parboil/benchmarks/lbm/src/opencl_nvidia/Makefile
+++ /dev/null
@@ -1,8 +0,0 @@
-# (c) 2010 The Board of Trustees of the University of Illinois.
-
-LANGUAGE=opencl
-SRCDIR_OBJS=main.o lbm.o ocl.o
-APP_CUDALDFLAGS=-lm
-APP_CFLAGS=-ffast-math -O3
-APP_CXXFLAGS=-ffast-math -O3
-KERNEL_OBJS=kernel_offline.nvptx.s
diff --git a/hpvm/test/parboil/benchmarks/lbm/src/opencl_nvidia/kernel.cl b/hpvm/test/parboil/benchmarks/lbm/src/opencl_nvidia/kernel.cl
deleted file mode 100644
index 3f34ea5ef25943ac7eeb18eead429f70ec4cf807..0000000000000000000000000000000000000000
--- a/hpvm/test/parboil/benchmarks/lbm/src/opencl_nvidia/kernel.cl
+++ /dev/null
@@ -1,176 +0,0 @@
-/***************************************************************************
- *cr
- *cr (C) Copyright 2010 The Board of Trustees of the
- *cr University of Illinois
- *cr All Rights Reserved
- *cr
- ***************************************************************************/
-
-#ifndef LBM_KERNEL_CL
-#define LBM_KERNEL_CL
-
-#include "layout_config.h"
-#include "lbm_macros.h"
-/******************************************************************************/
-
-__kernel void performStreamCollide_kernel( __global float* srcGrid, __global float* dstGrid )
-{
- srcGrid += MARGIN;
- dstGrid += MARGIN;
-
-
- //Using some predefined macros here. Consider this the declaration
- // and initialization of the variables SWEEP_X, SWEEP_Y and SWEEP_Z
-
- SWEEP_VAR
- SWEEP_X = get_local_id(0);
- SWEEP_Y = get_group_id(0);
- SWEEP_Z = get_group_id(1);
-
- float temp_swp, tempC, tempN, tempS, tempE, tempW, tempT, tempB;
- float tempNE, tempNW, tempSE, tempSW, tempNT, tempNB, tempST ;
- float tempSB, tempET, tempEB, tempWT, tempWB ;
-
- //Load all of the input fields
- //This is a gather operation of the SCATTER preprocessor variable
- // is undefined in layout_config.h, or a "local" read otherwise
- tempC = SRC_C(srcGrid);
-
- tempN = SRC_N(srcGrid);
- tempS = SRC_S(srcGrid);
- tempE = SRC_E(srcGrid);
- tempW = SRC_W(srcGrid);
- tempT = SRC_T(srcGrid);
- tempB = SRC_B(srcGrid);
-
- tempNE = SRC_NE(srcGrid);
- tempNW = SRC_NW(srcGrid);
- tempSE = SRC_SE(srcGrid);
- tempSW = SRC_SW(srcGrid);
- tempNT = SRC_NT(srcGrid);
- tempNB = SRC_NB(srcGrid);
- tempST = SRC_ST(srcGrid);
- tempSB = SRC_SB(srcGrid);
- tempET = SRC_ET(srcGrid);
- tempEB = SRC_EB(srcGrid);
- tempWT = SRC_WT(srcGrid);
- tempWB = SRC_WB(srcGrid);
-
- //Test whether the cell is fluid or obstacle
- if(as_uint(LOCAL(srcGrid,FLAGS)) & (OBSTACLE)) {
-
- //Swizzle the inputs: reflect any fluid coming into this cell
- // back to where it came from
- temp_swp = tempN ; tempN = tempS ; tempS = temp_swp ;
- temp_swp = tempE ; tempE = tempW ; tempW = temp_swp;
- temp_swp = tempT ; tempT = tempB ; tempB = temp_swp;
- temp_swp = tempNE; tempNE = tempSW ; tempSW = temp_swp;
- temp_swp = tempNW; tempNW = tempSE ; tempSE = temp_swp;
- temp_swp = tempNT ; tempNT = tempSB ; tempSB = temp_swp;
- temp_swp = tempNB ; tempNB = tempST ; tempST = temp_swp;
- temp_swp = tempET ; tempET= tempWB ; tempWB = temp_swp;
- temp_swp = tempEB ; tempEB = tempWT ; tempWT = temp_swp;
- }
- else {
-
- //The math meat of LBM: ignore for optimization
- float ux, uy, uz, rho, u2;
- float temp1, temp2, temp_base;
- rho = tempC + tempN
- + tempS + tempE
- + tempW + tempT
- + tempB + tempNE
- + tempNW + tempSE
- + tempSW + tempNT
- + tempNB + tempST
- + tempSB + tempET
- + tempEB + tempWT
- + tempWB;
-
- ux = + tempE - tempW
- + tempNE - tempNW
- + tempSE - tempSW
- + tempET + tempEB
- - tempWT - tempWB;
-
- uy = + tempN - tempS
- + tempNE + tempNW
- - tempSE - tempSW
- + tempNT + tempNB
- - tempST - tempSB;
-
- uz = + tempT - tempB
- + tempNT - tempNB
- + tempST - tempSB
- + tempET - tempEB
- + tempWT - tempWB;
-
- ux /= rho;
- uy /= rho;
- uz /= rho;
-
- if(as_uint(LOCAL(srcGrid,FLAGS)) & (ACCEL)) {
-
- ux = 0.005f;
- uy = 0.002f;
- uz = 0.000f;
- }
-
- u2 = 1.5f * (ux*ux + uy*uy + uz*uz) - 1.0f;
- temp_base = OMEGA*rho;
- temp1 = DFL1*temp_base;
-
- //Put the output values for this cell in the shared memory
- temp_base = OMEGA*rho;
- temp1 = DFL1*temp_base;
- temp2 = 1.0f-OMEGA;
- tempC = temp2*tempC + temp1*( - u2);
- temp1 = DFL2*temp_base;
- tempN = temp2*tempN + temp1*( uy*(4.5f*uy + 3.0f) - u2);
- tempS = temp2*tempS + temp1*( uy*(4.5f*uy - 3.0f) - u2);
- tempT = temp2*tempT + temp1*( uz*(4.5f*uz + 3.0f) - u2);
- tempB = temp2*tempB + temp1*( uz*(4.5f*uz - 3.0f) - u2);
- tempE = temp2*tempE + temp1*( ux*(4.5f*ux + 3.0f) - u2);
- tempW = temp2*tempW + temp1*( ux*(4.5f*ux - 3.0f) - u2);
- temp1 = DFL3*temp_base;
- tempNT= temp2*tempNT + temp1 *( (+uy+uz)*(4.5f*(+uy+uz) + 3.0f) - u2);
- tempNB= temp2*tempNB + temp1 *( (+uy-uz)*(4.5f*(+uy-uz) + 3.0f) - u2);
- tempST= temp2*tempST + temp1 *( (-uy+uz)*(4.5f*(-uy+uz) + 3.0f) - u2);
- tempSB= temp2*tempSB + temp1 *( (-uy-uz)*(4.5f*(-uy-uz) + 3.0f) - u2);
- tempNE = temp2*tempNE + temp1 *( (+ux+uy)*(4.5f*(+ux+uy) + 3.0f) - u2);
- tempSE = temp2*tempSE + temp1 *((+ux-uy)*(4.5f*(+ux-uy) + 3.0f) - u2);
- tempET = temp2*tempET + temp1 *( (+ux+uz)*(4.5f*(+ux+uz) + 3.0f) - u2);
- tempEB = temp2*tempEB + temp1 *( (+ux-uz)*(4.5f*(+ux-uz) + 3.0f) - u2);
- tempNW = temp2*tempNW + temp1 *( (-ux+uy)*(4.5f*(-ux+uy) + 3.0f) - u2);
- tempSW = temp2*tempSW + temp1 *( (-ux-uy)*(4.5f*(-ux-uy) + 3.0f) - u2);
- tempWT = temp2*tempWT + temp1 *( (-ux+uz)*(4.5f*(-ux+uz) + 3.0f) - u2);
- tempWB = temp2*tempWB + temp1 *( (-ux-uz)*(4.5f*(-ux-uz) + 3.0f) - u2);
- }
-
- //Write the results computed above
- //This is a scatter operation of the SCATTER preprocessor variable
- // is defined in layout_config.h, or a "local" write otherwise
- DST_C ( dstGrid ) = tempC;
-
- DST_N ( dstGrid ) = tempN;
- DST_S ( dstGrid ) = tempS;
- DST_E ( dstGrid ) = tempE;
- DST_W ( dstGrid ) = tempW;
- DST_T ( dstGrid ) = tempT;
- DST_B ( dstGrid ) = tempB;
-
- DST_NE( dstGrid ) = tempNE;
- DST_NW( dstGrid ) = tempNW;
- DST_SE( dstGrid ) = tempSE;
- DST_SW( dstGrid ) = tempSW;
- DST_NT( dstGrid ) = tempNT;
- DST_NB( dstGrid ) = tempNB;
- DST_ST( dstGrid ) = tempST;
- DST_SB( dstGrid ) = tempSB;
- DST_ET( dstGrid ) = tempET;
- DST_EB( dstGrid ) = tempEB;
- DST_WT( dstGrid ) = tempWT;
- DST_WB( dstGrid ) = tempWB;
-}
-
-#endif // LBM_KERNEL_CL
diff --git a/hpvm/test/parboil/benchmarks/lbm/src/opencl_nvidia/kernel_offline.cl b/hpvm/test/parboil/benchmarks/lbm/src/opencl_nvidia/kernel_offline.cl
deleted file mode 100644
index 3f34ea5ef25943ac7eeb18eead429f70ec4cf807..0000000000000000000000000000000000000000
--- a/hpvm/test/parboil/benchmarks/lbm/src/opencl_nvidia/kernel_offline.cl
+++ /dev/null
@@ -1,176 +0,0 @@
-/***************************************************************************
- *cr
- *cr (C) Copyright 2010 The Board of Trustees of the
- *cr University of Illinois
- *cr All Rights Reserved
- *cr
- ***************************************************************************/
-
-#ifndef LBM_KERNEL_CL
-#define LBM_KERNEL_CL
-
-#include "layout_config.h"
-#include "lbm_macros.h"
-/******************************************************************************/
-
-__kernel void performStreamCollide_kernel( __global float* srcGrid, __global float* dstGrid )
-{
- srcGrid += MARGIN;
- dstGrid += MARGIN;
-
-
- //Using some predefined macros here. Consider this the declaration
- // and initialization of the variables SWEEP_X, SWEEP_Y and SWEEP_Z
-
- SWEEP_VAR
- SWEEP_X = get_local_id(0);
- SWEEP_Y = get_group_id(0);
- SWEEP_Z = get_group_id(1);
-
- float temp_swp, tempC, tempN, tempS, tempE, tempW, tempT, tempB;
- float tempNE, tempNW, tempSE, tempSW, tempNT, tempNB, tempST ;
- float tempSB, tempET, tempEB, tempWT, tempWB ;
-
- //Load all of the input fields
- //This is a gather operation of the SCATTER preprocessor variable
- // is undefined in layout_config.h, or a "local" read otherwise
- tempC = SRC_C(srcGrid);
-
- tempN = SRC_N(srcGrid);
- tempS = SRC_S(srcGrid);
- tempE = SRC_E(srcGrid);
- tempW = SRC_W(srcGrid);
- tempT = SRC_T(srcGrid);
- tempB = SRC_B(srcGrid);
-
- tempNE = SRC_NE(srcGrid);
- tempNW = SRC_NW(srcGrid);
- tempSE = SRC_SE(srcGrid);
- tempSW = SRC_SW(srcGrid);
- tempNT = SRC_NT(srcGrid);
- tempNB = SRC_NB(srcGrid);
- tempST = SRC_ST(srcGrid);
- tempSB = SRC_SB(srcGrid);
- tempET = SRC_ET(srcGrid);
- tempEB = SRC_EB(srcGrid);
- tempWT = SRC_WT(srcGrid);
- tempWB = SRC_WB(srcGrid);
-
- //Test whether the cell is fluid or obstacle
- if(as_uint(LOCAL(srcGrid,FLAGS)) & (OBSTACLE)) {
-
- //Swizzle the inputs: reflect any fluid coming into this cell
- // back to where it came from
- temp_swp = tempN ; tempN = tempS ; tempS = temp_swp ;
- temp_swp = tempE ; tempE = tempW ; tempW = temp_swp;
- temp_swp = tempT ; tempT = tempB ; tempB = temp_swp;
- temp_swp = tempNE; tempNE = tempSW ; tempSW = temp_swp;
- temp_swp = tempNW; tempNW = tempSE ; tempSE = temp_swp;
- temp_swp = tempNT ; tempNT = tempSB ; tempSB = temp_swp;
- temp_swp = tempNB ; tempNB = tempST ; tempST = temp_swp;
- temp_swp = tempET ; tempET= tempWB ; tempWB = temp_swp;
- temp_swp = tempEB ; tempEB = tempWT ; tempWT = temp_swp;
- }
- else {
-
- //The math meat of LBM: ignore for optimization
- float ux, uy, uz, rho, u2;
- float temp1, temp2, temp_base;
- rho = tempC + tempN
- + tempS + tempE
- + tempW + tempT
- + tempB + tempNE
- + tempNW + tempSE
- + tempSW + tempNT
- + tempNB + tempST
- + tempSB + tempET
- + tempEB + tempWT
- + tempWB;
-
- ux = + tempE - tempW
- + tempNE - tempNW
- + tempSE - tempSW
- + tempET + tempEB
- - tempWT - tempWB;
-
- uy = + tempN - tempS
- + tempNE + tempNW
- - tempSE - tempSW
- + tempNT + tempNB
- - tempST - tempSB;
-
- uz = + tempT - tempB
- + tempNT - tempNB
- + tempST - tempSB
- + tempET - tempEB
- + tempWT - tempWB;
-
- ux /= rho;
- uy /= rho;
- uz /= rho;
-
- if(as_uint(LOCAL(srcGrid,FLAGS)) & (ACCEL)) {
-
- ux = 0.005f;
- uy = 0.002f;
- uz = 0.000f;
- }
-
- u2 = 1.5f * (ux*ux + uy*uy + uz*uz) - 1.0f;
- temp_base = OMEGA*rho;
- temp1 = DFL1*temp_base;
-
- //Put the output values for this cell in the shared memory
- temp_base = OMEGA*rho;
- temp1 = DFL1*temp_base;
- temp2 = 1.0f-OMEGA;
- tempC = temp2*tempC + temp1*( - u2);
- temp1 = DFL2*temp_base;
- tempN = temp2*tempN + temp1*( uy*(4.5f*uy + 3.0f) - u2);
- tempS = temp2*tempS + temp1*( uy*(4.5f*uy - 3.0f) - u2);
- tempT = temp2*tempT + temp1*( uz*(4.5f*uz + 3.0f) - u2);
- tempB = temp2*tempB + temp1*( uz*(4.5f*uz - 3.0f) - u2);
- tempE = temp2*tempE + temp1*( ux*(4.5f*ux + 3.0f) - u2);
- tempW = temp2*tempW + temp1*( ux*(4.5f*ux - 3.0f) - u2);
- temp1 = DFL3*temp_base;
- tempNT= temp2*tempNT + temp1 *( (+uy+uz)*(4.5f*(+uy+uz) + 3.0f) - u2);
- tempNB= temp2*tempNB + temp1 *( (+uy-uz)*(4.5f*(+uy-uz) + 3.0f) - u2);
- tempST= temp2*tempST + temp1 *( (-uy+uz)*(4.5f*(-uy+uz) + 3.0f) - u2);
- tempSB= temp2*tempSB + temp1 *( (-uy-uz)*(4.5f*(-uy-uz) + 3.0f) - u2);
- tempNE = temp2*tempNE + temp1 *( (+ux+uy)*(4.5f*(+ux+uy) + 3.0f) - u2);
- tempSE = temp2*tempSE + temp1 *((+ux-uy)*(4.5f*(+ux-uy) + 3.0f) - u2);
- tempET = temp2*tempET + temp1 *( (+ux+uz)*(4.5f*(+ux+uz) + 3.0f) - u2);
- tempEB = temp2*tempEB + temp1 *( (+ux-uz)*(4.5f*(+ux-uz) + 3.0f) - u2);
- tempNW = temp2*tempNW + temp1 *( (-ux+uy)*(4.5f*(-ux+uy) + 3.0f) - u2);
- tempSW = temp2*tempSW + temp1 *( (-ux-uy)*(4.5f*(-ux-uy) + 3.0f) - u2);
- tempWT = temp2*tempWT + temp1 *( (-ux+uz)*(4.5f*(-ux+uz) + 3.0f) - u2);
- tempWB = temp2*tempWB + temp1 *( (-ux-uz)*(4.5f*(-ux-uz) + 3.0f) - u2);
- }
-
- //Write the results computed above
- //This is a scatter operation of the SCATTER preprocessor variable
- // is defined in layout_config.h, or a "local" write otherwise
- DST_C ( dstGrid ) = tempC;
-
- DST_N ( dstGrid ) = tempN;
- DST_S ( dstGrid ) = tempS;
- DST_E ( dstGrid ) = tempE;
- DST_W ( dstGrid ) = tempW;
- DST_T ( dstGrid ) = tempT;
- DST_B ( dstGrid ) = tempB;
-
- DST_NE( dstGrid ) = tempNE;
- DST_NW( dstGrid ) = tempNW;
- DST_SE( dstGrid ) = tempSE;
- DST_SW( dstGrid ) = tempSW;
- DST_NT( dstGrid ) = tempNT;
- DST_NB( dstGrid ) = tempNB;
- DST_ST( dstGrid ) = tempST;
- DST_SB( dstGrid ) = tempSB;
- DST_ET( dstGrid ) = tempET;
- DST_EB( dstGrid ) = tempEB;
- DST_WT( dstGrid ) = tempWT;
- DST_WB( dstGrid ) = tempWB;
-}
-
-#endif // LBM_KERNEL_CL
diff --git a/hpvm/test/parboil/benchmarks/lbm/src/opencl_nvidia/layout_config.h b/hpvm/test/parboil/benchmarks/lbm/src/opencl_nvidia/layout_config.h
deleted file mode 100644
index d44088661d313eeca6d44612549337b5a2630e04..0000000000000000000000000000000000000000
--- a/hpvm/test/parboil/benchmarks/lbm/src/opencl_nvidia/layout_config.h
+++ /dev/null
@@ -1,86 +0,0 @@
-/***************************************************************************
- *cr
- *cr (C) Copyright 2010 The Board of Trustees of the
- *cr University of Illinois
- *cr All Rights Reserved
- *cr
- ***************************************************************************/
-
-/*############################################################################*/
-
-#ifndef _LAYOUT_CONFIG_H_
-#define _LAYOUT_CONFIG_H_
-
-/*############################################################################*/
-
-// Unchangeable settings: volume simulation size for the given example
-#define SIZE_X (120)
-#define SIZE_Y (120)
-#define SIZE_Z (150)
-
-// Changeable settings
-// Padding in each dimension
-#define PADDING_X (8)
-#define PADDING_Y (0)
-#define PADDING_Z (4)
-
-// Pitch in each dimension
-#define PADDED_X (SIZE_X + PADDING_X)
-#define PADDED_Y (SIZE_Y + PADDING_Y)
-#define PADDED_Z (SIZE_Z + PADDING_Z)
-
-#define TOTAL_CELLS (SIZE_X * SIZE_Y * SIZE_Z)
-#define TOTAL_PADDED_CELLS (PADDED_X * PADDED_Y * PADDED_Z)
-
-// Flattening function
-// This macro will be used to map a 3-D index and element to a value
-#define CALC_INDEX(x, y, z, e) \
- (TOTAL_PADDED_CELLS * e + ((x) + (y)*PADDED_X + (z)*PADDED_X * PADDED_Y))
-
-#define MARGIN (CALC_INDEX(0, 0, 2, 0) - CALC_INDEX(0, 0, 0, 0))
-
-// Set this value to 1 for GATHER, or 0 for SCATTER
-#if 1
-#define GATHER
-#else
-#define SCATTER
-#endif
-
-// OpenCL block size (not trivially changeable here)
-#define BLOCK_SIZE SIZE_X
-
-/*############################################################################*/
-
-typedef enum {
- C = 0,
- N,
- S,
- E,
- W,
- T,
- B,
- NE,
- NW,
- SE,
- SW,
- NT,
- NB,
- ST,
- SB,
- ET,
- EB,
- WT,
- WB,
- FLAGS,
- N_CELL_ENTRIES
-} CELL_ENTRIES;
-
-#define N_DISTR_FUNCS FLAGS
-
-typedef enum {
- OBSTACLE = 1 << 0,
- ACCEL = 1 << 1,
- IN_OUT_FLOW = 1 << 2
-} CELL_FLAGS;
-
-#endif /* _CONFIG_H_ */
diff --git a/hpvm/test/parboil/benchmarks/lbm/src/opencl_nvidia/lbm.c b/hpvm/test/parboil/benchmarks/lbm/src/opencl_nvidia/lbm.c
deleted file mode 100644
index 14ffa4211b3763d7c1c6538e693a76be61a0b158..0000000000000000000000000000000000000000
--- a/hpvm/test/parboil/benchmarks/lbm/src/opencl_nvidia/lbm.c
+++ /dev/null
@@ -1,321 +0,0 @@
-/***************************************************************************
- *cr
- *cr (C) Copyright 2010 The Board of Trustees of the
- *cr University of Illinois
- *cr All Rights Reserved
- *cr
- ***************************************************************************/
-
-/*############################################################################*/
-
-// includes, system
-#include <CL/cl.h>
-#include <float.h>
-#include <math.h>
-#include <stdio.h>
-#include <stdlib.h>
-#include <string.h>
-
-// includes, project
-#include "layout_config.h"
-#include "lbm.h"
-#include "lbm_macros.h"
-#include "ocl.h"
-
-/******************************************************************************/
-
-void OpenCL_LBM_performStreamCollide(const OpenCL_Param *prm, cl_mem srcGrid,
- cl_mem dstGrid) {
-
- cl_int clStatus;
-
- clStatus = clSetKernelArg(prm->clKernel, 0, sizeof(cl_mem), (void *)&srcGrid);
- CHECK_ERROR("clSetKernelArg")
-
- clStatus = clSetKernelArg(prm->clKernel, 1, sizeof(cl_mem), (void *)&dstGrid);
- CHECK_ERROR("clSetKernelArg")
-
- size_t dimBlock[3] = {SIZE_X, 1, 1};
- size_t dimGrid[3] = {SIZE_X * SIZE_Y, SIZE_Z, 1};
- clStatus = clEnqueueNDRangeKernel(prm->clCommandQueue, prm->clKernel, 3, NULL,
- dimGrid, dimBlock, 0, NULL, NULL);
- CHECK_ERROR("clEnqueueNDRangeKernel")
-
- clStatus = clFinish(prm->clCommandQueue);
- CHECK_ERROR("clFinish")
-}
-/*############################################################################*/
-
-void LBM_allocateGrid(float **ptr) {
- const size_t size = TOTAL_PADDED_CELLS * N_CELL_ENTRIES * sizeof(float);
-
- *ptr = (float *)malloc(size);
- if (!*ptr) {
- printf("LBM_allocateGrid: could not allocate %.1f MByte\n",
- size / (1024.0 * 1024.0));
- exit(1);
- }
-
- memset(*ptr, 0, size);
-
- printf("LBM_allocateGrid: allocated %.1f MByte\n", size / (1024.0 * 1024.0));
-
- *ptr += MARGIN;
-}
-
-/******************************************************************************/
-
-void OpenCL_LBM_allocateGrid(const OpenCL_Param *prm, cl_mem *ptr) {
- const size_t size = TOTAL_PADDED_CELLS * N_CELL_ENTRIES * sizeof(float);
- cl_int clStatus;
- *ptr =
- clCreateBuffer(prm->clContext, CL_MEM_READ_WRITE, size, NULL, &clStatus);
- CHECK_ERROR("clCreateBuffer")
-}
-
-/*############################################################################*/
-
-void LBM_freeGrid(float **ptr) {
- free(*ptr - MARGIN);
- *ptr = NULL;
-}
-
-/******************************************************************************/
-
-void OpenCL_LBM_freeGrid(cl_mem ptr) { clReleaseMemObject(ptr); }
-
-/*############################################################################*/
-
-void LBM_initializeGrid(LBM_Grid grid) {
- SWEEP_VAR
-
- SWEEP_START(0, 0, 0, 0, 0, SIZE_Z)
- SRC_C(grid) = DFL1;
- SRC_N(grid) = DFL2;
- SRC_S(grid) = DFL2;
- SRC_E(grid) = DFL2;
- SRC_W(grid) = DFL2;
- SRC_T(grid) = DFL2;
- SRC_B(grid) = DFL2;
- SRC_NE(grid) = DFL3;
- SRC_NW(grid) = DFL3;
- SRC_SE(grid) = DFL3;
- SRC_SW(grid) = DFL3;
- SRC_NT(grid) = DFL3;
- SRC_NB(grid) = DFL3;
- SRC_ST(grid) = DFL3;
- SRC_SB(grid) = DFL3;
- SRC_ET(grid) = DFL3;
- SRC_EB(grid) = DFL3;
- SRC_WT(grid) = DFL3;
- SRC_WB(grid) = DFL3;
-
- CLEAR_ALL_FLAGS_SWEEP(grid);
- SWEEP_END
-}
-
-/******************************************************************************/
-
-void OpenCL_LBM_initializeGrid(const OpenCL_Param *prm, cl_mem d_grid,
- LBM_Grid h_grid) {
- const size_t size = TOTAL_PADDED_CELLS * N_CELL_ENTRIES * sizeof(float);
- cl_int clStatus;
- clStatus = clEnqueueWriteBuffer(prm->clCommandQueue, d_grid, CL_TRUE, 0, size,
- h_grid - MARGIN, 0, NULL, NULL);
- CHECK_ERROR("clEnqueueWriteBuffer")
-}
-
-void OpenCL_LBM_getDeviceGrid(const OpenCL_Param *prm, cl_mem d_grid,
- LBM_Grid h_grid) {
- const size_t size = TOTAL_PADDED_CELLS * N_CELL_ENTRIES * sizeof(float);
- cl_int clStatus;
- clStatus = clEnqueueReadBuffer(prm->clCommandQueue, d_grid, CL_TRUE, 0, size,
- h_grid - MARGIN, 0, NULL, NULL);
- CHECK_ERROR("clEnqueueReadBuffer")
-}
-
-/*############################################################################*/
-
-void LBM_swapGrids(cl_mem *grid1, cl_mem *grid2) {
- cl_mem aux = *grid1;
- *grid1 = *grid2;
- *grid2 = aux;
-}
-
-/*############################################################################*/
-
-void LBM_loadObstacleFile(LBM_Grid grid, const char *filename) {
- int x, y, z;
-
- FILE *file = fopen(filename, "rb");
-
- for (z = 0; z < SIZE_Z; z++) {
- for (y = 0; y < SIZE_Y; y++) {
- for (x = 0; x < SIZE_X; x++) {
- if (fgetc(file) != '.')
- SET_FLAG(grid, x, y, z, OBSTACLE);
- }
- fgetc(file);
- }
- fgetc(file);
- }
-
- fclose(file);
-}
-
-/*############################################################################*/
-
-void LBM_initializeSpecialCellsForLDC(LBM_Grid grid) {
- int x, y, z;
-
- for (z = -2; z < SIZE_Z + 2; z++) {
- for (y = 0; y < SIZE_Y; y++) {
- for (x = 0; x < SIZE_X; x++) {
- if (x == 0 || x == SIZE_X - 1 || y == 0 || y == SIZE_Y - 1 || z == 0 ||
- z == SIZE_Z - 1) {
- SET_FLAG(grid, x, y, z, OBSTACLE);
- } else {
- if ((z == 1 || z == SIZE_Z - 2) && x > 1 && x < SIZE_X - 2 && y > 1 &&
- y < SIZE_Y - 2) {
- SET_FLAG(grid, x, y, z, ACCEL);
- }
- }
- }
- }
- }
-}
-
-/*############################################################################*/
-
-void LBM_showGridStatistics(LBM_Grid grid) {
- int nObstacleCells = 0, nAccelCells = 0, nFluidCells = 0;
- float ux, uy, uz;
- float minU2 = 1e+30, maxU2 = -1e+30, u2;
- float minRho = 1e+30, maxRho = -1e+30, rho;
- float mass = 0;
-
- SWEEP_VAR
-
- SWEEP_START(0, 0, 0, 0, 0, SIZE_Z)
- rho = LOCAL(grid, C) + LOCAL(grid, N) + LOCAL(grid, S) + LOCAL(grid, E) +
- LOCAL(grid, W) + LOCAL(grid, T) + LOCAL(grid, B) + LOCAL(grid, NE) +
- LOCAL(grid, NW) + LOCAL(grid, SE) + LOCAL(grid, SW) + LOCAL(grid, NT) +
- LOCAL(grid, NB) + LOCAL(grid, ST) + LOCAL(grid, SB) + LOCAL(grid, ET) +
- LOCAL(grid, EB) + LOCAL(grid, WT) + LOCAL(grid, WB);
-
- if (rho < minRho)
- minRho = rho;
- if (rho > maxRho)
- maxRho = rho;
- mass += rho;
-
- if (TEST_FLAG_SWEEP(grid, OBSTACLE)) {
- nObstacleCells++;
- } else {
- if (TEST_FLAG_SWEEP(grid, ACCEL))
- nAccelCells++;
- else
- nFluidCells++;
-
- ux = +LOCAL(grid, E) - LOCAL(grid, W) + LOCAL(grid, NE) - LOCAL(grid, NW) +
- LOCAL(grid, SE) - LOCAL(grid, SW) + LOCAL(grid, ET) + LOCAL(grid, EB) -
- LOCAL(grid, WT) - LOCAL(grid, WB);
- uy = +LOCAL(grid, N) - LOCAL(grid, S) + LOCAL(grid, NE) + LOCAL(grid, NW) -
- LOCAL(grid, SE) - LOCAL(grid, SW) + LOCAL(grid, NT) + LOCAL(grid, NB) -
- LOCAL(grid, ST) - LOCAL(grid, SB);
- uz = +LOCAL(grid, T) - LOCAL(grid, B) + LOCAL(grid, NT) - LOCAL(grid, NB) +
- LOCAL(grid, ST) - LOCAL(grid, SB) + LOCAL(grid, ET) - LOCAL(grid, EB) +
- LOCAL(grid, WT) - LOCAL(grid, WB);
- u2 = (ux * ux + uy * uy + uz * uz) / (rho * rho);
- if (u2 < minU2)
- minU2 = u2;
- if (u2 > maxU2)
- maxU2 = u2;
- }
- SWEEP_END
-
- printf("LBM_showGridStatistics:\n"
- "\tnObstacleCells: %7i nAccelCells: %7i nFluidCells: %7i\n"
- "\tminRho: %8.4f maxRho: %8.4f mass: %e\n"
- "\tminU: %e maxU: %e\n\n",
- nObstacleCells, nAccelCells, nFluidCells, minRho, maxRho, mass,
- sqrt(minU2), sqrt(maxU2));
-}
-
-/*############################################################################*/
-
-static void storeValue(FILE *file, OUTPUT_PRECISION *v) {
- const int litteBigEndianTest = 1;
- if ((*((unsigned char *)&litteBigEndianTest)) == 0) { /* big endian */
- const char *vPtr = (char *)v;
- char buffer[sizeof(OUTPUT_PRECISION)];
- int i;
-
- for (i = 0; i < sizeof(OUTPUT_PRECISION); i++)
- buffer[i] = vPtr[sizeof(OUTPUT_PRECISION) - i - 1];
-
- fwrite(buffer, sizeof(OUTPUT_PRECISION), 1, file);
- } else { /* little endian */
- fwrite(v, sizeof(OUTPUT_PRECISION), 1, file);
- }
-}
-
-/*############################################################################*/
-
-static void loadValue(FILE *file, OUTPUT_PRECISION *v) {
- const int litteBigEndianTest = 1;
- if ((*((unsigned char *)&litteBigEndianTest)) == 0) { /* big endian */
- char *vPtr = (char *)v;
- char buffer[sizeof(OUTPUT_PRECISION)];
- int i;
-
- fread(buffer, sizeof(OUTPUT_PRECISION), 1, file);
-
- for (i = 0; i < sizeof(OUTPUT_PRECISION); i++)
- vPtr[i] = buffer[sizeof(OUTPUT_PRECISION) - i - 1];
- } else { /* little endian */
- fread(v, sizeof(OUTPUT_PRECISION), 1, file);
- }
-}
-
-/*############################################################################*/
-
-void LBM_storeVelocityField(LBM_Grid grid, const char *filename,
- const int binary) {
- OUTPUT_PRECISION rho, ux, uy, uz;
-
- FILE *file = fopen(filename, (binary ? "wb" : "w"));
-
- SWEEP_VAR
- SWEEP_START(0, 0, 0, SIZE_X, SIZE_Y, SIZE_Z)
- rho = +SRC_C(grid) + SRC_N(grid) + SRC_S(grid) + SRC_E(grid) + SRC_W(grid) +
- SRC_T(grid) + SRC_B(grid) + SRC_NE(grid) + SRC_NW(grid) + SRC_SE(grid) +
- SRC_SW(grid) + SRC_NT(grid) + SRC_NB(grid) + SRC_ST(grid) +
- SRC_SB(grid) + SRC_ET(grid) + SRC_EB(grid) + SRC_WT(grid) +
- SRC_WB(grid);
- ux = +SRC_E(grid) - SRC_W(grid) + SRC_NE(grid) - SRC_NW(grid) + SRC_SE(grid) -
- SRC_SW(grid) + SRC_ET(grid) + SRC_EB(grid) - SRC_WT(grid) - SRC_WB(grid);
- uy = +SRC_N(grid) - SRC_S(grid) + SRC_NE(grid) + SRC_NW(grid) - SRC_SE(grid) -
- SRC_SW(grid) + SRC_NT(grid) + SRC_NB(grid) - SRC_ST(grid) - SRC_SB(grid);
- uz = +SRC_T(grid) - SRC_B(grid) + SRC_NT(grid) - SRC_NB(grid) + SRC_ST(grid) -
- SRC_SB(grid) + SRC_ET(grid) - SRC_EB(grid) + SRC_WT(grid) - SRC_WB(grid);
- ux /= rho;
- uy /= rho;
- uz /= rho;
-
- if (binary) {
- /*
- fwrite( &ux, sizeof( ux ), 1, file );
- fwrite( &uy, sizeof( uy ), 1, file );
- fwrite( &uz, sizeof( uz ), 1, file );
- */
- storeValue(file, &ux);
- storeValue(file, &uy);
- storeValue(file, &uz);
- } else
- fprintf(file, "%e %e %e\n", ux, uy, uz);
-
- SWEEP_END;
-
- fclose(file);
-}
diff --git a/hpvm/test/parboil/benchmarks/lbm/src/opencl_nvidia/lbm.h b/hpvm/test/parboil/benchmarks/lbm/src/opencl_nvidia/lbm.h
deleted file mode 100644
index 64a617feb862bdffdcb0c6aa57b0f1b09c26debb..0000000000000000000000000000000000000000
--- a/hpvm/test/parboil/benchmarks/lbm/src/opencl_nvidia/lbm.h
+++ /dev/null
@@ -1,39 +0,0 @@
-/***************************************************************************
- *cr
- *cr (C) Copyright 2010 The Board of Trustees of the
- *cr University of Illinois
- *cr All Rights Reserved
- *cr
- ***************************************************************************/
-
-/*############################################################################*/
-
-#ifndef _LBM_H_
-#define _LBM_H_
-
-/*############################################################################*/
-
-void LBM_allocateGrid(float **ptr);
-void LBM_freeGrid(float **ptr);
-void LBM_initializeGrid(LBM_Grid grid);
-void LBM_initializeSpecialCellsForLDC(LBM_Grid grid);
-void LBM_loadObstacleFile(LBM_Grid grid, const char *filename);
-void LBM_swapGrids(cl_mem *grid1, cl_mem *grid2);
-void LBM_showGridStatistics(LBM_Grid Grid);
-void LBM_storeVelocityField(LBM_Grid grid, const char *filename,
- const BOOL binary);
-
-/* OpenCL *********************************************************************/
-
-void OpenCL_LBM_allocateGrid(const OpenCL_Param *prm, cl_mem *ptr);
-void OpenCL_LBM_freeGrid(cl_mem ptr);
-void OpenCL_LBM_initializeGrid(const OpenCL_Param *prm, cl_mem d_grid,
- LBM_Grid h_grid);
-void OpenCL_LBM_getDeviceGrid(const OpenCL_Param *prm, cl_mem d_grid,
- LBM_Grid h_grid);
-void OpenCL_LBM_performStreamCollide(const OpenCL_Param *prm, cl_mem srcGrid,
- cl_mem dstGrid);
-
-/*############################################################################*/
-
-#endif /* _LBM_H_ */
diff --git a/hpvm/test/parboil/benchmarks/lbm/src/opencl_nvidia/lbm_macros.h b/hpvm/test/parboil/benchmarks/lbm/src/opencl_nvidia/lbm_macros.h
deleted file mode 100644
index 99c50c048a14bb47bb3659b61f088db95706bb0c..0000000000000000000000000000000000000000
--- a/hpvm/test/parboil/benchmarks/lbm/src/opencl_nvidia/lbm_macros.h
+++ /dev/null
@@ -1,198 +0,0 @@
-/***************************************************************************
- *cr
- *cr (C) Copyright 2010 The Board of Trustees of the
- *cr University of Illinois
- *cr All Rights Reserved
- *cr
- ***************************************************************************/
-
-#ifndef _LBM_MARCOS_H
-#define _LBM_MACROS_H_
-
-#define OMEGA (1.95f)
-
-#define OUTPUT_PRECISION float
-
-#define BOOL int
-#define TRUE (-1)
-#define FALSE (0)
-
-#define DFL1 (1.0f / 3.0f)
-#define DFL2 (1.0f / 18.0f)
-#define DFL3 (1.0f / 36.0f)
-
-/*############################################################################*/
-
-typedef float
- *LBM_Grid; // float LBM_Grid[PADDED_Z*PADDED_Y*PADDED_X*N_CELL_ENTRIES];
-typedef LBM_Grid *LBM_GridPtr;
-
-/*############################################################################*/
-
-#define SWEEP_X __temp_x__
-#define SWEEP_Y __temp_y__
-#define SWEEP_Z __temp_z__
-#define SWEEP_VAR int __temp_x__, __temp_y__, __temp_z__;
-
-#define SWEEP_START(x1, y1, z1, x2, y2, z2) \
- for (__temp_z__ = z1; __temp_z__ < z2; __temp_z__++) { \
- for (__temp_y__ = 0; __temp_y__ < SIZE_Y; __temp_y__++) { \
- for (__temp_x__ = 0; __temp_x__ < SIZE_X; __temp_x__++) {
-
-#define SWEEP_END \
- } \
- } \
- }
-
-#define GRID_ENTRY(g, x, y, z, e) ((g)[CALC_INDEX(x, y, z, e)])
-#define GRID_ENTRY_SWEEP(g, dx, dy, dz, e) \
- ((g)[CALC_INDEX((dx) + SWEEP_X, (dy) + SWEEP_Y, (dz) + SWEEP_Z, e)])
-
-#define LOCAL(g, e) (GRID_ENTRY_SWEEP(g, 0, 0, 0, e))
-#define NEIGHBOR_C(g, e) (GRID_ENTRY_SWEEP(g, 0, 0, 0, e))
-#define NEIGHBOR_N(g, e) (GRID_ENTRY_SWEEP(g, 0, +1, 0, e))
-#define NEIGHBOR_S(g, e) (GRID_ENTRY_SWEEP(g, 0, -1, 0, e))
-#define NEIGHBOR_E(g, e) (GRID_ENTRY_SWEEP(g, +1, 0, 0, e))
-#define NEIGHBOR_W(g, e) (GRID_ENTRY_SWEEP(g, -1, 0, 0, e))
-#define NEIGHBOR_T(g, e) (GRID_ENTRY_SWEEP(g, 0, 0, +1, e))
-#define NEIGHBOR_B(g, e) (GRID_ENTRY_SWEEP(g, 0, 0, -1, e))
-#define NEIGHBOR_NE(g, e) (GRID_ENTRY_SWEEP(g, +1, +1, 0, e))
-#define NEIGHBOR_NW(g, e) (GRID_ENTRY_SWEEP(g, -1, +1, 0, e))
-#define NEIGHBOR_SE(g, e) (GRID_ENTRY_SWEEP(g, +1, -1, 0, e))
-#define NEIGHBOR_SW(g, e) (GRID_ENTRY_SWEEP(g, -1, -1, 0, e))
-#define NEIGHBOR_NT(g, e) (GRID_ENTRY_SWEEP(g, 0, +1, +1, e))
-#define NEIGHBOR_NB(g, e) (GRID_ENTRY_SWEEP(g, 0, +1, -1, e))
-#define NEIGHBOR_ST(g, e) (GRID_ENTRY_SWEEP(g, 0, -1, +1, e))
-#define NEIGHBOR_SB(g, e) (GRID_ENTRY_SWEEP(g, 0, -1, -1, e))
-#define NEIGHBOR_ET(g, e) (GRID_ENTRY_SWEEP(g, +1, 0, +1, e))
-#define NEIGHBOR_EB(g, e) (GRID_ENTRY_SWEEP(g, +1, 0, -1, e))
-#define NEIGHBOR_WT(g, e) (GRID_ENTRY_SWEEP(g, -1, 0, +1, e))
-#define NEIGHBOR_WB(g, e) (GRID_ENTRY_SWEEP(g, -1, 0, -1, e))
-
-#ifdef SCATTER
-
-#define SRC_C(g) (LOCAL(g, C))
-#define SRC_N(g) (LOCAL(g, N))
-#define SRC_S(g) (LOCAL(g, S))
-#define SRC_E(g) (LOCAL(g, E))
-#define SRC_W(g) (LOCAL(g, W))
-#define SRC_T(g) (LOCAL(g, T))
-#define SRC_B(g) (LOCAL(g, B))
-#define SRC_NE(g) (LOCAL(g, NE))
-#define SRC_NW(g) (LOCAL(g, NW))
-#define SRC_SE(g) (LOCAL(g, SE))
-#define SRC_SW(g) (LOCAL(g, SW))
-#define SRC_NT(g) (LOCAL(g, NT))
-#define SRC_NB(g) (LOCAL(g, NB))
-#define SRC_ST(g) (LOCAL(g, ST))
-#define SRC_SB(g) (LOCAL(g, SB))
-#define SRC_ET(g) (LOCAL(g, ET))
-#define SRC_EB(g) (LOCAL(g, EB))
-#define SRC_WT(g) (LOCAL(g, WT))
-#define SRC_WB(g) (LOCAL(g, WB))
-
-#define DST_C(g) (NEIGHBOR_C(g, C))
-#define DST_N(g) (NEIGHBOR_N(g, N))
-#define DST_S(g) (NEIGHBOR_S(g, S))
-#define DST_E(g) (NEIGHBOR_E(g, E))
-#define DST_W(g) (NEIGHBOR_W(g, W))
-#define DST_T(g) (NEIGHBOR_T(g, T))
-#define DST_B(g) (NEIGHBOR_B(g, B))
-#define DST_NE(g) (NEIGHBOR_NE(g, NE))
-#define DST_NW(g) (NEIGHBOR_NW(g, NW))
-#define DST_SE(g) (NEIGHBOR_SE(g, SE))
-#define DST_SW(g) (NEIGHBOR_SW(g, SW))
-#define DST_NT(g) (NEIGHBOR_NT(g, NT))
-#define DST_NB(g) (NEIGHBOR_NB(g, NB))
-#define DST_ST(g) (NEIGHBOR_ST(g, ST))
-#define DST_SB(g) (NEIGHBOR_SB(g, SB))
-#define DST_ET(g) (NEIGHBOR_ET(g, ET))
-#define DST_EB(g) (NEIGHBOR_EB(g, EB))
-#define DST_WT(g) (NEIGHBOR_WT(g, WT))
-#define DST_WB(g) (NEIGHBOR_WB(g, WB))
-
-#else /* GATHER */
-
-#define SRC_C(g) (NEIGHBOR_C(g, C))
-#define SRC_N(g) (NEIGHBOR_S(g, N))
-#define SRC_S(g) (NEIGHBOR_N(g, S))
-#define SRC_E(g) (NEIGHBOR_W(g, E))
-#define SRC_W(g) (NEIGHBOR_E(g, W))
-#define SRC_T(g) (NEIGHBOR_B(g, T))
-#define SRC_B(g) (NEIGHBOR_T(g, B))
-#define SRC_NE(g) (NEIGHBOR_SW(g, NE))
-#define SRC_NW(g) (NEIGHBOR_SE(g, NW))
-#define SRC_SE(g) (NEIGHBOR_NW(g, SE))
-#define SRC_SW(g) (NEIGHBOR_NE(g, SW))
-#define SRC_NT(g) (NEIGHBOR_SB(g, NT))
-#define SRC_NB(g) (NEIGHBOR_ST(g, NB))
-#define SRC_ST(g) (NEIGHBOR_NB(g, ST))
-#define SRC_SB(g) (NEIGHBOR_NT(g, SB))
-#define SRC_ET(g) (NEIGHBOR_WB(g, ET))
-#define SRC_EB(g) (NEIGHBOR_WT(g, EB))
-#define SRC_WT(g) (NEIGHBOR_EB(g, WT))
-#define SRC_WB(g) (NEIGHBOR_ET(g, WB))
-
-#define DST_C(g) (LOCAL(g, C))
-#define DST_N(g) (LOCAL(g, N))
-#define DST_S(g) (LOCAL(g, S))
-#define DST_E(g) (LOCAL(g, E))
-#define DST_W(g) (LOCAL(g, W))
-#define DST_T(g) (LOCAL(g, T))
-#define DST_B(g) (LOCAL(g, B))
-#define DST_NE(g) (LOCAL(g, NE))
-#define DST_NW(g) (LOCAL(g, NW))
-#define DST_SE(g) (LOCAL(g, SE))
-#define DST_SW(g) (LOCAL(g, SW))
-#define DST_NT(g) (LOCAL(g, NT))
-#define DST_NB(g) (LOCAL(g, NB))
-#define DST_ST(g) (LOCAL(g, ST))
-#define DST_SB(g) (LOCAL(g, SB))
-#define DST_ET(g) (LOCAL(g, ET))
-#define DST_EB(g) (LOCAL(g, EB))
-#define DST_WT(g) (LOCAL(g, WT))
-#define DST_WB(g) (LOCAL(g, WB))
-
-#endif /* GATHER */
-
-#define MAGIC_CAST(v) ((unsigned int *)((void *)(&(v))))
-#define FLAG_VAR(v) unsigned int *_aux_ = MAGIC_CAST(v)
-
-#define TEST_FLAG_SWEEP(g, f) ((*MAGIC_CAST(LOCAL(g, FLAGS))) & (f))
-#define SET_FLAG_SWEEP(g, f) \
- { \
- FLAG_VAR(LOCAL(g, FLAGS)); \
- (*_aux_) |= (f); \
- }
-#define CLEAR_FLAG_SWEEP(g, f) \
- { \
- FLAG_VAR(LOCAL(g, FLAGS)); \
- (*_aux_) &= ~(f); \
- }
-#define CLEAR_ALL_FLAGS_SWEEP(g) \
- { \
- FLAG_VAR(LOCAL(g, FLAGS)); \
- (*_aux_) = 0; \
- }
-
-#define TEST_FLAG(g, x, y, z, f) \
- ((*MAGIC_CAST(GRID_ENTRY(g, x, y, z, FLAGS))) & (f))
-#define SET_FLAG(g, x, y, z, f) \
- { \
- FLAG_VAR(GRID_ENTRY(g, x, y, z, FLAGS)); \
- (*_aux_) |= (f); \
- }
-#define CLEAR_FLAG(g, x, y, z, f) \
- { \
- FLAG_VAR(GRID_ENTRY(g, x, y, z, FLAGS)); \
- (*_aux_) &= ~(f); \
- }
-#define CLEAR_ALL_FLAGS(g, x, y, z) \
- { \
- FLAG_VAR(GRID_ENTRY(g, x, y, z, FLAGS)); \
- (*_aux_) = 0; \
- }
-
-/*############################################################################*/
-
-#endif /* _CONFIG_H_ */
diff --git a/hpvm/test/parboil/benchmarks/lbm/src/opencl_nvidia/main.c b/hpvm/test/parboil/benchmarks/lbm/src/opencl_nvidia/main.c
deleted file mode 100644
index 18320b7394e5d499339ee820a992b00acd9b368e..0000000000000000000000000000000000000000
--- a/hpvm/test/parboil/benchmarks/lbm/src/opencl_nvidia/main.c
+++ /dev/null
@@ -1,288 +0,0 @@
-/***************************************************************************
- *cr
- *cr (C) Copyright 2010 The Board of Trustees of the
- *cr University of Illinois
- *cr All Rights Reserved
- *cr
- ***************************************************************************/
-
-/*############################################################################*/
-
-#include <CL/cl.h>
-#include <parboil.h>
-#include <stdio.h>
-#include <stdlib.h>
-#include <sys/stat.h>
-
-#include "layout_config.h"
-#include "lbm.h"
-#include "lbm_macros.h"
-#include "main.h"
-#include "ocl.h"
-
-/*############################################################################*/
-
-static cl_mem OpenCL_srcGrid, OpenCL_dstGrid;
-
-/*############################################################################*/
-
-struct pb_TimerSet timers;
-int main(int nArgs, char *arg[]) {
- MAIN_Param param;
- int t;
-
- OpenCL_Param prm;
-
- struct pb_Parameters *params;
- params = pb_ReadParameters(&nArgs, arg);
-
- // Setup TEMP datastructures
- MAIN_parseCommandLine(nArgs, arg, ¶m, params);
- MAIN_printInfo(¶m);
-
- /*MAIN_initialize( ¶m, &prm ); */ // This has been inlined
-
- static LBM_Grid TEMP_srcGrid, TEMP_dstGrid;
-
- /*pb_SwitchToTimer(&timers, pb_TimerID_COMPUTE);*/
- // Setup TEMP datastructures
- LBM_allocateGrid((float **)&TEMP_srcGrid);
- LBM_allocateGrid((float **)&TEMP_dstGrid);
- LBM_initializeGrid(TEMP_srcGrid);
- LBM_initializeGrid(TEMP_dstGrid);
-
- /*pb_SwitchToTimer(&timers, pb_TimerID_IO);*/
- if (param.obstacleFilename != NULL) {
- LBM_loadObstacleFile(TEMP_srcGrid, param.obstacleFilename);
- LBM_loadObstacleFile(TEMP_dstGrid, param.obstacleFilename);
- }
-
- /*pb_SwitchToTimer(&timers, pb_TimerID_COMPUTE);*/
- LBM_initializeSpecialCellsForLDC(TEMP_srcGrid);
- LBM_initializeSpecialCellsForLDC(TEMP_dstGrid);
-
- /*pb_SwitchToTimer(&timers, pb_TimerID_COMPUTE);*/
- LBM_showGridStatistics(TEMP_srcGrid);
-
- pb_InitializeTimerSet(&timers);
- pb_SwitchToTimer(&timers, pb_TimerID_COMPUTE);
-
- OpenCL_initialize(&prm);
-
- // Setup DEVICE datastructures
- pb_SwitchToTimer(&timers, pb_TimerID_COPY);
- OpenCL_LBM_allocateGrid(&prm, &OpenCL_srcGrid);
- OpenCL_LBM_allocateGrid(&prm, &OpenCL_dstGrid);
-
- // Initialize DEVICE datastructures
- OpenCL_LBM_initializeGrid(&prm, OpenCL_srcGrid, TEMP_srcGrid);
- OpenCL_LBM_initializeGrid(&prm, OpenCL_dstGrid, TEMP_dstGrid);
-
- for (t = 1; t <= param.nTimeSteps; t++) {
- pb_SwitchToTimer(&timers, visc_TimerID_COMPUTATION);
- OpenCL_LBM_performStreamCollide(&prm, OpenCL_srcGrid, OpenCL_dstGrid);
- pb_SwitchToTimer(&timers, pb_TimerID_COMPUTE);
- LBM_swapGrids(&OpenCL_srcGrid, &OpenCL_dstGrid);
-
- /*if( (t & 63) == 0 ) {*/
- /*printf( "timestep: %i\n", t );*/
-#if 0
- CUDA_LBM_getDeviceGrid((float**)&CUDA_srcGrid, (float**)&TEMP_srcGrid);
- LBM_showGridStatistics( *TEMP_srcGrid );
-#endif
- /*}*/
- }
-
- /*MAIN_finalize( ¶m, &prm );*/ // inlined
-
- pb_SwitchToTimer(&timers, pb_TimerID_COPY);
- OpenCL_LBM_getDeviceGrid(&prm, OpenCL_srcGrid, TEMP_srcGrid);
-
- pb_SwitchToTimer(&timers, pb_TimerID_COMPUTE);
- OpenCL_LBM_freeGrid(OpenCL_srcGrid);
- OpenCL_LBM_freeGrid(OpenCL_dstGrid);
-
- clReleaseProgram(prm.clProgram);
- clReleaseKernel(prm.clKernel);
- clReleaseCommandQueue(prm.clCommandQueue);
- clReleaseContext(prm.clContext);
-
- pb_SwitchToTimer(&timers, pb_TimerID_NONE);
- pb_PrintTimerSet(&timers);
-
- /*pb_SwitchToTimer(&timers, pb_TimerID_IO);*/
- LBM_showGridStatistics(TEMP_srcGrid);
- LBM_storeVelocityField(TEMP_srcGrid, param.resultFilename, TRUE);
-
- /*pb_SwitchToTimer(&timers, pb_TimerID_COMPUTE);*/
- LBM_freeGrid((float **)&TEMP_srcGrid);
- LBM_freeGrid((float **)&TEMP_dstGrid);
-
- pb_FreeParameters(params);
- return 0;
-}
-
-/*############################################################################*/
-
-void MAIN_parseCommandLine(int nArgs, char *arg[], MAIN_Param *param,
- struct pb_Parameters *params) {
- struct stat fileStat;
-
- if (nArgs < 2) {
- printf("syntax: lbm <time steps>\n");
- exit(1);
- }
-
- param->nTimeSteps = atoi(arg[1]);
-
- if (params->inpFiles[0] != NULL) {
- param->obstacleFilename = params->inpFiles[0];
-
- if (stat(param->obstacleFilename, &fileStat) != 0) {
- printf("MAIN_parseCommandLine: cannot stat obstacle file '%s'\n",
- param->obstacleFilename);
- exit(1);
- }
- if (fileStat.st_size != SIZE_X * SIZE_Y * SIZE_Z + (SIZE_Y + 1) * SIZE_Z) {
- printf("MAIN_parseCommandLine:\n"
- "\tsize of file '%s' is %i bytes\n"
- "\texpected size is %i bytes\n",
- param->obstacleFilename, (int)fileStat.st_size,
- SIZE_X * SIZE_Y * SIZE_Z + (SIZE_Y + 1) * SIZE_Z);
- exit(1);
- }
- } else
- param->obstacleFilename = NULL;
-
- param->resultFilename = params->outFile;
-}
-
-/*############################################################################*/
-
-void MAIN_printInfo(const MAIN_Param *param) {
- printf("MAIN_printInfo:\n"
- "\tgrid size : %i x %i x %i = %.2f * 10^6 Cells\n"
- "\tnTimeSteps : %i\n"
- "\tresult file : %s\n"
- "\taction : %s\n"
- "\tsimulation type: %s\n"
- "\tobstacle file : %s\n\n",
- SIZE_X, SIZE_Y, SIZE_Z, 1e-6 * SIZE_X * SIZE_Y * SIZE_Z,
- param->nTimeSteps, param->resultFilename, "store", "lid-driven cavity",
- (param->obstacleFilename == NULL) ? "<none>"
- : param->obstacleFilename);
-}
-
-/*############################################################################*/
-
-void MAIN_initialize(const MAIN_Param *param, const OpenCL_Param *prm) {
- static LBM_Grid TEMP_srcGrid, TEMP_dstGrid;
-
- pb_SwitchToTimer(&timers, pb_TimerID_COMPUTE);
- // Setup TEMP datastructures
- LBM_allocateGrid((float **)&TEMP_srcGrid);
- LBM_allocateGrid((float **)&TEMP_dstGrid);
- LBM_initializeGrid(TEMP_srcGrid);
- LBM_initializeGrid(TEMP_dstGrid);
-
- pb_SwitchToTimer(&timers, pb_TimerID_IO);
- if (param->obstacleFilename != NULL) {
- LBM_loadObstacleFile(TEMP_srcGrid, param->obstacleFilename);
- LBM_loadObstacleFile(TEMP_dstGrid, param->obstacleFilename);
- }
-
- pb_SwitchToTimer(&timers, pb_TimerID_COMPUTE);
- LBM_initializeSpecialCellsForLDC(TEMP_srcGrid);
- LBM_initializeSpecialCellsForLDC(TEMP_dstGrid);
-
- // Setup DEVICE datastructures
- OpenCL_LBM_allocateGrid(prm, &OpenCL_srcGrid);
- OpenCL_LBM_allocateGrid(prm, &OpenCL_dstGrid);
-
- // Initialize DEVICE datastructures
- pb_SwitchToTimer(&timers, pb_TimerID_COPY);
- OpenCL_LBM_initializeGrid(prm, OpenCL_srcGrid, TEMP_srcGrid);
- OpenCL_LBM_initializeGrid(prm, OpenCL_dstGrid, TEMP_dstGrid);
-
- pb_SwitchToTimer(&timers, pb_TimerID_COMPUTE);
- LBM_showGridStatistics(TEMP_srcGrid);
-
- LBM_freeGrid((float **)&TEMP_srcGrid);
- LBM_freeGrid((float **)&TEMP_dstGrid);
-}
-
-/*############################################################################*/
-
-void MAIN_finalize(const MAIN_Param *param, const OpenCL_Param *prm) {
- pb_SwitchToTimer(&timers, pb_TimerID_COMPUTE);
- LBM_Grid TEMP_srcGrid;
-
- // Setup TEMP datastructures
- LBM_allocateGrid((float **)&TEMP_srcGrid);
-
- pb_SwitchToTimer(&timers, pb_TimerID_COPY);
- OpenCL_LBM_getDeviceGrid(prm, OpenCL_srcGrid, TEMP_srcGrid);
-
- pb_SwitchToTimer(&timers, pb_TimerID_COMPUTE);
- LBM_showGridStatistics(TEMP_srcGrid);
-
- LBM_storeVelocityField(TEMP_srcGrid, param->resultFilename, TRUE);
-
- LBM_freeGrid((float **)&TEMP_srcGrid);
-
- pb_SwitchToTimer(&timers, pb_TimerID_COMPUTE);
- OpenCL_LBM_freeGrid(OpenCL_srcGrid);
- OpenCL_LBM_freeGrid(OpenCL_dstGrid);
-
- clReleaseProgram(prm->clProgram);
- clReleaseKernel(prm->clKernel);
- clReleaseCommandQueue(prm->clCommandQueue);
- clReleaseContext(prm->clContext);
-}
-
-void OpenCL_initialize(OpenCL_Param *prm) {
- cl_int clStatus;
-
- clStatus = clGetPlatformIDs(1, &(prm->clPlatform), NULL);
- CHECK_ERROR("clGetPlatformIDs")
-
- prm->clCps[0] = CL_CONTEXT_PLATFORM;
- prm->clCps[1] = (cl_context_properties)(prm->clPlatform);
- prm->clCps[2] = 0;
-
- clStatus = clGetDeviceIDs(prm->clPlatform, CL_DEVICE_TYPE_GPU, 1,
- &(prm->clDevice), NULL);
- CHECK_ERROR("clGetDeviceIDs")
-
- prm->clContext = clCreateContextFromType(prm->clCps, CL_DEVICE_TYPE_GPU, NULL,
- NULL, &clStatus);
- CHECK_ERROR("clCreateContextFromType")
-
- prm->clCommandQueue = clCreateCommandQueue(
- prm->clContext, prm->clDevice, CL_QUEUE_PROFILING_ENABLE, &clStatus);
- CHECK_ERROR("clCreateCommandQueue")
-
- pb_SetOpenCL(&(prm->clContext), &(prm->clCommandQueue));
-
- const char *clSource[] = {readFile("src/opencl_nvidia/kernel.cl")};
- prm->clProgram =
- clCreateProgramWithSource(prm->clContext, 1, clSource, NULL, &clStatus);
- CHECK_ERROR("clCreateProgramWithSource")
-
- char clOptions[100];
- sprintf(clOptions, "-I src/opencl_nvidia");
-
- clStatus = clBuildProgram(prm->clProgram, 1, &(prm->clDevice), clOptions,
- NULL, NULL);
- CHECK_ERROR("clBuildProgram")
-
- prm->clKernel =
- clCreateKernel(prm->clProgram, "performStreamCollide_kernel", &clStatus);
- CHECK_ERROR("clCreateKernel")
-
- free((void *)clSource[0]);
-
- /*pb_CreateAndBuildKernelFromBinary("build/opencl_nvidia_default/kernel_offline.nvptx.s",
- * "performStreamCollide_kernel", &prm->clContext, &prm->clDevice,
- * &prm->clProgram, &prm->clKernel);*/
-}
diff --git a/hpvm/test/parboil/benchmarks/lbm/src/opencl_nvidia/main.h b/hpvm/test/parboil/benchmarks/lbm/src/opencl_nvidia/main.h
deleted file mode 100644
index 9d8e145c93b37488a3826e77b964c56699377d2a..0000000000000000000000000000000000000000
--- a/hpvm/test/parboil/benchmarks/lbm/src/opencl_nvidia/main.h
+++ /dev/null
@@ -1,32 +0,0 @@
-/***************************************************************************
- *cr
- *cr (C) Copyright 2010 The Board of Trustees of the
- *cr University of Illinois
- *cr All Rights Reserved
- *cr
- ***************************************************************************/
-
-#ifndef _MAIN_H_
-#define _MAIN_H_
-
-/*############################################################################*/
-
-typedef struct {
- int nTimeSteps;
- char *resultFilename;
- char *obstacleFilename;
-} MAIN_Param;
-
-/*############################################################################*/
-
-void MAIN_parseCommandLine(int nArgs, char *arg[], MAIN_Param *param,
- struct pb_Parameters *);
-void MAIN_printInfo(const MAIN_Param *param);
-void MAIN_initialize(const MAIN_Param *param, const OpenCL_Param *prm);
-void MAIN_finalize(const MAIN_Param *param, const OpenCL_Param *prm);
-
-void OpenCL_initialize(OpenCL_Param *prm);
-
-/*############################################################################*/
-
-#endif /* _MAIN_H_ */
diff --git a/hpvm/test/parboil/benchmarks/lbm/src/opencl_nvidia/ocl.c b/hpvm/test/parboil/benchmarks/lbm/src/opencl_nvidia/ocl.c
deleted file mode 100644
index 4f232db0d9776f4f2d0eb4b2444036f35ff27257..0000000000000000000000000000000000000000
--- a/hpvm/test/parboil/benchmarks/lbm/src/opencl_nvidia/ocl.c
+++ /dev/null
@@ -1,36 +0,0 @@
-#include "ocl.h"
-#include <CL/cl.h>
-#include <stdio.h>
-#include <stdlib.h>
-
-char *readFile(char *fileName) {
- FILE *fp;
- fp = fopen(fileName, "r");
-
- if (fp == NULL) {
- printf("Error 1!\n");
- return NULL;
- }
-
- fseek(fp, 0, SEEK_END);
- long size = ftell(fp);
- rewind(fp);
-
- char *buffer = malloc(sizeof(char) * (size + 1));
- if (buffer == NULL) {
- printf("Error 2!\n");
- fclose(fp);
- return NULL;
- }
-
- size_t res = fread(buffer, 1, size, fp);
- if (res != size) {
- printf("Error 3!\n");
- fclose(fp);
- return NULL;
- }
-
- buffer[size] = 0;
- fclose(fp);
- return buffer;
-}
diff --git a/hpvm/test/parboil/benchmarks/lbm/src/opencl_nvidia/ocl.h b/hpvm/test/parboil/benchmarks/lbm/src/opencl_nvidia/ocl.h
deleted file mode 100644
index 5d5d984ba698d6ac71af3e51de3e6724a79135aa..0000000000000000000000000000000000000000
--- a/hpvm/test/parboil/benchmarks/lbm/src/opencl_nvidia/ocl.h
+++ /dev/null
@@ -1,23 +0,0 @@
-#ifndef __OCLH__
-#define __OCLH__
-
-typedef struct {
- cl_platform_id clPlatform;
- cl_context_properties clCps[3];
- cl_device_id clDevice;
- cl_context clContext;
- cl_command_queue clCommandQueue;
- cl_program clProgram;
- cl_kernel clKernel;
-} OpenCL_Param;
-
-#define CHECK_ERROR(errorMessage) \
- if (clStatus != CL_SUCCESS) { \
- printf("Error: %s!\n", errorMessage); \
- printf("Line: %d\n", __LINE__); \
- exit(1); \
- }
-
-char *readFile(char *);
-
-#endif
diff --git a/hpvm/test/parboil/benchmarks/sgemm/src/opencl_base/Makefile b/hpvm/test/parboil/benchmarks/sgemm/src/opencl_base/Makefile
deleted file mode 100644
index 36b421ec6f1359114ea0035d21048ab0b95bf30e..0000000000000000000000000000000000000000
--- a/hpvm/test/parboil/benchmarks/sgemm/src/opencl_base/Makefile
+++ /dev/null
@@ -1,8 +0,0 @@
-# (c) 2010 The Board of Trustees of the University of Illinois.
-
-LANGUAGE=opencl
-SRCDIR_OBJS=main.o io.o #compute_gold.o
-APP_CUDALDFLAGS=-lm -lstdc++
-APP_CFLAGS=-ffast-math -O3
-APP_CXXFLAGS=-ffast-math -O3
-KERNEL_OBJS=kernel_offline.nvptx.s
diff --git a/hpvm/test/parboil/benchmarks/sgemm/src/opencl_base/io.cc b/hpvm/test/parboil/benchmarks/sgemm/src/opencl_base/io.cc
deleted file mode 100644
index 04744f404ebaf6e669c2bbe91600519742b57dc9..0000000000000000000000000000000000000000
--- a/hpvm/test/parboil/benchmarks/sgemm/src/opencl_base/io.cc
+++ /dev/null
@@ -1,84 +0,0 @@
-/***************************************************************************
- *cr
- *cr (C) Copyright 2010 The Board of Trustees of the
- *cr University of Illinois
- *cr All Rights Reserved
- *cr
- ***************************************************************************/
-
-/* I/O routines for reading and writing matrices in column-major
- * layout
- */
-
-#include <fstream>
-#include <iostream>
-#include <vector>
-
-char *readFile(const char *fileName) {
- std::fstream f(fileName, std::fstream::in);
- if (!f.good()) {
- std::cerr << "Error Reading File!!" << std::endl;
- return NULL;
- }
-
- f.seekg(0, std::ios::end);
- int length = f.tellg();
- f.seekg(0, std::ios::beg);
-
- char *buffer;
-
- if (length > 0) {
- buffer = new char[length];
- f.read(buffer, length);
- buffer[length - 1] = 0;
- } else {
- buffer = new char;
- buffer[0] = 0;
- }
-
- f.close();
-
- return buffer;
-}
-
-bool readColMajorMatrixFile(const char *fn, int &nr_row, int &nr_col,
- std::vector<float> &v) {
- std::cerr << "Opening file:" << fn << std::endl;
- std::fstream f(fn, std::fstream::in);
- if (!f.good()) {
- return false;
- }
-
- // Read # of rows and cols
- f >> nr_row;
- f >> nr_col;
-
- float data;
- std::cerr << "Matrix dimension: " << nr_row << "x" << nr_col << std::endl;
- while (f.good()) {
- f >> data;
- v.push_back(data);
- }
- v.pop_back(); // remove the duplicated last element
- return true;
-}
-
-bool writeColMajorMatrixFile(const char *fn, int nr_row, int nr_col,
- std::vector<float> &v) {
- std::cerr << "Opening file:" << fn << " for write." << std::endl;
- std::fstream f(fn, std::fstream::out);
- if (!f.good()) {
- return false;
- }
-
- // Read # of rows and cols
- f << nr_row << " " << nr_col << " ";
-
- float data;
- std::cerr << "Matrix dimension: " << nr_row << "x" << nr_col << std::endl;
- for (int i = 0; i < v.size(); ++i) {
- f << v[i] << ' ';
- }
- f << "\n";
- return true;
-}
diff --git a/hpvm/test/parboil/benchmarks/sgemm/src/opencl_base/kernel.cl b/hpvm/test/parboil/benchmarks/sgemm/src/opencl_base/kernel.cl
deleted file mode 100644
index 5ee2fd5df0a57bc59c1f714bc3efb6b3670b0386..0000000000000000000000000000000000000000
--- a/hpvm/test/parboil/benchmarks/sgemm/src/opencl_base/kernel.cl
+++ /dev/null
@@ -1,25 +0,0 @@
-/***************************************************************************
- *cr
- *cr (C) Copyright 2010 The Board of Trustees of the
- *cr University of Illinois
- *cr All Rights Reserved
- *cr
- ***************************************************************************/
-
-/*
- * Kernel of dense matrix-matrix multiplication kernel.
- */
-
-__kernel void mysgemmNT( __global const float *A, int lda, __global const float *B, int ldb, __global float* C, int ldc, int k, float alpha, float beta )
-{
- float c = 0.0f;
- int m = get_global_id(0);
- int n = get_global_id(1);
-
- /*for (int i = 0; i < k; ++i) {*/
- /*float a = A[m + i * lda]; */
- /*float b = B[n + i * ldb];*/
- /*c += a * b;*/
- /*}*/
- C[m+n*ldc] = beta + alpha * c;
-}
diff --git a/hpvm/test/parboil/benchmarks/sgemm/src/opencl_base/kernel_offline.cl b/hpvm/test/parboil/benchmarks/sgemm/src/opencl_base/kernel_offline.cl
deleted file mode 100644
index f376a27d90003e3c7c18dafb9f64a8b459a40029..0000000000000000000000000000000000000000
--- a/hpvm/test/parboil/benchmarks/sgemm/src/opencl_base/kernel_offline.cl
+++ /dev/null
@@ -1,25 +0,0 @@
-/***************************************************************************
- *cr
- *cr (C) Copyright 2010 The Board of Trustees of the
- *cr University of Illinois
- *cr All Rights Reserved
- *cr
- ***************************************************************************/
-
-/*
- * Kernel of dense matrix-matrix multiplication kernel.
- */
-
-__kernel void mysgemmNT( __global const float *A, int lda, __global const float *B, int ldb, __global float* C, int ldc, int k, float alpha, float beta )
-{
- float c = 0.0f;
- int m = get_global_id(0);
- int n = get_global_id(1);
-
- for (int i = 0; i < k; ++i) {
- float a = A[m + i * lda];
- float b = B[n + i * ldb];
- c += a * b;
- }
- C[m+n*ldc] = C[m+n*ldc] * beta + alpha * c;
-}
diff --git a/hpvm/test/parboil/benchmarks/sgemm/src/opencl_base/main.cc b/hpvm/test/parboil/benchmarks/sgemm/src/opencl_base/main.cc
deleted file mode 100644
index 5489f6a55ce6e8ba3676b0c98ad4b37ac7f4a7fd..0000000000000000000000000000000000000000
--- a/hpvm/test/parboil/benchmarks/sgemm/src/opencl_base/main.cc
+++ /dev/null
@@ -1,242 +0,0 @@
-/***************************************************************************
- *cr
- *cr (C) Copyright 2010 The Board of Trustees of the
- *cr University of Illinois
- *cr All Rights Reserved
- *cr
- ***************************************************************************/
-
-/*
- * Main entry of dense matrix-matrix multiplication kernel
- */
-
-#include <CL/cl.h>
-#include <iostream>
-#include <malloc.h>
-#include <math.h>
-#include <parboil.h>
-#include <stdio.h>
-#include <stdlib.h>
-#include <string.h>
-#include <sys/time.h>
-#include <vector>
-
-// I/O routines
-extern bool readColMajorMatrixFile(const char *fn, int &nr_row, int &nr_col,
- std::vector<float> &v);
-extern bool writeColMajorMatrixFile(const char *fn, int, int,
- std::vector<float> &);
-extern char *readFile(const char *);
-
-// Parameters of tile sizes
-#define TILE_SZ 16
-
-#define CHECK_ERROR(errorMessage) \
- if (clStatus != CL_SUCCESS) { \
- std::cout << errorMessage << ": " << clStatus << " Error!\n"; \
- std::cout << "Line: " << __LINE__ << "\n"; \
- exit(1); \
- }
-
-void basicSgemm(char transa, char transb, int m, int n, int k, float alpha,
- cl_mem A, int lda, cl_mem B, int ldb, float beta, cl_mem C,
- int ldc, cl_kernel clKernel, cl_command_queue clCommandQueue) {
- if ((transa != 'N') && (transa != 'n')) {
- std::cerr << "unsupported value of 'transa' in regtileSgemm()" << std::endl;
- return;
- }
-
- if ((transb != 'T') && (transb != 't')) {
- std::cerr << "unsupported value of 'transb' in regtileSgemm()" << std::endl;
- return;
- }
-
- // In this code we assume the matrix sizes are multiple of tile size
- if ((m % TILE_SZ) || (n % TILE_SZ)) {
- std::cerr << "unsupported size of matrix. m should be multiple of "
- << TILE_SZ << "; n should be multiple of " << TILE_SZ
- << std::endl;
- }
-
- size_t db[2] = {TILE_SZ, TILE_SZ};
- size_t dg[2] = {m / TILE_SZ * db[0], n / TILE_SZ * db[1]};
-
- cl_int clStatus;
-
- clStatus = clSetKernelArg(clKernel, 0, sizeof(cl_mem), (void *)&A);
- clStatus = clSetKernelArg(clKernel, 1, sizeof(int), (void *)&lda);
- clStatus = clSetKernelArg(clKernel, 2, sizeof(cl_mem), (void *)&B);
- clStatus = clSetKernelArg(clKernel, 3, sizeof(int), (void *)&ldb);
- clStatus = clSetKernelArg(clKernel, 4, sizeof(cl_mem), (void *)&C);
- clStatus = clSetKernelArg(clKernel, 5, sizeof(int), (void *)&ldc);
- clStatus = clSetKernelArg(clKernel, 6, sizeof(int), (void *)&k);
- clStatus = clSetKernelArg(clKernel, 7, sizeof(float), (void *)&alpha);
- clStatus = clSetKernelArg(clKernel, 8, sizeof(float), (void *)&beta);
- CHECK_ERROR("clSetKernelArg")
-
- clStatus = clEnqueueNDRangeKernel(clCommandQueue, clKernel, 2, NULL, dg, db,
- 0, NULL, NULL);
- CHECK_ERROR("clEnqueueNDRangeKernel")
-
- clStatus = clFinish(clCommandQueue);
- CHECK_ERROR("clFinish")
-}
-
-int main(int argc, char *argv[]) {
-
- struct pb_Parameters *params;
- struct pb_TimerSet timers;
-
- size_t A_sz, B_sz, C_sz;
- int matArow, matAcol;
- int matBrow, matBcol;
- std::vector<float> matA, matBT;
-
- /* Read command line. Expect 3 inputs: A, B and B^T
- in column-major layout*/
- params = pb_ReadParameters(&argc, argv);
- if ((params->inpFiles[0] == NULL) || (params->inpFiles[1] == NULL) ||
- (params->inpFiles[2] == NULL) || (params->inpFiles[3] != NULL)) {
- fprintf(stderr, "Expecting three input filenames\n");
- exit(-1);
- }
-
- /* Read in data */
- // load A
- readColMajorMatrixFile(params->inpFiles[0], matArow, matAcol, matA);
- // load B^T
- readColMajorMatrixFile(params->inpFiles[2], matBcol, matBrow, matBT);
-
- pb_InitializeTimerSet(&timers);
-
- pb_SwitchToTimer(&timers, visc_TimerID_SETUP);
- cl_int clStatus;
- cl_platform_id clPlatform;
- clStatus = clGetPlatformIDs(1, &clPlatform, NULL);
- CHECK_ERROR("clGetPlatformIDs")
-
- cl_context_properties clCps[3] = {CL_CONTEXT_PLATFORM,
- (cl_context_properties)clPlatform, 0};
- cl_context clContext =
- clCreateContextFromType(clCps, CL_DEVICE_TYPE_GPU, NULL, NULL, &clStatus);
- CHECK_ERROR("clCreateContextFromType")
-
- cl_device_id clDevice;
- clStatus = clGetDeviceIDs(clPlatform, CL_DEVICE_TYPE_GPU, 1, &clDevice, NULL);
- CHECK_ERROR("clGetDeviceIDs")
-
- cl_command_queue clCommandQueue = clCreateCommandQueue(
- clContext, clDevice, CL_QUEUE_PROFILING_ENABLE, &clStatus);
- CHECK_ERROR("clCreateCommandQueue")
-
- pb_SetOpenCL(&clContext, &clCommandQueue);
-
- const char *clSource[] = {readFile("src/opencl_base/kernel_offline.nvptx.s")};
- // const char* clSource[] = {readFile("src/opencl_base/kernel.cl")};
- // cl_program clProgram =
- // clCreateProgramWithSource(clContext,1,clSource,NULL,&clStatus);
- cl_kernel clKernel;
- cl_program clProgram;
- pb_CreateAndBuildKernelFromBinary(
- "build/opencl_base_default/kernel_offline.nvptx.s", "mysgemmNT",
- &clContext, &clDevice, &clProgram, &clKernel);
- // cl_program clProgram =
- // clCreateProgramWithSource(clContext,1,clSource,NULL,&clStatus);
- // CHECK_ERROR("clCreateProgramWithSource")
-
- // char clOptions[50];
- // sprintf(clOptions,"");
-
- // clStatus = clBuildProgram(clProgram,1,&clDevice,clOptions,NULL,NULL);
- // CHECK_ERROR("clBuildProgram")
-
- // size_t binarySizes = 0;
- // clStatus = clGetProgramInfo(clProgram, CL_PROGRAM_BINARY_SIZES,
- // sizeof(size_t), &binarySizes, NULL); CHECK_ERROR("clGetProgramInfo")
-
- // std::cout << "Binary Size = " << binarySizes << "\n";
-
- // unsigned char* binaries = (unsigned char*) malloc(binarySizes);
- // clStatus = clGetProgramInfo(clProgram, CL_PROGRAM_BINARIES, binarySizes,
- // &binaries, NULL); CHECK_ERROR("clGetProgramInfo")
-
- // std::cout << "Binary = \n" << binaries << "\n";
-
- // cl_kernel clKernel = clCreateKernel(clProgram,"mysgemmNT",&clStatus);
- // CHECK_ERROR("clCreateKernel")
-
- pb_SwitchToTimer(&timers, pb_TimerID_COMPUTE);
- // copy A to device memory
- A_sz = matArow * matAcol * sizeof(float);
- B_sz = matBrow * matBcol * sizeof(float);
-
- // allocate space for C
- C_sz = matArow * matBcol * sizeof(float);
-
- // OpenCL memory allocation
- std::vector<float> matC(matArow * matBcol);
-
- pb_SwitchToTimer(&timers, pb_TimerID_COPY);
- cl_mem dA =
- clCreateBuffer(clContext, CL_MEM_READ_ONLY, A_sz, NULL, &clStatus);
- CHECK_ERROR("clCreateBuffer")
- cl_mem dB =
- clCreateBuffer(clContext, CL_MEM_READ_ONLY, B_sz, NULL, &clStatus);
- CHECK_ERROR("clCreateBuffer")
- cl_mem dC =
- clCreateBuffer(clContext, CL_MEM_WRITE_ONLY, C_sz, NULL, &clStatus);
- CHECK_ERROR("clCreateBuffer")
-
- // Copy A and B^T into device memory
- clStatus = clEnqueueWriteBuffer(clCommandQueue, dA, CL_FALSE, 0, A_sz,
- &matA.front(), 0, NULL, NULL);
- CHECK_ERROR("clEnqueueWriteBuffer")
- clStatus = clEnqueueWriteBuffer(clCommandQueue, dB, CL_FALSE, 0, B_sz,
- &matBT.front(), 0, NULL, NULL);
- CHECK_ERROR("clEnqueueWriteBuffer")
-
- for (int i = 0; i < matC.size(); i++)
- matC[i] = 0.0f;
-
- clStatus = clEnqueueWriteBuffer(clCommandQueue, dC, CL_TRUE, 0, C_sz,
- &matC.front(), 0, NULL, NULL);
- CHECK_ERROR("clEnqueueWriteBuffer")
-
- pb_SwitchToTimer(&timers, pb_TimerID_KERNEL);
-
- // Use standard sgemm interface
- basicSgemm('N', 'T', matArow, matBcol, matAcol, 1.0f, dA, matArow, dB,
- matBcol, 0.0f, dC, matArow, clKernel, clCommandQueue);
-
- pb_SwitchToTimer(&timers, pb_TimerID_COPY);
- clEnqueueReadBuffer(clCommandQueue, dC, CL_TRUE, 0, C_sz, &matC.front(), 0,
- NULL, NULL);
-
- pb_SwitchToTimer(&timers, visc_TimerID_SETUP);
- clStatus = clReleaseKernel(clKernel);
- clStatus = clReleaseProgram(clProgram);
- clStatus = clReleaseMemObject(dA);
- clStatus = clReleaseMemObject(dB);
- clStatus = clReleaseMemObject(dC);
- clStatus = clReleaseCommandQueue(clCommandQueue);
- clStatus = clReleaseContext(clContext);
-
- pb_SwitchToTimer(&timers, pb_TimerID_NONE);
- pb_PrintTimerSet(&timers);
-
- if (params->outFile) {
-
- /* Write C to file */
- // pb_SwitchToTimer(&timers, pb_TimerID_IO);
- writeColMajorMatrixFile(params->outFile, matArow, matBcol, matC);
- }
-
- double GPUtime = pb_GetElapsedTime(&(timers.timers[pb_TimerID_KERNEL]));
- std::cout << "GFLOPs = " << 2. * matArow * matBcol * matAcol / GPUtime / 1e9
- << std::endl;
- pb_FreeParameters(params);
-
- // free((void*)clSource[0]);
-
- return 0;
-}
diff --git a/hpvm/test/parboil/benchmarks/spmv/src/opencl_nvidia/Makefile b/hpvm/test/parboil/benchmarks/spmv/src/opencl_nvidia/Makefile
deleted file mode 100644
index 5774b2827bd0588b16cfc32b2ac787346c5b7b1d..0000000000000000000000000000000000000000
--- a/hpvm/test/parboil/benchmarks/spmv/src/opencl_nvidia/Makefile
+++ /dev/null
@@ -1,11 +0,0 @@
-# (c) 2010 The Board of Trustees of the University of Illinois.
-
-LANGUAGE=opencl
-TOOLS_SRC=common_src/convert-dataset
-SRCDIR_OBJS=main.o gpu_info.o file.o ocl.o
-APP_CUDALDFLAGS=-lm
-APP_CFLAGS=-ffast-math -O3 -I$(TOOLS_SRC)
-APP_CXXFLAGS=-ffast-math -O3
-KERNEL_OBJS=kernel_offline.nvptx.s
-
-include $(TOOLS_SRC)/commontools.mk
diff --git a/hpvm/test/parboil/benchmarks/spmv/src/opencl_nvidia/file.c b/hpvm/test/parboil/benchmarks/spmv/src/opencl_nvidia/file.c
deleted file mode 100644
index 22397498f7f43b3f60926bf51c2ddbff91529787..0000000000000000000000000000000000000000
--- a/hpvm/test/parboil/benchmarks/spmv/src/opencl_nvidia/file.c
+++ /dev/null
@@ -1,72 +0,0 @@
-/***************************************************************************
- *cr
- *cr (C) Copyright 2007 The Board of Trustees of the
- *cr University of Illinois
- *cr All Rights Reserved
- *cr
- ***************************************************************************/
-
-#include <endian.h>
-#include <inttypes.h>
-#include <malloc.h>
-#include <stdio.h>
-#include <stdlib.h>
-
-#if __BYTE_ORDER != __LITTLE_ENDIAN
-#error "File I/O is not implemented for this system: wrong endianness."
-#endif
-
-void inputData(char *fName, int *len, int *depth, int *dim, int *nzcnt_len,
- int *pad, float **h_data, int **h_indices, int **h_ptr,
- int **h_perm, int **h_nzcnt) {
- FILE *fid = fopen(fName, "rb");
-
- if (fid == NULL) {
- fprintf(stderr, "Cannot open input file\n");
- exit(-1);
- }
-
- fscanf(fid, "%d %d %d %d %d\n", len, depth, nzcnt_len, dim, pad);
- int _len = len[0];
- int _depth = depth[0];
- int _dim = dim[0];
- int _pad = pad[0];
- int _nzcnt_len = nzcnt_len[0];
-
- *h_data = (float *)malloc(_len * sizeof(float));
- fread(*h_data, sizeof(float), _len, fid);
-
- *h_indices = (int *)malloc(_len * sizeof(int));
- fread(*h_indices, sizeof(int), _len, fid);
-
- *h_ptr = (int *)malloc(_depth * sizeof(int));
- fread(*h_ptr, sizeof(int), _depth, fid);
-
- *h_perm = (int *)malloc(_dim * sizeof(int));
- fread(*h_perm, sizeof(int), _dim, fid);
-
- *h_nzcnt = (int *)malloc(_nzcnt_len * sizeof(int));
- fread(*h_nzcnt, sizeof(int), _nzcnt_len, fid);
-
- fclose(fid);
-}
-
-void input_vec(char *fName, float *h_vec, int dim) {
- FILE *fid = fopen(fName, "rb");
- fread(h_vec, sizeof(float), dim, fid);
- fclose(fid);
-}
-
-void outputData(char *fName, float *h_Ax_vector, int dim) {
- FILE *fid = fopen(fName, "w");
- uint32_t tmp32;
- if (fid == NULL) {
- fprintf(stderr, "Cannot open output file\n");
- exit(-1);
- }
- tmp32 = dim;
- fwrite(&tmp32, sizeof(uint32_t), 1, fid);
- fwrite(h_Ax_vector, sizeof(float), dim, fid);
-
- fclose(fid);
-}
diff --git a/hpvm/test/parboil/benchmarks/spmv/src/opencl_nvidia/file.h b/hpvm/test/parboil/benchmarks/spmv/src/opencl_nvidia/file.h
deleted file mode 100644
index 5e38a6875e9e5f8be4d01b68569d80adf8c49548..0000000000000000000000000000000000000000
--- a/hpvm/test/parboil/benchmarks/spmv/src/opencl_nvidia/file.h
+++ /dev/null
@@ -1,18 +0,0 @@
-/***************************************************************************
- *cr
- *cr (C) Copyright 2007 The Board of Trustees of the
- *cr University of Illinois
- *cr All Rights Reserved
- *cr
- ***************************************************************************/
-#ifndef __FILEH__
-#define __FILEH__
-
-void inputData(char *fName, int *len, int *depth, int *dim, int *nzcnt_len,
- int *pad, float **h_data, int **h_indices, int **h_ptr,
- int **h_perm, int **h_nzcnt);
-
-void input_vec(char *fNanme, float *h_vec, int dim);
-void outputData(char *fName, float *h_Ax_vector, int dim);
-
-#endif
\ No newline at end of file
diff --git a/hpvm/test/parboil/benchmarks/spmv/src/opencl_nvidia/gpu_info.c b/hpvm/test/parboil/benchmarks/spmv/src/opencl_nvidia/gpu_info.c
deleted file mode 100644
index 90beedd747480ede3fd1e5da4017ed0051e043be..0000000000000000000000000000000000000000
--- a/hpvm/test/parboil/benchmarks/spmv/src/opencl_nvidia/gpu_info.c
+++ /dev/null
@@ -1,44 +0,0 @@
-/***************************************************************************
- *cr
- *cr (C) Copyright 2010 The Board of Trustees of the
- *cr University of Illinois
- *cr All Rights Reserved
- *cr
- ***************************************************************************/
-#include <endian.h>
-#include <inttypes.h>
-#include <malloc.h>
-#include <stdio.h>
-#include <stdlib.h>
-
-#include "gpu_info.h"
-
-void compute_active_thread(size_t *thread, size_t *grid, int task, int pad,
- int major, int minor, int sm) {
- int max_thread;
- int max_block = 8;
- if (major == 1) {
- if (minor >= 2)
- max_thread = 1024;
- else
- max_thread = 768;
- } else if (major == 2)
- max_thread = 1536;
- else
- // newer GPU //keep using 2.0
- max_thread = 1536;
-
- int _grid;
- int _thread;
-
- if (task * pad > sm * max_thread) {
- _thread = max_thread / max_block;
- _grid = ((task * pad + _thread - 1) / _thread) * _thread;
- } else {
- _thread = pad;
- _grid = task * pad;
- }
-
- thread[0] = _thread;
- grid[0] = _grid;
-}
diff --git a/hpvm/test/parboil/benchmarks/spmv/src/opencl_nvidia/gpu_info.h b/hpvm/test/parboil/benchmarks/spmv/src/opencl_nvidia/gpu_info.h
deleted file mode 100644
index ab1af7d0b8ba92f87c643582171e48cee0a9b95e..0000000000000000000000000000000000000000
--- a/hpvm/test/parboil/benchmarks/spmv/src/opencl_nvidia/gpu_info.h
+++ /dev/null
@@ -1,15 +0,0 @@
-/***************************************************************************
- *cr
- *cr (C) Copyright 2010 The Board of Trustees of the
- *cr University of Illinois
- *cr All Rights Reserved
- *cr
- ***************************************************************************/
-
-#ifndef __GPUINFOH__
-#define __GPUINFOH__
-
-void compute_active_thread(size_t *thread, size_t *grid, int task, int pad,
- int major, int minor, int sm);
-
-#endif
diff --git a/hpvm/test/parboil/benchmarks/spmv/src/opencl_nvidia/kernel.cl b/hpvm/test/parboil/benchmarks/spmv/src/opencl_nvidia/kernel.cl
deleted file mode 100644
index 8c19a22511cd4c65364cb55c5df1a22d12b182b7..0000000000000000000000000000000000000000
--- a/hpvm/test/parboil/benchmarks/spmv/src/opencl_nvidia/kernel.cl
+++ /dev/null
@@ -1,73 +0,0 @@
-/***************************************************************************
- *cr
- *cr (C) Copyright 2010 The Board of Trustees of the
- *cr University of Illinois
- *cr All Rights Reserved
- *cr
- ***************************************************************************/
-#define WARP_BITS 5
-
-__kernel void spmv_jds(__global float *dst_vector, __global float *d_data,
- __global int *d_index, __global int *d_perm,
- __global float *x_vec, const int dim,
- __constant int *jds_ptr_int,
- __constant int *sh_zcnt_int)
-{
- int ix = get_global_id(0);
- int warp_id=ix>>WARP_BITS;
-
- if(ix<dim)
- {
- float sum=0.0f;
- int bound=sh_zcnt_int[warp_id];
- //prefetch 0
- int j=jds_ptr_int[0]+ix;
- float d = d_data[j];
- int i = d_index[j];
- float t = x_vec[i];
-
- if (bound>1) //bound >=2
- {
- //prefetch 1
- j=jds_ptr_int[1]+ix;
- i = d_index[j];
- int in;
- float dn;
- float tn;
- for(int k=2;k<bound;k++ )
- {
- //prefetch k-1
- dn = d_data[j];
- //prefetch k
- j=jds_ptr_int[k]+ix;
- in = d_index[j];
- //prefetch k-1
- tn = x_vec[i];
-
- //compute k-2
- sum += d*t;
- //sweep to k
- i = in;
- //sweep to k-1
- d = dn;
- t =tn;
- }
-
- //fetch last
- dn = d_data[j];
- tn = x_vec[i];
-
- //compute last-1
- sum += d*t;
- //sweep to last
- d=dn;
- t=tn;
- }
- //compute last
- sum += d*t; // 3 3
-
- //write out data
- dst_vector[d_perm[ix]]=sum;
- }
-}
-
diff --git a/hpvm/test/parboil/benchmarks/spmv/src/opencl_nvidia/kernel_offline.cl b/hpvm/test/parboil/benchmarks/spmv/src/opencl_nvidia/kernel_offline.cl
deleted file mode 100644
index 9a17a299afdd032fb9c07d8ff5b559055704a573..0000000000000000000000000000000000000000
--- a/hpvm/test/parboil/benchmarks/spmv/src/opencl_nvidia/kernel_offline.cl
+++ /dev/null
@@ -1,73 +0,0 @@
-/***************************************************************************
- *cr
- *cr (C) Copyright 2010 The Board of Trustees of the
- *cr University of Illinois
- *cr All Rights Reserved
- *cr
- ***************************************************************************/
-#define WARP_BITS 5
-
-__kernel void spmv_jds(__global float *dst_vector, __global float *d_data,
- __global int *d_index, __global int *d_perm,
- __global float *x_vec, const int dim,
- __constant int *jds_ptr_int,
- __constant int *sh_zcnt_int)
-{
- int ix = get_global_id(0);
- int warp_id=ix>>WARP_BITS;
-
- if(ix<dim)
- {
- float sum=0.0f;
- int bound=sh_zcnt_int[warp_id];
- //prefetch 0
- int j=jds_ptr_int[0]+ix;
- float d = d_data[j];
- int i = d_index[j];
- float t = x_vec[i];
-
- if (bound>1) //bound >=2
- {
- //prefetch 1
- j=jds_ptr_int[1]+ix;
- i = d_index[j];
- int in;
- float dn;
- float tn;
- for(int k=2;k<bound;k++ )
- {
- //prefetch k-1
- dn = d_data[j];
- //prefetch k
- j=jds_ptr_int[k]+ix;
- in = d_index[j];
- //prefetch k-1
- tn = x_vec[i];
-
- //compute k-2
- sum += d*t;
- //sweep to k
- i = in;
- //sweep to k-1
- d = dn;
- t =tn;
- }
-
- //fetch last
- dn = d_data[j];
- tn = x_vec[i];
-
- //compute last-1
- sum += d*t;
- //sweep to last
- d=dn;
- t=tn;
- }
- //compute last
- sum += d*t; // 3 3
-
- //write out data
- dst_vector[d_perm[ix]]=sum;
- }
-}
-
diff --git a/hpvm/test/parboil/benchmarks/spmv/src/opencl_nvidia/main.c b/hpvm/test/parboil/benchmarks/spmv/src/opencl_nvidia/main.c
deleted file mode 100644
index 343814149aa74139930380c2178e2f447c64e806..0000000000000000000000000000000000000000
--- a/hpvm/test/parboil/benchmarks/spmv/src/opencl_nvidia/main.c
+++ /dev/null
@@ -1,283 +0,0 @@
-/***************************************************************************
- *cr
- *cr (C) Copyright 2010 The Board of Trustees of the
- *cr University of Illinois
- *cr All Rights Reserved
- *cr
- ***************************************************************************/
-
-#include <CL/cl.h>
-#include <CL/cl_ext.h>
-#include <parboil.h>
-#include <stdio.h>
-#include <stdlib.h>
-
-#include "convert_dataset.h"
-#include "file.h"
-#include "gpu_info.h"
-#include "ocl.h"
-
-static int generate_vector(float *x_vector, int dim) {
- srand(54321);
- int i;
- for (i = 0; i < dim; i++) {
- x_vector[i] = (rand() / (float)RAND_MAX);
- }
- return 0;
-}
-
-int main(int argc, char **argv) {
- struct pb_TimerSet timers;
- struct pb_Parameters *parameters;
-
- printf("OpenCL accelerated sparse matrix vector multiplication****\n");
- printf("Li-Wen Chang <lchang20@illinois.edu> and Shengzhao "
- "Wu<wu14@illinois.edu>\n");
- parameters = pb_ReadParameters(&argc, argv);
-
- if ((parameters->inpFiles[0] == NULL) || (parameters->inpFiles[1] == NULL)) {
- fprintf(stderr, "Expecting one two filenames\n");
- exit(-1);
- }
-
- // load matrix from files
- /*pb_SwitchToTimer(&timers, pb_TimerID_IO);*/
- // inputData(parameters->inpFiles[0], &len, &depth, &dim,&nzcnt_len,&pad,
- // &h_data, &h_indices, &h_ptr,
- // &h_perm, &h_nzcnt);
- int col_count;
-
- // parameters declaration
- int len;
- int depth;
- int dim;
- int pad = 32;
- int nzcnt_len;
-
- // host memory allocation
- // matrix
- float *h_data;
- int *h_indices;
- int *h_ptr;
- int *h_perm;
- int *h_nzcnt;
-
- // vector
- float *h_Ax_vector;
- float *h_x_vector;
-
- coo_to_jds(parameters->inpFiles[0], // bcsstk32.mtx, fidapm05.mtx, jgl009.mtx
- 1, // row padding
- pad, // warp size
- 1, // pack size
- 1, // is mirrored?
- 0, // binary matrix
- 1, // debug level [0:2]
- &h_data, &h_ptr, &h_nzcnt, &h_indices, &h_perm, &col_count, &dim,
- &len, &nzcnt_len, &depth);
-
- h_Ax_vector = (float *)malloc(sizeof(float) * dim);
- h_x_vector = (float *)malloc(sizeof(float) * dim);
- input_vec(parameters->inpFiles[1], h_x_vector, dim);
-
- pb_InitializeTimerSet(&timers);
- pb_SwitchToTimer(&timers, pb_TimerID_COMPUTE);
-
- cl_int clStatus;
- cl_platform_id clPlatform;
- clStatus = clGetPlatformIDs(1, &clPlatform, NULL);
- CHECK_ERROR("clGetPlatformIDs")
-
- cl_context_properties clCps[3] = {CL_CONTEXT_PLATFORM,
- (cl_context_properties)clPlatform, 0};
-
- cl_device_id clDevice;
- clStatus = clGetDeviceIDs(clPlatform, CL_DEVICE_TYPE_GPU, 1, &clDevice, NULL);
- CHECK_ERROR("clGetDeviceIDs")
-
- cl_context clContext =
- clCreateContextFromType(clCps, CL_DEVICE_TYPE_GPU, NULL, NULL, &clStatus);
- CHECK_ERROR("clCreateContextFromType")
-
- cl_command_queue clCommandQueue = clCreateCommandQueue(
- clContext, clDevice, CL_QUEUE_PROFILING_ENABLE, &clStatus);
- CHECK_ERROR("clCreateCommandQueue")
-
- pb_SetOpenCL(&clContext, &clCommandQueue);
-
- const char *clSource[] = {readFile("src/opencl_nvidia/kernel.cl")};
- cl_program clProgram =
- clCreateProgramWithSource(clContext, 1, clSource, NULL, &clStatus);
- CHECK_ERROR("clCreateProgramWithSource")
-
- char clOptions[50];
- sprintf(clOptions, "");
- clStatus = clBuildProgram(clProgram, 1, &clDevice, clOptions, NULL, NULL);
- CHECK_ERROR("clBuildProgram")
-
- cl_kernel clKernel = clCreateKernel(clProgram, "spmv_jds", &clStatus);
- CHECK_ERROR("clCreateKernel")
- /*cl_kernel clKernel;*/
- /*cl_program clProgram;*/
- /*pb_CreateAndBuildKernelFromBinary("build/opencl_nvidia_default/kernel_offline.nvptx.s",
- * "spmv_jds", &clContext, &clDevice, &clProgram, &clKernel);*/
-
- pb_SwitchToTimer(&timers, pb_TimerID_COMPUTE);
- // device memory allocation
- // matrix
- cl_mem d_data;
- cl_mem d_indices;
- cl_mem d_perm;
- cl_mem d_Ax_vector;
- cl_mem d_x_vector;
-
- cl_mem jds_ptr_int;
- cl_mem sh_zcnt_int;
-
- OpenCLDeviceProp clDeviceProp;
- clStatus = clGetDeviceInfo(clDevice, CL_DEVICE_COMPUTE_CAPABILITY_MAJOR_NV,
- sizeof(cl_uint), &(clDeviceProp.major), NULL);
- CHECK_ERROR("clGetDeviceInfo")
- clStatus = clGetDeviceInfo(clDevice, CL_DEVICE_COMPUTE_CAPABILITY_MINOR_NV,
- sizeof(cl_uint), &(clDeviceProp.minor), NULL);
- CHECK_ERROR("clGetDeviceInfo")
- clStatus =
- clGetDeviceInfo(clDevice, CL_DEVICE_MAX_COMPUTE_UNITS, sizeof(cl_uint),
- &(clDeviceProp.multiProcessorCount), NULL);
- CHECK_ERROR("clGetDeviceInfo")
-
- pb_SwitchToTimer(&timers, pb_TimerID_COPY);
-
- // memory allocation
- d_data = clCreateBuffer(clContext, CL_MEM_READ_ONLY, len * sizeof(float),
- NULL, &clStatus);
- CHECK_ERROR("clCreateBuffer")
- d_indices = clCreateBuffer(clContext, CL_MEM_READ_ONLY, len * sizeof(int),
- NULL, &clStatus);
- CHECK_ERROR("clCreateBuffer")
- d_perm = clCreateBuffer(clContext, CL_MEM_READ_ONLY, dim * sizeof(int), NULL,
- &clStatus);
- CHECK_ERROR("clCreateBuffer")
- d_x_vector = clCreateBuffer(clContext, CL_MEM_READ_ONLY, dim * sizeof(float),
- NULL, &clStatus);
- CHECK_ERROR("clCreateBuffer")
- d_Ax_vector = clCreateBuffer(clContext, CL_MEM_WRITE_ONLY,
- dim * sizeof(float), NULL, &clStatus);
- CHECK_ERROR("clCreateBuffer")
-
- jds_ptr_int = clCreateBuffer(clContext, CL_MEM_READ_ONLY, depth * sizeof(int),
- NULL, &clStatus);
- CHECK_ERROR("clCreateBuffer")
- sh_zcnt_int = clCreateBuffer(clContext, CL_MEM_READ_ONLY,
- nzcnt_len * sizeof(int), NULL, &clStatus);
- CHECK_ERROR("clCreateBuffer")
-
- clMemSet(clCommandQueue, d_Ax_vector, 0, dim * sizeof(float));
-
- // memory copy
- clStatus = clEnqueueWriteBuffer(clCommandQueue, d_data, CL_FALSE, 0,
- len * sizeof(float), h_data, 0, NULL, NULL);
- CHECK_ERROR("clEnqueueWriteBuffer")
- clStatus = clEnqueueWriteBuffer(clCommandQueue, d_indices, CL_FALSE, 0,
- len * sizeof(int), h_indices, 0, NULL, NULL);
- CHECK_ERROR("clEnqueueWriteBuffer")
- clStatus = clEnqueueWriteBuffer(clCommandQueue, d_perm, CL_FALSE, 0,
- dim * sizeof(int), h_perm, 0, NULL, NULL);
- CHECK_ERROR("clEnqueueWriteBuffer")
- clStatus = clEnqueueWriteBuffer(clCommandQueue, d_x_vector, CL_FALSE, 0,
- dim * sizeof(int), h_x_vector, 0, NULL, NULL);
- CHECK_ERROR("clEnqueueWriteBuffer")
-
- clStatus = clEnqueueWriteBuffer(clCommandQueue, jds_ptr_int, CL_FALSE, 0,
- depth * sizeof(int), h_ptr, 0, NULL, NULL);
- CHECK_ERROR("clEnqueueWriteBuffer")
- clStatus =
- clEnqueueWriteBuffer(clCommandQueue, sh_zcnt_int, CL_TRUE, 0,
- nzcnt_len * sizeof(int), h_nzcnt, 0, NULL, NULL);
- CHECK_ERROR("clEnqueueWriteBuffer")
-
- pb_SwitchToTimer(&timers, pb_TimerID_COMPUTE);
-
- size_t grid;
- size_t block;
-
- compute_active_thread(&block, &grid, nzcnt_len, pad, clDeviceProp.major,
- clDeviceProp.minor, clDeviceProp.multiProcessorCount);
-
- clStatus = clSetKernelArg(clKernel, 0, sizeof(cl_mem), &d_Ax_vector);
- CHECK_ERROR("clSetKernelArg")
- clStatus = clSetKernelArg(clKernel, 1, sizeof(cl_mem), &d_data);
- CHECK_ERROR("clSetKernelArg")
- clStatus = clSetKernelArg(clKernel, 2, sizeof(cl_mem), &d_indices);
- CHECK_ERROR("clSetKernelArg")
- clStatus = clSetKernelArg(clKernel, 3, sizeof(cl_mem), &d_perm);
- CHECK_ERROR("clSetKernelArg")
- clStatus = clSetKernelArg(clKernel, 4, sizeof(cl_mem), &d_x_vector);
- CHECK_ERROR("clSetKernelArg")
- clStatus = clSetKernelArg(clKernel, 5, sizeof(int), &dim);
- CHECK_ERROR("clSetKernelArg")
-
- clStatus = clSetKernelArg(clKernel, 6, sizeof(cl_mem), &jds_ptr_int);
- CHECK_ERROR("clSetKernelArg")
- clStatus = clSetKernelArg(clKernel, 7, sizeof(cl_mem), &sh_zcnt_int);
- CHECK_ERROR("clSetKernelArg")
-
- // main execution
-
- pb_SwitchToTimer(&timers, visc_TimerID_COMPUTATION);
- int i;
- for (i = 0; i < 50; i++) {
- clStatus = clEnqueueNDRangeKernel(clCommandQueue, clKernel, 1, NULL, &grid,
- &block, 0, NULL, NULL);
- CHECK_ERROR("clEnqueueNDRangeKernel")
- }
-
- clStatus = clFinish(clCommandQueue);
- CHECK_ERROR("clFinish")
-
- pb_SwitchToTimer(&timers, pb_TimerID_COPY);
-
- // HtoD memory copy
- clStatus =
- clEnqueueReadBuffer(clCommandQueue, d_Ax_vector, CL_TRUE, 0,
- dim * sizeof(float), h_Ax_vector, 0, NULL, NULL);
- CHECK_ERROR("clEnqueueReadBuffer")
-
- pb_SwitchToTimer(&timers, pb_TimerID_COMPUTE);
- clStatus = clReleaseKernel(clKernel);
- clStatus = clReleaseProgram(clProgram);
-
- clStatus = clReleaseMemObject(d_data);
- clStatus = clReleaseMemObject(d_indices);
- clStatus = clReleaseMemObject(d_perm);
- clStatus = clReleaseMemObject(sh_zcnt_int);
- clStatus = clReleaseMemObject(jds_ptr_int);
- clStatus = clReleaseMemObject(d_x_vector);
- clStatus = clReleaseMemObject(d_Ax_vector);
-
- clStatus = clReleaseCommandQueue(clCommandQueue);
- clStatus = clReleaseContext(clContext);
-
- pb_SwitchToTimer(&timers, pb_TimerID_NONE);
-
- pb_PrintTimerSet(&timers);
- if (parameters->outFile) {
- /*pb_SwitchToTimer(&timers, pb_TimerID_IO);*/
- outputData(parameters->outFile, h_Ax_vector, dim);
- }
- /*pb_SwitchToTimer(&timers, pb_TimerID_COMPUTE);*/
-
- // free((void*)clSource[0]);
-
- free(h_data);
- free(h_indices);
- free(h_ptr);
- free(h_perm);
- free(h_nzcnt);
- free(h_Ax_vector);
- free(h_x_vector);
-
- pb_FreeParameters(parameters);
-
- return 0;
-}
diff --git a/hpvm/test/parboil/benchmarks/spmv/src/opencl_nvidia/ocl.c b/hpvm/test/parboil/benchmarks/spmv/src/opencl_nvidia/ocl.c
deleted file mode 100644
index 2990031255acae7fe480b0fe7cdc79db7cb08287..0000000000000000000000000000000000000000
--- a/hpvm/test/parboil/benchmarks/spmv/src/opencl_nvidia/ocl.c
+++ /dev/null
@@ -1,46 +0,0 @@
-#include "ocl.h"
-#include <CL/cl.h>
-#include <stdio.h>
-#include <string.h>
-
-char *readFile(const char *fileName) {
- FILE *fp;
- fp = fopen(fileName, "r");
- if (fp == NULL) {
- printf("Error 1!\n");
- exit(1);
- }
-
- fseek(fp, 0, SEEK_END);
- long size = ftell(fp);
- rewind(fp);
-
- char *buffer = (char *)malloc(sizeof(char) * (size + 1));
- if (buffer == NULL) {
- printf("Error 2!\n");
- fclose(fp);
- exit(1);
- }
-
- size_t res = fread(buffer, 1, size, fp);
- if (res != size) {
- printf("Error 3!\n");
- fclose(fp);
- exit(1);
- }
-
- buffer[size] = 0;
- fclose(fp);
- return buffer;
-}
-
-void clMemSet(cl_command_queue clCommandQueue, cl_mem buf, int val,
- size_t size) {
- cl_int clStatus;
- char *temp = (char *)malloc(size);
- memset(temp, val, size);
- clStatus = clEnqueueWriteBuffer(clCommandQueue, buf, CL_TRUE, 0, size, temp,
- 0, NULL, NULL);
- CHECK_ERROR("clEnqueueWriteBuffer")
- free(temp);
-}
diff --git a/hpvm/test/parboil/benchmarks/spmv/src/opencl_nvidia/ocl.h b/hpvm/test/parboil/benchmarks/spmv/src/opencl_nvidia/ocl.h
deleted file mode 100644
index 42ff7b4d1059550293b56325d0cce2afea6c004b..0000000000000000000000000000000000000000
--- a/hpvm/test/parboil/benchmarks/spmv/src/opencl_nvidia/ocl.h
+++ /dev/null
@@ -1,20 +0,0 @@
-#ifndef __OCLH__
-#define __OCLH__
-
-typedef struct {
- cl_uint major;
- cl_uint minor;
- cl_uint multiProcessorCount;
-} OpenCLDeviceProp;
-
-void clMemSet(cl_command_queue, cl_mem, int, size_t);
-char *readFile(const char *);
-
-#define CHECK_ERROR(errorMessage) \
- if (clStatus != CL_SUCCESS) { \
- printf("Error: %s!\n", errorMessage); \
- printf("Line: %d\n", __LINE__); \
- exit(1); \
- }
-
-#endif
diff --git a/hpvm/test/parboil/benchmarks/stencil/src/opencl_base/Makefile b/hpvm/test/parboil/benchmarks/stencil/src/opencl_base/Makefile
deleted file mode 100644
index c4ef6c0dac3a1646d63aca309b2d44138abb55d3..0000000000000000000000000000000000000000
--- a/hpvm/test/parboil/benchmarks/stencil/src/opencl_base/Makefile
+++ /dev/null
@@ -1,8 +0,0 @@
-# (c) 2010 The Board of Trustees of the University of Illinois.
-
-LANGUAGE=opencl
-SRCDIR_OBJS=main.o file.o
-APP_CUDALDFLAGS=-lm
-APP_CFLAGS=-ffast-math -O3
-APP_CXXFLAGS=-ffast-math -O3
-KERNEL_OBJS=kernel_offline.nvptx.s
diff --git a/hpvm/test/parboil/benchmarks/stencil/src/opencl_base/common.h b/hpvm/test/parboil/benchmarks/stencil/src/opencl_base/common.h
deleted file mode 100644
index 1a682890b3619ef712c5e5e3a7313e325935ec6f..0000000000000000000000000000000000000000
--- a/hpvm/test/parboil/benchmarks/stencil/src/opencl_base/common.h
+++ /dev/null
@@ -1,12 +0,0 @@
-/***************************************************************************
- *cr
- *cr (C) Copyright 2010 The Board of Trustees of the
- *cr University of Illinois
- *cr All Rights Reserved
- *cr
- ***************************************************************************/
-
-#ifndef _COMMON_H_
-#define _COMMON_H_
-#define Index3D(_nx, _ny, _i, _j, _k) ((_i) + _nx * ((_j) + _ny * (_k)))
-#endif
diff --git a/hpvm/test/parboil/benchmarks/stencil/src/opencl_base/file.cc b/hpvm/test/parboil/benchmarks/stencil/src/opencl_base/file.cc
deleted file mode 100644
index 95cd65c4a0e013c60c6edd43077346a7efdad1ae..0000000000000000000000000000000000000000
--- a/hpvm/test/parboil/benchmarks/stencil/src/opencl_base/file.cc
+++ /dev/null
@@ -1,76 +0,0 @@
-/***************************************************************************
- *cr
- *cr (C) Copyright 2010 The Board of Trustees of the
- *cr University of Illinois
- *cr All Rights Reserved
- *cr
- ***************************************************************************/
-
-#include <endian.h>
-#include <inttypes.h>
-#include <malloc.h>
-#include <stdio.h>
-#include <stdlib.h>
-
-#if __BYTE_ORDER != __LITTLE_ENDIAN
-#error "File I/O is not implemented for this system: wrong endianness."
-#endif
-
-extern "C" void inputData(char *fName, int *nx, int *ny, int *nz) {
- FILE *fid = fopen(fName, "r");
-
- if (fid == NULL) {
- fprintf(stderr, "Cannot open input file\n");
- exit(-1);
- }
-
- fread(nx, sizeof(int), 1, fid);
- fread(ny, sizeof(int), 1, fid);
- fread(nz, sizeof(int), 1, fid);
- fclose(fid);
-}
-
-extern "C" void outputData(char *fName, float *h_A0, int nx, int ny, int nz) {
- FILE *fid = fopen(fName, "w");
- uint32_t tmp32;
- if (fid == NULL) {
- fprintf(stderr, "Cannot open output file\n");
- exit(-1);
- }
- tmp32 = nx * ny * nz;
- fwrite(&tmp32, sizeof(uint32_t), 1, fid);
- fwrite(h_A0, sizeof(float), tmp32, fid);
-
- fclose(fid);
-}
-
-extern "C" char *readFile(const char *fileName) {
- FILE *fp;
- fp = fopen(fileName, "r");
- if (fp == NULL) {
- printf("Error 1!\n");
- exit(1);
- }
-
- fseek(fp, 0, SEEK_END);
- long size = ftell(fp);
- rewind(fp);
-
- char *buffer = (char *)malloc(sizeof(char) * (size + 1));
- if (buffer == NULL) {
- printf("Error 2!\n");
- fclose(fp);
- exit(1);
- }
-
- size_t res = fread(buffer, 1, size, fp);
- if (res != size) {
- printf("Error 3!\n");
- fclose(fp);
- exit(1);
- }
-
- buffer[size] = 0;
- fclose(fp);
- return buffer;
-}
diff --git a/hpvm/test/parboil/benchmarks/stencil/src/opencl_base/file.h b/hpvm/test/parboil/benchmarks/stencil/src/opencl_base/file.h
deleted file mode 100644
index b45c42371bbde3c3a39d88277adf39a8f537baab..0000000000000000000000000000000000000000
--- a/hpvm/test/parboil/benchmarks/stencil/src/opencl_base/file.h
+++ /dev/null
@@ -1,23 +0,0 @@
-/***************************************************************************
- *cr
- *cr (C) Copyright 2010 The Board of Trustees of the
- *cr University of Illinois
- *cr All Rights Reserved
- *cr
- ***************************************************************************/
-#ifndef __FILEH__
-#define __FILEH__
-
-#ifdef __cplusplus
-extern "C" {
-#endif
-
-void inputData(char *fName, int *nx, int *ny, int *nz);
-void outputData(char *fName, float *h_A0, int nx, int ny, int nz);
-char *readFile(const char *fileName);
-
-#ifdef __cplusplus
-}
-#endif
-
-#endif
diff --git a/hpvm/test/parboil/benchmarks/stencil/src/opencl_base/kernel.cl b/hpvm/test/parboil/benchmarks/stencil/src/opencl_base/kernel.cl
deleted file mode 100644
index 4c5d1263db5948e4e61ea2baa27261613cd0ea06..0000000000000000000000000000000000000000
--- a/hpvm/test/parboil/benchmarks/stencil/src/opencl_base/kernel.cl
+++ /dev/null
@@ -1,28 +0,0 @@
-/***************************************************************************
- *cr
- *cr (C) Copyright 2010 The Board of Trustees of the
- *cr University of Illinois
- *cr All Rights Reserved
- *cr
- ***************************************************************************/
-
-#include "common.h"
-
-__kernel void naive_kernel(float c0,float c1,__global float* A0,__global float *Anext,int nx,int ny,int nz)
-{
- int i = get_global_id(0)+1;
- int j = get_global_id(1)+1;
- int k = get_global_id(2)+1;
-
-if(i<nx-1)
-{
- Anext[Index3D (nx, ny, i, j, k)] = c1 *
- ( A0[Index3D (nx, ny, i, j, k + 1)] +
- A0[Index3D (nx, ny, i, j, k - 1)] +
- A0[Index3D (nx, ny, i, j + 1, k)] +
- A0[Index3D (nx, ny, i, j - 1, k)] +
- A0[Index3D (nx, ny, i + 1, j, k)] +
- A0[Index3D (nx, ny, i - 1, j, k)] )
- - A0[Index3D (nx, ny, i, j, k)] * c0;
-}
-}
diff --git a/hpvm/test/parboil/benchmarks/stencil/src/opencl_base/kernel_offline.cl b/hpvm/test/parboil/benchmarks/stencil/src/opencl_base/kernel_offline.cl
deleted file mode 100644
index 4c5d1263db5948e4e61ea2baa27261613cd0ea06..0000000000000000000000000000000000000000
--- a/hpvm/test/parboil/benchmarks/stencil/src/opencl_base/kernel_offline.cl
+++ /dev/null
@@ -1,28 +0,0 @@
-/***************************************************************************
- *cr
- *cr (C) Copyright 2010 The Board of Trustees of the
- *cr University of Illinois
- *cr All Rights Reserved
- *cr
- ***************************************************************************/
-
-#include "common.h"
-
-__kernel void naive_kernel(float c0,float c1,__global float* A0,__global float *Anext,int nx,int ny,int nz)
-{
- int i = get_global_id(0)+1;
- int j = get_global_id(1)+1;
- int k = get_global_id(2)+1;
-
-if(i<nx-1)
-{
- Anext[Index3D (nx, ny, i, j, k)] = c1 *
- ( A0[Index3D (nx, ny, i, j, k + 1)] +
- A0[Index3D (nx, ny, i, j, k - 1)] +
- A0[Index3D (nx, ny, i, j + 1, k)] +
- A0[Index3D (nx, ny, i, j - 1, k)] +
- A0[Index3D (nx, ny, i + 1, j, k)] +
- A0[Index3D (nx, ny, i - 1, j, k)] )
- - A0[Index3D (nx, ny, i, j, k)] * c0;
-}
-}
diff --git a/hpvm/test/parboil/benchmarks/stencil/src/opencl_base/main.c b/hpvm/test/parboil/benchmarks/stencil/src/opencl_base/main.c
deleted file mode 100644
index ec47c22227648df094cbf03ea1b667943207207e..0000000000000000000000000000000000000000
--- a/hpvm/test/parboil/benchmarks/stencil/src/opencl_base/main.c
+++ /dev/null
@@ -1,231 +0,0 @@
-
-/***************************************************************************
- *cr
- *cr (C) Copyright 2010 The Board of Trustees of the
- *cr University of Illinois
- *cr All Rights Reserved
- *cr
- ***************************************************************************/
-
-#include <CL/cl.h>
-#include <parboil.h>
-#include <stdio.h>
-#include <stdlib.h>
-#include <string.h>
-
-#include "common.h"
-#include "file.h"
-
-#define CHECK_ERROR(errorMessage) \
- if (clStatus != CL_SUCCESS) { \
- printf("Error: %s!\n", errorMessage); \
- printf("Line: %d\n", __LINE__); \
- exit(1); \
- }
-
-static int read_data(float *A0, int nx, int ny, int nz, FILE *fp) {
- int s = 0;
- int i, j, k;
- for (i = 0; i < nz; i++) {
- for (j = 0; j < ny; j++) {
- for (k = 0; k < nx; k++) {
- fread(A0 + s, sizeof(float), 1, fp);
- s++;
- }
- }
- }
- return 0;
-}
-
-int main(int argc, char **argv) {
- struct pb_TimerSet timers;
- struct pb_Parameters *parameters;
-
- printf("OpenCL accelerated 7 points stencil codes****\n");
- printf("Author: Li-Wen Chang <lchang20@illinois.edu>\n");
- parameters = pb_ReadParameters(&argc, argv);
-
- // declaration
- unsigned nx, ny, nz;
- unsigned size;
- int iteration;
- float c0 = 1.0f / 6.0f;
- float c1 = 1.0f / 6.0f / 6.0f;
-
- if (argc < 5) {
- printf("Usage: probe nx ny nz t\n"
- "nx: the grid size x\n"
- "ny: the grid size y\n"
- "nz: the grid size z\n"
- "t: the iteration time\n");
- return -1;
- }
-
- nx = atoi(argv[1]);
- if (nx < 1)
- return -1;
- ny = atoi(argv[2]);
- if (ny < 1)
- return -1;
- nz = atoi(argv[3]);
- if (nz < 1)
- return -1;
- iteration = atoi(argv[4]);
- if (iteration < 1)
- return -1;
-
- // host data
- float *h_A0;
- float *h_Anext;
- // load data from files
-
- size = nx * ny * nz;
-
- h_A0 = (float *)malloc(sizeof(float) * size);
- h_Anext = (float *)malloc(sizeof(float) * size);
- /*pb_SwitchToTimer(&timers, pb_TimerID_IO);*/
- FILE *fp = fopen(parameters->inpFiles[0], "rb");
- read_data(h_A0, nx, ny, nz, fp);
- fclose(fp);
-
- pb_InitializeTimerSet(&timers);
- pb_SwitchToTimer(&timers, pb_TimerID_COMPUTE);
- cl_int clStatus;
- cl_platform_id clPlatform;
- clStatus = clGetPlatformIDs(1, &clPlatform, NULL);
- CHECK_ERROR("clGetPlaformIDs")
-
- cl_context_properties clCps[3] = {CL_CONTEXT_PLATFORM,
- (cl_context_properties)clPlatform, 0};
-
- cl_device_id clDevice;
- clStatus = clGetDeviceIDs(clPlatform, CL_DEVICE_TYPE_GPU, 1, &clDevice, NULL);
- CHECK_ERROR("clGetDeviceIDs")
-
- cl_context clContext =
- clCreateContextFromType(clCps, CL_DEVICE_TYPE_GPU, NULL, NULL, &clStatus);
- CHECK_ERROR("clCreateContextFromType")
-
- cl_command_queue clCommandQueue = clCreateCommandQueue(
- clContext, clDevice, CL_QUEUE_PROFILING_ENABLE, &clStatus);
- CHECK_ERROR("clCreateCommandQueue")
-
- pb_SetOpenCL(&clContext, &clCommandQueue);
-
- /*cl_program clProgram;*/
- /*cl_kernel clKernel;*/
-
- /*pb_CreateAndBuildKernelFromBinary("build/opencl_base_default/kernel_offline.nvptx.s",
- * "naive_kernel", &clContext, &clDevice, &clProgram, &clKernel);*/
- const char *clSource[] = {readFile("src/opencl_base/kernel.cl")};
- cl_program clProgram =
- clCreateProgramWithSource(clContext, 1, clSource, NULL, &clStatus);
- CHECK_ERROR("clCreateProgramWithSource")
-
- char clOptions[50];
- sprintf(clOptions, "-I src/opencl_base");
- clStatus = clBuildProgram(clProgram, 1, &clDevice, clOptions, NULL, NULL);
- CHECK_ERROR("clBuildProgram")
-
- cl_kernel clKernel = clCreateKernel(clProgram, "naive_kernel", &clStatus);
- CHECK_ERROR("clCreateKernel")
-
- // device
- cl_mem d_A0;
- cl_mem d_Anext;
-
- memcpy(h_Anext, h_A0, sizeof(float) * size);
-
- // memory allocation
- d_A0 = clCreateBuffer(clContext, CL_MEM_READ_WRITE, size * sizeof(float),
- NULL, &clStatus);
- CHECK_ERROR("clCreateBuffer")
- d_Anext = clCreateBuffer(clContext, CL_MEM_READ_WRITE, size * sizeof(float),
- NULL, &clStatus);
- CHECK_ERROR("clCreateBuffer")
-
- pb_SwitchToTimer(&timers, pb_TimerID_COPY);
- // memory copy
- clStatus = clEnqueueWriteBuffer(clCommandQueue, d_A0, CL_FALSE, 0,
- size * sizeof(float), h_A0, 0, NULL, NULL);
- CHECK_ERROR("clEnqueueWriteBuffer")
- clStatus = clEnqueueWriteBuffer(clCommandQueue, d_Anext, CL_TRUE, 0,
- size * sizeof(float), h_Anext, 0, NULL, NULL);
- CHECK_ERROR("clEnqueueWriteBuffer")
-
- pb_SwitchToTimer(&timers, pb_TimerID_COMPUTE);
-
- // only use 1D thread block
- unsigned tx = 256;
- size_t block[3] = {tx, 1, 1};
- size_t grid[3] = {(nx - 2 + tx - 1) / tx * tx, ny - 2, nz - 2};
- // size_t grid[3] = {nx-2,ny-2,nz-2};
- size_t offset[3] = {1, 1, 1};
- // printf("block x is %d and y is %d z \n",block[0],block[1]);
- // printf("grid x is %d and y is %d\n",grid[0],grid[1]);
-
- clStatus = clSetKernelArg(clKernel, 0, sizeof(float), (void *)&c0);
- clStatus = clSetKernelArg(clKernel, 1, sizeof(float), (void *)&c1);
- clStatus = clSetKernelArg(clKernel, 2, sizeof(cl_mem), (void *)&d_A0);
- clStatus = clSetKernelArg(clKernel, 3, sizeof(cl_mem), (void *)&d_Anext);
- clStatus = clSetKernelArg(clKernel, 4, sizeof(int), (void *)&nx);
- clStatus = clSetKernelArg(clKernel, 5, sizeof(int), (void *)&ny);
- clStatus = clSetKernelArg(clKernel, 6, sizeof(int), (void *)&nz);
- CHECK_ERROR("clSetKernelArg")
-
- // main execution
- pb_SwitchToTimer(&timers, visc_TimerID_COMPUTATION);
- int t;
- for (t = 0; t < iteration; t++) {
- clStatus = clEnqueueNDRangeKernel(clCommandQueue, clKernel, 3, NULL, grid,
- block, 0, NULL, NULL);
- // printf("iteration %d\n",t)
- CHECK_ERROR("clEnqueueNDRangeKernel")
-
- cl_mem d_temp = d_A0;
- d_A0 = d_Anext;
- d_Anext = d_temp;
-
- clStatus = clSetKernelArg(clKernel, 2, sizeof(cl_mem), (void *)&d_A0);
- clStatus = clSetKernelArg(clKernel, 3, sizeof(cl_mem), (void *)&d_Anext);
- }
-
- cl_mem d_temp = d_A0;
- d_A0 = d_Anext;
- d_Anext = d_temp;
-
- clStatus = clFinish(clCommandQueue);
- CHECK_ERROR("clFinish")
-
- pb_SwitchToTimer(&timers, pb_TimerID_COPY);
- clStatus = clEnqueueReadBuffer(clCommandQueue, d_Anext, CL_TRUE, 0,
- size * sizeof(float), h_Anext, 0, NULL, NULL);
- CHECK_ERROR("clEnqueueReadBuffer")
-
- pb_SwitchToTimer(&timers, pb_TimerID_COMPUTE);
- clStatus = clReleaseMemObject(d_A0);
- clStatus = clReleaseMemObject(d_Anext);
- clStatus = clReleaseKernel(clKernel);
- clStatus = clReleaseProgram(clProgram);
- clStatus = clReleaseCommandQueue(clCommandQueue);
- clStatus = clReleaseContext(clContext);
- CHECK_ERROR("clReleaseContext")
-
- pb_SwitchToTimer(&timers, pb_TimerID_NONE);
-
- pb_PrintTimerSet(&timers);
-
- if (parameters->outFile) {
- /*pb_SwitchToTimer(&timers, pb_TimerID_IO);*/
- outputData(parameters->outFile, h_Anext, nx, ny, nz);
- }
- /*pb_SwitchToTimer(&timers, pb_TimerID_COMPUTE);*/
-
- // free((void*)clSource[0]);
-
- free(h_A0);
- free(h_Anext);
- pb_FreeParameters(parameters);
-
- return 0;
-}
diff --git a/hpvm/test/template/Makefile b/hpvm/test/template/Makefile
index 82d7c61e2b77947da770aab506c45919ee139a63..3aa4bd1d6f2c3f7bb2be07ba5e662c5b6faf1655 100644
--- a/hpvm/test/template/Makefile
+++ b/hpvm/test/template/Makefile
@@ -41,8 +41,8 @@ OBJS_CFLAGS = $(APP_CFLAGS) $(PLATFORM_CFLAGS)
CXXFLAGS = $(APP_CXXFLAGS) $(PLATFORM_CXXFLAGS)
LDFLAGS= $(APP_LDFLAGS) $(PLATFORM_LDFLAGS)
-VISC_RT_PATH = $(LLVM_SRC_ROOT)/tools/hpvm/projects/visc-rt
-VISC_RT_LIB = $(VISC_RT_PATH)/visc-rt.ll
+VISC_RT_PATH = $(LLVM_BUILD_ROOT)/tools/hpvm/projects/visc-rt
+VISC_RT_LIB = $(VISC_RT_PATH)/visc-rt.bc
TESTGEN_OPTFLAGS = -load LLVMGenVISC.so -genvisc -globaldce
@@ -69,9 +69,7 @@ OBJS = $(call INBUILDDIR,$(SRCDIR_OBJS))
TEST_OBJS = $(call INBUILDDIR,$(VISC_OBJS))
KERNEL = $(TEST_OBJS).kernels.ll
-ifeq ($(TARGET),seq)
-else
- KERNEL_LINKED = $(BUILD_DIR)/$(APP).kernels.linked.ll
+ifeq ($(TARGET),gpu)
KERNEL_OCL = $(TEST_OBJS).kernels.cl
endif
@@ -88,7 +86,7 @@ endif
default: $(FAILSAFE) $(BUILD_DIR) $(KERNEL_OCL) $(EXE)
$(KERNEL_OCL) : $(KERNEL)
- $(OCLBE) -debug $< -o $@
+ $(OCLBE) $< -o $@
$(EXE) : $(HOST_LINKED)
$(CXX) -O3 $(LDFLAGS) $< -o $@
@@ -96,11 +94,8 @@ $(EXE) : $(HOST_LINKED)
$(HOST_LINKED) : $(HOST) $(OBJS) $(VISC_RT_LIB)
$(LLVM_LINK) $^ -S -o $@
-$(VISC_RT_LIB) : $(VISC_RT_PATH)/visc-rt.cpp
- make -C $(LLVM_LIB_PATH)
-
$(HOST) $(KERNEL): $(BUILD_DIR)/$(VISC_OBJS)
- $(OPT) -debug $(VISC_OPTFLAGS) -S $< -o $(HOST)
+ $(OPT) $(VISC_OPTFLAGS) -S $< -o $(HOST)
$(BUILD_DIR):
mkdir -p $(BUILD_DIR)
@@ -112,6 +107,6 @@ $(BUILD_DIR)/main.ll : $(SRC_DIR)/main.cc
$(CC) $(CXXFLAGS) -emit-llvm -S -o $@ $<
$(BUILD_DIR)/main.visc.ll : $(BUILD_DIR)/main.ll
- $(OPT) -debug-only=genvisc $(TESTGEN_OPTFLAGS) $< -S -o $@
+ $(OPT) $(TESTGEN_OPTFLAGS) $< -S -o $@
## END HPVM MAKEFILE
diff --git a/hpvm/test/template/Makefile.config.example b/hpvm/test/template/Makefile.config.example
index 269f0b7df273c958f0cd20a0f935716a329e00ae..9b182a3b7998f3a302a758796a9fe950740b1366 100644
--- a/hpvm/test/template/Makefile.config.example
+++ b/hpvm/test/template/Makefile.config.example
@@ -1,23 +1,23 @@
-CUDA_PATH=/usr/local/cuda
-CUDA_LIB_PATH=/usr/local/cuda/lib64
-OPENCL_PATH=/opt/intelFPGA_pro/18.0/hld/host/linux64
-OPENCL_LIB_PATH=$(OPENCL_PATH)/lib
+# These paths should be set to your local directories of CUDA and OpenCL
+CUDA_PATH=/software/cuda-9.1
+CUDA_LIB_PATH=$(CUDA_PATH)/lib64
+OPENCL_PATH=$(CUDA_PATH)
+OPENCL_LIB_PATH=$(OPENCL_PATH)/lib64
-# NOTE: You may need to configure this based on your root path.
-VISC_SRC_ROOT=$(LLVM_SRC_ROOT)
+# These paths should be set based on your HPVM build and source locations.
-VISC_BUILD_DIR =$(VISC_SRC_ROOT)/build
-CC = $(VISC_BUILD_DIR)/bin/clang
-PLATFORM_CFLAGS = -I$(LLVM_SRC_ROOT)/include -I$(VISC_BUILD_DIR)/include
+LLVM_BUILD_DIR =$(LLVM_SRC_ROOT)/../build
+CC = $(LLVM_BUILD_DIR)/bin/clang
+PLATFORM_CFLAGS = -I$(LLVM_SRC_ROOT)/include -I$(OPENCL_PATH)/include/CL/ -I$(LLVM_BUILD_DIR)/include
-CXX = $(VISC_BUILD_DIR)/bin/clang++
-PLATFORM_CXXFLAGS = -I$(LLVM_SRC_ROOT)/include -I$(VISC_BUILD_DIR)/include
+CXX = $(LLVM_BUILD_DIR)/bin/clang++
+PLATFORM_CXXFLAGS = -I$(LLVM_SRC_ROOT)/include -I$(OPENCL_PATH)/include/CL/ -I$(LLVM_BUILD_DIR)/include
-LINKER = $(VISC_BUILD_DIR)/bin/clang++
+LINKER = $(LLVM_BUILD_DIR)/bin/clang++
PLATFORM_LDFLAGS = -lm -lpthread -lrt -lOpenCL -L$(OPENCL_LIB_PATH)
-LLVM_LIB_PATH = $(VISC_BUILD_DIR)/lib
-LLVM_BIN_PATH = $(VISC_BUILD_DIR)/bin
+LLVM_LIB_PATH = $(LLVM_BUILD_DIR)/lib
+LLVM_BIN_PATH = $(LLVM_BUILD_DIR)/bin
OPT = $(LLVM_BIN_PATH)/opt
LLVM_LINK = $(LLVM_BIN_PATH)/llvm-link
diff --git a/hpvm/test/template/README.md b/hpvm/test/template/README.md
new file mode 100644
index 0000000000000000000000000000000000000000..198604817d5a8463e555451a1188b426ec4e31cd
--- /dev/null
+++ b/hpvm/test/template/README.md
@@ -0,0 +1,37 @@
+Compilation of a project with HPVM is a multi-step process.
+Let's look at the compilation of the `pipeline` test for gpu as an example.
+
+`clang` is used to produce an LLVM IR file that contains the HPVM intrinsics in the form of function calls.
+```
+/.../hpvm/build/bin/clang -Isrc/ -I -I/.../hpvm/llvm/include -I../include -I/.../hpvm/build/include -ffast-math -O3 -fno-lax-vector-conversions -fno-vectorize -fno-slp-vectorize -I/.../hpvm/llvm/include -I../include -I/.../hpvm/build/include -DDEVICE=GPU_TARGET -emit-llvm -S -o build/main.ll src/main.cc
+```
+
+`opt` is used to invoke the GenVISC pass, which converts the HPVM function calls to LLVM intrinsics.
+```
+/.../hpvm/build/bin/opt -debug-only=genvisc -load LLVMGenVISC.so -genvisc -globaldce -visc-timers-gen build/main.ll -S -o build/main.visc.ll
+```
+
+`opt` is used again to invoke the BuildDFG pass, which converts the textual representation to the internal HPVM representation.
+```
+/.../hpvm/build/bin/opt -debug -load LLVMBuildDFG.so -load LLVMLocalMem.so -load LLVMDFG2LLVM_NVPTX.so -load LLVMDFG2LLVM_X86.so -load LLVMClearDFG.so -localmem -dfg2llvm-nvptx -dfg2llvm-x86 -clearDFG -visc-timers-x86 -visc-timers-ptx -S build/main.visc.ll -o build/pipeline-gpu.host.ll
+```
+
+`llvm-cbe` is a C backend for LLVM. It is used here to create the OpenCL kernel.
+```
+/.../hpvm/build/bin/llvm-cbe -debug build/gpu/main.visc.ll.kernels.ll -o build/gpu/main.visc.ll.kernels.cl
+```
+
+`clang` is used again to compile a separate source file that contains I/O code.
+```
+/.../hpvm/build/bin/clang -Isrc/ -I -I/.../hpvm/llvm/include -I../include -I/.../hpvm/build/include -ffast-math -O3 -fno-lax-vector-conversions -fno-vectorize -fno-slp-vectorize -I/.../hpvm/llvm/include -I../include -I/.../hpvm/build/include -emit-llvm -S -o build/gpu/io.ll src/io.cc
+```
+
+`llvm-link` is used to link against the VISC runtime.
+```
+/.../hpvm/build/bin/llvm-link build/gpu/pipeline-gpu.host.ll build/gpu/io.ll /.../hpvm/llvm/tools/hpvm/projects/visc-rt/visc-rt.ll -S -o build/gpu/pipeline-gpu.linked.ll
+```
+
+`clang++` is used to do the final linking against OpenCL and emit the binary.
+```
+/.../hpvm/build/bin/clang++ -O3 `pkg-config opencv --libs` -lm -lpthread -lrt -lOpenCL -L/software/cuda-9.1/lib64 build/gpu/pipeline-gpu.linked.ll -o pipeline-gpu
+```
\ No newline at end of file