 	while(param_ptr->validity & ACTIVE){
 		if( param_ptr->validity & (FINISH + ABSFIN) ){  break; }
 
-		printf("FFT status EW\r");
+		cudaMemset( cuPowerSpec, 0, MAX_IF* (MAX_seg_len/2+1)* sizeof(float));	// Clear Power Spec
+
+		printf("FFT status EW\n");
 		sops.sem_num = (ushort)8; sops.sem_op = (short)-4; sops.sem_flg = (short)0;
         semop( param_ptr->sem_data_id, &sops, 1);
 
-		printf("FFT status ES\r");
-		cudaMemcpy( cuReal_data, segdata_ptr, SEGDATA_SIZE/2, cudaMemcpyHostToDevice);
-		printf("FFT status EF\r");
+		printf("FFT status ES\n");
+		cudaMemcpy( cuReal_data, segdata_ptr, MAX_IF* MAX_seg_sec/2* MAX_seg_len * sizeof(cufftReal), cudaMemcpyHostToDevice);
 		cufftExecR2C(cufft_plan, cuReal_data, cuSpec_data);
+		complexPowerVec<<<NFFT/2+1,  NSEG* MAX_IF/2>>>((float2 *)cuSpec_data, cuReal_data, (NFFT/2 + 1)* NSEG* MAX_IF/2);
+		printf("FFT status EF\n");
+		//---- Accumulate power spectra
+		for(seg_index=0; seg_index<MAX_seg_sec/2; seg_index++){
+			accumVector<<<NFFT/2+1, 1>>>( cuPowerSpec, &cuReal_data[seg_index* (NFFT/2+1)], NFFT/2+1); 
+		}
 
-		printf("FFT status OW\r");
+		printf("FFT status OW\n");
 		segdata_ptr += MAX_IF* MAX_seg_len* MAX_seg_sec / 2;
 		sops.sem_num = (ushort)9; sops.sem_op = (short)-4; sops.sem_flg = (short)0;
         semop( param_ptr->sem_data_id, &sops, 1);
-		printf("FFT status OS\r");
-		cudaMemcpy( cuReal_data, segdata_ptr, SEGDATA_SIZE/2, cudaMemcpyHostToDevice);
+		printf("FFT status OS\n");
+		cudaMemcpy( cuReal_data, segdata_ptr, MAX_IF* MAX_seg_sec/2* MAX_seg_len * sizeof(cufftReal), cudaMemcpyHostToDevice);
 		segdata_ptr -= MAX_IF* MAX_seg_len* MAX_seg_sec / 2;
-		printf("FFT status OF\r");
+		cufftExecR2C(cufft_plan, cuReal_data, cuSpec_data);
+		complexPowerVec<<<NFFT/2+1,  NSEG* MAX_IF/2>>>((float2 *)cuSpec_data, cuReal_data, (NFFT/2 + 1)* NSEG* MAX_IF/2);
+		printf("FFT status OF\n");
+		for(seg_index=0; seg_index<MAX_seg_sec/2; seg_index++){
+			accumVector<<<NFFT/2+1, 1>>>( cuPowerSpec, &cuReal_data[seg_index* (NFFT/2+1)], NFFT/2+1); 
+		}
 
-		cufftExecR2C(cufft_plan, cuReal_data, cuSpec_data);
+		cudaMemcpy( xspec_ptr, segdata_ptr, (NFFT/2+1)* sizeof(float), cudaMemcpyDeviceToHost);
 	}
 /*
 -------------------------------------------- RELEASE the SHM


__device__ float complexMod( float2 a )				// output |a|^2
{
	return  a.x* a.x + a.y* a.y 
}

__global__ void complexMultConjVec(		// calculate a x b*

{
	int tid = blockIdx.x* blockDim.x + threadIdx.x;
	if((tid >= 0) && (tid < length)){
		vec_out[tid] = complexMod(vec_in[tid]);
	}
}

__global__ void accumVector(	// a <- a + b
	float	*vec_in_a,		// Accumuration Results
	float	*vec_in_b,		// to be accumulated
	int		length)
{
    int tid = blockIdx.x* blockDim.x + threadIdx.x;
    if((tid >= 0) && (tid < length)){
        vec_in_a[tid] += vec_in_b[tid];
    }
}


