Skip to content

Commit

Permalink
TMP commit
Browse files Browse the repository at this point in the history
Extra checks and printouts
  • Loading branch information
Vakho Tsulaia committed Dec 24, 2024
1 parent d704197 commit 77fd950
Show file tree
Hide file tree
Showing 2 changed files with 33 additions and 2 deletions.
15 changes: 13 additions & 2 deletions tests/unit_tests/device/cuda/detector_cuda.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -143,6 +143,16 @@ TEST(detector_cuda, detector_alignment) {
transform_t{shifted, tf.x(), tf.y(), tf.z()});
}

std::cout << "Comparing translations STEP1" << std::endl;
typename decltype(det_host)::transform_container::context_type ctx{};
for (unsigned int i = 0u; i < 20; i++) {
auto trstatic = det_host.transform_store().at(i,ctx).translation();
auto traligned = tf_store_aligned_host.at(i,ctx).translation();
auto trdiff = traligned - trstatic;
EXPECT_POINT3_NEAR(trdiff, shift, 1e-6);
}


// copy the vector of aligned transforms to the device
// again, use synchronous copy and fixed size buffers
auto tf_buff_aligned =
Expand Down Expand Up @@ -173,10 +183,11 @@ TEST(detector_cuda, detector_alignment) {
surfacexf_data_static, surfacexf_data_aligned);

// check that the relevant transforms have been properly shifted
for (unsigned int i = 0u; i < surfacexf_device_static.size(); i++) {
std::cout << "Comparing translations STEP3" << std::endl;
for (unsigned int i = 0u; i < 20; i++) {
auto translation_static = surfacexf_device_static[i].translation();
auto translation_aligned = surfacexf_device_aligned[i].translation();
auto translation_diff = translation_aligned - translation_static;
EXPECT_POINT3_NEAR(translation_diff, shift, 1e-5);
EXPECT_POINT3_NEAR(translation_diff, shift, 1e-6);
}
}
20 changes: 20 additions & 0 deletions tests/unit_tests/device/cuda/detector_cuda_kernel.cu
Original file line number Diff line number Diff line change
Expand Up @@ -128,6 +128,26 @@ __global__ void detector_alignment_test_kernel(
vecmem::device_vector<transform_t> surfacexf_device_aligned(
surfacexf_data_aligned);

printf("Comparing translations STEP2\n");
for (unsigned int i = 0u; i < 20; i++) {
const auto sf = tracking_surface{det_device_static,
det_device_static.surfaces()[i]};
auto trstat = sf.transform(ctx).translation();
const auto af = tracking_surface{det_device_aligned,
det_device_aligned.surfaces()[i]};
auto tralig = af.transform(ctx).translation();
if(fabs(tralig[0]-trstat[0]-.1f)>1e-6

Check warning on line 139 in tests/unit_tests/device/cuda/detector_cuda_kernel.cu

View workflow job for this annotation

GitHub Actions / device-container (CUDA, C++20)

FP64 instructions emitted

Instruction(s) generated are 2 × `cvt.f64.f32` and 2 × `setp.gt.f64` in translation unit(s) `detector_cuda_kernel.ptx` and `detector_cuda_kernel.ptx`.
|| fabs(tralig[1]-trstat[1]-.2f)>1e-6

Check warning on line 140 in tests/unit_tests/device/cuda/detector_cuda_kernel.cu

View workflow job for this annotation

GitHub Actions / device-container (CUDA, C++20)

FP64 instructions emitted

Instruction(s) generated are 2 × `cvt.f64.f32` and 2 × `setp.gt.f64` in translation unit(s) `detector_cuda_kernel.ptx` and `detector_cuda_kernel.ptx`.
|| fabs(tralig[2]-trstat[2]-.3f)>1e-6) {

Check warning on line 141 in tests/unit_tests/device/cuda/detector_cuda_kernel.cu

View workflow job for this annotation

GitHub Actions / device-container (CUDA, C++20)

FP64 instructions emitted

Instruction(s) generated are 2 × `cvt.f64.f32` and 2 × `setp.gt.f64` in translation unit(s) `detector_cuda_kernel.ptx` and `detector_cuda_kernel.ptx`.
printf("[%f,%f,%f] - [%f,%f,%f] = [%f,%f,%f]\n",
tralig[0],tralig[1],tralig[2],
trstat[0],trstat[1],trstat[2],
fabs(tralig[0]-trstat[0]),
fabs(tralig[1]-trstat[1]),
fabs(tralig[2]-trstat[2]));
}

}
// copy surface transforms into relevant vectors
for (unsigned int i = 0u; i < det_device_static.surfaces().size(); i++) {
const auto sf = tracking_surface{det_device_static,
Expand Down

0 comments on commit 77fd950

Please sign in to comment.