diff --git a/cpp/src/mip/feasibility_jump/feasibility_jump_kernels.cu b/cpp/src/mip/feasibility_jump/feasibility_jump_kernels.cu index 03da7f363e..7ff901ca73 100644 --- a/cpp/src/mip/feasibility_jump/feasibility_jump_kernels.cu +++ b/cpp/src/mip/feasibility_jump/feasibility_jump_kernels.cu @@ -1153,7 +1153,9 @@ __global__ void select_variable_kernel(typename fj_t::climber_data_t:: raft::random::PCGenerator rng( fj.settings->seed, *fj.iterations * fj.settings->parameters.max_sampled_moves, 0); - __shared__ typename fj_t::move_score_t shmem[2 * raft::WarpSize]; + using move_score_t = typename fj_t::move_score_t; + __shared__ alignas(move_score_t) char shmem_storage[2 * raft::WarpSize * sizeof(move_score_t)]; + auto* const shmem = (move_score_t*)shmem_storage; auto th_best_score = fj_t::move_score_t::invalid(); i_t th_selected_var = std::numeric_limits::max(); @@ -1300,7 +1302,9 @@ DI thrust::tuple::move_score_t> gridwide_reduc cg::this_grid().sync(); if (blockIdx.x == 0) { - __shared__ typename fj_t::move_score_t shmem[2 * raft::WarpSize]; + using move_score_t = typename fj_t::move_score_t; + __shared__ alignas(move_score_t) char shmem_storage[2 * raft::WarpSize * sizeof(move_score_t)]; + auto* const shmem = (move_score_t*)shmem_storage; auto th_best_score = fj_t::move_score_t::invalid(); i_t th_best_block = 0;