Skip to content

Commit 67b734a

Browse files
umar456umar456
committed
Update OpenCL interop page so they discuss deleting of memory
* Created snippets for examples in the document
1 parent bcdf0ba commit 67b734a

5 files changed

Lines changed: 222 additions & 121 deletions

File tree

‎docs/pages/interop_opencl.md‎

Lines changed: 3 additions & 119 deletions
Original file line numberDiff line numberDiff line change
@@ -64,68 +64,7 @@ synchronization operations.
6464

6565
This process is best illustrated with a fully worked example:
6666

67-
~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~{.cpp}
68-
#include <arrayfire.h>
69-
// 1. Add the af/opencl.h include to your project
70-
#include <af/opencl.h>
71-
72-
int main() {
73-
size_t length = 10;
74-
75-
// Create ArrayFire array objects:
76-
af::array A = af::randu(length, f32);
77-
af::array B = af::constant(0, length, f32);
78-
79-
// ... additional ArrayFire operations here
80-
81-
// 2. Obtain the device, context, and queue used by ArrayFire
82-
static cl_context af_context = afcl::getContext();
83-
static cl_device_id af_device_id = afcl::getDeviceId();
84-
static cl_command_queue af_queue = afcl::getQueue();
85-
86-
// 3. Obtain cl_mem references to af::array objects
87-
cl_mem * d_A = A.device<cl_mem>();
88-
cl_mem * d_B = B.device<cl_mem>();
89-
90-
// 4. Load, build, and use your kernels.
91-
// For the sake of readability, we have omitted error checking.
92-
int status = CL_SUCCESS;
93-
94-
// A simple copy kernel, uses C++11 syntax for multi-line strings.
95-
const char * kernel_name = "copy_kernel";
96-
const char * source = R"(
97-
void __kernel
98-
copy_kernel(__global float * gA, __global float * gB)
99-
{
100-
int id = get_global_id(0);
101-
gB[id] = gA[id];
102-
}
103-
)";
104-
105-
// Create the program, build the executable, and extract the entry point
106-
// for the kernel.
107-
cl_program program = clCreateProgramWithSource(af_context, 1, &source, NULL, &status);
108-
status = clBuildProgram(program, 1, &af_device_id, NULL, NULL, NULL);
109-
cl_kernel kernel = clCreateKernel(program, kernel_name, &status);
110-
111-
// Set arguments and launch your kernels
112-
clSetKernelArg(kernel, 0, sizeof(cl_mem), d_A);
113-
clSetKernelArg(kernel, 1, sizeof(cl_mem), d_B);
114-
clEnqueueNDRangeKernel(af_queue, kernel, 1, NULL, &length, NULL, 0, NULL, NULL);
115-
116-
// 5. Return control of af::array memory to ArrayFire
117-
A.unlock();
118-
B.unlock();
119-
120-
// ... resume ArrayFire operations
121-
122-
// Because the device pointers, d_x and d_y, were returned to ArrayFire's
123-
// control by the unlock function, there is no need to free them using
124-
// clReleaseMemObject()
125-
126-
return 0;
127-
}
128-
~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
67+
\snippet test/interop_opencl_custom_kernel_snippet.cpp interop_opencl_custom_kernel_snippet
12968

13069
If your kernels needs to operate in their own OpenCL queue, the process is
13170
essentially identical, except you need to instruct ArrayFire to complete
@@ -187,64 +126,9 @@ so, please be cautious not to call `clReleaseMemObj` on a `cl_mem` when
187126
ArrayFire might be using it!
188127

