Exemple #1
0
/**
* In this tutorial we do not need additional auxiliary functions, allowing us to start right with main():
**/
int main()
{
  //Change this type definition to double if your gpu supports that
  typedef float       ScalarType;

  /**
  * <h2> Scalar Operations </h2>
  *
  * Although usually not very efficient because of PCI-Express latency, ViennaCL enables you to directly manipulate individual scalar values.
  * As such, a viennacl::scalar<double> behaves very similar to a normal `double`.
  *
  * Let us define a few CPU and ViennaCL scalars:
  *
  **/
  ScalarType s1 = ScalarType(3.1415926);   //note: writing ScalarType s1 = 3.1415926; leads to warnings with some compilers if ScalarType is 'float'.
  ScalarType s2 = ScalarType(2.71763);
  ScalarType s3 = ScalarType(42.0);

  viennacl::scalar<ScalarType> vcl_s1;
  viennacl::scalar<ScalarType> vcl_s2 = ScalarType(1.0);
  viennacl::scalar<ScalarType> vcl_s3 = ScalarType(1.0);

  /**
  * CPU scalars can be transparently assigned to GPU scalars and vice versa:
  **/
  std::cout << "Copying a few scalars..." << std::endl;
  vcl_s1 = s1;
  s2 = vcl_s2;
  vcl_s3 = s3;

  /**
  * Operations between GPU scalars work just as for CPU scalars:
  * (Note that such single compute kernels on the GPU are considerably slower than on the CPU)
  **/

  std::cout << "Manipulating a few scalars..." << std::endl;
  std::cout << "operator +=" << std::endl;
  s1 += s2;
  vcl_s1 += vcl_s2;

  std::cout << "operator *=" << std::endl;
  s1 *= s2;
  vcl_s1 *= vcl_s2;

  std::cout << "operator -=" << std::endl;
  s1 -= s2;
  vcl_s1 -= vcl_s2;

  std::cout << "operator /=" << std::endl;
  s1 /= s2;
  vcl_s1 /= vcl_s2;

  std::cout << "operator +" << std::endl;
  s1 = s2 + s3;
  vcl_s1 = vcl_s2 + vcl_s3;

  std::cout << "multiple operators" << std::endl;
  s1 = s2 + s3 * s2 - s3 / s1;
  vcl_s1 = vcl_s2 + vcl_s3 * vcl_s2 - vcl_s3 / vcl_s1;


  /**
  * Operations can also be mixed:
  **/
  std::cout << "mixed operations" << std::endl;
  vcl_s1 = s1 * vcl_s2 + s3 - vcl_s3;


  /**
  * The output stream is overloaded as well for direct printing to e.g. a terminal:
  **/

  std::cout << "CPU scalar s3: " << s3 << std::endl;
  std::cout << "GPU scalar vcl_s3: " << vcl_s3 << std::endl;


  /**
  *  <h2>Vector Operations
  *
  * Define a few vectors (from STL and plain C) and viennacl::vectors
  **/
  std::vector<ScalarType>      std_vec1(10);
  std::vector<ScalarType>      std_vec2(10);
  ScalarType                   plain_vec3[10];  //plain C array

  viennacl::vector<ScalarType> vcl_vec1(10);
  viennacl::vector<ScalarType> vcl_vec2(10);
  viennacl::vector<ScalarType> vcl_vec3(10);

  /**
  * Let us fill the CPU vectors with random values:
  * (random<> is a helper function from Random.hpp)
  **/

  for (unsigned int i = 0; i < 10; ++i)
  {
    std_vec1[i] = random<ScalarType>();
    vcl_vec2(i) = random<ScalarType>();  //also works for GPU vectors, but is MUCH slower (approx. factor 10.000) than the CPU analogue
    plain_vec3[i] = random<ScalarType>();
  }

  /**
  * Copy the CPU vectors to the GPU vectors and vice versa
  **/
  viennacl::copy(std_vec1.begin(), std_vec1.end(), vcl_vec1.begin()); //either the STL way
  viennacl::copy(vcl_vec2.begin(), vcl_vec2.end(), std_vec2.begin()); //either the STL way
  viennacl::copy(vcl_vec2, std_vec2);                                 //using the short hand notation for objects that provide .begin() and .end() members
  viennacl::copy(vcl_vec2.begin(), vcl_vec2.end(), plain_vec3);       //copy to plain C vector

  /**
  * Also partial copies by providing the corresponding iterators are possible:
  **/
  viennacl::copy(std_vec1.begin() + 4, std_vec1.begin() + 8, vcl_vec1.begin() + 4);   //cpu to gpu
  viennacl::copy(vcl_vec1.begin() + 4, vcl_vec1.begin() + 8, vcl_vec2.begin() + 1);   //gpu to gpu
  viennacl::copy(vcl_vec1.begin() + 4, vcl_vec1.begin() + 8, std_vec1.begin() + 1);   //gpu to cpu

  /**
  * Compute the inner product of two GPU vectors and write the result to either CPU or GPU
  **/
  vcl_s1 = viennacl::linalg::inner_prod(vcl_vec1, vcl_vec2);
  s1 = viennacl::linalg::inner_prod(vcl_vec1, vcl_vec2);
  s2 = viennacl::linalg::inner_prod(std_vec1, std_vec2); //inner prod can also be used with std::vector (computations are carried out on CPU then)

  /**
  * Compute norms:
  **/
  s1 = viennacl::linalg::norm_1(vcl_vec1);
  vcl_s2 = viennacl::linalg::norm_2(vcl_vec2);
  s3 = viennacl::linalg::norm_inf(vcl_vec3);


  /**
  * Plane rotation of two vectors:
  **/
  viennacl::linalg::plane_rotation(vcl_vec1, vcl_vec2, 1.1f, 2.3f);

  /**
  * Use viennacl::vector via the overloaded operators just as you would write it on paper:
  **/

  //simple expression:
  vcl_vec1 = vcl_s1 * vcl_vec2 / vcl_s3;

  //more complicated expression:
  vcl_vec1 = vcl_vec2 / vcl_s3 + vcl_s2 * (vcl_vec1 - vcl_s2 * vcl_vec2);


  /**
  * Swap the content of two vectors without a temporary vector:
  **/
  viennacl::swap(vcl_vec1, vcl_vec2);  //swaps all entries in memory
  viennacl::fast_swap(vcl_vec1, vcl_vec2); //swaps OpenCL memory handles only

  /**
  * The vectors can also be cleared directly:
  **/
  vcl_vec1.clear();
  vcl_vec2.clear();

  /**
  *  That's it, the tutorial is completed.
  **/
  std::cout << "!!!! TUTORIAL COMPLETED SUCCESSFULLY !!!!" << std::endl;

  return EXIT_SUCCESS;
}
Exemple #2
0
int run_vector_benchmark(test_config & config, viennacl::io::parameter_database& paras)
{
  typedef viennacl::scalar<ScalarType>   VCLScalar;
  typedef viennacl::vector<ScalarType>   VCLVector;
   
  ////////////////////////////////////////////////////////////////////
  //set up a little bit of data to play with:
  //ScalarType std_result = 0;
   
  ScalarType std_factor1 = static_cast<ScalarType>(3.1415);
  ScalarType std_factor2 = static_cast<ScalarType>(42.0);
  viennacl::scalar<ScalarType> vcl_factor1(std_factor1);
  viennacl::scalar<ScalarType> vcl_factor2(std_factor2);
  
  std::vector<ScalarType> std_vec1(BENCHMARK_VECTOR_SIZE);  //used to set all values to zero
  VCLVector vcl_vec1(BENCHMARK_VECTOR_SIZE);
  VCLVector vcl_vec2(BENCHMARK_VECTOR_SIZE); 
  VCLVector vcl_vec3(BENCHMARK_VECTOR_SIZE); 

  viennacl::copy(std_vec1, vcl_vec1); //initialize vectors with all zeros (no need to worry about overflows then)
  viennacl::copy(std_vec1, vcl_vec2); //initialize vectors with all zeros (no need to worry about overflows then)
  
  typedef test_data<VCLScalar, VCLVector>   TestDataType;
  test_data<VCLScalar, VCLVector> data(vcl_factor1, vcl_vec1, vcl_vec2, vcl_vec3);

  //////////////////////////////////////////////////////////
  ///////////// Start parameter recording  /////////////////
  //////////////////////////////////////////////////////////
  
  typedef std::map< double, std::pair<unsigned int, unsigned int> >   TimingType;
  std::map< std::string, TimingType > all_timings;
  
  // vector addition  
  std::cout << "------- Related to vector addition ----------" << std::endl;
  config.kernel_name("add");                    optimize_full(paras, all_timings[config.kernel_name()], vector_add<TestDataType>, config, data);
  config.kernel_name("inplace_add");            optimize_full(paras, all_timings[config.kernel_name()], vector_inplace_add<TestDataType>, config, data);
  config.kernel_name("mul_add");                optimize_full(paras, all_timings[config.kernel_name()], vector_mul_add<TestDataType>, config, data);
  config.kernel_name("cpu_mul_add");            optimize_full(paras, all_timings[config.kernel_name()], vector_cpu_mul_add<TestDataType>, config, data);
  config.kernel_name("inplace_mul_add");        optimize_full(paras, all_timings[config.kernel_name()], vector_inplace_mul_add<TestDataType>, config, data);
  config.kernel_name("cpu_inplace_mul_add");    optimize_full(paras, all_timings[config.kernel_name()], vector_cpu_inplace_mul_add<TestDataType>, config, data);
  config.kernel_name("inplace_div_add");        optimize_full(paras, all_timings[config.kernel_name()], vector_inplace_div_add<TestDataType>, config, data);

  std::cout << "------- Related to vector subtraction ----------" << std::endl;
  config.kernel_name("sub");                    optimize_full(paras, all_timings[config.kernel_name()], vector_sub<TestDataType>, config, data);
  config.kernel_name("inplace_sub");            optimize_full(paras, all_timings[config.kernel_name()], vector_inplace_sub<TestDataType>, config, data);
  config.kernel_name("mul_sub");                optimize_full(paras, all_timings[config.kernel_name()], vector_mul_sub<TestDataType>, config, data);
  config.kernel_name("inplace_mul_sub");        optimize_full(paras, all_timings[config.kernel_name()], vector_inplace_mul_sub<TestDataType>, config, data);
  config.kernel_name("inplace_div_sub");        optimize_full(paras, all_timings[config.kernel_name()], vector_inplace_div_sub<TestDataType>, config, data);

  std::cout << "------- Related to vector scaling (mult/div) ----------" << std::endl;
  config.kernel_name("mult");                   optimize_full(paras, all_timings[config.kernel_name()], vector_mult<TestDataType>, config, data);
  config.kernel_name("inplace_mult");           optimize_full(paras, all_timings[config.kernel_name()], vector_inplace_mult<TestDataType>, config, data);
  config.kernel_name("cpu_mult");               optimize_full(paras, all_timings[config.kernel_name()], vector_cpu_mult<TestDataType>, config, data);
  config.kernel_name("cpu_inplace_mult");       optimize_full(paras, all_timings[config.kernel_name()], vector_cpu_inplace_mult<TestDataType>, config, data);
  config.kernel_name("divide");                 optimize_full(paras, all_timings[config.kernel_name()], vector_divide<TestDataType>, config, data);
  config.kernel_name("inplace_divide");         optimize_full(paras, all_timings[config.kernel_name()], vector_inplace_divide<TestDataType>, config, data);
 
  std::cout << "------- Others ----------" << std::endl;
  config.kernel_name("inner_prod");             optimize_full(paras, all_timings[config.kernel_name()], vector_inner_prod<TestDataType>, config, data);
  config.kernel_name("swap");                   optimize_full(paras, all_timings[config.kernel_name()], vector_swap<TestDataType>, config, data);
  config.kernel_name("clear");                  optimize_full(paras, all_timings[config.kernel_name()], vector_clear<TestDataType>, config, data);
  config.kernel_name("plane_rotation");         optimize_full(paras, all_timings[config.kernel_name()], vector_plane_rotation<TestDataType>, config, data);
  
  //config.max_work_groups(32); //otherwise failures on 8500 GT
  config.kernel_name("norm_1");                 optimize_restricted(paras, all_timings[config.kernel_name()], vector_norm_1<TestDataType>, config, data);
  config.kernel_name("norm_2");                 optimize_restricted(paras, all_timings[config.kernel_name()], vector_norm_2<TestDataType>, config, data);
  config.kernel_name("norm_inf");               optimize_restricted(paras, all_timings[config.kernel_name()], vector_norm_inf<TestDataType>, config, data);

  
  //restricted optimizations:
  config.kernel_name("index_norm_inf");         optimize_restricted(paras, all_timings[config.kernel_name()], vector_index_norm_inf<TestDataType>, config, data);
  
  
  return 0;
}
Exemple #3
0
int run_matrix_benchmark(test_config & config, viennacl::io::parameter_database& paras)
{
  typedef viennacl::scalar<ScalarType>   VCLScalar;
  typedef viennacl::vector<ScalarType>   VCLVector;
  typedef viennacl::matrix<ScalarType>   VCLMatrix;
   
  ////////////////////////////////////////////////////////////////////
  //set up a little bit of data to play with:
  //ScalarType std_result = 0;
   
  ScalarType std_factor1 = static_cast<ScalarType>(3.1415);
  ScalarType std_factor2 = static_cast<ScalarType>(42.0);
  viennacl::scalar<ScalarType> vcl_factor1(std_factor1);
  viennacl::scalar<ScalarType> vcl_factor2(std_factor2);
  
  std::vector<ScalarType> std_vec1(BENCHMARK_MATRIX_SIZE);  //used to set all values to zero
  std::vector< std::vector<ScalarType> > stl_mat(BENCHMARK_MATRIX_SIZE);  //store identity matrix here
  VCLVector vcl_vec1(BENCHMARK_MATRIX_SIZE);
  VCLVector vcl_vec2(BENCHMARK_MATRIX_SIZE);
  VCLMatrix vcl_mat(BENCHMARK_MATRIX_SIZE, BENCHMARK_MATRIX_SIZE);
  
  for (int i=0; i<BENCHMARK_MATRIX_SIZE; ++i)
  {
    stl_mat[i].resize(BENCHMARK_MATRIX_SIZE);
    stl_mat[i][i] = 1.0;
  }

  copy(std_vec1, vcl_vec1); //initialize vectors with all zeros (no need to worry about overflows then)
  copy(std_vec1, vcl_vec2); //initialize vectors with all zeros (no need to worry about overflows then)
  copy(stl_mat, vcl_mat);
  
  typedef test_data<VCLScalar, VCLVector, VCLMatrix>   TestDataType;
  test_data<VCLScalar, VCLVector, VCLMatrix> data(vcl_factor1, vcl_vec1, vcl_vec2, vcl_mat);

  //////////////////////////////////////////////////////////
  ///////////// Start parameter recording  /////////////////
  //////////////////////////////////////////////////////////
  
  typedef std::map< double, std::pair<unsigned int, unsigned int> >   TimingType;
  std::map< std::string, TimingType > all_timings;
  
  std::cout << "------- Related to direct solution algorithms ----------" << std::endl;
  config.kernel_name("trans_lower_triangular_substitute_inplace");
  optimize_restricted(paras, all_timings[config.kernel_name()],
                      matrix_trans_lower_triangular_substitute_inplace<TestDataType>, config, data);

  config.kernel_name("lower_triangular_substitute_inplace");
  optimize_restricted(paras, all_timings[config.kernel_name()],
                      matrix_lower_triangular_substitute_inplace<TestDataType>, config, data);

  config.kernel_name("unit_lower_triangular_substitute_inplace");
  optimize_restricted(paras, all_timings[config.kernel_name()],
                      matrix_unit_lower_triangular_substitute_inplace<TestDataType>, config, data);

  config.kernel_name("upper_triangular_substitute_inplace");
  optimize_restricted(paras, all_timings[config.kernel_name()],
                      matrix_upper_triangular_substitute_inplace<TestDataType>, config, data);

  config.kernel_name("trans_upper_triangular_substitute_inplace");
  optimize_restricted(paras, all_timings[config.kernel_name()],
                      matrix_trans_upper_triangular_substitute_inplace<TestDataType>, config, data);

  config.kernel_name("unit_upper_triangular_substitute_inplace");
  optimize_restricted(paras, all_timings[config.kernel_name()],
                      matrix_unit_upper_triangular_substitute_inplace<TestDataType>, config, data);

  config.kernel_name("lu_factorize");
  optimize_restricted(paras, all_timings[config.kernel_name()],
                      matrix_lu_factorize<TestDataType>, config, data);

  //other kernels:
  std::cout << "------- Related to other operations ----------" << std::endl;
  config.kernel_name("rank1_update");
  optimize_restricted(paras, all_timings[config.kernel_name()],
                      matrix_rank1_update<TestDataType>, config, data);

  config.kernel_name("scaled_rank1_update");
  optimize_restricted(paras, all_timings[config.kernel_name()],
                      matrix_scaled_rank1_update<TestDataType>, config, data);

  config.kernel_name("vec_mul");
  optimize_restricted(paras, all_timings[config.kernel_name()],
                      matrix_vec_mul<TestDataType>, config, data);

  config.kernel_name("trans_vec_mul");
  optimize_restricted(paras, all_timings[config.kernel_name()],
                      matrix_trans_vec_mul<TestDataType>, config, data);

  return 0;
}
int run_matrix_benchmark(test_config & config, viennacl::io::parameter_database& paras)
{
  typedef viennacl::scalar<ScalarType>   VCLScalar;
  typedef viennacl::vector<ScalarType>   VCLVector;
  typedef viennacl::compressed_matrix<ScalarType>   VCLMatrix;
   
  ////////////////////////////////////////////////////////////////////
  //set up a little bit of data to play with:
  //ScalarType std_result = 0;
   
  ScalarType std_factor1 = static_cast<ScalarType>(3.1415);
  ScalarType std_factor2 = static_cast<ScalarType>(42.0);
  viennacl::scalar<ScalarType> vcl_factor1(std_factor1);
  viennacl::scalar<ScalarType> vcl_factor2(std_factor2);
  
  std::vector<ScalarType> std_vec1(BENCHMARK_MATRIX_SIZE);  //used to set all values to zero
  std::vector< std::map< unsigned int, ScalarType> > stl_mat(BENCHMARK_MATRIX_SIZE);  //store identity matrix here
  VCLVector vcl_vec1(BENCHMARK_MATRIX_SIZE);
  VCLVector vcl_vec2(BENCHMARK_MATRIX_SIZE);
  VCLMatrix vcl_mat(BENCHMARK_MATRIX_SIZE, BENCHMARK_MATRIX_SIZE);
  
  for (int i=0; i<BENCHMARK_MATRIX_SIZE; ++i)
  {
      if (i > 10)
      {
          stl_mat[i][i - 10] = 1.0;
          stl_mat[i][i - 7] = 1.0;
          stl_mat[i][i - 4] = 1.0;
          stl_mat[i][i - 2] = 1.0;
      }
      stl_mat[i][i] = 1.0;
      if (i + 10 < BENCHMARK_MATRIX_SIZE)
      {
          stl_mat[i][i + 5] = 1.0;
          stl_mat[i][i + 7] = 1.0;
          stl_mat[i][i + 9] = 1.0;
          stl_mat[i][i + 10] = 1.0;
      }
  }

  viennacl::copy(std_vec1, vcl_vec1); //initialize vectors with all zeros (no need to worry about overflows then)
  viennacl::copy(std_vec1, vcl_vec2); //initialize vectors with all zeros (no need to worry about overflows then)
  viennacl::copy(stl_mat, vcl_mat);
  
  typedef test_data<VCLScalar, VCLVector, VCLMatrix>   TestDataType;
  test_data<VCLScalar, VCLVector, VCLMatrix> data(vcl_factor1, vcl_vec1, vcl_vec2, vcl_mat);

  //////////////////////////////////////////////////////////
  ///////////// Start parameter recording  /////////////////
  //////////////////////////////////////////////////////////
  
  typedef std::map< double, std::pair<unsigned int, unsigned int> >   TimingType;
  std::map< std::string, TimingType > all_timings;
  

  //other kernels:
  std::cout << "------- Related to other operations ----------" << std::endl;

  config.kernel_name("vec_mul");
  optimize_full(paras, all_timings[config.kernel_name()],
                      matrix_vec_mul<TestDataType>, config, data);


  return 0;
}
int main()
{
  typedef float       ScalarType;

  viennacl::vector<ScalarType> vcl_vec1(10);
  viennacl::vector<ScalarType> vcl_vec2(10);
  viennacl::vector<ScalarType> vcl_vec3(10);

  //
  // Let us fill the CPU vectors with random values:
  // (random<> is a helper function from Random.hpp)
  //

  for (unsigned int i = 0; i < 10; ++i)
  {
    vcl_vec1[i] = ScalarType(i);
    vcl_vec2[i] = ScalarType(10 - i);
  }

  //
  // Build expression graph for the operation vcl_vec3 = vcl_vec1 + vcl_vec2
  //
  // This requires the following expression graph:
  //
  //             ( = )
  //            /      |
  //    vcl_vec3      ( + )
  //                 /     |
  //           vcl_vec1    vcl_vec2
  //
  // One expression node consists of two leaves and the operation connecting the two.
  // Here we thus need two nodes: One for {vcl_vec3, = , link}, where 'link' points to the second node
  // {vcl_vec1, +, vcl_vec2}.
  //
  // The following is the lowest level on which one could build the expression tree.
  // Even for a C API one would introduce some additional convenience layer such as add_vector_float_to_lhs(...); etc.
  //
  typedef viennacl::scheduler::statement::container_type   NodeContainerType;   // this is just std::vector<viennacl::scheduler::statement_node>
  NodeContainerType expression_nodes(2);                                        //container with two nodes

  ////// First node //////

  // specify LHS of first node, i.e. vcl_vec3:
  expression_nodes[0].lhs.type_family  = viennacl::scheduler::VECTOR_TYPE_FAMILY;   // family of vectors
  expression_nodes[0].lhs.subtype      = viennacl::scheduler::DENSE_VECTOR_TYPE;    // a dense vector
  expression_nodes[0].lhs.numeric_type = viennacl::scheduler::FLOAT_TYPE;           // vector consisting of floats
  expression_nodes[0].lhs.vector_float = &vcl_vec3;                                 // provide pointer to vcl_vec3;

  // specify assignment operation for this node:
  expression_nodes[0].op.type_family   = viennacl::scheduler::OPERATION_BINARY_TYPE_FAMILY; // this is a binary operation, so both LHS and RHS operands are important
  expression_nodes[0].op.type          = viennacl::scheduler::OPERATION_BINARY_ASSIGN_TYPE; // assignment operation: '='

  // specify RHS: Just refer to the second node:
  expression_nodes[0].rhs.type_family  = viennacl::scheduler::COMPOSITE_OPERATION_FAMILY; // this links to another node (no need to set .subtype and .numeric_type)
  expression_nodes[0].rhs.node_index   = 1;                                               // index of the other node

  ////// Second node //////

  // LHS
  expression_nodes[1].lhs.type_family  = viennacl::scheduler::VECTOR_TYPE_FAMILY;   // family of vectors
  expression_nodes[1].lhs.subtype      = viennacl::scheduler::DENSE_VECTOR_TYPE;    // a dense vector
  expression_nodes[1].lhs.numeric_type = viennacl::scheduler::FLOAT_TYPE;           // vector consisting of floats
  expression_nodes[1].lhs.vector_float = &vcl_vec1;                                 // provide pointer to vcl_vec1

  // OP
  expression_nodes[1].op.type_family   = viennacl::scheduler::OPERATION_BINARY_TYPE_FAMILY; // this is a binary operation, so both LHS and RHS operands are important
  expression_nodes[1].op.type          = viennacl::scheduler::OPERATION_BINARY_ADD_TYPE;    // addition operation: '+'

  // RHS
  expression_nodes[1].rhs.type_family  = viennacl::scheduler::VECTOR_TYPE_FAMILY;  // family of vectors
  expression_nodes[1].rhs.subtype      = viennacl::scheduler::DENSE_VECTOR_TYPE;   // a dense vector
  expression_nodes[1].rhs.numeric_type = viennacl::scheduler::FLOAT_TYPE;          // vector consisting of floats
  expression_nodes[1].rhs.vector_float = &vcl_vec2;                                // provide pointer to vcl_vec2


  // create the full statement (aka. single line of code such as vcl_vec3 = vcl_vec1 + vcl_vec2):
  viennacl::scheduler::statement vec_addition(expression_nodes);

  // print it
  std::cout << vec_addition << std::endl;

  // run it
  viennacl::scheduler::execute(vec_addition);

  // print vectors
  std::cout << "vcl_vec1: " << vcl_vec1 << std::endl;
  std::cout << "vcl_vec2: " << vcl_vec2 << std::endl;
  std::cout << "vcl_vec3: " << vcl_vec3 << std::endl;


  std::cout << "!!!! TUTORIAL COMPLETED SUCCESSFULLY !!!!" << std::endl;

  return EXIT_SUCCESS;
}
/**
* With this let us go right to main():
**/
int main()
{
  typedef float       ScalarType;


  /**
  * <h2>Part 1: Set up a custom context</h2>
  *
  * The following is rather lengthy because OpenCL is a fairly low-level framework.
  * For comparison, the subsequent code explicitly performs the OpenCL setup that is done
  * in the background within the 'custom_kernels'-tutorial
  **/

  //manually set up a custom OpenCL context:
  std::vector<cl_device_id> device_id_array;

  //get all available devices
  viennacl::ocl::platform pf;
  std::cout << "Platform info: " << pf.info() << std::endl;
  std::vector<viennacl::ocl::device> devices = pf.devices(CL_DEVICE_TYPE_DEFAULT);
  std::cout << devices[0].name() << std::endl;
  std::cout << "Number of devices for custom context: " << devices.size() << std::endl;

  //set up context using all found devices:
  for (std::size_t i=0; i<devices.size(); ++i)
  {
      device_id_array.push_back(devices[i].id());
  }

  std::cout << "Creating context..." << std::endl;
  cl_int err;
  cl_context my_context = clCreateContext(0, cl_uint(device_id_array.size()), &(device_id_array[0]), NULL, NULL, &err);
  VIENNACL_ERR_CHECK(err);


  //create two Vectors:
  unsigned int vector_size = 10;
  std::vector<ScalarType> vec1(vector_size);
  std::vector<ScalarType> vec2(vector_size);
  std::vector<ScalarType> result(vector_size);

  //
  // fill the operands vec1 and vec2:
  //
  for (unsigned int i=0; i<vector_size; ++i)
  {
    vec1[i] = static_cast<ScalarType>(i);
    vec2[i] = static_cast<ScalarType>(vector_size-i);
  }

  //
  // create memory in OpenCL context:
  //
  cl_mem mem_vec1 = clCreateBuffer(my_context, CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR, vector_size * sizeof(ScalarType), &(vec1[0]), &err);
  VIENNACL_ERR_CHECK(err);
  cl_mem mem_vec2 = clCreateBuffer(my_context, CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR, vector_size * sizeof(ScalarType), &(vec2[0]), &err);
  VIENNACL_ERR_CHECK(err);
  cl_mem mem_result = clCreateBuffer(my_context, CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR, vector_size * sizeof(ScalarType), &(result[0]), &err);
  VIENNACL_ERR_CHECK(err);

  //
  // create a command queue for each device:
  //

  std::vector<cl_command_queue> queues(devices.size());
  for (std::size_t i=0; i<devices.size(); ++i)
  {
    queues[i] = clCreateCommandQueue(my_context, devices[i].id(), 0, &err);
    VIENNACL_ERR_CHECK(err);
  }

  //
  // create and build a program in the context:
  //
  std::size_t source_len = std::string(my_compute_program).length();
  cl_program my_prog = clCreateProgramWithSource(my_context, 1, &my_compute_program, &source_len, &err);
  err = clBuildProgram(my_prog, 0, NULL, NULL, NULL, NULL);

/*            char buffer[1024];
            cl_build_status status;
            clGetProgramBuildInfo(my_prog, devices[1].id(), CL_PROGRAM_BUILD_STATUS, sizeof(cl_build_status), &status, NULL);
            clGetProgramBuildInfo(my_prog, devices[1].id(), CL_PROGRAM_BUILD_LOG, sizeof(char)*1024, &buffer, NULL);
            std::cout << "Build Scalar: Err = " << err << " Status = " << status << std::endl;
            std::cout << "Log: " << buffer << std::endl;*/

  VIENNACL_ERR_CHECK(err);

  //
  // create a kernel from the program:
  //
  const char * kernel_name = "elementwise_prod";
  cl_kernel my_kernel = clCreateKernel(my_prog, kernel_name, &err);
  VIENNACL_ERR_CHECK(err);


  //
  // Execute elementwise_prod kernel on first queue: result = vec1 .* vec2;
  //
  err = clSetKernelArg(my_kernel, 0, sizeof(cl_mem), (void*)&mem_vec1);
  VIENNACL_ERR_CHECK(err);
  err = clSetKernelArg(my_kernel, 1, sizeof(cl_mem), (void*)&mem_vec2);
  VIENNACL_ERR_CHECK(err);
  err = clSetKernelArg(my_kernel, 2, sizeof(cl_mem), (void*)&mem_result);
  VIENNACL_ERR_CHECK(err);
  err = clSetKernelArg(my_kernel, 3, sizeof(unsigned int), (void*)&vector_size);
  VIENNACL_ERR_CHECK(err);
  std::size_t global_size = vector_size;
  std::size_t local_size = vector_size;
  err = clEnqueueNDRangeKernel(queues[0], my_kernel, 1, NULL, &global_size, &local_size, 0, NULL, NULL);
  VIENNACL_ERR_CHECK(err);


  //
  // Read and output result:
  //
  err = clEnqueueReadBuffer(queues[0], mem_vec1, CL_TRUE, 0, sizeof(ScalarType)*vector_size, &(vec1[0]), 0, NULL, NULL);
  VIENNACL_ERR_CHECK(err);
  err = clEnqueueReadBuffer(queues[0], mem_result, CL_TRUE, 0, sizeof(ScalarType)*vector_size, &(result[0]), 0, NULL, NULL);
  VIENNACL_ERR_CHECK(err);

  std::cout << "vec1  : ";
  for (std::size_t i=0; i<vec1.size(); ++i)
    std::cout << vec1[i] << " ";
  std::cout << std::endl;

  std::cout << "vec2  : ";
  for (std::size_t i=0; i<vec2.size(); ++i)
    std::cout << vec2[i] << " ";
  std::cout << std::endl;

  std::cout << "result: ";
  for (std::size_t i=0; i<result.size(); ++i)
    std::cout << result[i] << " ";
  std::cout << std::endl;

  /**
  * <h2>Part 2: Reuse Custom OpenCL Context with ViennaCL</h2>
  *
  * To let ViennaCL reuse the previously created context, we need to make it known to ViennaCL \em before any ViennaCL objects are created.
  * We inject the custom context as the context with default id '0' when using viennacl::ocl::switch_context().
  **/
  viennacl::ocl::setup_context(0, my_context, device_id_array, queues);
  viennacl::ocl::switch_context(0); //activate the new context (only mandatory with context-id not equal to zero)

  /**
  * Check that ViennaCL really uses the new context:
  **/
  std::cout << "Existing context: " << my_context << std::endl;
  std::cout << "ViennaCL uses context: " << viennacl::ocl::current_context().handle().get() << std::endl;

  /**
  * Wrap existing OpenCL objects into ViennaCL:
  **/
  viennacl::vector<ScalarType> vcl_vec1(mem_vec1, vector_size);
  viennacl::vector<ScalarType> vcl_vec2(mem_vec2, vector_size);
  viennacl::vector<ScalarType> vcl_result(mem_result, vector_size);
  viennacl::scalar<ScalarType> vcl_s = 2.0;

  std::cout << "Standard vector operations within ViennaCL:" << std::endl;
  vcl_result = vcl_s * vcl_vec1 + vcl_vec2;

  std::cout << "vec1  : ";
  std::cout << vcl_vec1 << std::endl;

  std::cout << "vec2  : ";
  std::cout << vcl_vec2 << std::endl;

  std::cout << "result: ";
  std::cout << vcl_result << std::endl;

  /**
  * We can also reuse the existing elementwise_prod kernel.
  * Therefore, we first have to make the existing program known to ViennaCL
  * For more details on the three lines, see tutorial 'custom-kernels'
  **/
  std::cout << "Using existing kernel within the OpenCL backend of ViennaCL:" << std::endl;
  viennacl::ocl::program & my_vcl_prog = viennacl::ocl::current_context().add_program(my_prog, "my_compute_program");
  viennacl::ocl::kernel & my_vcl_kernel = my_vcl_prog.add_kernel(my_kernel, "elementwise_prod");
  viennacl::ocl::enqueue(my_vcl_kernel(vcl_vec1, vcl_vec2, vcl_result, static_cast<cl_uint>(vcl_vec1.size())));  //Note that std::size_t might differ between host and device. Thus, a cast to cl_uint is necessary here.

  std::cout << "vec1  : ";
  std::cout << vcl_vec1 << std::endl;

  std::cout << "vec2  : ";
  std::cout << vcl_vec2 << std::endl;

  std::cout << "result: ";
  std::cout << vcl_result << std::endl;


  /**
  * Since a linear piece of memory can be interpreted in several ways,
  * we will now create a 3x3 row-major matrix out of the linear memory in mem_vec1/
  * The first three entries in vcl_vec2 and vcl_result are used to carry out matrix-vector products:
  **/
  viennacl::matrix<ScalarType> vcl_matrix(mem_vec1, 3, 3);

  vcl_vec2.resize(3);   //note that the resize operation leads to new memory, thus vcl_vec2 is now at a different memory location (values are copied)
  vcl_result.resize(3); //note that the resize operation leads to new memory, thus vcl_vec2 is now at a different memory location (values are copied)
  vcl_result = viennacl::linalg::prod(vcl_matrix, vcl_vec2);

  std::cout << "result of matrix-vector product: ";
  std::cout << vcl_result << std::endl;

  /**
  *  Any further operations can be carried out in the same way.
  *  Just keep in mind that any resizing of vectors or matrices leads to a reallocation of the underlying memory buffer, through which the 'wrapper' is lost.
  **/
  std::cout << "!!!! TUTORIAL COMPLETED SUCCESSFULLY !!!!" << std::endl;

  return EXIT_SUCCESS;
}
int main(int argc, char * argv[])
{
  TAU_INIT(&argc, &argv)
  TAU_PROFILE_TIMER(orio_maintimer, "main()", "int (int, char **)", TAU_USER);
  TAU_PROFILE_START(orio_maintimer);
  TAU_PROFILE_SET_NODE(0);
    std::vector<viennacl::ocl::device> devices = viennacl::ocl::platform().devices();
    std::vector<cl_device_id> my_devices;
    my_devices.push_back(devices[0].id());
    viennacl::ocl::setup_context(0L, my_devices);

 
  //Change this type definition to double if your gpu supports that
  typedef double       ScalarType;
  
  /////////////////////////////////////////////////
  ///////////// Vector operations /////////////////
  /////////////////////////////////////////////////
  
  void * orio_profiler;
  TAU_PROFILER_CREATE(orio_profiler, "orio_generated_code", "", TAU_USER);

  for(int i = 0; i < 3; ++i) {

  //
  // Define a few vectors (from STL and plain C) and viennacl::vectors
  //
  std::vector<ScalarType>      std_vec1(1000000);
  std::vector<ScalarType>      std_vec2(1000000);

  viennacl::vector<ScalarType> vcl_vec1(1000000);
  viennacl::vector<ScalarType> vcl_vec2(1000000);
  viennacl::scalar<ScalarType> vcl_s1 = ScalarType(5.0);

  //
  // Let us fill the CPU vectors with random values:
  // (random<> is a helper function from Random.hpp)
  //
  
  for (unsigned int i = 0; i < 1000000; ++i)
  {
    std_vec1[i] = random<ScalarType>(); 
    std_vec2[i] = 0.0;
  }
  
  //
  // Copy the CPU vectors to the GPU vectors and vice versa
  //
    TAU_PROFILER_START(orio_profiler);   
  viennacl::copy(std_vec1.begin(), std_vec1.end(), vcl_vec1.begin()); //either the STL way
  viennacl::copy(std_vec2.begin(), std_vec2.end(), vcl_vec2.begin()); //either the STL way

  vcl_vec2 += vcl_s1 * vcl_vec1;

  viennacl::copy(vcl_vec2.begin(), vcl_vec2.end(), std_vec2.begin());

    TAU_PROFILER_STOP(orio_profiler);
  }
  double orio_inclusive[TAU_MAX_COUNTERS];
  TAU_PROFILER_GET_INCLUSIVE_VALUES(orio_profiler, &orio_inclusive);
  printf("{'/*@ coordinate @*/' : %g}\n", orio_inclusive[0]);
  return EXIT_SUCCESS;
}
Exemple #8
0
int main()
{
  //Change this type definition to double if your gpu supports that
  typedef float       ScalarType;

  // Choose the Phi (WORKS)
  viennacl::ocl::set_context_device_type(0, viennacl::ocl::accelerator_tag());

  /////////////////////////////////////////////////
  ///////////// Scalar operations /////////////////
  /////////////////////////////////////////////////

  //
  // Define a few CPU scalars:
  //
  ScalarType s1 = static_cast<ScalarType>(3.1415926);
  ScalarType s2 = static_cast<ScalarType>(2.71763);
  ScalarType s3 = static_cast<ScalarType>(42.0);

  //
  // ViennaCL scalars are defined in the same way:
  //
  viennacl::scalar<ScalarType> vcl_s1;
  viennacl::scalar<ScalarType> vcl_s2 = static_cast<ScalarType>(1.0);
  viennacl::scalar<ScalarType> vcl_s3 = static_cast<ScalarType>(1.0);

  //
  // CPU scalars can be transparently assigned to GPU scalars and vice versa:
  //
  vcl_s1 = s1;
  s2 = vcl_s2;
  vcl_s3 = s3;

  //
  // Operations between GPU scalars work just as for CPU scalars:
  // (Note that such single compute kernels on the GPU are considerably slower than on the CPU)
  //

  s1 += s2;
  vcl_s1 += vcl_s2;

  s1 *= s2;
  vcl_s1 *= vcl_s2;

  s1 -= s2;
  vcl_s1 -= vcl_s2;

  s1 /= s2;
  vcl_s1 /= vcl_s2;

  s1 = s2 + s3;
  vcl_s1 = vcl_s2 + vcl_s3;

  s1 = s2 + s3 * s2 - s3 / s1;
  vcl_s1 = vcl_s2 + vcl_s3 * vcl_s2 - vcl_s3 / vcl_s1;


  //
  // Operations can also be mixed:
  //

  vcl_s1 = s1 * vcl_s2 + s3 - vcl_s3;


  //
  // Output stream is overloaded as well:
  //

  std::cout << "CPU scalar s2: " << s2 << std::endl;
  std::cout << "GPU scalar vcl_s2: " << vcl_s2 << std::endl;

  std::vector< viennacl::ocl::device > devices = viennacl::ocl::platform().devices();

  for (int i = 0; i < devices.size(); i++) {
      std::cout << devices[i].info() << "\n";
  }

  std::cout << "SELECTED DEVICE: \n";
  std::cout << viennacl::ocl::current_context().current_device().info() << "\n";

  /////////////////////////////////////////////////
  ///////////// Vector operations /////////////////
  /////////////////////////////////////////////////

  //
  // Define a few vectors (from STL and plain C) and viennacl::vectors
  //
  std::vector<ScalarType>      std_vec1(10);
  std::vector<ScalarType>      std_vec2(10);
  ScalarType                   plain_vec3[10];  //plain C array

  viennacl::vector<ScalarType> vcl_vec1(10);
  viennacl::vector<ScalarType> vcl_vec2(10);
  viennacl::vector<ScalarType> vcl_vec3(10);

  //
  // Let us fill the CPU vectors with random values:
  // (random<> is a helper function from Random.hpp)
  //

  for (unsigned int i = 0; i < 10; ++i)
  {
    std_vec1[i] = random<ScalarType>();
    vcl_vec2(i) = random<ScalarType>();  //also works for GPU vectors, but is MUCH slower (approx. factor 10.000) than the CPU analogue
    plain_vec3[i] = random<ScalarType>();
  }

  //
  // Copy the CPU vectors to the GPU vectors and vice versa
  //
  copy(std_vec1.begin(), std_vec1.end(), vcl_vec1.begin()); //either the STL way
  copy(vcl_vec2.begin(), vcl_vec2.end(), std_vec2.begin()); //either the STL way
  copy(vcl_vec2, std_vec2);                                 //using the short hand notation for objects that provide .begin() and .end() members
  copy(vcl_vec2.begin(), vcl_vec2.end(), plain_vec3);       //copy to plain C vector

  //
  // Compute the inner product of two GPU vectors and write the result to either CPU or GPU
  //

  vcl_s1 = viennacl::linalg::inner_prod(vcl_vec1, vcl_vec2);
  s1 = viennacl::linalg::inner_prod(vcl_vec1, vcl_vec2);

  //
  // Compute norms:
  //

  s1 = viennacl::linalg::norm_1(vcl_vec1);
  vcl_s2 = viennacl::linalg::norm_2(vcl_vec2);
  s3 = viennacl::linalg::norm_inf(vcl_vec3);

  //
  // Plane rotation of two vectors:
  //

  viennacl::linalg::plane_rotation(vcl_vec1, vcl_vec2, 1.1f, 2.3f);

  //
  // Use viennacl::vector via the overloaded operators just as you would write it on paper:
  //

  //simple expression:
  vcl_vec1 = vcl_s1 * vcl_vec2 / vcl_s3;

  //more complicated expression:
  vcl_vec1 = vcl_vec2 / vcl_s1 + vcl_s2 * (vcl_vec1 - vcl_s2 * vcl_vec2);

  //
  // Swap the content of two vectors without a temporary vector:
  //

  swap(vcl_vec1, vcl_vec2);


  //
  //  That's it. Move on to the second tutorial, where dense matrices are explained.
  //
  std::cout << "!!!! TUTORIAL 1 COMPLETED SUCCESSFULLY !!!!" << std::endl;

  exit(EXIT_SUCCESS);
  return 0;
}