Track/test/TrajectoryStateSOA_t.h

0001 #include "Geometry/CommonTopologies/interface/SimplePixelTopology.h"
0002 #include "CUDADataFormats/Track/interface/PixelTrackUtilities.h"
0003 #include "CUDADataFormats/Track/interface/TrackSoAHeterogeneousHost.h"
0004 #include "CUDADataFormats/Track/interface/TrackSoAHeterogeneousDevice.h"
0005
0006 using Vector5d = Eigen::Matrix<double, 5, 1>;
0007 using Matrix5d = Eigen::Matrix<double, 5, 5>;
0008 using helper = TracksUtilities<pixelTopology::Phase1>;
0009
0010 __host__ __device__ Matrix5d loadCov(Vector5d const& e) {
0011   Matrix5d cov;
0012   for (int i = 0; i < 5; ++i)
0013     cov(i, i) = e(i) * e(i);
0014   for (int i = 0; i < 5; ++i) {
0015     for (int j = 0; j < i; ++j) {
0016       double v = 0.3 * std::sqrt(cov(i, i) * cov(j, j));  // this makes the matrix pos defined
0017       cov(i, j) = (i + j) % 2 ? -0.4 * v : 0.1 * v;
0018       cov(j, i) = cov(i, j);
0019     }
0020   }
0021   return cov;
0022 }
0023
0024 template <typename TrackerTraits>
0025 __global__ void testTSSoA(TrackSoAView<TrackerTraits> ts) {
0026   Vector5d par0;
0027   par0 << 0.2, 0.1, 3.5, 0.8, 0.1;
0028   Vector5d e0;
0029   e0 << 0.01, 0.01, 0.035, -0.03, -0.01;
0030   auto cov0 = loadCov(e0);
0031
0032   int first = threadIdx.x + blockIdx.x * blockDim.x;
0033
0034   for (int i = first; i < ts.metadata().size(); i += blockDim.x * gridDim.x) {
0035     helper::copyFromDense(ts, par0, cov0, i);
0036     Vector5d par1;
0037     Matrix5d cov1;
0038     helper::copyToDense(ts, par1, cov1, i);
0039     Vector5d delV = par1 - par0;
0040     Matrix5d delM = cov1 - cov0;
0041     for (int j = 0; j < 5; ++j) {
0042       assert(std::abs(delV(j)) < 1.e-5);
0043       for (auto k = j; k < 5; ++k) {
0044         assert(cov0(k, j) == cov0(j, k));
0045         assert(cov1(k, j) == cov1(j, k));
0046         assert(std::abs(delM(k, j)) < 1.e-5);
0047       }
0048     }
0049   }
0050 }
0051
0052 #ifdef __CUDACC__
0053 #include "HeterogeneousCore/CUDAUtilities/interface/requireDevices.h"
0054 #include "HeterogeneousCore/CUDAUtilities/interface/cudaCheck.h"
0055 #endif
0056
0057 int main() {
0058 #ifdef __CUDACC__
0059   cms::cudatest::requireDevices();
0060   cudaStream_t stream;
0061   cudaCheck(cudaStreamCreateWithFlags(&stream, cudaStreamNonBlocking));
0062 #endif
0063
0064 #ifdef __CUDACC__
0065   // Since we are going to copy data from ts_d to ts_h, we
0066   // need to initialize the Host collection with a stream.
0067   TrackSoAHeterogeneousHost<pixelTopology::Phase1> ts_h(stream);
0068   TrackSoAHeterogeneousDevice<pixelTopology::Phase1> ts_d(stream);
0069 #else
0070   // If CUDA is not available, Host collection must not be initialized
0071   // with a stream.
0072   TrackSoAHeterogeneousHost<pixelTopology::Phase1> ts_h;
0073 #endif
0074
0075 #ifdef __CUDACC__
0076   testTSSoA<pixelTopology::Phase1><<<1, 64, 0, stream>>>(ts_d.view());
0077   cudaCheck(cudaGetLastError());
0078   cudaCheck(cudaMemcpyAsync(
0079       ts_h.buffer().get(), ts_d.const_buffer().get(), ts_d.bufferSize(), cudaMemcpyDeviceToHost, stream));
0080   cudaCheck(cudaGetLastError());
0081   cudaCheck(cudaStreamSynchronize(stream));
0082 #else
0083   testTSSoA<pixelTopology::Phase1>(ts_h.view());
0084 #endif
0085 }