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

Code quality: made the backend static

Browse Source
This commit is contained in:
Philippe Tillet
2015-07-30 21:02:30 -07:00
parent 21a2566904
commit 29bdf7f546
13 changed files with 65 additions and 65 deletions
Download Patch File Download Diff File Expand all files Collapse all files

10
bench/blas.cpp
View File

@@ -171,7 +171,7 @@ void bench(isc::numeric_type dtype, std::string operation)
} }
unsigned int dtsize = isc::size_of(dtype); unsigned int dtsize = isc::size_of(dtype);
isc::driver::CommandQueue & queue = isc::driver::queues.default_queues()[0]; isc::driver::CommandQueue & queue = isc::driver::backend::default_queues()[0];
std::map<std::string, std::string> metric{ {"axpy", "GB/s"}, {"dot", "GB/s"}, {"gemv", "GB/s"}, {"gemm", "GFLOPS"}}; std::map<std::string, std::string> metric{ {"axpy", "GB/s"}, {"dot", "GB/s"}, {"gemv", "GB/s"}, {"gemm", "GFLOPS"}};
isc::array flush(1e6, dtype); isc::array flush(1e6, dtype);
std::cout << "#" << operation << " (" << metric[operation] << ")" << std::endl; std::cout << "#" << operation << " (" << metric[operation] << ")" << std::endl;
@@ -400,10 +400,10 @@ int main(int argc, char* argv[])
#ifdef BENCH_CLBLAS #ifdef BENCH_CLBLAS
clblasSetup(); clblasSetup();
#endif #endif
isc::driver::queues.queue_properties = CL_QUEUE_PROFILING_ENABLE; isc::driver::backend::queue_properties = CL_QUEUE_PROFILING_ENABLE;
int device_idx = 0; int device_idx = 0;
isc::driver::queues_type::container_type queues = isc::driver::queues.contexts(); isc::driver::backend::container_type queues = isc::driver::backend::contexts();
std::string operation; std::string operation;
if(queues.size() > 1) if(queues.size() > 1)
@@ -413,7 +413,7 @@ int main(int argc, char* argv[])
std::cerr << "usage : blas-bench DEVICE_IDX OPERATION" << std::endl; std::cerr << "usage : blas-bench DEVICE_IDX OPERATION" << std::endl;
std::cout << "Devices available: " << std::endl; std::cout << "Devices available: " << std::endl;
unsigned int current=0; unsigned int current=0;
for(isc::driver::queues_type::container_type::const_iterator it = queues.begin() ; it != queues.end() ; ++it) for(isc::driver::backend::container_type::const_iterator it = queues.begin() ; it != queues.end() ; ++it)
{ {
isc::driver::Device device = it->first.device(); isc::driver::Device device = it->first.device();
std::cout << current++ << ": " << device.name() << " on " << device.platform().name() << " " << device.platform().version() << std::endl; std::cout << current++ << ": " << device.name() << " on " << device.platform().name() << " " << device.platform().version() << std::endl;
@@ -433,7 +433,7 @@ int main(int argc, char* argv[])
operation = args[1]; operation = args[1];
} }
isc::driver::queues.default_device = device_idx; isc::driver::backend::default_device = device_idx;
std::cout << "#Benchmark : BLAS" << std::endl; std::cout << "#Benchmark : BLAS" << std::endl;
std::cout << "#----------------" << std::endl; std::cout << "#----------------" << std::endl;
bench<float>(isc::FLOAT_TYPE, operation); bench<float>(isc::FLOAT_TYPE, operation);

2
bench/overhead.cpp
View File

