From 1e178dab227cbbc70332417e33a8cca53f3e83df Mon Sep 17 00:00:00 2001
From: Philippe Tillet <ptillet@g.harvard.edu>
Date: Sat, 2 Jul 2016 12:06:05 -0700
Subject: [PATCH] Code quality: shortened parameter names in JIT code generator

---
 include/isaac/jit/generation/base.h           |  14 +-
 include/isaac/jit/generation/elementwise_1d.h |   6 +-
 include/isaac/jit/generation/elementwise_2d.h |   6 +-
 include/isaac/jit/generation/matrix_product.h |  24 +--
 include/isaac/jit/generation/reduce_1d.h      |   8 +-
 include/isaac/jit/generation/reduce_2d.h      |  12 +-
 lib/jit/generation/base.cpp                   |  18 +-
 lib/jit/generation/elementwise_1d.cpp         |  30 +--
 lib/jit/generation/elementwise_2d.cpp         |  22 +--
 lib/jit/generation/matrix_product.cpp         | 185 +++++++++---------
 lib/jit/generation/reduce_1d.cpp              |  44 ++---
 lib/jit/generation/reduce_2d.cpp              |  38 ++--
 lib/jit/generation/tools/loop.hpp             |  10 +-
 lib/jit/generation/tools/vector_types.hpp     |  28 +--
 lib/runtime/inference/profiles.cpp            |   2 +-
 15 files changed, 222 insertions(+), 225 deletions(-)

diff --git a/include/isaac/jit/generation/base.h b/include/isaac/jit/generation/base.h
index e98799c9e..dab65f2b9 100644
--- a/include/isaac/jit/generation/base.h
+++ b/include/isaac/jit/generation/base.h
@@ -39,7 +39,7 @@ namespace isaac
 namespace templates
 {
 
-enum fetching_policy_type
+enum fetch_type
 {
   FETCH_FROM_LOCAL,
   FETCH_FROM_GLOBAL_STRIDED,
@@ -75,10 +75,10 @@ class base
 public:
   struct parameters_type
   {
-    parameters_type(unsigned int _simd_width, int_t _local_size_1, int_t _local_size_2, int_t _num_kernels);
-    unsigned int simd_width;
-    unsigned int local_size_0;
-    unsigned int local_size_1;
+    parameters_type(unsigned int _vwidth, int_t _ls0, int_t _ls1, int_t _num_kernels);
+    unsigned int vwidth;
+    unsigned int ls0;
+    unsigned int ls1;
     unsigned int num_kernels;
   };
 private:
@@ -107,8 +107,8 @@ private:
 public:
   typedef ParametersType parameters_type;
   base_impl(parameters_type const & parameters, fusion_policy_t fusion_policy);
-  unsigned int local_size_0() const;
-  unsigned int local_size_1() const;
+  unsigned int ls0() const;
+  unsigned int ls1() const;
   std::shared_ptr<base> clone() const;
   /** @brief returns whether or not the profile has undefined behavior on particular device */
   int is_invalid(expression_tree const & expressions, driver::Device const & device) const;
diff --git a/include/isaac/jit/generation/elementwise_1d.h b/include/isaac/jit/generation/elementwise_1d.h
index 0c3dc7ec1..ed183a54b 100644
--- a/include/isaac/jit/generation/elementwise_1d.h
+++ b/include/isaac/jit/generation/elementwise_1d.h
@@ -32,9 +32,9 @@ namespace templates
 class elementwise_1d_parameters : public base::parameters_type
 {
 public:
-  elementwise_1d_parameters(unsigned int _simd_width, unsigned int _group_size, unsigned int _num_groups, fetching_policy_type _fetching_policy);
+  elementwise_1d_parameters(unsigned int _vwidth, unsigned int _group_size, unsigned int _num_groups, fetch_type _fetch);
   unsigned int num_groups;
-  fetching_policy_type fetching_policy;
+  fetch_type fetch;
 };
 
 class elementwise_1d : public base_impl<elementwise_1d, elementwise_1d_parameters>
@@ -44,7 +44,7 @@ private:
   std::string generate_impl(std::string const & suffix, expression_tree const  & expressions, driver::Device const & device, symbolic::symbols_table const & symbols) const;
 public:
   elementwise_1d(elementwise_1d::parameters_type const & parameters, fusion_policy_t fusion_policy = FUSE_INDEPENDENT);
-  elementwise_1d(unsigned int _simd_width, unsigned int _group_size, unsigned int _num_groups, fetching_policy_type _fetching_policy, fusion_policy_t fusion_policy = FUSE_INDEPENDENT);
+  elementwise_1d(unsigned int _vwidth, unsigned int _group_size, unsigned int _num_groups, fetch_type _fetch, fusion_policy_t fusion_policy = FUSE_INDEPENDENT);
   std::vector<int_t> input_sizes(expression_tree const  & expressions) const;
   void enqueue(driver::CommandQueue & queue, driver::Program const & program, std::string const & suffix, runtime::execution_handler const &);
 };
diff --git a/include/isaac/jit/generation/elementwise_2d.h b/include/isaac/jit/generation/elementwise_2d.h
index 43599da5e..966f1ec0d 100644
--- a/include/isaac/jit/generation/elementwise_2d.h
+++ b/include/isaac/jit/generation/elementwise_2d.h
@@ -33,11 +33,11 @@ namespace templates
 class elementwise_2d_parameters : public base::parameters_type
 {
 public:
-  elementwise_2d_parameters(unsigned int _simd_width, unsigned int _local_size_0, unsigned int _local_size_1, unsigned int _num_groups_0, unsigned int _num_groups_1, fetching_policy_type _fetching_policy);
+  elementwise_2d_parameters(unsigned int _vwidth, unsigned int _ls0, unsigned int _ls1, unsigned int _num_groups_0, unsigned int _num_groups_1, fetch_type _fetch);
 
   unsigned int num_groups_0;
   unsigned int num_groups_1;
-  fetching_policy_type fetching_policy;
+  fetch_type fetch;
 };
 
 class elementwise_2d : public base_impl<elementwise_2d, elementwise_2d_parameters>
@@ -47,7 +47,7 @@ private:
   std::string generate_impl(std::string const & suffix, expression_tree const  & expressions, driver::Device const & device, symbolic::symbols_table const & mapping) const;
 public:
   elementwise_2d(parameters_type const & parameters, fusion_policy_t fusion_policy = FUSE_INDEPENDENT);
-  elementwise_2d(unsigned int simd, unsigned int ls1, unsigned int ls2,  unsigned int ng1, unsigned int ng2, fetching_policy_type fetch, fusion_policy_t bind = FUSE_INDEPENDENT);
+  elementwise_2d(unsigned int simd, unsigned int ls1, unsigned int ls2,  unsigned int ng1, unsigned int ng2, fetch_type fetch, fusion_policy_t bind = FUSE_INDEPENDENT);
   std::vector<int_t> input_sizes(expression_tree const  & expressions) const;
   void enqueue(driver::CommandQueue & queue, driver::Program const & program, std::string const & suffix, runtime::execution_handler const &);
 };