189128
The eight steps above are best illustrated using a fully-worked example. Below we
190-
use the OpenCL 2.0 C++ API and omit error checking to keep the code readable.
191-
192-
~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~{.cpp}
193-
#include <CL/cl2.hpp>
194-
195-
// 1. Add arrayfire.h and af/opencl.h to your application
196-
#include "arrayfire.h"
197-
#include "af/opencl.h"
198-
199-
#include <cstdio>
200-
#include <vector>
201-
202-
int main() {
203-
204-
// Set up the OpenCL context, device, and queues
205-
cl::Context context(CL_DEVICE_TYPE_ALL);
206-
vector<cl::Device> devices = context.getInfo<CL_CONTEXT_DEVICES>();
207-
cl::Device device = devices[0];
208-
cl::CommandQueue queue(context, device);
209-
210-
// Create a buffer of size 10 filled with ones, copy it to the device
211-
int length = 10;
212-
vector<float> h_A(length, 1);
213-
cl::Buffer cl_A(context, CL_MEM_READ_WRITE, length * sizeof(float), h_A.data());
129+
use the OpenCL C++ API and omit error checking to keep the code readable.
214130

215-
// 2. Instruct OpenCL to complete its operations using clFinish (or similar)
216-
queue.finish();
217-
218-
// 3. Instruct ArrayFire to use the user-created context
219-
// First, create a device from the current OpenCL device + context + queue
220-
afcl::addDevice(device(), context(), queue());
221-
// Next switch ArrayFire to the device using the device and context as
222-
// identifiers:
223-
afcl::setDevice(device(), context());
224-
225-
// 4. Create ArrayFire arrays from OpenCL memory objects
226-
af::array af_A = afcl::array(length, cl_A(), f32, true);
227-
228-
// 5. Perform ArrayFire operations on the Arrays
229-
af_A = af_A + af::randu(length);
230-
231-
// NOTE: ArrayFire does not perform the above transaction using in-place memory,
232-
// thus the underlying OpenCL buffers containing the memory containing memory to
233-
// probably have changed
234-
235-
// 6. Instruct ArrayFire to finish operations using af::sync
236-
af::sync();
237-
238-
// 7. Obtain cl_mem references for important memory
239-
cl_A = *af_A.device<cl_mem>();
240-
241-
// 8. Continue your OpenCL application
242-
243-
// ...
244-
245-
return 0;
246-
}
247-
~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
131+
\snippet test/interop_opencl_external_context_snippet.cpp interop_opencl_external_context_snippet
248132

249133
# Using multiple devices
250134

‎include/af/array.h‎

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -725,6 +725,8 @@ namespace af
725725
726726
The device memory returned by this function is not freed until unlock() is called.
727727
728+
/note When using the OpenCL backend and using the cl_mem template argument, the
729+
delete function should be called on the pointer returned by this function.
728730
*/
729731
template<typename T> T* device() const;
730732

‎test/CMakeLists.txt‎

