diff --git a/psrdada_cpp/effelsberg/edd/SpectralKurtosisCuda.cuh b/psrdada_cpp/effelsberg/edd/SpectralKurtosisCuda.cuh index ab4bdcb8c8ceec3df3b3a349fa6a30d79b0d4897..653d233b0741b2365f4bcdbd12197cb0a4f4a695 100644 --- a/psrdada_cpp/effelsberg/edd/SpectralKurtosisCuda.cuh +++ b/psrdada_cpp/effelsberg/edd/SpectralKurtosisCuda.cuh @@ -45,6 +45,15 @@ public: */ void compute_sk(thrust::device_vector<thrust::complex<float>> const& data, RFIStatistics &stats); + /** + * @brief computes spectral kurtosis (using optimized kernel fucntion) for the given data and returns its rfi statistics. + * + * @param(in) data input data + * @param(out) stats RFI statistics + * + */ + void compute_sk_k(thrust::device_vector<thrust::complex<float>> &data, RFIStatistics &stats); + private: /** * @brief initializes the data members of the class. diff --git a/psrdada_cpp/effelsberg/edd/src/SpectralKurtosisCuda.cu b/psrdada_cpp/effelsberg/edd/src/SpectralKurtosisCuda.cu index cb6bc52847074358ce8c7062401e92d93dba4383..7becb782c93ebfdad2efb0a5ac160e5908ec3e58 100755 --- a/psrdada_cpp/effelsberg/edd/src/SpectralKurtosisCuda.cu +++ b/psrdada_cpp/effelsberg/edd/src/SpectralKurtosisCuda.cu @@ -49,9 +49,9 @@ __global__ void compute_sk_kernel(thrust::complex<float> *data, std::size_t samp int iend = ibegin + window_size; for(std::size_t i = ibegin; i < iend; i++){ float power = thrust::abs(data[i]) * thrust::abs(data[i]); - float power_sq = power * power; - s1 = s1 + power; - s2 = s2 + power_sq; + float power_sq = power * power; + s1 = s1 + power; + s2 = s2 + power_sq; } sk = ((window_size + 1) / (window_size - 1)) *(((window_size * s2) / (s1 * s1)) - 1); rfi_status[id] = (int) ((sk < sk_min) || (sk > sk_max)); @@ -81,7 +81,6 @@ void SpectralKurtosisCuda::init() _nwindows = _sample_size /_window_size; _d_s1.resize(_nwindows); _d_s2.resize(_nwindows); - _rfi_status.resize(_nwindows); } void SpectralKurtosisCuda::compute_sk(const thrust::device_vector<thrust::complex<float>> &data, RFIStatistics &stats){ @@ -118,15 +117,16 @@ void SpectralKurtosisCuda::compute_sk_k(thrust::device_vector<thrust::complex<fl _sample_size = data.size(); BOOST_LOG_TRIVIAL(debug) << "Computing SK_k for sample_size " << _sample_size << " and window_size " << _window_size <<".\n"; - //initializing class variables - init(); + _nwindows = _sample_size / _window_size; + stats.rfi_status.resize(_nwindows); thrust::complex<float> *k_data = thrust::raw_pointer_cast(data.data()); - int *k_rfi_status = thrust::raw_pointer_cast(_rfi_status.data()); - //compute_sk_kernel<<<1, 1>>> (k_data, _sample_size, _window_size, _sk_max, _sk_min, k_rfi_status); + int *k_rfi_status = thrust::raw_pointer_cast(stats.rfi_status.data()); + compute_sk_kernel<<<1, _nwindows>>> (k_data, _sample_size, _window_size, _sk_max, _sk_min, k_rfi_status); - float rfi_fraction = thrust::reduce(_rfi_status.begin(), _rfi_status.end(), 0.0f) / _nwindows; - BOOST_LOG_TRIVIAL(info) << "RFI fraction: " << rfi_fraction; + //compute_sk_kernel<<<_nwindows, 1>>> (k_data, _sample_size, _window_size, _sk_max, _sk_min, k_rfi_status); //works. + stats.rfi_fraction = thrust::reduce(stats.rfi_status.begin(), stats.rfi_status.end(), 0.0f) / _nwindows; + BOOST_LOG_TRIVIAL(info) << "RFI fraction: " << stats.rfi_fraction; nvtxRangePop(); } diff --git a/psrdada_cpp/effelsberg/edd/test/src/SpectralKurtosisCudaTester.cu b/psrdada_cpp/effelsberg/edd/test/src/SpectralKurtosisCudaTester.cu index 5076a7bd587895d372a7682318fee233e33d5241..3c36b78038a12dc75d9e78df0aee5662092b3315 100644 --- a/psrdada_cpp/effelsberg/edd/test/src/SpectralKurtosisCudaTester.cu +++ b/psrdada_cpp/effelsberg/edd/test/src/SpectralKurtosisCudaTester.cu @@ -81,7 +81,6 @@ TEST_F(SpectralKurtosisCudaTester, sk_withRFI) TEST_F(SpectralKurtosisCudaTester, sk_RFIreplacement) { - //std::size_t sample_size = (1 * 1024 * 1024 * 1024) / 8; std::size_t sample_size = 1024 * 1024; //std::size_t sample_size = 160000000; std::size_t window_size = 1024 * 2; @@ -115,7 +114,6 @@ TEST_F(SpectralKurtosisCudaTester, sk_RFIreplacement) TEST_F(SpectralKurtosisCudaTester, sk_kernel) { - //std::size_t sample_size = (1 * 1024 * 1024 * 1024) / 8; std::size_t sample_size = 1024 * 1024; //std::size_t sample_size = 160000000; std::size_t window_size = 1024 * 2;