diff --git a/benchmarks/opencl/hotspot/README b/benchmarks/opencl/hotspot/README new file mode 100644 index 00000000..e69de29b diff --git a/benchmarks/opencl/kmeans/README b/benchmarks/opencl/kmeans/README new file mode 100644 index 00000000..e69de29b diff --git a/benchmarks/opencl/sgemm/Makefile b/benchmarks/opencl/sgemm/Makefile new file mode 100644 index 00000000..7fa16540 --- /dev/null +++ b/benchmarks/opencl/sgemm/Makefile @@ -0,0 +1,54 @@ + +RISCV_TOOL_PATH=$(wildcard ~/dev/riscv-gnu-toolchain/drops) + +POCL_CC_PATH=$(wildcard ~/dev/pocl/drops_riscv_cc) +POCL_RT_PATH=$(wildcard ~/dev/pocl/drops_riscv_rt) + +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 +NEWLIB_PATH=$(RISCV_TOOL_PATH)/riscv32-unknown-elf/lib + +VX_NEWLIB = $(VX_RT_PATH)/newlib/newlib.c +VX_STR = $(VX_RT_PATH)/startup/vx_start.s +VX_INT = $(VX_RT_PATH)/intrinsics/vx_intrinsics.s +VX_IO = $(VX_RT_PATH)/io/vx_io.s $(VX_RT_PATH)/io/vx_io.c +VX_FIO = $(VX_RT_PATH)/fileio/fileio.s +VX_API = $(VX_RT_PATH)/vx_api/vx_api.c + +VX_SRCS = $(VX_STR) $(VX_FIO) $(NEWLIB) $(VX_INT) $(VX_IO) $(VX_API) $(VX_TEST) + +CXXFLAGS = -g -O0 -nostartfiles -Wl,-Bstatic,-T,$(VX_RT_PATH)/mains/nativevecadd/linker.ld -march=rv32im -mabi=ilp32 +CXXFLAGS += -ffreestanding # program may not begin at main() +CXXFLAGS += -fno-rtti -fno-non-call-exceptions # disable RTTI and exceptions +CXXFLAGS += -Wl,--gc-sections # enable garbage collection of unused input sections + +LIBS = -lOpenCL + +#$(NEWLIB_PATH)/libc.a $(NEWLIB_PATH)/libstdc++.a -static-libgcc -lgcc + +PROJECT=sgemm + +all: $(PROJECT).dump $(PROJECT).hex + +libsgemm.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 kernel.pocl kernel.cl + +$(PROJECT).elf: main.cc libsgemm.a + $(CXX) $(CXXFLAGS) -I$(POCL_RT_PATH)/include -L$(POCL_RT_PATH)/lib/static -L. $(VX_SRCS) main.cc -Wl,--whole-archive -lsgemm -Wl,--no-whole-archive $(LIBS) -o $(PROJECT).elf + +$(PROJECT).hex: $(PROJECT).elf + $(HEX) -O ihex $(PROJECT).elf $(PROJECT).hex + +$(PROJECT).dump: $(PROJECT).elf + $(DMP) -D $(PROJECT).elf > $(PROJECT).dump + +run: + $(VX_SIMX_PATH)/Vcache_simX -E -a rv32i --core $(PROJECT).hex -s -b 1> emulator.debug + +clean: + rm -rf *.elf *.dump *.hex *.a *.pocl diff --git a/benchmarks/opencl/sgemm/kernel.cl b/benchmarks/opencl/sgemm/kernel.cl new file mode 100644 index 00000000..17ece1d1 --- /dev/null +++ b/benchmarks/opencl/sgemm/kernel.cl @@ -0,0 +1,9 @@ +__kernel void sgemm(__global float *A, __global float *B, __global float *C, int ldc) +{ + long i = get_global_id(0); + long m = get_global_id(1); + long n = get_global_id(2); + float a = A[m+n*ldc]; + float b = B[m*ldc+i]; + C[i+n*ldc] = C[i+n*ldc] + a * b; +} diff --git a/benchmarks/opencl/sgemm/kernel.pocl b/benchmarks/opencl/sgemm/kernel.pocl new file mode 100644 index 00000000..6a643fae Binary files /dev/null and b/benchmarks/opencl/sgemm/kernel.pocl differ diff --git a/benchmarks/opencl/sgemm/libsgemm.a b/benchmarks/opencl/sgemm/libsgemm.a new file mode 100644 index 00000000..71719860 Binary files /dev/null and b/benchmarks/opencl/sgemm/libsgemm.a differ diff --git a/benchmarks/opencl/sgemm/main.cc b/benchmarks/opencl/sgemm/main.cc new file mode 100644 index 00000000..561a0f30 --- /dev/null +++ b/benchmarks/opencl/sgemm/main.cc @@ -0,0 +1,500 @@ +/* + * Simple OpenCL demo program + * + * Copyright (C) 2009 Clifford Wolf + * + * This program is free software; you can redistribute it and/or modify + * it under the terms of the GNU General Public License as published by + * the Free Software Foundation; either version 2 of the License, or + * (at your option) any later version. + * + * This program is distributed in the hope that it will be useful, + * but WITHOUT ANY WARRANTY; without even the implied warranty of + * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the + * GNU General Public License for more details. + * + * You should have received a copy of the GNU General Public License + * along with this program; if not, write to the Free Software + * Foundation, Inc., 59 Temple Place, Suite 330, Boston, MA 02111-1307 USA + * + * gcc -o cldemo -std=gnu99 -Wall -I/usr/include/nvidia-current cldemo.c -lOpenCL + * + */ + +#include +#include +#include +#include +#include +#include +#include +#include +#include + +#define NUM_DATA 64 + +#define CL_CHECK(_expr) \ + do { \ + cl_int _err = _expr; \ + if (_err == CL_SUCCESS) \ + break; \ + fprintf(stderr, "OpenCL Error: '%s' returned %d!\n", #_expr, (int)_err); \ + abort(); \ + } while (0) + +#define CL_CHECK_ERR(_expr) \ + ({ \ + cl_int _err = CL_INVALID_VALUE; \ + typeof(_expr) _ret = _expr; \ + if (_err != CL_SUCCESS) { \ + fprintf(stderr, "OpenCL Error: '%s' returned %d!\n", #_expr, (int)_err); \ + abort(); \ + } \ + _ret; \ + }) + +void pfn_notify(const char *errinfo, const void *private_info, size_t cb, void *user_data) +{ + fprintf(stderr, "OpenCL Error (via pfn_notify): %s\n", errinfo); +} + +/// +// Create an OpenCL program from the kernel source file +// +cl_program CreateProgram(cl_context context, cl_device_id device, const char* fileName) +{ + cl_int errNum; + cl_program program; + + std::ifstream kernelFile(fileName, std::ios::in); + if (!kernelFile.is_open()) + { + std::cerr << "Failed to open file for reading: " << fileName << std::endl; + return NULL; + } + + std::ostringstream oss; + oss << kernelFile.rdbuf(); + + std::string srcStdStr = oss.str(); + const char *srcStr = srcStdStr.c_str(); + program = clCreateProgramWithSource(context, 1, + (const char**)&srcStr, + NULL, NULL); + if (program == NULL) + { + std::cerr << "Failed to create CL program from source." << std::endl; + return NULL; + } + + errNum = clBuildProgram(program, 0, NULL, NULL, NULL, NULL); + if (errNum != CL_SUCCESS) + { + // Determine the reason for the error + char buildLog[16384]; + clGetProgramBuildInfo(program, device, CL_PROGRAM_BUILD_LOG, + sizeof(buildLog), buildLog, NULL); + + std::cerr << "Error in kernel: " << std::endl; + std::cerr << buildLog; + clReleaseProgram(program); + return NULL; + } + + return program; +} + +// +/// +// Retreive program binary for all of the devices attached to the +// program an and store the one for the device passed in +// +bool SaveProgramBinary(cl_program program, cl_device_id device, const char* fileName) +{ + //cl_uint numDevices = malloc(sizeof(cl_uint)); + //cl_uint* numDevices = malloc(sizeof(cl_uint)); + cl_int errNum; + + printf("try getting program info\n"); + // 1 - Query for number of devices attached to program + /*errNum = clGetProgramInfo(program, CL_PROGRAM_NUM_DEVICES, sizeof(cl_uint), + &numDevices, NULL); + printf("Got program_num_devices\n"); + if (errNum != CL_SUCCESS) + { + std::cerr << "Error querying for number of devices." << std::endl; + return false; + }*/ + + // 2 - Get all of the Device IDs + cl_device_id *devices = new cl_device_id[1]; + errNum = clGetProgramInfo(program, CL_PROGRAM_DEVICES, + sizeof(cl_device_id) * 1, + devices, NULL); + printf("Got program_devices\n"); + if (errNum != CL_SUCCESS) + { + std::cerr << "Error querying for devices." << std::endl; + delete [] devices; + return false; + } + + // 3 - Determine the size of each program binary + size_t *programBinarySizes = new size_t [1]; + errNum = clGetProgramInfo(program, CL_PROGRAM_BINARY_SIZES, + sizeof(size_t) * 1, + programBinarySizes, NULL); + printf("Got program_binary_sizes\n"); + if (errNum != CL_SUCCESS) + { + std::cerr << "Error querying for program binary sizes." << std::endl; + delete [] devices; + delete [] programBinarySizes; + return false; + } + + unsigned char **programBinaries = new unsigned char*[1]; + for (cl_uint i = 0; i < 1; i++) + { + programBinaries[i] = new unsigned char[programBinarySizes[i]]; + } + + // 4 - Get all of the program binaries + errNum = clGetProgramInfo(program, CL_PROGRAM_BINARIES, sizeof(unsigned char*) * 1, + programBinaries, NULL); + printf("Got program_binarys\n"); + if (errNum != CL_SUCCESS) + { + std::cerr << "Error querying for program binaries" << std::endl; + + delete [] devices; + delete [] programBinarySizes; + for (cl_uint i = 0; i < 1; i++) + { + delete [] programBinaries[i]; + } + delete [] programBinaries; + return false; + } + + // 5 - Finally store the binaries for the device requested out to disk for future reading. + for (cl_uint i = 0; i < 1; i++) + { + // Store the binary just for the device requested. In a scenario where + // multiple devices were being used you would save all of the binaries out here. + if (devices[i] == device) + { + FILE *fp = fopen(fileName, "wb"); + if(fp ==NULL){ + delete [] devices; + delete [] programBinarySizes; + for (cl_uint i = 0; i < 1; i++) + { + delete [] programBinaries[i]; + } + delete [] programBinaries; + return false; + } + printf("Opened file\n"); + fwrite(programBinaries[i], 1, programBinarySizes[i], fp); + printf("wrote file\n"); + fclose(fp); + printf("close file\n"); + break; + } + } + + // Cleanup + delete [] devices; + delete [] programBinarySizes; + for (cl_uint i = 0; i < 1; i++) + { + delete [] programBinaries[i]; + } + delete [] programBinaries; + return true; +} + +/// +// Attempt to create the program object from a cached binary. Note that +// on first run this will fail because the binary has not yet been created. +// +cl_program CreateProgramFromBinary(cl_context context, cl_device_id device, const char* fileName) +{ + FILE *fp = fopen(fileName, "rb"); + if (fp == NULL) + { + return NULL; + } + + // Determine the size of the binary + size_t binarySize; + fseek(fp, 0, SEEK_END); + binarySize = ftell(fp); + rewind(fp); + + unsigned char *programBinary = new unsigned char[binarySize]; + fread(programBinary, 1, binarySize, fp); + fclose(fp); + + cl_int errNum = 0; + cl_program program; + cl_int binaryStatus; + + program = clCreateProgramWithBinary(context, + 1, + &device, + &binarySize, + (const unsigned char**)&programBinary, + &binaryStatus, + &errNum); + delete [] programBinary; + if (errNum != CL_SUCCESS) + { + std::cerr << "Error loading program binary." << std::endl; + return NULL; + } + + if (binaryStatus != CL_SUCCESS) + { + std::cerr << "Invalid binary for device" << std::endl; + return NULL; + } + + errNum = clBuildProgram(program, 1, &device, NULL, NULL, NULL); + if (errNum != CL_SUCCESS) + { + printf("build errNum:%d\n", errNum); + // Determine the reason for the error + char buildLog[16384]; + clGetProgramBuildInfo(program, device, CL_PROGRAM_BUILD_LOG, + sizeof(buildLog), buildLog, NULL); + + std::cerr << "Error in program: " << std::endl; + std::cerr << buildLog << std::endl; + clReleaseProgram(program); + return NULL; + } + + return program; +} + +/// +// Cleanup any created OpenCL resources +// +void Cleanup(cl_context context, cl_command_queue commandQueue, + cl_program program, cl_kernel kernel, cl_mem memObjects[3]) +{ + for (int i = 0; i < 3; i++) + { + if (memObjects[i] != 0) + clReleaseMemObject(memObjects[i]); + } + if (commandQueue != 0) + clReleaseCommandQueue(commandQueue); + + if (kernel != 0) + clReleaseKernel(kernel); + + if (program != 0) + clReleaseProgram(program); + + if (context != 0) + clReleaseContext(context); + +} + +int main(int argc, char **argv) +{ + printf("enter demo main\n"); + fflush(stdout); + putenv("POCL_VERBOSE=1"); + putenv("POCL_DEVICES=basic"); + putenv("POCL_LEAVE_TEMP_DIRS=1"); + putenv("POCL_LEAVE_KERNEL_COMPILER_TEMP_FILES=1"); + putenv("POCL_TEMP_DIR=pocl"); + putenv("POCL_CACHE_DIR=pocl"); + putenv("POCL_WORK_GROUP_METHOD=spmd"); + if(argc >= 2){ + printf("argv[1]:%s:\n",argv[1]); + if(!strcmp(argv[1], "h")) + putenv("POCL_WORK_GROUP_METHOD=spmd"); + if(!strcmp(argv[1], "c")) + putenv("POCL_CROSS_COMPILE=1"); + } + if(argc >= 3){ + printf("argv[2]:%s:\n",argv[2]); + if(!strcmp(argv[2], "h")) + putenv("POCL_WORK_GROUP_METHOD=spmd"); + if(!strcmp(argv[2], "c")) + putenv("POCL_CROSS_COMPILE=1"); + } + + //putenv("LD_LIBRARY_PATH=/scratch/colins/build/linux/fs/lib"); + //putenv("LTDL_LIBRARY_PATH=/scratch/colins/build/linux/fs/lib"); + //lt_dlsetsearchpath("/scratch/colins/build/linux/fs/lib"); + //printf("SEARCH_PATH:%s\n",lt_dlgetsearchpath()); + cl_platform_id platforms[100]; + cl_uint platforms_n = 0; + CL_CHECK(clGetPlatformIDs(100, platforms, &platforms_n)); + + printf("=== %d OpenCL platform(s) found: ===\n", platforms_n); + for (int i=0; i