gaspersic/triton
1
0
Fork 0
Code Issues Pull Requests Packages Projects Releases Wiki Activity

Driver: no longer instantiating events when unnecessary ; general cleaning of the driver module

Browse Source
This commit is contained in:
Philippe Tillet
2015-12-21 18:43:05 -05:00
parent 05af3ccc10
commit f2c3aa0acb
12 changed files with 104 additions and 20 deletions
Download Patch File Download Diff File Expand all files Collapse all files

13
include/isaac/driver/backend.h
View File

@@ -42,6 +42,8 @@ class Buffer;
class CommandQueue;
class Context;
class Platform;
class Program;
class Kernel;
class ProgramCache;
class ISAACAPI backend
@@ -69,7 +71,18 @@ public:
DISABLE_MSVC_WARNING_C4251
static std::map<std::tuple<CommandQueue, expression_type, numeric_type>, ProgramCache * > cache_;
RESTORE_MSVC_WARNING_C4251
};
class ISAACAPI kernels
{
friend class backend;
public:
static void release();
static Kernel & get(Program const & program, std::string const & name);
private:
DISABLE_MSVC_WARNING_C4251
static std::map<std::tuple<Program, std::string>, Kernel * > cache_;
RESTORE_MSVC_WARNING_C4251
};
class ISAACAPI contexts

10
include/isaac/driver/buffer.h
View File

@@ -37,6 +37,10 @@ namespace driver
// Buffer
class ISAACAPI Buffer
{
public:
typedef HANDLE_TYPE(cl_mem, CUdeviceptr) handle_type;
private:
friend class CommandQueue;
friend class Kernel;
@@ -54,12 +58,12 @@ public:
Context const & context() const;
bool operator<(Buffer const &) const;
bool operator==(Buffer const &) const;
HANDLE_TYPE(cl_mem, CUdeviceptr)& handle();
HANDLE_TYPE(cl_mem, CUdeviceptr) const & handle() const;
handle_type& handle();
handle_type const & handle() const;
private:
backend_type backend_;
Context context_;
HANDLE_TYPE(cl_mem, CUdeviceptr) h_;
handle_type h_;
};
}

10
include/isaac/driver/command_queue.h
View File

@@ -43,25 +43,29 @@ class Buffer;
// Command Queue
class ISAACAPI CommandQueue
{
public:
typedef HANDLE_TYPE(cl_command_queue, CUstream) handle_type;
public:
CommandQueue(cl_command_queue const & queue, bool take_ownership = true);
CommandQueue(Context const & context, Device const & device, cl_command_queue_properties properties = 0);
backend_type backend() const;
Context const & context() const;
Device const & device() const;
void synchronize();
void enable_profiling();
void disable_profiling();
Event enqueue(Kernel const & kernel, NDRange global, driver::NDRange local, std::vector<Event> const *);
void enqueue(Kernel const & kernel, NDRange global, driver::NDRange local, std::vector<Event> const *, Event *event);
void write(Buffer const & buffer, bool blocking, std::size_t offset, std::size_t size, void const* ptr);
void read(Buffer const & buffer, bool blocking, std::size_t offset, std::size_t size, void* ptr);
bool operator==(CommandQueue const & other) const;
bool operator<(CommandQueue const & other) const;
HANDLE_TYPE(cl_command_queue, CUstream)& handle();
handle_type& handle();
private:
backend_type backend_;
Context context_;
Device device_;
HANDLE_TYPE(cl_command_queue, CUstream) h_;
handle_type h_;
};

9
include/isaac/driver/context.h
View File

@@ -40,6 +40,11 @@ class ISAACAPI Context
friend class Program;
friend class CommandQueue;
friend class Buffer;
public:
typedef HANDLE_TYPE(cl_context, CUcontext) handle_type;
private:
static std::string cache_path();
static CUdevice device(CUcontext)
@@ -59,13 +64,13 @@ public:
bool operator==(Context const &) const;
bool operator<(Context const &) const;
HANDLE_TYPE(cl_context, CUcontext) const & handle() const { return h_; }
handle_type const & handle() const { return h_; }
private:
DISABLE_MSVC_WARNING_C4251
backend_type backend_;
Device device_;
std::string cache_path_;
HANDLE_TYPE(cl_context, CUcontext) h_;
handle_type h_;
RESTORE_MSVC_WARNING_C4251
};

4
include/isaac/driver/device.h
View File

@@ -41,6 +41,8 @@ private:
friend class CommandQueue;
public:
typedef HANDLE_TYPE(cl_device_id, CUdevice) handle_type;
enum Type
{
GPU = CL_DEVICE_TYPE_GPU,
@@ -114,7 +116,7 @@ public:
private:
backend_type backend_;
HANDLE_TYPE(cl_device_id, CUdevice) h_;
handle_type h_;
};
}

9
include/isaac/driver/event.h
View File

@@ -35,15 +35,20 @@ namespace driver
// Event
class ISAACAPI Event
{
private:
friend class CommandQueue;
public:
typedef HANDLE_TYPE(cl_event, cu_event_t) handle_type;
public:
Event(cl_event const & event, bool take_ownership = true);
Event(backend_type backend);
long elapsed_time() const;
HANDLE_TYPE(cl_event, cu_event_t)& handle();
handle_type& handle();
private:
backend_type backend_;
HANDLE_TYPE(cl_event, cu_event_t) h_;
handle_type h_;
};
}

5
include/isaac/driver/kernel.h
View File

