github-mirrors/vortex
1
0
Fork 0
mirror of https://github.com/vortexgpgpu/vortex.git synced 2025-04-24 13:57:17 -04:00
Code Issues Projects Releases Packages Wiki Activity Actions

constant integration updates

Browse source
This commit is contained in:
Blaise Tine 2020-11-15 21:30:31 -08:00
parent 60f5c15bc8
commit 2904f6441d
19 changed files with 58 additions and 61 deletions
Download patch file Download diff file Expand all files Collapse all files

2
benchmarks/opencl/bfs/CLHelper.h
View file

@ -259,7 +259,7 @@ free(allPlatforms);*/
std::abort();
oclHandles.program = clCreateProgramWithBinary(
oclHandles.context, 1, &oclHandles.devices[DEVICE_ID_INUSED], &kernel_size, &kernel_bin, &binary_status, &resultCL);
oclHandles.context, 1, &oclHandles.devices[DEVICE_ID_INUSED], &kernel_size, (const uint8_t**)&kernel_bin, &binary_status, &resultCL);
free(kernel_bin);
if ((resultCL != CL_SUCCESS) || (oclHandles.program == NULL))

3
benchmarks/opencl/bfs/Makefile
View file

@ -11,7 +11,8 @@ K_LLCFLAGS += "-O3 -march=riscv32 -target-abi=ilp32f -mcpu=generic-rv32 -mattr=+
K_CFLAGS += "-v -O3 --sysroot=$(SYSROOT) --gcc-toolchain=$(RISCV_TOOLCHAIN_PATH) -march=rv32imf -mabi=ilp32f -I$(VORTEX_RT_PATH)/include -fno-rtti -fno-exceptions -ffreestanding -nostartfiles -Wl,--gc-sections -Wl,-Bstatic,-T$(VORTEX_RT_PATH)/linker/vx_link.ld"
K_LDFLAGS += "$(VORTEX_RT_PATH)/libvortexrt.a -lm"
CXXFLAGS += -std=c++11 -O0 -g -fpermissive -Wall -Wextra -Wfatal-errors
CXXFLAGS += -std=c++11 -O2 -Wall -Wextra -pedantic -Wfatal-errors
#CXXFLAGS += -std=c++11 -O0 -g -Wall -Wextra -pedantic -Wfatal-errors
CXXFLAGS += -I$(POCL_RT_PATH)/include

3
benchmarks/opencl/convolution/Makefile
View file

@ -11,7 +11,8 @@ K_LLCFLAGS += "-O3 -march=riscv32 -target-abi=ilp32f -mcpu=generic-rv32 -mattr=+
K_CFLAGS += "-v -O3 --sysroot=$(SYSROOT) --gcc-toolchain=$(RISCV_TOOLCHAIN_PATH) -march=rv32imf -mabi=ilp32f -I$(VORTEX_RT_PATH)/include -fno-rtti -fno-exceptions -ffreestanding -nostartfiles -Wl,--gc-sections -Wl,-Bstatic,-T$(VORTEX_RT_PATH)/linker/vx_link.ld"
K_LDFLAGS += "$(VORTEX_RT_PATH)/libvortexrt.a -lm"
CXXFLAGS += -std=c++11 -O0 -g -fpermissive -Wall -Wextra -pedantic -Wfatal-errors
CXXFLAGS += -std=c++11 -O2 -Wall -Wextra -pedantic -Wfatal-errors
#CXXFLAGS += -std=c++11 -O0 -g -Wall -Wextra -pedantic -Wfatal-errors
CXXFLAGS += -I$(POCL_RT_PATH)/include

4
benchmarks/opencl/guassian/Makefile
View file

