 		a5 = coefs[5];
 		a6 = coefs[6];
 		a7 = coefs[7];
-//		printf("\npc_block (numactive == 8) %d %d\n", lim, num);
-		for ( j = lim; j < num; j++ )			// <---------------------- make parallel
-		{
-			LOOP_ALIGN
-
-			top = in[j - lim];
-			pin = in + j - 1;
-
-			b0 = top - (*pin--);
-			b1 = top - (*pin--);
-			b2 = top - (*pin--);
-			b3 = top - (*pin--);
-			b4 = top - (*pin--);
-			b5 = top - (*pin--);
-			b6 = top - (*pin--);
-			b7 = top - (*pin);
-			pin += 8;
-
-			sum1 = (denhalf - a0 * b0 - a1 * b1 - a2 * b2 - a3 * b3
-					- a4 * b4 - a5 * b5 - a6 * b6 - a7 * b7) >> denshift;
-
-			del = in[j] - top - sum1;
-			del = (del << chanshift) >> chanshift;
-			pc1[j] = del;	     
-			del0 = del;
-
-			sg = sign_of_int(del);
-			if ( sg > 0 )
-			{
-				sgn = sign_of_int( b7 );
-				a7 -= sgn;
-				del0 -= 1 * ((sgn * b7) >> denshift);
-				if ( del0 <= 0 )
-					continue;
-				
-				sgn = sign_of_int( b6 );
-				a6 -= sgn;
-				del0 -= 2 * ((sgn * b6) >> denshift);
-				if ( del0 <= 0 )
-					continue;
-				
-				sgn = sign_of_int( b5 );
-				a5 -= sgn;
-				del0 -= 3 * ((sgn * b5) >> denshift);
-				if ( del0 <= 0 )
-					continue;
-
-				sgn = sign_of_int( b4 );
-				a4 -= sgn;
-				del0 -= 4 * ((sgn * b4) >> denshift);
-				if ( del0 <= 0 )
-					continue;
-				
-				sgn = sign_of_int( b3 );
-				a3 -= sgn;
-				del0 -= 5 * ((sgn * b3) >> denshift);
-				if ( del0 <= 0 )
-					continue;
-				
-				sgn = sign_of_int( b2 );
-				a2 -= sgn;
-				del0 -= 6 * ((sgn * b2) >> denshift);
-				if ( del0 <= 0 )
-					continue;
-				
-				sgn = sign_of_int( b1 );
-				a1 -= sgn;
-				del0 -= 7 * ((sgn * b1) >> denshift);
-				if ( del0 <= 0 )
-					continue;
-
-				a0 -= sign_of_int( b0 );
-			}
-			else if ( sg < 0 )
-			{
-				// note: to avoid unnecessary negations, we flip the value of "sgn"
-				sgn = -sign_of_int( b7 );
-				a7 -= sgn;
-				del0 -= 1 * ((sgn * b7) >> denshift);
-				if ( del0 >= 0 )
-					continue;
-				
-				sgn = -sign_of_int( b6 );
-				a6 -= sgn;
-				del0 -= 2 * ((sgn * b6) >> denshift);
-				if ( del0 >= 0 )
-					continue;
-				
-				sgn = -sign_of_int( b5 );
-				a5 -= sgn;
-				del0 -= 3 * ((sgn * b5) >> denshift);
-				if ( del0 >= 0 )
-					continue;
-
-				sgn = -sign_of_int( b4 );
-				a4 -= sgn;
-				del0 -= 4 * ((sgn * b4) >> denshift);
-				if ( del0 >= 0 )
-					continue;
-				
-				sgn = -sign_of_int( b3 );
-				a3 -= sgn;
-				del0 -= 5 * ((sgn * b3) >> denshift);
-				if ( del0 >= 0 )
-					continue;
-				
-				sgn = -sign_of_int( b2 );
-				a2 -= sgn;
-				del0 -= 6 * ((sgn * b2) >> denshift);
-				if ( del0 >= 0 )
-					continue;
-				
-				sgn = -sign_of_int( b1 );
-				a1 -= sgn;
-				del0 -= 7 * ((sgn * b1) >> denshift);
-				if ( del0 >= 0 )
-					continue;
-
-				a0 += sign_of_int( b0 );
-			}
+
+		int32_t *d_pin, *d_in, *d_pc1;
+		int16_t	*d_a0, *d_a1, *d_a2, *d_a3, *d_a4, *d_a5, *d_a6, *d_a7;
+
+		cudaMalloc(&d_in, num * sizeof(int32_t));
+		cudaMalloc(&d_pc1, num * sizeof(int32_t));
+
+		cudaMalloc((void**)&d_a0, sizeof(int16_t));
+		cudaMalloc((void**)&d_a1, sizeof(int16_t));
+		cudaMalloc((void**)&d_a2, sizeof(int16_t));
+		cudaMalloc((void**)&d_a3, sizeof(int16_t));
+		cudaMalloc((void**)&d_a4, sizeof(int16_t));
+		cudaMalloc((void**)&d_a5, sizeof(int16_t));
+		cudaMalloc((void**)&d_a6, sizeof(int16_t));
+		cudaMalloc((void**)&d_a7, sizeof(int16_t));
+
+		cudaMemcpy(d_in, in, num * sizeof(int32_t), cudaMemcpyHostToDevice);
+		cudaMemcpy(d_pc1, pc1, num * sizeof(int32_t), cudaMemcpyHostToDevice);
+
+		cudaMemcpy(d_a0, &a0, sizeof(int16_t), cudaMemcpyHostToDevice);
+		cudaMemcpy(d_a1, &a1, sizeof(int16_t), cudaMemcpyHostToDevice);
+		cudaMemcpy(d_a2, &a2, sizeof(int16_t), cudaMemcpyHostToDevice);
+		cudaMemcpy(d_a3, &a3, sizeof(int16_t), cudaMemcpyHostToDevice);
+		cudaMemcpy(d_a4, &a4, sizeof(int16_t), cudaMemcpyHostToDevice);
+		cudaMemcpy(d_a5, &a5, sizeof(int16_t), cudaMemcpyHostToDevice);
+		cudaMemcpy(d_a6, &a6, sizeof(int16_t), cudaMemcpyHostToDevice);
+		cudaMemcpy(d_a7, &a7, sizeof(int16_t), cudaMemcpyHostToDevice);
+
+		int n = num;
+		for (int i = 0; i < (num + SIZE - 1) / SIZE; i++){
+			gpu_pc_block_2 << <1, SIZE >> >(d_pin, d_in, d_pc1, lim, num, sum1, denhalf, denshift, d_a0, d_a1, d_a2, d_a3, d_a4, d_a5, d_a6, d_a7, del, sg, del0, chanshift, sgn, i, n);
+			n -= 1024;
 		}
 
+		cudaMemcpy(&a0, d_a0, sizeof(int16_t), cudaMemcpyDeviceToHost);
+		cudaMemcpy(&a1, d_a1, sizeof(int16_t), cudaMemcpyDeviceToHost);
+		cudaMemcpy(&a2, d_a2, sizeof(int16_t), cudaMemcpyDeviceToHost);
+		cudaMemcpy(&a3, d_a3, sizeof(int16_t), cudaMemcpyDeviceToHost);
+		cudaMemcpy(&a4, d_a4, sizeof(int16_t), cudaMemcpyDeviceToHost);
+		cudaMemcpy(&a5, d_a5, sizeof(int16_t), cudaMemcpyDeviceToHost);
+		cudaMemcpy(&a6, d_a6, sizeof(int16_t), cudaMemcpyDeviceToHost);
+		cudaMemcpy(&a7, d_a7, sizeof(int16_t), cudaMemcpyDeviceToHost);
+
+		cudaMemcpy(in, d_in, num * sizeof(int32_t), cudaMemcpyDeviceToHost);
+		cudaMemcpy(pc1, d_pc1, num * sizeof(int32_t), cudaMemcpyDeviceToHost);
+
+		cudaFree(d_in);
+		cudaFree(d_pc1);
+
+		cudaFree(d_a0);
+		cudaFree(d_a1);
+		cudaFree(d_a2);
+		cudaFree(d_a3);
+		cudaFree(d_a4);
+		cudaFree(d_a5);
+		cudaFree(d_a6);
+		cudaFree(d_a7);
+		
 		coefs[0] = a0;
 		coefs[1] = a1;
 		coefs[2] = a2;

		
		int n = num;
		for (int i = 0; i < (num + SIZE - 1) / SIZE; i++){
			gpu_pc_block << <1, SIZE >> >(d_pin, d_in, d_pc1, lim, num, sum1, denhalf, denshift, d_a0, d_a1, d_a2, d_a3, del, sg, del0, chanshift, sgn, i, n);
			n -= 1024;
		}


	}
	else if ( numactive == 8 )
	{

		// optimization for numactive == 8
		register int16_t	a4, a5, a6, a7;
		register int32_t	b4, b5, b6, b7;
