Commit 1f4cb421 authored by Daniel Campora's avatar Daniel Campora
Browse files

Better Velo consolidation. Some fixes to performance plot.

parent 87e0df64
Pipeline #935471 passed with stage
in 1 minute and 15 seconds
......@@ -48,7 +48,7 @@ def main(argv):
# Execute sequence with nvprof and save output
print("Executing script...")
filename = os.getcwd() + "/" + output_path + "/" + "nvprof_output.txt"
return_value = os.system("nvprof " + " ".join(argv) + " > " + filename +
return_value = os.system("/usr/local/cuda-10.0/bin/nvprof " + " ".join(argv) + " > " + filename +
" 2>&1")
if return_value != 0:
print("Script returned an error message.\nCheck " + filename)
......@@ -122,7 +122,7 @@ def main(argv):
"SciFi": {
"algorithms": scifi_algorithms,
"value": 0,
"color": colors[10]
"color": colors[11]
},
"Muon": {
"algorithms": muon_algorithms,
......@@ -165,13 +165,13 @@ def main(argv):
# Some default parameters for the figure
figure_scale = 1.5
scale = 2.5
fig = plt.figure(figsize=(12 * figure_scale, 12 * figure_scale))
fig = plt.figure(figsize=(6 * figure_scale, 12 * figure_scale))
ax = plt.axes()
bar_width = 0.85
opacity = 0.8
ax.xaxis.grid(True, linestyle="dashed")
# ax.xaxis.grid(True, linestyle="dashed")
keylist_colors = []
for k in keylist:
......@@ -199,7 +199,7 @@ def main(argv):
plt.tick_params(axis='both', which='major', labelsize=6 * scale)
plt.xlabel(
'Fraction of Allen sequence (%)', fontdict={'fontsize': 12 * scale})
'Fraction of Allen sequence (%)', fontdict={'fontsize': 6 * scale})
# Make the bar plot
labels = [a for a in labels_order]
......@@ -219,8 +219,8 @@ def main(argv):
framealpha=0.8,
loc='right',
ncol=1,
fontsize=8 * scale)
fontsize=6 * scale)
"""
t = plt.text(
0.95,
0.05,
......@@ -231,6 +231,12 @@ def main(argv):
'fontsize': 8 * scale,
'style': 'italic'
})
"""
ax.set_xticks(range(0, int(ax.get_xlim()[1]), 1), minor=True)
ax.xaxis.grid(which='both', linestyle="dashed")
ax.xaxis.grid(which='minor', alpha=0.4)
ax.xaxis.grid(which='major', alpha=1.0)
output_filename = "allen_timing_fractions_no_prefix_sum"
print("Producing plots in " + output_path + "/" + output_filename)
......
......@@ -9,14 +9,6 @@
#include "ArgumentsVelo.cuh"
#include <cstdint>
__device__ VeloState means_square_fit(
Velo::Consolidated::Hits& consolidated_hits,
const float* hit_Xs,
const float* hit_Ys,
const float* hit_Zs,
const uint* hit_IDs,
const Velo::TrackHits& track);
__global__ void consolidate_velo_tracks(
int* dev_atomics_velo,
const Velo::TrackHits* dev_tracks,
......
......@@ -5,10 +5,6 @@
*/
__device__ VeloState means_square_fit(
Velo::Consolidated::Hits& consolidated_hits,
const float* hit_Xs,
const float* hit_Ys,
const float* hit_Zs,
const uint* hit_IDs,
const Velo::TrackHits& track)
{
VeloState state;
......@@ -21,10 +17,9 @@ __device__ VeloState means_square_fit(
// Iterate over hits
for (unsigned short h = 0; h < track.hitsNum; ++h) {
const auto hit_index = track.hits[h];
const auto x = hit_Xs[hit_index];
const auto y = hit_Ys[hit_index];
const auto z = hit_Zs[hit_index];
const auto x = consolidated_hits.x[h];
const auto y = consolidated_hits.y[h];
const auto z = consolidated_hits.z[h];
const auto wx = Velo::Tracking::param_w;
const auto wx_t_x = wx * x;
......@@ -45,67 +40,17 @@ __device__ VeloState means_square_fit(
uz2 += wy_t_z * z;
}
{
// Calculate tx, ty and backward
const auto dens = 1.0f / (sz2 * s0 - sz * sz);
state.tx = (sxz * s0 - sx * sz) * dens;
state.x = (sx * sz2 - sxz * sz) * dens;
// Calculate tx, ty and backward
const auto dens = 1.0f / (sz2 * s0 - sz * sz);
state.tx = (sxz * s0 - sx * sz) * dens;
state.x = (sx * sz2 - sxz * sz) * dens;
const auto denu = 1.0f / (uz2 * u0 - uz * uz);
state.ty = (uyz * u0 - uy * uz) * denu;
state.y = (uy * uz2 - uyz * uz) * denu;
const auto denu = 1.0f / (uz2 * u0 - uz * uz);
state.ty = (uyz * u0 - uy * uz) * denu;
state.y = (uy * uz2 - uyz * uz) * denu;
state.z = -(state.x * state.tx + state.y * state.ty) / (state.tx * state.tx + state.ty * state.ty);
state.backward = state.z > consolidated_hits.z[0];
}
// Not needed by any later algorithm
// {
// // Covariance
// const auto m00 = s0;
// const auto m11 = u0;
// const auto m20 = sz - state.z * s0;
// const auto m31 = uz - state.z * u0;
// const auto m22 = sz2 - 2 * state.z * sz + state.z * state.z * s0;
// const auto m33 = uz2 - 2 * state.z * uz + state.z * state.z * u0;
// const auto den20 = 1.0f / (m22 * m00 - m20 * m20);
// const auto den31 = 1.0f / (m33 * m11 - m31 * m31);
// state.c00 = m22 * den20;
// state.c20 = -m20 * den20;
// state.c22 = m00 * den20;
// state.c11 = m33 * den31;
// state.c31 = -m31 * den31;
// state.c33 = m11 * den31;
// }
// {
// //=========================================================================
// // Chi2 / degrees-of-freedom of straight-line fit
// //=========================================================================
// float ch = 0.0f;
// int nDoF = -4;
// for (uint h = 0; h < track.hitsNum; ++h) {
// const auto z = consolidated_hits.z[h];
// const auto x = state.x + state.tx * z;
// const auto y = state.y + state.ty * z;
// const auto dx = x - consolidated_hits.x[h];
// const auto dy = y - consolidated_hits.y[h];
// ch += dx * dx * Velo::Tracking::param_w + dy * dy * Velo::Tracking::param_w;
// // Nice :)
// // TODO: We can get rid of the X and Y read here
// // float sum_w_xzi_2 = CL_Velo::Tracking::param_w * x; // for each hit
// // float sum_w_xi_2 = CL_Velo::Tracking::param_w * hit_Xs[hitno]; // for each hit
// // ch = (sum_w_xzi_2 - sum_w_xi_2) + (sum_w_yzi_2 - sum_w_yi_2);
// nDoF += 2;
// }
// state.chi2 = ch / nDoF;
// }
state.z = -(state.x * state.tx + state.y * state.ty) / (state.tx * state.tx + state.ty * state.ty);
state.backward = state.z > consolidated_hits.z[0];
state.x = state.x + state.tx * state.z;
state.y = state.y + state.ty * state.z;
......@@ -113,6 +58,18 @@ __device__ VeloState means_square_fit(
return state;
}
template<typename T>
__device__ void populate(
const Velo::TrackHits& track,
T* __restrict__ a,
const T* __restrict__ b)
{
for (int i = 0; i < track.hitsNum; ++i) {
const auto hit_index = track.hits[i];
a[i] = b[hit_index];
}
}
__global__ void consolidate_velo_tracks(
int* dev_atomics_velo,
const Velo::TrackHits* dev_tracks,
......@@ -141,29 +98,23 @@ __global__ void consolidate_velo_tracks(
const uint hit_offset = module_hitStarts[0];
// Order has changed since SortByPhi
const float* hit_Ys = (float*) (dev_velo_cluster_container + hit_offset);
const float* hit_Zs = (float*) (dev_velo_cluster_container + number_of_hits + hit_offset);
const float* hit_Xs = (float*) (dev_velo_cluster_container + 5 * number_of_hits + hit_offset);
const uint32_t* hit_IDs = (uint32_t*) (dev_velo_cluster_container + 2 * number_of_hits + hit_offset);
const float* float_dev_velo_cluster_container = (float*) (dev_velo_cluster_container);
// const uint32_t* hit_IDs = (uint32_t*) (dev_velo_cluster_container + 2 * number_of_hits + hit_offset);
// const float* hit_Xs = (float*) (dev_velo_cluster_container + 5 * number_of_hits + hit_offset);
// const float* hit_Ys = (float*) (dev_velo_cluster_container + hit_offset);
// const float* hit_Zs = (float*) (dev_velo_cluster_container + number_of_hits + hit_offset);
for (uint i = threadIdx.x; i < number_of_tracks_event; i += blockDim.x) {
Velo::Consolidated::Hits consolidated_hits = velo_tracks.get_hits(dev_velo_track_hits, i);
const Velo::TrackHits track = event_tracks[i];
auto populate = [&track](uint32_t* __restrict__ a, uint32_t* __restrict__ b) {
for (int i = 0; i < track.hitsNum; ++i) {
const auto hit_index = track.hits[i];
a[i] = b[hit_index];
}
};
populate((uint32_t*) consolidated_hits.x, (uint32_t*) hit_Xs);
populate((uint32_t*) consolidated_hits.y, (uint32_t*) hit_Ys);
populate((uint32_t*) consolidated_hits.z, (uint32_t*) hit_Zs);
populate((uint32_t*) consolidated_hits.LHCbID, (uint32_t*) hit_IDs);
populate<float>(track, consolidated_hits.x, float_dev_velo_cluster_container + 5 * number_of_hits + hit_offset);
populate<float>(track, consolidated_hits.y, float_dev_velo_cluster_container + hit_offset);
populate<float>(track, consolidated_hits.z, float_dev_velo_cluster_container + number_of_hits + hit_offset);
populate<uint32_t>(track, consolidated_hits.LHCbID, (uint32_t*) dev_velo_cluster_container + 2 * number_of_hits + hit_offset);
// Calculate and store fit in consolidated container
VeloState beam_state = means_square_fit(consolidated_hits, hit_Xs, hit_Ys, hit_Zs, hit_IDs, track);
VeloState beam_state = means_square_fit(consolidated_hits, track);
velo_states.set(event_tracks_offset + i, beam_state);
}
}
Supports Markdown
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