132 lines
4.2 KiB
C++
132 lines
4.2 KiB
C++
/* Copyright 2015-2017 Philippe Tillet
|
|
*
|
|
* Permission is hereby granted, free of charge, to any person obtaining
|
|
* a copy of this software and associated documentation files
|
|
* (the "Software"), to deal in the Software without restriction,
|
|
* including without limitation the rights to use, copy, modify, merge,
|
|
* publish, distribute, sublicense, and/or sell copies of the Software,
|
|
* and to permit persons to whom the Software is furnished to do so,
|
|
* subject to the following conditions:
|
|
*
|
|
* The above copyright notice and this permission notice shall be
|
|
* included in all copies or substantial portions of the Software.
|
|
*
|
|
* THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND,
|
|
* EXPRESS OR IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF
|
|
* MERCHANTABILITY, FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.
|
|
* IN NO EVENT SHALL THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY
|
|
* CLAIM, DAMAGES OR OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT,
|
|
* TORT OR OTHERWISE, ARISING FROM, OUT OF OR IN CONNECTION WITH THE
|
|
* SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE.
|
|
*/
|
|
|
|
#include <cassert>
|
|
#include <algorithm>
|
|
#include <string>
|
|
|
|
#include "isaac/array.h"
|
|
#include "isaac/jit/generation/engine/keywords.h"
|
|
#include "isaac/jit/generation/elementwise_1d.h"
|
|
#include "isaac/jit/generation/reduce_1d.h"
|
|
#include "isaac/jit/generation/elementwise_2d.h"
|
|
#include "isaac/jit/generation/reduce_2d.h"
|
|
#include "isaac/jit/generation/gemm.h"
|
|
#include "isaac/jit/generation/base.h"
|
|
#include "isaac/exception/api.h"
|
|
#include "isaac/jit/syntax/engine/process.h"
|
|
#include "isaac/tools/cpp/string.hpp"
|
|
|
|
namespace isaac
|
|
{
|
|
namespace templates
|
|
{
|
|
|
|
base::base()
|
|
{}
|
|
|
|
unsigned int base::lmem_usage(expression_tree const &) const
|
|
{ return 0; }
|
|
|
|
unsigned int base::registers_usage(expression_tree const &) const
|
|
{ return 0; }
|
|
|
|
unsigned int base::temporary_workspace(expression_tree const &) const
|
|
{ return 0; }
|
|
|
|
base::~base()
|
|
{ }
|
|
|
|
std::string base::generate(std::string const & suffix, expression_tree const & expression, driver::Device const & device)
|
|
{
|
|
int err = is_invalid(expression, device);
|
|
if(err != 0)
|
|
throw operation_not_supported_exception("The supplied parameters for this template are invalid : err " + tools::to_string(err));
|
|
|
|
//Create mapping
|
|
symbolic::symbols_table mapping = symbolic::symbolize(expression);
|
|
return generate_impl(suffix, expression, device, mapping);
|
|
}
|
|
|
|
|
|
/* External base */
|
|
external_base::external_base()
|
|
{ }
|
|
|
|
std::string external_base::generate_impl(std::string const &, expression_tree const &, driver::Device const &, symbolic::symbols_table const &) const
|
|
{ return ""; }
|
|
|
|
unsigned int external_base::temporary_workspace(expression_tree const &) const
|
|
{ return 0; }
|
|
|
|
unsigned int external_base::lmem_usage(expression_tree const &) const
|
|
{ return 0; }
|
|
|
|
unsigned int external_base::registers_usage(expression_tree const &) const
|
|
{ return 0; }
|
|
|
|
/* Parameterized base */
|
|
int parameterized_base::is_invalid_impl(driver::Device const &, expression_tree const &) const
|
|
{ return TEMPLATE_VALID; }
|
|
|
|
parameterized_base::parameterized_base(unsigned int vwidth, int_t ls0, int_t ls1): vwidth_(vwidth), ls0_(ls0), ls1_(ls1)
|
|
{ }
|
|
|
|
unsigned int parameterized_base::ls0() const
|
|
{ return ls0_; }
|
|
|
|
unsigned int parameterized_base::ls1() const
|
|
{ return ls1_; }
|
|
|
|
int parameterized_base::is_invalid(expression_tree const & expressions, driver::Device const & device) const
|
|
{
|
|
//Query device informations
|
|
size_t lmem_available = device.local_mem_size();
|
|
size_t lmem_used = lmem_usage(expressions);
|
|
if (lmem_used>lmem_available)
|
|
return TEMPLATE_LOCAL_MEMORY_OVERFLOW;
|
|
|
|
//Invalid work group size
|
|
size_t max_workgroup_size = device.max_work_group_size();
|
|
std::vector<size_t> max_work_item_sizes = device.max_work_item_sizes();
|
|
|
|
if (ls0_*ls1_ > max_workgroup_size)
|
|
return TEMPLATE_WORK_GROUP_SIZE_OVERFLOW;
|
|
if (ls0_ > max_work_item_sizes[0])
|
|
return TEMPLATE_LOCAL_SIZE_0_OVERFLOW;
|
|
|
|
if (ls1_ > max_work_item_sizes[1])
|
|
return TEMPLATE_LOCAL_SIZE_1_OVERFLOW;
|
|
|
|
//Invalid SIMD Width
|
|
if (vwidth_!=1 && vwidth_!=2 && vwidth_!=3 && vwidth_!=4)
|
|
return TEMPLATE_INVALID_SIMD_WIDTH;
|
|
|
|
return is_invalid_impl(device, expressions);
|
|
}
|
|
|
|
std::shared_ptr<base> base::getptr()
|
|
{ return shared_from_this(); }
|
|
|
|
}
|
|
}
|