add OpenCL kernels as string literals

Author: huanzhang12

src/treelearner/gpu_tree_learner.cpp

@@ -295,12 +295,12 @@ void GPUTreeLearner::InitGPU(int platform_id, int device_id) {
   std::string kernel_source;
   std::string kernel_name;
   if (max_num_bin_ <= 64) {
-    kernel_source = "histogram64.cl";
+    kernel_source = kernel64_src_;
     kernel_name = "histogram64";
     device_bin_size_ = 64;
   }
   else if ( max_num_bin_ <= 256) {
-    kernel_source = "histogram256.cl";
+    kernel_source = kernel256_src_;
     kernel_name = "histogram256";
     device_bin_size_ = 256;
   }
@@ -332,9 +332,10 @@ void GPUTreeLearner::InitGPU(int platform_id, int device_id) {
   device_histogram_outputs_ = boost::compute::buffer(ctx_, num_dense_feature4_ * 4 * device_bin_size_ * sizeof(GPUHistogramBinEntry), 
                            boost::compute::memory_object::write_only, nullptr);
   Log::Info("Using GPU Device: %s, Vendor: %s", dev_.name().c_str(), dev_.vendor().c_str());
-  Log::Info("Compiling OpenCL Kernel from %s...", kernel_source.c_str());
+  Log::Info("Compiling OpenCL Kernel with %d bins...", device_bin_size_);
+  // Log::Info("Compiling OpenCL Kernel:\n%s", kernel_source.c_str());
   for (int i = 0; i <= max_exp_workgroups_per_feature_; ++i) {
-    auto program = boost::compute::program::create_with_source_file(kernel_source, ctx_);
+    auto program = boost::compute::program::create_with_source(kernel_source, ctx_);
     std::ostringstream opts;
     // FIXME: sparse data
     opts << "-D FEATURE_SIZE=" << num_data_ << " -D POWER_FEATURE_WORKGROUPS=" << i
@@ -358,7 +359,7 @@ void GPUTreeLearner::InitGPU(int platform_id, int device_id) {
   }
   // create the OpenCL kernel for the root node (all data)
   int full_exp_workgroups_per_feature = GetNumWorkgroupsPerFeature(num_data_);
-  auto program = boost::compute::program::create_with_source_file(kernel_source, ctx_);
+  auto program = boost::compute::program::create_with_source(kernel_source, ctx_);
   std::ostringstream opts;
   // FIXME: sparse data
   opts << "-D FEATURE_SIZE=" << num_data_ << " -D POWER_FEATURE_WORKGROUPS=" << full_exp_workgroups_per_feature

src/treelearner/gpu_tree_learner.h

@@ -195,6 +195,16 @@ class GPUTreeLearner: public TreeLearner {
   boost::compute::device dev_;
   boost::compute::context ctx_;
   boost::compute::command_queue queue_;
+  /*! \brief GPU kernel for 256 bins */
+  const char *kernel256_src_ = 
+  #include "ocl/histogram256.cl"
+  ;
+  /*! \brief GPU kernel for 64 bins */
+  const char *kernel64_src_ = 
+  #include "ocl/histogram64.cl"
+  ;
+  /*! \brief GPU kernel for 64 bins */
+
   /*! \brief a array of histogram kernels with different number
      of workgroups per feature */
   std::vector<boost::compute::kernel> histogram_kernels_;

histogram256.cl src/treelearner/ocl/histogram256.cl

@@ -1,3 +1,18 @@
+// this file can either be read and passed to an OpenCL compiler directly,
+// or included in a C++11 source file as a string literal
+#ifndef __OPENCL_VERSION__
+// If we are including this file in C++,
+// the entire source file following (except the last #endif) will become
+// a raw string literal. The extra ")" is just for mathcing parentheses
+// to make the editor happy. The extra ")" and extra endif will be skipped.
+// DO NOT add anything between here and the next #ifdef, otherwise you need
+// to modify the skip count at the end of this file.
+R""()
+#endif
+
+#ifndef _HISTOGRAM_256_KERNEL_
+#define _HISTOGRAM_256_KERNEL_
+
 #pragma OPENCL EXTENSION cl_khr_local_int32_base_atomics : enable
 #pragma OPENCL EXTENSION cl_khr_local_int32_extended_atomics : enable
 
@@ -551,3 +566,8 @@ __kernel void histogram256(__global const uchar4* feature_data_base,
     }
 }
 
+// The following line ends the string literal, adds an extra #endif at the end
+// the +9 skips extra characters ")", newline, "#endif" and newline at the beginning
+// )"" "\n#endif" + 9
+#endif
+

histogram64.cl src/treelearner/ocl/histogram64.cl

@@ -1,3 +1,18 @@
+// this file can either be read and passed to an OpenCL compiler directly,
+// or included in a C++11 source file as a string literal
+#ifndef __OPENCL_VERSION__
+// If we are including this file in C++,
+// the entire source file following (except the last #endif) will become
+// a raw string literal. The extra ")" is just for mathcing parentheses
+// to make the editor happy. The extra ")" and extra endif will be skipped.
+// DO NOT add anything between here and the next #ifdef, otherwise you need
+// to modify the skip count at the end of this file.
+R""()
+#endif
+
+#ifndef _HISTOGRAM_64_KERNEL_
+#define _HISTOGRAM_64_KERNEL_
+
 #pragma OPENCL EXTENSION cl_khr_local_int32_base_atomics : enable
 #pragma OPENCL EXTENSION cl_khr_local_int32_extended_atomics : enable
 
@@ -503,3 +518,8 @@ __kernel void histogram64(__global const uchar4* feature_data_base,
     }
 }
 
+// The following line ends the string literal, adds an extra #endif at the end
+// the +9 skips extra characters ")", newline, "#endif" and newline at the beginning
+// )"" "\n#endif" + 9
+#endif
+