Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Dense: conversion to complex and extraction of real and imaginary part #658

Merged
merged 7 commits into from
Nov 6, 2020
Merged

Dense: conversion to complex and extraction of real and imaginary part #658

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
7 commits
Select commit Hold shift + click to select a range
974b176
Add conversion to complex and extraction of real and imaginary part t…
fritzgoebel Oct 22, 2020
39b39e3
Fix wrong parameter types in gpu kernels
fritzgoebel Oct 26, 2020
a07a28d
Apply suggestions from code review
fritzgoebel Oct 27, 2020
db511c5
Add conversions into given matrix
fritzgoebel Oct 28, 2020
a63cb7b
Use temporary clone in conversions with given result
fritzgoebel Oct 29, 2020
339a065
Add dpcpp GKO_NOT_IMPLEMENTED kernels, code review
fritzgoebel Nov 5, 2020
c393e0b
Fix missed dpcpp kernel
fritzgoebel Nov 5, 2020
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
44 changes: 44 additions & 0 deletions common/matrix/dense_kernels.hpp.inc
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -539,4 +539,48 @@ __global__ __launch_bounds__(default_block_size) void outplace_absolute_dense(
}


template <typename ValueType, typename ComplexType>
__global__ __launch_bounds__(default_block_size) void make_complex(
size_type num_rows, size_type num_cols, const ValueType *__restrict__ in,
size_type stride_in, ComplexType *__restrict__ out, size_type stride_out)
Copy link
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

If this ComplexType is only for complex, I prefer use one template and use to_complex to represent complex.
Before Ginkgo 2.0.0, I think we need to discuss the precision(ValueType) conversion...

Copy link
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

This doesn't work, since to_complex returns the CPU complex type, not thrust::complex. An alternative would be thrust::complex<remove_complex<ValueType>>

Copy link
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

you are right. maybe does another way cuda_type<to_complex<ValueType>> work?

{
const auto tidx = thread::get_thread_id_flat();
auto row = tidx / num_cols;
auto col = tidx % num_cols;
if (row < num_rows) {
out[row * stride_out + col] = in[row * stride_in + col];
}
}


template <typename ValueType>
__global__ __launch_bounds__(default_block_size) void get_real(
size_type num_rows, size_type num_cols, const ValueType *__restrict__ in,
size_type stride_in, remove_complex<ValueType> *__restrict__ out,
size_type stride_out)
{
const auto tidx = thread::get_thread_id_flat();
auto row = tidx / num_cols;
auto col = tidx % num_cols;
if (row < num_rows) {
out[row * stride_out + col] = real(in[row * stride_in + col]);
}
}


template <typename ValueType>
__global__ __launch_bounds__(default_block_size) void get_imag(
size_type num_rows, size_type num_cols, const ValueType *__restrict__ in,
size_type stride_in, remove_complex<ValueType> *__restrict__ out,
size_type stride_out)
{
const auto tidx = thread::get_thread_id_flat();
auto row = tidx / num_cols;
auto col = tidx % num_cols;
if (row < num_rows) {
out[row * stride_out + col] = imag(in[row * stride_in + col]);
}
}


} // namespace kernel
15 changes: 15 additions & 0 deletions core/device_hooks/common_kernels.inc.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -256,6 +256,21 @@ GKO_DECLARE_OUTPLACE_ABSOLUTE_DENSE_KERNEL(ValueType)
GKO_NOT_COMPILED(GKO_HOOK_MODULE);
GKO_INSTANTIATE_FOR_EACH_VALUE_TYPE(GKO_DECLARE_OUTPLACE_ABSOLUTE_DENSE_KERNEL);

template <typename ValueType>
GKO_DECLARE_MAKE_COMPLEX_KERNEL(ValueType)
GKO_NOT_COMPILED(GKO_HOOK_MODULE);
GKO_INSTANTIATE_FOR_EACH_VALUE_TYPE(GKO_DECLARE_MAKE_COMPLEX_KERNEL);