@ -11,8 +11,8 @@ K_LLCFLAGS += "-O3 -march=riscv32 -target-abi=ilp32f -mcpu=generic-rv32 -mattr=+
K_CFLAGS += "-v -O3 --sysroot=$(SYSROOT) --gcc-toolchain=$(RISCV_TOOLCHAIN_PATH) -march=rv32imf -mabi=ilp32f -I$(VORTEX_RT_PATH)/include -fno-rtti -fno-exceptions -ffreestanding -nostartfiles -Wl,--gc-sections -Wl,-Bstatic,-T$(VORTEX_RT_PATH)/linker/vx_link.ld"
K_LDFLAGS += "$(VORTEX_RT_PATH)/libvortexrt.a -lm"
CXXFLAGS += -std=c++11 -O2 -fpermissive -Wall -Wextra -pedantic -Wfatal-errors
#CXXFLAGS += -std=c++11 -O0 -g -fpermissive -Wall -Wextra -pedantic -Wfatal-errors
CXXFLAGS += -std=c++11 -O2 -Wall -Wextra -pedantic -Wfatal-errors
#CXXFLAGS += -std=c++11 -O0 -g -Wall -Wextra -pedantic -Wfatal-errors
CXXFLAGS += -I$(POCL_RT_PATH)/include

2
benchmarks/opencl/guassian/clutils.cpp
View file

@ -868,7 +868,7 @@ cl_program cl_compileProgram(char* kernelPath, char* compileoptions, bool verbos
status = read_kernel_file("kernel.pocl", &kernel_bin, &kernel_size);
cl_errChk(status, "read_kernel_file", true);
cl_program clProgramReturn = clCreateProgramWithBinary(
context, 1, &device, &kernel_size, &kernel_bin, &binary_status, &status);
context, 1, &device, &kernel_size, (const uint8_t**)&kernel_bin, &binary_status, &status);
free(kernel_bin);
cl_errChk(status, "Creating program", true);

3
benchmarks/opencl/kmeans/Makefile
View file

@ -11,7 +11,8 @@ K_LLCFLAGS += "-O3 -march=riscv32 -target-abi=ilp32f -mcpu=generic-rv32 -mattr=+
K_CFLAGS += "-v -O3 --sysroot=$(SYSROOT) --gcc-toolchain=$(RISCV_TOOLCHAIN_PATH) -march=rv32imf -mabi=ilp32f -I$(VORTEX_RT_PATH)/include -fno-rtti -fno-exceptions -ffreestanding -nostartfiles -Wl,--gc-sections -Wl,-Bstatic,-T$(VORTEX_RT_PATH)/linker/vx_link.ld"
K_LDFLAGS += "$(VORTEX_RT_PATH)/libvortexrt.a -lm"
CXXFLAGS += -std=c++11 -O0 -g -fpermissive -Wall -Wextra -pedantic -Wfatal-errors
CXXFLAGS += -std=c++11 -O2 -Wall -Wextra -pedantic -Wfatal-errors
#CXXFLAGS += -std=c++11 -O0 -g -Wall -Wextra -pedantic -Wfatal-errors
CXXFLAGS += -I$(POCL_RT_PATH)/include

4
benchmarks/opencl/nearn/Makefile
View file

@ -11,8 +11,8 @@ K_LLCFLAGS += "-O3 -march=riscv32 -target-abi=ilp32f -mcpu=generic-rv32 -mattr=+
K_CFLAGS += "-v -O3 --sysroot=$(SYSROOT) --gcc-toolchain=$(RISCV_TOOLCHAIN_PATH) -march=rv32imf -mabi=ilp32f -I$(VORTEX_RT_PATH)/include -fno-rtti -fno-exceptions -ffreestanding -nostartfiles -Wl,--gc-sections -Wl,-Bstatic,-T$(VORTEX_RT_PATH)/linker/vx_link.ld"
K_LDFLAGS += "$(VORTEX_RT_PATH)/libvortexrt.a -lm"
CXXFLAGS += -std=c++11 -O2 -fpermissive -Wall -Wextra -pedantic -Wfatal-errors
#CXXFLAGS += -std=c++11 -O0 -g -Wall -fpermissive -Wextra -pedantic -Wfatal-errors
CXXFLAGS += -std=c++11 -O2 -Wall -Wextra -pedantic -Wfatal-errors
#CXXFLAGS += -std=c++11 -O0 -g -Wall -Wextra -pedantic -Wfatal-errors
CXXFLAGS += -I$(POCL_RT_PATH)/include

2
benchmarks/opencl/nearn/clutils.cpp
View file

