Commit 76224e8f authored by Daniel Campora's avatar Daniel Campora
Browse files

Full sequence working.

parent 30a66e14
This diff is collapsed.
......@@ -266,7 +266,7 @@ void SequenceVisitor::visit<saxpy_t>(
// Copy memory from host to device
cudaCheck(cudaMemcpyAsync(
arguments.offset<dev_x>(),
arguments.begin<dev_x>(),
host_buffers.host_x,
saxpy_N * sizeof(float),
cudaMemcpyHostToDevice,
......@@ -274,7 +274,7 @@ void SequenceVisitor::visit<saxpy_t>(
));
cudaCheck(cudaMemcpyAsync(
arguments.offset<dev_y>(),
arguments.begin<dev_y>(),
host_buffers.host_y,
saxpy_N * sizeof(float),
cudaMemcpyHostToDevice,
......@@ -286,8 +286,8 @@ void SequenceVisitor::visit<saxpy_t>(
// Setup arguments for kernel call
state.set_arguments(
arguments.offset<dev_x>(),
arguments.offset<dev_y>(),
arguments.begin<dev_x>(),
arguments.begin<dev_y>(),
saxpy_N,
2.0f
);
......@@ -298,7 +298,7 @@ void SequenceVisitor::visit<saxpy_t>(
// Retrieve result
cudaCheck(cudaMemcpyAsync(
host_buffers.host_y,
arguments.offset<dev_y>(),
arguments.begin<dev_y>(),
arguments.size<dev_y>(),
cudaMemcpyDeviceToHost,
cuda_stream
......
......@@ -46,12 +46,12 @@ namespace pv_beamline_calculate_denom {
cudaEvent_t& cuda_generic_event) const {
function(dim3(value<host_number_of_selected_events_t>(arguments)), block_dimension(), cuda_stream)(
Parameters{
offset<dev_offsets_all_velo_tracks_t>(arguments),
offset<dev_offsets_velo_track_hit_number_t>(arguments),
offset<dev_pvtracks_t>(arguments),
offset<dev_pvtracks_denom_t>(arguments),
offset<dev_zpeaks_t>(arguments),
offset<dev_number_of_zpeaks_t>(arguments)});
begin<dev_offsets_all_velo_tracks_t>(arguments),
begin<dev_offsets_velo_track_hit_number_t>(arguments),
begin<dev_pvtracks_t>(arguments),
begin<dev_pvtracks_denom_t>(arguments),
begin<dev_zpeaks_t>(arguments),
begin<dev_number_of_zpeaks_t>(arguments)});
}
};
}
\ No newline at end of file
......@@ -46,29 +46,29 @@ namespace pv_beamline_cleanup {
cudaEvent_t& cuda_generic_event) const
{
cudaCheck(cudaMemsetAsync(
offset<dev_number_of_multi_final_vertices_t>(arguments),
begin<dev_number_of_multi_final_vertices_t>(arguments),
0,
size<dev_number_of_multi_final_vertices_t>(arguments),
cuda_stream));
function(dim3(value<host_number_of_selected_events_t>(arguments)), block_dimension(), cuda_stream)(
Parameters {offset<dev_multi_fit_vertices_t>(arguments),
offset<dev_number_of_multi_fit_vertices_t>(arguments),
offset<dev_multi_final_vertices_t>(arguments),
offset<dev_number_of_multi_final_vertices_t>(arguments)});
Parameters {begin<dev_multi_fit_vertices_t>(arguments),
begin<dev_number_of_multi_fit_vertices_t>(arguments),
begin<dev_multi_final_vertices_t>(arguments),
begin<dev_number_of_multi_final_vertices_t>(arguments)});
if (runtime_options.do_check) {
// Retrieve result
cudaCheck(cudaMemcpyAsync(
host_buffers.host_reconstructed_multi_pvs,
offset<dev_multi_final_vertices_t>(arguments),
begin<dev_multi_final_vertices_t>(arguments),
size<dev_multi_final_vertices_t>(arguments),
cudaMemcpyDeviceToHost,
cuda_stream));
cudaCheck(cudaMemcpyAsync(
host_buffers.host_number_of_multivertex,
offset<dev_number_of_multi_final_vertices_t>(arguments),
begin<dev_number_of_multi_final_vertices_t>(arguments),
size<dev_number_of_multi_final_vertices_t>(arguments),
cudaMemcpyDeviceToHost,
cuda_stream));
......
......@@ -46,11 +46,11 @@ namespace pv_beamline_extrapolate {
cudaEvent_t& cuda_generic_event) const {
function(dim3(value<host_number_of_selected_events_t>(arguments)), block_dimension(), cuda_stream)(
Parameters{
offset<dev_velo_kalman_beamline_states_t>(arguments),
offset<dev_offsets_all_velo_tracks_t>(arguments),
offset<dev_offsets_velo_track_hit_number_t>(arguments),
offset<dev_pvtracks_t>(arguments),
offset<dev_pvtrack_z_t>(arguments)
begin<dev_velo_kalman_beamline_states_t>(arguments),
begin<dev_offsets_all_velo_tracks_t>(arguments),
begin<dev_offsets_velo_track_hit_number_t>(arguments),
begin<dev_pvtracks_t>(arguments),
begin<dev_pvtrack_z_t>(arguments)
});
}
};
......
......@@ -43,10 +43,10 @@ namespace pv_beamline_histo {
cudaStream_t& cuda_stream,
cudaEvent_t& cuda_generic_event) const {
function(dim3(value<host_number_of_selected_events_t>(arguments)), block_dimension(), cuda_stream)(
Parameters {offset<dev_offsets_all_velo_tracks_t>(arguments),
offset<dev_offsets_velo_track_hit_number_t>(arguments),
offset<dev_pvtracks_t>(arguments),
offset<dev_zhisto_t>(arguments)},
Parameters {begin<dev_offsets_all_velo_tracks_t>(arguments),
begin<dev_offsets_velo_track_hit_number_t>(arguments),
begin<dev_pvtracks_t>(arguments),
begin<dev_zhisto_t>(arguments)},
constants.dev_beamline.data());
}
};
......
......@@ -53,21 +53,21 @@ namespace pv_beamline_multi_fitter {
cudaStream_t& cuda_stream,
cudaEvent_t& cuda_generic_event) const {
cudaCheck(cudaMemsetAsync(
offset<dev_number_of_multi_fit_vertices_t>(arguments),
begin<dev_number_of_multi_fit_vertices_t>(arguments),
0,
size<dev_number_of_multi_fit_vertices_t>(arguments),
cuda_stream));
function(dim3(value<host_number_of_selected_events_t>(arguments)), block_dimension(), cuda_stream)(
Parameters {offset<dev_offsets_all_velo_tracks_t>(arguments),
offset<dev_offsets_velo_track_hit_number_t>(arguments),
offset<dev_pvtracks_t>(arguments),
offset<dev_pvtracks_denom_t>(arguments),
offset<dev_zpeaks_t>(arguments),
offset<dev_number_of_zpeaks_t>(arguments),
offset<dev_multi_fit_vertices_t>(arguments),
offset<dev_number_of_multi_fit_vertices_t>(arguments),
offset<dev_pvtrack_z_t>(arguments)},
Parameters {begin<dev_offsets_all_velo_tracks_t>(arguments),
begin<dev_offsets_velo_track_hit_number_t>(arguments),
begin<dev_pvtracks_t>(arguments),
begin<dev_pvtracks_denom_t>(arguments),
begin<dev_zpeaks_t>(arguments),
begin<dev_number_of_zpeaks_t>(arguments),
begin<dev_multi_fit_vertices_t>(arguments),
begin<dev_number_of_multi_fit_vertices_t>(arguments),
begin<dev_pvtrack_z_t>(arguments)},
constants.dev_beamline.data());
}
};
......
......@@ -49,7 +49,7 @@ namespace pv_beamline_peak {
function(grid_dim, PV::num_threads_pv_beamline_peak_t, cuda_stream)(
Parameters {
offset<dev_zhisto_t>(arguments), offset<dev_zpeaks_t>(arguments), offset<dev_number_of_zpeaks_t>(arguments)},
begin<dev_zhisto_t>(arguments), begin<dev_zpeaks_t>(arguments), begin<dev_number_of_zpeaks_t>(arguments)},
value<host_number_of_selected_events_t>(arguments));
}
};
......
......@@ -56,26 +56,26 @@ namespace fit_seeds {
cudaEvent_t& cuda_generic_event) const
{
function(dim3(value<host_number_of_selected_events_t>(arguments)), block_dimension(), cuda_stream)(
Parameters {offset<dev_vertex_t>(arguments),
offset<dev_number_vertex_t>(arguments),
offset<dev_seeds_t>(arguments),
offset<dev_number_seeds_t>(arguments),
offset<dev_velo_kalman_beamline_states_t>(arguments),
offset<dev_atomics_velo_t>(arguments),
offset<dev_velo_track_hit_number_t>(arguments)});
Parameters {begin<dev_vertex_t>(arguments),
begin<dev_number_vertex_t>(arguments),
begin<dev_seeds_t>(arguments),
begin<dev_number_seeds_t>(arguments),
begin<dev_velo_kalman_beamline_states_t>(arguments),
begin<dev_atomics_velo_t>(arguments),
begin<dev_velo_track_hit_number_t>(arguments)});
if (runtime_options.do_check) {
// Retrieve result
cudaCheck(cudaMemcpyAsync(
host_buffers.host_reconstructed_pvs,
offset<dev_vertex_t>(arguments),
begin<dev_vertex_t>(arguments),
size<dev_vertex_t>(arguments),
cudaMemcpyDeviceToHost,
cuda_stream));
cudaCheck(cudaMemcpyAsync(
host_buffers.host_number_of_vertex,
offset<dev_number_vertex_t>(arguments),
begin<dev_number_vertex_t>(arguments),
size<dev_number_vertex_t>(arguments),
cudaMemcpyDeviceToHost,
cuda_stream));
......
......@@ -47,16 +47,16 @@ namespace pv_get_seeds {
cudaEvent_t& cuda_generic_event) const
{
function(dim3(value<host_number_of_selected_events_t>(arguments)), block_dimension(), cuda_stream)(
Parameters {offset<dev_velo_kalman_beamline_states_t>(arguments),
offset<dev_atomics_velo_t>(arguments),
offset<dev_velo_track_hit_number_t>(arguments),
offset<dev_seeds_t>(arguments),
offset<dev_number_seeds_t>(arguments)});
Parameters {begin<dev_velo_kalman_beamline_states_t>(arguments),
begin<dev_atomics_velo_t>(arguments),
begin<dev_velo_track_hit_number_t>(arguments),
begin<dev_seeds_t>(arguments),
begin<dev_number_seeds_t>(arguments)});
if (runtime_options.do_check) {
cudaCheck(cudaMemcpyAsync(
host_buffers.host_number_of_seeds,
offset<dev_number_seeds_t>(arguments),
begin<dev_number_seeds_t>(arguments),
size<dev_number_seeds_t>(arguments),
cudaMemcpyDeviceToHost,
cuda_stream));
......
......@@ -56,52 +56,52 @@ namespace scifi_consolidate_tracks {
cudaEvent_t& cuda_generic_event) const
{
function(dim3(value<host_number_of_selected_events_t>(arguments)), block_dimension(), cuda_stream)(
Parameters {offset<dev_scifi_hits_t>(arguments),
offset<dev_scifi_hit_offsets_t>(arguments),
offset<dev_scifi_track_hits_t>(arguments),
offset<dev_offsets_forward_tracks_t>(arguments),
offset<dev_offsets_scifi_track_hit_number>(arguments),
offset<dev_scifi_qop_t>(arguments),
offset<dev_scifi_states_t>(arguments),
offset<dev_scifi_track_ut_indices_t>(arguments),
offset<dev_offsets_ut_tracks_t>(arguments),
offset<dev_offsets_ut_track_hit_number_t>(arguments),
offset<dev_scifi_tracks_t>(arguments),
offset<dev_scifi_lf_parametrization_consolidate_t>(arguments)});
Parameters {begin<dev_scifi_hits_t>(arguments),
begin<dev_scifi_hit_offsets_t>(arguments),
begin<dev_scifi_track_hits_t>(arguments),
begin<dev_offsets_forward_tracks_t>(arguments),
begin<dev_offsets_scifi_track_hit_number>(arguments),
begin<dev_scifi_qop_t>(arguments),
begin<dev_scifi_states_t>(arguments),
begin<dev_scifi_track_ut_indices_t>(arguments),
begin<dev_offsets_ut_tracks_t>(arguments),
begin<dev_offsets_ut_track_hit_number_t>(arguments),
begin<dev_scifi_tracks_t>(arguments),
begin<dev_scifi_lf_parametrization_consolidate_t>(arguments)});
// Transmission device to host of Scifi consolidated tracks
if (runtime_options.do_check) {
cudaCheck(cudaMemcpyAsync(
host_buffers.host_atomics_scifi,
offset<dev_offsets_forward_tracks_t>(arguments),
begin<dev_offsets_forward_tracks_t>(arguments),
size<dev_offsets_forward_tracks_t>(arguments),
cudaMemcpyDeviceToHost,
cuda_stream));
cudaCheck(cudaMemcpyAsync(
host_buffers.host_scifi_track_hit_number,
offset<dev_offsets_scifi_track_hit_number>(arguments),
begin<dev_offsets_scifi_track_hit_number>(arguments),
size<dev_offsets_scifi_track_hit_number>(arguments),
cudaMemcpyDeviceToHost,
cuda_stream));
cudaCheck(cudaMemcpyAsync(
host_buffers.host_scifi_track_hits,
offset<dev_scifi_track_hits_t>(arguments),
begin<dev_scifi_track_hits_t>(arguments),
size<dev_scifi_track_hits_t>(arguments),
cudaMemcpyDeviceToHost,
cuda_stream));
cudaCheck(cudaMemcpyAsync(
host_buffers.host_scifi_qop,
offset<dev_scifi_qop_t>(arguments),
begin<dev_scifi_qop_t>(arguments),
size<dev_scifi_qop_t>(arguments),
cudaMemcpyDeviceToHost,
cuda_stream));
cudaCheck(cudaMemcpyAsync(
host_buffers.host_scifi_track_ut_indices,
offset<dev_scifi_track_ut_indices_t>(arguments),
begin<dev_scifi_track_ut_indices_t>(arguments),
size<dev_scifi_track_ut_indices_t>(arguments),
cudaMemcpyDeviceToHost,
cuda_stream));
......
......@@ -42,10 +42,10 @@ namespace scifi_copy_track_hit_number {
cudaEvent_t& cuda_generic_event) const
{
function(dim3(value<host_number_of_selected_events_t>(arguments)), block_dimension(), cuda_stream)(
Parameters {offset<dev_offsets_ut_tracks_t>(arguments),
offset<dev_scifi_tracks_t>(arguments),
offset<dev_offsets_forward_tracks_t>(arguments),
offset<dev_scifi_track_hit_number_t>(arguments)});
Parameters {begin<dev_offsets_ut_tracks_t>(arguments),
begin<dev_scifi_tracks_t>(arguments),
begin<dev_offsets_forward_tracks_t>(arguments),
begin<dev_scifi_track_hit_number_t>(arguments)});
}
};
} // namespace scifi_copy_track_hit_number
\ No newline at end of file
......@@ -6,7 +6,6 @@
__global__ void scifi_copy_track_hit_number::scifi_copy_track_hit_number(
scifi_copy_track_hit_number::Parameters parameters)
{
const auto number_of_events = gridDim.x;
const auto event_number = blockIdx.x;
const auto ut_event_tracks_offset = parameters.dev_atomics_ut[event_number];
......
......@@ -55,18 +55,18 @@ namespace lf_calculate_parametrization {
cudaEvent_t& cuda_generic_event) const
{
function(dim3(value<host_number_of_selected_events_t>(arguments)), block_dimension(), cuda_stream)(
Parameters {offset<dev_scifi_hits_t>(arguments),
offset<dev_scifi_hit_offsets_t>(arguments),
offset<dev_offsets_all_velo_tracks_t>(arguments),
offset<dev_offsets_velo_track_hit_number_t>(arguments),
offset<dev_velo_states_t>(arguments),
offset<dev_offsets_ut_tracks_t>(arguments),
offset<dev_offsets_ut_track_hit_number_t>(arguments),
offset<dev_ut_track_velo_indices_t>(arguments),
offset<dev_ut_qop_t>(arguments),
offset<dev_scifi_lf_tracks_t>(arguments),
offset<dev_scifi_lf_atomics_t>(arguments),
offset<dev_scifi_lf_parametrization_t>(arguments)},
Parameters {begin<dev_scifi_hits_t>(arguments),
begin<dev_scifi_hit_offsets_t>(arguments),
begin<dev_offsets_all_velo_tracks_t>(arguments),
begin<dev_offsets_velo_track_hit_number_t>(arguments),
begin<dev_velo_states_t>(arguments),
begin<dev_offsets_ut_tracks_t>(arguments),
begin<dev_offsets_ut_track_hit_number_t>(arguments),
begin<dev_ut_track_velo_indices_t>(arguments),
begin<dev_ut_qop_t>(arguments),
begin<dev_scifi_lf_tracks_t>(arguments),
begin<dev_scifi_lf_atomics_t>(arguments),
begin<dev_scifi_lf_parametrization_t>(arguments)},
constants.dev_looking_forward_constants);
}
};
......
......@@ -45,15 +45,15 @@ namespace lf_extend_tracks_uv {
cudaEvent_t& cuda_generic_event) const
{
function(dim3(value<host_number_of_selected_events_t>(arguments)), block_dimension(), cuda_stream)(
Parameters {offset<dev_scifi_hits_t>(arguments),
offset<dev_scifi_hit_offsets_t>(arguments),
offset<dev_offsets_ut_tracks_t>(arguments),
offset<dev_offsets_ut_track_hit_number_t>(arguments),
offset<dev_scifi_lf_tracks_t>(arguments),
offset<dev_scifi_lf_atomics_t>(arguments),
offset<dev_ut_states_t>(arguments),
offset<dev_scifi_lf_initial_windows_t>(arguments),
offset<dev_scifi_lf_parametrization_t>(arguments)},
Parameters {begin<dev_scifi_hits_t>(arguments),
begin<dev_scifi_hit_offsets_t>(arguments),
begin<dev_offsets_ut_tracks_t>(arguments),
begin<dev_offsets_ut_track_hit_number_t>(arguments),
begin<dev_scifi_lf_tracks_t>(arguments),
begin<dev_scifi_lf_atomics_t>(arguments),
begin<dev_ut_states_t>(arguments),
begin<dev_scifi_lf_initial_windows_t>(arguments),
begin<dev_scifi_lf_parametrization_t>(arguments)},
constants.dev_looking_forward_constants);
}
};
......
......@@ -43,14 +43,14 @@ namespace lf_extend_tracks_x {
cudaEvent_t& cuda_generic_event) const
{
function(dim3(value<host_number_of_selected_events_t>(arguments)), block_dimension(), cuda_stream)(
Parameters {offset<dev_scifi_hits_t>(arguments),
offset<dev_scifi_hit_offsets_t>(arguments),
offset<dev_offsets_ut_tracks_t>(arguments),
offset<dev_offsets_ut_track_hit_number_t>(arguments),
offset<dev_scifi_lf_tracks_t>(arguments),
offset<dev_scifi_lf_atomics_t>(arguments),
offset<dev_scifi_lf_initial_windows_t>(arguments),
offset<dev_scifi_lf_parametrization_t>(arguments)},
Parameters {begin<dev_scifi_hits_t>(arguments),
begin<dev_scifi_hit_offsets_t>(arguments),
begin<dev_offsets_ut_tracks_t>(arguments),
begin<dev_offsets_ut_track_hit_number_t>(arguments),
begin<dev_scifi_lf_tracks_t>(arguments),
begin<dev_scifi_lf_atomics_t>(arguments),
begin<dev_scifi_lf_initial_windows_t>(arguments),
begin<dev_scifi_lf_parametrization_t>(arguments)},
constants.dev_looking_forward_constants);
}
};
......
......@@ -40,12 +40,12 @@ namespace lf_least_mean_square_fit {
cudaEvent_t& cuda_generic_event) const
{
function(dim3(value<host_number_of_selected_events_t>(arguments)), block_dimension(), cuda_stream)(
Parameters {offset<dev_scifi_hits_t>(arguments),
offset<dev_scifi_hit_count_t>(arguments),
offset<dev_atomics_ut_t>(arguments),
offset<dev_scifi_tracks_t>(arguments),
offset<dev_atomics_scifi_t>(arguments),
offset<dev_scifi_lf_parametrization_x_filter_t>(arguments)},
Parameters {begin<dev_scifi_hits_t>(arguments),
begin<dev_scifi_hit_count_t>(arguments),
begin<dev_atomics_ut_t>(arguments),
begin<dev_scifi_tracks_t>(arguments),
begin<dev_atomics_scifi_t>(arguments),
begin<dev_scifi_lf_parametrization_x_filter_t>(arguments)},
constants.dev_looking_forward_constants);
}
};
......
......@@ -70,39 +70,39 @@ namespace lf_quality_filter {
cudaEvent_t& cuda_generic_event) const
{
cudaCheck(
cudaMemsetAsync(offset<dev_atomics_scifi_t>(arguments), 0, size<dev_atomics_scifi_t>(arguments), cuda_stream));
cudaMemsetAsync(begin<dev_atomics_scifi_t>(arguments), 0, size<dev_atomics_scifi_t>(arguments), cuda_stream));
function(dim3(value<host_number_of_selected_events_t>(arguments)), block_dimension(), cuda_stream)(
Parameters {offset<dev_scifi_hits_t>(arguments),
offset<dev_scifi_hit_offsets_t>(arguments),
offset<dev_offsets_ut_tracks_t>(arguments),
offset<dev_offsets_ut_track_hit_number_t>(arguments),
offset<dev_scifi_lf_length_filtered_tracks_t>(arguments),
offset<dev_scifi_lf_length_filtered_atomics_t>(arguments),
offset<dev_atomics_scifi_t>(arguments),
offset<dev_scifi_tracks_t>(arguments),
offset<dev_scifi_lf_parametrization_length_filter_t>(arguments),
offset<dev_scifi_lf_y_parametrization_length_filter_t>(arguments),
offset<dev_scifi_lf_parametrization_consolidate_t>(arguments),
offset<dev_ut_states_t>(arguments),
offset<dev_velo_states_t>(arguments),
offset<dev_offsets_all_velo_tracks_t>(arguments),
offset<dev_offsets_velo_track_hit_number_t>(arguments),
offset<dev_ut_track_velo_indices_t>(arguments)},
Parameters {begin<dev_scifi_hits_t>(arguments),
begin<dev_scifi_hit_offsets_t>(arguments),
begin<dev_offsets_ut_tracks_t>(arguments),
begin<dev_offsets_ut_track_hit_number_t>(arguments),
begin<dev_scifi_lf_length_filtered_tracks_t>(arguments),
begin<dev_scifi_lf_length_filtered_atomics_t>(arguments),
begin<dev_atomics_scifi_t>(arguments),
begin<dev_scifi_tracks_t>(arguments),
begin<dev_scifi_lf_parametrization_length_filter_t>(arguments),
begin<dev_scifi_lf_y_parametrization_length_filter_t>(arguments),
begin<dev_scifi_lf_parametrization_consolidate_t>(arguments),
begin<dev_ut_states_t>(arguments),
begin<dev_velo_states_t>(arguments),
begin<dev_offsets_all_velo_tracks_t>(arguments),
begin<dev_offsets_velo_track_hit_number_t>(arguments),
begin<dev_ut_track_velo_indices_t>(arguments)},
constants.dev_looking_forward_constants,
constants.dev_magnet_polarity.data());
if (runtime_options.do_check) {
cudaCheck(cudaMemcpyAsync(
host_buffers.host_atomics_scifi,
offset<dev_atomics_scifi_t>(arguments),
begin<dev_atomics_scifi_t>(arguments),
size<dev_atomics_scifi_t>(arguments),
cudaMemcpyDeviceToHost,
cuda_stream));
cudaCheck(cudaMemcpyAsync(
host_buffers.host_scifi_tracks,
offset<dev_scifi_tracks_t>(arguments),
begin<dev_scifi_tracks_t>(arguments),
size<dev_scifi_tracks_t>(arguments),
cudaMemcpyDeviceToHost,
cuda_stream));
......
......@@ -54,20 +54,20 @@ namespace lf_quality_filter_length {
cudaEvent_t& cuda_generic_event) const
{
cudaCheck(cudaMemsetAsync(
offset<dev_scifi_lf_length_filtered_atomics_t>(arguments),
begin<dev_scifi_lf_length_filtered_atomics_t>(arguments),
0,
size<dev_scifi_lf_length_filtered_atomics_t>(arguments),
cuda_stream));
function(dim3(value<host_number_of_selected_events_t>(arguments)), block_dimension(), cuda_stream)(
Parameters {offset<dev_offsets_ut_tracks_t>(arguments),
offset<dev_offsets_ut_track_hit_number_t>(arguments),
offset<dev_scifi_lf_tracks_t>(arguments),
offset<dev_scifi_lf_atomics_t>(arguments),
offset<dev_scifi_lf_length_filtered_tracks_t>(arguments),
offset<dev_scifi_lf_length_filtered_atomics_t>(arguments),
offset<dev_scifi_lf_parametrization_t>(arguments),
offset<dev_scifi_lf_parametrization_length_filter_t>(arguments)});
Parameters {begin<dev_offsets_ut_tracks_t>(arguments),
begin<dev_offsets_ut_track_hit_number_t>(arguments),
begin<dev_scifi_lf_tracks_t>(arguments),
begin<dev_scifi_lf_atomics_t>(arguments),
begin<dev_scifi_lf_length_filtered_tracks_t>(arguments),
begin<dev_scifi_lf_length_filtered_atomics_t>(arguments),
begin<dev_scifi_lf_parametrization_t>(arguments),
begin<dev_scifi_lf_parametrization_length_filter_t>(arguments)});
}
};
} // namespace lf_quality_filter_length
......@@ -53,19 +53,19 @@ namespace lf_quality_filter_x {
cudaEvent_t& cuda_generic_event) const
{
cudaCheck(cudaMemsetAsync(
offset<dev_scifi_lf_x_filtered_atomics_t>(arguments),
begin<dev_scifi_lf_x_filtered_atomics_t>(arguments),
0,
size<dev_scifi_lf_x_filtered_atomics_t>(arguments),
cuda_stream));
function(dim3(value<host_number_of_selected_events_t>(arguments), 24), block_dimension(), cuda_stream)(
Parameters {offset<dev_atomics_ut_t>(arguments),
offset<dev_scifi_lf_tracks_t>(arguments),
offset<dev_scifi_lf_atomics_t>(arguments),
offset<dev_scifi_lf_x_filtered_tracks_t>(arguments),
offset<dev_scifi_lf_x_filtered_atomics_t>(arguments),
offset<dev_scifi_lf_parametrization_t>(arguments),
offset<dev_scifi_lf_parametrization_x_filter_t>(arguments)});
Parameters {begin<dev_atomics_ut_t>(arguments),
begin<dev_scifi_lf_tracks_t>(arguments),
begin<dev_scifi_lf_atomics_t>(arguments),
begin<dev_scifi_lf_x_filtered_tracks_t>(arguments),
begin<dev_scifi_lf_x_filtered_atomics_t>(arguments),
begin<dev_scifi_lf_parametrization_t>(arguments),
begin<dev_scifi_lf_parametrization_x_filter_t>(arguments)});
}
};
} // namespace lf_quality_filter_x
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