template <typename ValueType>
GKO_DECLARE_GET_REAL_KERNEL(ValueType)
GKO_NOT_COMPILED(GKO_HOOK_MODULE);
GKO_INSTANTIATE_FOR_EACH_VALUE_TYPE(GKO_DECLARE_GET_REAL_KERNEL);

template <typename ValueType>
GKO_DECLARE_GET_IMAG_KERNEL(ValueType)
GKO_NOT_COMPILED(GKO_HOOK_MODULE);
GKO_INSTANTIATE_FOR_EACH_VALUE_TYPE(GKO_DECLARE_GET_IMAG_KERNEL);


} // namespace dense

Expand Down
81 changes: 81 additions & 0 deletions core/matrix/dense.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -88,6 +88,9 @@ GKO_REGISTER_OPERATION(convert_to_sparsity_csr, dense::convert_to_sparsity_csr);
GKO_REGISTER_OPERATION(extract_diagonal, dense::extract_diagonal);
GKO_REGISTER_OPERATION(inplace_absolute_dense, dense::inplace_absolute_dense);
GKO_REGISTER_OPERATION(outplace_absolute_dense, dense::outplace_absolute_dense);
GKO_REGISTER_OPERATION(make_complex, dense::make_complex);
GKO_REGISTER_OPERATION(get_real, dense::get_real);
GKO_REGISTER_OPERATION(get_imag, dense::get_imag);


} // namespace dense
Expand Down Expand Up @@ -795,6 +798,84 @@ Dense<ValueType>::compute_absolute() const
}


template <typename ValueType>
std::unique_ptr<typename Dense<ValueType>::complex_type>
Dense<ValueType>::make_complex() const
{
auto exec = this->get_executor();

auto complex_dense = complex_type::create(exec, this->get_size());

exec->run(dense::make_make_complex(this, complex_dense.get()));

return complex_dense;
}


template <typename ValueType>
void Dense<ValueType>::make_complex(Dense<to_complex<ValueType>> *result) const
{
auto exec = this->get_executor();

GKO_ASSERT_EQUAL_DIMENSIONS(this, result);

exec->run(dense::make_make_complex(
this, make_temporary_clone(exec, result).get()));
}


template <typename ValueType>
std::unique_ptr<typename Dense<ValueType>::absolute_type>
Dense<ValueType>::get_real() const
{
auto exec = this->get_executor();

auto real_dense = absolute_type::create(exec, this->get_size());

exec->run(dense::make_get_real(this, real_dense.get()));

return real_dense;
}


template <typename ValueType>
void Dense<ValueType>::get_real(Dense<remove_complex<ValueType>> *result) const
{
auto exec = this->get_executor();

GKO_ASSERT_EQUAL_DIMENSIONS(this, result);

exec->run(
dense::make_get_real(this, make_temporary_clone(exec, result).get()));
}


template <typename ValueType>
std::unique_ptr<typename Dense<ValueType>::absolute_type>
Dense<ValueType>::get_imag() const
{
auto exec = this->get_executor();

auto imag_dense = absolute_type::create(exec, this->get_size());

exec->run(dense::make_get_imag(this, imag_dense.get()));

return imag_dense;
}


template <typename ValueType>
void Dense<ValueType>::get_imag(Dense<remove_complex<ValueType>> *result) const
{
auto exec = this->get_executor();

GKO_ASSERT_EQUAL_DIMENSIONS(this, result);

exec->run(
dense::make_get_imag(this, make_temporary_clone(exec, result).get()));
}


#define GKO_DECLARE_DENSE_MATRIX(_type) class Dense<_type>
GKO_INSTANTIATE_FOR_EACH_VALUE_TYPE(GKO_DECLARE_DENSE_MATRIX);

Expand Down
23 changes: 22 additions & 1 deletion core/matrix/dense_kernels.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -182,6 +182,21 @@ namespace kernels {
const matrix::Dense<_vtype> *source, \
matrix::Dense<remove_complex<_vtype>> *result)

#define GKO_DECLARE_MAKE_COMPLEX_KERNEL(_vtype) \
void make_complex(std::shared_ptr<const DefaultExecutor> exec, \
const matrix::Dense<_vtype> *source, \
matrix::Dense<to_complex<_vtype>> *result)

