Merge pull request #19 from listenlink/upstream

intelblas_gemm clean patch
triton-lang · Feb 16, 2017 · 9b97c56 · 9b97c56
2 parents 74e9a03 + 461b83d
commit 9b97c56
Show file tree

Hide file tree

Showing 18 changed files with 219,389 additions and 57,504 deletions.
diff --git a/CMakeLists.txt b/CMakeLists.txt
@@ -47,8 +47,6 @@ foreach(FILE ${LIBISAAC_SRC})
     set(LIBISAAC_SRC_STR "${_TMP} ${LIBISAAC_SRC_STR}")
 endforeach()
 
-
-
 #Include directories
 set(INCLUDE_DIRECTORIES_STR)
 get_property(INCLUDE_DIRECTORIES_LST DIRECTORY ${CMAKE_CURRENT_SOURCE_DIR} PROPERTY INCLUDE_DIRECTORIES)

diff --git a/include/isaac/driver/dispatch.h b/include/isaac/driver/dispatch.h
@@ -116,8 +116,11 @@ class dispatch
     static cl_int clGetKernelWorkGroupInfo(cl_kernel, cl_device_id, cl_kernel_work_group_info, size_t, void *, size_t *);
     static cl_kernel clCreateKernel(cl_program, const char *, cl_int *);
     static cl_mem clCreateBuffer(cl_context, cl_mem_flags, size_t, void *, cl_int *);
+    static cl_mem clCreateImage(cl_context, cl_mem_flags, const cl_image_format *, const cl_image_desc *, void *, cl_int *);
     static cl_program clCreateProgramWithSource(cl_context, cl_uint, const char **, const size_t *, cl_int *);
     static cl_int clReleaseKernel(cl_kernel);
+    static cl_int clEnqueueCopyBufferToImage(cl_command_queue, cl_mem, cl_mem, size_t, const size_t *, const size_t *, cl_uint, const cl_event *, cl_event *);
+    static cl_int clSetEventCallback(cl_event, cl_int, void (CL_CALLBACK * /* pfn_notify */)(cl_event, cl_int, void *), void *);
 
     //CUDA
     static CUresult cuCtxDestroy_v2(CUcontext ctx);
@@ -202,8 +205,11 @@ class dispatch
     static void* clGetKernelWorkGroupInfo_;
     static void* clCreateKernel_;
     static void* clCreateBuffer_;
+    static void* clCreateImage_;
     static void* clCreateProgramWithSource_;
     static void* clReleaseKernel_;
+    static void* clEnqueueCopyBufferToImage_;
+    static void* clSetEventCallback_;
 
     //CUDA
     static void* cuCtxDestroy_v2_;

diff --git a/include/isaac/jit/generation/gemm.h b/include/isaac/jit/generation/gemm.h
@@ -47,6 +47,39 @@ class cublas_gemm : public external_base
   bool init_;
 };
 
+class intelblas_gemm : public external_base
+{
+  bool init();
+public:
+  intelblas_gemm(char A_trans, char B_trans);
+  int is_invalid(expression_tree const  &, driver::Device const &) const;
+  std::vector<int_t> input_sizes(expression_tree const & expressions) const;
+  void enqueue(driver::CommandQueue & queue, driver::Program const &, std::string const &, runtime::execution_handler const & h);
+  expression_type type() const;
+private:
+  std::string generate_impl(std::string const & suffix, expression_tree const &, driver::Device const & device, symbolic::symbols_table const &) const;
+  const char A_trans_;
+  const char B_trans_;
+  bool init_;
+};
+
+class intelblas_gemm_image : public external_base
+{
+  bool init();
+public:
+  intelblas_gemm_image(char A_trans, char B_trans);
+  int is_invalid(expression_tree const  &, driver::Device const &) const;
+  std::vector<int_t> input_sizes(expression_tree const & expressions) const;
+  void enqueue(driver::CommandQueue & queue, driver::Program const &, std::string const &, runtime::execution_handler const & h);
+  expression_type type() const;
+private:
+  std::string generate_impl(std::string const & suffix, expression_tree const &, driver::Device const & device, symbolic::symbols_table const &) const;
+  const char A_trans_;
+  const char B_trans_;
+  bool init_;
+};
+
+
 class gemm : public parameterized_base
 {
 private:

diff --git a/lib/api/blas/clBLAS.cpp b/lib/api/blas/clBLAS.cpp
@@ -278,8 +278,8 @@ extern "C"
         }\
         sc::int_t As1 = (sc::int_t)M, As2 = (sc::int_t)K;\
         sc::int_t Bs1 = (sc::int_t)K, Bs2 = (sc::int_t)N;\
