Skip to content

Commit

Permalink
Skip CUDA-related tests if no GPU is present (#252)
Browse files Browse the repository at this point in the history 
Make unit tests that require a CUDA device skip the test and exit
succesfully if the CUDA runtime is not available, or no CUDA devices
are available.
  • Loading branch information
@fwyzard
fwyzard authored Jan 17, 2019
1 parent f437a80 commit 73d20e3
Show file tree
Hide file tree
Showing 37 changed files with 389 additions and 445 deletions.
13 changes: 7 additions & 6 deletions DataFormats/CaloRecHit/test/BuildFile.xml
View file
Original file line number Diff line number Diff line change
@@ -1,13 +1,14 @@
<bin name="testCaloCluster" file="testRunner.cpp,testCaloCluster.cppunit.cc">

<use name="DataFormats/CaloRecHit"/>
<use name="DataFormats/Math"/>
<use name="cppunit"/>
<bin name="testCaloCluster" file="testRunner.cpp,testCaloCluster.cppunit.cc">
<use name="cppunit"/>
<use name="DataFormats/CaloRecHit"/>
<use name="DataFormats/Math"/>
</bin>

<iftool name="cuda-gcc-support">
<bin name="test_calo_rechit" file="test_calo_rechit.cu">
<use name="cuda"/>
<use name="DataFormats/CaloRecHit"/>
<use name="DataFormats/DetId"/>
<use name="cuda"/>
<use name="HeterogeneousCore/CUDAUtilities"/>
</bin>
</iftool>
14 changes: 6 additions & 8 deletions DataFormats/CaloRecHit/test/test_calo_rechit.cu
View file
Original file line number Diff line number Diff line change
@@ -1,10 +1,11 @@
#include <cassert>
#include <iostream>

#include <cuda.h>
#include <cuda_runtime.h>

#include <iostream>
#include <cassert>

#include "DataFormats/CaloRecHit/interface/CaloRecHit.h"
#include "HeterogeneousCore/CUDAUtilities/interface/exitSansCUDADevices.h"

__global__ void kernel_test_calo_rechit(CaloRecHit* other) {
CaloRecHit rh{DetId(0), 10, 1, 0, 0};
Expand Down Expand Up @@ -42,12 +43,9 @@ void test_calo_rechit() {
}

int main(int argc, char** argv) {
int nDevices;
cudaGetDeviceCount(&nDevices);
std::cout << "nDevices = " << nDevices << std::endl;
exitSansCUDADevices();

if (nDevices > 0)
test_calo_rechit();
test_calo_rechit();

std::cout << "all good!" << std::endl;
return 0;
Expand Down
3 changes: 2 additions & 1 deletion DataFormats/DetId/test/BuildFile.xml
View file
Original file line number Diff line number Diff line change
@@ -1,6 +1,7 @@
<iftool name="cuda-gcc-support">
<bin name="test_detid" file="test_detid.cu">
<use name="DataFormats/DetId"/>
<use name="cuda"/>
<use name="DataFormats/DetId"/>
<use name="HeterogeneousCore/CUDAUtilities"/>
</bin>
</iftool>
16 changes: 7 additions & 9 deletions DataFormats/DetId/test/test_detid.cu
View file
Original file line number Diff line number Diff line change
@@ -1,10 +1,12 @@
#include <cuda_runtime.h>
#include <cassert>
#include <iostream>

#include <cuda.h>
#include <cuda_runtime.h>

#include <iostream>
#include <assert.h>
#include "DataFormats/DetId/interface/DetId.h"
#include "DataFormats/HcalDetId/interface/HcalDetId.h"
#include "HeterogeneousCore/CUDAUtilities/interface/exitSansCUDADevices.h"

__global__ void test_gen_detid(DetId* id, uint32_t const rawid) {
DetId did{rawid};
Expand All @@ -27,11 +29,7 @@ void test_detid() {
}

int main(int argc, char** argv) {
int nDevices;
cudaGetDeviceCount(&nDevices);
std::cout << "nDevices = " << nDevices << std::endl;
exitSansCUDADevices();

// test det id functionality
if (nDevices > 0)
test_detid();
test_detid();
}
2 changes: 2 additions & 0 deletions DataFormats/GeometrySurface/test/BuildFile.xml
View file
Original file line number Diff line number Diff line change
Expand Up @@ -18,12 +18,14 @@
<bin file="gpuFrameTransformKernel.cu, gpuFrameTransformTest.cpp" name="gpuFrameTransformTest">
<use name="cuda"/>
<use name="cuda-api-wrappers"/>
<use name="HeterogeneousCore/CUDAUtilities"/>
<flags CXXFLAGS="-g"/>
</bin>

<bin file="gpuFrameTransformKernel.cu, gpuFrameTransformTest.cpp" name="gpuFrameTransformTestRep">
<use name="cuda"/>
<use name="cuda-api-wrappers"/>
<use name="HeterogeneousCore/CUDAUtilities"/>
<flags CXXFLAGS="-ffp-contract=off"/>
<flags CUDA_FLAGS="-fmad=false -ftz=false -prec-div=true -prec-sqrt=true"/>
</bin>
39 changes: 16 additions & 23 deletions DataFormats/GeometrySurface/test/gpuFrameTransformKernel.cu
View file
Original file line number Diff line number Diff line change
@@ -1,45 +1,38 @@
#include "DataFormats/GeometrySurface/interface/SOARotation.h"
#include <cstdint>
#include <iostream>
#include <iomanip>

#include "cuda/api_wrappers.h"

#include "DataFormats/GeometrySurface/interface/SOARotation.h"

__global__
void toGlobal(SOAFrame<float> const * frame,
float const * xl, float const * yl,
float * x, float * y, float * z,
float const * xl, float const * yl,
float * x, float * y, float * z,
float const * le, float * ge,
uint32_t n)
uint32_t n)
{
int i = blockDim.x * blockIdx.x + threadIdx.x;
if (i >= n) return;

frame[0].toGlobal(xl[i],yl[i],x[i],y[i],z[i]);
frame[0].toGlobal(le[3*i],le[3*i+1],le[3*i+2],ge+6*i);

}

#include<iostream>
#include <iomanip>
#include "cuda/api_wrappers.h"


void toGlobalWrapper(SOAFrame<float> const * frame,
float const * xl, float const * yl,
float * x, float * y, float * z,
float const * le, float * ge,
uint32_t n) {

float const * xl, float const * yl,
float * x, float * y, float * z,
float const * le, float * ge,
uint32_t n)
{
int threadsPerBlock = 256;
int blocksPerGrid = (n + threadsPerBlock - 1) / threadsPerBlock;
std::cout
<< "CUDA toGlobal kernel launch with " << blocksPerGrid
<< " blocks of " << threadsPerBlock << " threads" << std::endl;

cuda::launch(
toGlobal,
{ blocksPerGrid, threadsPerBlock },
frame, xl, yl, x, y, z, le, ge,
n
);

cuda::launch(toGlobal,
{ blocksPerGrid, threadsPerBlock },
frame, xl, yl, x, y, z, le, ge, n);
}

111 changes: 42 additions & 69 deletions DataFormats/GeometrySurface/test/gpuFrameTransformTest.cpp
View file
Original file line number Diff line number Diff line change
@@ -1,140 +1,113 @@
#include "DataFormats/GeometrySurface/interface/SOARotation.h"
#include <algorithm>
#include <cassert>
#include <cmath>
#include <cstdint>



void toGlobalWrapper(SOAFrame<float> const * frame,
float const * xl, float const * yl,
float * x, float * y, float * z,
float const * le, float * ge,
uint32_t n);



#include "DataFormats/GeometrySurface/interface/TkRotation.h"
#include "DataFormats/GeometrySurface/interface/GloballyPositioned.h"


#include "cuda/api_wrappers.h"


#include <cstdint>
#include <iomanip>
#include <iostream>
#include <memory>
#include <algorithm>
#include<numeric>
#include<cmath>
#include<cassert>
#include <numeric>

#include <cuda/api_wrappers.h>

#include "DataFormats/GeometrySurface/interface/GloballyPositioned.h"
#include "DataFormats/GeometrySurface/interface/SOARotation.h"
#include "DataFormats/GeometrySurface/interface/TkRotation.h"
#include "HeterogeneousCore/CUDAUtilities/interface/exitSansCUDADevices.h"

#include<iostream>
#include <iomanip>
void toGlobalWrapper(SOAFrame<float> const * frame,
float const * xl, float const * yl,
float * x, float * y, float * z,
float const * le, float * ge,
uint32_t n);

int main(void)
{

typedef float T;
typedef TkRotation<T> Rotation;
typedef SOARotation<T> SRotation;
typedef GloballyPositioned<T> Frame;
typedef SOAFrame<T> SFrame;
typedef typename Frame::PositionType Position;
typedef typename Frame::GlobalVector GlobalVector;
typedef typename Frame::GlobalPoint GlobalPoint;
typedef typename Frame::LocalVector LocalVector;
typedef typename Frame::LocalPoint LocalPoint;



if (cuda::device::count() == 0) {
std::cerr << "No CUDA devices on this system" << "\n";
exit(EXIT_FAILURE);
}


exitSansCUDADevices();

typedef float T;
typedef TkRotation<T> Rotation;
typedef SOARotation<T> SRotation;
typedef GloballyPositioned<T> Frame;
typedef SOAFrame<T> SFrame;
typedef typename Frame::PositionType Position;
typedef typename Frame::GlobalVector GlobalVector;
typedef typename Frame::GlobalPoint GlobalPoint;
typedef typename Frame::LocalVector LocalVector;
typedef typename Frame::LocalPoint LocalPoint;

constexpr uint32_t size = 10000;
constexpr uint32_t size32 = size*sizeof(float);


float xl[size],yl[size];
float x[size],y[size],z[size];

// errors
float le[3*size];
float ge[6*size];


auto current_device = cuda::device::current::get();
auto d_xl = cuda::memory::device::make_unique<float[]>(current_device, size);
auto d_yl = cuda::memory::device::make_unique<float[]>(current_device, size);

auto d_x = cuda::memory::device::make_unique<float[]>(current_device, size);
auto d_y = cuda::memory::device::make_unique<float[]>(current_device, size);
auto d_z = cuda::memory::device::make_unique<float[]>(current_device, size);

auto d_le = cuda::memory::device::make_unique<float[]>(current_device, 3*size);
auto d_ge = cuda::memory::device::make_unique<float[]>(current_device, 6*size);


double a = 0.01;
double ca = std::cos(a);
double sa = std::sin(a);

Rotation r1(ca, sa, 0,
-sa, ca, 0,
0, 0, 1);
-sa, ca, 0,
0, 0, 1);
Frame f1(Position(2,3,4), r1);
std::cout << "f1.position() " << f1.position() << std::endl;
std::cout << "f1.rotation() " << '\n' << f1.rotation() << std::endl;

SFrame sf1(f1.position().x(),
f1.position().y(),
f1.position().z(),
f1.rotation()
);
f1.position().y(),
f1.position().z(),
f1.rotation()
);


// auto d_sf = cuda::memory::device::make_unique<SFrame[]>(current_device, 1);
auto d_sf = cuda::memory::device::make_unique<char[]>(current_device, sizeof(SFrame));
cuda::memory::copy(d_sf.get(), &sf1, sizeof(SFrame));




for (auto i=0U; i<size; ++i) {
xl[i]=yl[i] = 0.1f*float(i)-float(size/2);
le[3*i] = 0.01f; le[3*i+2] = (i>size/2) ? 1.f : 0.04f;
le[2*i+1]=0.;
}
std::random_shuffle(xl,xl+size);
std::random_shuffle(yl,yl+size);

cuda::memory::copy(d_xl.get(), xl, size32);
cuda::memory::copy(d_yl.get(), yl, size32);
cuda::memory::copy(d_le.get(), le, 3*size32);


toGlobalWrapper((SFrame const *)(d_sf.get()), d_xl.get(), d_yl.get(), d_x.get(), d_y.get(), d_z.get(),
d_le.get(), d_ge.get(), size
);

cuda::memory::copy(x,d_x.get(), size32);
cuda::memory::copy(y,d_y.get(), size32);
cuda::memory::copy(z,d_z.get(), size32);
cuda::memory::copy(ge,d_ge.get(), 6*size32);


float eps=0.;
for (auto i=0U; i<size; ++i) {
auto gp = f1.toGlobal(LocalPoint(xl[i],yl[i]));
eps = std::max(eps,std::abs(x[i]-gp.x()));
eps = std::max(eps,std::abs(y[i]-gp.y()));
eps = std::max(eps,std::abs(z[i]-gp.z()));
}
}

std::cout << "max eps " << eps << std::endl;



return 0;
}

