eic-opticks/qudarap/QTexRotate.cc

0001 #include <cuda_runtime.h>
0002 #include "cudaCheckErrors.h"
0003
0004 #include "QTex.hh"
0005 #include "QTexRotate.hh"
0006
0007 template<typename T>
0008 QTexRotate<T>::QTexRotate(const QTex<T>* tex_)
0009     :
0010     tex(tex_),
0011     rotate_dst(nullptr),
0012     d_rotate_dst(nullptr)
0013 {
0014 }
0015
0016 template<typename T>
0017 QTexRotate<T>::~QTexRotate()
0018 {
0019     delete[] rotate_dst ;
0020     cudaFree(d_rotate_dst);
0021 }
0022
0023 extern "C" void QTex_uchar4_rotate_kernel(dim3 dimGrid, dim3 dimBlock, uchar4* d_output, cudaTextureObject_t texObj,  size_t width, size_t height, float theta );
0024
0025 template<typename T>
0026 void QTexRotate<T>::rotate(float theta)
0027 {
0028     unsigned width = tex->width ;
0029     unsigned height = tex->height ;
0030
0031
0032     cudaMalloc(&d_rotate_dst, width*height*sizeof(T));
0033
0034     dim3 threadsPerBlock(16, 16);
0035     dim3 numBlocks((width + threadsPerBlock.x - 1) / threadsPerBlock.x, (height + threadsPerBlock.y - 1) / threadsPerBlock.y);
0036
0037     QTex_uchar4_rotate_kernel( numBlocks, threadsPerBlock, d_rotate_dst, tex->texObj, width, height, theta );
0038
0039     cudaDeviceSynchronize();
0040     cudaCheckErrors("cudaDeviceSynchronize");
0041     // Fatal error: cudaDeviceSynchronize (linear filtering not supported for non-float type at SIMGStandaloneTest.cu:123)
0042
0043     if(rotate_dst == nullptr)
0044     {
0045         rotate_dst = new T[width*height] ;
0046     }
0047
0048     cudaMemcpy(rotate_dst, d_rotate_dst, width*height*sizeof(T), cudaMemcpyDeviceToHost);
0049 }
0050
0051
0052
0053 template struct QUDARAP_API QTexRotate<uchar4>;
0054
0055