diff --git a/benchmarks/opencl/bfs/CLHelper.h b/benchmarks/opencl/bfs/CLHelper.h index 4ea9b747..94536d7c 100755 --- a/benchmarks/opencl/bfs/CLHelper.h +++ b/benchmarks/opencl/bfs/CLHelper.h @@ -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)) diff --git a/benchmarks/opencl/bfs/Makefile b/benchmarks/opencl/bfs/Makefile index 7719f7bf..10f2e231 100644 --- a/benchmarks/opencl/bfs/Makefile +++ b/benchmarks/opencl/bfs/Makefile @@ -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 diff --git a/benchmarks/opencl/convolution/Makefile b/benchmarks/opencl/convolution/Makefile index a35e63d6..931bdefe 100644 --- a/benchmarks/opencl/convolution/Makefile +++ b/benchmarks/opencl/convolution/Makefile @@ -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 diff --git a/benchmarks/opencl/guassian/Makefile b/benchmarks/opencl/guassian/Makefile index 97db2f7a..26b32dc2 100644 --- a/benchmarks/opencl/guassian/Makefile +++ b/benchmarks/opencl/guassian/Makefile @@ -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 diff --git a/benchmarks/opencl/guassian/clutils.cpp b/benchmarks/opencl/guassian/clutils.cpp index 32feef52..37cabb7d 100755 --- a/benchmarks/opencl/guassian/clutils.cpp +++ b/benchmarks/opencl/guassian/clutils.cpp @@ -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); diff --git a/benchmarks/opencl/kmeans/Makefile b/benchmarks/opencl/kmeans/Makefile index 1c2eb6fc..d5af8a02 100644 --- a/benchmarks/opencl/kmeans/Makefile +++ b/benchmarks/opencl/kmeans/Makefile @@ -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 diff --git a/benchmarks/opencl/nearn/Makefile b/benchmarks/opencl/nearn/Makefile index 44131346..80f1bc90 100644 --- a/benchmarks/opencl/nearn/Makefile +++ b/benchmarks/opencl/nearn/Makefile @@ -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 diff --git a/benchmarks/opencl/nearn/clutils.cpp b/benchmarks/opencl/nearn/clutils.cpp index 6bc42304..3c433c0c 100755 --- a/benchmarks/opencl/nearn/clutils.cpp +++ b/benchmarks/opencl/nearn/clutils.cpp @@ -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); diff --git a/benchmarks/opencl/saxpy/Makefile b/benchmarks/opencl/saxpy/Makefile index bf244c96..18be0652 100644 --- a/benchmarks/opencl/saxpy/Makefile +++ b/benchmarks/opencl/saxpy/Makefile @@ -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 diff --git a/benchmarks/opencl/saxpy/main.cc b/benchmarks/opencl/saxpy/main.cc index d1d95bda..e5f66415 100644 --- a/benchmarks/opencl/saxpy/main.cc +++ b/benchmarks/opencl/saxpy/main.cc @@ -29,6 +29,7 @@ #include #include #include +#include //#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(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])); diff --git a/benchmarks/opencl/sfilter/Makefile b/benchmarks/opencl/sfilter/Makefile index b7cf8222..6b497101 100644 --- a/benchmarks/opencl/sfilter/Makefile +++ b/benchmarks/opencl/sfilter/Makefile @@ -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 diff --git a/benchmarks/opencl/sfilter/main.cc b/benchmarks/opencl/sfilter/main.cc index 98c644ca..894c32f3 100644 --- a/benchmarks/opencl/sfilter/main.cc +++ b/benchmarks/opencl/sfilter/main.cc @@ -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])); diff --git a/benchmarks/opencl/sgemm/Makefile b/benchmarks/opencl/sgemm/Makefile index b98fb1a0..692bb8aa 100644 --- a/benchmarks/opencl/sgemm/Makefile +++ b/benchmarks/opencl/sgemm/Makefile @@ -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 diff --git a/benchmarks/opencl/sgemm/main.cc b/benchmarks/opencl/sgemm/main.cc index 81775741..11251f6a 100644 --- a/benchmarks/opencl/sgemm/main.cc +++ b/benchmarks/opencl/sgemm/main.cc @@ -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; diff --git a/benchmarks/opencl/transpose/Makefile b/benchmarks/opencl/transpose/Makefile index 14bad041..e4837cdc 100644 --- a/benchmarks/opencl/transpose/Makefile +++ b/benchmarks/opencl/transpose/Makefile @@ -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 diff --git a/benchmarks/opencl/transpose/main.cc b/benchmarks/opencl/transpose/main.cc index f72cb851..55ec3f6d 100644 --- a/benchmarks/opencl/transpose/main.cc +++ b/benchmarks/opencl/transpose/main.cc @@ -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); diff --git a/benchmarks/opencl/vecadd/Makefile b/benchmarks/opencl/vecadd/Makefile index 7bd9bdc0..a69f21ce 100644 --- a/benchmarks/opencl/vecadd/Makefile +++ b/benchmarks/opencl/vecadd/Makefile @@ -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 diff --git a/benchmarks/opencl/vecadd/main.cc b/benchmarks/opencl/vecadd/main.cc index 8a1081e7..81f7bff4 100644 --- a/benchmarks/opencl/vecadd/main.cc +++ b/benchmarks/opencl/vecadd/main.cc @@ -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; diff --git a/hw/rtl/cache/VX_bank.v b/hw/rtl/cache/VX_bank.v index 31fc23fd..cf21bce8 100644 --- a/hw/rtl/cache/VX_bank.v +++ b/hw/rtl/cache/VX_bank.v @@ -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 #(