Commit c928f9f6 authored by Daniel Campora's avatar Daniel Campora
Browse files

Replaced all global function stream by contexts.

parent 29cc0c1b
......@@ -244,15 +244,14 @@ void saxpy::saxpy_t::operator()(
{
global_function(saxpy)(
dim3(1),
property<block_dim_t>(),
stream)(arguments);
property<block_dim_t>(), context)(arguments);
}
```
In order to invoke host and global functions, wrapper methods `host_function` and `global_function` should be used. The syntax is as follows:
host_function(<host_function_identifier>)(<parameters of function>)
global_function(<global_function_identifier>)(<grid_size>, <block_size>, <stream>)(<parameters of function>)
global_function(<global_function_identifier>)(<grid_size>, <block_size>, context)(<parameters of function>)
`global_function` wraps a function identifier, such as `saxpy`. The object it returns can be used to invoke a _global kernel_ following a syntax that is similar to [CUDA's kernel invocation syntax](https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#kernels). It expects:
......@@ -308,8 +307,7 @@ Typically, events are processed by independent blocks of execution. When that's
```c++
global_function(kernel)(
size<dev_event_list_t>(),
property<block_dim_t>(),
stream)(arguments);
property<block_dim_t>(), context)(arguments);
```
Then, in the kernel itself, in order to access the event under execution, the following idiom is used:
......
......@@ -20,7 +20,7 @@ void pv_beamline_calculate_denom::pv_beamline_calculate_denom_t::operator()(
const Allen::Context& context) const
{
global_function(pv_beamline_calculate_denom)(
dim3(size<dev_event_list_t>(arguments)), property<block_dim_t>(), stream)(arguments);
dim3(size<dev_event_list_t>(arguments)), property<block_dim_t>(), context)(arguments);
}
__global__ void pv_beamline_calculate_denom::pv_beamline_calculate_denom(
......
......@@ -22,7 +22,7 @@ void pv_beamline_cleanup::pv_beamline_cleanup_t::operator()(
{
initialize<dev_number_of_multi_final_vertices_t>(arguments, 0, context);
global_function(pv_beamline_cleanup)(dim3(size<dev_event_list_t>(arguments)), property<block_dim_t>(), stream)(
global_function(pv_beamline_cleanup)(dim3(size<dev_event_list_t>(arguments)), property<block_dim_t>(), context)(
arguments);
// Retrieve result
......
......@@ -20,7 +20,7 @@ void pv_beamline_extrapolate::pv_beamline_extrapolate_t::operator()(
HostBuffers&,
const Allen::Context& context) const
{
global_function(pv_beamline_extrapolate)(dim3(size<dev_event_list_t>(arguments)), property<block_dim_t>(), stream)(
global_function(pv_beamline_extrapolate)(dim3(size<dev_event_list_t>(arguments)), property<block_dim_t>(), context)(
arguments);
}
......
......@@ -22,7 +22,7 @@ void pv_beamline_histo::pv_beamline_histo_t::operator()(
HostBuffers&,
const Allen::Context& context) const
{
global_function(pv_beamline_histo)(dim3(size<dev_event_list_t>(arguments)), property<block_dim_t>(), stream)(
global_function(pv_beamline_histo)(dim3(size<dev_event_list_t>(arguments)), property<block_dim_t>(), context)(
arguments, constants.dev_beamline.data());
}
......
......@@ -29,7 +29,7 @@ void pv_beamline_multi_fitter::pv_beamline_multi_fitter_t::operator()(
const auto block_dimension = dim3(1, property<block_dim_y_t>());
#endif
global_function(pv_beamline_multi_fitter)(dim3(size<dev_event_list_t>(arguments)), block_dimension, stream)(
global_function(pv_beamline_multi_fitter)(dim3(size<dev_event_list_t>(arguments)), block_dimension, context)(
arguments, constants.dev_beamline.data());
}
......
......@@ -23,7 +23,7 @@ void pv_beamline_peak::pv_beamline_peak_t::operator()(
const auto grid_dim =
dim3((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)(
global_function(pv_beamline_peak)(grid_dim, property<block_dim_x_t>().get(), context)(
arguments, size<dev_event_list_t>(arguments));
}
......
......@@ -20,7 +20,7 @@ void fit_seeds::pv_fit_seeds_t::operator()(
HostBuffers& host_buffers,
const Allen::Context& context) const
{
global_function(fit_seeds)(dim3(first<host_number_of_events_t>(arguments)), property<block_dim_t>(), stream)(
global_function(fit_seeds)(dim3(first<host_number_of_events_t>(arguments)), property<block_dim_t>(), context)(
arguments);
if (runtime_options.do_check) {
......
......@@ -20,7 +20,7 @@ void pv_get_seeds::pv_get_seeds_t::operator()(
HostBuffers& host_buffers,
const Allen::Context& context) const
{
global_function(pv_get_seeds)(dim3(first<host_number_of_events_t>(arguments)), property<block_dim_t>(), stream)(
global_function(pv_get_seeds)(dim3(first<host_number_of_events_t>(arguments)), property<block_dim_t>(), context)(
arguments);
if (runtime_options.do_check) {
......
......@@ -23,7 +23,7 @@ void scifi_consolidate_tracks::scifi_consolidate_tracks_t::operator()(
HostBuffers& host_buffers,
const Allen::Context& context) const
{
global_function(scifi_consolidate_tracks)(dim3(size<dev_event_list_t>(arguments)), property<block_dim_t>(), stream)(
global_function(scifi_consolidate_tracks)(dim3(size<dev_event_list_t>(arguments)), property<block_dim_t>(), context)(
arguments);
// Transmission device to host of Scifi consolidated tracks
......
......@@ -20,7 +20,7 @@ void scifi_copy_track_hit_number::scifi_copy_track_hit_number_t::operator()(
const Allen::Context& context) const
{
global_function(scifi_copy_track_hit_number)(
dim3(first<host_number_of_events_t>(arguments)), property<block_dim_t>(), stream)(arguments);
dim3(first<host_number_of_events_t>(arguments)), property<block_dim_t>(), context)(arguments);
}
/**
......
......@@ -33,14 +33,14 @@ void lf_create_tracks::lf_create_tracks_t::operator()(
initialize<dev_scifi_lf_atomics_t>(arguments, 0, context);
global_function(lf_triplet_keep_best)(
dim3(size<dev_event_list_t>(arguments)), property<triplet_keep_best_block_dim_t>(), stream)(
dim3(size<dev_event_list_t>(arguments)), property<triplet_keep_best_block_dim_t>(), context)(
arguments, constants.dev_looking_forward_constants);
global_function(lf_calculate_parametrization)(
dim3(size<dev_event_list_t>(arguments)), property<calculate_parametrization_block_dim_t>(), stream)(
dim3(size<dev_event_list_t>(arguments)), property<calculate_parametrization_block_dim_t>(), context)(
arguments, constants.dev_looking_forward_constants);
global_function(lf_extend_tracks)(
dim3(size<dev_event_list_t>(arguments)), property<extend_tracks_block_dim_t>(), stream)(
dim3(size<dev_event_list_t>(arguments)), property<extend_tracks_block_dim_t>(), context)(
arguments, constants.dev_looking_forward_constants);
}
\ No newline at end of file
......@@ -18,7 +18,7 @@ void lf_least_mean_square_fit::lf_least_mean_square_fit_t::operator()(
const Allen::Context& context) const
{
global_function(lf_least_mean_square_fit)(
dim3(first<host_number_of_events_t>(arguments)), property<block_dim_t>(), stream)(
dim3(first<host_number_of_events_t>(arguments)), property<block_dim_t>(), context)(
arguments, constants.dev_looking_forward_constants);
}
......
......@@ -35,7 +35,7 @@ void lf_quality_filter::lf_quality_filter_t::operator()(
{
initialize<dev_atomics_scifi_t>(arguments, 0, context);
global_function(lf_quality_filter)(dim3(size<dev_event_list_t>(arguments)), property<block_dim_t>(), stream)(
global_function(lf_quality_filter)(dim3(size<dev_event_list_t>(arguments)), property<block_dim_t>(), context)(
arguments, constants.dev_looking_forward_constants, constants.dev_magnet_polarity.data());
if (runtime_options.do_check) {
......
......@@ -30,7 +30,7 @@ void lf_quality_filter_length::lf_quality_filter_length_t::operator()(
{
initialize<dev_scifi_lf_length_filtered_atomics_t>(arguments, 0, context);
global_function(lf_quality_filter_length)(dim3(size<dev_event_list_t>(arguments)), property<block_dim_t>(), stream)(
global_function(lf_quality_filter_length)(dim3(size<dev_event_list_t>(arguments)), property<block_dim_t>(), context)(
arguments);
}
......
......@@ -30,7 +30,7 @@ void lf_search_initial_windows::lf_search_initial_windows_t::operator()(
{
initialize<dev_scifi_lf_initial_windows_t>(arguments, 0, context);
global_function(lf_search_initial_windows)(dim3(size<dev_event_list_t>(arguments)), property<block_dim_t>(), stream)(
global_function(lf_search_initial_windows)(dim3(size<dev_event_list_t>(arguments)), property<block_dim_t>(), context)(
arguments,
constants.dev_scifi_geometry,
constants.dev_looking_forward_constants,
......
......@@ -32,7 +32,7 @@ void lf_triplet_seeding::lf_triplet_seeding_t::operator()(
initialize<dev_scifi_lf_number_of_found_triplets_t>(arguments, 0, context);
global_function(lf_triplet_seeding)(
dim3(size<dev_event_list_t>(arguments)), dim3(LookingForward::triplet_seeding_block_dim_x, 2), stream)(
dim3(size<dev_event_list_t>(arguments)), dim3(LookingForward::triplet_seeding_block_dim_x, context), stream)(
arguments, constants.dev_looking_forward_constants);
}
......
......@@ -25,12 +25,12 @@ void scifi_calculate_cluster_count_v4::scifi_calculate_cluster_count_v4_t::opera
if (runtime_options.mep_layout) {
global_function(scifi_calculate_cluster_count_v4_mep)(
dim3(size<dev_event_list_t>(arguments)), property<block_dim_t>(), stream)(
dim3(size<dev_event_list_t>(arguments)), property<block_dim_t>(), context)(
arguments, constants.dev_scifi_geometry);
}
else {
global_function(scifi_calculate_cluster_count_v4)(
dim3(size<dev_event_list_t>(arguments)), property<block_dim_t>(), stream)(
dim3(size<dev_event_list_t>(arguments)), property<block_dim_t>(), context)(
arguments, constants.dev_scifi_geometry);
}
}
......
......@@ -25,12 +25,12 @@ void scifi_calculate_cluster_count_v6::scifi_calculate_cluster_count_v6_t::opera
if (runtime_options.mep_layout) {
global_function(scifi_calculate_cluster_count_v6_mep)(
dim3(size<dev_event_list_t>(arguments)), dim3(SciFi::SciFiRawBankParams::NbBanks), stream)(
dim3(size<dev_event_list_t>(arguments)), dim3(SciFi::SciFiRawBankParams::NbBanks), context)(
arguments, constants.dev_scifi_geometry);
}
else {
global_function(scifi_calculate_cluster_count_v6)(
dim3(size<dev_event_list_t>(arguments)), dim3(SciFi::SciFiRawBankParams::NbBanks), stream)(
dim3(size<dev_event_list_t>(arguments)), dim3(SciFi::SciFiRawBankParams::NbBanks), context)(
arguments, constants.dev_scifi_geometry);
}
}
......
......@@ -24,12 +24,12 @@ void scifi_pre_decode_v4::scifi_pre_decode_v4_t::operator()(
{
if (runtime_options.mep_layout) {
global_function(scifi_pre_decode_v4_mep)(
dim3(size<dev_event_list_t>(arguments)), dim3(SciFi::SciFiRawBankParams::NbBanks), stream)(
dim3(size<dev_event_list_t>(arguments)), dim3(SciFi::SciFiRawBankParams::NbBanks), context)(
arguments, constants.dev_scifi_geometry);
}
else {
global_function(scifi_pre_decode_v4)(
dim3(size<dev_event_list_t>(arguments)), dim3(SciFi::SciFiRawBankParams::NbBanks), stream)(
dim3(size<dev_event_list_t>(arguments)), dim3(SciFi::SciFiRawBankParams::NbBanks), context)(
arguments, constants.dev_scifi_geometry);
}
}
......
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