1 change: 1 addition & 0 deletions DataFormats/HcalDetId/test/BuildFile.xml
View file
Original file line number Diff line number Diff line change
Expand Up @@ -2,5 +2,6 @@
<bin name="test_hcal_detid" file="test_hcal_detid.cu">
<use name="DataFormats/DetId"/>
<use name="DataFormats/HcalDetId"/>
<use name="HeterogeneousCore/CUDAUtilities"/>
</bin>
</iftool>
18 changes: 8 additions & 10 deletions DataFormats/HcalDetId/test/test_hcal_detid.cu
View file
Original file line number Diff line number Diff line change
@@ -1,10 +1,12 @@
#include <cuda_runtime.h>
#include <cassert>
#include <iostream>

#include <cuda.h>
#include <cuda_runtime.h>

#include <iostream>
#include <assert.h>
#include "DataFormats/DetId/interface/DetId.h"
#include "DataFormats/HcalDetId/interface/HcalDetId.h"
#include "HeterogeneousCore/CUDAUtilities/interface/exitSansCUDADevices.h"

__global__ void test_gen_detid(DetId* id) {
DetId did;
Expand Down Expand Up @@ -63,17 +65,13 @@ void test_hcal_detid() {
}

int main(int argc, char** argv) {
int nDevices;
cudaGetDeviceCount(&nDevices);
std::cout << "nDevices = " << nDevices << std::endl;
exitSansCUDADevices();

// test det id functionality
if (nDevices>0)
test_detid();
test_detid();

// test hcal det ids
if (nDevices>0)
test_hcal_detid();
test_hcal_detid();

return 0;
}
Loading

0 comments on commit 73d20e3

Please sign in to comment.