Skip to content

Update OpenCL interop page so they discuss deleting of memory #2918

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
9prady9 merged 1 commit into from
Jun 3, 2020
Merged

Update OpenCL interop page so they discuss deleting of memory #2918

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
122 changes: 3 additions & 119 deletions docs/pages/interop_opencl.md
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -64,68 +64,7 @@ synchronization operations.

This process is best illustrated with a fully worked example:

~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~{.cpp}
#include <arrayfire.h>
// 1. Add the af/opencl.h include to your project
#include <af/opencl.h>

int main() {
size_t length = 10;

// Create ArrayFire array objects:
af::array A = af::randu(length, f32);
af::array B = af::constant(0, length, f32);

// ... additional ArrayFire operations here

// 2. Obtain the device, context, and queue used by ArrayFire
static cl_context af_context = afcl::getContext();
static cl_device_id af_device_id = afcl::getDeviceId();
static cl_command_queue af_queue = afcl::getQueue();

// 3. Obtain cl_mem references to af::array objects
cl_mem * d_A = A.device<cl_mem>();
cl_mem * d_B = B.device<cl_mem>();

// 4. Load, build, and use your kernels.
// For the sake of readability, we have omitted error checking.
int status = CL_SUCCESS;

// A simple copy kernel, uses C++11 syntax for multi-line strings.
const char * kernel_name = "copy_kernel";
const char * source = R"(
void __kernel
copy_kernel(__global float * gA, __global float * gB)
{
int id = get_global_id(0);
gB[id] = gA[id];
}
)";

// Create the program, build the executable, and extract the entry point
// for the kernel.
cl_program program = clCreateProgramWithSource(af_context, 1, &source, NULL, &status);
status = clBuildProgram(program, 1, &af_device_id, NULL, NULL, NULL);
cl_kernel kernel = clCreateKernel(program, kernel_name, &status);

// Set arguments and launch your kernels
clSetKernelArg(kernel, 0, sizeof(cl_mem), d_A);
clSetKernelArg(kernel, 1, sizeof(cl_mem), d_B);
clEnqueueNDRangeKernel(af_queue, kernel, 1, NULL, &length, NULL, 0, NULL, NULL);

// 5. Return control of af::array memory to ArrayFire
A.unlock();
B.unlock();

// ... resume ArrayFire operations

// Because the device pointers, d_x and d_y, were returned to ArrayFire's
// control by the unlock function, there is no need to free them using
// clReleaseMemObject()

return 0;
}
~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
\snippet test/interop_opencl_custom_kernel_snippet.cpp interop_opencl_custom_kernel_snippet

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

The eight steps above are best illustrated using a fully-worked example. Below we
use the OpenCL 2.0 C++ API and omit error checking to keep the code readable.

~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~{.cpp}
#include <CL/cl2.hpp>

// 1. Add arrayfire.h and af/opencl.h to your application
#include "arrayfire.h"
#include "af/opencl.h"

#include <cstdio>
#include <vector>

int main() {

// Set up the OpenCL context, device, and queues
cl::Context context(CL_DEVICE_TYPE_ALL);
vector<cl::Device> devices = context.getInfo<CL_CONTEXT_DEVICES>();
cl::Device device = devices[0];
cl::CommandQueue queue(context, device);

// Create a buffer of size 10 filled with ones, copy it to the device
int length = 10;
vector<float> h_A(length, 1);
cl::Buffer cl_A(context, CL_MEM_READ_WRITE, length * sizeof(float), h_A.data());
use the OpenCL C++ API and omit error checking to keep the code readable.

// 2. Instruct OpenCL to complete its operations using clFinish (or similar)
queue.finish();

// 3. Instruct ArrayFire to use the user-created context
// First, create a device from the current OpenCL device + context + queue
afcl::addDevice(device(), context(), queue());
// Next switch ArrayFire to the device using the device and context as
// identifiers:
afcl::setDevice(device(), context());

// 4. Create ArrayFire arrays from OpenCL memory objects
af::array af_A = afcl::array(length, cl_A(), f32, true);

// 5. Perform ArrayFire operations on the Arrays
af_A = af_A + af::randu(length);

// NOTE: ArrayFire does not perform the above transaction using in-place memory,
// thus the underlying OpenCL buffers containing the memory containing memory to
// probably have changed

// 6. Instruct ArrayFire to finish operations using af::sync
af::sync();

// 7. Obtain cl_mem references for important memory
cl_A = *af_A.device<cl_mem>();

// 8. Continue your OpenCL application

// ...

return 0;
}
~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
\snippet test/interop_opencl_external_context_snippet.cpp interop_opencl_external_context_snippet