@@ -12,7 +12,7 @@ __global__ void dummy(){}
int main() int main()
{ {
for(isc::driver::queues_type::data_type::const_iterator it = isc::driver::queues.data().begin() ; it != isc::driver::queues.data().end() ; ++it) for(isc::driver::backend::data_type::const_iterator it = isc::driver::queues.data().begin() ; it != isc::driver::queues.data().end() ; ++it)
{ {
cl::CommandQueue queue = it->second[0]; cl::CommandQueue queue = it->second[0];
cl::Context context = it->first; cl::Context context = it->first;

18
include/isaac/array.h
View File

@@ -21,22 +21,22 @@ protected:
public: public:
//1D Constructors //1D Constructors
array(int_t size1, numeric_type dtype, driver::Context context = driver::queues.default_context()); array(int_t size1, numeric_type dtype, driver::Context context = driver::backend::default_context());
array(int_t size1, numeric_type dtype, driver::Buffer data, int_t start, int_t inc); array(int_t size1, numeric_type dtype, driver::Buffer data, int_t start, int_t inc);
template<typename DT> template<typename DT>
array(std::vector<DT> const & data, driver::Context context = driver::queues.default_context()); array(std::vector<DT> const & data, driver::Context context = driver::backend::default_context());
array(array & v, slice const & s1); array(array & v, slice const & s1);
//2D Constructors //2D Constructors
array(int_t size1, int_t size2, numeric_type dtype, driver::Context context = driver::queues.default_context()); array(int_t size1, int_t size2, numeric_type dtype, driver::Context context = driver::backend::default_context());
array(int_t size1, int_t size2, numeric_type dtype, driver::Buffer data, int_t start, int_t ld); array(int_t size1, int_t size2, numeric_type dtype, driver::Buffer data, int_t start, int_t ld);
template<typename DT> template<typename DT>
array(int_t size1, int_t size2, std::vector<DT> const & data, driver::Context context = driver::queues.default_context()); array(int_t size1, int_t size2, std::vector<DT> const & data, driver::Context context = driver::backend::default_context());
array(array & M, slice const & s1, slice const & s2); array(array & M, slice const & s1, slice const & s2);
//3D Constructors //3D Constructors
array(int_t size1, int_t size2, int_t size3, numeric_type dtype, driver::Context context = driver::queues.default_context()); array(int_t size1, int_t size2, int_t size3, numeric_type dtype, driver::Context context = driver::backend::default_context());
//General constructor //General constructor
@@ -112,8 +112,8 @@ private:
template<class T> T cast() const; template<class T> T cast() const;
public: public:
explicit scalar(numeric_type dtype, driver::Buffer const & data, int_t offset); explicit scalar(numeric_type dtype, driver::Buffer const & data, int_t offset);
explicit scalar(value_scalar value, driver::Context context = driver::queues.default_context()); explicit scalar(value_scalar value, driver::Context context = driver::backend::default_context());
explicit scalar(numeric_type dtype, driver::Context context = driver::queues.default_context()); explicit scalar(numeric_type dtype, driver::Context context = driver::backend::default_context());
scalar(array_expression const & proxy); scalar(array_expression const & proxy);
scalar& operator=(value_scalar const &); scalar& operator=(value_scalar const &);
// scalar& operator=(scalar const & s); // scalar& operator=(scalar const & s);
@@ -230,8 +230,8 @@ ISAAC_DECLARE_DOT((max))
ISAAC_DECLARE_DOT((min)) ISAAC_DECLARE_DOT((min))
ISAAC_DECLARE_DOT(argmin) ISAAC_DECLARE_DOT(argmin)
ISAACAPI array_expression eye(std::size_t, std::size_t, isaac::numeric_type, driver::Context context = driver::queues.default_context()); ISAACAPI array_expression eye(std::size_t, std::size_t, isaac::numeric_type, driver::Context context = driver::backend::default_context());
ISAACAPI array_expression zeros(std::size_t M, std::size_t N, numeric_type dtype, driver::Context context = driver::queues.default_context()); ISAACAPI array_expression zeros(std::size_t M, std::size_t N, numeric_type dtype, driver::Context context = driver::backend::default_context());
ISAACAPI array_expression reshape(array const &, int_t, int_t); ISAACAPI array_expression reshape(array const &, int_t, int_t);
// //

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

@@ -13,35 +13,31 @@ namespace isaac
namespace driver namespace driver
{ {
class ISAACAPI queues_type class ISAACAPI backend
{ {
public: public:
typedef std::list<std::pair<Context, std::vector<CommandQueue> > > container_type; typedef std::list<std::pair<Context, std::vector<CommandQueue> > > container_type;
private: private:
std::vector<CommandQueue> & append( Context const & context); static std::vector<CommandQueue> & append( Context const & context);
void cuinit(); static void cuinit();
void clinit(); static void clinit();
void init(); static void init();
public: public:
queues_type(); static container_type const & contexts();
container_type const & contexts(); static Context default_context();
Context default_context(); static std::vector<CommandQueue> & default_queues();
std::vector<CommandQueue> & default_queues(); static std::vector<CommandQueue> & queues(Context const &);
std::vector<CommandQueue> & operator[](Context const &);
private: private:
container_type data_; static container_type data_;
public: public:
unsigned int default_device; static unsigned int default_device;
cl_command_queue_properties queue_properties; static cl_command_queue_properties queue_properties;
}; };
ISAACAPI void synchronize(std::vector<CommandQueue> const &); ISAACAPI void synchronize(std::vector<CommandQueue> const &);
ISAACAPI void synchronize(Context const &); ISAACAPI void synchronize(Context const &);
extern ISAACAPI queues_type queues;
} }
} }

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

@@ -238,7 +238,7 @@ struct execution_options_type
{ {
if(queue_) if(queue_)
return *queue_; return *queue_;
return driver::queues[context][queue_id_]; return driver::backend::queues(context)[queue_id_];
} }
std::list<driver::Event>* events; std::list<driver::Event>* events;

16
lib/array.cpp
View File

@@ -284,9 +284,9 @@ namespace detail
{ {
template<class T> template<class T>
void copy(driver::Context & ctx, driver::Buffer const & data, T value) void copy(driver::Context & context, driver::Buffer const & data, T value)
{ {
driver::queues[ctx][0].write(data, CL_TRUE, 0, sizeof(T), (void*)&value); driver::backend::queues(context)[0].write(data, CL_TRUE, 0, sizeof(T), (void*)&value);
} }
} }
@@ -323,7 +323,7 @@ void scalar::inject(values_holder & v) const
int_t dtsize = size_of(dtype_); int_t dtsize = size_of(dtype_);
#define HANDLE_CASE(DTYPE, VAL) \ #define HANDLE_CASE(DTYPE, VAL) \
case DTYPE:\ case DTYPE:\
driver::queues[context_][0].read(data_, CL_TRUE, start_[0]*dtsize, dtsize, (void*)&v.VAL); break;\ driver::backend::queues(context_)[0].read(data_, CL_TRUE, start_[0]*dtsize, dtsize, (void*)&v.VAL); break;\
switch(dtype_) switch(dtype_)
{ {
@@ -370,7 +370,7 @@ TYPE scalar::cast() const
scalar& scalar::operator=(value_scalar const & s) scalar& scalar::operator=(value_scalar const & s)
{ {
driver::CommandQueue& queue = driver::queues[context_][0]; driver::CommandQueue& queue = driver::backend::queues(context_)[0];
int_t dtsize = size_of(dtype_); int_t dtsize = size_of(dtype_);
#define HANDLE_CASE(TYPE, CLTYPE) case TYPE:\ #define HANDLE_CASE(TYPE, CLTYPE) case TYPE:\
@@ -831,10 +831,10 @@ void copy(array const & x, void* data, driver::CommandQueue & queue, bool blocki
} }
void copy(void const *data, array &x, bool blocking) void copy(void const *data, array &x, bool blocking)
{ copy(data, x, driver::queues[x.context()][0], blocking); } { copy(data, x, driver::backend::queues(x.context())[0], blocking); }
void copy(array const & x, void* data, bool blocking) void copy(array const & x, void* data, bool blocking)
{ copy(x, data, driver::queues[x.context()][0], blocking); } { copy(x, data, driver::backend::queues(x.context())[0], blocking); }
//std::vector<> //std::vector<>
template<class T> template<class T>
@@ -859,11 +859,11 @@ void copy(array const & x, std::vector<T> & cx, driver::CommandQueue & queue, bo
template<class T> template<class T>
void copy(std::vector<T> const & cx, array & x, bool blocking) void copy(std::vector<T> const & cx, array & x, bool blocking)
{ copy(cx, x, driver::queues[x.context()][0], blocking); } { copy(cx, x, driver::backend::queues(x.context())[0], blocking); }
template<class T> template<class T>
void copy(array const & x, std::vector<T> & cx, bool blocking) void copy(array const & x, std::vector<T> & cx, bool blocking)
{ copy(x, cx, driver::queues[x.context()][0], blocking); } { copy(x, cx, driver::backend::queues(x.context())[0], blocking); }
#define INSTANTIATE(T) \ #define INSTANTIATE(T) \
template void ISAACAPI copy<T>(std::vector<T> const &, array &, driver::CommandQueue&, bool);\ template void ISAACAPI copy<T>(std::vector<T> const &, array &, driver::CommandQueue&, bool);\

32
lib/driver/backend.cpp
View File

@@ -9,16 +9,13 @@ namespace isaac
namespace driver namespace driver
{ {
queues_type::queues_type(): default_device(0), queue_properties(0) std::vector<CommandQueue> & backend::append(Context const & context)
{}
std::vector<CommandQueue> & queues_type::append(Context const & context)
{ {
data_.push_back(std::make_pair(context, std::vector<CommandQueue>(1, CommandQueue(context, context.device(), queue_properties)))); data_.push_back(std::make_pair(context, std::vector<CommandQueue>(1, CommandQueue(context, context.device(), queue_properties))));
return data_.back().second; return data_.back().second;
} }
void queues_type::cuinit() void backend::cuinit()
{ {
#ifdef ISAAC_WITH_CUDA #ifdef ISAAC_WITH_CUDA
cuda::check(cuInit(0)); cuda::check(cuInit(0));
@@ -29,7 +26,7 @@ void queues_type::cuinit()
#endif #endif
} }
void queues_type::clinit() void backend::clinit()
{ {
cl_uint nplatforms; cl_uint nplatforms;
ocl::check(clGetPlatformIDs(0, NULL, &nplatforms)); ocl::check(clGetPlatformIDs(0, NULL, &nplatforms));
@@ -46,7 +43,7 @@ void queues_type::clinit()
} }
} }
void queues_type::init() void backend::init()
{ {
if(data_.empty()) if(data_.empty())
{ {
@@ -55,7 +52,7 @@ void queues_type::init()
} }
} }
std::vector<CommandQueue> & queues_type::operator[](Context const & context) std::vector<CommandQueue> & backend::queues(Context const & context)
{ {
init(); init();
for(auto & x : data_) for(auto & x : data_)
@@ -63,7 +60,7 @@ std::vector<CommandQueue> & queues_type::operator[](Context const & context)
return append(context); return append(context);
} }
Context queues_type::default_context() Context backend::default_context()
{ {
init(); init();
container_type::iterator it = data_.begin(); container_type::iterator it = data_.begin();
@@ -71,11 +68,11 @@ Context queues_type::default_context()
return it->first; return it->first;
} }
std::vector<CommandQueue> & queues_type::default_queues() std::vector<CommandQueue> & backend::default_queues()
{ return (*this)[default_context()]; } { return backend::queues(default_context()); }
queues_type::container_type const & queues_type::contexts() backend::container_type const & backend::contexts()
{ {
init(); init();
return data_; return data_;
@@ -89,9 +86,16 @@ ISAACAPI void synchronize(std::vector<CommandQueue > & queues)
} }
ISAACAPI void synchronize(Context const & context) ISAACAPI void synchronize(Context const & context)
{ synchronize(queues[context]); } { synchronize(backend::queues(context)); }
ISAACAPI queues_type queues;
//Static variables
unsigned int backend::default_device = 0;
cl_command_queue_properties backend::queue_properties = 0;
backend::container_type backend::data_ = backend::container_type();
} }

4
python/src/wrap/driver.cpp
View File

@@ -146,7 +146,7 @@ void export_driver()
bp::class_<isc::driver::Context>("context", bp::no_init) bp::class_<isc::driver::Context>("context", bp::no_init)
.def("__init__", bp::make_constructor(&detail::make_context)) .def("__init__", bp::make_constructor(&detail::make_context))
.add_property("queues", bp::make_function(static_cast<std::vector<isc::driver::CommandQueue> & (*)(const isc::driver::Context&)>( [](const isc::driver::Context & ctx) -> std::vector<isc::driver::CommandQueue> & { return isc::driver::queues[ctx]; }) , bp::return_internal_reference<>())) .add_property("queues", bp::make_function(static_cast<std::vector<isc::driver::CommandQueue> & (*)(const isc::driver::Context&)>( [](const isc::driver::Context & ctx) -> std::vector<isc::driver::CommandQueue> & { return isc::driver::backend[ctx]; }) , bp::return_internal_reference<>()))
.add_property("backend", &isc::driver::Context::backend) .add_property("backend", &isc::driver::Context::backend)
; ;
@@ -167,7 +167,7 @@ void export_driver()
bp::def("enqueue", &detail::enqueue, (bp::arg("expression"), bp::arg("queue_id") = 0, bp::arg("dependencies")=bp::list(), bp::arg("tune") = false, bp::arg("label")=-1, bp::arg("program_name")="", bp::arg("recompile") = false)); bp::def("enqueue", &detail::enqueue, (bp::arg("expression"), bp::arg("queue_id") = 0, bp::arg("dependencies")=bp::list(), bp::arg("tune") = false, bp::arg("label")=-1, bp::arg("program_name")="", bp::arg("recompile") = false));
bp::class_<state_type>("state_type") bp::class_<state_type>("state_type")
.def_readwrite("queue_properties",&isc::driver::queues.queue_properties) .def_readwrite("queue_properties",&isc::driver::backend::queue_properties)
; ;
bp::scope().attr("state") = bp::object(bp::ptr(&state)); bp::scope().attr("state") = bp::object(bp::ptr(&state));

4
tests/linalg/axpy.cpp
View File

@@ -16,7 +16,7 @@ void test_element_wise_vector(T epsilon, simple_vector_base<T> & cx, simple_vect
int failure_count = 0; int failure_count = 0;
isc::numeric_type dtype = x.dtype(); isc::numeric_type dtype = x.dtype();
isc::driver::Context const & ctx = x.context(); isc::driver::Context const & ctx = x.context();
isc::driver::CommandQueue queue = isc::driver::queues[ctx][0]; isc::driver::CommandQueue queue = isc::driver::backend::queues(ctx)[0];
cl_command_queue clqueue = queue.handle().cl(); cl_command_queue clqueue = queue.handle().cl();
int_t N = cz.size(); int_t N = cz.size();
@@ -140,7 +140,7 @@ void test_impl(T epsilon, isc::driver::Context const & ctx)
int main() int main()
{ {
clblasSetup(); clblasSetup();
auto data = isc::driver::queues.contexts(); auto data = isc::driver::backend::contexts();
for(const auto & elem : data) for(const auto & elem : data)
{ {
isc::driver::Device device = elem.second[0].device(); isc::driver::Device device = elem.second[0].device();

4
tests/linalg/dot.cpp
View File

@@ -15,7 +15,7 @@ void test_reduction(T epsilon, simple_vector_base<T> & cx, simple_vector_base<T
using namespace std; using namespace std;
isc::driver::Context const & ctx = x.context(); isc::driver::Context const & ctx = x.context();
int_t N = cx.size(); int_t N = cx.size();
isc::driver::CommandQueue queue = isc::driver::queues[ctx][0]; isc::driver::CommandQueue queue = isc::driver::backend::queues(ctx)[0];
cl_command_queue clqueue = queue.handle().cl(); cl_command_queue clqueue = queue.handle().cl();
isc::array scratch(N, x.dtype()); isc::array scratch(N, x.dtype());
@@ -90,7 +90,7 @@ void test_impl(T epsilon, isc::driver::Context const & ctx)
int main() int main()
{ {
clblasSetup(); clblasSetup();
auto data = isc::driver::queues.contexts(); auto data = isc::driver::backend::contexts();
for(const auto & elem : data) for(const auto & elem : data)
{ {
isc::driver::Device device = elem.second[0].device(); isc::driver::Device device = elem.second[0].device();

4
tests/linalg/gemm.cpp
View File

@@ -20,7 +20,7 @@ void test_impl(T epsilon, simple_matrix_base<T> & cC, simple_matrix_base<T> cons
T alpha = 1; T alpha = 1;
T beta = 0; T beta = 0;
isc::driver::CommandQueue queue = isc::driver::queues[C.context()][0]; isc::driver::CommandQueue queue = isc::driver::backend::queues(C.context())[0];
for(int i = 0 ; i < M ; ++i) for(int i = 0 ; i < M ; ++i)
{ {
@@ -125,7 +125,7 @@ void test_impl(T epsilon, isc::driver::Context const & ctx)
int main() int main()
{ {
auto data = isc::driver::queues.contexts(); auto data = isc::driver::backend::contexts();
clblasSetup(); clblasSetup();
for(const auto & elem : data) for(const auto & elem : data)
{ {

4
tests/linalg/gemv.cpp
View File

@@ -20,7 +20,7 @@ void test_row_wise_reduction(T epsilon, simple_vector_base<T> & cy, simple_matri
T alpha = 4.2, beta = 5.6; T alpha = 4.2, beta = 5.6;
isc::driver::CommandQueue queue = isc::driver::queues[y.context()][0]; isc::driver::CommandQueue queue = isc::driver::backend::queues(y.context())[0];
T yi = 0, xi = 0; T yi = 0, xi = 0;
#define TEST_OPERATION(NAME, SIZE1, SIZE2, REDUCTION, ASSIGNMENT, GPU_REDUCTION, RES, BUF, CRES)\ #define TEST_OPERATION(NAME, SIZE1, SIZE2, REDUCTION, ASSIGNMENT, GPU_REDUCTION, RES, BUF, CRES)\
@@ -108,7 +108,7 @@ void test_impl(T epsilon, isc::driver::Context const & ctx)
int main() int main()
{ {
clblasSetup(); clblasSetup();
auto data = isc::driver::queues.contexts(); auto data = isc::driver::backend::contexts();
for(const auto & elem : data) for(const auto & elem : data)
{ {
isc::driver::Device device = elem.second[0].device(); isc::driver::Device device = elem.second[0].device();

2
tests/linalg/ger.cpp
View File

@@ -124,7 +124,7 @@ void test_impl(T epsilon, isc::driver::Context const & ctx)
int main() int main()
{ {
auto data = isc::driver::queues.contexts(); auto data = isc::driver::backend::contexts();
for(const auto & elem : data) for(const auto & elem : data)
{ {
isc::driver::Device device = elem.second[0].device(); isc::driver::Device device = elem.second[0].device();