Lines changed: 17 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -115,7 +115,7 @@ target_compile_definitions(arrayfire_test
115115
# 'BACKENDS' Backends to target for this test. If not set then the test will
116116
# compiled againat all backends
117117
function(make_test)
118-
set(options CXX11 SERIAL USE_MMIO)
118+
set(options CXX11 SERIAL USE_MMIO NO_ARRAYFIRE_TEST)
119119
set(single_args SRC)
120120
set(multi_args LIBRARIES DEFINITIONS BACKENDS)
121121
cmake_parse_arguments(mt_args "${options}" "${single_args}" "${multi_args}" ${ARGN})
@@ -127,7 +127,12 @@ function(make_test)
127127
continue()
128128
endif()
129129
set(target "test_${src_name}_${backend}")
130-
add_executable(${target} ${mt_args_SRC} $<TARGET_OBJECTS:arrayfire_test>)
130+
131+
if (${mt_args_NO_ARRAYFIRE_TEST})
132+
add_executable(${target} ${mt_args_SRC})
133+
else()
134+
add_executable(${target} ${mt_args_SRC} $<TARGET_OBJECTS:arrayfire_test>)
135+
endif()
131136
target_include_directories(${target}
132137
PRIVATE
133138
${ArrayFire_SOURCE_DIR}/extern/half/include
@@ -285,6 +290,16 @@ if(OpenCL_FOUND)
285290
LIBRARIES OpenCL::OpenCL
286291
BACKENDS "opencl"
287292
CXX11)
293+
make_test(SRC interop_opencl_custom_kernel_snippet.cpp
294+
LIBRARIES OpenCL::OpenCL
295+
BACKENDS "opencl"
296+
NO_ARRAYFIRE_TEST
297+
CXX11)
298+
make_test(SRC interop_opencl_external_context_snippet.cpp
299+
LIBRARIES OpenCL::OpenCL
300+
BACKENDS "opencl"
301+
NO_ARRAYFIRE_TEST
302+
CXX11)
288303
endif()
289304

290305
if(CUDA_FOUND)

‎test/interop_opencl_custom_kernel_snippet.cpp‎

Lines changed: 96 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,96 @@
1+
/*******************************************************
2+
* Copyright (c) 2020, ArrayFire
3+
* All rights reserved.
4+
*
5+
* This file is distributed under 3-clause BSD license.
6+
* The complete license agreement can be obtained at:
7+
* http://arrayfire.com/licenses/BSD-3-Clause
8+
********************************************************/
9+
10+
// clang-format off
11+
// ![interop_opencl_custom_kernel_snippet]
12+
#include <arrayfire.h>
13+
// 1. Add the af/opencl.h include to your project
14+
#include <af/opencl.h>
15+
16+
#include <cassert>
17+
18+
#define OCL_CHECK(call) \
19+
if (cl_int err = (call) != CL_SUCCESS) { \
20+
fprintf(stderr, __FILE__ "(%d):Returned error code %d\n", __LINE__, \
21+
err); \
22+
}
23+
24+
int main() {
25+
size_t length = 10;
26+
27+
// Create ArrayFire array objects:
28+
af::array A = af::randu(length, f32);
29+
af::array B = af::constant(0, length, f32);
30+
31+
// ... additional ArrayFire operations here
32+
33+
// 2. Obtain the device, context, and queue used by ArrayFire
34+
static cl_context af_context = afcl::getContext();
35+
static cl_device_id af_device_id = afcl::getDeviceId();
36+
static cl_command_queue af_queue = afcl::getQueue();
37+
38+
// 3. Obtain cl_mem references to af::array objects
39+
cl_mem* d_A = A.device<cl_mem>();
40+
cl_mem* d_B = B.device<cl_mem>();
41+
42+
// 4. Load, build, and use your kernels.
43+
// For the sake of readability, we have omitted error checking.
44+
int status = CL_SUCCESS;
45+
46+
// A simple copy kernel, uses C++11 syntax for multi-line strings.
47+
const char* kernel_name = "copy_kernel";
48+
const char* source = R"(
49+
void __kernel
50+
copy_kernel(__global float* gA, __global float* gB) {
51+
int id = get_global_id(0);
52+
gB[id] = gA[id];
53+
}
54+
)";
55+
56+
// Create the program, build the executable, and extract the entry point
57+
// for the kernel.
58+
cl_program program = clCreateProgramWithSource(af_context, 1, &source, NULL, &status);
59+
OCL_CHECK(status);
60+
OCL_CHECK(clBuildProgram(program, 1, &af_device_id, NULL, NULL, NULL));
61+
cl_kernel kernel = clCreateKernel(program, kernel_name, &status);
62+
OCL_CHECK(status);
63+
64+
// Set arguments and launch your kernels
65+
OCL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), d_A));
66+
OCL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_mem), d_B));
67+
OCL_CHECK(clEnqueueNDRangeKernel(af_queue, kernel, 1, NULL, &length, NULL,
68+
0, NULL, NULL));
69+
70+
// 5. Return control of af::array memory to ArrayFire
71+
A.unlock();
72+
B.unlock();
73+
74+
/// A and B should not be the same because of the copy_kernel user code
75+
assert(af::allTrue<bool>(A == B));
76+
77+
// Delete the pointers returned by the device function. This does NOT
78+
// delete the cl_mem memory and only deletes the pointers
79+
delete d_A;
80+
delete d_B;
81+
82+
// ... resume ArrayFire operations
83+
84+
// Because the device pointers, d_x and d_y, were returned to ArrayFire's
85+
// control by the unlock function, there is no need to free them using
86+
// clReleaseMemObject()
87+
88+
// Free the kernel and program objects because they are created in user
89+
// code
90+
OCL_CHECK(clReleaseKernel(kernel));
91+
OCL_CHECK(clReleaseProgram(program));
92+
93+
return 0;
94+
}
95+
// ![interop_opencl_custom_kernel_snippet]
96+
// clang-format on

‎test/interop_opencl_external_context_snippet.cpp‎

