Warning, /CUDADataFormats/TrackingRecHit/test/TrackingRecHitSoA_test.cu is written in an unsupported language. File is not indexed.
0001 #include "CUDADataFormats/TrackingRecHit/interface/TrackingRecHitsUtilities.h"
0002 #include "CUDADataFormats/TrackingRecHit/interface/TrackingRecHitSoADevice.h"
0003
0004 namespace testTrackingRecHitSoA {
0005
0006 template <typename TrackerTraits>
0007 __global__ void fill(TrackingRecHitSoAView<TrackerTraits> soa) {
0008 int i = threadIdx.x;
0009 int j = blockIdx.x;
0010 if (i == 0 and j == 0) {
0011 soa.offsetBPIX2() = 22;
0012 soa[10].xLocal() = 1.11;
0013 }
0014
0015 soa[i].iphi() = i % 10;
0016 soa.hitsLayerStart()[j] = j;
0017 __syncthreads();
0018 }
0019
0020 template <typename TrackerTraits>
0021 __global__ void show(TrackingRecHitSoAView<TrackerTraits> soa) {
0022 int i = threadIdx.x;
0023 int j = blockIdx.x;
0024
0025 if (i == 0 and j == 0) {
0026 printf("nbins = %d \n", soa.phiBinner().nbins());
0027 printf("offsetBPIX %d ->%d \n", i, soa.offsetBPIX2());
0028 printf("nHits %d ->%d \n", i, soa.nHits());
0029 printf("hitsModuleStart %d ->%d \n", i, soa.hitsModuleStart().at(28));
0030 }
0031
0032 if (i < 10) // can be increased to soa.nHits() for debugging
0033 printf("iPhi %d ->%d \n", i, soa[i].iphi());
0034
0035 if (j * blockDim.x + i < 10) // can be increased to soa.phiBinner().nbins() for debugging
0036 printf(">bin size %d ->%d \n", j * blockDim.x + i, soa.phiBinner().size(j * blockDim.x + i));
0037 __syncthreads();
0038 }
0039
0040 template <typename TrackerTraits>
0041 void runKernels(TrackingRecHitSoADevice<TrackerTraits>& hits, cudaStream_t stream) {
0042 printf("> RUN!\n");
0043 fill<TrackerTraits><<<10, 100, 0, stream>>>(hits.view());
0044
0045 cudaCheck(cudaDeviceSynchronize());
0046 cms::cuda::fillManyFromVector(&(hits.view().phiBinner()),
0047 10,
0048 hits.view().iphi(),
0049 hits.view().hitsLayerStart().data(),
0050 2000,
0051 256,
0052 hits.view().phiBinnerStorage(),
0053 stream);
0054 cudaCheck(cudaDeviceSynchronize());
0055 show<TrackerTraits><<<10, 1000, 0, stream>>>(hits.view());
0056 cudaCheck(cudaDeviceSynchronize());
0057 }
0058
0059 template void runKernels<pixelTopology::Phase1>(TrackingRecHitSoADevice<pixelTopology::Phase1>& hits,
0060 cudaStream_t stream);
0061 template void runKernels<pixelTopology::Phase2>(TrackingRecHitSoADevice<pixelTopology::Phase2>& hits,
0062 cudaStream_t stream);
0063
0064 } // namespace testTrackingRecHitSoA
