All Classes Namespaces Files Functions Variables Typedefs Enumerations Enumerator Friends Macros Pages
misc_operations.hpp
Go to the documentation of this file.
#ifndef VIENNACL_LINALG_CUDA_MISC_OPERATIONS_HPP_
#define VIENNACL_LINALG_CUDA_MISC_OPERATIONS_HPP_
/* =========================================================================
Copyright (c) 2010-2016, Institute for Microelectronics,
Institute for Analysis and Scientific Computing,
TU Wien.
Portions of this software are copyright by UChicago Argonne, LLC.
-----------------
ViennaCL - The Vienna Computing Library
-----------------
Project Head: Karl Rupp rupp@iue.tuwien.ac.at
(A list of authors and contributors can be found in the manual)
License: MIT (X11), see file LICENSE in the base directory
============================================================================= */
#include "viennacl/forwards.h"
#include "viennacl/scalar.hpp"
#include "viennacl/vector.hpp"
#include "viennacl/tools/tools.hpp"
#include "viennacl/linalg/cuda/common.hpp"
namespace viennacl
{
namespace linalg
{
namespace cuda
{
namespace detail
{
template<typename NumericT>
__global__ void level_scheduling_substitute_kernel(
const unsigned int * row_index_array,
const unsigned int * row_indices,
const unsigned int * column_indices,
const NumericT * elements,
NumericT * vec,
unsigned int size)
{
for (unsigned int row = blockDim.x * blockIdx.x + threadIdx.x;
row < size;
row += gridDim.x * blockDim.x)
{
unsigned int eq_row = row_index_array[row];
NumericT vec_entry = vec[eq_row];
unsigned int row_end = row_indices[row+1];
for (unsigned int j = row_indices[row]; j < row_end; ++j)
vec_entry -= vec[column_indices[j]] * elements[j];
vec[eq_row] = vec_entry;
}
}
template<typename NumericT>
void level_scheduling_substitute(vector<NumericT> & vec,
viennacl::backend::mem_handle const & row_index_array,
viennacl::backend::mem_handle const & row_buffer,
viennacl::backend::mem_handle const & col_buffer,
viennacl::backend::mem_handle const & element_buffer,
vcl_size_t num_rows
)
{
level_scheduling_substitute_kernel<<<128, 128>>>(viennacl::cuda_arg<unsigned int>(row_index_array),
viennacl::cuda_arg<unsigned int>(row_buffer),
viennacl::cuda_arg<unsigned int>(col_buffer),
viennacl::cuda_arg<NumericT>(element_buffer),
viennacl::cuda_arg(vec),
static_cast<unsigned int>(num_rows)
);
}
} //namespace detail
} //namespace cuda
} //namespace linalg
} //namespace viennacl
#endif
tools.hpp
Various little tools used here and there in ViennaCL.
forwards.h
This file provides the forward declarations for the main types used within ViennaCL.
NumericT
float NumericT
Definition: bisect.cpp:40
viennacl::traits::size
vcl_size_t size(VectorType const &vec)
Generic routine for obtaining the size of a vector (ViennaCL, uBLAS, etc.)
Definition: size.hpp:239
viennacl::vcl_size_t
std::size_t vcl_size_t
Definition: forwards.h:75
viennacl::vector< NumericT >
viennacl::row
vector_expression< const matrix_base< NumericT, F >, const unsigned int, op_row > row(const matrix_base< NumericT, F > &A, unsigned int i)
Definition: matrix.hpp:910
common.hpp
Common routines for CUDA execution.
vector.hpp
The vector type with operator-overloads and proxy classes is defined here. Linear algebra operations ...
viennacl::linalg::cuda::detail::level_scheduling_substitute
void level_scheduling_substitute(vector< NumericT > &vec, viennacl::backend::mem_handle const &row_index_array, viennacl::backend::mem_handle const &row_buffer, viennacl::backend::mem_handle const &col_buffer, viennacl::backend::mem_handle const &element_buffer, vcl_size_t num_rows)
Definition: misc_operations.hpp:68
viennacl::backend::mem_handle
Main abstraction class for multiple memory domains. Represents a buffer in either main RAM...
Definition: mem_handle.hpp:89
viennacl::linalg::cuda::detail::level_scheduling_substitute_kernel
__global__ void level_scheduling_substitute_kernel(const unsigned int *row_index_array, const unsigned int *row_indices, const unsigned int *column_indices, const NumericT *elements, NumericT *vec, unsigned int size)
Definition: misc_operations.hpp:42
viennacl::cuda_arg
NumericT * cuda_arg(scalar< NumericT > &obj)
Convenience helper function for extracting the CUDA handle from a ViennaCL scalar. Non-const version.
Definition: common.hpp:39
scalar.hpp
Implementation of the ViennaCL scalar class.