#ifndef RTS_COMPLEXFIELD_H #define RTS_COMPLEXFIELD_H #include "cublas_v2.h" #include #include "../math/field.cuh" #include "../math/complex.h" #include "../math/realfield.cuh" namespace stim{ template __global__ void gpu_complexfield_mag(T* dest, complex* source, unsigned int r0, unsigned int r1){ int iu = blockIdx.x * blockDim.x + threadIdx.x; int iv = blockIdx.y * blockDim.y + threadIdx.y; //make sure that the thread indices are in-bounds if(iu >= r0 || iv >= r1) return; //compute the index into the field int i = iv*r0 + iu; //calculate and store the result dest[i] = source[i].abs(); } /*This class stores functions for saving images of complex fields */ template class complexfield : public field< stim::complex, D >{ using field< stim::complex, D >::R; using field< stim::complex, D >::X; using field< stim::complex, D >::shape; using field< stim::complex, D >::cuda_params; public: //find the maximum value of component n stim::complex find_max(unsigned int n){ cublasStatus_t stat; cublasHandle_t handle; //create a CUBLAS handle stat = cublasCreate(&handle); if(stat != CUBLAS_STATUS_SUCCESS){ std::cout<<"CUBLAS Error: initialization failed"< result; if(sizeof(T) == 8) stat = cublasIcamax(handle, L, (const cuComplex*)X[n], 1, &index); else stat = cublasIzamax(handle, L, (const cuDoubleComplex*)X[n], 1, &index); index -= 1; //adjust for 1-based indexing //if there was a GPU error, terminate if(stat != CUBLAS_STATUS_SUCCESS){ std::cout<<"CUBLAS Error: failure finding maximum value."<), cudaMemcpyDeviceToHost)); return result; } public: enum attribute {magnitude, real, imaginary}; //constructor (no parameters) complexfield() : field, D>(){}; //constructor (resolution specified) complexfield(unsigned int r0, unsigned int r1) : field, D>(r0, r1){}; //assignment from a field of complex values complexfield & operator=(const field< stim::complex, D > rhs){ field< complex, D >::operator=(rhs); return *this; } //assignment operator (scalar value) complexfield & operator= (const complex rhs){ field< complex, D >::operator=(rhs); return *this; } //assignment operator (vector value) complexfield & operator= (const vec< complex, D > rhs){ field< complex, D >::operator=(rhs); return *this; } //cropping complexfield crop(unsigned int width, unsigned int height){ complexfield result; result = field< complex, D>::crop(width, height); return result; } void toImage(std::string filename, attribute type = magnitude, unsigned int n=0){ field rf(R[0], R[1]); //get cuda parameters dim3 blocks, grids; cuda_params(grids, blocks); if(type == magnitude){ gpu_complexfield_mag <<>> (rf.ptr(), X[n], R[0], R[1]); rf.toImage(filename, n, true); } } }; } //end namespace rts #endif