eic-opticks/qudarap/QCerenkov.cc

0001 #include <sstream>
0002 #include <algorithm>
0003 #include <cuda_runtime.h>
0004
0005 #include "SLOG.hh"
0006 #include "spath.h"
0007 #include "scuda.h"
0008 #include "squad.h"
0009 #include "sphoton.h"
0010 #include "scerenkov.h"
0011
0012 #include "qrng.h"
0013 #include "qcerenkov.h"
0014
0015 #include "NP.hh"
0016
0017 #include "QUDA_CHECK.h"
0018 #include "QBase.hh"
0019 #include "QCerenkov.hh"
0020 #include "QTex.hh"
0021 #include "QTexMaker.hh"
0022 #include "QTexLookup.hh"
0023 #include "QBnd.hh"
0024 #include "QProp.hh"
0025 #include "QU.hh"
0026 #include "QCK.hh"
0027
0028
0029 const plog::Severity QCerenkov::LEVEL = SLOG::EnvLevel("QCerenkov", "DEBUG");
0030
0031 const QCerenkov* QCerenkov::INSTANCE = nullptr ;
0032 const QCerenkov* QCerenkov::Get(){ return INSTANCE ;  }
0033
0034
0035 const char* QCerenkov::DEFAULT_FOLD = "$TMP/QCerenkovIntegralTest/test_makeICDF_SplitBin" ;
0036
0037 NP* QCerenkov::Load(const char* fold, const char* name)  // static
0038 {
0039     const char* path = spath::Resolve(fold);
0040     NP* a = NP::Load(path);
0041     LOG_IF(fatal, !a) << " failed to load array from path " << path ;
0042     return a ;
0043 }
0044
0045 /**
0046 QCerenkov::MakeTex
0047 --------------------
0048
0049 **/
0050
0051 QTex<float4>* QCerenkov::MakeTex(const NP* icdf, char filterMode, bool normalizedCoords ) // static
0052 {
0053     LOG(info) << " icdf " << icdf << " icdf.sstr " << icdf->sstr() << " filterMode " << filterMode ;
0054     QTex<float4>* tex = QTexMaker::Make2d_f4(icdf, filterMode, normalizedCoords );
0055     LOG(info) << " tex " << tex ;
0056     return tex ;
0057 }
0058
0059 /**
0060 QCerenkov::MakeInstance
0061 ------------------------
0062
0063 QProp was assuming a saved GGeo and IDPath and access to "$IDPath/GScintillatorLib/LS_ori/RINDEX.npy"
0064 see GGeo::convertSim_Prop
0065
0066 **/
0067
0068 qcerenkov* QCerenkov::MakeInstance() // static
0069 {
0070     const QBase* base = QBase::Get();
0071     assert( base );
0072
0073     const QBnd* bnd = QBnd::Get();
0074     assert( bnd );
0075
0076     const QProp<float>* prop = QProp<float>::Get();
0077     // assert(prop);
0078
0079     qcerenkov* ck = new qcerenkov ;
0080     ck->base = base->d_base ;
0081     ck->bnd = bnd->d_qb ;
0082     ck->prop = prop ? prop->d_prop : nullptr ;
0083     return ck ;
0084 }
0085
0086
0087
0088 /**
0089 QCerenkov::QCerenkov
0090 ----------------------
0091
0092 Currently relies on icdf created by QCerenkovIntegralTest
0093
0094 TODO: formalize icdf creation into the pre-cache geometry conversion and icdf location into geocache
0095
0096 **/
0097
0098 QCerenkov::QCerenkov(const char* fold_ )
0099     :
0100     fold(fold_),
0101     icdf_( nullptr ),
0102     icdf( nullptr ),
0103     filterMode('P'),    // CAUTION: P: is Point mode with interpolation disabled
0104     normalizedCoords(true),
0105     tex(nullptr),
0106     look(nullptr),
0107     cerenkov(MakeInstance()),
0108     d_cerenkov(QU::UploadArray<qcerenkov>(cerenkov, 1, "QCerenkov::QCerenkov/d_cerenkov.1"))
0109 {
0110     init();
0111 }
0112
0113 QCerenkov::QCerenkov()
0114     :
0115     fold(nullptr),  // HMM: whats the point of this with no fold ?
0116     icdf_(nullptr),
0117     icdf(nullptr),
0118     filterMode('P'),
0119     normalizedCoords(true),
0120     tex(nullptr),
0121     look(nullptr),
0122     cerenkov(MakeInstance()),
0123     d_cerenkov(QU::UploadArray<qcerenkov>(cerenkov, 1,"QCerenkov::QCerenkov/d_cerenkov.0"))
0124 {
0125     init();
0126 }
0127
0128
0129
0130 /**
0131 QCerenkov::init
0132 ----------------
0133
0134 TODO: move most of this into statics
0135
0136 TODO: treat icdf just like bnd, base, prop ?
0137
0138 **/
0139
0140
0141 void QCerenkov::init()
0142 {
0143     INSTANCE = this ;
0144     if(fold == nullptr) return ;
0145
0146     icdf_ = Load(fold, "icdf.npy");
0147     if( icdf_ == nullptr )
0148     {
0149         LOG(fatal) << "failed to load icdf.npy from fold " << fold ;
0150         return ;
0151     }
0152
0153     LOG(info)
0154         << " icdf_ " << icdf_
0155         << " icdf_.sstr " << icdf_->sstr()
0156         << " icdf_.meta " << icdf_->meta
0157         ;
0158
0159     icdf = icdf_->ebyte == 4 ? icdf_ : NP::MakeNarrow(icdf_) ;
0160
0161     LOG(info)
0162         << " icdf " << icdf
0163         << " icdf.sstr " << icdf->sstr()
0164         << " icdf.meta " << icdf->meta
0165         ;
0166
0167     tex = MakeTex(icdf, filterMode, normalizedCoords ) ;
0168
0169     LOG(info) << " tex " << tex ;
0170
0171     look = new QTexLookup<float4>( tex ) ;
0172
0173     LOG(info) << " look " << look ;
0174 }
0175
0176
0177 NP* QCerenkov::lookup()
0178 {
0179     return look ? look->lookup() : nullptr ;
0180 }
0181
0182
0183 std::string QCerenkov::desc() const
0184 {
0185     std::stringstream ss ;
0186     ss << "QCerenkov"
0187        << " fold " << ( fold ? fold : "-" )
0188        << " icdf_ " << ( icdf_ ? icdf_->sstr() : "-" )
0189        << " icdf " << ( icdf ? icdf->sstr() : "-" )
0190        << " tex " << tex
0191        ;
0192
0193     std::string s = ss.str();
0194     return s ;
0195 }
0196
0197
0198
0199
0200 extern "C" void QCerenkov_check(dim3 numBlocks, dim3 threadsPerBlock, unsigned width, unsigned height  );
0201
0202 template <typename T>
0203 extern void QCerenkov_lookup(dim3 numBlocks, dim3 threadsPerBlock, cudaTextureObject_t texObj, quad4* meta, T* lookup, unsigned num_lookup, unsigned width, unsigned height  );
0204
0205
0206 void QCerenkov::configureLaunch( dim3& numBlocks, dim3& threadsPerBlock, unsigned width, unsigned height )
0207 {
0208     threadsPerBlock.x = 512 ;
0209     threadsPerBlock.y = 1 ;
0210     threadsPerBlock.z = 1 ;
0211
0212     numBlocks.x = (width + threadsPerBlock.x - 1) / threadsPerBlock.x ;
0213     numBlocks.y = (height + threadsPerBlock.y - 1) / threadsPerBlock.y ;
0214     numBlocks.z = 1 ;
0215
0216     LOG(LEVEL)
0217         << " width " << std::setw(7) << width
0218         << " height " << std::setw(7) << height
0219         << " width*height " << std::setw(7) << width*height
0220         << " threadsPerBlock"
0221         << "("
0222         << std::setw(3) << threadsPerBlock.x << " "
0223         << std::setw(3) << threadsPerBlock.y << " "
0224         << std::setw(3) << threadsPerBlock.z << " "
0225         << ")"
0226         << " numBlocks "
0227         << "("
0228         << std::setw(3) << numBlocks.x << " "
0229         << std::setw(3) << numBlocks.y << " "
0230         << std::setw(3) << numBlocks.z << " "
0231         << ")"
0232         ;
0233 }
0234
0235 void QCerenkov::check()
0236 {
0237     unsigned width = tex->width ;
0238     unsigned height = tex->height ;
0239
0240     LOG(LEVEL)
0241         << " width " << width
0242         << " height " << height
0243         ;
0244
0245     dim3 numBlocks ;
0246     dim3 threadsPerBlock ;
0247     configureLaunch( numBlocks, threadsPerBlock, width, height );
0248     QCerenkov_check(numBlocks, threadsPerBlock, width, height );
0249
0250     cudaDeviceSynchronize();
0251 }
0252
0253
0254 void QCerenkov::dump( float* lookup, unsigned num_lookup, unsigned edgeitems  )
0255 {
0256     LOG(LEVEL);
0257     for(unsigned i=0 ; i < num_lookup ; i++)
0258     {
0259         if( i < edgeitems || i > num_lookup - edgeitems )
0260         std::cout
0261             << std::setw(6) << i
0262             << std::setw(10) << std::fixed << std::setprecision(3) << lookup[i]
0263             << std::endl
0264             ;
0265     }
0266 }
0267