   void operator()(agency::parallel_group<agency::cuda::concurrent_agent>& self)
   {
     printf("hello from {%d, %d}\n", (int)self.outer().index(), (int)self.inner().index());
+  }
+
+  __device__
+  void operator()(agency::parallel_group_2d<agency::cuda::concurrent_agent_2d>& self)
+  {
   }
 };
 

#include <agency/cuda/execution_policy.hpp>
#include <cstdio>

struct kernel
{

}


//using parallel_execution_policy_2d = agency::cuda::detail::basic_execution_policy<agency::cuda::parallel_agent_2d, agency::cuda::parallel_executor>;
//const parallel_execution_policy_2d par_2d{};
//
//using concurrent_execution_policy_2d = agency::cuda::detail::basic_execution_policy<agency::cuda::concurrent_agent_2d, agency::cuda::block_executor>;
//const concurrent_execution_policy_2d con_2d{};
//
//// XXX eliminate the need for these specializations
//namespace agency
//{
//
//template<>
//struct is_execution_policy<parallel_execution_policy_2d> : std::true_type {};
//
//template<>
//struct is_execution_policy<concurrent_execution_policy_2d> : std::true_type {};
//
//}
//
//
//auto grid(agency::size2 num_blocks, agency::size2 num_threads)
//  -> decltype(
//       par_2d(num_blocks, con_2d(num_threads))
//     )
//{
//  return par_2d(num_blocks, con_2d(num_threads));
//}


int main()
{
  bulk_invoke(grid(2,32), kernel());

  // XXX this takes too long to compile
  //agency::cuda::bulk_invoke(grid({1,2}, {1,32}), kernel());

  return 0;
}