@ -868,7 +868,7 @@ cl_program cl_compileProgram(char* kernelPath, char* compileoptions, bool verbos
// Create the program object
//cl_program clProgramReturn = clCreateProgramWithSource(context, 1, (const char **)&source, NULL, &status);
cl_program clProgramReturn = clCreateProgramWithBinary(
context, 1, devices, &kernel_size, &kernel_bin, &binary_status, &status);
context, 1, devices, &kernel_size, (const uint8_t**)&kernel_bin, &binary_status, &status);
free(kernel_bin);
cl_errChk(status, "Creating program", true);

4
benchmarks/opencl/saxpy/Makefile
View file

@ -11,8 +11,8 @@ K_LLCFLAGS += "-O3 -march=riscv32 -target-abi=ilp32f -mcpu=generic-rv32 -mattr=+
K_CFLAGS += "-v -O3 --sysroot=$(SYSROOT) --gcc-toolchain=$(RISCV_TOOLCHAIN_PATH) -march=rv32imf -mabi=ilp32f -I$(VORTEX_RT_PATH)/include -fno-rtti -fno-exceptions -ffreestanding -nostartfiles -Wl,--gc-sections -Wl,-Bstatic,-T$(VORTEX_RT_PATH)/linker/vx_link.ld"
K_LDFLAGS += "$(VORTEX_RT_PATH)/libvortexrt.a -lm"
CXXFLAGS += -std=c++11 -O2 -fpermissive -Wall -Wextra -pedantic -Wfatal-errors
#CXXFLAGS += -std=c++11 -O0 -g -fpermissive -Wall -Wextra -pedantic -Wfatal-errors
CXXFLAGS += -std=c++11 -O2 -Wall -Wextra -pedantic -Wfatal-errors
#CXXFLAGS += -std=c++11 -O0 -g -Wall -Wextra -pedantic -Wfatal-errors
CXXFLAGS += -I$(POCL_RT_PATH)/include

57
benchmarks/opencl/saxpy/main.cc
View file

