Skip to content

Commit 49f667d

Browse files
authored
Fix build with Clang (#116)
* fix build with clang and remove pragma to ignore -Wintel-compat * fix heap corruption in global FFTs after real_real merge
1 parent 6338eb7 commit 49f667d

File tree

3 files changed

+6
-9
lines changed

3 files changed

+6
-9
lines changed

src/portfft/common/global.hpp

Lines changed: 0 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -360,8 +360,6 @@ static void dispatch_transpose_kernel_impl(const Scalar* input, Scalar* output,
360360
});
361361
}
362362

363-
#pragma clang diagnostic push
364-
#pragma clang diagnostic ignored "-Wintel-compat"
365363
/**
366364
* Prepares the launch of transposition at a particular level
367365
* @tparam Scalar Scalar type
@@ -523,7 +521,6 @@ std::vector<sycl::event> compute_level(
523521
}
524522
return events;
525523
}
526-
#pragma clang diagnostic pop
527524
} // namespace detail
528525
} // namespace portfft
529526

src/portfft/descriptor.hpp

Lines changed: 4 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -179,18 +179,18 @@ class committed_descriptor {
179179
template <typename Scalar1, domain Domain1, direction Dir, detail::layout LayoutIn, detail::layout LayoutOut,
180180
Idx SubgroupSize, typename TIn>
181181
friend std::vector<sycl::event> detail::compute_level(
182-
const typename committed_descriptor<Scalar1, Domain1>::kernel_data_struct& kd_struct, TIn input, Scalar* output,
183-
const Scalar* twiddles_ptr, const IdxGlobal* factors_triple, Scalar scale_factor,
182+
const typename committed_descriptor<Scalar1, Domain1>::kernel_data_struct& kd_struct, TIn input, Scalar1* output,
183+
const Scalar1* twiddles_ptr, const IdxGlobal* factors_triple, Scalar1 scale_factor,
184184
IdxGlobal intermediate_twiddle_offset, IdxGlobal subimpl_twiddle_offset, IdxGlobal input_global_offset,
185185
IdxGlobal committed_size, Idx num_batches_in_l2, IdxGlobal n_transforms, IdxGlobal batch_start, Idx factor_id,
186186
Idx total_factors, const std::vector<sycl::event>& dependencies, sycl::queue& queue);
187187

188188
template <typename Scalar1, domain Domain1, typename TOut>
189189
friend sycl::event detail::transpose_level(
190-
const typename committed_descriptor<Scalar1, Domain1>::kernel_data_struct& kd_struct, const Scalar* input,
190+
const typename committed_descriptor<Scalar1, Domain1>::kernel_data_struct& kd_struct, const Scalar1* input,
191191
TOut output, const IdxGlobal* factors_triple, IdxGlobal committed_size, Idx num_batches_in_l2,
192192
IdxGlobal n_transforms, IdxGlobal batch_start, Idx factor_num, Idx total_factors, IdxGlobal output_offset,
193-
sycl::queue& queue, std::shared_ptr<Scalar>& ptr1, std::shared_ptr<Scalar>& ptr2,
193+
sycl::queue& queue, std::shared_ptr<Scalar1>& ptr1, std::shared_ptr<Scalar1>& ptr2,
194194
const std::vector<sycl::event>& events);
195195

196196
descriptor<Scalar, Domain> params;

src/portfft/dispatcher/global_dispatcher.hpp

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -309,7 +309,7 @@ struct committed_descriptor<Scalar, Domain>::run_kernel_struct<Dir, LayoutIn, La
309309
detail::layout::BATCH_INTERLEAVED, SubgroupSize>(
310310
desc.dimensions.at(0).kernels.at(0), in, desc.scratch_ptr_1.get(), twiddles_ptr, factors_and_scan,
311311
scale_factor, intermediate_twiddles_offset, impl_twiddle_offset,
312-
2 * static_cast<IdxGlobal>(i) * committed_size + 2 * input_offset, committed_size,
312+
2 * static_cast<IdxGlobal>(i) * committed_size + input_offset, committed_size,
313313
static_cast<Idx>(max_batches_in_l2), static_cast<IdxGlobal>(num_batches), static_cast<IdxGlobal>(i), 0,
314314
desc.dimensions.at(0).num_factors, {event}, desc.queue);
315315
intermediate_twiddles_offset += 2 * desc.dimensions.at(0).kernels.at(0).batch_size *
@@ -358,7 +358,7 @@ struct committed_descriptor<Scalar, Domain>::run_kernel_struct<Dir, LayoutIn, La
358358
desc.dimensions.at(0).kernels.at(static_cast<std::size_t>(num_factors)),
359359
static_cast<const Scalar*>(desc.scratch_ptr_1.get()), out, factors_and_scan, committed_size,
360360
static_cast<Idx>(max_batches_in_l2), n_transforms, static_cast<IdxGlobal>(i), 0, num_factors,
361-
2 * static_cast<IdxGlobal>(i) * committed_size + 2 * output_offset, desc.queue, desc.scratch_ptr_1,
361+
2 * static_cast<IdxGlobal>(i) * committed_size + output_offset, desc.queue, desc.scratch_ptr_1,
362362
desc.scratch_ptr_2, {event});
363363
}
364364
return event;

0 commit comments

Comments
 (0)