Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

GetXSecList performance improvements #2

Merged
merged 1 commit into from
Apr 14, 2016
Merged
Show file tree
Hide file tree
Changes from all commits
Commits
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
65 changes: 35 additions & 30 deletions geant4.10.02/source/externals/cuda/src/G4ParticleHPVector_CUDA.cu
Original file line number Diff line number Diff line change
Expand Up @@ -3,36 +3,38 @@
#include <cuda.h>
#include <cuda_runtime.h>
#include "G4ParticleHPVector_CUDA.hh"
#include <thrust/device_vector.h>
#include <stdio.h>
// #include "G4NeutronHPDataPoint.cu"
#include <iostream>
#include <math.h>

/***********************************************
* Device Methods
***********************************************/
__global__ void GetMinIndices_CUDA(G4ParticleHPDataPoint *d_theData, int nEntries,
double *d_energiesIn_xSecsOut, int numQueries, int *d_minIndices) {
const int idx = blockDim.x * blockIdx.x + threadIdx.x;
const int stepSize = (int)sqrt((float)nEntries);
__global__ void SetArrayTo(int *resArray, int querySize, int setValue)
{
int idx = blockDim.x*blockIdx.x + threadIdx.x; // determine threads ID
if(idx < querySize) {
resArray[idx] = setValue;
}
}

if (idx < numQueries) {
int i = 0;
double e = d_energiesIn_xSecsOut[idx];

for (i = 0; i < nEntries; i += stepSize) {
if (d_theData[i].energy >= e) {
break;
}
}

i = (i - (stepSize - 1) >= 0) ? i - (stepSize - 1) : 0;
for (; i < nEntries; i++) {
if (d_theData[i].energy >= e) {
break;
__global__ void findMinArray2(G4ParticleHPDataPoint *theData_d, G4double *queryArray_d, int *resArray_d, int numThreads, int querySize, int nEntries)
{
int idx = blockDim.x*blockIdx.x + threadIdx.x; // determine threads ID
for (int i = 0; i < querySize; i++){
G4double queryEnergy = queryArray_d[i];
for(int j = idx; j <= nEntries; j+= numThreads){// check threads designated chunk of data
if(theData_d[j].energy > queryEnergy){
atomicMin(&resArray_d[i], j);
}
}

d_minIndices[idx] = i;
}
}



/***********************************************
* Device Methods
***********************************************/
void G4ParticleHPVector_CUDA::SetInterpolationManager(G4InterpolationManager & aManager) {
theManager = aManager;
}
Expand Down Expand Up @@ -69,13 +71,16 @@ void G4ParticleHPVector_CUDA::GetXsecList(G4double* energiesIn_xSecsOut, G4int n
cudaMemcpy(d_theData, theData, sizeof(G4ParticleHPDataPoint) * nEntries, cudaMemcpyHostToDevice);
cudaMemcpy(d_energiesIn_xSecsOut, energiesIn_xSecsOut, sizeof(G4double) * numQueries, cudaMemcpyHostToDevice);

// need to add 1 block if doesn't divide evenly (e.g 32 T_P_B, 36 numQueries we need 1+1=2 blocks to get those last 4 queries)
int numBlocksSingleElement = numQueries/THREADS_PER_BLOCK + ((numQueries % THREADS_PER_BLOCK == 0) ? 0 : 1);

GetMinIndices_CUDA <<<numBlocksSingleElement, THREADS_PER_BLOCK>>>
(d_theData, nEntries, d_energiesIn_xSecsOut, numQueries, d_minIndices);

cudaMemcpy(minIndices, d_minIndices, sizeof(G4int) * numQueries, cudaMemcpyDeviceToHost);
// NEW IMPL =========================================================================
int queryBlocks = numQueries/THREADS_PER_BLOCK + (numQueries % THREADS_PER_BLOCK == 0 ? 0:1);
G4double resultVal = 0;
int dataChunk = 32;
int threadNum = nEntries/dataChunk;
int arrayBlocks = threadNum/THREADS_PER_BLOCK + (threadNum % THREADS_PER_BLOCK == 0 ? 0:1);

SetArrayTo <<< queryBlocks, THREADS_PER_BLOCK >>>(d_minIndices, numQueries, nEntries-1);
findMinArray2 <<< arrayBlocks, THREADS_PER_BLOCK >>> (d_theData, d_energiesIn_xSecsOut, d_minIndices, threadNum, numQueries, nEntries);
cudaMemcpy(minIndices, d_minIndices, numQueries * sizeof(G4int), cudaMemcpyDeviceToHost);

for (int i = 0; i < numQueries; i++) {
int minIndex = minIndices[i];
Expand Down
Original file line number Diff line number Diff line change
Expand Up @@ -68,7 +68,7 @@ hash: 17710924904126496405
@numQueries=10000
hash: 7506977939755255107
@numQueries=100000
hash: 1638135657286475825
hash: 10300846530834065403
#void Init(std::istream & aDataFile, G4int total, G4double ux=1., G4double uy=1.)_5
theData xVals hash: 7171123492023199933
theData yVals hash: 9268461341921925588
Expand Down
Original file line number Diff line number Diff line change
@@ -0,0 +1,42 @@
Method Signature,Case Number,nEntries,Input,CPU Time,GPU Time
void G4ParticleHPVector_CUDA::GetXsecBuffer(G4double * queryList. G4int length),0,0,numQueries=10,4.05312e-06,2.14577e-06
void G4ParticleHPVector_CUDA::GetXsecBuffer(G4double * queryList. G4int length),0,0,numQueries=50,0,0
void G4ParticleHPVector_CUDA::GetXsecBuffer(G4double * queryList. G4int length),0,0,numQueries=100,9.53674e-07,9.53674e-07
void G4ParticleHPVector_CUDA::GetXsecBuffer(G4double * queryList. G4int length),0,0,numQueries=10000,0.000500917,5.00679e-05
void G4ParticleHPVector_CUDA::GetXsecBuffer(G4double * queryList. G4int length),0,0,numQueries=100000,0.000389099,0.00041914
void Init(std::istream & aDataFile. G4int total. G4double ux=1.. G4double uy=1.),1,,,0.000202894,0.000185966
void G4ParticleHPVector_CUDA::GetXsecBuffer(G4double * queryList. G4int length),1,80,numQueries=10,8.82149e-06,7.15256e-06
void G4ParticleHPVector_CUDA::GetXsecBuffer(G4double * queryList. G4int length),1,80,numQueries=50,4.05312e-06,1.90735e-06
void G4ParticleHPVector_CUDA::GetXsecBuffer(G4double * queryList. G4int length),1,80,numQueries=100,6.19888e-06,4.05312e-06
void G4ParticleHPVector_CUDA::GetXsecBuffer(G4double * queryList. G4int length),1,80,numQueries=10000,0.000139952,0.000111818
void G4ParticleHPVector_CUDA::GetXsecBuffer(G4double * queryList. G4int length),1,80,numQueries=100000,0.00121808,0.00114107
void Init(std::istream & aDataFile. G4int total. G4double ux=1.. G4double uy=1.),2,,,0.00348091,0.00426793
void G4ParticleHPVector_CUDA::GetXsecBuffer(G4double * queryList. G4int length),2,1509,numQueries=10,2.14577e-06,1.90735e-06
void G4ParticleHPVector_CUDA::GetXsecBuffer(G4double * queryList. G4int length),2,1509,numQueries=50,2.14577e-06,1.90735e-06
void G4ParticleHPVector_CUDA::GetXsecBuffer(G4double * queryList. G4int length),2,1509,numQueries=100,5.00679e-06,5.00679e-06
void G4ParticleHPVector_CUDA::GetXsecBuffer(G4double * queryList. G4int length),2,1509,numQueries=10000,0.000267029,0.000273943
void G4ParticleHPVector_CUDA::GetXsecBuffer(G4double * queryList. G4int length),2,1509,numQueries=100000,0.00193715,0.00187492
void Init(std::istream & aDataFile. G4int total. G4double ux=1.. G4double uy=1.),3,,,0.023423,0.0248361
void G4ParticleHPVector_CUDA::GetXsecBuffer(G4double * queryList. G4int length),3,8045,numQueries=10,5.00679e-06,1.90735e-06
void G4ParticleHPVector_CUDA::GetXsecBuffer(G4double * queryList. G4int length),3,8045,numQueries=50,2.14577e-06,1.90735e-06
void G4ParticleHPVector_CUDA::GetXsecBuffer(G4double * queryList. G4int length),3,8045,numQueries=100,5.00679e-06,5.00679e-06
void G4ParticleHPVector_CUDA::GetXsecBuffer(G4double * queryList. G4int length),3,8045,numQueries=10000,0.000792027,0.000674963
void G4ParticleHPVector_CUDA::GetXsecBuffer(G4double * queryList. G4int length),3,8045,numQueries=100000,0.00246,0.00214505
void Init(std::istream & aDataFile. G4int total. G4double ux=1.. G4double uy=1.),4,,,0.105233,0.110869
void G4ParticleHPVector_CUDA::GetXsecBuffer(G4double * queryList. G4int length),4,41854,numQueries=10,2.86102e-06,1.90735e-06
void G4ParticleHPVector_CUDA::GetXsecBuffer(G4double * queryList. G4int length),4,41854,numQueries=50,4.05312e-06,2.86102e-06
void G4ParticleHPVector_CUDA::GetXsecBuffer(G4double * queryList. G4int length),4,41854,numQueries=100,7.15256e-06,5.00679e-06
void G4ParticleHPVector_CUDA::GetXsecBuffer(G4double * queryList. G4int length),4,41854,numQueries=10000,0.00113606,0.0011971
void G4ParticleHPVector_CUDA::GetXsecBuffer(G4double * queryList. G4int length),4,41854,numQueries=100000,0.00560093,0.0082221
void Init(std::istream & aDataFile. G4int total. G4double ux=1.. G4double uy=1.),5,,,0.171404,0.169494
void G4ParticleHPVector_CUDA::GetXsecBuffer(G4double * queryList. G4int length),5,98995,numQueries=10,2.14577e-06,1.90735e-06
void G4ParticleHPVector_CUDA::GetXsecBuffer(G4double * queryList. G4int length),5,98995,numQueries=50,2.86102e-06,1.90735e-06
void G4ParticleHPVector_CUDA::GetXsecBuffer(G4double * queryList. G4int length),5,98995,numQueries=100,5.00679e-06,5.00679e-06
void G4ParticleHPVector_CUDA::GetXsecBuffer(G4double * queryList. G4int length),5,98995,numQueries=10000,0.000799894,0.00079608
void G4ParticleHPVector_CUDA::GetXsecBuffer(G4double * queryList. G4int length),5,98995,numQueries=100000,0.0110049,0.0108991
void Init(std::istream & aDataFile. G4int total. G4double ux=1.. G4double uy=1.),6,,,0.527347,0.508139
void G4ParticleHPVector_CUDA::GetXsecBuffer(G4double * queryList. G4int length),6,242594,numQueries=10,2.14577e-06,2.14577e-06
void G4ParticleHPVector_CUDA::GetXsecBuffer(G4double * queryList. G4int length),6,242594,numQueries=50,1.90735e-06,1.90735e-06
void G4ParticleHPVector_CUDA::GetXsecBuffer(G4double * queryList. G4int length),6,242594,numQueries=100,5.96046e-06,5.00679e-06
void G4ParticleHPVector_CUDA::GetXsecBuffer(G4double * queryList. G4int length),6,242594,numQueries=10000,0.000939131,0.000840187
void G4ParticleHPVector_CUDA::GetXsecBuffer(G4double * queryList. G4int length),6,242594,numQueries=100000,0.014972,0.011513
Original file line number Diff line number Diff line change
@@ -0,0 +1,42 @@
Method Signature,Case Number,nEntries,Input,CPU Time,GPU Time
void G4ParticleHPVector_CUDA::GetXsecBuffer(G4double * queryList. G4int length),0,0,numQueries=10,4.05312e-06,1.90735e-06
void G4ParticleHPVector_CUDA::GetXsecBuffer(G4double * queryList. G4int length),0,0,numQueries=50,0,0
void G4ParticleHPVector_CUDA::GetXsecBuffer(G4double * queryList. G4int length),0,0,numQueries=100,9.53674e-07,0
void G4ParticleHPVector_CUDA::GetXsecBuffer(G4double * queryList. G4int length),0,0,numQueries=10000,0.000500917,5.00679e-05
void G4ParticleHPVector_CUDA::GetXsecBuffer(G4double * queryList. G4int length),0,0,numQueries=100000,0.000389099,0.000422955
void Init(std::istream & aDataFile. G4int total. G4double ux=1.. G4double uy=1.),1,,,0.000202894,0.000238895
void G4ParticleHPVector_CUDA::GetXsecBuffer(G4double * queryList. G4int length),1,80,numQueries=10,8.82149e-06,9.05991e-06
void G4ParticleHPVector_CUDA::GetXsecBuffer(G4double * queryList. G4int length),1,80,numQueries=50,4.05312e-06,1.90735e-06
void G4ParticleHPVector_CUDA::GetXsecBuffer(G4double * queryList. G4int length),1,80,numQueries=100,6.19888e-06,4.05312e-06
void G4ParticleHPVector_CUDA::GetXsecBuffer(G4double * queryList. G4int length),1,80,numQueries=10000,0.000139952,0.000111818
void G4ParticleHPVector_CUDA::GetXsecBuffer(G4double * queryList. G4int length),1,80,numQueries=100000,0.00121808,0.00127101
void Init(std::istream & aDataFile. G4int total. G4double ux=1.. G4double uy=1.),2,,,0.00348091,0.00367308
void G4ParticleHPVector_CUDA::GetXsecBuffer(G4double * queryList. G4int length),2,1509,numQueries=10,2.14577e-06,1.19209e-06
void G4ParticleHPVector_CUDA::GetXsecBuffer(G4double * queryList. G4int length),2,1509,numQueries=50,2.14577e-06,3.09944e-06
void G4ParticleHPVector_CUDA::GetXsecBuffer(G4double * queryList. G4int length),2,1509,numQueries=100,5.00679e-06,5.00679e-06
void G4ParticleHPVector_CUDA::GetXsecBuffer(G4double * queryList. G4int length),2,1509,numQueries=10000,0.000267029,0.000267982
void G4ParticleHPVector_CUDA::GetXsecBuffer(G4double * queryList. G4int length),2,1509,numQueries=100000,0.00193715,0.00184894
void Init(std::istream & aDataFile. G4int total. G4double ux=1.. G4double uy=1.),3,,,0.023423,0.026448
void G4ParticleHPVector_CUDA::GetXsecBuffer(G4double * queryList. G4int length),3,8045,numQueries=10,5.00679e-06,3.09944e-06
void G4ParticleHPVector_CUDA::GetXsecBuffer(G4double * queryList. G4int length),3,8045,numQueries=50,2.14577e-06,3.09944e-06
void G4ParticleHPVector_CUDA::GetXsecBuffer(G4double * queryList. G4int length),3,8045,numQueries=100,5.00679e-06,5.96046e-06
void G4ParticleHPVector_CUDA::GetXsecBuffer(G4double * queryList. G4int length),3,8045,numQueries=10000,0.000792027,0.000694036
void G4ParticleHPVector_CUDA::GetXsecBuffer(G4double * queryList. G4int length),3,8045,numQueries=100000,0.00246,0.0020721
void Init(std::istream & aDataFile. G4int total. G4double ux=1.. G4double uy=1.),4,,,0.105233,0.107445
void G4ParticleHPVector_CUDA::GetXsecBuffer(G4double * queryList. G4int length),4,41854,numQueries=10,2.86102e-06,1.90735e-06
void G4ParticleHPVector_CUDA::GetXsecBuffer(G4double * queryList. G4int length),4,41854,numQueries=50,4.05312e-06,1.40667e-05
void G4ParticleHPVector_CUDA::GetXsecBuffer(G4double * queryList. G4int length),4,41854,numQueries=100,7.15256e-06,5.00679e-06
void G4ParticleHPVector_CUDA::GetXsecBuffer(G4double * queryList. G4int length),4,41854,numQueries=10000,0.00113606,0.000864983
void G4ParticleHPVector_CUDA::GetXsecBuffer(G4double * queryList. G4int length),4,41854,numQueries=100000,0.00560093,0.00549889
void Init(std::istream & aDataFile. G4int total. G4double ux=1.. G4double uy=1.),5,,,0.171404,0.174278
void G4ParticleHPVector_CUDA::GetXsecBuffer(G4double * queryList. G4int length),5,98995,numQueries=10,2.14577e-06,1.90735e-06
void G4ParticleHPVector_CUDA::GetXsecBuffer(G4double * queryList. G4int length),5,98995,numQueries=50,2.86102e-06,1.90735e-06
void G4ParticleHPVector_CUDA::GetXsecBuffer(G4double * queryList. G4int length),5,98995,numQueries=100,5.00679e-06,5.00679e-06
void G4ParticleHPVector_CUDA::GetXsecBuffer(G4double * queryList. G4int length),5,98995,numQueries=10000,0.000799894,0.00879192
void G4ParticleHPVector_CUDA::GetXsecBuffer(G4double * queryList. G4int length),5,98995,numQueries=100000,0.0110049,0.01284
void Init(std::istream & aDataFile. G4int total. G4double ux=1.. G4double uy=1.),6,,,0.527347,0.535003
void G4ParticleHPVector_CUDA::GetXsecBuffer(G4double * queryList. G4int length),6,242594,numQueries=10,2.14577e-06,3.09944e-06
void G4ParticleHPVector_CUDA::GetXsecBuffer(G4double * queryList. G4int length),6,242594,numQueries=50,1.90735e-06,2.14577e-06
void G4ParticleHPVector_CUDA::GetXsecBuffer(G4double * queryList. G4int length),6,242594,numQueries=100,5.96046e-06,5.96046e-06
void G4ParticleHPVector_CUDA::GetXsecBuffer(G4double * queryList. G4int length),6,242594,numQueries=10000,0.000939131,0.000936985
void G4ParticleHPVector_CUDA::GetXsecBuffer(G4double * queryList. G4int length),6,242594,numQueries=100000,0.014972,0.01388
Original file line number Diff line number Diff line change
@@ -0,0 +1,42 @@
Method Signature,Case Number,nEntries,Input,CPU Time,GPU Time
void G4ParticleHPVector_CUDA::GetXsecBuffer(G4double * queryList. G4int length),0,0,numQueries=10,4.05312e-06,1.90735e-06
void G4ParticleHPVector_CUDA::GetXsecBuffer(G4double * queryList. G4int length),0,0,numQueries=50,0,9.53674e-07
void G4ParticleHPVector_CUDA::GetXsecBuffer(G4double * queryList. G4int length),0,0,numQueries=100,9.53674e-07,9.53674e-07
void G4ParticleHPVector_CUDA::GetXsecBuffer(G4double * queryList. G4int length),0,0,numQueries=10000,0.000500917,5.00679e-05
void G4ParticleHPVector_CUDA::GetXsecBuffer(G4double * queryList. G4int length),0,0,numQueries=100000,0.000389099,0.000420094
void Init(std::istream & aDataFile. G4int total. G4double ux=1.. G4double uy=1.),1,,,0.000202894,0.00019908
void G4ParticleHPVector_CUDA::GetXsecBuffer(G4double * queryList. G4int length),1,80,numQueries=10,8.82149e-06,9.05991e-06
void G4ParticleHPVector_CUDA::GetXsecBuffer(G4double * queryList. G4int length),1,80,numQueries=50,4.05312e-06,3.8147e-06
void G4ParticleHPVector_CUDA::GetXsecBuffer(G4double * queryList. G4int length),1,80,numQueries=100,6.19888e-06,4.05312e-06
void G4ParticleHPVector_CUDA::GetXsecBuffer(G4double * queryList. G4int length),1,80,numQueries=10000,0.000139952,0.000117779
void G4ParticleHPVector_CUDA::GetXsecBuffer(G4double * queryList. G4int length),1,80,numQueries=100000,0.00121808,0.00111699
void Init(std::istream & aDataFile. G4int total. G4double ux=1.. G4double uy=1.),2,,,0.00348091,0.00348496
void G4ParticleHPVector_CUDA::GetXsecBuffer(G4double * queryList. G4int length),2,1509,numQueries=10,2.14577e-06,9.53674e-07
void G4ParticleHPVector_CUDA::GetXsecBuffer(G4double * queryList. G4int length),2,1509,numQueries=50,2.14577e-06,3.09944e-06
void G4ParticleHPVector_CUDA::GetXsecBuffer(G4double * queryList. G4int length),2,1509,numQueries=100,5.00679e-06,5.00679e-06
void G4ParticleHPVector_CUDA::GetXsecBuffer(G4double * queryList. G4int length),2,1509,numQueries=10000,0.000267029,0.000267029
void G4ParticleHPVector_CUDA::GetXsecBuffer(G4double * queryList. G4int length),2,1509,numQueries=100000,0.00193715,0.00184393
void Init(std::istream & aDataFile. G4int total. G4double ux=1.. G4double uy=1.),3,,,0.023423,0.02549
void G4ParticleHPVector_CUDA::GetXsecBuffer(G4double * queryList. G4int length),3,8045,numQueries=10,5.00679e-06,4.05312e-06
void G4ParticleHPVector_CUDA::GetXsecBuffer(G4double * queryList. G4int length),3,8045,numQueries=50,2.14577e-06,2.14577e-06
void G4ParticleHPVector_CUDA::GetXsecBuffer(G4double * queryList. G4int length),3,8045,numQueries=100,5.00679e-06,5.96046e-06
void G4ParticleHPVector_CUDA::GetXsecBuffer(G4double * queryList. G4int length),3,8045,numQueries=10000,0.000792027,0.000828028
void G4ParticleHPVector_CUDA::GetXsecBuffer(G4double * queryList. G4int length),3,8045,numQueries=100000,0.00246,0.00222898
void Init(std::istream & aDataFile. G4int total. G4double ux=1.. G4double uy=1.),4,,,0.105233,0.113319
void G4ParticleHPVector_CUDA::GetXsecBuffer(G4double * queryList. G4int length),4,41854,numQueries=10,2.86102e-06,1.90735e-06
void G4ParticleHPVector_CUDA::GetXsecBuffer(G4double * queryList. G4int length),4,41854,numQueries=50,4.05312e-06,2.86102e-06
void G4ParticleHPVector_CUDA::GetXsecBuffer(G4double * queryList. G4int length),4,41854,numQueries=100,7.15256e-06,5.00679e-06
void G4ParticleHPVector_CUDA::GetXsecBuffer(G4double * queryList. G4int length),4,41854,numQueries=10000,0.00113606,0.000825882
void G4ParticleHPVector_CUDA::GetXsecBuffer(G4double * queryList. G4int length),4,41854,numQueries=100000,0.00560093,0.00593495
void Init(std::istream & aDataFile. G4int total. G4double ux=1.. G4double uy=1.),5,,,0.171404,0.174828
void G4ParticleHPVector_CUDA::GetXsecBuffer(G4double * queryList. G4int length),5,98995,numQueries=10,2.14577e-06,1.90735e-06
void G4ParticleHPVector_CUDA::GetXsecBuffer(G4double * queryList. G4int length),5,98995,numQueries=50,2.86102e-06,3.09944e-06
void G4ParticleHPVector_CUDA::GetXsecBuffer(G4double * queryList. G4int length),5,98995,numQueries=100,5.00679e-06,5.00679e-06
void G4ParticleHPVector_CUDA::GetXsecBuffer(G4double * queryList. G4int length),5,98995,numQueries=10000,0.000799894,0.000813007
void G4ParticleHPVector_CUDA::GetXsecBuffer(G4double * queryList. G4int length),5,98995,numQueries=100000,0.0110049,0.011142
void Init(std::istream & aDataFile. G4int total. G4double ux=1.. G4double uy=1.),6,,,0.527347,0.507052
void G4ParticleHPVector_CUDA::GetXsecBuffer(G4double * queryList. G4int length),6,242594,numQueries=10,2.14577e-06,2.14577e-06
void G4ParticleHPVector_CUDA::GetXsecBuffer(G4double * queryList. G4int length),6,242594,numQueries=50,1.90735e-06,3.09944e-06
void G4ParticleHPVector_CUDA::GetXsecBuffer(G4double * queryList. G4int length),6,242594,numQueries=100,5.96046e-06,5.00679e-06
void G4ParticleHPVector_CUDA::GetXsecBuffer(G4double * queryList. G4int length),6,242594,numQueries=10000,0.000939131,0.000813961
void G4ParticleHPVector_CUDA::GetXsecBuffer(G4double * queryList. G4int length),6,242594,numQueries=100000,0.014972,0.011879
Loading