-        if(transA==clblasTrans) std::swap(As1, As2);\
-        if(transB==clblasTrans) std::swap(Bs1, Bs2);\
+        if(transA==clblasTrans || transA==clblasConjTrans) std::swap(As1, As2);\
+        if(transB==clblasTrans || transB==clblasConjTrans) std::swap(Bs1, Bs2);\
         /*Struct*/\
         sc::array A(As1, As2, TYPE_ISAAC, sc::driver::Buffer(mA, false), (sc::int_t)offA, (sc::int_t)lda);\
         sc::array B(Bs1, Bs2, TYPE_ISAAC, sc::driver::Buffer(mB, false), (sc::int_t)offB, (sc::int_t)ldb);\

diff --git a/lib/driver/device.cpp b/lib/driver/device.cpp
@@ -301,6 +301,9 @@ std::string Device::infos() const
   return oss.str();
 }
 
+Device::handle_type const & Device::handle() const
+{ return h_; }
+
 // Properties
 #define WRAP_ATTRIBUTE(ret, fname, CUNAME, CLNAME) \
   ret Device::fname() const\

diff --git a/lib/driver/dispatch.cpp b/lib/driver/dispatch.cpp
@@ -142,6 +142,9 @@ cl_int dispatch::clBuildProgram(cl_program a, cl_uint b, const cl_device_id * c,
 cl_context dispatch::clCreateContext(const cl_context_properties * a, cl_uint b, const cl_device_id * c, void (*d)(const char *, const void *, size_t, void *), void * e, cl_int * f)
 { return f_impl<dispatch::clinit>(opencl_, dispatch::clCreateContext, dispatch::clCreateContext_, "clCreateContext", a, b, c, d, e, f); }
 
+cl_int dispatch::clSetEventCallback(cl_event event, cl_int a, void(CL_CALLBACK *pfn_notify)(cl_event, cl_int, void *), void * arg)
+{ return f_impl<dispatch::clinit>(opencl_, dispatch::clSetEventCallback, dispatch::clSetEventCallback_, "clSetEventCallback", event, a, pfn_notify, arg); }
+
 OCL_DEFINE9(cl_int, clEnqueueNDRangeKernel, cl_command_queue, cl_kernel, cl_uint, const size_t*, const size_t*, const size_t*,  cl_uint, const cl_event*, cl_event*)
 OCL_DEFINE4(cl_int, clSetKernelArg, cl_kernel, cl_uint, size_t, const void *)
 OCL_DEFINE1(cl_int, clReleaseMemObject, cl_mem)
@@ -171,8 +174,10 @@ OCL_DEFINE5(cl_int, clGetKernelInfo, cl_kernel, cl_kernel_info, size_t, void *,
 OCL_DEFINE6(cl_int, clGetKernelWorkGroupInfo, cl_kernel, cl_device_id, cl_kernel_work_group_info, size_t, void *, size_t *)
 OCL_DEFINE3(cl_kernel, clCreateKernel, cl_program, const char *, cl_int *)
 OCL_DEFINE5(cl_mem, clCreateBuffer, cl_context, cl_mem_flags, size_t, void *, cl_int *)
+OCL_DEFINE6(cl_mem, clCreateImage, cl_context, cl_mem_flags, const cl_image_format *, const cl_image_desc *, void *, cl_int *) 
 OCL_DEFINE5(cl_program, clCreateProgramWithSource, cl_context, cl_uint, const char **, const size_t *, cl_int *)
 OCL_DEFINE1(cl_int, clReleaseKernel, cl_kernel)
+OCL_DEFINE9(cl_int, clEnqueueCopyBufferToImage, cl_command_queue, cl_mem, cl_mem, size_t, const size_t *, const size_t *, cl_uint, const cl_event *, cl_event *)
 
 //CUDA
 CUDA_DEFINE1(CUresult, cuCtxDestroy_v2, CUcontext)
@@ -291,8 +296,11 @@ void* dispatch::clGetKernelInfo_;
 void* dispatch::clGetKernelWorkGroupInfo_;
 void* dispatch::clCreateKernel_;
 void* dispatch::clCreateBuffer_;
+void* dispatch::clCreateImage_;
 void* dispatch::clCreateProgramWithSource_;
 void* dispatch::clReleaseKernel_;
+void* dispatch::clEnqueueCopyBufferToImage_;
+void* dispatch::clSetEventCallback_;
 
 //CUDA
 void* dispatch::cuCtxDestroy_v2_;

diff --git a/lib/driver/kernel.cpp b/lib/driver/kernel.cpp
@@ -127,6 +127,9 @@ void Kernel::setSizeArg(unsigned int index, size_t N)
   }
 }
 
+Kernel::handle_type const & Kernel::handle() const
+{ return h_; }
+
 }
 
 }

diff --git a/lib/jit/generation/elementwise_2d.cpp b/lib/jit/generation/elementwise_2d.cpp
@@ -28,7 +28,6 @@
 #include "tools/loop.hpp"
 #include "tools/vector_types.hpp"
 
-
 namespace isaac
 {
 namespace templates