# Using multiple devices

Expand Down
2 changes: 2 additions & 0 deletions include/af/array.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -725,6 +725,8 @@ namespace af

The device memory returned by this function is not freed until unlock() is called.

/note When using the OpenCL backend and using the cl_mem template argument, the
delete function should be called on the pointer returned by this function.
*/
template<typename T> T* device() const;

Expand Down
19 changes: 17 additions & 2 deletions test/CMakeLists.txt
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -115,7 +115,7 @@ target_compile_definitions(arrayfire_test
# 'BACKENDS' Backends to target for this test. If not set then the test will
# compiled againat all backends
function(make_test)
set(options CXX11 SERIAL USE_MMIO)
set(options CXX11 SERIAL USE_MMIO NO_ARRAYFIRE_TEST)
set(single_args SRC)
set(multi_args LIBRARIES DEFINITIONS BACKENDS)
cmake_parse_arguments(mt_args "${options}" "${single_args}" "${multi_args}" ${ARGN})
Expand All @@ -127,7 +127,12 @@ function(make_test)
continue()
endif()
set(target "test_${src_name}_${backend}")
add_executable(${target} ${mt_args_SRC} $<TARGET_OBJECTS:arrayfire_test>)

if (${mt_args_NO_ARRAYFIRE_TEST})
add_executable(${target} ${mt_args_SRC})
else()
add_executable(${target} ${mt_args_SRC} $<TARGET_OBJECTS:arrayfire_test>)
endif()
target_include_directories(${target}
PRIVATE
${ArrayFire_SOURCE_DIR}/extern/half/include
Expand Down Expand Up @@ -285,6 +290,16 @@ if(OpenCL_FOUND)
LIBRARIES OpenCL::OpenCL
BACKENDS "opencl"
CXX11)
make_test(SRC interop_opencl_custom_kernel_snippet.cpp
LIBRARIES OpenCL::OpenCL
BACKENDS "opencl"
NO_ARRAYFIRE_TEST
CXX11)
make_test(SRC interop_opencl_external_context_snippet.cpp
LIBRARIES OpenCL::OpenCL
BACKENDS "opencl"
NO_ARRAYFIRE_TEST
CXX11)
endif()

if(CUDA_FOUND)
Expand Down
96 changes: 96 additions & 0 deletions test/interop_opencl_custom_kernel_snippet.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,96 @@
/*******************************************************
* Copyright (c) 2020, ArrayFire
* All rights reserved.
*
* This file is distributed under 3-clause BSD license.
* The complete license agreement can be obtained at:
* http://arrayfire.com/licenses/BSD-3-Clause
********************************************************/

// clang-format off
// ![interop_opencl_custom_kernel_snippet]
#include <arrayfire.h>
// 1. Add the af/opencl.h include to your project
#include <af/opencl.h>

#include <cassert>

#define OCL_CHECK(call) \
if (cl_int err = (call) != CL_SUCCESS) { \
fprintf(stderr, __FILE__ "(%d):Returned error code %d\n", __LINE__, \
err); \
}

int main() {
size_t length = 10;

// Create ArrayFire array objects:
af::array A = af::randu(length, f32);
af::array B = af::constant(0, length, f32);

// ... additional ArrayFire operations here

// 2. Obtain the device, context, and queue used by ArrayFire
static cl_context af_context = afcl::getContext();
static cl_device_id af_device_id = afcl::getDeviceId();
static cl_command_queue af_queue = afcl::getQueue();

// 3. Obtain cl_mem references to af::array objects
cl_mem* d_A = A.device<cl_mem>();
cl_mem* d_B = B.device<cl_mem>();

// 4. Load, build, and use your kernels.
// For the sake of readability, we have omitted error checking.
int status = CL_SUCCESS;

// A simple copy kernel, uses C++11 syntax for multi-line strings.
const char* kernel_name = "copy_kernel";
const char* source = R"(
void __kernel
copy_kernel(__global float* gA, __global float* gB) {
int id = get_global_id(0);
gB[id] = gA[id];
}
)";

// Create the program, build the executable, and extract the entry point
// for the kernel.
cl_program program = clCreateProgramWithSource(af_context, 1, &source, NULL, &status);
OCL_CHECK(status);
OCL_CHECK(clBuildProgram(program, 1, &af_device_id, NULL, NULL, NULL));
cl_kernel kernel = clCreateKernel(program, kernel_name, &status);
OCL_CHECK(status);

// Set arguments and launch your kernels
OCL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), d_A));
OCL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_mem), d_B));
OCL_CHECK(clEnqueueNDRangeKernel(af_queue, kernel, 1, NULL, &length, NULL,
0, NULL, NULL));

