/*
comm2.cu (CUDA+MPI)
*/

#include "oth.h"
#include "oth_mpi.h"
#include "oth_cuda.h"

#ifdef _MPI
#include "mpi.h"
#endif

extern void cuda_memcpy(int, void *, const void *, size_t, cudaMemcpyKind);

// share boundary
void comm_boundary(cuDoubleComplex *v)
{
#ifdef _MPI
	MPI_Status status;
	const int tag = 0;

	// -X direction
	if (commRank > 0) {
		// copy to send buffer
		cuda_memcpy(GPU, Sendcomplex_m, &v[Send0_m], Sendcount_m * sizeof(cuDoubleComplex), cudaMemcpyDeviceToHost);
		for (int64_t n = 0; n < Sendcount_m; n++) {
			Sendbuf_m[(2 * n) + 0] = Sendcomplex_m[n].x;
			Sendbuf_m[(2 * n) + 1] = Sendcomplex_m[n].y;
		}
		// MPI
		int dst = commRank - 1;
		MPI_Sendrecv(Sendbuf_m, 2 * Sendcount_m, MPI_DOUBLE, dst, tag,
		             Recvbuf_m, 2 * Recvcount_m, MPI_DOUBLE, dst, tag, MPI_COMM_WORLD, &status);
		// copy from recv buffer
		for (int64_t n = 0; n < Recvcount_m; n++) {
			 Recvcomplex_m[n].x = Recvbuf_m[(2 * n) + 0];
			 Recvcomplex_m[n].y = Recvbuf_m[(2 * n) + 1];
		}
		cuda_memcpy(GPU, &v[Recv0_m], Recvcomplex_m, Recvcount_m * sizeof(cuDoubleComplex), cudaMemcpyHostToDevice);
	}

	// +X direction
	if (commRank < commSize - 1) {
		// copy to send buffer
		cuda_memcpy(GPU, Sendcomplex_p, &v[Send0_p], Sendcount_p * sizeof(cuDoubleComplex), cudaMemcpyDeviceToHost);
		for (int64_t n = 0; n < Sendcount_p; n++) {
			Sendbuf_p[(2 * n) + 0] = Sendcomplex_p[n].x;
			Sendbuf_p[(2 * n) + 1] = Sendcomplex_p[n].y;
		}
		// MPI
		int dst = commRank + 1;
		MPI_Sendrecv(Sendbuf_p, 2 * Sendcount_p, MPI_DOUBLE, dst, tag,
		             Recvbuf_p, 2 * Recvcount_p, MPI_DOUBLE, dst, tag, MPI_COMM_WORLD, &status);
		// copy from recv buffer
		for (int64_t n = 0; n < Recvcount_p; n++) {
			 Recvcomplex_p[n].x = Recvbuf_p[(2 * n) + 0];
			 Recvcomplex_p[n].y = Recvbuf_p[(2 * n) + 1];
		}
		cuda_memcpy(GPU, &v[Recv0_p], Recvcomplex_p, Recvcount_p * sizeof(cuDoubleComplex), cudaMemcpyHostToDevice);
	}
#endif
}

// X-PBC
void comm_Xpbc(cuDoubleComplex *v)
{
#ifdef _MPI
	MPI_Status status;
	const int tag = 0;

	// -X boundary
	if (commRank == 0) {
		// copy to send buffer
		cuda_memcpy(GPU, Sendcomplex_Xpbc_m, &v[Send0_Xpbc_m], Sendcount_Xpbc_m * sizeof(cuDoubleComplex), cudaMemcpyDeviceToHost);
		for (int64_t n = 0; n < Sendcount_Xpbc_m; n++) {
			Sendbuf_Xpbc_m[(2 * n) + 0] = Sendcomplex_Xpbc_m[n].x;
			Sendbuf_Xpbc_m[(2 * n) + 1] = Sendcomplex_Xpbc_m[n].y;
		}
		// MPI
		int dst = commSize - 1;
		MPI_Sendrecv(Sendbuf_Xpbc_m, 2 * Sendcount_Xpbc_m, MPI_DOUBLE, dst, tag,
		             Recvbuf_Xpbc_m, 2 * Recvcount_Xpbc_m, MPI_DOUBLE, dst, tag, MPI_COMM_WORLD, &status);
		// copy from recv buffer
		for (int64_t n = 0; n < Recvcount_Xpbc_m; n++) {
			Recvcomplex_Xpbc_m[n].x = Recvbuf_Xpbc_m[(2 * n) + 0];
			Recvcomplex_Xpbc_m[n].y = Recvbuf_Xpbc_m[(2 * n) + 1];
		}
		cuda_memcpy(GPU, &v[Recv0_Xpbc_m], Recvcomplex_Xpbc_m, Recvcount_Xpbc_m * sizeof(cuDoubleComplex), cudaMemcpyHostToDevice);
	}

	// +X boundary
	if (commRank == commSize - 1) {
		// copy to send buffer
		cuda_memcpy(GPU, Sendcomplex_Xpbc_p, &v[Send0_Xpbc_p], Sendcount_Xpbc_p * sizeof(cuDoubleComplex), cudaMemcpyDeviceToHost);
		for (int64_t n = 0; n < Sendcount_Xpbc_p; n++) {
			Sendbuf_Xpbc_p[(2 * n) + 0] = Sendcomplex_Xpbc_p[n].x;
			Sendbuf_Xpbc_p[(2 * n) + 1] = Sendcomplex_Xpbc_p[n].y;
		}
		// MPI
		int dst = 0;
		MPI_Sendrecv(Sendbuf_Xpbc_p, 2 * Sendcount_Xpbc_p, MPI_DOUBLE, dst, tag,
		             Recvbuf_Xpbc_p, 2 * Recvcount_Xpbc_p, MPI_DOUBLE, dst, tag, MPI_COMM_WORLD, &status);
		// copy from recv buffer
		for (int64_t n = 0; n < Recvcount_Xpbc_p; n++) {
			Recvcomplex_Xpbc_p[n].x = Recvbuf_Xpbc_p[(2 * n) + 0];
			Recvcomplex_Xpbc_p[n].y = Recvbuf_Xpbc_p[(2 * n) + 1];
		}
		cuda_memcpy(GPU, &v[Recv0_Xpbc_p], Recvcomplex_Xpbc_p, Recvcount_Xpbc_p * sizeof(cuDoubleComplex), cudaMemcpyHostToDevice);
	}
#endif
}
