Commit d38973a0 authored by sakthipriyas's avatar sakthipriyas
Browse files

udated kernel code for gridSize = 1 or blockSize = 1

parent a8d1c533
...@@ -45,6 +45,15 @@ public: ...@@ -45,6 +45,15 @@ public:
*/ */
void compute_sk(thrust::device_vector<thrust::complex<float>> const& data, RFIStatistics &stats); 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: private:
/** /**
* @brief initializes the data members of the class. * @brief initializes the data members of the class.
......
...@@ -49,9 +49,9 @@ __global__ void compute_sk_kernel(thrust::complex<float> *data, std::size_t samp ...@@ -49,9 +49,9 @@ __global__ void compute_sk_kernel(thrust::complex<float> *data, std::size_t samp
int iend = ibegin + window_size; int iend = ibegin + window_size;
for(std::size_t i = ibegin; i < iend; i++){ for(std::size_t i = ibegin; i < iend; i++){
float power = thrust::abs(data[i]) * thrust::abs(data[i]); float power = thrust::abs(data[i]) * thrust::abs(data[i]);
float power_sq = power * power; float power_sq = power * power;
s1 = s1 + power; s1 = s1 + power;
s2 = s2 + power_sq; s2 = s2 + power_sq;
} }
sk = ((window_size + 1) / (window_size - 1)) *(((window_size * s2) / (s1 * s1)) - 1); sk = ((window_size + 1) / (window_size - 1)) *(((window_size * s2) / (s1 * s1)) - 1);
rfi_status[id] = (int) ((sk < sk_min) || (sk > sk_max)); rfi_status[id] = (int) ((sk < sk_min) || (sk > sk_max));
...@@ -81,7 +81,6 @@ void SpectralKurtosisCuda::init() ...@@ -81,7 +81,6 @@ void SpectralKurtosisCuda::init()
_nwindows = _sample_size /_window_size; _nwindows = _sample_size /_window_size;
_d_s1.resize(_nwindows); _d_s1.resize(_nwindows);
_d_s2.resize(_nwindows); _d_s2.resize(_nwindows);
_rfi_status.resize(_nwindows);
} }
void SpectralKurtosisCuda::compute_sk(const thrust::device_vector<thrust::complex<float>> &data, RFIStatistics &stats){ 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 ...@@ -118,15 +117,16 @@ void SpectralKurtosisCuda::compute_sk_k(thrust::device_vector<thrust::complex<fl
_sample_size = data.size(); _sample_size = data.size();
BOOST_LOG_TRIVIAL(debug) << "Computing SK_k for sample_size " << _sample_size BOOST_LOG_TRIVIAL(debug) << "Computing SK_k for sample_size " << _sample_size
<< " and window_size " << _window_size <<".\n"; << " 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()); thrust::complex<float> *k_data = thrust::raw_pointer_cast(data.data());
int *k_rfi_status = thrust::raw_pointer_cast(_rfi_status.data()); int *k_rfi_status = thrust::raw_pointer_cast(stats.rfi_status.data());
//compute_sk_kernel<<<1, 1>>> (k_data, _sample_size, _window_size, _sk_max, _sk_min, k_rfi_status);
compute_sk_kernel<<<1, _nwindows>>> (k_data, _sample_size, _window_size, _sk_max, _sk_min, k_rfi_status); 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; //compute_sk_kernel<<<_nwindows, 1>>> (k_data, _sample_size, _window_size, _sk_max, _sk_min, k_rfi_status); //works.
BOOST_LOG_TRIVIAL(info) << "RFI fraction: " << rfi_fraction; 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(); nvtxRangePop();
} }
......
...@@ -81,7 +81,6 @@ TEST_F(SpectralKurtosisCudaTester, sk_withRFI) ...@@ -81,7 +81,6 @@ TEST_F(SpectralKurtosisCudaTester, sk_withRFI)
TEST_F(SpectralKurtosisCudaTester, sk_RFIreplacement) 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 = 1024 * 1024;
//std::size_t sample_size = 160000000; //std::size_t sample_size = 160000000;
std::size_t window_size = 1024 * 2; std::size_t window_size = 1024 * 2;
...@@ -115,7 +114,6 @@ TEST_F(SpectralKurtosisCudaTester, sk_RFIreplacement) ...@@ -115,7 +114,6 @@ TEST_F(SpectralKurtosisCudaTester, sk_RFIreplacement)
TEST_F(SpectralKurtosisCudaTester, sk_kernel) 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 = 1024 * 1024;
//std::size_t sample_size = 160000000; //std::size_t sample_size = 160000000;
std::size_t window_size = 1024 * 2; std::size_t window_size = 1024 * 2;
......
Supports Markdown
0% or .
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment