update

benchmarks/opencl/nearn/Makefile

@@ -0,0 +1,68 @@
+RISCV_TOOL_PATH = $(wildcard ~/dev/riscv-gnu-toolchain/drops)
+POCL_CC_PATH = $(wildcard ~/dev/pocl/drops_riscv_cc)
+POCL_INC_PATH = $(wildcard ../include)
+POCL_LIB_PATH = $(wildcard ../lib)
+VX_RT_PATH = $(wildcard ../../../runtime)
+VX_SIMX_PATH = $(wildcard ../../../simX/obj_dir)
+
+CC  = $(RISCV_TOOL_PATH)/bin/riscv32-unknown-elf-gcc
+CXX = $(RISCV_TOOL_PATH)/bin/riscv32-unknown-elf-g++
+DMP = $(RISCV_TOOL_PATH)/bin/riscv32-unknown-elf-objdump
+HEX = $(RISCV_TOOL_PATH)/bin/riscv32-unknown-elf-objcopy
+GDB = $(RISCV_TOOL_PATH)/bin/riscv32-unknown-elf-gdb
+
+VX_SRCS =  $(VX_RT_PATH)/newlib/newlib.c
+VX_SRCS += $(VX_RT_PATH)/startup/vx_start.s
+VX_SRCS += $(VX_RT_PATH)/intrinsics/vx_intrinsics.s
+VX_SRCS += $(VX_RT_PATH)/io/vx_io.s $(VX_RT_PATH)/io/vx_io.c
+VX_SRCS += $(VX_RT_PATH)/fileio/fileio.s
+VX_SRCS += $(VX_RT_PATH)/tests/tests.c
+VX_SRCS += $(VX_RT_PATH)/vx_api/vx_api.c
+VX_SRCS += $(VX_STR) $(VX_FIO) $(VX_NEWLIB) $(VX_INT) $(VX_IO) $(VX_API) $(VX_TEST)
+
+VX_CFLAGS = -nostartfiles -Wl,-Bstatic,-T,$(VX_RT_PATH)/mains/vortex_link.ld
+
+CXXFLAGS =  -g -O0 -march=rv32im -mabi=ilp32 
+CXXFLAGS += -ffreestanding # program may not begin at main()
+CXXFLAGS += -Wl,--gc-sections # enable garbage collection of unused input sections
+CXXFLAGS += -fno-rtti -fno-non-call-exceptions # disable RTTI and exceptions
+CXXFLAGS += -I$(POCL_INC_PATH)
+
+VX_LIBS = -Wl,--whole-archive lib$(PROJECT).a -Wl,--no-whole-archive $(POCL_LIB_PATH)/libOpenCL.a
+QEMU_LIBS = -Wl,--whole-archive lib$(PROJECT).a -Wl,--no-whole-archive $(POCL_LIB_PATH)/qemu/libOpenCL.a
+
+PROJECT = nearn
+
+SRCS = main.cc clutils.cpp utils.cpp
+
+all: $(PROJECT).dump $(PROJECT).hex
+
+lib$(PROJECT).a: kernel.cl
+	POCL_DEBUG=all POCL_DEBUG_LLVM_PASSES=1 LD_LIBRARY_PATH=$(RISCV_TOOL_PATH)/lib:$(POCL_CC_PATH)/lib $(POCL_CC_PATH)/bin/poclcc -o lib$(PROJECT).a kernel.cl
+
+$(PROJECT).elf: $(SRCS) lib$(PROJECT).a
+	$(CXX) $(CXXFLAGS) $(VX_CFLAGS) $(VX_SRCS) $(SRCS) $(VX_LIBS) -o $(PROJECT).elf
+
+$(PROJECT).qemu: $(SRCS) lib$(PROJECT).a
+	$(CXX) $(CXXFLAGS) $(SRCS)  $(QEMU_LIBS) -o $(PROJECT).qemu
+
+$(PROJECT).hex: $(PROJECT).elf
+	$(HEX) -O ihex $(PROJECT).elf $(PROJECT).hex
+
+$(PROJECT).dump: $(PROJECT).elf
+	$(DMP) -D $(PROJECT).elf > $(PROJECT).dump
+
+run: $(PROJECT).hex
+	POCL_DEBUG=all $(VX_SIMX_PATH)/Vcache_simX -E -a rv32i --core $(PROJECT).hex -s -b 1> emulator.debug
+
+qemu: $(PROJECT).qemu
+	POCL_DEBUG=all $(RISCV_TOOL_PATH)/bin/qemu-riscv32 -d in_asm -D debug.log $(PROJECT).qemu
+
+gdb-s: $(PROJECT).qemu
+	POCL_DEBUG=all $(RISCV_TOOL_PATH)/bin/qemu-riscv32 -g 1234 -d in_asm -D debug.log $(PROJECT).qemu
+
+gdb-c: $(PROJECT).qemu
+	$(GDB) $(PROJECT).qemu
+
+clean:
+	rm -rf *.o *.elf *.dump *.hex *.qemu *.log *.debug

benchmarks/opencl/nearn/README.txt

