[SYCL] Restore SYCL_USE_KERNEL_SPV env var handling to load SPIRV program from a file.

Signed-off-by: Vladimir Lazarev <vladimir.lazarev@intel.com>
Signed-off-by: Konstantin S Bobrovsky <konstantin.s.bobrovsky@intel.com>

Author: kbobrovs

sycl/source/detail/program_manager/program_manager.cpp

@@ -252,7 +252,7 @@ cnri_program ProgramManager::loadProgram(OSModuleHandle M,
     Img->Kind = SYCL_OFFLOAD_KIND;
     Img->Format = CNRI_IMG_NONE;
     Img->DeviceTargetSpec = CNRI_TGT_STR_UNKNOWN;
-    Img->BuildOptions = nullptr;
+    Img->BuildOptions = "";
     Img->ManifestStart = nullptr;
     Img->ManifestEnd = nullptr;
     Img->ImageStart = Data;
@@ -291,13 +291,13 @@ cnri_program ProgramManager::loadProgram(OSModuleHandle M,
   assert(Img->ImageEnd >= Img->ImageStart);
   size_t ImgSize = static_cast<size_t>(Img->ImageEnd - Img->ImageStart);
 
-  // Determine the kind of the image if not set already
-  if (Img->Kind == CNRI_IMG_NONE) {
+  // Determine the format of the image if not set already
+  if (Img->Format == CNRI_IMG_NONE) {
     struct {
       cnri_device_image_format Fmt;
-      const int32_t Magic;
+      const uint32_t Magic;
     } Fmts[] = {{CNRI_IMG_SPIRV, 0x07230203},
-                {CNRI_IMG_LLVMIR_BITCODE, 0x4243C0DE}};
+                {CNRI_IMG_LLVMIR_BITCODE, 0xDEC04342}};
     if (ImgSize >= sizeof(Fmts[0].Magic)) {
       std::remove_const<decltype(Fmts[0].Magic)>::type Hdr = 0;
       std::copy(Img->ImageStart, Img->ImageStart + sizeof(Hdr),
@@ -308,7 +308,8 @@ cnri_program ProgramManager::loadProgram(OSModuleHandle M,
           Img->Format = Fmt.Fmt;
 
           if (DbgProgMgr > 1) {
-            std::cerr << "determined image format: " << Img->Format;
+            std::cerr << "determined image format: " << (int)Img->Format
+              << "\n";
           }
           break;
         }
@@ -321,9 +322,9 @@ cnri_program ProgramManager::loadProgram(OSModuleHandle M,
     Fname += Img->DeviceTargetSpec;
     std::string Ext;
 
-    if (Img->Kind == CNRI_IMG_SPIRV) {
+    if (Img->Format == CNRI_IMG_SPIRV) {
       Ext = ".spv";
-    } else if (Img->Kind == CNRI_IMG_LLVMIR_BITCODE) {
+    } else if (Img->Format == CNRI_IMG_LLVMIR_BITCODE) {
       Ext = ".bc";
     } else {
       Ext = ".bin";

sycl/test/kernel_from_file/hw.cpp

@@ -0,0 +1,46 @@
+// RUN: %clang -std=c++11 --sycl -fno-sycl-use-bitcode -Xclang -fsycl-int-header=%t.h -c %s -o %t.spv
+// RUN: %clang -std=c++11 -include %t.h -g %s -o %t.out -lOpenCL -lsycl -lstdc++
+// RUN: env SYCL_USE_KERNEL_SPV=%t.spv %t.out | FileCheck %s
+// CHECK: Passed
+
+
+#include <CL/sycl.hpp>
+#include <iostream>
+
+using namespace cl::sycl;
+
+int main(int argc, char **argv) {
+  int data = 5;
+
+  try {
+    queue myQueue;
+    buffer<int, 1> buf(&data, range<1>(1));
+
+    event e = myQueue.submit([&](handler& cgh) {
+      auto ptr = buf.get_access<access::mode::read_write>(cgh);
+
+      cgh.single_task<class my_kernel>([=]() {
+        ptr[0]++;
+      });
+    });
+    e.wait_and_throw();
+
+  } catch (cl::sycl::exception const& e) {
+    std::cerr << "SYCL exception caught:\n";
+    std::cerr << e.what() << "\n";
+    return 2;
+  }
+  catch (...) {
+    std::cerr << "unknown exception caught\n";
+    return 1;
+  }
+
+  if (data == 6) {
+    std::cout << "Passed\n";
+    return 0;
+  } else {
+    std::cout << "Failed: " << data << "!= 6(gold)\n";
+    return 1;
+  }
+}
+