diff --git a/include/isaac/jit/generation/matrix_product.h b/include/isaac/jit/generation/matrix_product.h
index a59b3943e..a9d179940 100644
--- a/include/isaac/jit/generation/matrix_product.h
+++ b/include/isaac/jit/generation/matrix_product.h
@@ -33,11 +33,11 @@ namespace templates
 
 struct matrix_product_parameters : public base::parameters_type
 {
-  matrix_product_parameters(unsigned int simd_width
-                            , unsigned int local_size_0, unsigned int KL, unsigned int local_size_1, unsigned int D
+  matrix_product_parameters(unsigned int vwidth
+                            , unsigned int ls0, unsigned int KL, unsigned int ls1, unsigned int D
                             , unsigned int ms, unsigned int ks, unsigned int ns
-                            , fetching_policy_type A_fetching_policy, fetching_policy_type B_fetching_policy
-                            , unsigned int local_fetch_0, unsigned int local_fetch_1);
+                            , fetch_type Afetch, fetch_type Bfetch
+                            , unsigned int lf0, unsigned int lf1);
 
   unsigned int kL;
   unsigned int depth;
@@ -46,11 +46,11 @@ struct matrix_product_parameters : public base::parameters_type
   unsigned int kS;
   unsigned int nS;
 
-  fetching_policy_type A_fetching_policy;
-  fetching_policy_type B_fetching_policy;
+  fetch_type Afetch;
+  fetch_type Bfetch;
 
-  unsigned int local_fetch_0;
-  unsigned int local_fetch_1;
+  unsigned int lf0;
+  unsigned int lf1;
 
   unsigned int mL;
   unsigned int nL;
@@ -84,7 +84,7 @@ class matrix_product_nn : public matrix_product
 {
 public:
   matrix_product_nn(unsigned int simd, int_t ls0, int_t KL, int_t ls1, int_t D
-                      , int_t ms, int_t ks, int_t ns, fetching_policy_type Afetch , fetching_policy_type Bfetch
+                      , int_t ms, int_t ks, int_t ns, fetch_type Afetch , fetch_type Bfetch
                       , int_t lfetch0, int_t lfetch1);
 };
 
@@ -92,7 +92,7 @@ class matrix_product_tn : public matrix_product
 {
 public:
   matrix_product_tn(unsigned int simd, int_t ls0, int_t KL, int_t ls1, int_t D
-                      , int_t ms, int_t ks, int_t ns, fetching_policy_type Afetch , fetching_policy_type Bfetch
+                      , int_t ms, int_t ks, int_t ns, fetch_type Afetch , fetch_type Bfetch
                       , int_t lfetch0, int_t lfetch1);
 };
 
@@ -101,7 +101,7 @@ class matrix_product_nt : public matrix_product
 {
 public:
   matrix_product_nt(unsigned int simd, int_t ls0, int_t KL, int_t ls1, int_t D
-                      , int_t ms, int_t ks, int_t ns, fetching_policy_type Afetch , fetching_policy_type Bfetch
+                      , int_t ms, int_t ks, int_t ns, fetch_type Afetch , fetch_type Bfetch
                       , int_t lfetch0, int_t lfetch1);
 };
 
@@ -110,7 +110,7 @@ class matrix_product_tt : public matrix_product
 {
 public:
   matrix_product_tt(unsigned int simd, int_t ls0, int_t KL, int_t ls1, int_t D
-                      , int_t ms, int_t ks, int_t ns, fetching_policy_type Afetch , fetching_policy_type Bfetch
+                      , int_t ms, int_t ks, int_t ns, fetch_type Afetch , fetch_type Bfetch
                       , int_t lfetch0, int_t lfetch1);
 };
 
diff --git a/include/isaac/jit/generation/reduce_1d.h b/include/isaac/jit/generation/reduce_1d.h
index 11d1b820d..60c24360b 100644
--- a/include/isaac/jit/generation/reduce_1d.h
+++ b/include/isaac/jit/generation/reduce_1d.h
@@ -31,11 +31,11 @@ namespace templates
 
 struct reduce_1d_parameters : public base::parameters_type
 {
-  reduce_1d_parameters(unsigned int _simd_width,
+  reduce_1d_parameters(unsigned int _vwidth,
                        unsigned int _group_size, unsigned int _num_groups,
-                       fetching_policy_type _fetching_policy);
+                       fetch_type _fetch);
   unsigned int num_groups;
-  fetching_policy_type fetching_policy;
+  fetch_type fetch;
 };
 
 class reduce_1d : public base_impl<reduce_1d, reduce_1d_parameters>
@@ -50,7 +50,7 @@ private:
 
 public:
   reduce_1d(reduce_1d::parameters_type const & parameters, fusion_policy_t fusion_policy = FUSE_INDEPENDENT);
-  reduce_1d(unsigned int simd, unsigned int ls, unsigned int ng, fetching_policy_type fetch, fusion_policy_t bind = FUSE_INDEPENDENT);
+  reduce_1d(unsigned int simd, unsigned int ls, unsigned int ng, fetch_type fetch, fusion_policy_t bind = FUSE_INDEPENDENT);
   std::vector<int_t> input_sizes(expression_tree const  & expressions) const;
   void enqueue(driver::CommandQueue & queue, driver::Program const & program, std::string const & suffix, runtime::execution_handler const &);
 private:
diff --git a/include/isaac/jit/generation/reduce_2d.h b/include/isaac/jit/generation/reduce_2d.h
index 90bdfc08b..4fb28346a 100644
--- a/include/isaac/jit/generation/reduce_2d.h
+++ b/include/isaac/jit/generation/reduce_2d.h
@@ -33,12 +33,12 @@ namespace templates
 {
 struct reduce_2d_parameters : public base::parameters_type
 {
-  reduce_2d_parameters(unsigned int _simd_width,
-                                unsigned int _local_size_0, unsigned int _local_size_1,
-                                unsigned int _num_groups_0, unsigned int _num_groups_1, fetching_policy_type _fetch_policy);
+  reduce_2d_parameters(unsigned int _vwidth,
+                                unsigned int _ls0, unsigned int _ls1,
+                                unsigned int _num_groups_0, unsigned int _num_groups_1, fetch_type _fetch_policy);
   unsigned int num_groups_0;
   unsigned int num_groups_1;
-  fetching_policy_type fetch_policy;
+  fetch_type fetch_policy;
 };
 
 
@@ -62,14 +62,14 @@ class reduce_2d_rows : public reduce_2d
 {
 public:
   reduce_2d_rows(reduce_2d::parameters_type  const &, fusion_policy_t fusion_policy = FUSE_INDEPENDENT);
-  reduce_2d_rows(unsigned int simd, unsigned int ls1, unsigned int ls2, unsigned int ng1, unsigned int ng2, fetching_policy_type fetch, fusion_policy_t bind = FUSE_INDEPENDENT);
+  reduce_2d_rows(unsigned int simd, unsigned int ls1, unsigned int ls2, unsigned int ng1, unsigned int ng2, fetch_type fetch, fusion_policy_t bind = FUSE_INDEPENDENT);
 };
 
 class reduce_2d_cols : public reduce_2d
 {
 public:
   reduce_2d_cols(reduce_2d::parameters_type  const &, fusion_policy_t fusion_policy = FUSE_INDEPENDENT);
-  reduce_2d_cols(unsigned int simd, unsigned int ls1, unsigned int ls2, unsigned int ng1, unsigned int ng2, fetching_policy_type fetch, fusion_policy_t bind = FUSE_INDEPENDENT);
+  reduce_2d_cols(unsigned int simd, unsigned int ls1, unsigned int ls2, unsigned int ng1, unsigned int ng2, fetch_type fetch, fusion_policy_t bind = FUSE_INDEPENDENT);
 };
 
 }
diff --git a/lib/jit/generation/base.cpp b/lib/jit/generation/base.cpp
index 6873cbd27..12ea14a4e 100644
--- a/lib/jit/generation/base.cpp
+++ b/lib/jit/generation/base.cpp
@@ -40,7 +40,7 @@ namespace isaac
 namespace templates
 {
 
-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)
+base::parameters_type::parameters_type(unsigned int _vwidth, int_t _ls0, int_t _ls1, int_t _num_kernels) : vwidth(_vwidth), ls0(_ls0), ls1(_ls1), num_kernels(_num_kernels)
 { }
 
 base::base(fusion_policy_t fusion_policy) : fusion_policy_(fusion_policy)
@@ -79,12 +79,12 @@ base_impl<TType, PType>::base_impl(parameters_type const & parameters, fusion_po
 { }
 
 template<class TType, class PType>
-unsigned int base_impl<TType, PType>::local_size_0() const
-{ return p_.local_size_0; }
+unsigned int base_impl<TType, PType>::ls0() const
+{ return p_.ls0; }
 
 template<class TType, class PType>
-unsigned int base_impl<TType, PType>::local_size_1() const
-{ return p_.local_size_1; }
+unsigned int base_impl<TType, PType>::ls1() const
+{ return p_.ls1; }
 
 template<class TType, class PType>
 std::shared_ptr<base> base_impl<TType, PType>::clone() const
@@ -102,16 +102,16 @@ int base_impl<TType, PType>::is_invalid(expression_tree const  & expressions, dr
   //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 (p_.local_size_0*p_.local_size_1 > max_workgroup_size)
+  if (p_.ls0*p_.ls1 > max_workgroup_size)
     return TEMPLATE_WORK_GROUP_SIZE_OVERFLOW;
-  if (p_.local_size_0 > max_work_item_sizes[0])
+  if (p_.ls0 > max_work_item_sizes[0])
     return TEMPLATE_LOCAL_SIZE_0_OVERFLOW;
 
-  if (p_.local_size_1 > max_work_item_sizes[1])
+  if (p_.ls1 > max_work_item_sizes[1])
     return TEMPLATE_LOCAL_SIZE_1_OVERFLOW;
 
   //Invalid SIMD Width
-  if (p_.simd_width!=1 && p_.simd_width!=2 && p_.simd_width!=3 && p_.simd_width!=4)
+  if (p_.vwidth!=1 && p_.vwidth!=2 && p_.vwidth!=3 && p_.vwidth!=4)
     return TEMPLATE_INVALID_SIMD_WIDTH;
 
   return is_invalid_impl(device, expressions);
diff --git a/lib/jit/generation/elementwise_1d.cpp b/lib/jit/generation/elementwise_1d.cpp
index bd74d72b6..db9633b32 100644
--- a/lib/jit/generation/elementwise_1d.cpp
+++ b/lib/jit/generation/elementwise_1d.cpp
@@ -36,17 +36,17 @@ namespace isaac
 namespace templates
 {
 
-elementwise_1d_parameters::elementwise_1d_parameters(unsigned int _simd_width,
+elementwise_1d_parameters::elementwise_1d_parameters(unsigned int _vwidth,
                        unsigned int _group_size, unsigned int _num_groups,
-                       fetching_policy_type _fetching_policy) :
-      base::parameters_type(_simd_width, _group_size, 1, 1), num_groups(_num_groups), fetching_policy(_fetching_policy)
+                       fetch_type _fetch) :
+      base::parameters_type(_vwidth, _group_size, 1, 1), num_groups(_num_groups), fetch(_fetch)
 {
 }
 
 
 int elementwise_1d::is_invalid_impl(driver::Device const &, expression_tree const &) const
 {
-  if (p_.fetching_policy==FETCH_FROM_LOCAL)
+  if (p_.fetch==FETCH_FROM_LOCAL)
     return TEMPLATE_INVALID_FETCHING_POLICY_TYPE;
   return TEMPLATE_VALID;
 }
@@ -65,7 +65,7 @@ std::string elementwise_1d::generate_impl(std::string const & suffix, expression
     case driver::CUDA:
       stream << "#include  \"vector.h\"" << std::endl; break;
     case driver::OPENCL:
-      stream << " __attribute__((reqd_work_group_size(" << p_.local_size_0 << "," << p_.local_size_1 << ",1)))" << std::endl; break;
+      stream << " __attribute__((reqd_work_group_size(" << p_.ls0 << "," << p_.ls1 << ",1)))" << std::endl; break;
   }
 
   stream << "$KERNEL void elementwise_1d" << suffix << "($SIZE_T N, " << tools::join(kernel_arguments(device, symbols, tree), ", ") << ")";
@@ -83,9 +83,9 @@ std::string elementwise_1d::generate_impl(std::string const & suffix, expression
     stream.inc_tab();
   }
 
-  element_wise_loop_1D(stream, p_.fetching_policy, p_.simd_width, "i", "N", "$GLOBAL_IDX_0", "$GLOBAL_SIZE_0", device, [&](unsigned int simd_width)
+  element_wise_loop_1D(stream, p_.fetch, p_.vwidth, "i", "N", "$GLOBAL_IDX_0", "$GLOBAL_SIZE_0", device, [&](unsigned int vwidth)
   {
-    std::string dtype = append_width("#scalartype",simd_width);
+    std::string dtype = append_width("#scalartype",vwidth);
 
     //Declares register to store results
     for(symbolic::leaf* sym: symbolic::extract<symbolic::leaf>(tree, symbols, assignments_lhs, false))
@@ -93,17 +93,17 @@ std::string elementwise_1d::generate_impl(std::string const & suffix, expression
 
     //Load to registers
     for(symbolic::leaf* sym: symbolic::extract<symbolic::leaf>(tree, symbols, assignments_rhs, false))
-      stream << sym->process(dtype + " #name = " + append_width("loadv", simd_width) + "(i);") << std::endl;
+      stream << sym->process(dtype + " #name = " + append_width("loadv", vwidth) + "(i);") << std::endl;
 
     //Compute
     for(size_t idx: assignments)
-      for(unsigned int s = 0 ; s < simd_width ; ++s)
-         stream << symbols.at(idx)->evaluate({{"leaf", access_vector_type("#name", s, simd_width)}}) << ";" << std::endl;
+      for(unsigned int s = 0 ; s < vwidth ; ++s)
+         stream << symbols.at(idx)->evaluate({{"leaf", access_vector_type("#name", s, vwidth)}}) << ";" << std::endl;
 
     //Writes back
     for(symbolic::leaf* sym: symbolic::extract<symbolic::leaf>(tree, symbols, assignments_lhs, false))
-      for(unsigned int s = 0 ; s < simd_width ; ++s)
-          stream << sym->process("at(i+" + tools::to_string(s)+") = " + access_vector_type("#name", s, simd_width) + ";") << std::endl;
+      for(unsigned int s = 0 ; s < vwidth ; ++s)
+          stream << sym->process("at(i+" + tools::to_string(s)+") = " + access_vector_type("#name", s, vwidth) + ";") << std::endl;
   });
   //Close user-provided for-loops
   if(sfors.size()){
@@ -124,7 +124,7 @@ elementwise_1d::elementwise_1d(elementwise_1d_parameters const & parameters,
 {}
 
 elementwise_1d::elementwise_1d(unsigned int simd, unsigned int ls, unsigned int ng,
-                               fetching_policy_type fetch, fusion_policy_t bind):
+                               fetch_type fetch, fusion_policy_t bind):
     base_impl<elementwise_1d, elementwise_1d_parameters>(elementwise_1d_parameters(simd,ls,ng,fetch), bind)
 {}
 
@@ -144,8 +144,8 @@ void elementwise_1d::enqueue(driver::CommandQueue &, driver::Program const & pro
   name += suffix;
   driver::Kernel kernel(program, name.c_str());
   //NDRange
-  driver::NDRange global(p_.local_size_0*p_.num_groups);
-  driver::NDRange local(p_.local_size_0);
+  driver::NDRange global(p_.ls0*p_.num_groups);
+  driver::NDRange local(p_.ls0);
   //Arguments
   unsigned int current_arg = 0;
   kernel.setSizeArg(current_arg++, size);
diff --git a/lib/jit/generation/elementwise_2d.cpp b/lib/jit/generation/elementwise_2d.cpp
index 80c463394..db02451b9 100644
--- a/lib/jit/generation/elementwise_2d.cpp
+++ b/lib/jit/generation/elementwise_2d.cpp
@@ -33,18 +33,18 @@ namespace isaac
 namespace templates
 {
 
-elementwise_2d_parameters::elementwise_2d_parameters(unsigned int _simd_width,
-                          unsigned int _local_size_0, unsigned int _local_size_1,
+elementwise_2d_parameters::elementwise_2d_parameters(unsigned int _vwidth,
+                          unsigned int _ls0, unsigned int _ls1,
                           unsigned int _num_groups_0, unsigned int _num_groups_1,
-                          fetching_policy_type _fetching_policy) : base::parameters_type(_simd_width, _local_size_0, _local_size_1, 1), num_groups_0(_num_groups_0), num_groups_1(_num_groups_1), fetching_policy(_fetching_policy){ }
+                          fetch_type _fetch) : base::parameters_type(_vwidth, _ls0, _ls1, 1), num_groups_0(_num_groups_0), num_groups_1(_num_groups_1), fetch(_fetch){ }
 
 
 
 int elementwise_2d::is_invalid_impl(driver::Device const &, expression_tree const  &) const
 {
-  if (p_.simd_width>1)
+  if (p_.vwidth>1)
     return TEMPLATE_INVALID_SIMD_WIDTH;
-  if(p_.fetching_policy==FETCH_FROM_LOCAL)
+  if(p_.fetch==FETCH_FROM_LOCAL)
     return TEMPLATE_INVALID_FETCHING_POLICY_TYPE;
   return TEMPLATE_VALID;
 }
@@ -67,7 +67,7 @@ std::string elementwise_2d::generate_impl(std::string const & suffix, expression
     case driver::CUDA:
       stream << "#include  \"vector.h\"" << std::endl; break;
     case driver::OPENCL:
-      stream << " __attribute__((reqd_work_group_size(" << p_.local_size_0 << "," << p_.local_size_1 << ",1)))" << std::endl; break;
+      stream << " __attribute__((reqd_work_group_size(" << p_.ls0 << "," << p_.ls1 << ",1)))" << std::endl; break;
   }
 
   stream << "$KERNEL void elementwise_2d" << suffix << "($SIZE_T M, $SIZE_T N, " << tools::join(kernel_arguments(device, symbols, tree), ", ") << ")" << std::endl;
@@ -75,11 +75,11 @@ std::string elementwise_2d::generate_impl(std::string const & suffix, expression
   stream.inc_tab();
 
 
-  fetching_loop_info(p_.fetching_policy, "M", stream, init0, upper_bound0, inc0,  "$GLOBAL_IDX_0", "$GLOBAL_SIZE_0", device);
+  fetching_loop_info(p_.fetch, "M", stream, init0, upper_bound0, inc0,  "$GLOBAL_IDX_0", "$GLOBAL_SIZE_0", device);
   stream << "for($SIZE_T i = " << init0 << "; i < " << upper_bound0 << "; i += " << inc0 << ")" << std::endl;
   stream << "{" << std::endl;
   stream.inc_tab();
-  fetching_loop_info(p_.fetching_policy, "N", stream, init1, upper_bound1, inc1, "$GLOBAL_IDX_1", "$GLOBAL_SIZE_1", device);
+  fetching_loop_info(p_.fetch, "N", stream, init1, upper_bound1, inc1, "$GLOBAL_IDX_1", "$GLOBAL_SIZE_1", device);
   stream << "for($SIZE_T j = " << init1 << "; j < " << upper_bound1 << "; j += " << inc1 << ")" << std::endl;
   stream << "{" << std::endl;
   stream.inc_tab();
@@ -115,7 +115,7 @@ elementwise_2d::elementwise_2d(parameters_type const & parameters, fusion_policy
   base_impl<elementwise_2d, elementwise_2d_parameters>(parameters, fusion_policy){ }
 
 elementwise_2d::elementwise_2d(unsigned int simd, unsigned int ls1, unsigned int ls2,
-                               unsigned int ng1, unsigned int ng2, fetching_policy_type fetch,
+                               unsigned int ng1, unsigned int ng2, fetch_type fetch,
                                fusion_policy_t bind):
     base_impl<elementwise_2d, elementwise_2d_parameters>(elementwise_2d_parameters(simd, ls1, ls2, ng1, ng2, fetch), bind)
 {}
@@ -130,8 +130,8 @@ void elementwise_2d::enqueue(driver::CommandQueue & /*queue*/, driver::Program c
   std::string name = "elementwise_2d";
   name +=suffix;
   driver::Kernel kernel(program, name.c_str());
-  driver::NDRange global(p_.local_size_0*p_.num_groups_0, p_.local_size_1*p_.num_groups_1);
-  driver::NDRange local(p_.local_size_0, p_.local_size_1);
+  driver::NDRange global(p_.ls0*p_.num_groups_0, p_.ls1*p_.num_groups_1);
+  driver::NDRange local(p_.ls0, p_.ls1);
   unsigned int current_arg = 0;
   std::vector<int_t> MN = input_sizes(expressions);
   kernel.setSizeArg(current_arg++, MN[0]);
diff --git a/lib/jit/generation/matrix_product.cpp b/lib/jit/generation/matrix_product.cpp
index 20e894f67..bbb310a0c 100644
--- a/lib/jit/generation/matrix_product.cpp
+++ b/lib/jit/generation/matrix_product.cpp
@@ -37,14 +37,14 @@ namespace isaac
 namespace templates
 {
 
-matrix_product_parameters::matrix_product_parameters(unsigned int simd_width
-                          , unsigned int local_size_0, unsigned int KL, unsigned int local_size_1, unsigned int D
+matrix_product_parameters::matrix_product_parameters(unsigned int vwidth
+                          , unsigned int ls0, unsigned int KL, unsigned int ls1, unsigned int D
                           , unsigned int ms, unsigned int ks, unsigned int ns
-                          , fetching_policy_type A_fetching_policy, fetching_policy_type B_fetching_policy
-                          , unsigned int local_fetch_0, unsigned int local_fetch_1): base::parameters_type(simd_width, local_size_0, local_size_1, 1),
-  kL(KL), depth(D), mS(ms), kS(ks), nS(ns), A_fetching_policy(A_fetching_policy), B_fetching_policy(B_fetching_policy),
-  local_fetch_0(local_fetch_0), local_fetch_1(local_fetch_1),
-  mL(ms*local_size_0), nL(ns*local_size_1)
+                          , fetch_type Afetch, fetch_type Bfetch
+                          , unsigned int lf0, unsigned int lf1): base::parameters_type(vwidth, ls0, ls1, 1),
+  kL(KL), depth(D), mS(ms), kS(ks), nS(ns), Afetch(Afetch), Bfetch(Bfetch),
+  lf0(lf0), lf1(lf1),
+  mL(ms*ls0), nL(ns*ls1)
 {
 }
 
@@ -74,13 +74,10 @@ matrix_product_parameters::matrix_product_parameters(unsigned int simd_width
 
   int matrix_product::is_invalid_impl(driver::Device const &, expression_tree const &) const
   {
-//    if(device.vendor()==driver::Device::Vendor::NVIDIA && p_.simd_width > 1)
-//      return TEMPLATE_INVALID_SIMD_WIDTH;
-
-    if(p_.A_fetching_policy!=FETCH_FROM_LOCAL || p_.B_fetching_policy!=FETCH_FROM_LOCAL)
+    if(p_.Afetch!=FETCH_FROM_LOCAL || p_.Bfetch!=FETCH_FROM_LOCAL)
       return TEMPLATE_INVALID_FETCHING_POLICY_TYPE;
 
-    if ((p_.mS % p_.simd_width) > 0 || (p_.nS % p_.simd_width) > 0)
+    if ((p_.mS % p_.vwidth) > 0 || (p_.nS % p_.vwidth) > 0)
       return TEMPLATE_MS_NS_MUST_BE_SIMD_WIDTH_MULTIPLE;
 
     if(p_.mL > 256 || p_.nL > 256)
@@ -89,32 +86,32 @@ matrix_product_parameters::matrix_product_parameters(unsigned int simd_width
     if ( p_.kS % p_.kL == 0)
       return TEMPLATE_KS_MUST_BE_SMALLER_THAN_KL;
 
-    if (p_.A_fetching_policy==FETCH_FROM_LOCAL || p_.B_fetching_policy==FETCH_FROM_LOCAL){
-      if ((p_.local_fetch_0*p_.local_fetch_1) !=(p_.local_size_0*p_.local_size_1))
+    if (p_.Afetch==FETCH_FROM_LOCAL || p_.Bfetch==FETCH_FROM_LOCAL){
+      if ((p_.lf0*p_.lf1) !=(p_.ls0*p_.ls1))
         return TEMPLATE_LOCAL_FETCH_PRODUCT_MUST_MATCH_LOCAL_SIZE_PRODUCT;
     }
 
-    if (p_.A_fetching_policy==FETCH_FROM_LOCAL)
+    if (p_.Afetch==FETCH_FROM_LOCAL)
     {
       unsigned int bound1 = (A_trans_=='N')?p_.kL:p_.mL;
       unsigned int bound0 = (A_trans_=='N')?p_.mL:p_.kL;
 
-      if (p_.local_fetch_1>0 && (bound1 % p_.local_fetch_1)> 0)
+      if (p_.lf1>0 && (bound1 % p_.lf1)> 0)
         return A_trans_=='N'?TEMPLATE_LOCAL_FETCH_1_MUST_BE_KL_MULTIPLE:TEMPLATE_LOCAL_FETCH_1_MUST_BE_ML_MULTIPLE;
 
-      if (p_.local_fetch_0>0 && (bound0 % (p_.local_fetch_0*p_.simd_width)) > 0)
+      if (p_.lf0>0 && (bound0 % (p_.lf0*p_.vwidth)) > 0)
         return A_trans_=='N'?TEMPLATE_LOCAL_FETCH_0_MUST_BE_NL_MULTIPLE:TEMPLATE_LOCAL_FETCH_0_MUST_BE_KL_MULTIPLE;
 
     }
-    if (p_.B_fetching_policy==FETCH_FROM_LOCAL)
+    if (p_.Bfetch==FETCH_FROM_LOCAL)
     {
       unsigned int bound1 = (B_trans_=='T')?p_.kL:p_.nL;
       unsigned int bound0 = (B_trans_=='T')?p_.nL:p_.kL;
 
-      if (p_.local_fetch_1>0 && (bound1 % p_.local_fetch_1)> 0)
+      if (p_.lf1>0 && (bound1 % p_.lf1)> 0)
         return B_trans_=='T'?TEMPLATE_LOCAL_FETCH_1_MUST_BE_KL_MULTIPLE:TEMPLATE_LOCAL_FETCH_1_MUST_BE_ML_MULTIPLE;
 
-      if (p_.local_fetch_0>0 && (bound0 % (p_.local_fetch_0*p_.simd_width)) > 0)
+      if (p_.lf0>0 && (bound0 % (p_.lf0*p_.vwidth)) > 0)
         return B_trans_=='T'?TEMPLATE_LOCAL_FETCH_1_MUST_BE_KL_MULTIPLE:TEMPLATE_LOCAL_FETCH_1_MUST_BE_ML_MULTIPLE;
 
     }
@@ -129,9 +126,9 @@ matrix_product_parameters::matrix_product_parameters(unsigned int simd_width
 
     driver::backend_type backend = device.backend();
     bool has_depth = p_.depth > 1;
-#define VLOAD(offset, ptr) vload(p_.simd_width, sdtype, offset, ptr, "1", backend, true)
-#define VLOAD_MISALIGNED(offset, ptr) vload(p_.simd_width, sdtype, offset, ptr, "1", backend, false)
-#define VSTORE(value, offset, ptr) vstore(p_.simd_width, sdtype, value, offset, ptr, "1", backend)
+#define VLOAD(offset, ptr) vload(p_.vwidth, sdtype, offset, ptr, "1", backend, true)
+#define VLOAD_MISALIGNED(offset, ptr) vload(p_.vwidth, sdtype, offset, ptr, "1", backend, false)
+#define VSTORE(value, offset, ptr) vstore(p_.vwidth, sdtype, value, offset, ptr, "1", backend)
 
     symbolic::preset::matrix_product::args args;
     infos(tree, args);
@@ -145,7 +142,7 @@ matrix_product_parameters::matrix_product_parameters(unsigned int simd_width
     kernel_generation_stream stream(backend);
     numeric_type dtype = tree.dtype();
     std::string sdtype = to_string(dtype);
-    std::string vdtype = append_width(sdtype, p_.simd_width);
+    std::string vdtype = append_width(sdtype, p_.vwidth);
 
     //////////////////
     /// DECLARATIONS
@@ -159,7 +156,7 @@ matrix_product_parameters::matrix_product_parameters(unsigned int simd_width
     switch(backend)
     {
       case driver::OPENCL:
-        stream << " __attribute__((reqd_work_group_size(" << p_.local_size_0 << "," << p_.local_size_1 << ",1)))" << std::endl;
+        stream << " __attribute__((reqd_work_group_size(" << p_.ls0 << "," << p_.ls1 << ",1)))" << std::endl;
         break;
       default:
         break;
@@ -178,8 +175,8 @@ matrix_product_parameters::matrix_product_parameters(unsigned int simd_width
     ///Declare
     stream << "//blocks" << std::endl;
     stream << sdtype << " rC[" << p_.mS << "][" << p_.nS << "] = {{0}};" << std::endl;
-    stream << vdtype << " rA[" << p_.kS << "][" << p_.mS/p_.simd_width << "];" << std::endl;
-    stream << vdtype << " rB[" << p_.kS << "][" << p_.nS/p_.simd_width << "];" << std::endl;
+    stream << vdtype << " rA[" << p_.kS << "][" << p_.mS/p_.vwidth << "];" << std::endl;
+    stream << vdtype << " rB[" << p_.kS << "][" << p_.nS/p_.vwidth << "];" << std::endl;
     stream << std::endl;
 
     stream << "//pointers" << std::endl;
@@ -187,8 +184,8 @@ matrix_product_parameters::matrix_product_parameters(unsigned int simd_width
     size_t lldb = (B_trans_=='T')?p_.nL:p_.kL;
     stream << "$LOCAL " << sdtype << " lA[" << p_.kL*p_.mL << "];" << std::endl;
     stream << "$LOCAL " << sdtype << " lB[" << p_.kL*p_.nL << "];" << std::endl;
-    unsigned int npA = p_.mL/(A_trans_=='N'?p_.local_fetch_0*p_.simd_width:p_.local_fetch_1);
-    unsigned int npB = p_.nL/(B_trans_=='T'?p_.local_fetch_0*p_.simd_width:p_.local_fetch_1);
+    unsigned int npA = p_.mL/(A_trans_=='N'?p_.lf0*p_.vwidth:p_.lf1);
+    unsigned int npB = p_.nL/(B_trans_=='T'?p_.lf0*p_.vwidth:p_.lf1);
     stream << "$GLOBAL " << sdtype << "* Ai[" << npA << "];" << std::endl;
     stream << "$GLOBAL " << sdtype << "* Bi[" << npB << "];" << std::endl;
     stream << std::endl;
@@ -218,15 +215,15 @@ matrix_product_parameters::matrix_product_parameters(unsigned int simd_width
       stream << "K = min(K - div*gidz, ($SIZE_T)div);" << std::endl;
     }
 
-    stream << "idt = " << p_.local_size_0 << "*ids.w + ids.z;" << std::endl;
-    stream << "idT.y = idt/" << p_.local_fetch_0 << ";" << std::endl;
-    stream << "idT.x = idt - " << p_.local_fetch_0 << "*idT.y;" << std::endl;
+    stream << "idt = " << p_.ls0 << "*ids.w + ids.z;" << std::endl;
+    stream << "idT.y = idt/" << p_.lf0 << ";" << std::endl;
+    stream << "idT.x = idt - " << p_.lf0 << "*idT.y;" << std::endl;
     stream << std::endl;
 
     stream << "//Adjust pointers and bounds per work-item" << std::endl;
     stream << "ids.x *= " << p_.mL << ";" << std::endl;
     stream << "ids.y *= " << p_.nL << ";" << std::endl;
-    stream << "idT.x *= " << p_.simd_width << ";" << std::endl;
+    stream << "idT.x *= " << p_.vwidth << ";" << std::endl;
 
     stream << "M -= ids.x;" << std::endl;
     if(A_trans_=='N')
@@ -289,15 +286,15 @@ matrix_product_parameters::matrix_product_parameters(unsigned int simd_width
 
     for(unsigned int i = 0 ; i < npA ; i++ )
         if (A_trans_=='N')
-          stream << "Ai[" << i << "] += " << Select(backend, to_string(i*p_.local_fetch_0*p_.simd_width) + " < M", "(int)((idT.x + " + to_string(i*p_.local_fetch_0*p_.simd_width) + ")" + ASTRIDE1 + ")", "0") << ";" << std::endl;
+          stream << "Ai[" << i << "] += " << Select(backend, to_string(i*p_.lf0*p_.vwidth) + " < M", "(int)((idT.x + " + to_string(i*p_.lf0*p_.vwidth) + ")" + ASTRIDE1 + ")", "0") << ";" << std::endl;
         else
-          stream << "Ai[" << i << "] += " << Select(backend, to_string(i*p_.local_fetch_1) + " < M", "(int)((idT.y + " + to_string(i*p_.local_fetch_1) + ")*lda)", "0") << ";" << std::endl;
+          stream << "Ai[" << i << "] += " << Select(backend, to_string(i*p_.lf1) + " < M", "(int)((idT.y + " + to_string(i*p_.lf1) + ")*lda)", "0") << ";" << std::endl;
 
     for(unsigned int i = 0 ; i < npB ; i++ )
         if (B_trans_=='T')
-            stream << "Bi[" << i << "] += " << Select(backend, to_string(i*p_.local_fetch_0*p_.simd_width) + " < N", "(int)((idT.x + " + to_string(i*p_.local_fetch_0*p_.simd_width) + ")" + BSTRIDE1 + ")", "0") << ";" << std::endl;
+            stream << "Bi[" << i << "] += " << Select(backend, to_string(i*p_.lf0*p_.vwidth) + " < N", "(int)((idT.x + " + to_string(i*p_.lf0*p_.vwidth) + ")" + BSTRIDE1 + ")", "0") << ";" << std::endl;
         else
-            stream << "Bi[" << i << "] += " << Select(backend, to_string(i*p_.local_fetch_1) + " < N", "(int)((idT.y + " + to_string(i*p_.local_fetch_1) + ")*ldb)", "0") << ";" << std::endl;
+            stream << "Bi[" << i << "] += " << Select(backend, to_string(i*p_.lf1) + " < N", "(int)((idT.y + " + to_string(i*p_.lf1) + ")*ldb)", "0") << ";" << std::endl;
 
     stream << std::endl;
     stream << "//Outer loop" << std::endl;
@@ -315,13 +312,13 @@ matrix_product_parameters::matrix_product_parameters(unsigned int simd_width
         stream << "//Fetch A to local memory" << std::endl;
         if (A_trans_=='N')
         {
-          for(unsigned int k = 0; k < p_.kL; k += p_.local_fetch_1)
-            for(unsigned int m = 0; m < p_.mL; m += p_.local_fetch_0*p_.simd_width)
+          for(unsigned int k = 0; k < p_.kL; k += p_.lf1)
+            for(unsigned int m = 0; m < p_.mL; m += p_.lf0*p_.vwidth)
             {
-              std::string mm = to_string(m/(p_.simd_width*p_.local_fetch_0));
+              std::string mm = to_string(m/(p_.vwidth*p_.lf0));
               std::string kk = to_string(k);
               if(last_iteration)
-                  for(unsigned int s = 0 ; s < p_.simd_width ; ++s)
+                  for(unsigned int s = 0 ; s < p_.vwidth ; ++s)
                       stream << "ldsA[" << k*llda + m + s << "] = (condy" << k << " && " << s << "< M)? Ai[" << mm << "][" << k << "*lda + " << s << "] : 0;" << std::endl;
               else
                 stream << VSTORE(VLOAD_MISALIGNED("0" ,"&Ai[" + mm +"][" + kk + "*lda]"), "0", "ldsA + " + to_string(k*llda+m)) << ";" << std::endl;
@@ -329,13 +326,13 @@ matrix_product_parameters::matrix_product_parameters(unsigned int simd_width
         }
         else
         {
-            for(unsigned int k = 0; k < p_.kL; k += p_.local_fetch_0*p_.simd_width)
-            for(unsigned int m = 0; m < p_.mL; m += p_.local_fetch_1)
+            for(unsigned int k = 0; k < p_.kL; k += p_.lf0*p_.vwidth)
+            for(unsigned int m = 0; m < p_.mL; m += p_.lf1)
               {
-                std::string mm = to_string(m/p_.local_fetch_1);
+                std::string mm = to_string(m/p_.lf1);
                 std::string kk = to_string(k);
                 if(last_iteration)
-                    for(unsigned int s = 0 ; s < p_.simd_width ; ++s)
+                    for(unsigned int s = 0 ; s < p_.vwidth ; ++s)
                         stream << "ldsA[" << m*llda + k + s << "] = condx" << k + s << "? Ai[" << mm << "][" << k + s << ASTRIDE1 << "] : 0;" << std::endl;
 
                 else
@@ -346,13 +343,13 @@ matrix_product_parameters::matrix_product_parameters(unsigned int simd_width
         stream << "//Fetch B to local memory" << std::endl;
         if (B_trans_=='T')
         {
-          for(unsigned int k = 0; k < p_.kL; k += p_.local_fetch_1)
-            for(unsigned int n = 0; n < p_.nL; n += p_.local_fetch_0*p_.simd_width)
+          for(unsigned int k = 0; k < p_.kL; k += p_.lf1)
+            for(unsigned int n = 0; n < p_.nL; n += p_.lf0*p_.vwidth)
             {
-              std::string nn = to_string(n/(p_.simd_width*p_.local_fetch_0));
+              std::string nn = to_string(n/(p_.vwidth*p_.lf0));
               std::string kk = to_string(k);
               if(last_iteration)
-                  for(unsigned int s = 0 ; s < p_.simd_width ; ++s)
+                  for(unsigned int s = 0 ; s < p_.vwidth ; ++s)
                       stream << "ldsB[" << k*lldb + n + s << "] = (condy" << k << " && " << s << "< N)? Bi[" <<  nn << "][" << kk << "*ldb +" << s << "] : 0;" << std::endl;
               else
                 stream << VSTORE(VLOAD_MISALIGNED("0" ,"&Bi[" + nn +"][" + kk + "*ldb]"), "0", "ldsB + " + to_string(k*lldb+n)) << ";" << std::endl;
@@ -360,13 +357,13 @@ matrix_product_parameters::matrix_product_parameters(unsigned int simd_width
         }
         else
         {
-          for(unsigned int k = 0; k < p_.kL; k += p_.local_fetch_0*p_.simd_width)
-            for(unsigned int n = 0; n < p_.nL; n += p_.local_fetch_1)
+          for(unsigned int k = 0; k < p_.kL; k += p_.lf0*p_.vwidth)
+            for(unsigned int n = 0; n < p_.nL; n += p_.lf1)
             {
-              std::string nn = to_string(n/p_.local_fetch_1);
+              std::string nn = to_string(n/p_.lf1);
               std::string kk = to_string(k);
               if(last_iteration)
-                  for(unsigned int s = 0 ; s < p_.simd_width ; ++s)
+                  for(unsigned int s = 0 ; s < p_.vwidth ; ++s)
                       stream << "ldsB[" << n*lldb + k + s << "] = condx" << k + s << "? Bi[" << nn << "][" << k + s << BSTRIDE1 << "] : 0;" << std::endl;
 
               else
@@ -375,14 +372,14 @@ matrix_product_parameters::matrix_product_parameters(unsigned int simd_width
         }
 
         if(A_trans_=='N')
-            stream << "ldsA = lA + ids.z*" << p_.simd_width << ";" << std::endl;
+            stream << "ldsA = lA + ids.z*" << p_.vwidth << ";" << std::endl;
         else
-            stream << "ldsA = lA + ids.z*" << llda*p_.simd_width << ";" << std::endl;
+            stream << "ldsA = lA + ids.z*" << llda*p_.vwidth << ";" << std::endl;
 
         if(B_trans_=='T')
-            stream << "ldsB = lB + ids.w*" << p_.simd_width << ";" << std::endl;
+            stream << "ldsB = lB + ids.w*" << p_.vwidth << ";" << std::endl;
         else
-            stream << "ldsB = lB + ids.w*" << lldb*p_.simd_width << ";" << std::endl;
+            stream << "ldsB = lB + ids.w*" << lldb*p_.vwidth << ";" << std::endl;
 
         stream << "$LOCAL_BARRIER;" << std::endl;
 
@@ -393,19 +390,19 @@ matrix_product_parameters::matrix_product_parameters(unsigned int simd_width
         stream << "//Fetch A to registers" << std::endl;
         stream << "#pragma unroll" << std::endl;
         stream << "for(unsigned int kk = 0; kk < " << p_.kS << "; kk++)" << std::endl;
-        stream << "#pragma unroll " << p_.mS/p_.simd_width << std::endl;
-        stream << "for(unsigned int mm = 0; mm < " << p_.mS/p_.simd_width << "; mm++)" << std::endl;
+        stream << "#pragma unroll " << p_.mS/p_.vwidth << std::endl;
+        stream << "for(unsigned int mm = 0; mm < " << p_.mS/p_.vwidth << "; mm++)" << std::endl;
         stream << "{" << std::endl;
         stream.inc_tab();
         if(A_trans_=='N')
-            stream << "rA[kk][mm] = "  << VLOAD("0", "ldsA + k*" + to_string(llda) + " + mm*" + to_string(p_.local_size_0*p_.simd_width) + "+ kk*" + to_string(llda)) << ";" << std::endl;
+            stream << "rA[kk][mm] = "  << VLOAD("0", "ldsA + k*" + to_string(llda) + " + mm*" + to_string(p_.ls0*p_.vwidth) + "+ kk*" + to_string(llda)) << ";" << std::endl;
         else
         {
-            if(p_.simd_width==1)
-                stream << "rA[kk][mm] = ldsA[k + mm*" << p_.local_size_0*llda <<  "+ kk"  << "];" << std::endl;
+            if(p_.vwidth==1)
+                stream << "rA[kk][mm] = ldsA[k + mm*" << p_.ls0*llda <<  "+ kk"  << "];" << std::endl;
             else
-                for(unsigned int s = 0 ; s < p_.simd_width ; ++s)
-                    stream << access_vector_type("rA[kk][mm]", s) << " = ldsA[k + (mm*" << p_.simd_width*p_.local_size_0 << " + " << s << ")*" << llda <<  "+ kk];" << std::endl;
+                for(unsigned int s = 0 ; s < p_.vwidth ; ++s)
+                    stream << access_vector_type("rA[kk][mm]", s) << " = ldsA[k + (mm*" << p_.vwidth*p_.ls0 << " + " << s << ")*" << llda <<  "+ kk];" << std::endl;
         }
 
         stream.dec_tab();
@@ -414,19 +411,19 @@ matrix_product_parameters::matrix_product_parameters(unsigned int simd_width
         stream << "//Fetch B to registers" << std::endl;
         stream << "#pragma unroll " << p_.kS << std::endl;
         stream << "for(unsigned int kk = 0; kk < " << p_.kS << "; kk++)" << std::endl;
-        stream << "#pragma unroll " << p_.nS/p_.simd_width << std::endl;
-        stream << "for(unsigned int nn = 0; nn < " << p_.nS/p_.simd_width << "; nn++)" << std::endl;
+        stream << "#pragma unroll " << p_.nS/p_.vwidth << std::endl;
+        stream << "for(unsigned int nn = 0; nn < " << p_.nS/p_.vwidth << "; nn++)" << std::endl;
         stream << "{" << std::endl;
         stream.inc_tab();
         if(B_trans_=='T')
-            stream << "rB[kk][nn] = " << VLOAD("0", "ldsB + k*" + to_string(lldb) + " + nn*" + to_string(p_.local_size_1*p_.simd_width)  + "+ kk*" + to_string(lldb)) << ";" << std::endl;
+            stream << "rB[kk][nn] = " << VLOAD("0", "ldsB + k*" + to_string(lldb) + " + nn*" + to_string(p_.ls1*p_.vwidth)  + "+ kk*" + to_string(lldb)) << ";" << std::endl;
         else
         {
-            if(p_.simd_width==1)
-                stream << "rB[kk][nn] = ldsB[k"  << " + nn*" << p_.local_size_1*lldb <<  "+ kk"  << "];" << std::endl;
+            if(p_.vwidth==1)
+                stream << "rB[kk][nn] = ldsB[k"  << " + nn*" << p_.ls1*lldb <<  "+ kk"  << "];" << std::endl;
             else
-                for(unsigned int s = 0 ; s < p_.simd_width ; ++s)
-                    stream << access_vector_type("rB[kk][nn]", s) << " = ldsB[k"  << " + (nn*" << p_.simd_width*p_.local_size_1 << " + " << s << ")*" << lldb <<  "+ kk];" << std::endl;
+                for(unsigned int s = 0 ; s < p_.vwidth ; ++s)
+                    stream << access_vector_type("rB[kk][nn]", s) << " = ldsB[k"  << " + (nn*" << p_.vwidth*p_.ls1 << " + " << s << ")*" << lldb <<  "+ kk];" << std::endl;
         }
         stream.dec_tab();
         stream << "}" << std::endl;
@@ -437,14 +434,14 @@ matrix_product_parameters::matrix_product_parameters(unsigned int simd_width
         for(unsigned int mm=0; mm < p_.mS; ++mm){
           string res_str, lhs_str, rhs_str;
           res_str = "rC[" + to_string(mm) + "][" + to_string(nn) + "]";
-          if (p_.simd_width==1)
+          if (p_.vwidth==1)
             lhs_str = "rA[" + to_string(kk) + "][" + to_string(mm) + "]";
           else
-            lhs_str = access_vector_type("rA[" + to_string(kk) + "][" + to_string(mm/p_.simd_width) + "]", mm%p_.simd_width);
-          if (p_.simd_width==1)
+            lhs_str = access_vector_type("rA[" + to_string(kk) + "][" + to_string(mm/p_.vwidth) + "]", mm%p_.vwidth);
+          if (p_.vwidth==1)
             rhs_str = "rB[" + to_string(kk) + "]["+to_string(nn)+"]";
           else
-            rhs_str = access_vector_type("rB[" + to_string(kk) + "]["+to_string(nn/p_.simd_width)+"]", nn%p_.simd_width);
+            rhs_str = access_vector_type("rB[" + to_string(kk) + "]["+to_string(nn/p_.vwidth)+"]", nn%p_.vwidth);
           stream << res_str << "= $MAD(" << lhs_str << "," << rhs_str << "," << res_str << ");" << std::endl;
         }
 
@@ -476,15 +473,15 @@ matrix_product_parameters::matrix_product_parameters(unsigned int simd_width
     if(A_trans_=='N' || B_trans_=='T')
     {
         stream << "int Ky = K - idT.y;" << std::endl;
-        for(unsigned int k = 0; k < p_.kL; k += p_.local_fetch_1)
+        for(unsigned int k = 0; k < p_.kL; k += p_.lf1)
             stream << "int condy" << k << " = " << k << " < Ky;" << std::endl;
     }
 
     if(A_trans_=='T' || B_trans_=='N')
     {
         stream << "int Kx = K - idT.x;" << std::endl;
-        for(unsigned int k = 0 ; k < p_.kL ; k += p_.local_fetch_0*p_.simd_width)
-            for(unsigned int s = 0 ; s < p_.simd_width ; ++s)
+        for(unsigned int k = 0 ; k < p_.kL ; k += p_.lf0*p_.vwidth)
+            for(unsigned int s = 0 ; s < p_.vwidth ; ++s)
                 stream << "int condx" << k + s << " = " << k + s << " < Kx;" << std::endl;
     }
     fetch_to_lds(true);
@@ -503,35 +500,35 @@ matrix_product_parameters::matrix_product_parameters(unsigned int simd_width
     stream << "N += ids.y;" << std::endl;
 
     stream << "C += ids.x" << CSTRIDE1 << ";" << std::endl;
-    stream << "C += ids.z*" << p_.simd_width << CSTRIDE1 << ";" << std::endl;
+    stream << "C += ids.z*" << p_.vwidth << CSTRIDE1 << ";" << std::endl;
     stream << "C += ids.y*ldc;" << std::endl;
-    stream << "C += ids.w*" << p_.simd_width << "*ldc;" << std::endl;
+    stream << "C += ids.w*" << p_.vwidth << "*ldc;" << std::endl;
     if(has_depth)
         stream << "C += gidz*ldc*N;" << std::endl;
 
     stream << "M -= ids.x;" << std::endl;
-    stream << "M -= ids.z*" << p_.simd_width << ";" << std::endl;
+    stream << "M -= ids.z*" << p_.vwidth << ";" << std::endl;
 
     stream << "N -= ids.y;" << std::endl;
-    stream << "N -= ids.w*" << p_.simd_width <<  ";" << std::endl;
+    stream << "N -= ids.w*" << p_.vwidth <<  ";" << std::endl;
 
     for(unsigned int n=0; n < p_.nS; ++n)
     {
-        string Cj = to_string((n/p_.simd_width)*(p_.local_size_1*p_.simd_width) + n%p_.simd_width);
+        string Cj = to_string((n/p_.vwidth)*(p_.ls1*p_.vwidth) + n%p_.vwidth);
         stream << "if(" << Cj << " >= N) return;" << std::endl;
         for(unsigned int m=0; m < p_.mS; ++m)
             stream << "rC[" << m << "][" << n << "] *= alpha;" << std::endl;
         for(unsigned int m=0; m < p_.mS; ++m)
         {
-            string Ci = to_string((m/p_.simd_width)*(p_.local_size_0*p_.simd_width) + m%p_.simd_width);
+            string Ci = to_string((m/p_.vwidth)*(p_.ls0*p_.vwidth) + m%p_.vwidth);
             stream << "if(" << Ci << "< M) ";
             if(has_depth)
                 stream << "C[" << Ci << CSTRIDE1 << "] = rC[" << m << "][" << n << "];" << std::endl;
             else
                 stream << "C[" << Ci << CSTRIDE1 << "] = rC[" << m << "][" << n << "] + ((beta != (" << sdtype << ")0)?(beta*" << "C[" << Ci << CSTRIDE1 << "]):0);" << std::endl;
         }
-        if((n+1)%p_.simd_width==0){
-            stream << "C += ldc*" << p_.local_size_1*p_.simd_width - p_.simd_width + 1 << ";" << std::endl;
+        if((n+1)%p_.vwidth==0){
+            stream << "C += ldc*" << p_.ls1*p_.vwidth - p_.vwidth + 1 << ";" << std::endl;
         }
         else{
             stream << "C += ldc;" << std::endl;
@@ -599,8 +596,8 @@ matrix_product_parameters::matrix_product_parameters(unsigned int simd_width
     reduce_name += suffix;
 
     driver::Kernel matrix_product(program, matrix_product_name.c_str());
-    driver::NDRange local(p_.local_size_0, p_.local_size_1, 1);
-    driver::NDRange global(align(align(M,p_.mS)/p_.mS, p_.local_size_0), align(align(N,p_.nS)/p_.nS, p_.local_size_1), p_.depth);
+    driver::NDRange local(p_.ls0, p_.ls1, 1);
+    driver::NDRange global(align(align(M,p_.mS)/p_.mS, p_.ls0), align(align(N,p_.nS)/p_.nS, p_.ls1), p_.depth);
 
     unsigned int current_arg = 0;
 
@@ -651,8 +648,8 @@ matrix_product_parameters::matrix_product_parameters(unsigned int simd_width
     {
       unsigned int current_arg = 0;
       driver::Kernel reduce(program, reduce_name.c_str());
-      driver::NDRange local(p_.local_size_0, p_.local_size_1);
-      driver::NDRange global(align(M, p_.local_size_0), align(N, p_.local_size_1));
+      driver::NDRange local(p_.ls0, p_.ls1);
+      driver::NDRange global(align(M, p_.ls0), align(N, p_.ls1));
       reduce.setSizeArg(current_arg++, M);
       reduce.setSizeArg(current_arg++, N);
       reduce.setSizeArg(current_arg++, p_.depth);
@@ -717,7 +714,7 @@ matrix_product_parameters::matrix_product_parameters(unsigned int simd_width
   matrix_product_nn::matrix_product_nn(unsigned int simd
                            , int_t ls0, int_t KL, int_t ls1, int_t D
                            , int_t ms, int_t ks, int_t ns
-                           , fetching_policy_type Afetch , fetching_policy_type Bfetch
+                           , fetch_type Afetch , fetch_type Bfetch
                            , int_t lfetch0, int_t lfetch1) :
     matrix_product(matrix_product_parameters(simd, ls0, KL, ls1, D, ms, ks, ns, Afetch, Bfetch, lfetch0, lfetch1), 'N', 'N')
   {
@@ -727,7 +724,7 @@ matrix_product_parameters::matrix_product_parameters(unsigned int simd_width
   matrix_product_tn::matrix_product_tn(unsigned int simd
                            , int_t ls0, int_t KL, int_t ls1, int_t D
                            , int_t ms, int_t ks, int_t ns
-                           , fetching_policy_type Afetch , fetching_policy_type Bfetch
+                           , fetch_type Afetch , fetch_type Bfetch
                            , int_t lfetch0, int_t lfetch1) :
     matrix_product(matrix_product_parameters(simd, ls0, KL, ls1, D, ms, ks, ns, Afetch, Bfetch, lfetch0, lfetch1), 'T', 'N')
   { }
@@ -736,7 +733,7 @@ matrix_product_parameters::matrix_product_parameters(unsigned int simd_width
   matrix_product_nt::matrix_product_nt(unsigned int simd
                            , int_t ls0, int_t KL, int_t ls1, int_t D
                            , int_t ms, int_t ks, int_t ns
-                           , fetching_policy_type Afetch , fetching_policy_type Bfetch
+                           , fetch_type Afetch , fetch_type Bfetch
                            , int_t lfetch0, int_t lfetch1) :
     matrix_product(matrix_product_parameters(simd, ls0, KL, ls1, D, ms, ks, ns, Afetch, Bfetch, lfetch0, lfetch1), 'N', 'T')
   { }
@@ -745,7 +742,7 @@ matrix_product_parameters::matrix_product_parameters(unsigned int simd_width
   matrix_product_tt::matrix_product_tt(unsigned int simd
                            , int_t ls0, int_t KL, int_t ls1, int_t D
                            , int_t ms, int_t ks, int_t ns
-                           , fetching_policy_type Afetch , fetching_policy_type Bfetch
+                           , fetch_type Afetch , fetch_type Bfetch
                            , int_t lfetch0, int_t lfetch1) :
     matrix_product(matrix_product_parameters(simd, ls0, KL, ls1, D, ms, ks, ns, Afetch, Bfetch, lfetch0, lfetch1), 'T', 'T')
   { }
diff --git a/lib/jit/generation/reduce_1d.cpp b/lib/jit/generation/reduce_1d.cpp
index 62080850b..90badfcbb 100644
--- a/lib/jit/generation/reduce_1d.cpp
+++ b/lib/jit/generation/reduce_1d.cpp
@@ -35,19 +35,19 @@ namespace isaac
 {
 namespace templates
 {
-reduce_1d_parameters::reduce_1d_parameters(unsigned int _simd_width,
+reduce_1d_parameters::reduce_1d_parameters(unsigned int _vwidth,
                      unsigned int _group_size, unsigned int _num_groups,
-                     fetching_policy_type _fetching_policy) : base::parameters_type(_simd_width, _group_size, 1, 2), num_groups(_num_groups), fetching_policy(_fetching_policy)
+                     fetch_type _fetch) : base::parameters_type(_vwidth, _group_size, 1, 2), num_groups(_num_groups), fetch(_fetch)
 { }
 
 unsigned int reduce_1d::lmem_usage(expression_tree const  & x) const
 {
-  return p_.local_size_0*size_of(x.dtype());
+  return p_.ls0*size_of(x.dtype());
 }
 
 int reduce_1d::is_invalid_impl(driver::Device const &, expression_tree const  &) const
 {
-  if (p_.fetching_policy==FETCH_FROM_LOCAL)
+  if (p_.fetch==FETCH_FROM_LOCAL)
     return TEMPLATE_INVALID_FETCHING_POLICY_TYPE;
   return TEMPLATE_VALID;
 }
@@ -122,7 +122,7 @@ std::string reduce_1d::generate_impl(std::string const & suffix, expression_tree
     case driver::CUDA:
       stream << "#include  \"vector.h\"" << std::endl; break;
     case driver::OPENCL:
-      stream << " __attribute__((reqd_work_group_size(" << p_.local_size_0 << ",1,1)))" << std::endl; break;
+      stream << " __attribute__((reqd_work_group_size(" << p_.ls0 << ",1,1)))" << std::endl; break;
   }
   stream << "$KERNEL void prod" << suffix << "($SIZE_T N, $GLOBAL char* tmp," << tools::join(kernel_arguments(device, symbols, tree), ", ") << ")" << std::endl;
   stream << "{" << std::endl;
@@ -139,33 +139,33 @@ std::string reduce_1d::generate_impl(std::string const & suffix, expression_tree
   {
     if(is_indexing(rd->op().type))
     {
-      stream << rd->process("$LOCAL #scalartype #name_buf_value[" + tools::to_string(p_.local_size_0) + "];") << std::endl;
+      stream << rd->process("$LOCAL #scalartype #name_buf_value[" + tools::to_string(p_.ls0) + "];") << std::endl;
       stream << rd->process("#scalartype #name_acc_value = " + neutral_element(rd->op(), backend, "#scalartype") + ";") << std::endl;
-      stream << rd->process("$LOCAL unsigned int #name_buf[" + tools::to_string(p_.local_size_0) + "];") << std::endl;
+      stream << rd->process("$LOCAL unsigned int #name_buf[" + tools::to_string(p_.ls0) + "];") << std::endl;
       stream << rd->process("unsigned int #name_acc = 0;") << std::endl;
     }
     else
     {
-      stream << rd->process("$LOCAL #scalartype #name_buf[" + tools::to_string(p_.local_size_0) + "];") << std::endl;
+      stream << rd->process("$LOCAL #scalartype #name_buf[" + tools::to_string(p_.ls0) + "];") << std::endl;
       stream << rd->process("#scalartype #name_acc = " + neutral_element(rd->op(), backend, "#scalartype") + ";") << std::endl;
     }
   }
-  element_wise_loop_1D(stream, p_.fetching_policy, p_.simd_width, "i", "N", "$GLOBAL_IDX_0", "$GLOBAL_SIZE_0", device, [&](unsigned int simd_width)
+  element_wise_loop_1D(stream, p_.fetch, p_.vwidth, "i", "N", "$GLOBAL_IDX_0", "$GLOBAL_SIZE_0", device, [&](unsigned int vwidth)
   {
-    std::string dtype = append_width("#scalartype",simd_width);
+    std::string dtype = append_width("#scalartype",vwidth);
     //Fetch vector entry
     std::set<std::string> fetched;
      for (symbolic::reduce_1d* rd : reductions)
        for(symbolic::leaf* leaf: symbolic::extract<symbolic::leaf>(tree, symbols, rd->root(), false))
           if(fetched.insert(leaf->process("#name")).second)
-            stream << leaf->process(dtype + " #name = " + append_width("loadv", simd_width) + "(i);") << std::endl;
+            stream << leaf->process(dtype + " #name = " + append_width("loadv", vwidth) + "(i);") << std::endl;
     //Update accumulators
     for (symbolic::reduce_1d* rd : reductions)
-      for (unsigned int s = 0; s < simd_width; ++s)
+      for (unsigned int s = 0; s < vwidth; ++s)
       {
-        std::string value = rd->lhs()->evaluate({{"leaf", access_vector_type("#name", s, simd_width)}});
+        std::string value = rd->lhs()->evaluate({{"leaf", access_vector_type("#name", s, vwidth)}});
         if (is_indexing(rd->op().type))
-          compute_index_reduce_1d(stream, rd->process("#name_acc"),  "i*" + tools::to_string(simd_width) + "+" + tools::to_string(s), rd->process("#name_acc_value"), value,rd->op());
+          compute_index_reduce_1d(stream, rd->process("#name_acc"),  "i*" + tools::to_string(vwidth) + "+" + tools::to_string(s), rd->process("#name_acc_value"), value,rd->op());
         else
           compute_reduce_1d(stream, rd->process("#name_acc"), value,rd->op());
       }
@@ -178,7 +178,7 @@ std::string reduce_1d::generate_impl(std::string const & suffix, expression_tree
     stream << rd->process("#name_buf[lid] = #name_acc;") << std::endl;
   }
   //Reduce local memory
-  reduce_1d_local_memory(stream, p_.local_size_0, reductions, "#name_buf", "#name_buf_value", backend);
+  reduce_1d_local_memory(stream, p_.ls0, reductions, "#name_buf", "#name_buf_value", backend);
   //Write to temporary buffers
   stream << "if (lid==0)" << std::endl;
   stream << "{" << std::endl;
@@ -209,14 +209,14 @@ std::string reduce_1d::generate_impl(std::string const & suffix, expression_tree
   {
     if (is_indexing(rd->op().type))
     {
-      stream << rd->process("$LOCAL unsigned int #name_buf[" + tools::to_string(p_.local_size_0) + "];");
+      stream << rd->process("$LOCAL unsigned int #name_buf[" + tools::to_string(p_.ls0) + "];");
       stream << rd->process("unsigned int #name_acc = 0;") << std::endl;
-      stream << rd->process("$LOCAL #scalartype #name_buf_value[" + tools::to_string(p_.local_size_0) + "];") << std::endl;
+      stream << rd->process("$LOCAL #scalartype #name_buf_value[" + tools::to_string(p_.ls0) + "];") << std::endl;
       stream << rd->process("#scalartype #name_acc_value = " + neutral_element(rd->op(), backend, "#scalartype") + ";");
     }
     else
     {
-      stream << rd->process("$LOCAL #scalartype #name_buf[" + tools::to_string(p_.local_size_0) + "];") << std::endl;
+      stream << rd->process("$LOCAL #scalartype #name_buf[" + tools::to_string(p_.ls0) + "];") << std::endl;
       stream << rd->process("#scalartype #name_acc = " + neutral_element(rd->op(), backend, "#scalartype") + ";");
     }
   }
@@ -238,7 +238,7 @@ std::string reduce_1d::generate_impl(std::string const & suffix, expression_tree
     stream << rd->process("#name_buf[lid] = #name_acc;") << std::endl;
   }
   //Local reduction
-  reduce_1d_local_memory(stream, p_.local_size_0, reductions, "#name_buf", "#name_buf_value", backend);
+  reduce_1d_local_memory(stream, p_.ls0, reductions, "#name_buf", "#name_buf_value", backend);
   //Write
   stream << "if (lid==0)" << std::endl;
   stream << "{" << std::endl;
@@ -258,7 +258,7 @@ reduce_1d::reduce_1d(reduce_1d::parameters_type const & parameters,
 { }
 
 reduce_1d::reduce_1d(unsigned int simd, unsigned int ls, unsigned int ng,
-                               fetching_policy_type fetch, fusion_policy_t bind):
+                               fetch_type fetch, fusion_policy_t bind):
     base_impl<reduce_1d, reduce_1d_parameters>(reduce_1d_parameters(simd,ls,ng,fetch), bind)
 {}
 
@@ -284,8 +284,8 @@ void reduce_1d::enqueue(driver::CommandQueue & queue, driver::Program const & pr
   driver::Kernel kernels[2] = { driver::Kernel(program,name[0].c_str()), driver::Kernel(program,name[1].c_str()) };
 
   //NDRange
-  driver::NDRange global[2] = { driver::NDRange(p_.local_size_0*p_.num_groups), driver::NDRange(p_.local_size_0) };
-  driver::NDRange local[2] = { driver::NDRange(p_.local_size_0), driver::NDRange(p_.local_size_0) };
+  driver::NDRange global[2] = { driver::NDRange(p_.ls0*p_.num_groups), driver::NDRange(p_.ls0) };
+  driver::NDRange local[2] = { driver::NDRange(p_.ls0), driver::NDRange(p_.ls0) };
   //Arguments
   for (auto & kernel : kernels)
   {
diff --git a/lib/jit/generation/reduce_2d.cpp b/lib/jit/generation/reduce_2d.cpp
index 93ec4723c..4b2ea7a36 100644
--- a/lib/jit/generation/reduce_2d.cpp
+++ b/lib/jit/generation/reduce_2d.cpp
@@ -39,9 +39,9 @@ namespace isaac
 namespace templates
 {
 
-reduce_2d_parameters::reduce_2d_parameters(unsigned int _simd_width,
-                              unsigned int _local_size_0, unsigned int _local_size_1,
-                              unsigned int _num_groups_0, unsigned int _num_groups_1, fetching_policy_type _fetch_policy): base::parameters_type(_simd_width, _local_size_0, _local_size_1, 1),
+reduce_2d_parameters::reduce_2d_parameters(unsigned int _vwidth,
+                              unsigned int _ls0, unsigned int _ls1,
+                              unsigned int _num_groups_0, unsigned int _num_groups_1, fetch_type _fetch_policy): base::parameters_type(_vwidth, _ls0, _ls1, 1),
 num_groups_0(_num_groups_0), num_groups_1(_num_groups_1), fetch_policy(_fetch_policy) { }
 
 
@@ -54,7 +54,7 @@ int reduce_2d::is_invalid_impl(driver::Device const &, expression_tree const &)
 
 unsigned int reduce_2d::lmem_usage(const expression_tree&) const
 {
-  return (p_.local_size_0+1)*p_.local_size_1;
+  return (p_.ls0+1)*p_.ls1;
 }
 
 unsigned int reduce_2d::temporary_workspace(expression_tree const & expressions) const
@@ -80,7 +80,7 @@ std::string reduce_2d::generate_impl(std::string const & suffix, expression_tree
   name[0] += suffix;
   name[1] += suffix;
 
-  unsigned int ldls = p_.local_size_0;
+  unsigned int ldls = p_.ls0;
   std::string ls0ldstr = to_string(ldls);
 
   auto unroll_tmp = [&]()
@@ -113,7 +113,7 @@ std::string reduce_2d::generate_impl(std::string const & suffix, expression_tree
       stream << "#include  \"vector.h\"" << std::endl;
       break;
     case driver::OPENCL:
-      stream << " __attribute__((reqd_work_group_size(" << p_.local_size_0 << "," << p_.local_size_1 << ",1)))" << std::endl;
+      stream << " __attribute__((reqd_work_group_size(" << p_.ls0 << "," << p_.ls1 << ",1)))" << std::endl;
       break;
   }
   stream << "$KERNEL void " << name[0] << "($SIZE_T M, $SIZE_T N, $GLOBAL char* tmp, " << tools::join(kernel_arguments(device, symbols, tree), ", ") << ")" << std::endl;
@@ -125,13 +125,13 @@ std::string reduce_2d::generate_impl(std::string const & suffix, expression_tree
   stream << "$SIZE_T lidy = $LOCAL_IDX_1;" << std::endl;
   //Loop r
   std::ostringstream upper;
-  upper << "(M +" << p_.local_size_1 - 1 << ")/" << p_.local_size_1 << "*" << p_.local_size_1;
+  upper << "(M +" << p_.ls1 - 1 << ")/" << p_.ls1 << "*" << p_.ls1;
 
-  element_wise_loop_1D(stream, p_.fetch_policy, (reduction_type_==REDUCE_ROWS)?p_.simd_width:1, "r", upper.str(), "$GLOBAL_IDX_1", "$GLOBAL_SIZE_1", device, [&](unsigned int cwidth)
+  element_wise_loop_1D(stream, p_.fetch_policy, (reduction_type_==REDUCE_ROWS)?p_.vwidth:1, "r", upper.str(), "$GLOBAL_IDX_1", "$GLOBAL_SIZE_1", device, [&](unsigned int cwidth)
   {
   //Declare Buffers
   for (symbolic::reduce_2d* rd : reductions)
-    stream << rd->process("$LOCAL " + append_width("#scalartype", cwidth) + " #name_buf[" + to_string(p_.local_size_1*ldls) + "];") << std::endl;
+    stream << rd->process("$LOCAL " + append_width("#scalartype", cwidth) + " #name_buf[" + to_string(p_.ls1*ldls) + "];") << std::endl;
 
   //Accumulators
   for (symbolic::reduce_2d* rd : reductions){
@@ -142,7 +142,7 @@ std::string reduce_2d::generate_impl(std::string const & suffix, expression_tree
   stream << "if (r < M)" << std::endl;
   stream << "{" << std::endl;
   stream.inc_tab();
-  element_wise_loop_1D(stream, p_.fetch_policy, (reduction_type_==REDUCE_COLUMNS)?p_.simd_width:1, "c", "N", "$GLOBAL_IDX_0", "$GLOBAL_SIZE_0", device, [&](unsigned int rwidth)
+  element_wise_loop_1D(stream, p_.fetch_policy, (reduction_type_==REDUCE_COLUMNS)?p_.vwidth:1, "c", "N", "$GLOBAL_IDX_0", "$GLOBAL_SIZE_0", device, [&](unsigned int rwidth)
   {
     std::string rdtype = append_width("#scalartype", rwidth);
     std::string cdtype = append_width("#scalartype", cwidth);
@@ -173,7 +173,7 @@ std::string reduce_2d::generate_impl(std::string const & suffix, expression_tree
     stream << rd->process("#name_buf[lidy*" + ls0ldstr + "+ lidx] = #name_acc;") << std::endl;
   //Reduce local memory
   stream << "#pragma unroll" << std::endl;
-  stream << "for($SIZE_T stride = " << p_.local_size_0/2 << "; stride >0; stride /=2)" << std::endl;
+  stream << "for($SIZE_T stride = " << p_.ls0/2 << "; stride >0; stride /=2)" << std::endl;
   stream << "{" << std::endl;
   stream.inc_tab();
   stream << "$LOCAL_BARRIER;" << std::endl;
@@ -220,14 +220,14 @@ std::string reduce_2d::generate_impl(std::string const & suffix, expression_tree
   if(p_.num_groups_0>1)
   {
   if(backend==driver::OPENCL)
-    stream << " __attribute__((reqd_work_group_size(" << p_.local_size_0 << "," << p_.local_size_1 << ",1)))" << std::endl;
+    stream << " __attribute__((reqd_work_group_size(" << p_.ls0 << "," << p_.ls1 << ",1)))" << std::endl;
   stream << "$KERNEL void " << name[1] << "($SIZE_T M, $SIZE_T N , $GLOBAL char* tmp, " << tools::join(kernel_arguments(device, symbols, tree), ", ") << ")" << std::endl;
   stream << "{" << std::endl;
   stream.inc_tab();
   unroll_tmp();
   for (symbolic::reduce_2d* rd : reductions)
-    stream << rd->process("$LOCAL #scalartype #name_buf[" + to_string(p_.local_size_1*ldls) + "];") << std::endl;
-  stream << "for($SIZE_T r = $GLOBAL_IDX_1; r < (M +" << p_.local_size_1 - 1 << ")/" << p_.local_size_1 << "*" << p_.local_size_1 << "; r += " << GlobalSize1(backend) << "){" << std::endl;
+    stream << rd->process("$LOCAL #scalartype #name_buf[" + to_string(p_.ls1*ldls) + "];") << std::endl;
+  stream << "for($SIZE_T r = $GLOBAL_IDX_1; r < (M +" << p_.ls1 - 1 << ")/" << p_.ls1 << "*" << p_.ls1 << "; r += " << GlobalSize1(backend) << "){" << std::endl;
   stream.inc_tab();
   stream << "$SIZE_T lidx = $LOCAL_IDX_0;" << std::endl;
   stream << "$SIZE_T lidy = $LOCAL_IDX_1;" << std::endl;
@@ -247,7 +247,7 @@ std::string reduce_2d::generate_impl(std::string const & suffix, expression_tree
   for (symbolic::reduce_2d* rd : reductions)
     stream << rd->process("#name_buf[lidy*" + ls0ldstr + "+ lidx] = #name_acc;") << std::endl;
   stream << "#pragma unroll" << std::endl;
-  stream << "for($SIZE_T stride = " << p_.local_size_0/2 << "; stride >0; stride /=2)" << std::endl;
+  stream << "for($SIZE_T stride = " << p_.ls0/2 << "; stride >0; stride /=2)" << std::endl;
   stream << "{" << std::endl;
   stream.inc_tab();
   stream << "$LOCAL_BARRIER;" << std::endl;
@@ -326,8 +326,8 @@ void reduce_2d::enqueue(driver::CommandQueue & queue, driver::Program const & pr
   }
 
   //NDRange
-  driver::NDRange global[2] = { driver::NDRange(p_.local_size_0*p_.num_groups_0, p_.local_size_1*p_.num_groups_1), driver::NDRange(p_.local_size_0, p_.local_size_1*p_.num_groups_1) };
-  driver::NDRange local[2] = { driver::NDRange(p_.local_size_0, p_.local_size_1), driver::NDRange(p_.local_size_0, p_.local_size_1) };
+  driver::NDRange global[2] = { driver::NDRange(p_.ls0*p_.num_groups_0, p_.ls1*p_.num_groups_1), driver::NDRange(p_.ls0, p_.ls1*p_.num_groups_1) };
+  driver::NDRange local[2] = { driver::NDRange(p_.ls0, p_.ls1), driver::NDRange(p_.ls0, p_.ls1) };
   for(unsigned int i = 0 ; i < nk ; ++i)
     control.execution_options().enqueue(program.context(), kernels[i], global[i], local[i]);
 }
@@ -335,12 +335,12 @@ void reduce_2d::enqueue(driver::CommandQueue & queue, driver::Program const & pr
 reduce_2d_rows::reduce_2d_rows(reduce_2d_parameters  const & parameters,fusion_policy_t fusion_policy): reduce_2d(parameters, REDUCE_ROWS, fusion_policy){}
 
 reduce_2d_rows::reduce_2d_rows(unsigned int simd, unsigned int ls1, unsigned int ls2,  unsigned int ng1, unsigned int ng2,
-               fetching_policy_type fetch, fusion_policy_t bind): reduce_2d(reduce_2d_parameters(simd, ls1, ls2, ng1, ng2, fetch), REDUCE_ROWS, bind) {}
+               fetch_type fetch, fusion_policy_t bind): reduce_2d(reduce_2d_parameters(simd, ls1, ls2, ng1, ng2, fetch), REDUCE_ROWS, bind) {}
 
 reduce_2d_cols::reduce_2d_cols(reduce_2d::parameters_type  const & parameters, fusion_policy_t fusion_policy): reduce_2d(parameters, REDUCE_COLUMNS, fusion_policy){}
 
 reduce_2d_cols::reduce_2d_cols(unsigned int simd, unsigned int ls1, unsigned int ls2, unsigned int ng1, unsigned int ng2,
-               fetching_policy_type fetch, fusion_policy_t bind): reduce_2d(reduce_2d_parameters(simd, ls1, ls2, ng1, ng2, fetch), REDUCE_COLUMNS, bind) {}
+               fetch_type fetch, fusion_policy_t bind): reduce_2d(reduce_2d_parameters(simd, ls1, ls2, ng1, ng2, fetch), REDUCE_COLUMNS, bind) {}
 
 
 }
diff --git a/lib/jit/generation/tools/loop.hpp b/lib/jit/generation/tools/loop.hpp
index a79479618..b66ae701c 100644
--- a/lib/jit/generation/tools/loop.hpp
+++ b/lib/jit/generation/tools/loop.hpp
@@ -29,7 +29,7 @@ namespace isaac
 namespace templates
 {
 
-inline void fetching_loop_info(fetching_policy_type policy, std::string const & bound, kernel_generation_stream & stream, std::string & init, std::string & upper_bound, std::string & inc, std::string const & domain_id, std::string const & domain_size, driver::Device const &)
+inline void fetching_loop_info(fetch_type policy, std::string const & bound, kernel_generation_stream & stream, std::string & init, std::string & upper_bound, std::string & inc, std::string const & domain_id, std::string const & domain_size, driver::Device const &)
 {
   if (policy==FETCH_FROM_GLOBAL_STRIDED)
   {
@@ -54,10 +54,10 @@ inline void fetching_loop_info(fetching_policy_type policy, std::string const &
 
 
 template<class Fun>
-inline void element_wise_loop_1D(kernel_generation_stream & stream, fetching_policy_type fetch, unsigned int simd_width,
+inline void element_wise_loop_1D(kernel_generation_stream & stream, fetch_type fetch, unsigned int vwidth,
                                  std::string const & i, std::string const & bound, std::string const & domain_id, std::string const & domain_size, driver::Device const & device, Fun const & generate_body)
 {
-  std::string strwidth = tools::to_string(simd_width);
+  std::string strwidth = tools::to_string(vwidth);
 
   std::string init, upper_bound, inc;
   fetching_loop_info(fetch, bound, stream, init, upper_bound, inc, domain_id, domain_size, device);
@@ -65,11 +65,11 @@ inline void element_wise_loop_1D(kernel_generation_stream & stream, fetching_pol
   stream << "for(unsigned int " << i << " = " << init << "*" << strwidth << "; " << i << " < " << boundround << "; " << i << " += " << inc << "*" << strwidth << ")" << std::endl;
   stream << "{" << std::endl;
   stream.inc_tab();
-  generate_body(simd_width);
+  generate_body(vwidth);
   stream.dec_tab();
   stream << "}" << std::endl;
 
-  if (simd_width>1)
+  if (vwidth>1)
   {
     stream << "for(unsigned int " << i << " = " << boundround << " + " << domain_id << "; " << i << " < " << bound << "; " << i << " += " + domain_size + ")" << std::endl;
     stream << "{" << std::endl;
diff --git a/lib/jit/generation/tools/vector_types.hpp b/lib/jit/generation/tools/vector_types.hpp
index 4f8cb4db0..6160b358b 100644
--- a/lib/jit/generation/tools/vector_types.hpp
+++ b/lib/jit/generation/tools/vector_types.hpp
@@ -43,9 +43,9 @@ inline std::string access_vector_type(std::string const & v, int i)
     }
 }
 
-inline std::string access_vector_type(std::string const & v, int i, unsigned int simd_width)
+inline std::string access_vector_type(std::string const & v, int i, unsigned int vwidth)
 {
-    if(simd_width==1)
+    if(vwidth==1)
       return v;
     else
       return access_vector_type(v, i);
@@ -59,40 +59,40 @@ inline std::string append_width(std::string const & str, unsigned int width)
 }
 
 
-inline std::string vstore(unsigned int simd_width, std::string const & dtype, std::string const & value, std::string const & offset, std::string const & ptr, std::string const & stride, driver::backend_type backend, bool aligned = true)
+inline std::string vstore(unsigned int vwidth, std::string const & dtype, std::string const & value, std::string const & offset, std::string const & ptr, std::string const & stride, driver::backend_type backend, bool aligned = true)
 {
-    std::string vdtype = append_width(dtype,simd_width);
-    if (simd_width==1)
+    std::string vdtype = append_width(dtype,vwidth);
+    if (vwidth==1)
       return "(" + ptr + ")[" + offset + "] = " + value;
     else
     {
         if(backend == driver::CUDA && stride == "1" && aligned)
           return "reinterpret_cast<" + vdtype + "*>(" + ptr + ")[" + offset + "] = " + value;
         else if(backend == driver::OPENCL && stride == "1")
-          return append_width("vstore", simd_width) + "(" + value + ", " + offset + ", " + ptr + ")";
+          return append_width("vstore", vwidth) + "(" + value + ", " + offset + ", " + ptr + ")";
         else
         {
           std::string stridestr = (stride=="1")?"":("*" + stride);
           std::string res;
-          for(unsigned int s = 0 ; s < simd_width ; ++s)
-              res +=  (s>0?";(":"(") + ptr + ")[" + offset + "*" + tools::to_string(simd_width) + " + " + tools::to_string(s) + stridestr + "] = " + access_vector_type(value, s);
+          for(unsigned int s = 0 ; s < vwidth ; ++s)
+              res +=  (s>0?";(":"(") + ptr + ")[" + offset + "*" + tools::to_string(vwidth) + " + " + tools::to_string(s) + stridestr + "] = " + access_vector_type(value, s);
           return res;
         }
     }
 }
 
 
-inline std::string vload(unsigned int simd_width, std::string const & dtype, std::string const & offset, std::string const & ptr, std::string const & stride, driver::backend_type backend, bool aligned = true)
+inline std::string vload(unsigned int vwidth, std::string const & dtype, std::string const & offset, std::string const & ptr, std::string const & stride, driver::backend_type backend, bool aligned = true)
 {
-    std::string vdtype = append_width(dtype,simd_width);
-    if (simd_width==1)
+    std::string vdtype = append_width(dtype,vwidth);
+    if (vwidth==1)
       return "(" + ptr + ")[" + offset + "]";
     else
     {
       if(backend == driver::CUDA && stride == "1" && aligned)
           return "reinterpret_cast<" + vdtype + "*>(" + ptr + ")[" + offset + "]";
       else if(backend == driver::OPENCL && stride == "1")
-          return append_width("vload", simd_width) + "(" + offset + ", " + ptr + ")";
+          return append_width("vload", vwidth) + "(" + offset + ", " + ptr + ")";
       else
       {
         std::string stridestr = (stride=="1")?"":("*" + stride);
@@ -101,8 +101,8 @@ inline std::string vload(unsigned int simd_width, std::string const & dtype, std
           res = "make_" + vdtype + "(";
         else
           res = "(" + vdtype + ")(";
-        for(unsigned int s = 0 ; s < simd_width ; ++s)
-            res += ((s>0)?",(":"(") + ptr + ")[" + offset + "*" + tools::to_string(simd_width) + " + " + tools::to_string(s) + stridestr  + "]";
+        for(unsigned int s = 0 ; s < vwidth ; ++s)
+            res += ((s>0)?",(":"(") + ptr + ")[" + offset + "*" + tools::to_string(vwidth) + " + " + tools::to_string(s) + stridestr  + "]";
         res += ")";
         return res;
       }
diff --git a/lib/runtime/inference/profiles.cpp b/lib/runtime/inference/profiles.cpp
index 66e912cd7..a15c20e2c 100644
--- a/lib/runtime/inference/profiles.cpp
+++ b/lib/runtime/inference/profiles.cpp
@@ -142,7 +142,7 @@ profiles::value_type::templates_container const & profiles::value_type::template
 
 std::shared_ptr<templates::base> profiles::create(std::string const & template_name, std::vector<int> const & x)
 {
-  templates::fetching_policy_type fetch[] = {templates::FETCH_FROM_LOCAL, templates::FETCH_FROM_GLOBAL_STRIDED, templates::FETCH_FROM_GLOBAL_CONTIGUOUS};
+  templates::fetch_type fetch[] = {templates::FETCH_FROM_LOCAL, templates::FETCH_FROM_GLOBAL_STRIDED, templates::FETCH_FROM_GLOBAL_CONTIGUOUS};
   if(template_name=="elementwise_1d")
     return std::shared_ptr<templates::base>(new templates::elementwise_1d(x[0], x[1], x[2], fetch[x[3]]));
   else if(template_name=="reduce_1d")