@@ -0,0 +1,33 @@
+The Nearest Neighbor application computes the nearest location to a specific 
+latitude and longitude for a number of hurricanes (data from: http://weather.unisys.com/hurricane/).
+
+The Makefile may need to be adjusted for different machines, but it was written for Mac OS X and
+Linux with either NVIDIA or AMD OpenCL SDKs.
+
+The hurricane data is located in a number of data files that are copied into the working
+directory by the Makefile.  A separate text file lists the names of the data files that
+will be used, and it is this text file that should be passed to the application (see usage, below).
+
+Nearest Neighbor Usage
+
+nearestNeighbor [filename] -r [int] -lat [float] -lng [float] [-hqt] [-p [int] -d [int]]
+
+example:
+$ ./nearestNeighbor filelist.txt -r 5 -lat 30 -lng 90
+
+filename     the filename that lists the data input files
+-r [int]     the number of records to return (default: 10)
+-lat [float] the latitude for nearest neighbors (default: 0)
+-lng [float] the longitude for nearest neighbors (default: 0)
+
+-h, --help   Display the help file
+-q           Quiet mode. Suppress all text output.
+-t           Print timing information.
+
+-p [int]     Choose the platform (must choose both platform and device)
+-d [int]     Choose the device (must choose both platform and device)
+
+
+Notes: 1. The filename is required as the first parameter.
+       2. If you declare either the device or the platform,
+          you must declare both.

benchmarks/opencl/nearn/clutils.h

@@ -0,0 +1,281 @@
+/****************************************************************************\ 
+ * Copyright (c) 2011, Advanced Micro Devices, Inc.                           *
+ * All rights reserved.                                                       *
+ *                                                                            *
+ * Redistribution and use in source and binary forms, with or without         *
+ * modification, are permitted provided that the following conditions         *
+ * are met:                                                                   *
+ *                                                                            *
+ * Redistributions of source code must retain the above copyright notice,     *
+ * this list of conditions and the following disclaimer.                      *
+ *                                                                            *
+ * Redistributions in binary form must reproduce the above copyright notice,  *
+ * this list of conditions and the following disclaimer in the documentation  *
+ * and/or other materials provided with the distribution.                     *
+ *                                                                            *
+ * Neither the name of the copyright holder nor the names of its contributors *
+ * may be used to endorse or promote products derived from this software      *
+ * without specific prior written permission.                                 *
+ *                                                                            *
+ * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS        *
+ * "AS IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED  *
+ * TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR *
+ * PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR          *
+ * CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL,      *
+ * EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO,        *
+ * PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR         *
+ * PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF     *
+ * LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING       *
+ * NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS         *
+ * SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.               *
+ *                                                                            *
+ * If you use the software (in whole or in part), you shall adhere to all     *
+ * applicable U.S., European, and other export laws, including but not        *
+ * limited to the U.S. Export Administration Regulations (“EAR”), (15 C.F.R.  *
+ * Sections 730 through 774), and E.U. Council Regulation (EC) No 1334/2000   *
+ * of 22 June 2000.  Further, pursuant to Section 740.6 of the EAR, you       *
+ * hereby certify that, except pursuant to a license granted by the United    *
+ * States Department of Commerce Bureau of Industry and Security or as        *
+ * otherwise permitted pursuant to a License Exception under the U.S. Export  *
+ * Administration Regulations ("EAR"), you will not (1) export, re-export or  *
+ * release to a national of a country in Country Groups D:1, E:1 or E:2 any   *
+ * restricted technology, software, or source code you receive hereunder,     *
+ * or (2) export to Country Groups D:1, E:1 or E:2 the direct product of such *
+ * technology or software, if such foreign produced direct product is subject *
+ * to national security controls as identified on the Commerce Control List   *
+ *(currently found in Supplement 1 to Part 774 of EAR).  For the most current *
+ * Country Group listings, or for additional information about the EAR or     *
+ * your obligations under those regulations, please refer to the U.S. Bureau  *
+ * of Industry and Security’s website at http://www.bis.doc.gov/.             *
+ \****************************************************************************/
+
+#ifndef __CL_UTILS_H__
+#define __CL_UTILS_H__
+
+#include <CL/cl.h>
+
+// The cl_time type is OS specific
+#ifdef _WIN32
+#include <tchar.h>
+#include <Windows.h>
+typedef __int64 cl_time; 
+#else
+#include <sys/time.h>
+typedef double cl_time;
+#endif
+
+//-------------------------------------------------------
+// Initialization and Cleanup
+//-------------------------------------------------------
+
+// Detects platforms and devices, creates context and command queue
+cl_context cl_init(char devicePreference='\0');
+
+// Creates a context given a platform and a device
+cl_context cl_init_context(int platform,int dev,int quiet=0);
+
+// Releases resources used by clutils
+void    cl_cleanup();
+
+// Releases a kernel object
+void    cl_freeKernel(cl_kernel kernel);
+
+// Releases a memory object
+void    cl_freeMem(cl_mem mem);
+
+// Releases a program object
+void    cl_freeProgram(cl_program program);
+
+// Returns the global command queue
+cl_command_queue cl_getCommandQueue();
+
+
+//-------------------------------------------------------
+// Synchronization functions
+//-------------------------------------------------------
+
+// Performs a clFinish on the command queue
+void    cl_sync();
+
+
+//-------------------------------------------------------
+// Memory allocation
+//-------------------------------------------------------
+
+// Allocates a regular buffer on the device
+cl_mem  cl_allocBuffer(size_t mem_size, 
+            cl_mem_flags flags = CL_MEM_READ_WRITE);
+
+// XXX I don't think this does exactly what we want it to do
+// Allocates a read-only buffer and transfers the data
+cl_mem  cl_allocBufferConst(size_t mem_size, void* host_ptr);
+
+// Allocates pinned memory on the host
+cl_mem  cl_allocBufferPinned(size_t mem_size);
+
+// Allocates an image on the device
+cl_mem  cl_allocImage(size_t height, size_t width, char type, 
+            cl_mem_flags flags = CL_MEM_READ_WRITE);
+
+
+
+//-------------------------------------------------------
+// Data transfers
+//-------------------------------------------------------
+
+// Copies a buffer from the device to pinned memory on the host and 
+// maps it so it can be read
+void*   cl_copyAndMapBuffer(cl_mem dst, cl_mem src, size_t size); 
+
+// Copies from one buffer to another
+void    cl_copyBufferToBuffer(cl_mem dst, cl_mem src, size_t size);
+
+// Copies data to a buffer on the device
+void    cl_copyBufferToDevice(cl_mem dst, void *src, size_t mem_size, 
+            cl_bool blocking = CL_TRUE);
+
+// Copies data to an image on the device
+void cl_copyImageToDevice(cl_mem dst, void* src, size_t height, size_t width);
+
+// Copies an image from the device to the host
+void    cl_copyImageToHost(void* dst, cl_mem src, size_t height, size_t width);
+
+// Copies data from a device buffer to the host
+void    cl_copyBufferToHost(void *dst, cl_mem src, size_t mem_size, 
+            cl_bool blocking = CL_TRUE);
+
+// Copies data from a buffer on the device to an image on the device
+void    cl_copyBufferToImage(cl_mem src, cl_mem dst, int height, int width);
+
+// Maps a buffer
+void*   cl_mapBuffer(cl_mem mem, size_t mem_size, cl_mem_flags flags);
+
+// Unmaps a buffer
+void    cl_unmapBuffer(cl_mem mem, void *ptr);
+
+// Writes data to a zero-copy buffer on the device
+void    cl_writeToZCBuffer(cl_mem mem, void* data, size_t size);
+
+//-------------------------------------------------------
+// Program and kernels
+//-------------------------------------------------------
+
+// Compiles a program
+cl_program  cl_compileProgram(char* kernelPath, char* compileoptions, 
+                bool verboseoptions = 0);
+
+// Creates a kernel
+cl_kernel   cl_createKernel(cl_program program, const char* kernelName);
+
+
+// Sets a kernel argument
+void        cl_setKernelArg(cl_kernel kernel, unsigned int index, size_t size, 
+                void* data);
+
+
+//-------------------------------------------------------
+// Profiling/events
+//-------------------------------------------------------
+
+// Computes the execution time (start to end) for an event
+double  cl_computeExecTime(cl_event);
+
+// Compute the elapsed time between two CPU timer values
+double  cl_computeTime(cl_time start, cl_time end); 
+
+// Creates an event from CPU timers
+void    cl_createUserEvent(cl_time start, cl_time end, char* desc);
+
+// Disable logging of events
+void    cl_disableEvents();
+
+// Enable logging of events
+void    cl_enableEvents();
+
+// Query the current system time
+void    cl_getTime(cl_time* time); 
+
+// Calls a function which prints events to the terminal
+void    cl_printEvents();
+
+// Calls a function which writes the events to a file
+void    cl_writeEventsToFile(char* path);
+
+
+//-------------------------------------------------------
+// Error handling
+//-------------------------------------------------------
+
+// Compare a status value to CL_SUCCESS and optionally exit on error
+int     cl_errChk(const cl_int status, const char *msg, bool exitOnErr);
+
+// Queries the supported image formats for the device and prints
+// them to the screen
+void    printSupportedImageFormats();
+
+//-------------------------------------------------------
+// Platform and device information
+//-------------------------------------------------------
+
+bool    cl_deviceIsAMD(cl_device_id dev=NULL);
+bool    cl_deviceIsNVIDIA(cl_device_id dev=NULL);
+bool    cl_platformIsNVIDIA(cl_platform_id plat=NULL);
+char*   cl_getDeviceDriverVersion(cl_device_id dev=NULL);
+char*   cl_getDeviceName(cl_device_id dev=NULL);
+char*   cl_getDeviceVendor(cl_device_id dev=NULL);
+char*   cl_getDeviceVersion(cl_device_id dev=NULL);
+char*   cl_getPlatformName(cl_platform_id platform);
+char*   cl_getPlatformVendor(cl_platform_id platform);
+
+//-------------------------------------------------------
+// Utility functions
+//-------------------------------------------------------
+
+char* catStringWithInt(const char* str, int integer);
+
+char* itoa_portable(int value, char* result, int base);
+
+//-------------------------------------------------------
+// Data types
+//-------------------------------------------------------
+typedef struct{
+    int x;
+    int y;
+} int2;
+
+typedef struct{
+    float x;
+    float y;
+}float2;
+
+typedef struct{
+    float x;
+    float y;
+    float z;
+    float w;
+}float4;
+
+//-------------------------------------------------------
+// Defines
+//-------------------------------------------------------
+
+#define MAX_ERR_VAL 64
+
+#define NUM_PROGRAMS 7
+
+#define NUM_KERNELS 13
+#define KERNEL_INIT_DET 0 
+#define KERNEL_BUILD_DET 1 
+#define KERNEL_SURF_DESC 2
+#define KERNEL_NORM_DESC 3
+#define KERNEL_NON_MAX_SUP 4
+#define KERNEL_GET_ORIENT1 5
+#define KERNEL_GET_ORIENT2 6
+#define KERNEL_NN 7
+#define KERNEL_SCAN 8
+#define KERNEL_SCAN4 9
+#define KERNEL_TRANSPOSE 10
+#define KERNEL_SCANIMAGE 11
+#define KERNEL_TRANSPOSEIMAGE 12
+
+#endif

benchmarks/opencl/nearn/filelist.txt

@@ -0,0 +1,4 @@
+cane4_0.db
+cane4_1.db
+cane4_2.db
+cane4_3.db

benchmarks/opencl/nearn/gettimeofday.cpp

@@ -0,0 +1,74 @@
+#include "stdio.h"
+#include <time.h>
+#include <windows.h>
+#include <iostream> 
+//using namespace System;
+using namespace std;
+ 
+#if defined(_MSC_VER) || defined(_MSC_EXTENSIONS)
+  #define DELTA_EPOCH_IN_MICROSECS  11644473600000000Ui64
+#else
+  #define DELTA_EPOCH_IN_MICROSECS  11644473600000000ULL
+#endif
+ 
+struct timezone
+{
+  int  tz_minuteswest; /* minutes W of Greenwich */
+  int  tz_dsttime;     /* type of dst correction */
+};
+ 
+
+// Definition of a gettimeofday function
+ int gettimeofday(struct timeval *tv, struct timezone *tz)
+{
+// Define a structure to receive the current Windows filetime
+  FILETIME ft;
+ 
+// Initialize the present time to 0 and the timezone to UTC
+  unsigned __int64 tmpres = 0;
+  static int tzflag = 0;
+ 
+  if (NULL != tv)
+  {
+    GetSystemTimeAsFileTime(&ft);
+ 
+// The GetSystemTimeAsFileTime returns the number of 100 nanosecond 
+// intervals since Jan 1, 1601 in a structure. Copy the high bits to 
+// the 64 bit tmpres, shift it left by 32 then or in the low 32 bits.
+    tmpres |= ft.dwHighDateTime;
+    tmpres <<= 32;
+    tmpres |= ft.dwLowDateTime;
+ 
+// Convert to microseconds by dividing by 10
+    tmpres /= 10;
+ 
+// The Unix epoch starts on Jan 1 1970.  Need to subtract the difference 
+// in seconds from Jan 1 1601.
+    tmpres -= DELTA_EPOCH_IN_MICROSECS;
+ 
+// Finally change microseconds to seconds and place in the seconds value. 
+// The modulus picks up the microseconds.
+    tv->tv_sec = (long)(tmpres / 1000000UL);
+    tv->tv_usec = (long)(tmpres % 1000000UL);
+  }
+ 
+  if (NULL != tz)
+  {
+    if (!tzflag)
+    {
+      _tzset();
+      tzflag++;
+    }
+  
+// Adjust for the timezone west of Greenwich
+	long seconds_diff;
+	_get_timezone(&seconds_diff);
+      tz->tz_minuteswest = seconds_diff / 60;
+	int hours_offset;
+	_get_daylight(&hours_offset);
+    tz->tz_dsttime = hours_offset;
+  }
+ 
+  return 0;
+}
+    

benchmarks/opencl/nearn/gettimeofday.h

@@ -0,0 +1,17 @@
+
+#ifdef _WIN32
+#include <WinSock.h>
+/**
+Based on code seen at.
+
+http://www.winehq.org/pipermail/wine-devel/2003-June/018082.html
+
+http://msdn.microsoft.com/en-us/library/ms740560
+
+*/
+int gettimeofday(struct timeval *tv, struct timezone *tz);
+#else
+#include <sys/time.h>
+#endif
+
+

benchmarks/opencl/nearn/ipoint.h

@@ -0,0 +1,29 @@
+/*********************************************************** 
+*  --- OpenSURF ---                                        *
+*  This library is distributed under the GNU GPL. Please   *
+*  contact chris.evans@irisys.co.uk for more information.  *
+*                                                          *
+*  C. Evans, Research Into Robust Visual Features,         *
+*  MSc University of Bristol, 2008.                        *
+*                                                          *
+************************************************************/
+
+#ifndef IPOINT_H
+#define IPOINT_H
+
+#include <vector>
+#include <math.h>
+
+
+
+//-------------------------------------------------------
+typedef struct{
+        int x;
+        int y;
+		float descriptor[64];
+	} Ipoint;
+
+//-------------------------------------------------------
+
+  typedef std::vector<Ipoint> IpVec;
+#endif

benchmarks/opencl/nearn/kernel.cl

@@ -0,0 +1,22 @@
+//#pragma OPENCL EXTENSION cl_khr_byte_addressable_store : enable
+
+typedef struct latLong
+    {
+        float lat;
+        float lng;
+    } LatLong;
+
+__kernel void NearestNeighbor(__global LatLong *d_locations,
+							  __global float *d_distances,
+							  const int numRecords,
+							  const float lat,
+							  const float lng) {
+	 int globalId = get_global_id(0);
+							  
+     if (globalId < numRecords) {
+         __global LatLong *latLong = d_locations+globalId;
+    
+         __global float *dist=d_distances+globalId;
+         *dist = (float)sqrt((lat-latLong->lat)*(lat-latLong->lat)+(lng-latLong->lng)*(lng-latLong->lng));
+	 }
+}

benchmarks/opencl/nearn/main.cc

@@ -0,0 +1,346 @@
+#ifndef __NEAREST_NEIGHBOR__
+#define __NEAREST_NEIGHBOR__
+
+#include "nearestNeighbor.h"
+
+cl_context context = NULL;
+
+int main(int argc, char *argv[]) {
+  std::vector<Record> records;
+  float *recordDistances;
+  // LatLong locations[REC_WINDOW];
+  std::vector<LatLong> locations;
+  int i;
+  // args
+  char filename[100];
+  int resultsCount = 5, quiet = 0, timing = 0, platform = -1, device = -1;
+  float lat = 30, lng = 90;
+
+  // parse command line
+  if (parseCommandline(argc, argv, filename, &resultsCount, &lat, &lng, &quiet,
+                       &timing, &platform, &device)) {
+    printUsage();
+    return 0;
+  }
+
+  int numRecords = loadData(filename, records, locations);
+
+  // for(i=0;i<numRecords;i++)
+  //    printf("%s, %f,
+  //    %f\n",(records[i].recString),locations[i].lat,locations[i].lng);
+
+  printf("Number of records: %d\n", numRecords);
+  printf("Finding the %d closest neighbors.\n", resultsCount);
+
+  if (resultsCount > numRecords)
+    resultsCount = numRecords;
+
+  context = cl_init_context(platform, device, quiet);
+  
+  recordDistances = OpenClFindNearestNeighbors(context, numRecords, locations,
+                                               lat, lng, timing);
+
+  // find the resultsCount least distances
+  findLowest(records, recordDistances, numRecords, resultsCount);
+
+  // print out results
+  if (!quiet)
+    for (i = 0; i < resultsCount; i++) {
+      printf("%s --> Distance=%f\n", records[i].recString, records[i].distance);
+    }
+  free(recordDistances);
+  return 0;
+}
+
+float *OpenClFindNearestNeighbors(cl_context context, int numRecords,
+                                  std::vector<LatLong> &locations, float lat,
+                                  float lng, int timing) {
+
+  // 1. set up kernel
+  cl_kernel NN_kernel;
+  cl_int status;
+  cl_program cl_NN_program;
+  cl_NN_program = cl_compileProgram((char *)"nearestNeighbor_kernel.cl", NULL);
+
+  NN_kernel = clCreateKernel(cl_NN_program, "NearestNeighbor", &status);
+  status =
+      cl_errChk(status, (char *)"Error Creating Nearest Neighbor kernel", true);
+  if (status)
+    exit(1);
+  // 2. set up memory on device and send ipts data to device
+  // copy ipts(1,2) to device
+  // also need to alloate memory for the distancePoints
+  cl_mem d_locations;
+  cl_mem d_distances;
+
+  cl_int error = 0;
+
+  d_locations = clCreateBuffer(context, CL_MEM_READ_ONLY,
+                               sizeof(LatLong) * numRecords, NULL, &error);
+
+  d_distances = clCreateBuffer(context, CL_MEM_READ_WRITE,
+                               sizeof(float) * numRecords, NULL, &error);
+
+  cl_command_queue command_queue = cl_getCommandQueue();
+  cl_event writeEvent, kernelEvent, readEvent;
+  error = clEnqueueWriteBuffer(command_queue, d_locations,
+                               1, // change to 0 for nonblocking write
+                               0, // offset
+                               sizeof(LatLong) * numRecords, &locations[0], 0,
+                               NULL, &writeEvent);
+
+  // 3. send arguments to device
+  cl_int argchk;
+  argchk = clSetKernelArg(NN_kernel, 0, sizeof(cl_mem), (void *)&d_locations);
+  argchk |= clSetKernelArg(NN_kernel, 1, sizeof(cl_mem), (void *)&d_distances);
+  argchk |= clSetKernelArg(NN_kernel, 2, sizeof(int), (void *)&numRecords);
+  argchk |= clSetKernelArg(NN_kernel, 3, sizeof(float), (void *)&lat);
+  argchk |= clSetKernelArg(NN_kernel, 4, sizeof(float), (void *)&lng);
+
+  cl_errChk(argchk, "ERROR in Setting Nearest Neighbor kernel args", true);
+
+  // 4. enqueue kernel
+  size_t globalWorkSize[1];
+  globalWorkSize[0] = numRecords;
+  if (numRecords % 64)
+    globalWorkSize[0] += 64 - (numRecords % 64);
+  // printf("Global Work Size: %zu\n",globalWorkSize[0]);
+
+  error = clEnqueueNDRangeKernel(command_queue, NN_kernel, 1, 0, globalWorkSize,
+                                 NULL, 0, NULL, &kernelEvent);
+
+  cl_errChk(error, "ERROR in Executing Kernel NearestNeighbor", true);
+
+  // 5. transfer data off of device
+
+  // create distances std::vector
+  float *distances = (float *)malloc(sizeof(float) * numRecords);
+
+  error = clEnqueueReadBuffer(command_queue, d_distances,
+                              1, // change to 0 for nonblocking write
+                              0, // offset
+                              sizeof(float) * numRecords, distances, 0, NULL,
+                              &readEvent);
+
+  cl_errChk(error, "ERROR with clEnqueueReadBuffer", true);
+  if (timing) {
+    clFinish(command_queue);
+    cl_ulong eventStart, eventEnd, totalTime = 0;
+    printf("# Records\tWrite(s) [size]\t\tKernel(s)\tRead(s)  "
+           "[size]\t\tTotal(s)\n");
+    printf("%d        \t", numRecords);
+    // Write Buffer
+    error = clGetEventProfilingInfo(writeEvent, CL_PROFILING_COMMAND_START,
+                                    sizeof(cl_ulong), &eventStart, NULL);
+    cl_errChk(error, "ERROR in Event Profiling (Write Start)", true);
+    error = clGetEventProfilingInfo(writeEvent, CL_PROFILING_COMMAND_END,
+                                    sizeof(cl_ulong), &eventEnd, NULL);
+    cl_errChk(error, "ERROR in Event Profiling (Write End)", true);
+
+    printf("%f [%.2fMB]\t", (float)((eventEnd - eventStart) / 1e9),
+           (float)((sizeof(LatLong) * numRecords) / 1e6));
+    totalTime += eventEnd - eventStart;
+    // Kernel
+    error = clGetEventProfilingInfo(kernelEvent, CL_PROFILING_COMMAND_START,
+                                    sizeof(cl_ulong), &eventStart, NULL);
+    cl_errChk(error, "ERROR in Event Profiling (Kernel Start)", true);
+    error = clGetEventProfilingInfo(kernelEvent, CL_PROFILING_COMMAND_END,
+                                    sizeof(cl_ulong), &eventEnd, NULL);
+    cl_errChk(error, "ERROR in Event Profiling (Kernel End)", true);
+
+    printf("%f\t", (float)((eventEnd - eventStart) / 1e9));
+    totalTime += eventEnd - eventStart;
+    // Read Buffer
+    error = clGetEventProfilingInfo(readEvent, CL_PROFILING_COMMAND_START,
+                                    sizeof(cl_ulong), &eventStart, NULL);
+    cl_errChk(error, "ERROR in Event Profiling (Read Start)", true);
+    error = clGetEventProfilingInfo(readEvent, CL_PROFILING_COMMAND_END,
+                                    sizeof(cl_ulong), &eventEnd, NULL);
+    cl_errChk(error, "ERROR in Event Profiling (Read End)", true);
+
+    printf("%f [%.2fMB]\t", (float)((eventEnd - eventStart) / 1e9),
+           (float)((sizeof(float) * numRecords) / 1e6));
+    totalTime += eventEnd - eventStart;
+
+    printf("%f\n\n", (float)(totalTime / 1e9));
+  }
+  // 6. return finalized data and release buffers
+  clReleaseMemObject(d_locations);
+  clReleaseMemObject(d_distances);
+  return distances;
+}
+
+int loadData(char *filename, std::vector<Record> &records,
+             std::vector<LatLong> &locations) {
+  FILE *flist, *fp;
+  int i = 0;
+  char dbname[64];
+  int recNum = 0;
+
+  /**Main processing **/
+
+  int q = 0;
+
+  flist = fopen(filename, "r");
+  while (!feof(flist)) {
+    /**
+    * Read in REC_WINDOW records of length REC_LENGTH
+    * If this is the last file in the filelist, then done
+    * else open next file to be read next iteration
+    */       
+    if (fscanf(flist, "%s\n", dbname) != 1) {
+      printf("error reading filelist\n");
+      exit(0);
+    }
+    printf("loading db: %s\n", dbname);
+    fp = fopen(dbname, "r");
+    if (!fp) {
+      printf("error opening a db\n");
+      exit(1);
+    }    
+    // read each record
+    while (!feof(fp)) {
+      Record record;
+      LatLong latLong;
+      fgets(record.recString, 49, fp);
+      fgetc(fp); // newline
+      if (feof(fp))
+        break;
+
+      // parse for lat and long
+      char substr[6];
+
+      for (i = 0; i < 5; i++)
+        substr[i] = *(record.recString + i + 28);
+      substr[5] = '\0';
+      latLong.lat = atof(substr);
+
+      for (i = 0; i < 5; i++)
+        substr[i] = *(record.recString + i + 33);
+      substr[5] = '\0';
+      latLong.lng = atof(substr);
+
+      locations.push_back(latLong);
+      records.push_back(record);
+      recNum++;
+      if (0 == (recNum % 500))
+        break;
+    }
+    
+    if (++q == 3)
+        break;
+    fclose(fp);
+  }
+  fclose(flist);
+  return recNum;
+}
+
+void findLowest(std::vector<Record> &records, float *distances, int numRecords,
+                int topN) {
+  int i, j;
+  float val;
+  int minLoc;
+  Record *tempRec;
+  float tempDist;
+
+  for (i = 0; i < topN; i++) {
+    minLoc = i;
+    for (j = i; j < numRecords; j++) {
+      val = distances[j];
+      if (val < distances[minLoc])
+        minLoc = j;
+    }
+    // swap locations and distances
+    tempRec = &records[i];
+    records[i] = records[minLoc];
+    records[minLoc] = *tempRec;
+
+    tempDist = distances[i];
+    distances[i] = distances[minLoc];
+    distances[minLoc] = tempDist;
+
+    // add distance to the min we just found
+    records[i].distance = distances[i];
+  }
+}
+
+int parseCommandline(int argc, char *argv[], char *filename, int *r, float *lat,
+                     float *lng, int *q, int *t, int *p, int *d) {
+  int i;
+  // if (argc < 2) return 1; // error
+  strncpy(filename, "filelist.txt", 100);
+  char flag;
+
+  for (i = 1; i < argc; i++) {
+    if (argv[i][0] == '-') { // flag
+      flag = argv[i][1];
+      switch (flag) {
+      case 'r': // number of results
+        i++;
+        *r = atoi(argv[i]);
+        break;
+      case 'l':                  // lat or lng
+        if (argv[i][2] == 'a') { // lat
+          *lat = atof(argv[i + 1]);
+        } else { // lng
+          *lng = atof(argv[i + 1]);
+        }
+        i++;
+        break;
+      case 'h': // help
+        return 1;
+        break;
+      case 'q': // quiet
+        *q = 1;
+        break;
+      case 't': // timing
+        *t = 1;
+        break;
+      case 'p': // platform
+        i++;
+        *p = atoi(argv[i]);
+        break;
+      case 'd': // device
+        i++;
+        *d = atoi(argv[i]);
+        break;
+      }
+    }
+  }
+  if ((*d >= 0 && *p < 0) ||
+      (*p >= 0 &&
+       *d < 0)) // both p and d must be specified if either are specified
+    return 1;
+  return 0;
+}
+
+void printUsage() {
+  printf("Nearest Neighbor Usage\n");
+  printf("\n");
+  printf("nearestNeighbor [filename] -r [int] -lat [float] -lng [float] [-hqt] "
+         "[-p [int] -d [int]]\n");
+  printf("\n");
+  printf("example:\n");
+  printf("$ ./nearestNeighbor filelist.txt -r 5 -lat 30 -lng 90\n");
+  printf("\n");
+  printf("filename     the filename that lists the data input files\n");
+  printf("-r [int]     the number of records to return (default: 10)\n");
+  printf("-lat [float] the latitude for nearest neighbors (default: 0)\n");
+  printf("-lng [float] the longitude for nearest neighbors (default: 0)\n");
+  printf("\n");
+  printf("-h, --help   Display the help file\n");
+  printf("-q           Quiet mode. Suppress all text output.\n");
+  printf("-t           Print timing information.\n");
+  printf("\n");
+  printf("-p [int]     Choose the platform (must choose both platform and "
+         "device)\n");
+  printf("-d [int]     Choose the device (must choose both platform and "
+         "device)\n");
+  printf("\n");
+  printf("\n");
+  printf("Notes: 1. The filename is required as the first parameter.\n");
+  printf("       2. If you declare either the device or the platform,\n");
+  printf("          you must declare both.\n\n");
+}
+
+#endif

benchmarks/opencl/nearn/nearestNeighbor.h

@@ -0,0 +1,50 @@
+#ifndef _NEARESTNEIGHBOR
+#define _NEARESTNEIGHBOR
+
+#include <iostream>
+#include <vector>
+#include <float.h>
+#include <string.h>
+#include <stdio.h>
+#include <stdlib.h>
+
+// All OpenCL headers
+#if defined (__APPLE__) || defined(MACOSX)
+    #include <OpenCL/opencl.h>
+#else
+    #include <CL/opencl.h>
+#endif 
+
+#include "clutils.h"
+//#include "utils.h"
+
+#include <algorithm>
+
+
+	
+#define REC_LENGTH 49 // size of a record in db
+
+typedef struct latLong
+{
+  float lat;
+  float lng;
+} LatLong;
+
+typedef struct record
+{
+  char recString[REC_LENGTH];
+  float distance;
+} Record;
+
+float *OpenClFindNearestNeighbors(
+	cl_context context,
+	int numRecords,
+	std::vector<LatLong> &locations,float lat,float lng,
+	int timing);
+
+int loadData(char *filename,std::vector<Record> &records,std::vector<LatLong> &locations);
+void findLowest(std::vector<Record> &records,float *distances,int numRecords,int topN);
+void printUsage();
+int parseCommandline(int argc, char *argv[], char* filename,int *r,float *lat,float *lng,
+                     int *q, int *t, int *p, int *d);
+#endif

benchmarks/opencl/nearn/run

@@ -0,0 +1 @@
+./nn filelist.txt -r 5 -lat 30 -lng 90

benchmarks/opencl/nearn/utils.cpp

@@ -0,0 +1,204 @@
+/****************************************************************************\ 
+ * Copyright (c) 2011, Advanced Micro Devices, Inc.                           *
+ * All rights reserved.                                                       *
+ *                                                                            *
+ * Redistribution and use in source and binary forms, with or without         *
+ * modification, are permitted provided that the following conditions         *
+ * are met:                                                                   *
+ *                                                                            *
+ * Redistributions of source code must retain the above copyright notice,     *
+ * this list of conditions and the following disclaimer.                      *
+ *                                                                            *
+ * Redistributions in binary form must reproduce the above copyright notice,  *
+ * this list of conditions and the following disclaimer in the documentation  *
+ * and/or other materials provided with the distribution.                     *
+ *                                                                            *
+ * Neither the name of the copyright holder nor the names of its contributors *
+ * may be used to endorse or promote products derived from this software      *
+ * without specific prior written permission.                                 *
+ *                                                                            *
+ * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS        *
+ * "AS IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED  *
+ * TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR *
+ * PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR          *
+ * CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL,      *
+ * EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO,        *
+ * PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR         *
+ * PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF     *
+ * LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING       *
+ * NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS         *
+ * SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.               *
+ *                                                                            *
+ * If you use the software (in whole or in part), you shall adhere to all     *
+ * applicable U.S., European, and other export laws, including but not        *
+ * limited to the U.S. Export Administration Regulations (“EAR”), (15 C.F.R.  *
+ * Sections 730 through 774), and E.U. Council Regulation (EC) No 1334/2000   *
+ * of 22 June 2000.  Further, pursuant to Section 740.6 of the EAR, you       *
+ * hereby certify that, except pursuant to a license granted by the United    *
+ * States Department of Commerce Bureau of Industry and Security or as        *
+ * otherwise permitted pursuant to a License Exception under the U.S. Export  *
+ * Administration Regulations ("EAR"), you will not (1) export, re-export or  *
+ * release to a national of a country in Country Groups D:1, E:1 or E:2 any   *
+ * restricted technology, software, or source code you receive hereunder,     *
+ * or (2) export to Country Groups D:1, E:1 or E:2 the direct product of such *
+ * technology or software, if such foreign produced direct product is subject *
+ * to national security controls as identified on the Commerce Control List   *
+ *(currently found in Supplement 1 to Part 774 of EAR).  For the most current *
+ * Country Group listings, or for additional information about the EAR or     *
+ * your obligations under those regulations, please refer to the U.S. Bureau  *
+ * of Industry and Security’s website at http://www.bis.doc.gov/.             *
+ \****************************************************************************/
+
+#include <stdio.h>
+#include <sys/stat.h>
+#include <string.h>
+#include <stdlib.h>
+
+#include "utils.h"
+
+static bool usingImages = true;
+
+//! A wrapper for malloc that checks the return value
+void* alloc(size_t size) {
+
+    void* ptr = NULL;
+    ptr = malloc(size);
+    if(ptr == NULL) {
+        perror("malloc");
+        exit(-1);
+    }
+
+    return ptr;
+}
+
+// This function checks to make sure a file exists before we open it
+void checkFile(char* filename) 
+{
+    
+    struct stat fileStatus;
+    if(stat(filename, &fileStatus) != 0) {
+        printf("Error opening file: %s\n", filename);
+        exit(-1);
+    }
+    else {
+        if(!(S_IFREG & fileStatus.st_mode)) {
+            printf("File %s is not a regular file\n", filename);
+            exit(-1);
+        }
+    }
+}
+
+
+// This function checks to make sure a directory exists 
+void checkDir(char* dirpath) 
+{
+    
+    struct stat fileStatus;
+    if(stat(dirpath, &fileStatus) != 0) {
+        printf("Directory does not exist: %s\n", dirpath);
+        exit(-1);
+    }
+    else {
+        if(!(S_IFDIR & fileStatus.st_mode)) {
+            printf("Directory was not provided: %s\n", dirpath);
+            exit(-1);
+        }
+    }
+}
+
+// Parse the command line arguments
+void parseArguments(int argc, char** argv, char** input, char** events, 
+    char** ipts, char* devicePref, bool* verifyResults) 
+{
+    
+    for(int i = 2; i < argc; i++) {
+        if(strcmp(argv[i], "-d") == 0) {   // Event dump found
+            if(i == argc-1) {
+                printf("Usage: -e Needs directory path\n");
+                exit(-1);
+            }
+            devicePref[0] = argv[i+1][0];
+            i++;
+            continue;
+        }
+        if(strcmp(argv[i], "-e") == 0) {   // Event dump found
+            if(i == argc-1) {
+                printf("Usage: -e Needs directory path\n");
+                exit(-1);
+            }
+            *events = argv[i+1];
+            i++;
+            continue;
+        }
+        if(strcmp(argv[i], "-i") == 0) {   // Input found
+            if(i == argc-1) {
+                printf("Usage: -i Needs directory path\n");
+                exit(-1);
+            }
+            *input = argv[i+1];
+            i++;
+            continue;
+        }
+        if(strcmp(argv[i], "-l") == 0) {   // Ipts dump found
+            if(i == argc-1) {
+                printf("Usage: -l Needs directory path\n");
+                exit(-1);
+            }
+            *ipts = argv[i+1];
+            i++;
+            continue;
+        }
+        if(strcmp(argv[i], "-n") == 0) {   // Don't use OpenCL images
+            setUsingImages(false);
+            continue;
+        }
+        if(strcmp(argv[i], "-v") == 0) {   // Verify results
+            *verifyResults = true;
+            continue;
+        }
+    }
+}
+
+
+// This function that takes a positive integer 'value' and returns
+// the nearest multiple of 'multiple' (used for padding columns)
+unsigned int roundUp(unsigned int value, unsigned int multiple) {
+
+   unsigned int remainder = value % multiple;
+   
+   // Make the value a multiple of multiple
+   if(remainder != 0) {
+      value += (multiple-remainder);
+   }
+
+   return value;
+}
+
+
+// Concatenate two strings and return a pointer to the new string
+char* smartStrcat(char* str1, char* str2) 
+{
+    char* newStr = NULL;
+
+    newStr = (char*)alloc((strlen(str1)+strlen(str2)+1)*sizeof(char));
+
+    strcpy(newStr, str1);
+    strcat(newStr, str2);
+
+    return newStr;
+}
+
+
+// Set the value of using images to true if they are being
+// used, or false if they are not
+void setUsingImages(bool val) 
+{
+    usingImages = val;
+}
+
+
+// Return whether or not images are being used
+bool isUsingImages() 
+{
+    return usingImages;
+}

benchmarks/opencl/nearn/utils.h

@@ -0,0 +1,84 @@
+/****************************************************************************\ 
+ * Copyright (c) 2011, Advanced Micro Devices, Inc.                           *
+ * All rights reserved.                                                       *
+ *                                                                            *
+ * Redistribution and use in source and binary forms, with or without         *
+ * modification, are permitted provided that the following conditions         *
+ * are met:                                                                   *
+ *                                                                            *
+ * Redistributions of source code must retain the above copyright notice,     *
+ * this list of conditions and the following disclaimer.                      *
+ *                                                                            *
+ * Redistributions in binary form must reproduce the above copyright notice,  *
+ * this list of conditions and the following disclaimer in the documentation  *
+ * and/or other materials provided with the distribution.                     *
+ *                                                                            *
+ * Neither the name of the copyright holder nor the names of its contributors *
+ * may be used to endorse or promote products derived from this software      *
+ * without specific prior written permission.                                 *
+ *                                                                            *
+ * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS        *
+ * "AS IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED  *
+ * TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR *
+ * PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR          *
+ * CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL,      *
+ * EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO,        *
+ * PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR         *
+ * PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF     *
+ * LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING       *
+ * NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS         *
+ * SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.               *
+ *                                                                            *
+ * If you use the software (in whole or in part), you shall adhere to all     *
+ * applicable U.S., European, and other export laws, including but not        *
+ * limited to the U.S. Export Administration Regulations (“EAR”), (15 C.F.R.  *
+ * Sections 730 through 774), and E.U. Council Regulation (EC) No 1334/2000   *
+ * of 22 June 2000.  Further, pursuant to Section 740.6 of the EAR, you       *
+ * hereby certify that, except pursuant to a license granted by the United    *
+ * States Department of Commerce Bureau of Industry and Security or as        *
+ * otherwise permitted pursuant to a License Exception under the U.S. Export  *
+ * Administration Regulations ("EAR"), you will not (1) export, re-export or  *
+ * release to a national of a country in Country Groups D:1, E:1 or E:2 any   *
+ * restricted technology, software, or source code you receive hereunder,     *
+ * or (2) export to Country Groups D:1, E:1 or E:2 the direct product of such *
+ * technology or software, if such foreign produced direct product is subject *
+ * to national security controls as identified on the Commerce Control List   *
+ *(currently found in Supplement 1 to Part 774 of EAR).  For the most current *
+ * Country Group listings, or for additional information about the EAR or     *
+ * your obligations under those regulations, please refer to the U.S. Bureau  *
+ * of Industry and Security’s website at http://www.bis.doc.gov/.             *
+ \****************************************************************************/
+
+#ifndef _UTILS_
+#define _UTILS_
+
+// Wrapper for malloc
+void* alloc(size_t size);
+
+// Checks for existence of directory
+void checkDir(char* dirpath);
+
+// Check for existence of file
+void checkFile(char* filename);
+
+// Parse the input command line options to the program
+void parseArguments(int argc, char** argv, char** input, char** events, 
+    char** ipts, char* devicePref, bool* verifyResults);
+
+
+// Print the program usage information
+void printUsage();
+
+// Rounds up size to the nearest multiple of multiple
+unsigned int roundUp(unsigned int value, unsigned int multiple);
+
+// Concatenate two strings, creating a new one
+char* smartStrcat(char* str1, char* str2);
+
+// Set the value of usingImages
+void setUsingImages(bool val);
+
+// Return whether or not images are being used
+bool isUsingImages(); 
+
+#endif