 #define HANDLE_ERROR(err) (handleError(err, __FILE__, __LINE__))
 
 void checkParam();
-__global__ void initLine();
+__global__ void initLine(float*, float*, int);
 __global__ void updateAll(float*, float*, float*, int);
 void printResult();
 

	printf("%ld\n", clock());
	printf("Initializing points on the line...\n");

	initLine<<<numOfBlocks, threadsPerBlock>>>();

	HANDLE_ERROR(cudaMemcpy(currVal, devCurrVal, allocPoints * sizeof(float), cudaMemcpyDeviceToHost));
	printResult();

	printf("Using points = %d, steps = %d\n", totalPoints, totalSteps);
}

__global__ void initLine(float *__devPrevVal, float *__devCurrVal, int __totalPoints) {
	int i = blockIdx.x * blockDim.x + threadIdx.x;
	if (i < __totalPoints) {
		float x = i / (__totalPoints - 1);
		__devPrevVal[i] = __devCurrVal[i] = __sinf(2.0 * PI * x);
	}
}

__global__ void updateAll(float *__devPrevVal, float *__devCurrVal, float *__devNextVal, int __totalPoints) {
	int i = blockIdx.x * blockDim.x + threadIdx.x;
	if (i < __totalPoints) {
		if ((i == 0) || (i == __totalPoints - 1))
			__devNextVal[i] = 0.0;
		else
			__devNextVal[i] = 2.0 * __devCurrVal[i] - __devPrevVal[i] + 0.09 * (-2.0) * __devCurrVal[i];
		__syncthreads();
		__devPrevVal[i] = __devCurrVal[i];
		__syncthreads();
		__devCurrVal[i] = __devNextVal[i];
	}
}

