2015-12-19 21:35:35 -05:00
|
|
|
/*
|
|
|
|
* Copyright (c) 2015, PHILIPPE TILLET. All rights reserved.
|
|
|
|
*
|
|
|
|
* This file is part of ISAAC.
|
|
|
|
*
|
|
|
|
* ISAAC is free software; you can redistribute it and/or
|
|
|
|
* modify it under the terms of the GNU Lesser General Public
|
|
|
|
* License as published by the Free Software Foundation; either
|
|
|
|
* version 2.1 of the License, or (at your option) any later version.
|
|
|
|
*
|
|
|
|
* This library is distributed in the hope that it will be useful,
|
|
|
|
* but WITHOUT ANY WARRANTY; without even the implied warranty of
|
|
|
|
* MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the GNU
|
|
|
|
* Lesser General Public License for more details.
|
|
|
|
*
|
|
|
|
* You should have received a copy of the GNU Lesser General Public
|
|
|
|
* License along with this library; if not, write to the Free Software
|
|
|
|
* Foundation, Inc., 51 Franklin Street, Fifth Floor, Boston,
|
|
|
|
* MA 02110-1301 USA
|
|
|
|
*/
|
2015-12-21 17:04:09 -05:00
|
|
|
|
2015-01-12 13:20:53 -05:00
|
|
|
#include <cassert>
|
2015-07-28 15:45:14 -07:00
|
|
|
#include <algorithm>
|
2015-08-06 20:20:08 -07:00
|
|
|
#include <string>
|
2015-01-12 13:20:53 -05:00
|
|
|
|
2015-04-29 15:50:57 -04:00
|
|
|
#include "isaac/array.h"
|
2015-09-30 15:31:41 -04:00
|
|
|
#include "isaac/tuple.h"
|
2015-08-04 20:56:05 -07:00
|
|
|
#include "isaac/kernels/keywords.h"
|
2015-12-12 18:32:06 -05:00
|
|
|
#include "isaac/kernels/templates/elementwise_1d.h"
|
|
|
|
#include "isaac/kernels/templates/reduce_1d.h"
|
|
|
|
#include "isaac/kernels/templates/elementwise_2d.h"
|
|
|
|
#include "isaac/kernels/templates/reduce_2d.h"
|
|
|
|
#include "isaac/kernels/templates/matrix_product.h"
|
2015-08-04 20:56:05 -07:00
|
|
|
#include "isaac/kernels/templates/base.h"
|
|
|
|
#include "isaac/kernels/parse.h"
|
2015-04-29 15:50:57 -04:00
|
|
|
#include "isaac/exception/unknown_datatype.h"
|
2015-08-17 16:30:21 -07:00
|
|
|
#include "isaac/exception/operation_not_supported.h"
|
2015-04-29 15:50:57 -04:00
|
|
|
#include "isaac/symbolic/io.h"
|
|
|
|
|
2015-08-06 16:14:33 -07:00
|
|
|
#include "tools/map.hpp"
|
2015-12-19 21:35:35 -05:00
|
|
|
#include "cpp/to_string.hpp"
|
2015-08-06 12:05:12 -07:00
|
|
|
|
2015-04-29 15:50:57 -04:00
|
|
|
namespace isaac
|
2015-01-12 13:20:53 -05:00
|
|
|
{
|
2015-07-11 09:36:01 -04:00
|
|
|
namespace templates
|
|
|
|
{
|
2015-01-12 13:20:53 -05:00
|
|
|
|
2015-01-17 10:48:02 -05:00
|
|
|
base::parameters_type::parameters_type(unsigned int _simd_width, int_t _local_size_1, int_t _local_size_2, int_t _num_kernels) : simd_width(_simd_width), local_size_0(_local_size_1), local_size_1(_local_size_2), num_kernels(_num_kernels)
|
2015-01-12 13:20:53 -05:00
|
|
|
{ }
|
|
|
|
|
|
|
|
|
2015-12-19 02:55:24 -05:00
|
|
|
bool base::requires_fallback(expression_tree const & expression)
|
2015-01-12 13:20:53 -05:00
|
|
|
{
|
2015-12-19 02:55:24 -05:00
|
|
|
for(expression_tree::node const & node: expression.tree())
|
2015-11-19 12:37:18 -05:00
|
|
|
if( (node.lhs.subtype==DENSE_ARRAY_TYPE && (node.lhs.array->stride()[0]>1 || node.lhs.array->start()>0))
|
|
|
|
|| (node.rhs.subtype==DENSE_ARRAY_TYPE && (node.rhs.array->stride()[0]>1 || node.rhs.array->start()>0)))
|
2015-09-30 15:31:41 -04:00
|
|
|
return true;
|
2015-01-12 13:20:53 -05:00
|
|
|
return false;
|
|
|
|
}
|
|
|
|
|
2015-12-19 02:55:24 -05:00
|
|
|
int_t base::vector_size(expression_tree::node const & node)
|
2015-01-12 13:20:53 -05:00
|
|
|
{
|
2015-12-19 02:04:39 -05:00
|
|
|
if (node.op.type==MATRIX_DIAG_TYPE)
|
2015-04-29 15:50:57 -04:00
|
|
|
return std::min<int_t>(node.lhs.array->shape()[0], node.lhs.array->shape()[1]);
|
2015-12-19 02:04:39 -05:00
|
|
|
else if (node.op.type==MATRIX_ROW_TYPE)
|
2015-04-29 15:50:57 -04:00
|
|
|
return node.lhs.array->shape()[1];
|
2015-12-19 02:04:39 -05:00
|
|
|
else if (node.op.type==MATRIX_COLUMN_TYPE)
|
2015-04-29 15:50:57 -04:00
|
|
|
return node.lhs.array->shape()[0];
|
2015-01-12 13:20:53 -05:00
|
|
|
else
|
2015-11-19 12:37:18 -05:00
|
|
|
return node.lhs.array->shape().max();
|
2015-01-12 13:20:53 -05:00
|
|
|
|
|
|
|
}
|
|
|
|
|
2015-12-19 02:55:24 -05:00
|
|
|
std::pair<int_t, int_t> base::matrix_size(expression_tree::container_type const & tree, expression_tree::node const & node)
|
2015-01-12 13:20:53 -05:00
|
|
|
{
|
2015-12-19 02:04:39 -05:00
|
|
|
if (node.op.type==VDIAG_TYPE)
|
2015-01-12 13:20:53 -05:00
|
|
|
{
|
2015-04-29 15:50:57 -04:00
|
|
|
int_t size = node.lhs.array->shape()[0];
|
2015-01-12 13:20:53 -05:00
|
|
|
return std::make_pair(size,size);
|
|
|
|
}
|
2015-12-19 02:04:39 -05:00
|
|
|
else if(node.op.type==REPEAT_TYPE)
|
2015-09-30 15:31:41 -04:00
|
|
|
{
|
|
|
|
size_t rep0 = tuple_get(tree, node.rhs.node_index, 0);
|
|
|
|
size_t rep1 = tuple_get(tree, node.rhs.node_index, 1);
|
|
|
|
std::cout << rep0 << " " << rep1 << std::endl;
|
|
|
|
return std::make_pair(node.lhs.array->shape()[0]*rep0, node.lhs.array->shape()[1]*rep1);
|
|
|
|
}
|
2015-01-12 13:20:53 -05:00
|
|
|
else
|
2015-04-29 15:50:57 -04:00
|
|
|
return std::make_pair(node.lhs.array->shape()[0],node.lhs.array->shape()[1]);
|
2015-01-12 13:20:53 -05:00
|
|
|
}
|
|
|
|
|
|
|
|
|
2015-01-17 10:48:02 -05:00
|
|
|
base::base(binding_policy_t binding_policy) : binding_policy_(binding_policy)
|
2015-01-12 13:20:53 -05:00
|
|
|
{}
|
|
|
|
|
2015-12-19 02:55:24 -05:00
|
|
|
unsigned int base::lmem_usage(expression_tree const &) const
|
2015-01-12 13:20:53 -05:00
|
|
|
{ return 0; }
|
|
|
|
|
2015-12-19 02:55:24 -05:00
|
|
|
unsigned int base::registers_usage(expression_tree const &) const
|
2015-01-12 13:20:53 -05:00
|
|
|
{ return 0; }
|
|
|
|
|
2015-12-19 02:55:24 -05:00
|
|
|
unsigned int base::temporary_workspace(expression_tree const &) const
|
2015-08-10 10:19:50 -07:00
|
|
|
{ return 0; }
|
|
|
|
|
2015-01-17 10:48:02 -05:00
|
|
|
base::~base()
|
2015-08-17 18:01:17 -07:00
|
|
|
{
|
|
|
|
}
|
2015-01-12 13:20:53 -05:00
|
|
|
|
2015-12-19 02:55:24 -05:00
|
|
|
std::string base::generate(std::string const & suffix, expression_tree const & expression, driver::Device const & device)
|
2015-01-12 13:20:53 -05:00
|
|
|
{
|
2015-09-30 15:31:41 -04:00
|
|
|
int err = is_invalid(expression, device);
|
2015-08-17 16:30:21 -07:00
|
|
|
if(err != 0)
|
|
|
|
throw operation_not_supported_exception("The supplied parameters for this template are invalid : err " + tools::to_string(err));
|
|
|
|
|
2015-01-12 13:20:53 -05:00
|
|
|
//Create mapping
|
2015-09-30 15:31:41 -04:00
|
|
|
mapping_type mapping;
|
2015-08-06 16:14:33 -07:00
|
|
|
std::unique_ptr<symbolic_binder> binder;
|
2015-09-30 15:31:41 -04:00
|
|
|
if (binding_policy_==BIND_SEQUENTIAL)
|
|
|
|
binder.reset(new bind_sequential());
|
2015-08-06 16:14:33 -07:00
|
|
|
else
|
2015-09-30 15:31:41 -04:00
|
|
|
binder.reset(new bind_independent());
|
2015-01-12 13:20:53 -05:00
|
|
|
|
2015-09-30 15:31:41 -04:00
|
|
|
traverse(expression, expression.root(), map_functor(*binder, mapping, device), true);
|
|
|
|
return generate_impl(suffix, expression, device, mapping);
|
2015-01-12 13:20:53 -05:00
|
|
|
}
|
|
|
|
|
|
|
|
template<class TType, class PType>
|
2015-12-19 02:55:24 -05:00
|
|
|
int base_impl<TType, PType>::is_invalid_impl(driver::Device const &, expression_tree const &) const
|
2015-01-12 13:20:53 -05:00
|
|
|
{ return TEMPLATE_VALID; }
|
|
|
|
|
|
|
|
template<class TType, class PType>
|
2015-01-17 10:48:02 -05:00
|
|
|
base_impl<TType, PType>::base_impl(parameters_type const & parameters, binding_policy_t binding_policy) : base(binding_policy), p_(parameters)
|
2015-01-12 13:20:53 -05:00
|
|
|
{ }
|
|
|
|
|
|
|
|
template<class TType, class PType>
|
2015-08-13 14:30:11 -07:00
|
|
|
unsigned int base_impl<TType, PType>::local_size_0() const
|
2015-01-12 13:20:53 -05:00
|
|
|
{ return p_.local_size_0; }
|
|
|
|
|
|
|
|
template<class TType, class PType>
|
2015-08-13 14:30:11 -07:00
|
|
|
unsigned int base_impl<TType, PType>::local_size_1() const
|
2015-01-12 13:20:53 -05:00
|
|
|
{ return p_.local_size_1; }
|
|
|
|
|
|
|
|
template<class TType, class PType>
|
2015-07-28 15:26:10 -07:00
|
|
|
std::shared_ptr<base> base_impl<TType, PType>::clone() const
|
|
|
|
{ return std::shared_ptr<base>(new TType(*dynamic_cast<TType const *>(this))); }
|
2015-01-12 13:20:53 -05:00
|
|
|
|
|
|
|
template<class TType, class PType>
|
2015-12-19 02:55:24 -05:00
|
|
|
int base_impl<TType, PType>::is_invalid(expression_tree const & expressions, driver::Device const & device) const
|
2015-01-12 13:20:53 -05:00
|
|
|
{
|
|
|
|
//Query device informations
|
2015-04-29 15:50:57 -04:00
|
|
|
size_t lmem_available = device.local_mem_size();
|
2015-02-01 22:28:49 -05:00
|
|
|
size_t lmem_used = lmem_usage(expressions);
|
2015-01-12 13:20:53 -05:00
|
|
|
if (lmem_used>lmem_available)
|
|
|
|
return TEMPLATE_LOCAL_MEMORY_OVERFLOW;
|
|
|
|
|
|
|
|
//Invalid work group size
|
2015-04-29 15:50:57 -04:00
|
|
|
size_t max_workgroup_size = device.max_work_group_size();
|
|
|
|
std::vector<size_t> max_work_item_sizes = device.max_work_item_sizes();
|
2015-01-12 13:20:53 -05:00
|
|
|
if (p_.local_size_0*p_.local_size_1 > max_workgroup_size)
|
|
|
|
return TEMPLATE_WORK_GROUP_SIZE_OVERFLOW;
|
|
|
|
if (p_.local_size_0 > max_work_item_sizes[0])
|
|
|
|
return TEMPLATE_LOCAL_SIZE_0_OVERFLOW;
|
|
|
|
|
|
|
|
if (p_.local_size_1 > max_work_item_sizes[1])
|
|
|
|
return TEMPLATE_LOCAL_SIZE_1_OVERFLOW;
|
|
|
|
|
|
|
|
//Invalid SIMD Width
|
2015-04-29 15:50:57 -04:00
|
|
|
if (p_.simd_width!=1 && p_.simd_width!=2 && p_.simd_width!=3 && p_.simd_width!=4)
|
2015-01-12 13:20:53 -05:00
|
|
|
return TEMPLATE_INVALID_SIMD_WIDTH;
|
|
|
|
|
2015-04-29 15:50:57 -04:00
|
|
|
return is_invalid_impl(device, expressions);
|
2015-01-12 13:20:53 -05:00
|
|
|
}
|
|
|
|
|
2015-12-12 18:32:06 -05:00
|
|
|
template class base_impl<elementwise_1d, elementwise_1d_parameters>;
|
|
|
|
template class base_impl<reduce_1d, reduce_1d_parameters>;
|
|
|
|
template class base_impl<elementwise_2d, elementwise_2d_parameters>;
|
|
|
|
template class base_impl<reduce_2d, reduce_2d_parameters>;
|
|
|
|
template class base_impl<matrix_product, matrix_product_parameters>;
|
2015-01-12 13:20:53 -05:00
|
|
|
|
|
|
|
}
|
2015-07-11 09:36:01 -04:00
|
|
|
}
|