Commit 52c200d8 authored by Daniel Campora's avatar Daniel Campora
Browse files

Fixed pv beamline peak bug by which it accessed data outside of the intended events.

parent 1d3dadb6
Pipeline #2054137 failed with stages
in 17 minutes and 46 seconds
......@@ -51,8 +51,7 @@ def PVSequence(initialize_lists, velo_copy_track_hit_number,
name="pv_beamline_peak",
host_number_of_events_t=initialize_lists.host_number_of_events_t(),
dev_zhisto_t=pv_beamline_histo.dev_zhisto_t(),
dev_event_list_t=initialize_lists.dev_event_list_t(),
dev_number_of_events_t=initialize_lists.dev_number_of_events_t())
dev_event_list_t=initialize_lists.dev_event_list_t())
pv_beamline_calculate_denom = pv_beamline_calculate_denom_t(
name="pv_beamline_calculate_denom",
......
......@@ -319,7 +319,6 @@ Then, in the kernel itself, in order to access the event under execution, the fo
```c++
__global__ void kernel(namespace::Parameters parameters) {
const unsigned event_number = parameters.dev_event_list[blockIdx.x];
const unsigned number_of_events = parameters.dev_number_of_events[0];
```
Configuring the algorithm in a sequence
......
......@@ -19,13 +19,12 @@ namespace pv_beamline_peak {
Parameters,
(HOST_INPUT(host_number_of_events_t, unsigned), host_number_of_events),
(DEVICE_INPUT(dev_event_list_t, unsigned), dev_event_list),
(DEVICE_INPUT(dev_number_of_events_t, unsigned), dev_number_of_events),
(DEVICE_INPUT(dev_zhisto_t, float), dev_zhisto),
(DEVICE_OUTPUT(dev_zpeaks_t, float), dev_zpeaks),
(DEVICE_OUTPUT(dev_number_of_zpeaks_t, unsigned), dev_number_of_zpeaks),
(PROPERTY(block_dim_x_t, "block_dim_x", "block dimension X", unsigned), block_dim_x))
__global__ void pv_beamline_peak(Parameters);
__global__ void pv_beamline_peak(Parameters, const unsigned event_list_size);
struct pv_beamline_peak_t : public DeviceAlgorithm, Parameters {
void set_arguments_size(
......
......@@ -22,18 +22,17 @@ void pv_beamline_peak::pv_beamline_peak_t::operator()(
cudaEvent_t&) const
{
const auto grid_dim = dim3(
(first<host_number_of_events_t>(arguments) + property<block_dim_x_t>().get() - 1) /
(size<dev_event_list_t>(arguments) + property<block_dim_x_t>().get() - 1) /
property<block_dim_x_t>().get());
global_function(pv_beamline_peak)(grid_dim, property<block_dim_x_t>().get(), stream)(arguments);
global_function(pv_beamline_peak)(grid_dim, property<block_dim_x_t>().get(), stream)(arguments, size<dev_event_list_t>(arguments));
}
__global__ void pv_beamline_peak::pv_beamline_peak(pv_beamline_peak::Parameters parameters)
__global__ void pv_beamline_peak::pv_beamline_peak(pv_beamline_peak::Parameters parameters, const unsigned event_list_size)
{
// At least parallelize over events, even if it's
// one event on each thread
const unsigned number_of_events = parameters.dev_number_of_events[0];
for (auto event_index = blockIdx.x * blockDim.x + threadIdx.x; event_index < number_of_events;
for (auto event_index = blockIdx.x * blockDim.x + threadIdx.x; event_index < event_list_size;
event_index += blockDim.x * gridDim.x) {
const unsigned event_number = parameters.dev_event_list[event_index];
......
Markdown is supported
0% or .
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment