adding sample for CSDK

cmfe/examples/Gen/sample/CMakeLists.txt

@@ -0,0 +1,87 @@
+cmake_minimum_required(VERSION 3.10)
+
+set(CMAKE_CXX_STANDARD 17)
+set(CMAKE_CXX_STANDARD_REQUIRED True)
+
+project(sample)
+
+if(NOT DEFINED ENV{CSDK_IGC})
+  message(FATAL_ERROR "CSDK_IGC environment variable is not set - did you run 'setenv.bat' ?")
+endif()
+
+# set paths
+set(CSDK_IGC $ENV{CSDK_IGC})
+set(CMEMU_PATH ${CSDK_IGC}/cmemu)
+
+# locate all relevant packages
+foreach(program CMOC OCLOC)
+  string(TOLOWER ${program} binary)
+  find_program(${program} ${binary} REQUIRED)
+  if(NOT ${program})
+    message(FATAL_ERROR "Unable to locate ${binary} executable - did you run 'setenv.bat' ?")
+  else()
+    message(INFO " using ${binary} from ${${program}}")
+  endif()
+endforeach(program)
+find_library(LIB_OPENCL NAMES Intel_OpenCL_ICD64 PATHS ${CSDK_IGC}/runtime/opencl/lib)
+find_library(LIB_LEVEL0 NAMES ze_loader PATHS ${CSDK_IGC}/runtime/level_zero/lib)
+
+# our sources
+set(KERNEL ${CMAKE_SOURCE_DIR}/kernel.cpp)
+set(HOST_OCL ${CMAKE_SOURCE_DIR}/host.cpp)
+set(HOST_L0  ${CMAKE_SOURCE_DIR}/host_l0.cpp)
+
+# os-specific
+if (CMAKE_HOST_SYSTEM_NAME MATCHES Windows)
+  set(dll ${CMAKE_SHARED_LIBRARY_SUFFIX})
+else()
+  set(dll)
+endif()
+set(INSTALL_DIR ${CMAKE_BINARY_DIR}/bin)
+
+
+
+
+#######
+# GPU
+set(out ${CMAKE_BINARY_DIR}/kernel.spv.skl ${CMAKE_BINARY_DIR}/kernel.skl)
+list(GET out 0 spirv)
+list(GET out 1 binary)
+string(REPLACE ${CMAKE_BINARY_DIR}/ "" kernel ${binary})
+add_custom_command(OUTPUT ${out}
+  COMMAND ${CMOC} -emit-spirv -fcmocl -mcpu=SKL -m64 ${KERNEL} -o ${spirv}
+  COMMAND ${OCLOC} -device skl -output_no_suffix -options "-cmc" -spirv_input -file ${spirv} -output ${binary}
+)
+add_custom_target(kernel_gpu DEPENDS ${out})
+install(FILES ${out} DESTINATION ${INSTALL_DIR})
+
+add_executable(vector.skl ${HOST_OCL})
+set_target_properties(vector.skl PROPERTIES COMPILE_FLAGS -DKERNEL=\\\"${kernel}\\\")
+target_include_directories(vector.skl PUBLIC ${CSDK_IGC}/runtime/opencl/include)
+target_link_libraries(vector.skl ${LIB_OPENCL})
+add_dependencies(vector.skl kernel_gpu)
+install(TARGETS vector.skl DESTINATION ${INSTALL_DIR})
+
+
+
+
+
+
+
+#######
+# GPU / L0
+set(kernel kernel.spv.skl)
+add_executable(vector.l0.skl ${HOST_L0})
+set_target_properties(vector.l0.skl PROPERTIES COMPILE_FLAGS -DKERNEL=\\\"${kernel}\\\")
+target_include_directories(vector.l0.skl PUBLIC ${CSDK_IGC}/runtime/level_zero/include)
+target_link_libraries(vector.l0.skl ${LIB_LEVEL0})
+add_dependencies(vector.l0.skl kernel_gpu)
+install(TARGETS vector.l0.skl DESTINATION ${INSTALL_DIR})
+
+
+
+
+
+# all targets to build
+add_custom_target(build ALL)
+add_dependencies(build vector.skl vector.l0.skl)

cmfe/examples/Gen/sample/host.cpp

@@ -0,0 +1,143 @@
+/*===================== begin_copyright_notice ==================================
+
+ Copyright (c) 2020, 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.
+======================= end_copyright_notice ==================================*/
+
+#include <iostream>
+#include <cassert>
+#include <math.h>
+#include <vector>
+
+#include <CL/cl.h>
+
+#define SZ 160
+#define KERNEL_SZ 16
+#define CHECK(a) do { \
+    err = (a); \
+    if (err != CL_SUCCESS) { \
+        fprintf(stderr, "FAIL: err=%d @ line=%d (%s)\n", err, __LINE__, (#a)); \
+        exit(err); \
+    } \
+}while (0)
+#define CHECK2(a) do { \
+    (a); \
+    if (err != CL_SUCCESS) { \
+        fprintf(stderr, "FAIL: err=%d @ line=%d (%s)\n", err, __LINE__, (#a)); \
+        exit(err); \
+    } \
+}while (0)
+#ifndef KERNEL
+#error "Error: KERNEL must be defined with location of kernel binary"
+#endif
+
+int main( int argc, char* argv[])
+{
+    // initialize data
+    int *src1 = (int *)malloc(sizeof(int)*SZ);
+    int *src2 = (int *)malloc(sizeof(int)*SZ);
+    int *dst  = (int *)malloc(sizeof(int)*SZ);
+
+    for (unsigned i=0; i<SZ; i++)
+    {
+        src1[i] = i;
+        src2[i] = i<<2;
+    }
+
+    // initialize GPU
+    cl_platform_id platform;  // OpenCL platform
+    cl_device_id device;      // device ID
+    cl_context context;       // context
+    cl_command_queue queue;   // command queue
+    cl_program program;       // program
+    cl_kernel kernel;         // kernel
+    cl_int err;
+
+    CHECK(clGetPlatformIDs(1, &platform, NULL));
+    CHECK(clGetDeviceIDs(platform, CL_DEVICE_TYPE_GPU, 1, &device, NULL));
+    CHECK2(context = clCreateContext(NULL, 1, &device, NULL, NULL, &err));
+    CHECK2(queue = clCreateCommandQueueWithProperties(context, device, 0, &err));
+
+    // diagnostic info
+    char name_buffer[256];
+    CHECK(clGetPlatformInfo(platform, CL_PLATFORM_NAME, sizeof(name_buffer), name_buffer, NULL));
+    fprintf(stderr, "INFO: using platform: %s\n", name_buffer);
+    CHECK(clGetDeviceInfo(device, CL_DEVICE_NAME, sizeof(name_buffer), name_buffer, NULL));
+    fprintf(stderr, "INFO: using device: %s\n", name_buffer);
+
+    // read in and initialize kernel
+    FILE *fp = fopen(KERNEL, "rb");
+    if (fp == NULL) {
+        fprintf(stderr, "FAIL: unable to open %s\n", KERNEL);
+        exit(-1);
+    }
+    fseek(fp, 0, SEEK_END);
+    size_t sz = ftell(fp);
+    rewind(fp);
+
+    unsigned char *code = (unsigned char *)malloc(sz);
+    fread(code, 1, sz, fp);
+    fclose(fp);
+
+    cl_int errNum = 0;
+    const unsigned char *codes[1] = {code};
+    size_t sizes[1] = {sz};
+    CHECK2(program = clCreateProgramWithBinary(context, 1, &device, sizes, codes, &err, &errNum));
+    CHECK(clBuildProgram(program, 0, NULL, NULL, NULL, NULL));
+    CHECK2(kernel = clCreateKernel(program, "vector_add", &err));
+
+    // kernel parameter initialization
+    cl_mem d_a, d_b, d_c;
+    size_t bytes = SZ*sizeof(int);
+    d_a = clCreateBuffer(context, CL_MEM_READ_ONLY, bytes, NULL, NULL);
+    d_b = clCreateBuffer(context, CL_MEM_READ_ONLY, bytes, NULL, NULL);
+    d_c = clCreateBuffer(context, CL_MEM_WRITE_ONLY, bytes, NULL, NULL);
+    CHECK(clEnqueueWriteBuffer(queue, d_a, CL_TRUE, 0, bytes, src1, 0, NULL, NULL));
+    CHECK(clEnqueueWriteBuffer(queue, d_b, CL_TRUE, 0, bytes, src2, 0, NULL, NULL));
+    CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), &d_a));
+    CHECK(clSetKernelArg(kernel, 1, sizeof(cl_mem), &d_b));
+    CHECK(clSetKernelArg(kernel, 2, sizeof(cl_mem), &d_c));
+
+    // send to GPU
+    size_t globalSize = SZ/KERNEL_SZ;
+    size_t localSize = 1;
+    CHECK(clEnqueueNDRangeKernel(queue, kernel, 1, NULL, &globalSize, &localSize, 0, NULL, NULL));
+    clFinish(queue);
+
+    // process output and cleanup
+    clEnqueueReadBuffer(queue, d_c, CL_TRUE, 0, bytes, dst, 0, NULL, NULL );
+    clReleaseMemObject(d_a);
+    clReleaseMemObject(d_b);
+    clReleaseMemObject(d_c);
+    clReleaseProgram(program);
+    clReleaseKernel(kernel);
+    clReleaseCommandQueue(queue);
+    clReleaseContext(context);
+
+    // verify results
+    for (unsigned i=0; i<SZ; i++)
+       if ((src1[i] + src2[i]) != dst[i]) {
+           fprintf(stderr, "FAIL: comparison at index[%d]: %d + %d => %d(host), but %d(gpu)\n", i, src1[i], src2[i], (src1[i]+src2[i]), dst[i]);
+           //exit(-1);
+       }
+    fprintf(stderr, "PASSED\n");
+    return 0;
+}