@@ -41,6 +41,9 @@ class Buffer;
class ISAACAPI Kernel
{
friend class CommandQueue;
public:
typedef HANDLE_TYPE(cl_kernel, CUfunction) handle_type;
public:
Kernel(Program const & program, const char * name);
void setArg(unsigned int index, std::size_t size, void* ptr);
@@ -53,7 +56,7 @@ private:
unsigned int address_bits_;
std::vector<std::shared_ptr<void> > cu_params_store_;
std::vector<void*> cu_params_;
HANDLE_TYPE(cl_kernel, CUfunction) h_;
handle_type h_;
};
}

8
include/isaac/driver/program.h
View File

@@ -39,16 +39,22 @@ class Device;
class ISAACAPI Program
{
public:
typedef HANDLE_TYPE(cl_program, CUmodule) handle_type;
private:
friend class Kernel;
public:
Program(Context const & context, std::string const & source);
Context const & context() const;
//Comparison operators
bool operator==(Program const & other) const;
bool operator<(Program const & other) const;
private:
DISABLE_MSVC_WARNING_C4251
backend_type backend_;
Context context_;
std::string source_;
HANDLE_TYPE(cl_program, CUmodule) h_;
handle_type h_;
RESTORE_MSVC_WARNING_C4251
};

8
include/isaac/symbolic/expression.h
View File

@@ -262,9 +262,15 @@ struct execution_options_type
void enqueue(driver::Context const & context, driver::Kernel const & kernel, driver::NDRange global, driver::NDRange local) const
{
driver::Event event = queue(context).enqueue(kernel, global, local, dependencies);
driver::CommandQueue & q = queue(context);
if(events)
{
driver::Event event(q.backend());
q.enqueue(kernel, global, local, dependencies, &event);
events->push_back(event);
}
else
q.enqueue(kernel, global, local, dependencies, NULL);
}
driver::CommandQueue & queue(driver::Context const & context) const

23
lib/driver/backend.cpp
View File

@@ -23,6 +23,7 @@
#include "isaac/driver/buffer.h"
#include "isaac/driver/context.h"
#include "isaac/driver/command_queue.h"
#include "isaac/driver/kernel.h"
#include "isaac/driver/program_cache.h"
#include <assert.h>
@@ -76,6 +77,27 @@ ProgramCache & backend::programs::get(CommandQueue const & queue, expression_typ
std::map<std::tuple<CommandQueue, expression_type, numeric_type>, ProgramCache * > backend::programs::cache_;
/*-----------------------------------*/
//----------- Kernels --------------*/
/*-----------------------------------*/
void backend::kernels::release()
{
for(auto & x: cache_)
delete x.second;
cache_.clear();
}
Kernel & backend::kernels::get(Program const & program, std::string const & name)
{
std::tuple<Program, std::string> key(program, name);
if(cache_.find(key)==cache_.end())
return *cache_.insert(std::make_pair(key, new Kernel(program, name.c_str()))).first->second;
return *cache_.at(key);
}
std::map<std::tuple<Program, std::string>, Kernel * > backend::kernels::cache_;
/*-----------------------------------*/
//------------ Queues --------------*/
/*-----------------------------------*/
@@ -219,6 +241,7 @@ void backend::synchronize(Context const & context)
void backend::release()
{
backend::kernels::release();
backend::programs::release();
backend::workspaces::release();
backend::queues::release();

19
lib/driver/command_queue.cpp
View File

@@ -67,6 +67,11 @@ CommandQueue::CommandQueue(Context const & context, Device const & device, cl_co
}
}
backend_type CommandQueue::backend() const
{
return backend_;
}
Context const & CommandQueue::context() const
{
return context_;
@@ -87,23 +92,25 @@ void CommandQueue::synchronize()
}
}
Event CommandQueue::enqueue(Kernel const & kernel, NDRange global, driver::NDRange local, std::vector<Event> const *)
void CommandQueue::enqueue(Kernel const & kernel, NDRange global, driver::NDRange local, std::vector<Event> const *, Event* event)
{
Event event(backend_);
switch(backend_)
{
case CUDA:
cuda::check(dispatch::cuEventRecord(event.h_.cu().first, h_.cu()));
if(event)
cuda::check(dispatch::cuEventRecord(event->h_.cu().first, h_.cu()));
cuda::check(dispatch::cuLaunchKernel(kernel.h_.cu(), global[0]/local[0], global[1]/local[1], global[2]/local[2],
local[0], local[1], local[2], 0, h_.cu(),(void**)&kernel.cu_params_[0], NULL));
cuda::check(dispatch::cuEventRecord(event.h_.cu().second, h_.cu()));
if(event)
cuda::check(dispatch::cuEventRecord(event->h_.cu().second, h_.cu()));
break;
case OPENCL:
ocl::check(dispatch::clEnqueueNDRangeKernel(h_.cl(), kernel.h_.cl(), global.dimension(), NULL, (const size_t *)global, (const size_t *) local, 0, NULL, &event.handle().cl()));
ocl::check(dispatch::clEnqueueNDRangeKernel(h_.cl(), kernel.h_.cl(), global.dimension(), NULL, (const size_t *)global, (const size_t *) local, 0, NULL, &event->h_.cl()));
break;
default: throw;
}
return event;
}
void CommandQueue::write(Buffer const & buffer, bool blocking, std::size_t offset, std::size_t size, void const* ptr)

6
lib/driver/program.cpp
View File

@@ -177,6 +177,12 @@ Program::Program(Context const & context, std::string const & source) : backend_
}
}
bool Program::operator==(Program const & other) const
{ return h_ == other.h_; }
bool Program::operator<(Program const & other) const
{ return h_ < other.h_; }
Context const & Program::context() const
{
return context_;