#define GKO_DECLARE_GET_REAL_KERNEL(_vtype) \
void get_real(std::shared_ptr<const DefaultExecutor> exec, \
const matrix::Dense<_vtype> *source, \
matrix::Dense<remove_complex<_vtype>> *result)

#define GKO_DECLARE_GET_IMAG_KERNEL(_vtype) \
void get_imag(std::shared_ptr<const DefaultExecutor> exec, \
const matrix::Dense<_vtype> *source, \
matrix::Dense<remove_complex<_vtype>> *result)


#define GKO_DECLARE_ALL_AS_TEMPLATES \
template <typename ValueType> \
Expand Down Expand Up @@ -235,7 +250,13 @@ namespace kernels {
template <typename ValueType> \
GKO_DECLARE_INPLACE_ABSOLUTE_DENSE_KERNEL(ValueType); \
template <typename ValueType> \
GKO_DECLARE_OUTPLACE_ABSOLUTE_DENSE_KERNEL(ValueType)
GKO_DECLARE_OUTPLACE_ABSOLUTE_DENSE_KERNEL(ValueType); \
template <typename ValueType> \
GKO_DECLARE_MAKE_COMPLEX_KERNEL(ValueType); \
template <typename ValueType> \
GKO_DECLARE_GET_REAL_KERNEL(ValueType); \
template <typename ValueType> \
GKO_DECLARE_GET_IMAG_KERNEL(ValueType)


namespace omp {
Expand Down
51 changes: 51 additions & 0 deletions cuda/matrix/dense_kernels.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -734,6 +734,57 @@ void outplace_absolute_dense(std::shared_ptr<const CudaExecutor> exec,
GKO_INSTANTIATE_FOR_EACH_VALUE_TYPE(GKO_DECLARE_OUTPLACE_ABSOLUTE_DENSE_KERNEL);


template <typename ValueType>
void make_complex(std::shared_ptr<const CudaExecutor> exec,
const matrix::Dense<ValueType> *source,
matrix::Dense<to_complex<ValueType>> *result)
{
auto dim = source->get_size();
const dim3 grid_dim = ceildiv(dim[0] * dim[1], default_block_size);

kernel::make_complex<<<grid_dim, default_block_size>>>(
dim[0], dim[1], as_cuda_type(source->get_const_values()),
source->get_stride(), as_cuda_type(result->get_values()),
result->get_stride());
}

GKO_INSTANTIATE_FOR_EACH_VALUE_TYPE(GKO_DECLARE_MAKE_COMPLEX_KERNEL);


template <typename ValueType>
void get_real(std::shared_ptr<const CudaExecutor> exec,
const matrix::Dense<ValueType> *source,
matrix::Dense<remove_complex<ValueType>> *result)
{
auto dim = source->get_size();
const dim3 grid_dim = ceildiv(dim[0] * dim[1], default_block_size);

kernel::get_real<<<grid_dim, default_block_size>>>(
dim[0], dim[1], as_cuda_type(source->get_const_values()),
source->get_stride(), as_cuda_type(result->get_values()),
result->get_stride());
}

GKO_INSTANTIATE_FOR_EACH_VALUE_TYPE(GKO_DECLARE_GET_REAL_KERNEL);


template <typename ValueType>
void get_imag(std::shared_ptr<const CudaExecutor> exec,
const matrix::Dense<ValueType> *source,
matrix::Dense<remove_complex<ValueType>> *result)
{
auto dim = source->get_size();
const dim3 grid_dim = ceildiv(dim[0] * dim[1], default_block_size);

kernel::get_imag<<<grid_dim, default_block_size>>>(
dim[0], dim[1], as_cuda_type(source->get_const_values()),
source->get_stride(), as_cuda_type(result->get_values()),
result->get_stride());
}

GKO_INSTANTIATE_FOR_EACH_VALUE_TYPE(GKO_DECLARE_GET_IMAG_KERNEL);


} // namespace dense
} // namespace cuda
} // namespace kernels
Expand Down
72 changes: 72 additions & 0 deletions cuda/test/matrix/dense_kernels.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -643,4 +643,76 @@ TEST_F(Dense, OutplaceAbsoluteMatrixIsEquivalentToRef)
}


TEST_F(Dense, MakeComplexIsEquivalentToRef)
{
set_up_apply_data();

auto complex_x = x->make_complex();
auto dcomplex_x = dx->make_complex();

GKO_ASSERT_MTX_NEAR(complex_x, dcomplex_x, 0);
}


TEST_F(Dense, MakeComplexWithGivenResultIsEquivalentToRef)
{
set_up_apply_data();

auto complex_x = ComplexMtx::create(ref, x->get_size());
x->make_complex(complex_x.get());
auto dcomplex_x = ComplexMtx::create(cuda, x->get_size());
dx->make_complex(dcomplex_x.get());

GKO_ASSERT_MTX_NEAR(complex_x, dcomplex_x, 0);
}


TEST_F(Dense, GetRealIsEquivalentToRef)
{
set_up_apply_data();

auto real_x = x->get_real();
auto dreal_x = dx->get_real();

GKO_ASSERT_MTX_NEAR(real_x, dreal_x, 0);
}


TEST_F(Dense, GetRealWithGivenResultIsEquivalentToRef)
{
set_up_apply_data();

auto real_x = Mtx::create(ref, x->get_size());
x->get_real(real_x.get());
auto dreal_x = Mtx::create(cuda, dx->get_size());
dx->get_real(dreal_x.get());

GKO_ASSERT_MTX_NEAR(real_x, dreal_x, 0);
}


TEST_F(Dense, GetImagIsEquivalentToRef)
{
set_up_apply_data();

auto imag_x = x->get_imag();
auto dimag_x = dx->get_imag();

GKO_ASSERT_MTX_NEAR(imag_x, dimag_x, 0);
}


TEST_F(Dense, GetImagWithGivenResultIsEquivalentToRef)
{
set_up_apply_data();

auto imag_x = Mtx::create(ref, x->get_size());
x->get_imag(imag_x.get());
auto dimag_x = Mtx::create(cuda, dx->get_size());
dx->get_imag(dimag_x.get());

GKO_ASSERT_MTX_NEAR(imag_x, dimag_x, 0);
}


} // namespace
27 changes: 27 additions & 0 deletions dpcpp/matrix/dense_kernels.dp.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -308,6 +308,33 @@ void outplace_absolute_dense(std::shared_ptr<const DpcppExecutor> exec,
GKO_INSTANTIATE_FOR_EACH_VALUE_TYPE(GKO_DECLARE_OUTPLACE_ABSOLUTE_DENSE_KERNEL);


template <typename ValueType>
void make_complex(std::shared_ptr<const DpcppExecutor> exec,
const matrix::Dense<ValueType> *source,
matrix::Dense<to_complex<ValueType>> *result)
GKO_NOT_IMPLEMENTED;

GKO_INSTANTIATE_FOR_EACH_VALUE_TYPE(GKO_DECLARE_MAKE_COMPLEX_KERNEL);


template <typename ValueType>
void get_real(std::shared_ptr<const DpcppExecutor> exec,
const matrix::Dense<ValueType> *source,
matrix::Dense<remove_complex<ValueType>> *result)
GKO_NOT_IMPLEMENTED;

GKO_INSTANTIATE_FOR_EACH_VALUE_TYPE(GKO_DECLARE_GET_REAL_KERNEL);


template <typename ValueType>
void get_imag(std::shared_ptr<const DpcppExecutor> exec,
const matrix::Dense<ValueType> *source,
matrix::Dense<remove_complex<ValueType>> *result)
GKO_NOT_IMPLEMENTED;

GKO_INSTANTIATE_FOR_EACH_VALUE_TYPE(GKO_DECLARE_GET_IMAG_KERNEL);


} // namespace dense
} // namespace dpcpp
} // namespace kernels
Expand Down
Loading