@ -29,6 +29,7 @@
#include <stdio.h>
#include <stdlib.h>
#include <string.h>
#include <chrono>
//#define NUM_DATA 65536
#define NUM_DATA 1024
@ -136,7 +137,7 @@ int main(int argc, char **argv) {
// and store the binary for future use.
std::cout << "Attempting to create program from binary..." << std::endl;
cl_program program = CL_CHECK_ERR(clCreateProgramWithBinary(
context, 1, &device_id, &kernel_size, &kernel_bin, &binary_status, &_err));
context, 1, &device_id, &kernel_size, (const uint8_t**)&kernel_bin, &binary_status, &_err));
if (program == NULL) {
std::cerr << "Failed to write program binary" << std::endl;
Cleanup(context, queue, program, kernel, memObjects);
@ -148,17 +149,17 @@ int main(int argc, char **argv) {
// Build program
CL_CHECK(clBuildProgram(program, 1, &device_id, NULL, NULL, NULL));
size_t nbytes = sizeof(float) * NUM_DATA;
printf("attempting to create input buffer\n");
fflush(stdout);
cl_mem input_buffer;
input_buffer = CL_CHECK_ERR(clCreateBuffer(
context, CL_MEM_READ_ONLY, sizeof(float) * NUM_DATA, NULL, &_err));
context, CL_MEM_READ_ONLY, nbytes, NULL, &_err));
printf("attempting to create output buffer\n");
fflush(stdout);
cl_mem output_buffer;
output_buffer = CL_CHECK_ERR(clCreateBuffer(
context, CL_MEM_WRITE_ONLY, sizeof(float) * NUM_DATA, NULL, &_err));
context, CL_MEM_WRITE_ONLY, nbytes, NULL, &_err));
memObjects[0] = input_buffer;
memObjects[1] = output_buffer;
@ -166,50 +167,42 @@ int main(int argc, char **argv) {
float factor = ((float)rand() / (float)(RAND_MAX)) * 100.0;
printf("attempting to create kernel\n");
fflush(stdout);
kernel = CL_CHECK_ERR(clCreateKernel(program, "saxpy", &_err));
printf("setting up kernel args cl_mem:%lx \n", input_buffer);
fflush(stdout);
CL_CHECK(clSetKernelArg(kernel, 0, sizeof(input_buffer), &input_buffer));
CL_CHECK(clSetKernelArg(kernel, 1, sizeof(output_buffer), &output_buffer));
CL_CHECK(clSetKernelArg(kernel, 2, sizeof(factor), &factor));
printf("attempting to enqueue write buffer\n");
fflush(stdout);
float* h_src = (float*)malloc(nbytes);
for (int i = 0; i < NUM_DATA; i++) {
float in = ((float)rand() / (float)(RAND_MAX)) * 100.0;
CL_CHECK(clEnqueueWriteBuffer(queue, input_buffer, CL_TRUE,
i * sizeof(float), 4, &in, 0, NULL, NULL));
h_src[i] = ((float)rand() / (float)(RAND_MAX)) * 100.0;
}
CL_CHECK(clEnqueueWriteBuffer(queue, input_buffer, CL_TRUE, 0, nbytes, h_src, 0, NULL, NULL));
free(h_src);
cl_event kernel_completion;
size_t global_work_size[] = {NUM_DATA/2,NUM_DATA/2};
size_t global_work_size[] = {NUM_DATA/2, NUM_DATA/2};
printf("attempting to enqueue kernel\n");
fflush(stdout);
auto time_start = std::chrono::high_resolution_clock::now();
CL_CHECK(clEnqueueNDRangeKernel(queue, kernel, 1, NULL, global_work_size,
NULL, 0, NULL, &kernel_completion));
printf("Enqueue'd kerenel\n");
fflush(stdout);
cl_ulong time_start, time_end;
CL_CHECK(clWaitForEvents(1, &kernel_completion));
CL_CHECK(clGetEventProfilingInfo(kernel_completion,
CL_PROFILING_COMMAND_START,
sizeof(time_start), &time_start, NULL));
CL_CHECK(clGetEventProfilingInfo(kernel_completion, CL_PROFILING_COMMAND_END,
sizeof(time_end), &time_end, NULL));
double elapsed = time_end - time_start;
printf("time(ns):%lg\n", elapsed);
CL_CHECK(clReleaseEvent(kernel_completion));
NULL, 0, NULL, NULL));
CL_CHECK(clFinish(queue));
auto time_end = std::chrono::high_resolution_clock::now();
double elapsed = std::chrono::duration_cast<std::chrono::milliseconds>(time_end - time_start).count();
printf("Elapsed time: %lg ms\n", elapsed);
printf("Download destination buffer\n");
float* h_dst = (float*)malloc(nbytes);
CL_CHECK(clEnqueueReadBuffer(queue, output_buffer, CL_TRUE, 0, nbytes, h_dst, 0, NULL, NULL));
printf("Result:");
for (int i = 0; i < NUM_DATA; i++) {
float data;
CL_CHECK(clEnqueueReadBuffer(queue, output_buffer, CL_TRUE,
i * sizeof(float), 4, &data, 0, NULL, NULL));
float data = h_dst[i];
// printf(" %f", data);
}
printf("\n");
printf("Passed!\n");
free(h_dst);
CL_CHECK(clReleaseMemObject(memObjects[0]));
CL_CHECK(clReleaseMemObject(memObjects[1]));

4
benchmarks/opencl/sfilter/Makefile
View file

@ -11,8 +11,8 @@ K_LLCFLAGS += "-O3 -march=riscv32 -target-abi=ilp32f -mcpu=generic-rv32 -mattr=+
K_CFLAGS += "-v -O3 --sysroot=$(SYSROOT) --gcc-toolchain=$(RISCV_TOOLCHAIN_PATH) -march=rv32imf -mabi=ilp32f -I$(VORTEX_RT_PATH)/include -fno-rtti -fno-exceptions -ffreestanding -nostartfiles -Wl,--gc-sections -Wl,-Bstatic,-T$(VORTEX_RT_PATH)/linker/vx_link.ld"
K_LDFLAGS += "$(VORTEX_RT_PATH)/libvortexrt.a -lm"
CXXFLAGS += -std=c++11 -O2 -fpermissive -Wall -Wextra -pedantic -Wfatal-errors
#CXXFLAGS += -std=c++11 -O0 -g -fpermissive -Wall -Wextra -pedantic -Wfatal-errors
CXXFLAGS += -std=c++11 -O2 -Wall -Wextra -pedantic -Wfatal-errors
#CXXFLAGS += -std=c++11 -O0 -g -Wall -Wextra -pedantic -Wfatal-errors
CXXFLAGS += -I$(POCL_RT_PATH)/include

