From 529a2a0242059407f13c2b473e7d8ad776d7a924 Mon Sep 17 00:00:00 2001 From: talubik Date: Tue, 30 Sep 2025 18:04:53 +0300 Subject: [PATCH 1/4] Add : test for copyBuffer --- .gitignore | 2 ++ tests/opencl/Makefile | 7 ++++- tests/opencl/copybuf/Makefile | 11 ++++++++ tests/opencl/copybuf/main.cpp | 48 +++++++++++++++++++++++++++++++++++ 4 files changed, 67 insertions(+), 1 deletion(-) create mode 100644 tests/opencl/copybuf/Makefile create mode 100644 tests/opencl/copybuf/main.cpp diff --git a/.gitignore b/.gitignore index 43388e9cb5..427852b54b 100644 --- a/.gitignore +++ b/.gitignore @@ -2,3 +2,5 @@ /.vscode *.cache *.code-workspace +/tests/opencl/atomics/ +/tests/opencl/op_binary/ \ No newline at end of file diff --git a/tests/opencl/Makefile b/tests/opencl/Makefile index e60cd6ec74..993faec1a0 100644 --- a/tests/opencl/Makefile +++ b/tests/opencl/Makefile @@ -22,6 +22,8 @@ all: $(MAKE) -C kmeans $(MAKE) -C blackscholes $(MAKE) -C bfs + $(MAKE) -C copybuf + run-simx: $(MAKE) -C vecadd run-simx @@ -43,6 +45,7 @@ run-simx: $(MAKE) -C kmeans run-simx $(MAKE) -C blackscholes run-simx $(MAKE) -C bfs run-simx + $(MAKE) -C copybuf run-simx run-rtlsim: $(MAKE) -C vecadd run-rtlsim @@ -64,6 +67,7 @@ run-rtlsim: $(MAKE) -C kmeans run-rtlsim $(MAKE) -C blackscholes run-rtlsim $(MAKE) -C bfs run-rtlsim + $(MAKE) -C copybuf run-rtlsim clean: $(MAKE) -C vecadd clean @@ -85,4 +89,5 @@ clean: $(MAKE) -C guassian clean $(MAKE) -C kmeans clean $(MAKE) -C blackscholes clean - $(MAKE) -C bfs clean \ No newline at end of file + $(MAKE) -C bfs clean + $(MAKE) -C copybuf clean \ No newline at end of file diff --git a/tests/opencl/copybuf/Makefile b/tests/opencl/copybuf/Makefile new file mode 100644 index 0000000000..26e76dce93 --- /dev/null +++ b/tests/opencl/copybuf/Makefile @@ -0,0 +1,11 @@ +ROOT_DIR := $(realpath ../../..) +include $(ROOT_DIR)/config.mk + +PROJECT := copybuf + +SRC_DIR := $(VORTEX_HOME)/tests/opencl/$(PROJECT) + +SRCS := $(SRC_DIR)/main.cpp + + +include ../common.mk diff --git a/tests/opencl/copybuf/main.cpp b/tests/opencl/copybuf/main.cpp new file mode 100644 index 0000000000..e13c42200b --- /dev/null +++ b/tests/opencl/copybuf/main.cpp @@ -0,0 +1,48 @@ +#define CL_TARGET_OPENCL_VERSION 120 +#define CL_HPP_TARGET_OPENCL_VERSION 120 +#define CL_HPP_MINIMUM_OPENCL_VERSION 120 + +#include +#include +#include + +int main() { + std::vector platforms; + cl::Platform::get(&platforms); + if (platforms.empty()) { + std::cerr << "No OpenCL platform\n"; + return 1; + } + + std::vector devices; + platforms[0].getDevices(CL_DEVICE_TYPE_ALL, &devices); + if (devices.empty()) { + std::cerr << "No OpenCL device\n"; + return 2; + } + + cl::Context context(devices[0]); + cl::CommandQueue queue(context, devices[0]); + + std::vector src = {10, 20, 30, 40, 50}; + std::vector dst(src.size(), 0); + + cl::Buffer bufSrc(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, + sizeof(int) * src.size(), src.data()); + cl::Buffer bufDst(context, CL_MEM_WRITE_ONLY, + sizeof(int) * dst.size()); + + queue.enqueueCopyBuffer(bufSrc, bufDst, 0, 0, sizeof(int) * src.size()); + + queue.enqueueReadBuffer(bufDst, CL_TRUE, 0, sizeof(int) * dst.size(), dst.data()); + + std::cout << "Source Buffer: "; + for (auto v : src) + std::cout << v << " "; + std::cout << "\nDestination buffer: "; + for (auto v : dst) + std::cout << v << " "; + std::cout << std::endl; + + return 0; +} From 279b2dd20943cd561366994d2ae40b6d35c7cd8c Mon Sep 17 00:00:00 2001 From: talubik Date: Tue, 30 Sep 2025 20:38:38 +0300 Subject: [PATCH 2/4] Refactor: test copybuf --- tests/opencl/copybuf/Makefile | 2 +- tests/opencl/copybuf/main.cc | 69 +++++++++++++++++++++++++++++++++++ tests/opencl/copybuf/main.cpp | 48 ------------------------ 3 files changed, 70 insertions(+), 49 deletions(-) create mode 100644 tests/opencl/copybuf/main.cc delete mode 100644 tests/opencl/copybuf/main.cpp diff --git a/tests/opencl/copybuf/Makefile b/tests/opencl/copybuf/Makefile index 26e76dce93..f10a02695a 100644 --- a/tests/opencl/copybuf/Makefile +++ b/tests/opencl/copybuf/Makefile @@ -5,7 +5,7 @@ PROJECT := copybuf SRC_DIR := $(VORTEX_HOME)/tests/opencl/$(PROJECT) -SRCS := $(SRC_DIR)/main.cpp +SRCS := $(SRC_DIR)/main.cc include ../common.mk diff --git a/tests/opencl/copybuf/main.cc b/tests/opencl/copybuf/main.cc new file mode 100644 index 0000000000..8a97dc0aa5 --- /dev/null +++ b/tests/opencl/copybuf/main.cc @@ -0,0 +1,69 @@ +#define CL_TARGET_OPENCL_VERSION 120 +#include +#include +#include + +int main() { + cl_int err; + cl_uint num_platforms = 0; + err = clGetPlatformIDs(0, NULL, &num_platforms); + if (err != CL_SUCCESS || num_platforms == 0) { + fprintf(stderr, "No OpenCL platform\n"); + return 1; + } + cl_platform_id *platforms = (cl_platform_id *)malloc(sizeof(cl_platform_id) * num_platforms); + clGetPlatformIDs(num_platforms, platforms, NULL); + cl_uint num_devices = 0; + err = clGetDeviceIDs(platforms[0], CL_DEVICE_TYPE_ALL, 0, NULL, &num_devices); + if (err != CL_SUCCESS || num_devices == 0) { + fprintf(stderr, "No OpenCL device\n"); + free(platforms); + return 2; + } + cl_device_id *devices = (cl_device_id *)malloc(sizeof(cl_device_id) * num_devices); + clGetDeviceIDs(platforms[0], CL_DEVICE_TYPE_ALL, num_devices, devices, NULL); + cl_context context = clCreateContext(NULL, 1, &devices[0], NULL, NULL, &err); + if (err != CL_SUCCESS) { + fprintf(stderr, "Failed to create context\n"); + return 3; + } + cl_command_queue queue = clCreateCommandQueue(context, devices[0], 0, &err); + if (err != CL_SUCCESS) { + fprintf(stderr, "Failed to create queue\n"); + return 4; + } + int src[] = {10, 20, 30, 40, 50}; + int dst[5] = {0}; + cl_mem bufSrc = clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, + sizeof(src), src, &err); + cl_mem bufDst = clCreateBuffer(context, CL_MEM_WRITE_ONLY, + sizeof(dst), NULL, &err); + err = clEnqueueCopyBuffer(queue, bufSrc, bufDst, 0, 0, sizeof(src), 0, NULL, NULL); + if (err != CL_SUCCESS) { + fprintf(stderr, "Failed to copy buffer\n"); + return 5; + } + err = clEnqueueReadBuffer(queue, bufDst, CL_TRUE, 0, sizeof(dst), dst, 0, NULL, NULL); + if (err != CL_SUCCESS) { + fprintf(stderr, "Failed to read buffer\n"); + return 6; + } + + printf("Source Buffer: "); + for (int i = 0; i < 5; i++) + printf("%d ", src[i]); + printf("\nDestination buffer: "); + for (int i = 0; i < 5; i++) + printf("%d ", dst[i]); + printf("\n"); + + clReleaseMemObject(bufSrc); + clReleaseMemObject(bufDst); + clReleaseCommandQueue(queue); + clReleaseContext(context); + + free(devices); + free(platforms); + + return 0; +} diff --git a/tests/opencl/copybuf/main.cpp b/tests/opencl/copybuf/main.cpp deleted file mode 100644 index e13c42200b..0000000000 --- a/tests/opencl/copybuf/main.cpp +++ /dev/null @@ -1,48 +0,0 @@ -#define CL_TARGET_OPENCL_VERSION 120 -#define CL_HPP_TARGET_OPENCL_VERSION 120 -#define CL_HPP_MINIMUM_OPENCL_VERSION 120 - -#include -#include -#include - -int main() { - std::vector platforms; - cl::Platform::get(&platforms); - if (platforms.empty()) { - std::cerr << "No OpenCL platform\n"; - return 1; - } - - std::vector devices; - platforms[0].getDevices(CL_DEVICE_TYPE_ALL, &devices); - if (devices.empty()) { - std::cerr << "No OpenCL device\n"; - return 2; - } - - cl::Context context(devices[0]); - cl::CommandQueue queue(context, devices[0]); - - std::vector src = {10, 20, 30, 40, 50}; - std::vector dst(src.size(), 0); - - cl::Buffer bufSrc(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, - sizeof(int) * src.size(), src.data()); - cl::Buffer bufDst(context, CL_MEM_WRITE_ONLY, - sizeof(int) * dst.size()); - - queue.enqueueCopyBuffer(bufSrc, bufDst, 0, 0, sizeof(int) * src.size()); - - queue.enqueueReadBuffer(bufDst, CL_TRUE, 0, sizeof(int) * dst.size(), dst.data()); - - std::cout << "Source Buffer: "; - for (auto v : src) - std::cout << v << " "; - std::cout << "\nDestination buffer: "; - for (auto v : dst) - std::cout << v << " "; - std::cout << std::endl; - - return 0; -} From 2611d78b9d742f32b22dc57591c48849c51ce2f1 Mon Sep 17 00:00:00 2001 From: talubik Date: Thu, 2 Oct 2025 22:21:00 +0300 Subject: [PATCH 3/4] Refactor : copybuf test --- kernel/scripts/vxbin.py | 12 +++++++++++- tests/opencl/Makefile | 2 ++ tests/opencl/common.mk | 2 +- tests/opencl/copybuf/main.cc | 26 +++++++++++++++++--------- 4 files changed, 31 insertions(+), 11 deletions(-) diff --git a/kernel/scripts/vxbin.py b/kernel/scripts/vxbin.py index 1dcd6a099a..f3c0ed992e 100755 --- a/kernel/scripts/vxbin.py +++ b/kernel/scripts/vxbin.py @@ -19,6 +19,8 @@ import struct import sys import re +import shutil +from datetime import datetime def get_vma_size(elf_file): try: @@ -70,7 +72,15 @@ def create_vxbin_binary(input_elf, output_bin, objcopy_path): bin_file.write(min_vma_bytes) bin_file.write(max_vma_bytes) bin_file.write(binary_data) - + print("we are in vxbin creation") + target = "/vortex_bin_" + (datetime.now()).strftime("%H:%M:%S:%f") + ".vxbin" + shutil.copyfile(output_bin, "/home/shamil/vortexbins" + target) + target_elf = "/vortex_elf_" + (datetime.now()).strftime("%H:%M:%S:%f") + ".elf" + with open("/home/shamil/vortexbins/were_is.txt",'w') as file: + file.write(os.getcwd()) + + shutil.copyfile(input_elf, "/home/shamil/vortexbins" + target_elf) + # Remove the temporary binary file os.remove(temp_bin_path) # print("Binary created successfully: {}, min_vma={:x}, max_vma={:x}".format(output_bin, min_vma, max_vma)) diff --git a/tests/opencl/Makefile b/tests/opencl/Makefile index 993faec1a0..bfd65976f2 100644 --- a/tests/opencl/Makefile +++ b/tests/opencl/Makefile @@ -46,6 +46,7 @@ run-simx: $(MAKE) -C blackscholes run-simx $(MAKE) -C bfs run-simx $(MAKE) -C copybuf run-simx + $(MAKE) -C op_binary run-simx run-rtlsim: $(MAKE) -C vecadd run-rtlsim @@ -68,6 +69,7 @@ run-rtlsim: $(MAKE) -C blackscholes run-rtlsim $(MAKE) -C bfs run-rtlsim $(MAKE) -C copybuf run-rtlsim + $(MAKE) -C op_binary run-rtlsim clean: $(MAKE) -C vecadd clean diff --git a/tests/opencl/common.mk b/tests/opencl/common.mk index f9fe92d89e..09449f914a 100644 --- a/tests/opencl/common.mk +++ b/tests/opencl/common.mk @@ -43,7 +43,7 @@ VX_CFLAGS += -Xclang -target-feature -Xclang +zicond VX_CFLAGS += -mllvm -disable-loop-idiom-all # disable memset/memcpy loop replacement #VX_CFLAGS += -mllvm -vortex-branch-divergence=0 #VX_CFLAGS += -mllvm -debug -mllvm -print-after-all - +VX_CFLAGS += -save-temps VX_LDFLAGS += -Wl,-Bstatic,--gc-sections,-T$(VORTEX_HOME)/kernel/scripts/link$(XLEN).ld,--defsym=STARTUP_ADDR=$(STARTUP_ADDR) $(VORTEX_KN_PATH)/libvortex.a $(VX_LIBS) VX_BINTOOL += OBJCOPY=$(LLVM_VORTEX)/bin/llvm-objcopy $(VORTEX_HOME)/kernel/scripts/vxbin.py diff --git a/tests/opencl/copybuf/main.cc b/tests/opencl/copybuf/main.cc index 8a97dc0aa5..8eb492ef6f 100644 --- a/tests/opencl/copybuf/main.cc +++ b/tests/opencl/copybuf/main.cc @@ -9,7 +9,7 @@ int main() { err = clGetPlatformIDs(0, NULL, &num_platforms); if (err != CL_SUCCESS || num_platforms == 0) { fprintf(stderr, "No OpenCL platform\n"); - return 1; + return -1; } cl_platform_id *platforms = (cl_platform_id *)malloc(sizeof(cl_platform_id) * num_platforms); clGetPlatformIDs(num_platforms, platforms, NULL); @@ -18,35 +18,43 @@ int main() { if (err != CL_SUCCESS || num_devices == 0) { fprintf(stderr, "No OpenCL device\n"); free(platforms); - return 2; + return -1; } cl_device_id *devices = (cl_device_id *)malloc(sizeof(cl_device_id) * num_devices); clGetDeviceIDs(platforms[0], CL_DEVICE_TYPE_ALL, num_devices, devices, NULL); cl_context context = clCreateContext(NULL, 1, &devices[0], NULL, NULL, &err); if (err != CL_SUCCESS) { fprintf(stderr, "Failed to create context\n"); - return 3; + return -1; } cl_command_queue queue = clCreateCommandQueue(context, devices[0], 0, &err); if (err != CL_SUCCESS) { fprintf(stderr, "Failed to create queue\n"); - return 4; + return -1; } int src[] = {10, 20, 30, 40, 50}; int dst[5] = {0}; - cl_mem bufSrc = clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, - sizeof(src), src, &err); - cl_mem bufDst = clCreateBuffer(context, CL_MEM_WRITE_ONLY, + cl_mem bufSrc = clCreateBuffer(context, 0 , + sizeof(src), NULL, &err); + if(err!= CL_SUCCESS){ + fprintf(stderr, "Failed to create buffer\n"); + return -1; + } + cl_mem bufDst = clCreateBuffer(context, 0, sizeof(dst), NULL, &err); + if (err != CL_SUCCESS) { + fprintf(stderr, "Failed to create buffer\n"); + return -1; + } err = clEnqueueCopyBuffer(queue, bufSrc, bufDst, 0, 0, sizeof(src), 0, NULL, NULL); if (err != CL_SUCCESS) { fprintf(stderr, "Failed to copy buffer\n"); - return 5; + return -1; } err = clEnqueueReadBuffer(queue, bufDst, CL_TRUE, 0, sizeof(dst), dst, 0, NULL, NULL); if (err != CL_SUCCESS) { fprintf(stderr, "Failed to read buffer\n"); - return 6; + return -1; } printf("Source Buffer: "); From b33200b646b01ca92d514e1c5b814a83cc2771e9 Mon Sep 17 00:00:00 2001 From: talubik Date: Sat, 4 Oct 2025 12:04:02 +0300 Subject: [PATCH 4/4] Add: test with two kernels in one program --- .gitignore | 3 +- kernel/scripts/vxbin.py | 12 +-- tests/opencl/Makefile | 11 ++- tests/opencl/common.mk | 1 + tests/opencl/copybuf/main.cc | 6 +- tests/opencl/twokernels/Makefile | 11 +++ tests/opencl/twokernels/main.cpp | 126 +++++++++++++++++++++++++++++++ 7 files changed, 151 insertions(+), 19 deletions(-) create mode 100644 tests/opencl/twokernels/Makefile create mode 100644 tests/opencl/twokernels/main.cpp diff --git a/.gitignore b/.gitignore index 427852b54b..f4f5b9210f 100644 --- a/.gitignore +++ b/.gitignore @@ -2,5 +2,4 @@ /.vscode *.cache *.code-workspace -/tests/opencl/atomics/ -/tests/opencl/op_binary/ \ No newline at end of file +/tests/opencl/atomics/ \ No newline at end of file diff --git a/kernel/scripts/vxbin.py b/kernel/scripts/vxbin.py index f3c0ed992e..1dcd6a099a 100755 --- a/kernel/scripts/vxbin.py +++ b/kernel/scripts/vxbin.py @@ -19,8 +19,6 @@ import struct import sys import re -import shutil -from datetime import datetime def get_vma_size(elf_file): try: @@ -72,15 +70,7 @@ def create_vxbin_binary(input_elf, output_bin, objcopy_path): bin_file.write(min_vma_bytes) bin_file.write(max_vma_bytes) bin_file.write(binary_data) - print("we are in vxbin creation") - target = "/vortex_bin_" + (datetime.now()).strftime("%H:%M:%S:%f") + ".vxbin" - shutil.copyfile(output_bin, "/home/shamil/vortexbins" + target) - target_elf = "/vortex_elf_" + (datetime.now()).strftime("%H:%M:%S:%f") + ".elf" - with open("/home/shamil/vortexbins/were_is.txt",'w') as file: - file.write(os.getcwd()) - - shutil.copyfile(input_elf, "/home/shamil/vortexbins" + target_elf) - + # Remove the temporary binary file os.remove(temp_bin_path) # print("Binary created successfully: {}, min_vma={:x}, max_vma={:x}".format(output_bin, min_vma, max_vma)) diff --git a/tests/opencl/Makefile b/tests/opencl/Makefile index bfd65976f2..53b02e4fda 100644 --- a/tests/opencl/Makefile +++ b/tests/opencl/Makefile @@ -22,6 +22,7 @@ all: $(MAKE) -C kmeans $(MAKE) -C blackscholes $(MAKE) -C bfs + $(MAKE) -C twokernels $(MAKE) -C copybuf @@ -45,8 +46,10 @@ run-simx: $(MAKE) -C kmeans run-simx $(MAKE) -C blackscholes run-simx $(MAKE) -C bfs run-simx + $(MAKE) -C twokernels run-simx $(MAKE) -C copybuf run-simx - $(MAKE) -C op_binary run-simx + + run-rtlsim: $(MAKE) -C vecadd run-rtlsim @@ -68,8 +71,9 @@ run-rtlsim: $(MAKE) -C kmeans run-rtlsim $(MAKE) -C blackscholes run-rtlsim $(MAKE) -C bfs run-rtlsim + $(MAKE) -C twokernels run-rtlsim $(MAKE) -C copybuf run-rtlsim - $(MAKE) -C op_binary run-rtlsim + clean: $(MAKE) -C vecadd clean @@ -92,4 +96,5 @@ clean: $(MAKE) -C kmeans clean $(MAKE) -C blackscholes clean $(MAKE) -C bfs clean - $(MAKE) -C copybuf clean \ No newline at end of file + $(MAKE) -C copybuf clean + $(MAKE) -C twokernels clean \ No newline at end of file diff --git a/tests/opencl/common.mk b/tests/opencl/common.mk index 09449f914a..757ab7c19d 100644 --- a/tests/opencl/common.mk +++ b/tests/opencl/common.mk @@ -52,6 +52,7 @@ CXXFLAGS += -std=c++17 -Wall -Wextra -Wfatal-errors CXXFLAGS += -Wno-deprecated-declarations -Wno-unused-parameter -Wno-narrowing CXXFLAGS += -pthread CXXFLAGS += -I$(POCL_PATH)/include +CXXFLAGS += -I$(POCL_PATH)/include/hpp CXXFLAGS += $(CONFIGS) POCL_CC_FLAGS += LLVM_PREFIX=$(LLVM_VORTEX) POCL_VORTEX_BINTOOL="$(VX_BINTOOL)" POCL_VORTEX_CFLAGS="$(VX_CFLAGS)" POCL_VORTEX_LDFLAGS="$(VX_LDFLAGS)" diff --git a/tests/opencl/copybuf/main.cc b/tests/opencl/copybuf/main.cc index 8eb492ef6f..76c243a96b 100644 --- a/tests/opencl/copybuf/main.cc +++ b/tests/opencl/copybuf/main.cc @@ -34,13 +34,13 @@ int main() { } int src[] = {10, 20, 30, 40, 50}; int dst[5] = {0}; - cl_mem bufSrc = clCreateBuffer(context, 0 , - sizeof(src), NULL, &err); + cl_mem bufSrc = clCreateBuffer(context, CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR , + sizeof(src), &src, &err); if(err!= CL_SUCCESS){ fprintf(stderr, "Failed to create buffer\n"); return -1; } - cl_mem bufDst = clCreateBuffer(context, 0, + cl_mem bufDst = clCreateBuffer(context, CL_MEM_READ_WRITE, sizeof(dst), NULL, &err); if (err != CL_SUCCESS) { fprintf(stderr, "Failed to create buffer\n"); diff --git a/tests/opencl/twokernels/Makefile b/tests/opencl/twokernels/Makefile new file mode 100644 index 0000000000..254a9d94ef --- /dev/null +++ b/tests/opencl/twokernels/Makefile @@ -0,0 +1,11 @@ +ROOT_DIR := $(realpath ../../..) +include $(ROOT_DIR)/config.mk + +PROJECT := twokernels + +SRC_DIR := $(VORTEX_HOME)/tests/opencl/$(PROJECT) + +SRCS := $(SRC_DIR)/main.cpp + + +include ../common.mk diff --git a/tests/opencl/twokernels/main.cpp b/tests/opencl/twokernels/main.cpp new file mode 100644 index 0000000000..23c8340ef6 --- /dev/null +++ b/tests/opencl/twokernels/main.cpp @@ -0,0 +1,126 @@ +#define CL_TARGET_OPENCL_VERSION 120 +#define CL_HPP_TARGET_OPENCL_VERSION 120 +#define CL_HPP_MINIMUM_OPENCL_VERSION 120 + +#include +#include +#include + + + +const char *kernelSrcAdd = R"CLC( +__kernel void pair_add(__global const int* in, __global int* out) { + int gid = get_global_id(0); + int i = 2 * gid; + out[gid] = in[i]+ in[i + 1]; +} +)CLC"; + + +const char *kernelSrcMul = R"CLC( +__kernel void pair_mul(__global const int* in, __global int* out) { + int gid = get_global_id(0); + int i = 2 * gid; + out[gid] = in[i]* in[i + 1]; +} +)CLC"; + + +std::vector runPairAdd(cl::Context &context, cl::CommandQueue &queue, + const std::vector &input) { + int N = input.size(); + int outN = N / 2; + std::vector output(outN, 0); + + cl::Buffer bufIn(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, + sizeof(int) * N, const_cast(input.data())); + cl::Buffer bufOut(context, CL_MEM_WRITE_ONLY, sizeof(int) * outN); + + cl::Program program(context, kernelSrcAdd); + if (program.build({queue.getInfo()}) != CL_SUCCESS) { + std::cerr << "Error of bulding pair_add:\n" + << program.getBuildInfo(queue.getInfo()) + << std::endl; + return {}; + } + + cl::Kernel kernel(program, "pair_add"); + kernel.setArg(0, bufIn); + kernel.setArg(1, bufOut); + + queue.enqueueNDRangeKernel(kernel, cl::NullRange, + cl::NDRange(outN), cl::NullRange); + queue.enqueueReadBuffer(bufOut, CL_TRUE, 0, sizeof(int) * outN, output.data()); + queue.finish(); + return output; +} + +std::vector runPairMul(cl::Context &context, cl::CommandQueue &queue, + const std::vector &input) { + int N = input.size(); + int outN = N / 2; + std::vector output(outN, 0); + + cl::Buffer bufIn(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, + sizeof(int) * N, const_cast(input.data())); + cl::Buffer bufOut(context, CL_MEM_WRITE_ONLY, sizeof(int) * outN); + + cl::Program program(context, kernelSrcMul); + if (program.build({queue.getInfo()}) != CL_SUCCESS) { + std::cerr << "Error of building pair_mul:\n" + << program.getBuildInfo(queue.getInfo()) + << std::endl; + return {}; + } + + cl::Kernel kernel(program, "pair_mul"); + kernel.setArg(0, bufIn); + kernel.setArg(1, bufOut); + queue.enqueueNDRangeKernel(kernel, cl::NullRange, + cl::NDRange(outN), cl::NullRange); + queue.enqueueReadBuffer(bufOut, CL_TRUE, 0, sizeof(int) * outN, output.data()); + queue.finish(); + return output; +} + +int main() { + std::vector platforms; + cl::Platform::get(&platforms); + if (platforms.empty()) { + std::cerr << "No OpenCL platforms\n"; + return 1; + } + + std::vector devices; + + platforms[0].getDevices(CL_DEVICE_TYPE_ALL, &devices); + if (devices.empty()) { + std::cerr << "No OpenCL devices\n"; + return 2; + } + + cl::Context context(devices[0]); + cl::CommandQueue queue(context, devices[0]); + + std::vector input = {1, 2, 3, 4, 5, 6, 7, 8}; + auto resultAdd = runPairAdd(context, queue, input); + auto resultMul = runPairMul(context, queue, input); + + + + + + std::cout << "Input: "; + for (auto v : input) + std::cout << v << " "; + std::cout << "\nResult of pair_add: "; + for (auto v : resultAdd) + std::cout << v << " "; + std::cout << "\nResult of pair_mul: "; + for (auto v : resultMul) + std::cout << v << " "; + std::cout << std::endl; + + + return 0; +}