Skip to content

Commit 42656d2

Browse files
sheralskumarsheralskumar
committed
cpp OpenCL and kernel code for vector addition
0 parents  commit 42656d2

3 files changed

Lines changed: 206 additions & 0 deletions

File tree

‎CMakeLists.txt‎

Lines changed: 15 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,15 @@
1+
cmake_minimum_required(VERSION 3.10)
2+
project(opencl_bench)
3+
4+
set(CMAKE_CXX_STANDARD 17) # Use C++17
5+
set(CMAKE_CXX_STANDARD_REQUIRED ON)
6+
7+
set(OpenCL_INCLUDE_DIR "")
8+
set(OpenCL_LIBRARY "")
9+
10+
11+
include_directories(${OpenCL_INCLUDE_DIR})
12+
13+
add_executable(benchmark src/benchmark.cpp)
14+
15+
target_link_libraries(benchmark ${OpenCL_LIBRARY})

‎src/benchmark.cpp‎

Lines changed: 185 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,185 @@
1+
2+
#include <CL/cl.h> // library for OpenCL
3+
#include <iostream>
4+
#include <fstream>
5+
#include <vector>
6+
#include <chrono>
7+
#include <filesystem>
8+
9+
// error checking macro for OpenCL calls to determine whether they were successful
10+
// CL_SUCCESS is defined in cl.h as 0
11+
#define CHECK(x) if((x) != CL_SUCCESS){std::cerr << "OpenCL error " << x << std::endl; exit(1);}
12+
13+
// function to load the contents of a file into a string
14+
// used to read OpenCL kernel source code from a file
15+
std::string load_file(const std::filesystem::path& path) {
16+
std::ifstream f(path);
17+
if (!f.is_open()) {
18+
std::cerr << "Failed to open " << path << std::endl;
19+
exit(1);
20+
}
21+
return std::string((std::istreambuf_iterator<char>(f)), std::istreambuf_iterator<char>());
22+
}
23+
24+
// main function
25+
int main(int argc, char** argv) {
26+
// read in command line arguments for device index and vector size
27+
int device_index = (argc > 1) ? atoi(argv[1]) : 0;
28+
// int N = (argc > 2) ? atoi(argv[2]) : (1 << 22);
29+
int N = (argc > 2) ? atoi(argv[2]) : (1 << 26);
30+
31+
// 1. Select openCL platform using the clGetPlatformIDs function
32+
// note platform vs device: platform is the vendor (e.g., NVIDIA, AMD, Intel), device is the actual hardware (e.g., GPU, CPU)
33+
cl_uint num_platforms; // cl_uint is OpenCL-defined (I think unsigned int32) to make sure code is platform independent
34+
35+
// when num_entries is 0 and platforms is NULL, clGetPlatformIDs returns the number of platforms available
36+
CHECK(clGetPlatformIDs(0, NULL, &num_platforms));
37+
// when num_entries is non-zero and platforms is not NULL, clGetPlatformIDs fills the platforms array with platform IDs
38+
std::vector<cl_platform_id> platforms(num_platforms);
39+
CHECK(clGetPlatformIDs(num_platforms, platforms.data(), NULL));
40+
41+
// 2. Now select a device on the platform using clGetDeviceIDs
42+
std::vector<cl_device_id> devices;
43+
for (auto p : platforms) {
44+
cl_uint dev_count; // store number of devices on the platform
45+
cl_uint err;
46+
47+
// get number of devices on a platform p, similar to the way we got number of platforms
48+
// CL_DEVICE_TYPE_ALL specifies we want all types of devices (CPU, GPU, etc.)
49+
// CL_DEVICE_TYPE_GPU could be used to select only GPU devices
50+
err = clGetDeviceIDs(p, CL_DEVICE_TYPE_ALL, 0, NULL, &dev_count);
51+
if (err == CL_DEVICE_NOT_FOUND) { continue; } // skip platform with no devices
52+
CHECK(err);
53+
54+
// we can now get the device IDs for platform p, similar to how we got platform IDs
55+
std::vector<cl_device_id> devs(dev_count);
56+
clGetDeviceIDs(p, CL_DEVICE_TYPE_ALL, dev_count, devs.data(), NULL);
57+
devices.insert(devices.end(), devs.begin(), devs.end()); // append this platform's devices to the overall devices list
58+
}
59+
60+
// error check to see if device index requested by the user is valid
61+
if (device_index >= devices.size()) {
62+
std::cerr << "Device index out of range.\n";
63+
return 1;
64+
}
65+
66+
// grab the device ID for the requested device
67+
cl_device_id dev = devices[device_index];
68+
69+
// get the name of the device and print it out
70+
char name[128];
71+
clGetDeviceInfo(dev, CL_DEVICE_NAME, 128, name, NULL);
72+
std::cout << "Using device: " << name << std::endl;
73+
74+
// 3. Create an OpenCL context and command queue on that device
75+
// the context is like a container for all OpenCL objects (buffers, programs, kernels, etc.) associated with the device
76+
// the queue is used to submit work (kernel executions, memory transfers, etc.) to the device
77+
// operations must be done through the command queue
78+
cl_int err;
79+
cl_context ctx = clCreateContext(NULL, 1, &dev, NULL, NULL, &err); // properties, num_devices, array of device ids (dev), call back for errors, user data for call back, error code
80+
CHECK(err);
81+
cl_command_queue queue = clCreateCommandQueueWithProperties(ctx, dev, 0, &err); // context, device, properties, error code
82+
CHECK(err);
83+
84+
// 4. Create and build the OpenCL program from source
85+
// need to know about program objects and kernel objects
86+
// program object is container for OpenCL code (kernels) that will be executed on the device (can contain multiple kernels)
87+
// kernel objects created/managed by program object represent individual functions (kernels) in the OpenCL code that can be executed on the device
88+
// for ex. an algebraic program might have separate kernels for vector addition, matrix multiplication, etc. in a single program object
89+
// to create a kernel, we need to prepare source code in OpenCL C (similar to C99) and load it into a string
90+
91+
// std::string src = load_file("../src/kernel.cl"); // load file wth the C source code as a string
92+
std::filesystem::path exe_path = std::filesystem::absolute(argv[0]).parent_path();
93+
std::filesystem::path kernel_path = exe_path.parent_path() / "src" / "kernel.cl";
94+
std::string src = load_file(kernel_path);
95+
96+
const char* csrc = src.c_str(); // convert it to C-style character array
97+
size_t len = src.size();
98+
// create the program object from the source code string
99+
cl_program program = clCreateProgramWithSource(ctx, 1, &csrc, &len, &err); // context, number of strings, array of strings, array of string lengths, error code
100+
CHECK(err);
101+
102+
// build the program object
103+
err = clBuildProgram(program, 1, &dev, NULL, NULL, NULL); // program, num_devices, array of device ids, build options, call back, user data
104+
if (err != CL_SUCCESS) { // if that fails, get and print the build log
105+
size_t log_size;
106+
clGetProgramBuildInfo(program, dev, CL_PROGRAM_BUILD_LOG, 0, NULL, &log_size);
107+
std::string log(log_size, '\0');
108+
clGetProgramBuildInfo(program, dev, CL_PROGRAM_BUILD_LOG, log_size, &log[0], NULL);
109+
std::cerr << "Build error:\n" << log << std::endl;
110+
return 1;
111+
}
112+
113+
// 5. Create kernel object from the built program, extract it by its name
114+
cl_kernel kernel = clCreateKernel(program, "vector_add", &err);
115+
CHECK(err);
116+
117+
// 6. Create device buffers and transfer input data to the device
118+
// buffers are memory allocations on the device that can be read from and written to by kernels
119+
// openCL defines three memory types: buffer, image, and pipe
120+
// buffer stores contiguous linear data (like arrays/vectors/matrices) and can be accessed on the device using pointers, created using clCreateBuffer
121+
// image is optimized for 2D/3D data and provides built-in functions for sampling and filtering (used in graphics/vision applications)
122+
// pipe is used for streaming data between kernels
123+
// here we use buffer memory type to create buffers for our input/output vectors
124+
125+
// make three buffers: two for input vectors and one for output vector, fill the input buffers with data (host)
126+
size_t bytes = N*sizeof(float);
127+
std::vector<float> A(N), B(N), C(N);
128+
for(int i=0;i<N;i++){
129+
A[i] = i;
130+
B[i] = 2*i;
131+
}
132+
133+
// create device buffers and copy input data from host to device
134+
// note bitwise or operator (|) is used to combine multiple memory flags
135+
cl_mem dA = clCreateBuffer(ctx, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, bytes, A.data(), &err); // context, flags, size in bytes, host pointer (data to copy), error code
136+
CHECK(err);
137+
cl_mem dB = clCreateBuffer(ctx, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, bytes, B.data(), &err);
138+
CHECK(err);
139+
cl_mem dC = clCreateBuffer(ctx, CL_MEM_WRITE_ONLY, bytes, NULL, &err);
140+
CHECK(err);
141+
142+
// 7. Set kernel arguments as the buffers created above
143+
clSetKernelArg(kernel, 0, sizeof(cl_mem), &dA); // kernel, argument index, size of argument, pointer to argument value
144+
clSetKernelArg(kernel, 1, sizeof(cl_mem), &dB);
145+
clSetKernelArg(kernel, 2, sizeof(cl_mem), &dC);
146+
clSetKernelArg(kernel, 3, sizeof(int), &N);
147+
148+
// 8. Kernel execution
149+
// specify the global and local work sizes. Global work size is total number of work items (threads) to execute the kernel
150+
// local work size is number of work items in a work group (like CUDA blocks), just a way to organize work items for better performance
151+
// 256 is a common choice for local work size
152+
size_t global = N;
153+
size_t local = 256;
154+
155+
auto start = std::chrono::high_resolution_clock::now(); // start timer for benchmarking
156+
CHECK(clEnqueueNDRangeKernel(queue, kernel, 1, NULL, &global, &local, 0, NULL, NULL)); // command queue, kernel, work dimension, global work offset, global work size, local work size, num events in wait list, event wait list, event
157+
clFinish(queue); // wait for all commands in the queue to finish
158+
auto end = std::chrono::high_resolution_clock::now(); // end timer
159+
160+
double ms = std::chrono::duration<double, std::milli>(end-start).count();
161+
std::cout << "Time: " << ms << " ms\n"; // print elapsed time
162+
163+
// read back the result from device to host, queue, buffer, blocking read, offset, size, host pointer, num events in wait list, event wait list, event
164+
clEnqueueReadBuffer(queue, dC, CL_TRUE, 0, bytes, C.data(), 0, NULL, NULL);
165+
166+
// 9. Verify the result
167+
bool ok = true;
168+
for (int i=0;i<10;i++) {
169+
if (C[i] != A[i] + B[i]) { ok = false; }
170+
}
171+
std::cout << "Correct: " << (ok ? "yes" : "no") << "\n";
172+
173+
// 10. Cleanup OpenCL resources
174+
// release device buffers
175+
clReleaseMemObject(dA);
176+
clReleaseMemObject(dB);
177+
clReleaseMemObject(dC);
178+
179+
// release kernel, program, queue, and context
180+
clReleaseKernel(kernel);
181+
clReleaseProgram(program);
182+
clReleaseCommandQueue(queue);
183+
clReleaseContext(ctx);
184+
return 0;
185+
}

‎src/kernel.cl‎

Lines changed: 6 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,6 @@
1+
__kernel void vector_add(__global const float* A, __global const float* B, __global float* C, int N)
2+
{
3+
int i = get_global_id(0);
4+
if (i < N)
5+
C[i] = A[i] + B[i];
6+
}

0 commit comments

Comments
 (0)