start preparations for QCerenkov GPU 2d(BetaInverse,u) float4 icdf te…

…xture for Cerenkov energy lookup

qudarap/CMakeLists.txt

@@ -34,6 +34,7 @@ set(SOURCES
 
     QTex.cc
     QTex.cu
+    QTexMaker.cc
 
     QBnd.cc
     QBnd.cu
@@ -77,6 +78,7 @@ SET(HEADERS
     QRng.hh
 
     QTex.hh
+    QTexMaker.hh
     QBnd.hh
     QScint.hh
     QCerenkov.hh

qudarap/QScint.cc

@@ -41,6 +41,14 @@ std::string QScint::desc() const
     return s ; 
 }
 
+/**
+QScint::MakeScintTex
+-----------------------
+
+TODO: move the hd_factor into payload instead of items for easier extension to 2d 
+
+**/
+
 QTex<float>* QScint::MakeScintTex(const NP* src, unsigned hd_factor )  // static 
 {
     assert( src->has_shape(1,4096,1) ||  src->has_shape(3,4096,1) ); 

qudarap/QTex.cc

@@ -12,15 +12,16 @@
 #include "QU.hh"
 #include "QTex.hh"
 
+
 template<typename T>
 QTex<T>::QTex(size_t width_, size_t height_ , const void* src_, char filterMode_ )
     :   
     width(width_),
     height(height_),
     src(src_),
     filterMode(filterMode_),
-    dst(new T[width*height]),
-    d_dst(nullptr),
+    rotate_dst(nullptr),
+    d_rotate_dst(nullptr),
     cuArray(nullptr),
     channelDesc(cudaCreateChannelDesc<T>()),
     texObj(0),
@@ -53,14 +54,14 @@ QTex<T>::~QTex()
 {
     cudaDestroyTextureObject(texObj);
     cudaFreeArray(cuArray);
-    delete[] dst ; 
-    cudaFree(d_dst);
+    delete[] rotate_dst ; 
+    cudaFree(d_rotate_dst);
 }
 
 template<typename T>
 void QTex<T>::init()
 {
-    createArray();
+    createArray();   // cudaMallocArray using channelDesc for T 
     uploadToArray();
     createTextureObject();
 
@@ -236,18 +237,23 @@ potentially with some spare threads at edge when workspace is not an exact multi
 template<typename T>
 void QTex<T>::rotate(float theta)
 {
-    cudaMalloc(&d_dst, width*height*sizeof(T));
+    cudaMalloc(&d_rotate_dst, width*height*sizeof(T));
 
     dim3 threadsPerBlock(16, 16);
     dim3 numBlocks((width + threadsPerBlock.x - 1) / threadsPerBlock.x, (height + threadsPerBlock.y - 1) / threadsPerBlock.y);
 
-    QTex_uchar4_rotate_kernel( numBlocks, threadsPerBlock, d_dst, texObj, width, height, theta );
+    QTex_uchar4_rotate_kernel( numBlocks, threadsPerBlock, d_rotate_dst, texObj, width, height, theta );
 
     cudaDeviceSynchronize();
     cudaCheckErrors("cudaDeviceSynchronize");
     // Fatal error: cudaDeviceSynchronize (linear filtering not supported for non-float type at SIMGStandaloneTest.cu:123)
 
-    cudaMemcpy(dst, d_dst, width*height*sizeof(T), cudaMemcpyDeviceToHost);
+    if(rotate_dst == nullptr)
+    {
+        rotate_dst = new T[width*height] ;  
+    }
+
+    cudaMemcpy(rotate_dst, d_rotate_dst, width*height*sizeof(T), cudaMemcpyDeviceToHost);
 }
 
 /**

qudarap/QTex.hh

@@ -22,8 +22,10 @@ struct QUDARAP_API QTex
     const void*  src ;
     char         filterMode ;  // 'L':cudaFilterModeLinear OR 'P':cudaFilterModePoint 
 
-    T*           dst ; 
-    T*           d_dst ; 
+    // TODO: split off rotation elsewhere "QTexRotate?", it is not typically needed 
+    T*           rotate_dst ; 
+    T*           d_rotate_dst ; 
+
 
     cudaArray*   cuArray ; 
     cudaChannelFormatDesc channelDesc ;

qudarap/QTexMaker.cc

@@ -0,0 +1,27 @@
+#include "QTexMaker.hh"
+
+#include "scuda.h"
+#include "NP.hh"
+#include "QTex.hh"
+
+QTex<float4>* QTexMaker::Make2d_f4( const NP* a, char filterMode )  // static 
+{
+    assert( a->ebyte == 4 && "need to narrow double precision arrays first ");  
+    unsigned ndim = a->shape.size(); 
+    assert( ndim == 3 );      
+
+    unsigned ni = a->shape[0] ; 
+    unsigned nj = a->shape[1] ; 
+    unsigned nk = a->shape[2] ; assert( nk == 4 ); 
+
+    size_t height = ni ; 
+    size_t  width = nj ; 
+    const void* src = (const void*)a->bytes(); 
+
+    // note that from the point of view of array content, saying (height, width) 
+    // is a more appropriate ordering than the usual contrary convention  
+
+    QTex<float4>* tex = new QTex<float4>( width, height, src, filterMode  );
+    return tex ; 
+}
+

qudarap/QTexMaker.hh

@@ -0,0 +1,14 @@
+#pragma once
+
+struct NP ; 
+struct float4 ; 
+template <typename T> struct QTex ;  
+
+#include "QUDARAP_API_EXPORT.hh"
+
+struct QUDARAP_API QTexMaker
+{
+    static QTex<float4>* Make2d_f4( const NP* a, char filterMode ); 
+};
+
+

qudarap/tests/CMakeLists.txt

@@ -19,6 +19,7 @@ set(TEST_SOURCES
    QEventTest.cc 
    QSimWithEventTest.cc 
    QUTest.cc
+   QTexMakerTest.cc
 
 )
 

qudarap/tests/QScintTest.cc

@@ -38,7 +38,7 @@ int main(int argc, char** argv)
     NP* icdf = slib->getBuf(); 
 #else
     const char* cfbase = SPath::Resolve(SSys::getenvvar("CFBASE", "$TMP/CSG_GGeo" ));
-    NP* icdf = NP::Load(cfbase, "CSGFoundry", "icdf.npy"); 
+    NP* icdf = NP::Load(cfbase, "CSGFoundry", "icdf.npy"); // HMM: this needs a more destinctive name/location  
     //icdf->dump(); 
 #endif
 

qudarap/tests/QTexMakerTest.cc

@@ -0,0 +1,45 @@
+#include "NP.hh"
+#include "QTexMaker.hh"
+#include "OPTICKS_LOG.hh"
+
+
+NP* make_array(unsigned ni, unsigned nj, unsigned nk)
+{
+    std::vector<float> src ; 
+    for(unsigned i=0 ; i < ni ; i++)
+    for(unsigned j=0 ; j < nj ; j++)
+    for(unsigned k=0 ; k < nk ; k++)
+    {
+        float val = float(i*100 + j*10 + k) ;  
+        src.push_back(val); 
+    }
+
+    NP* a = NP::Make<float>( ni, nj, nk ); 
+    a->read(src.data()) ; 
+    a->set_meta<unsigned>("hd_factor", 10); 
+
+    return a ; 
+}
+
+int main(int argc, char** argv)
+{
+    OPTICKS_LOG(argc, argv); 
+
+    const NP* a = make_array( 5, 10, 4 ); 
+    char filterMode = 'P' ; 
+    unsigned hd_factor = a->get_meta<unsigned>("hd_factor"); 
+    assert( hd_factor > 0 ); 
+
+    QTex<float4>* tex = QTexMaker::Make2d_f4(a, filterMode ); 
+    assert(tex); 
+
+    tex->setHDFactor(hd_factor); 
+    tex->uploadMeta(); 
+
+    // TODO: lookup checks on the GPU texture, like QScint 
+
+
+    return 0 ; 
+}
+
+

qudarap/tests/QTexRotateTest.cc

@@ -33,7 +33,7 @@ int main(int argc, char** argv)
 
     std::cout << "writing to " << opath << std::endl ;
 
-    SIMG img2(img.width, img.height, img.channels, (unsigned char*)qtex.dst );
+    SIMG img2(img.width, img.height, img.channels, (unsigned char*)qtex.rotate_dst );
     img2.writePNG(opath);
 
     return 0;