Simplify cudacompat layer to use a 1-dimensional grid (#586)

Remove the possibility of changing the grid size used by the
cms::cudacompat layer, and make it a constant equal to {1, 1, 1}.

This avoids a thread-related problem caused by TBB using worker threads
where the grid size had not been initialised.

The kernel for pixel clustering need to be rewritten to support a
one-dimensional grid to run on the CPU.
Currently they are only used on the GPU in the Patatrack workflows, but
they are exercised on the CPU by the gpuClustering_t tests; those tests
have been commented out until the kernels can be updated.

RecoPixelVertexing/PixelVertexFinding/plugins/gpuVertexFinderImpl.h

@@ -109,7 +109,6 @@ namespace gpuVertexFinder {
     loadTracks<<<numberOfBlocks, blockSize, 0, stream>>>(tksoa, soa, ws_d.get(), ptMin);
     cudaCheck(cudaGetLastError());
 #else
-    cms::cudacompat::resetGrid();
     init(soa, ws_d.get());
     loadTracks(tksoa, soa, ws_d.get(), ptMin);
 #endif
@@ -157,10 +156,7 @@ namespace gpuVertexFinder {
     // std::cout << "found " << (*ws_d).nvIntermediate << " vertices " << std::endl;
     fitVertices(soa, ws_d.get(), 50.);
     // one block per vertex!
-    blockIdx.x = 0;
-    gridDim.x = 1;
     splitVertices(soa, ws_d.get(), 9.f);
-    resetGrid();
     fitVertices(soa, ws_d.get(), 5000.);
     sortByPt2(soa, ws_d.get());
 #endif

RecoPixelVertexing/PixelVertexFinding/test/VertexFinder_t.h

@@ -266,10 +266,7 @@ int main() {
       cms::cuda::launch(gpuVertexFinder::splitVerticesKernel, {1024, 64}, onGPU_d.get(), ws_d.get(), 9.f);
       cudaCheck(cudaMemcpy(&nv, LOC_WS(nvIntermediate), sizeof(uint32_t), cudaMemcpyDeviceToHost));
 #else
-      gridDim.x = 1;
-      assert(blockIdx.x == 0);
       splitVertices(onGPU_d.get(), ws_d.get(), 9.f);
-      resetGrid();
       nv = ws_d->nvIntermediate;
 #endif
       std::cout << "after split " << nv << std::endl;