Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Jump to bottom

intelblas_gemm clean patch #19

Merged
merged 9 commits into from
Feb 16, 2017
Merged

intelblas_gemm clean patch #19

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
9 commits
Select commit Hold shift + click to select a range
ce1c65d
integrate intelblas_gemm.
wujunkai166 Nov 21, 2016
bf28d43
Added log printing of external kernels
Dec 16, 2016
f04f866
Modified external template bench print for more than one template for…
Dec 28, 2016
8702578
Add image support for gemm kernel.
wujunkai166 Dec 29, 2016
4b5cf48
Fix some minor issues and update json files for Intel platforms.
Jan 19, 2017
b882fe9
Introduce some new dispatch APIs for OCL.
Jan 20, 2017
1bc3be6
Add image block version of gemm kernel. No need to fall back to
wujunkai166 Jan 19, 2017
bc9121d
Added intelblas python expternal scripts and updated json files
Feb 8, 2017
461b83d
Ocl building warnings removed
Feb 16, 2017
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
2 changes: 0 additions & 2 deletions CMakeLists.txt
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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)
Expand Down
6 changes: 6 additions & 0 deletions include/isaac/driver/dispatch.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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);
Expand Down Expand Up @@ -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_;
Expand Down
33 changes: 33 additions & 0 deletions include/isaac/jit/generation/gemm.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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:
Expand Down
4 changes: 2 additions & 2 deletions lib/api/blas/clBLAS.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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);\
Expand Down
3 changes: 3 additions & 0 deletions lib/driver/device.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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\
Expand Down
8 changes: 8 additions & 0 deletions lib/driver/dispatch.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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)
Expand Down Expand Up @@ -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)
Expand Down Expand Up @@ -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_;
Expand Down
3 changes: 3 additions & 0 deletions lib/driver/kernel.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -127,6 +127,9 @@ void Kernel::setSizeArg(unsigned int index, size_t N)
}
}

Kernel::handle_type const & Kernel::handle() const
{ return h_; }

}

}
Expand Down
1 change: 0 additions & 1 deletion lib/jit/generation/elementwise_2d.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -28,7 +28,6 @@
#include "tools/loop.hpp"
#include "tools/vector_types.hpp"


namespace isaac
{
namespace templates
Expand Down
Loading