Lines changed: 104 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,104 @@
1+
/*******************************************************
2+
* Copyright (c) 2020, ArrayFire
3+
* All rights reserved.
4+
*
5+
* This file is distributed under 3-clause BSD license.
6+
* The complete license agreement can be obtained at:
7+
* http://arrayfire.com/licenses/BSD-3-Clause
8+
********************************************************/
9+
10+
#pragma GCC diagnostic push
11+
#pragma GCC diagnostic ignored "-Wunused-function"
12+
#pragma GCC diagnostic ignored "-Wunused-parameter"
13+
#pragma GCC diagnostic ignored "-Wignored-qualifiers"
14+
#pragma GCC diagnostic ignored "-Wignored-attributes"
15+
#pragma GCC diagnostic ignored "-Wdeprecated-declarations"
16+
#if __GNUC__ >= 8
17+
#pragma GCC diagnostic ignored "-Wcatch-value="
18+
#endif
19+
// ![interop_opencl_external_context_snippet]
20+
#include <arrayfire.h>
21+
// 1. Add the af/opencl.h include to your project
22+
#include <af/opencl.h>
23+
24+
#include <cassert>
25+
26+
// definitions required by cl2.hpp
27+
#define CL_HPP_ENABLE_EXCEPTIONS
28+
#define CL_HPP_TARGET_OPENCL_VERSION 120
29+
#define CL_HPP_MINIMUM_OPENCL_VERSION 120
30+
#include <CL/cl2.hpp>
31+
32+
// 1. Add arrayfire.h and af/opencl.h to your application
33+
#include "af/opencl.h"
34+
#include "arrayfire.h"
35+
36+
#include <cstdio>
37+
#include <vector>
38+
39+
using std::vector;
40+
41+
int main() {
42+
// 1. Set up the OpenCL context, device, and queues
43+
cl::Context context;
44+
try {
45+
context = cl::Context(CL_DEVICE_TYPE_ALL);
46+
} catch (const cl::Error& err) {
47+
fprintf(stderr, "Exiting creating context");
48+
return EXIT_FAILURE;
49+
}
50+
vector<cl::Device> devices = context.getInfo<CL_CONTEXT_DEVICES>();
51+
if (devices.empty()) {
52+
fprintf(stderr, "Exiting. No devices found");
53+
return EXIT_SUCCESS;
54+
}
55+
cl::Device device = devices[0];
56+
cl::CommandQueue queue(context, device);
57+
58+
// Create a buffer of size 10 filled with ones, copy it to the device
59+
int length = 10;
60+
vector<float> h_A(length, 1);
61+
cl::Buffer cl_A(context, CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR,
62+
length * sizeof(float), h_A.data());
63+
64+
// 2. Instruct OpenCL to complete its operations using clFinish (or similar)
65+
queue.finish();
66+
67+
// 3. Instruct ArrayFire to use the user-created context
68+
// First, create a device from the current OpenCL device + context +
69+
// queue
70+
afcl::addDevice(device(), context(), queue());
71+
// Next switch ArrayFire to the device using the device and context as
72+
// identifiers:
73+
afcl::setDevice(device(), context());
74+
75+
// 4. Create ArrayFire arrays from OpenCL memory objects
76+
af::array af_A = afcl::array(length, cl_A(), f32, true);
77+
clRetainMemObject(cl_A());
78+
79+
// 5. Perform ArrayFire operations on the Arrays
80+
af_A = af_A + af::randu(length);
81+
82+
// NOTE: ArrayFire does not perform the above transaction using in-place
83+
// memory, thus the underlying OpenCL buffers containing the memory
84+
// containing memory to probably have changed
85+
86+
// 6. Instruct ArrayFire to finish operations using af::sync
87+
af::sync();
88+
89+
// 7. Obtain cl_mem references for important memory
90+
cl_mem* af_mem = af_A.device<cl_mem>();
91+
cl_A = cl::Buffer(*af_mem, /*retain*/ true);
92+
93+
/// Delete the af_mem pointer. The buffer returned by the device pointer is
94+
/// still valid
95+
delete af_mem;
96+
97+
// 8. Continue your OpenCL application
98+
99+
// ...
100+
return EXIT_SUCCESS;
101+
}
102+
// ![interop_opencl_external_context_snippet]
103+
104+
#pragma GCC diagnostic pop

0 commit comments

Comments
 (0)