DPC++/OpenCL interoperability sample (oneapi-src#500)

* Add License

Signed-off-by: Karl Qi <karl.qi@intel.com>

* OpenCL DPC++ Interoperability Samples

* Update Readme.md

Spelling corrections

* Update README

Signed-off-by: Karl Qi <karl.qi@intel.com>

* README.md

* Typo

* Change to use cmake rather than make, and add more to README

Signed-off-by: Karl Qi <karl.qi@intel.com>

* Add DevCloud and Output Details

Signed-off-by: Karl Qi <karl.qi@intel.com>

* Change DevCloud URL

Signed-off-by: Karl Qi <karl.qi@intel.com>

* Use constexpr instead of #define and static const

Signed-off-by: Karl Qi <karl.qi@intel.com>

* Remove Windows as Supported Platform

Signed-off-by: Karl Qi <karl.qi@intel.com>

* Using SYCL 2020 style buffer

Signed-off-by: Karl Qi <karl.qi@intel.com>

* misc grammar and punctuation updates

Co-authored-by: JoeOster <52936608+JoeOster@users.noreply.github.com>

Author: karlqi2000

DirectProgramming/DPC++/OpenCLInterop/CMakeLists.txt

@@ -0,0 +1,13 @@
+cmake_minimum_required (VERSION 3.0)
+
+set(CMAKE_CXX_COMPILER dpcpp)
+
+# Set default build type to RelWithDebInfo if not specified
+if (NOT CMAKE_BUILD_TYPE)
+    message (STATUS "Default CMAKE_BUILD_TYPE not set using Release with Debug Info")
+    set (CMAKE_BUILD_TYPE "RelWithDebInfo" CACHE
+        STRING "Choose the type of build, options are: None Debug Release RelWithDebInfo MinSizeRel"
+        FORCE)
+endif()
+project (opencl_interop)
+add_subdirectory (src)

DirectProgramming/DPC++/OpenCLInterop/License.txt

@@ -0,0 +1,8 @@
+Copyright Intel Corporation
+
+Permission is hereby granted, free of charge, to any person obtaining a copy of this software and associated documentation files (the "Software"), to deal in the Software without restriction, including without limitation the rights to use, copy, modify, merge, publish, distribute, sublicense, and/or sell copies of the Software, and to permit persons to whom the Software is furnished to do so, subject to the following conditions:
+
+The above copyright notice and this permission notice shall be included in all copies or substantial portions of the Software.
+
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE.
+

DirectProgramming/DPC++/OpenCLInterop/README.md

@@ -0,0 +1,77 @@
+# DPC++ OpenCL™ Interoperability Example
+
+This examples demonstrate how DPC++ can interact with OpenCL™. This code shample will show  programmers to incrementally migrate from
+OpenCL to DPC++. Two usage scenarios are shown. First is a DPC++ program that compiles and runs an OpenCL kernel. The second program converts OpenCL objects to DPC++.
+
+For more information on migrating from OpenCL to DPC++, see [Migrating OpenCL Designs to DPC++](https://software.intel.com/content/www/us/en/develop/articles/migrating-opencl-designs-to-dpcpp.html).
+
+| Optimized for                       | Description
+|:---                               |:---
+| OS                                | Linux* Ubuntu* 18.04, 20
+| Hardware                          | Skylake or newer
+| Software                          | Intel® oneAPI DPC++/C++ Compiler, Intel Devcloud
+| What you will learn               | How OpenCL code can interact with DPC++ with the Intel® oneAPI DPC++/C++ Compiler
+| Time to complete                  | 10 minutes
+
+## Purpose
+For users migrating from OpenCL to DPC++, interoperability allows the migration to take place piecemeal so that the migration of all kernels does not have to occur simultaneously.
+ 
+## Key Implementation Details
+The common OpenCL to DPC++ conversion scenarios are covered.
+1. In dpcpp_with_opencl_kernel.dp.cpp, the DPC++ program compiles and runs an OpenCL kernel. (For this, OpenCL must be set as the backend and not Level 0, the environment variable SYCL_DEVICE_FILTER=OPENCL is used)
+2. In dpcpp_with_opencl_objects.dp.cpp, the program converts OpenCL objects (Memory Objects, Platform, Context, Program, Kernel) to DPC++ and execute the program. 
+
+## License  
+Code samples are licensed under the MIT license. See
+[License.txt](https://github.com/oneapi-src/oneAPI-samples/blob/master/License.txt) for details.
+
+Third party program Licenses can be found here: [third-party-programs.txt](https://github.com/oneapi-src/oneAPI-samples/blob/master/third-party-programs.txt)
+
+## Building the Program
+
+> Note: if you have not already done so, set up your CLI 
+> environment by sourcing  the setvars script located in 
+> the root of your oneAPI installation. 
+>
+> Linux Sudo: . /opt/intel/oneapi/setvars.sh  
+> Linux User: . ~/intel/oneapi/setvars.sh  
+> Windows: C:\Program Files(x86)\Intel\oneAPI\setvars.bat
+
+### Running Samples In DevCloud
+If running a sample in the Intel DevCloud, remember that you must specify the compute node (CPU, GPU, FPGA) and whether to run in batch or interactive mode. For more information, see the Intel® oneAPI Base Toolkit Get Started Guide (https://devcloud.intel.com/oneapi/get_started/baseToolkitSamples/)
+
+### On a Linux* System
+Perform the following steps:
+1. Build the program
+	```
+    $ mkdir build
+    $ cd build
+    $ cmake ..
+	$ make
+	```
+
+2. Run the program:
+    ```
+    make run_prog1
+    make run_prog2
+    ```
+
+3. Clean the program using:
+    ```
+    make clean
+    ```
+
+### Example of Output
+```
+Device: Intel(R) HD Graphics 630 [0x5912]
+PASSED!
+Built target run_prog1
+
+Kernel Loading Done
+Platforms Found: 3
+Using Platform: Intel(R) FPGA Emulation Platform for OpenCL(TM)
+Devices Found: 1
+Device: Intel(R) FPGA Emulation Device
+Passed!
+Built target run_prog2
+```

DirectProgramming/DPC++/OpenCLInterop/bin/.gitkeep

@@ -0,0 +1,27 @@
+{
+  "guid": "E3DE34BF-E5B6-44AF-8722-B2A5A6BE8D57",
+  "name": "DPC++ OpenCL Interoperability Samples",
+  "categories": [ "Toolkit/DirectProgramming/DPC++/opencl_interop/" ],
+  "description": "Samples showing DPC++ OpenCL Interoperability",
+  "toolchain": [ "dpcpp" ],
+  "languages": [ { "cpp": { "properties": { "projectOptions": [ { "projectType": "makefile" } ] } } } ],
+  "os": [ "linux" ],
+  "builder": [ "ide", "make" ],
+  "targetDevice": [ "CPU", "GPU" ],
+  "ciTests": {
+    "linux": [
+      {
+        "env": [ "source /opt/intel/oneapi/setvars.sh" ],
+        "steps": [
+          "mkdir build",
+          "cd build",
+          "cmake ..",
+          "make",
+          "make run_prog1",
+          "make run_prog2",
+          "make clean"
+        ]
+      }
+    ]
+  }
+}

DirectProgramming/DPC++/OpenCLInterop/sample.json

@@ -0,0 +1,11 @@
+set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -g -std=c++17")
+set(CMAKE_EXE_LINKER_FLAGS "${CMAKE_EXE_LINKER_FLAGS}")
+
+add_executable(prog1 dpcpp_with_opencl_kernel.dp.cpp)
+target_link_libraries(prog1 OpenCL sycl)
+add_custom_target(run_prog1 ./prog1)
+
+file(COPY ../src/vector_add_kernel.cl DESTINATION .)
+add_executable(prog2 dpcpp_with_opencl_objects.dp.cpp)
+target_link_libraries(prog2 OpenCL sycl)
+add_custom_target(run_prog2 ./prog2)

DirectProgramming/DPC++/OpenCLInterop/src/CMakeLists.txt

@@ -0,0 +1,72 @@
+//==============================================================
+// Copyright © 2021 Intel Corporation
+//
+// SPDX-License-Identifier: MIT
+// =============================================================
+
+#include <stdlib.h>
+#include <CL/sycl.hpp>
+using namespace sycl;
+
+constexpr int N = 1024;
+
+void dpcpp_code(int *a, int *b, int *c) {
+  queue q{default_selector()};  // Create Command Queue Targeting GPU
+  std::cout << "Device: " << q.get_device().get_info<info::device::name>()
+            << std::endl;
+
+  program p(q.get_context());  // Create program from the same context as q
+
+  // Compile OpenCL vecAdd kernel. which is expressed as a C++ Raw String
+  // as indicated by R”
+  p.build_with_source(R"( __kernel void vecAdd(__global int *a, 
+                                                __global int *b, 
+                                                __global int *c) 
+                            {
+                                int i=get_global_id(0);
+                                c[i] = a[i] + b[i]; 
+                            } )");
+
+  buffer buf_a(a, range(N));
+  buffer buf_b(b, range(N));
+  buffer buf_c(c, range(N));
+
+  q.submit([&](handler &h) {
+    accessor A(buf_a, h, read_only);
+    accessor B(buf_b, h, read_only);
+    accessor C(buf_c, h, write_only);
+    // Set buffers as arguments to the kernel
+    h.set_args(A, B, C);
+    // Launch vecAdd kernel from the p program object across N elements.
+    h.parallel_for(range(N), p.get_kernel("vecAdd"));
+  });
+}
+
+int main(int argc, char **argv) {
+  // Ensure to use OpenCL backend for OpenCL Kernel Compilation
+  putenv((char *)"SYCL_DEVICE_FILTER=OPENCL");
+
+  size_t bytes = sizeof(int) * N;
+
+  int *a = (int *)malloc(bytes);
+  int *b = (int *)malloc(bytes);
+  int *c = (int *)malloc(bytes);
+  for (int i = 0; i < N; ++i) {
+    a[i] = i;
+    b[i] = i * 2;
+  }
+  dpcpp_code(a, b, c);
+
+  for (int i = 0; i < N; ++i) {
+    if (c[i] != i * 3) {
+      std::cout << "FAILED!" << std::endl;
+      return -1;
+    }
+  }
+  std::cout << "PASSED!" << std::endl;
+
+  free(a);
+  free(b);
+  free(c);
+  return 0;
+}

DirectProgramming/DPC++/OpenCLInterop/src/dpcpp_with_opencl_kernel.dp.cpp

@@ -0,0 +1,127 @@
+//==============================================================
+// Copyright © 2021 Intel Corporation
+//
+// SPDX-License-Identifier: MIT
+// =============================================================
+
+#include <CL/opencl.h>
+#include <stdio.h>
+#include <CL/sycl.hpp>
+using namespace sycl;
+
+constexpr int MAX_SOURCE_SIZE = 0x100000;
+constexpr int N = 1024;
+
+int main(int argc, char **argv) {
+  size_t bytes = sizeof(float) * N;
+
+  cl_float *host_a = (cl_float *)malloc(bytes);
+  cl_float *host_b = (cl_float *)malloc(bytes);
+  cl_float *host_c = (cl_float *)malloc(bytes);
+  for (int i = 0; i < N; ++i) {
+    host_a[i] = i;
+    host_b[i] = i * 2;
+  }
+
+  FILE *fp;
+  char *source_str;
+  size_t source_size;
+  fp = fopen("vector_add_kernel.cl", "r");
+  if (!fp) {
+    std::cerr << "Failed to load kernel file." << std::endl;
+  }
+  source_str = (char *)malloc(MAX_SOURCE_SIZE);
+  source_size = fread(source_str, 1, MAX_SOURCE_SIZE, fp);
+  fclose(fp);
+  std::cout << "Kernel Loading Done" << std::endl;
+
+  // Get platform and device information
+  cl_device_id device_id = NULL;
+  cl_uint ret_num_devices;
+  cl_uint ret_num_platforms;
+  cl_int ret = clGetPlatformIDs(0, NULL, &ret_num_platforms);
+  std::cout << "Platforms Found: " << ret_num_platforms << std::endl;
+
+  cl_platform_id *ocl_platforms = (cl_platform_id *)malloc(ret_num_platforms);
+  ret = clGetPlatformIDs(ret_num_platforms, ocl_platforms, NULL);
+  // Set Platform to Use
+  int platform_index = 0;
+  platform sycl_platform(ocl_platforms[platform_index]);
+  std::cout << "Using Platform: "
+            << sycl_platform.get_info<info::platform::name>() << std::endl;
+
+  ret = clGetDeviceIDs(ocl_platforms[platform_index], CL_DEVICE_TYPE_ALL, 1,
+                       &device_id, &ret_num_devices);
+  std::cout << "Devices Found: " << ret_num_devices << std::endl;
+
+  // Create an OpenCL context and queue
+  cl_context ocl_context =
+      clCreateContext(NULL, 1, &device_id, NULL, NULL, &ret);
+  cl_command_queue ocl_queue =
+      clCreateCommandQueueWithProperties(ocl_context, device_id, 0, &ret);
+
+  // Create a program from the kernel source, and build it
+  cl_program ocl_program =
+      clCreateProgramWithSource(ocl_context, 1, (const char **)&source_str,
+                                (const size_t *)&source_size, &ret);
+  ret = clBuildProgram(ocl_program, 1, &device_id, NULL, NULL, NULL);
+
+  // OpenCL Kernel and Memory Objects
+  cl_kernel ocl_kernel = clCreateKernel(ocl_program, "vector_add", &ret);
+  cl_mem ocl_buf_a =
+      clCreateBuffer(ocl_context, CL_MEM_READ_ONLY, bytes, NULL, NULL);
+  cl_mem ocl_buf_b =
+      clCreateBuffer(ocl_context, CL_MEM_READ_ONLY, bytes, NULL, NULL);
+  cl_mem ocl_buf_c =
+      clCreateBuffer(ocl_context, CL_MEM_READ_ONLY, bytes, NULL, NULL);
+  clEnqueueWriteBuffer(ocl_queue, ocl_buf_a, CL_TRUE, 0, bytes, host_a, 0, NULL,
+                       NULL);
+  clEnqueueWriteBuffer(ocl_queue, ocl_buf_b, CL_TRUE, 0, bytes, host_b, 0, NULL,
+                       NULL);
+
+  {  // DPC++ Application Scope
+    // Construct SYCL versions of the context, queue, kernel, and buffers
+    context sycl_context(ocl_context);
+    queue sycl_queue(ocl_queue, sycl_context);
+    std::cout << "Device: "
+              << sycl_queue.get_device().get_info<info::device::name>()
+              << std::endl;
+    kernel sycl_kernel(ocl_kernel, sycl_context);
+    buffer<int, 1> sycl_buf_a(ocl_buf_a, sycl_context);
+    buffer<int, 1> sycl_buf_b(ocl_buf_b, sycl_context);
+    buffer<int, 1> sycl_buf_c(ocl_buf_c, sycl_context);
+    sycl_queue.submit([&](handler &h) {
+      // Create accessors for each of the buffers
+      accessor a_accessor(sycl_buf_a, h, read_only);
+      accessor b_accessor(sycl_buf_b, h, read_only);
+      accessor c_accessor(sycl_buf_c, h, write_only);
+      // Map kernel arguments to accessors
+      h.set_args(a_accessor, b_accessor, c_accessor);
+      // Launch Kernel
+      h.parallel_for(range<1>(N), sycl_kernel);
+    });
+  }
+  // Read buffer content back to host array
+  clEnqueueReadBuffer(ocl_queue, ocl_buf_c, CL_TRUE, 0, bytes, host_c, 0, NULL,
+                      NULL);
+
+  for (int i = 0; i < N; ++i) {
+    if (host_c[i] != i * 3) {
+      std::cout << "Failed!" << std::endl;
+      return -1;
+    }
+  }
+  std::cout << "Passed!" << std::endl;
+
+  ret = clReleaseCommandQueue(ocl_queue);
+  ret = clReleaseKernel(ocl_kernel);
+  ret = clReleaseProgram(ocl_program);
+  ret = clReleaseMemObject(ocl_buf_a);
+  ret = clReleaseMemObject(ocl_buf_b);
+  ret = clReleaseMemObject(ocl_buf_c);
+  ret = clReleaseContext(ocl_context);
+  free(host_a);
+  free(host_b);
+  free(host_c);
+  return 0;
+}

DirectProgramming/DPC++/OpenCLInterop/src/dpcpp_with_opencl_objects.dp.cpp

@@ -0,0 +1,13 @@
+//==============================================================
+// Copyright © 2021 Intel Corporation
+//
+// SPDX-License-Identifier: MIT
+// =============================================================
+
+__kernel void vector_add(__global const float *x, __global const float *y,
+                         __global float *restrict z) {
+  // get index of the work item
+  int index = get_global_id(0);
+  // add the vector elements
+  z[index] = x[index] + y[index];
+}