9
benchmarks/opencl/sfilter/main.cc
View file

@ -211,7 +211,7 @@ int main(int argc, char **argv) {
// and store the binary for future use.
std::cout << "Attempting to create program from binary..." << std::endl;
cl_program program = CL_CHECK_ERR(clCreateProgramWithBinary(
context, 1, &device_id, &kernel_size, &kernel_bin, &binary_status, &_err));
context, 1, &device_id, &kernel_size, (const uint8_t**)&kernel_bin, &binary_status, &_err));
if (program == NULL) {
std::cerr << "Failed to write program binary" << std::endl;
Cleanup(context, queue, program, kernel, memObjects);
@ -264,8 +264,7 @@ int main(int argc, char **argv) {
CL_CHECK(clSetKernelArg(kernel, 10, sizeof(m7), (&m7)));
CL_CHECK(clSetKernelArg(kernel, 11, sizeof(m8), (&m8)));
printf("attempting to enqueue write buffer\n");
printf("attempting to enqueue write buffer\n");
float* h_src = (float*)malloc(nbytes);
for (int i = 0; i < NUM_DATA * NUM_DATA; i++) {
h_src[i] = ((float)rand() / (float)(RAND_MAX)) * 100.0;
@ -294,8 +293,8 @@ int main(int argc, char **argv) {
float data = h_dst[i];
// printf(" %f", data);
}
printf("\n");
printf("Passed!\n");
free(h_dst);
CL_CHECK(clReleaseMemObject(memObjects[0]));
CL_CHECK(clReleaseMemObject(memObjects[1]));

4
benchmarks/opencl/sgemm/Makefile
View file

@ -11,8 +11,8 @@ K_LLCFLAGS += "-O3 -march=riscv32 -target-abi=ilp32f -mcpu=generic-rv32 -mattr=+
K_CFLAGS += "-v -O3 --sysroot=$(SYSROOT) --gcc-toolchain=$(RISCV_TOOLCHAIN_PATH) -march=rv32imf -mabi=ilp32f -I$(VORTEX_RT_PATH)/include -fno-rtti -fno-exceptions -ffreestanding -nostartfiles -Wl,--gc-sections -Wl,-Bstatic,-T$(VORTEX_RT_PATH)/linker/vx_link.ld"
K_LDFLAGS += "$(VORTEX_RT_PATH)/libvortexrt.a -lm"
CXXFLAGS += -std=c++11 -O2 -fpermissive -Wall -Wextra -pedantic -Wfatal-errors
#CXXFLAGS += -std=c++11 -O0 -g -fpermissive -Wall -Wextra -pedantic -Wfatal-errors
CXXFLAGS += -std=c++11 -O2 -Wall -Wextra -pedantic -Wfatal-errors
#CXXFLAGS += -std=c++11 -O0 -g -Wall -Wextra -pedantic -Wfatal-errors
CXXFLAGS += -I$(POCL_RT_PATH)/include

2
benchmarks/opencl/sgemm/main.cc
View file

@ -127,7 +127,7 @@ int main (int argc, char **argv) {
printf("Create program from kernel source\n");
program = CL_CHECK2(clCreateProgramWithBinary(
context, 1, &device_id, &kernel_size, &kernel_bin, &binary_status, &_err));
context, 1, &device_id, &kernel_size, (const uint8_t**)&kernel_bin, &binary_status, &_err));
if (program == NULL) {
cleanup();
return -1;

3
benchmarks/opencl/transpose/Makefile
View file

@ -11,7 +11,8 @@ K_LLCFLAGS += "-O3 -march=riscv32 -target-abi=ilp32f -mcpu=generic-rv32 -mattr=+
K_CFLAGS += "-v -O3 --sysroot=$(SYSROOT) --gcc-toolchain=$(RISCV_TOOLCHAIN_PATH) -march=rv32imf -mabi=ilp32f -I$(VORTEX_RT_PATH)/include -fno-rtti -fno-exceptions -ffreestanding -nostartfiles -Wl,--gc-sections -Wl,-Bstatic,-T$(VORTEX_RT_PATH)/linker/vx_link.ld"
K_LDFLAGS += "$(VORTEX_RT_PATH)/libvortexrt.a -lm"
CXXFLAGS += -std=c++11 -O0 -g -fpermissive -Wall -Wextra -pedantic -Wfatal-errors
CXXFLAGS += -std=c++11 -O2 -Wall -Wextra -pedantic -Wfatal-errors
#CXXFLAGS += -std=c++11 -O0 -g -Wall -Wextra -pedantic -Wfatal-errors
CXXFLAGS += -I$(POCL_RT_PATH)/include

3
benchmarks/opencl/transpose/main.cc
View file

@ -311,7 +311,8 @@ int runTest( const int argc, const char** argv)
cl_int binary_status = 0;
cl_device_id device_id;
// create the program
rv_program = clCreateProgramWithBinary(cxGPUContext, 1, &device_id, &kernel_size, &kernel_bin, &binary_status, NULL);
rv_program = clCreateProgramWithBinary(
cxGPUContext, 1, &device_id, &kernel_size, (const uint8_t**)&kernel_bin, &binary_status, NULL);
//rv_program = clCreateProgramWithSource(cxGPUContext, 1,
// (const char **)&source, &program_length, &ciErrNum);
//oclCheckError(ciErrNum, CL_SUCCESS);

4
benchmarks/opencl/vecadd/Makefile
View file

@ -11,8 +11,8 @@ K_LLCFLAGS += "-O3 -march=riscv32 -target-abi=ilp32f -mcpu=generic-rv32 -mattr=+
K_CFLAGS += "-v -O3 --sysroot=$(SYSROOT) --gcc-toolchain=$(RISCV_TOOLCHAIN_PATH) -march=rv32imf -mabi=ilp32f -I$(VORTEX_RT_PATH)/include -fno-rtti -fno-exceptions -ffreestanding -nostartfiles -Wl,--gc-sections -Wl,-Bstatic,-T$(VORTEX_RT_PATH)/linker/vx_link.ld"
K_LDFLAGS += "$(VORTEX_RT_PATH)/libvortexrt.a -lm"
CXXFLAGS += -std=c++11 -O2 -fpermissive -Wall -Wextra -pedantic -Wfatal-errors
#CXXFLAGS += -std=c++11 -O0 -g -fpermissive -Wall -Wextra -pedantic -Wfatal-errors
CXXFLAGS += -std=c++11 -O2 -Wall -Wextra -pedantic -Wfatal-errors
#CXXFLAGS += -std=c++11 -O0 -g -Wall -Wextra -pedantic -Wfatal-errors
CXXFLAGS += -I$(POCL_RT_PATH)/include

2
benchmarks/opencl/vecadd/main.cc
View file

@ -113,7 +113,7 @@ int main (int argc, char **argv) {
printf("Create program from kernel source\n");
program = CL_CHECK2(clCreateProgramWithBinary(
context, 1, &device_id, &kernel_size, &kernel_bin, &binary_status, &_err));
context, 1, &device_id, &kernel_size, (const uint8_t**)&kernel_bin, &binary_status, &_err));
if (program == NULL) {
cleanup();
return -1;

4
hw/rtl/cache/VX_bank.v vendored
View file

@ -604,8 +604,8 @@ module VX_bank #(
// mark msrq entry that match DRAM fill as 'ready'
wire update_ready_st0 = dfpq_pop;
// push missed requests as 'ready' is this was a forced missed
// this request will be queued behind prior requests so will only pop when the fill arrives.
// push missed requests as 'ready' if it was a forced miss but actually had a hit
// or the fill request is comming for the missed block
wire msrq_init_ready_state_st3 = !miss_st3 || incoming_fill;
VX_cache_miss_resrv #(