// 5. Return control of af::array memory to ArrayFire
A.unlock();
B.unlock();

/// A and B should not be the same because of the copy_kernel user code
assert(af::allTrue<bool>(A == B));

// Delete the pointers returned by the device function. This does NOT
// delete the cl_mem memory and only deletes the pointers
delete d_A;
delete d_B;

// ... resume ArrayFire operations

// Because the device pointers, d_x and d_y, were returned to ArrayFire's
// control by the unlock function, there is no need to free them using
// clReleaseMemObject()

// Free the kernel and program objects because they are created in user
// code
OCL_CHECK(clReleaseKernel(kernel));
OCL_CHECK(clReleaseProgram(program));

return 0;
}
// ![interop_opencl_custom_kernel_snippet]
// clang-format on
104 changes: 104 additions & 0 deletions test/interop_opencl_external_context_snippet.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,104 @@
/*******************************************************
* Copyright (c) 2020, ArrayFire
* All rights reserved.
*
* This file is distributed under 3-clause BSD license.
* The complete license agreement can be obtained at:
* http://arrayfire.com/licenses/BSD-3-Clause
********************************************************/

#pragma GCC diagnostic push
#pragma GCC diagnostic ignored "-Wunused-function"
#pragma GCC diagnostic ignored "-Wunused-parameter"
#pragma GCC diagnostic ignored "-Wignored-qualifiers"
#pragma GCC diagnostic ignored "-Wignored-attributes"
#pragma GCC diagnostic ignored "-Wdeprecated-declarations"
#if __GNUC__ >= 8
#pragma GCC diagnostic ignored "-Wcatch-value="
#endif
// ![interop_opencl_external_context_snippet]
#include <arrayfire.h>
// 1. Add the af/opencl.h include to your project
#include <af/opencl.h>

#include <cassert>

// definitions required by cl2.hpp
#define CL_HPP_ENABLE_EXCEPTIONS
#define CL_HPP_TARGET_OPENCL_VERSION 120
#define CL_HPP_MINIMUM_OPENCL_VERSION 120
#include <CL/cl2.hpp>

// 1. Add arrayfire.h and af/opencl.h to your application
#include "af/opencl.h"
#include "arrayfire.h"

#include <cstdio>
#include <vector>

using std::vector;

int main() {
// 1. Set up the OpenCL context, device, and queues
cl::Context context;
try {
context = cl::Context(CL_DEVICE_TYPE_ALL);
} catch (const cl::Error& err) {
fprintf(stderr, "Exiting creating context");
return EXIT_FAILURE;
}
vector<cl::Device> devices = context.getInfo<CL_CONTEXT_DEVICES>();
if (devices.empty()) {
fprintf(stderr, "Exiting. No devices found");
return EXIT_SUCCESS;
}
cl::Device device = devices[0];
cl::CommandQueue queue(context, device);

// Create a buffer of size 10 filled with ones, copy it to the device
int length = 10;
vector<float> h_A(length, 1);
cl::Buffer cl_A(context, CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR,
length * sizeof(float), h_A.data());

// 2. Instruct OpenCL to complete its operations using clFinish (or similar)
queue.finish();

// 3. Instruct ArrayFire to use the user-created context
// First, create a device from the current OpenCL device + context +
// queue
afcl::addDevice(device(), context(), queue());
// Next switch ArrayFire to the device using the device and context as
// identifiers:
afcl::setDevice(device(), context());

// 4. Create ArrayFire arrays from OpenCL memory objects
af::array af_A = afcl::array(length, cl_A(), f32, true);
clRetainMemObject(cl_A());

@9prady9 9prady9 Jun 3, 2020

Copy link
Copy Markdown
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Feel like this unnecessary extra retain.


// 5. Perform ArrayFire operations on the Arrays
af_A = af_A + af::randu(length);

// NOTE: ArrayFire does not perform the above transaction using in-place
// memory, thus the underlying OpenCL buffers containing the memory
// containing memory to probably have changed

// 6. Instruct ArrayFire to finish operations using af::sync
af::sync();

// 7. Obtain cl_mem references for important memory
cl_mem* af_mem = af_A.device<cl_mem>();
cl_A = cl::Buffer(*af_mem, /*retain*/ true);

/// Delete the af_mem pointer. The buffer returned by the device pointer is
/// still valid
delete af_mem;

// 8. Continue your OpenCL application

// ...
return EXIT_SUCCESS;
}
// ![interop_opencl_external_context_snippet]

#pragma GCC diagnostic pop