.. index:: pair: example; gpu_opencl_interop.cpp
.. _doxid-gpu_opencl_interop_8cpp-example:

gpu_opencl_interop.cpp
======================

Annotated version: :ref:`Getting started on GPU with OpenCL extensions API <doxid-gpu_opencl_interop_cpp>`



.. ref-code-block:: cpp

	/*******************************************************************************
	* Copyright 2019 Intel Corporation
	*
	* Licensed under the Apache License, Version 2.0 (the "License");
	* you may not use this file except in compliance with the License.
	* You may obtain a copy of the License at
	*
	*     http://www.apache.org/licenses/LICENSE-2.0
	*
	* Unless required by applicable law or agreed to in writing, software
	* distributed under the License is distributed on an "AS IS" BASIS,
	* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
	* See the License for the specific language governing permissions and
	* limitations under the License.
	*******************************************************************************/
	
	
	
	
	// [Prologue]
	#include <iostream>
	#include <numeric>
	#include <stdexcept>
	
	#include <CL/cl.h>
	
	#include "oneapi/dnnl/dnnl.hpp"
	#include "oneapi/dnnl/dnnl_ocl.hpp"
	
	#include "example_utils.hpp"
	
	using namespace :ref:`dnnl <doxid-namespacednnl>`;
	using namespace :ref:`std <doxid-namespacestd>`;
	// [Prologue]
	
	#define OCL_CHECK(x) \
	    do { \
	        cl_int s = (x); \
	        if (s != CL_SUCCESS) { \
	            std::cout << "[" << __FILE__ << ":" << __LINE__ << "] '" << #x \
	                      << "' failed (status code: " << s << ")." << std::endl; \
	            exit(1); \
	        } \
	    } while (0)
	
	cl_kernel create_init_opencl_kernel(
	        cl_context ocl_ctx, const char *kernel_name, const char *ocl_code) {
	    cl_int err;
	    const char *sources[] = {ocl_code};
	    cl_program ocl_program
	            = clCreateProgramWithSource(ocl_ctx, 1, sources, nullptr, &err);
	    OCL_CHECK(err);
	
	    OCL_CHECK(
	            clBuildProgram(ocl_program, 0, nullptr, nullptr, nullptr, nullptr));
	
	    cl_kernel ocl_kernel = clCreateKernel(ocl_program, kernel_name, &err);
	    OCL_CHECK(err);
	
	    OCL_CHECK(clReleaseProgram(ocl_program));
	    return ocl_kernel;
	}
	
	void gpu_opencl_interop_tutorial() {
	    // [Initialize engine]
	    :ref:`engine <doxid-structdnnl_1_1engine>` eng(validate_engine_kind(:ref:`engine::kind::gpu <doxid-structdnnl_1_1engine_1a2635da16314dcbdb9bd9ea431316bb1aa0aa0be2a866411d9ff03515227454947>`), 0);
	    // [Initialize engine]
	
	    // [Initialize stream]
	    :ref:`dnnl::stream <doxid-structdnnl_1_1stream>` strm(eng);
	    // [Initialize stream]
	
	    //  [memory alloc]
	    :ref:`memory::dims <doxid-structdnnl_1_1memory_1a7d9f4b6ad8caf3969f436cd9ff27e9bb>` tz_dims = {2, 3, 4, 5};
	    const size_t N = std::accumulate(tz_dims.begin(), tz_dims.end(), (size_t)1,
	            std::multiplies<size_t>());
	
	    :ref:`memory::desc <doxid-structdnnl_1_1memory_1_1desc>` mem_d(
	            tz_dims, :ref:`memory::data_type::f32 <doxid-structdnnl_1_1memory_1a8e83474ec3a50e08e37af76c8c075dcea512dc597be7ae761876315165dc8bd2e>`, :ref:`memory::format_tag::nhwc <doxid-structdnnl_1_1memory_1a8e71077ed6a5f7fb7b3e6e1a5a2ecf3fa763cbf7ba1b7b8793dcdc6e2157b5c42>`);
	
	    :ref:`memory <doxid-structdnnl_1_1memory>` mem(mem_d, eng);
	    //  [memory alloc]
	
	    //  [ocl kernel]
	    const char *ocl_code
	            = "__kernel void init(__global float *data) {"
	              "    int id = get_global_id(0);"
	              "    data[id] = (id % 2) ? -id : id;"
	              "}";
	    //  [ocl kernel]
	
	    // [oclkernel create]
	    const char *kernel_name = "init";
	    cl_kernel ocl_init_kernel = create_init_opencl_kernel(
	            :ref:`ocl_interop::get_context <doxid-namespacednnl_1_1ocl__interop_1a248df8106d035e5a7e1ac5fd196c93c3>`(eng), kernel_name, ocl_code);
	    //  [oclkernel create]
	
	    // [oclexecution]
	    cl_mem ocl_buf = :ref:`ocl_interop::get_mem_object <doxid-namespacednnl_1_1ocl__interop_1ac117d62fba9de220fe53b0eedb9671f9>`(mem);
	    OCL_CHECK(clSetKernelArg(ocl_init_kernel, 0, sizeof(ocl_buf), &ocl_buf));
	
	    cl_command_queue ocl_queue = :ref:`ocl_interop::get_command_queue <doxid-namespacednnl_1_1ocl__interop_1a14281f69db5178363ff0c971510d0452>`(strm);
	    OCL_CHECK(clEnqueueNDRangeKernel(ocl_queue, ocl_init_kernel, 1, nullptr, &N,
	            nullptr, 0, nullptr, nullptr));
	    // [oclexecution]
	
	    //  [relu creation]
	    auto relu_pd = :ref:`eltwise_forward::primitive_desc <doxid-structdnnl_1_1eltwise__forward_1_1primitive__desc>`(eng, :ref:`prop_kind::forward <doxid-group__dnnl__api__attributes_1ggac7db48f6583aa9903e54c2a39d65438fa965dbaac085fc891bfbbd4f9d145bbc8>`,
	            :ref:`algorithm::eltwise_relu <doxid-group__dnnl__api__attributes_1gga00377dd4982333e42e8ae1d09a309640aba09bebb742494255b90b43871c01c69>`, mem_d, mem_d, 0.0f);
	    auto relu = :ref:`eltwise_forward <doxid-structdnnl_1_1eltwise__forward>`(relu_pd);
	    //  [relu creation]
	
	    // [relu exec]
	    relu.execute(strm, {{:ref:`DNNL_ARG_SRC <doxid-group__dnnl__api__primitives__common_1gac37ad67b48edeb9e742af0e50b70fe09>`, mem}, {:ref:`DNNL_ARG_DST <doxid-group__dnnl__api__primitives__common_1ga3ca217e4a06d42a0ede3c018383c388f>`, mem}});
	    strm.:ref:`wait <doxid-structdnnl_1_1stream_1a59985fa8746436057cf51a820ef8929c>`();
	    // [relu exec]
	
	    // [Check the results]
	    std::vector<float> mem_data(N);
	    read_from_dnnl_memory(mem_data.data(), mem);
	    for (size_t i = 0; i < N; i++) {
	        float expected = (i % 2) ? 0.0f : (float)i;
	        if (mem_data[i] != expected) {
	            std::cout << "Expect " << expected << " but got " << mem_data[i]
	                      << "." << std::endl;
	            throw std::logic_error("Accuracy check failed.");
	        }
	    }
	    // [Check the results]
	
	    OCL_CHECK(clReleaseKernel(ocl_init_kernel));
	}
	
	int main(int argc, char **argv) {
	    return handle_example_errors(
	            {:ref:`engine::kind::gpu <doxid-structdnnl_1_1engine_1a2635da16314dcbdb9bd9ea431316bb1aa0aa0be2a866411d9ff03515227454947>`}, gpu_opencl_interop_tutorial);
	}