diff --git a/build/lib/libradar.so b/build/lib/libradar.so index 68d621af8f1b6759e039c6fb29e1b71d0fa47068..39c0358fd178ea28a893cb1d71a72ea8a34658fa 100755 Binary files a/build/lib/libradar.so and b/build/lib/libradar.so differ diff --git a/lib/lib.go b/lib/lib.go index b42dfa30c3ffccbea0b1d3fc8c5967c78181cd94..3df6fef6040401efbe1a1af7989f7b9e9bdd75e6 100644 --- a/lib/lib.go +++ b/lib/lib.go @@ -24,12 +24,14 @@ func CalculateReflectivity(inputData data.RawData) data.CalculatedData { zdb_c := make([]C.float, 512) zdr_c := make([]C.float, 512) - C.calculate_reflectivity( + C.calculate( + C.int(inputData.SectorIndex), C.int(inputData.ElevationIndex), &r_hh[0], &i_hh[0], &r_vv[0], &i_vv[0], &r_hv[0], &i_hv[0], - &zdb_c[0], &zdr_c[0], ) + C.fetch_ppi(C.int(inputData.SectorIndex), &zdb_c[0], &zdr_c[0]) + zdb := ConvertToGo(zdb_c) zdr := ConvertToGo(zdr_c) diff --git a/lib/radar.cu b/lib/radar.cu index 5d6b294694added860841a69b2222ab3a5847736..472bfbd8c376367c5003a74a7c04375c7d65c59b 100644 --- a/lib/radar.cu +++ b/lib/radar.cu @@ -6,6 +6,8 @@ #include <cufft.h> #include <sys/time.h> #include <assert.h> +#include <thrust/device_ptr.h> +#include <thrust/fill.h> using namespace std; @@ -13,6 +15,7 @@ using namespace std; #define k_calib 1941.05 #define RESULT_SIZE 2 +#define RESULT_SECTOR_COUNT 143 #define DEBUG @@ -356,7 +359,7 @@ extern "C" { iqhh = new cuFloatComplex[m*n]; iqvv = new cuFloatComplex[m*n]; iqhv = new cuFloatComplex[m*n]; - result = new float[(m/2)*RESULT_SIZE]; + result = new float[(m/2) * RESULT_SIZE * RESULT_SECTOR_COUNT]; float a, b; @@ -395,7 +398,11 @@ extern "C" { cudaMalloc(&d_iqvv, m*n*sizeof(cuFloatComplex)); cudaMalloc(&d_iqhv, m*n*sizeof(cuFloatComplex)); cudaMalloc(&d_sum, m*n*sizeof(cuFloatComplex)); - cudaMalloc(&d_result, (m/2)*RESULT_SIZE*sizeof(float)); + cudaMalloc(&d_result, (m/2) * RESULT_SIZE * RESULT_SECTOR_COUNT * sizeof(float)); + + // set initial value + thrust::device_ptr<float> devPtr(d_result); + thrust::fill(devPtr, devPtr + (m/2) * RESULT_SIZE * RESULT_SECTOR_COUNT, -INFINITY); cudaMemcpy(d_hamming, hamming_coef, m*n*sizeof(float), cudaMemcpyHostToDevice); // cudaMemcpy(d_ma, fft_ma, n*sizeof(cuFloatComplex), cudaMemcpyHostToDevice); @@ -430,16 +437,17 @@ extern "C" { delete[] iqhv; } - void calculate_reflectivity( - int* r_hh, - int* i_hh, - int* r_vv, - int* i_vv, - int* r_hv, - int* i_hv, - float* zdb, - float* zdr - ) { + void fetch_ppi(int sector_idx, float* zdb, float* zdr){ + const int offset = sector_idx * (m/2) * RESULT_SIZE; + + for (int i=0; i<m/2; i++) { + zdb[i] = result[offset + i*RESULT_SIZE]; + zdr[i] = result[offset + i*RESULT_SIZE + 1]; + } + } + + void calculate(int sector_idx, int elevation_idx, int* r_hh, int* i_hh, + int* r_vv, int* i_vv, int* r_hv, int* i_hv) { // tock(&tb, &te, "initialization"); float ms; // elapsed time in milliseconds @@ -453,7 +461,6 @@ extern "C" { cudaEventRecord(startEvent,0); - // tick(&tb); // Read 1 sector data @@ -545,14 +552,10 @@ extern "C" { __sum_inplace_v4<<<m/4,n,2*n*sizeof(cuFloatComplex)>>>(d_iqhv); // Calculate ZdB, Zdr - __calcresult_v2<<<1,m/2>>>(d_iqhh, d_iqvv, d_iqhv, d_result, n); - - cudaMemcpy(result, d_result, (m/2)*RESULT_SIZE*sizeof(float), cudaMemcpyDeviceToHost); + const int offset = sector_idx * (m/2) * RESULT_SIZE; + __calcresult_v2<<<1,m/2>>>(d_iqhh, d_iqvv, d_iqhv, &d_result[offset], n); - for (int i=0; i<m/2; i++) { - zdb[i] = result[i*RESULT_SIZE]; - zdr[i] = result[i*RESULT_SIZE + 1]; - } + cudaMemcpy(&result[offset], &d_result[offset], (m/2)*RESULT_SIZE*sizeof(float), cudaMemcpyDeviceToHost); cudaEventRecord(stopEvent, 0); cudaEventSynchronize(stopEvent); diff --git a/lib/radar.h b/lib/radar.h index c0909aabbc28a9b41cf5be39b5886c870efa5b30..8dc4208717266f8d046f7d0ebc723596e05f7599 100644 --- a/lib/radar.h +++ b/lib/radar.h @@ -7,8 +7,9 @@ extern "C" { #endif -void calculate_reflectivity(int* r_hh, int* i_hh, int* r_vv, int* i_vv, - int* r_hv, int* i_hv, float* zdb, float* zdr); +void calculate(int sector_idx, int elevation_idx, int* r_hh, int* i_hh, + int* r_vv, int* i_vv, int* r_hv, int* i_hv); +void fetch_ppi(int sector_idx, float* zdb, float* zdr); void initialize(); void cleanup();