Skip to content

Commit 24c758c

Browse files
author
Thomas Grützmacher
committed
Fix CUDA 9.0 compiler error
1 parent 8f4613d commit 24c758c

1 file changed

Lines changed: 12 additions & 4 deletions

File tree

common/solver/cb_gmres_kernels.hpp.inc

Lines changed: 12 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -88,8 +88,10 @@ __global__ __launch_bounds__(block_size) void initialize_2_1_kernel(
8888
size_type stride_residual_nc)
8989
{
9090
const auto global_id = thread::get_thread_id_flat();
91-
const auto krylov_stride =
91+
const auto krylov_stride_ar =
9292
helper_functions_accessor<Accessor3d>::get_stride(krylov_bases);
93+
const auto krylov_stride =
94+
reinterpret_cast<const size_type(&)[2]>(krylov_stride_ar);
9395
// krylov indices
9496
const auto krylov_idx = global_id / krylov_stride[0];
9597
const auto reminder = global_id % krylov_stride[0];
@@ -122,8 +124,10 @@ __global__ __launch_bounds__(block_size) void initialize_2_2_kernel(
122124
size_type *__restrict__ final_iter_nums)
123125
{
124126
const auto global_id = thread::get_thread_id_flat();
125-
const auto krylov_stride =
127+
const auto krylov_stride_ar =
126128
helper_functions_accessor<Accessor3d>::get_stride(krylov_bases);
129+
const auto krylov_stride =
130+
reinterpret_cast<const size_type(&)[2]>(krylov_stride_ar);
127131
const auto row_idx = global_id / krylov_stride[1];
128132
const auto col_idx = global_id % krylov_stride[1];
129133

@@ -728,8 +732,10 @@ __global__
728732
// CAREFUL: krylov_bases changes might not be done properly since this
729733
// kernel is currently not used!
730734
const auto global_id = thread::get_thread_id_flat();
731-
const auto krylov_stride =
735+
const auto krylov_stride_ar =
732736
helper_functions_accessor<Accessor3d>::get_stride(krylov_bases);
737+
const auto krylov_stride =
738+
reinterpret_cast<const size_type(&)[2]>(krylov_stride_ar);
733739
const auto k = global_id / krylov_stride[0];
734740
const auto reminder = global_id / krylov_stride[0];
735741
const auto row_idx = reminder / krylov_stride[1];
@@ -1323,8 +1329,10 @@ __global__ __launch_bounds__(block_size) void set_scale_kernel(
13231329
static_assert(!is_complex_s<RealValueType>::value,
13241330
"ValueType must not be complex!");
13251331
const auto global_id = thread::get_thread_id_flat();
1326-
const auto krylov_stride =
1332+
const auto krylov_stride_ar =
13271333
helper_functions_accessor<Accessor3d>::get_stride(krylov_bases);
1334+
const auto krylov_stride =
1335+
reinterpret_cast<const size_type(&)[2]>(krylov_stride_ar);
13281336
const auto blk_idx = global_id / krylov_stride[1];
13291337
const auto col_idx = global_id % krylov_stride[1];
13301338

0 commit comments

Comments
 (0)