CudaPmeSolverUtilKernel.h File Reference

#include "HipDefines.h"

Go to the source code of this file.

Classes
struct	TransposeBatch< T >

Functions
void	spread_charge (const float4 *atoms, const int numAtoms, const int nfftx, const int nffty, const int nfftz, const int xsize, const int ysize, const int zsize, const int xdim, const int y00, const int z00, const bool periodicY, const bool periodicZ, float *data, const int order, cudaStream_t stream)
void	scalar_sum (const bool orderXYZ, const int nfft1, const int nfft2, const int nfft3, const int size1, const int size2, const int size3, const double kappa, const float recip1x, const float recip1y, const float recip1z, const float recip2x, const float recip2y, const float recip2z, const float recip3x, const float recip3y, const float recip3z, const double volume, const float *prefac1, const float *prefac2, const float *prefac3, const int k2_00, const int k3_00, const bool doEnergyVirial, double *energy, double *virial, float2 *data, cudaStream_t stream)
void	gather_force (const float4 *atoms, const int numAtoms, const int nfftx, const int nffty, const int nfftz, const int xsize, const int ysize, const int zsize, const int xdim, const int y00, const int z00, const bool periodicY, const bool periodicZ, const float *data, const int order, float3 *force, const cudaTextureObject_t gridTexObj, cudaStream_t stream)
void	transpose_xyz_yzx (const int nx, const int ny, const int nz, const int xsize_in, const int ysize_in, const int ysize_out, const int zsize_out, const float2 *data_in, float2 *data_out, cudaStream_t stream)
void	batchTranspose_xyz_yzx (const int numBatches, TransposeBatch< float2 > *batches, const int max_nx, const int ny, const int nz, const int xsize_in, const int ysize_in, cudaStream_t stream)
void	transpose_xyz_zxy (const int nx, const int ny, const int nz, const int xsize_in, const int ysize_in, const int zsize_out, const int xsize_out, const float2 *data_in, float2 *data_out, cudaStream_t stream)
void	batchTranspose_xyz_zxy (const int numBatches, TransposeBatch< float2 > *batches, const int max_nx, const int ny, const int nz, const int xsize_in, const int ysize_in, cudaStream_t stream)

Function Documentation

void batchTranspose_xyz_yzx	(	const int	numBatches,
		TransposeBatch< float2 > *	batches,
		const int	max_nx,
		const int	ny,
		const int	nz,
		const int	xsize_in,
		const int	ysize_in,
		cudaStream_t	stream
	)

Definition at line 1390 of file CudaPmeSolverUtilKernel.cu.

References cudaCheck, TILEDIM, and TILEROWS.

Referenced by CudaPmeTranspose::transposeXYZtoYZX().

                                                               {

  dim3 numthread(TILEDIM, TILEROWS, 1);
  dim3 numblock((max_nx-1)/TILEDIM+1, (ny-1)/TILEDIM+1, nz*numBatches);

  batchTranspose_xyz_yzx_kernel<float2> <<< numblock, numthread, 0, stream >>> (batches, ny, nz, xsize_in, ysize_in);

  cudaCheck(cudaGetLastError());
}

void batchTranspose_xyz_zxy	(	const int	numBatches,
		TransposeBatch< float2 > *	batches,
		const int	max_nx,
		const int	ny,
		const int	nz,
		const int	xsize_in,
		const int	ysize_in,
		cudaStream_t	stream
	)

Definition at line 1425 of file CudaPmeSolverUtilKernel.cu.

References cudaCheck, TILEDIM, and TILEROWS.

Referenced by CudaPmeTranspose::transposeXYZtoZXY().

                       {

  dim3 numthread(TILEDIM, TILEROWS, 1);
  dim3 numblock((max_nx-1)/TILEDIM+1, (nz-1)/TILEDIM+1, ny*numBatches);

  batchTranspose_xyz_zxy_kernel<float2> <<< numblock, numthread, 0, stream >>> (batches, ny, nz, xsize_in, ysize_in);

  cudaCheck(cudaGetLastError());
}

void gather_force	(	const float4 *	atoms,
		const int	numAtoms,
		const int	nfftx,
		const int	nffty,
		const int	nfftz,
		const int	xsize,
		const int	ysize,
		const int	zsize,
		const int	xdim,
		const int	y00,
		const int	z00,
		const bool	periodicY,
		const bool	periodicZ,
		const float *	data,
		const int	order,
		float3 *	force,
		const cudaTextureObject_t	gridTexObj,
		cudaStream_t	stream
	)

Definition at line 1225 of file CudaPmeSolverUtilKernel.cu.

References atoms, cudaCheck, and cudaNAMD_bug().

                       {

  dim3 nthread(32, 2, 1);
  dim3 nblock((numAtoms - 1)/nthread.x + 1, 1, 1);
  // dim3 nblock(npatch, 1, 1);

  switch(order) {
    case 4:
    if (periodicY && periodicZ)
      gather_force<float, float3, 4, true, true> <<< nblock, nthread, 0, stream >>> (
        atoms, numAtoms, nfftx, nffty, nfftz, xsize, ysize, zsize, xdim, y00, z00,
        // recip11, recip22, recip33,
        data,
#ifdef NAMD_CUDA
        gridTexObj,
#endif
        1, force);
    else if (periodicY)
      gather_force<float, float3, 4, true, false> <<< nblock, nthread, 0, stream >>> (
        atoms, numAtoms, nfftx, nffty, nfftz, xsize, ysize, zsize, xdim, y00, z00,
        // recip11, recip22, recip33,
        data,
#ifdef NAMD_CUDA
        gridTexObj,
#endif
        1, force);
    else if (periodicZ)
      gather_force<float, float3, 4, false, true> <<< nblock, nthread, 0, stream >>> (
        atoms, numAtoms, nfftx, nffty, nfftz, xsize, ysize, zsize, xdim, y00, z00,
        // recip11, recip22, recip33,
        data,
#ifdef NAMD_CUDA
        gridTexObj,
#endif
        1, force);
    else
      gather_force<float, float3, 4, false, false> <<< nblock, nthread, 0, stream >>> (
        atoms, numAtoms, nfftx, nffty, nfftz, xsize, ysize, zsize, xdim, y00, z00,
        // recip11, recip22, recip33,
        data,
#ifdef NAMD_CUDA
        gridTexObj,
#endif
        1, force);
    break;

    case 6:
    if (periodicY && periodicZ)
      gather_force<float, float3, 6, true, true> <<< nblock, nthread, 0, stream >>> (
        atoms, numAtoms, nfftx, nffty, nfftz, xsize, ysize, zsize, xdim, y00, z00,
        // recip11, recip22, recip33,
        data,
#ifdef NAMD_CUDA
        gridTexObj,
#endif
        1, force);
    else if (periodicY)
      gather_force<float, float3, 6, true, false> <<< nblock, nthread, 0, stream >>> (
        atoms, numAtoms, nfftx, nffty, nfftz, xsize, ysize, zsize, xdim, y00, z00,
        // recip11, recip22, recip33,
        data,
#ifdef NAMD_CUDA
        gridTexObj,
#endif
        1, force);
    else if (periodicZ)
      gather_force<float, float3, 6, false, true> <<< nblock, nthread, 0, stream >>> (
        atoms, numAtoms, nfftx, nffty, nfftz, xsize, ysize, zsize, xdim, y00, z00,
        // recip11, recip22, recip33,
        data,
#ifdef NAMD_CUDA
        gridTexObj,
#endif
        1, force);
    else
      gather_force<float, float3, 6, false, false> <<< nblock, nthread, 0, stream >>> (
        atoms, numAtoms, nfftx, nffty, nfftz, xsize, ysize, zsize, xdim, y00, z00,
        // recip11, recip22, recip33,
        data,
#ifdef NAMD_CUDA
        gridTexObj,
#endif
        1, force);
    break;
 
    case 8:
    if (periodicY && periodicZ)
      gather_force<float, float3, 8, true, true> <<< nblock, nthread, 0, stream >>> (
        atoms, numAtoms, nfftx, nffty, nfftz, xsize, ysize, zsize, xdim, y00, z00,
        // recip11, recip22, recip33,
        data,
#ifdef NAMD_CUDA
        gridTexObj,
#endif
        1, force);
    else if (periodicY)
      gather_force<float, float3, 8, true, false> <<< nblock, nthread, 0, stream >>> (
        atoms, numAtoms, nfftx, nffty, nfftz, xsize, ysize, zsize, xdim, y00, z00,
        // recip11, recip22, recip33,
        data,
#ifdef NAMD_CUDA
        gridTexObj,
#endif
        1, force);
    else if (periodicZ)
      gather_force<float, float3, 8, false, true> <<< nblock, nthread, 0, stream >>> (
        atoms, numAtoms, nfftx, nffty, nfftz, xsize, ysize, zsize, xdim, y00, z00,
        // recip11, recip22, recip33,
        data,
#ifdef NAMD_CUDA
        gridTexObj,
#endif
        1, force);
    else
      gather_force<float, float3, 8, false, false> <<< nblock, nthread, 0, stream >>> (
        atoms, numAtoms, nfftx, nffty, nfftz, xsize, ysize, zsize, xdim, y00, z00,
        // recip11, recip22, recip33,
        data,
#ifdef NAMD_CUDA
        gridTexObj,
#endif
        1, force);
    break;

    default:
    char str[128];
    sprintf(str, "gather_force, order %d not implemented",order);
    cudaNAMD_bug(str);
  }
  cudaCheck(cudaGetLastError());

}

void scalar_sum	(	const bool	orderXYZ,
		const int	nfft1,
		const int	nfft2,
		const int	nfft3,
		const int	size1,
		const int	size2,
		const int	size3,
		const double	kappa,
		const float	recip1x,
		const float	recip1y,
		const float	recip1z,
		const float	recip2x,
		const float	recip2y,
		const float	recip2z,
		const float	recip3x,
		const float	recip3y,
		const float	recip3z,
		const double	volume,
		const float *	prefac1,
		const float *	prefac2,
		const float *	prefac3,
		const int	k2_00,
		const int	k3_00,
		const bool	doEnergyVirial,
		double *	energy,
		double *	virial,
		float2 *	data,
		cudaStream_t	stream
	)

Definition at line 1158 of file CudaPmeSolverUtilKernel.cu.

References cudaCheck, M_PI, and WARPSIZE.

Referenced by CudaPmeKSpaceCompute::solve().

                       {
#ifdef NAMD_CUDA
  int nthread = 1024;
  int nblock = 64;
#else
  int nthread = 256;
  int nblock = 256;
#endif

  int shmem_size = sizeof(float)*(nfft1 + nfft2 + nfft3);
  if (doEnergyVirial) {    
    shmem_size = max(shmem_size, (int)((nthread/WARPSIZE)*sizeof(RecipVirial_t)));
  }

  float piv_inv = (float)(1.0/(M_PI*volume));
  float fac = (float)(M_PI*M_PI/(kappa*kappa));

  if (doEnergyVirial) {
    if (orderXYZ) {
      scalar_sum_kernel<float, float2, true, true, false> <<< nblock, nthread, shmem_size, stream >>> (
        nfft1, nfft2, nfft3, size1, size2, size3,
        recip1x, recip1y, recip1z,
        recip2x, recip2y, recip2z,
        recip3x, recip3y, recip3z,
        prefac1, prefac2, prefac3,
        fac, piv_inv, k2_00, k3_00, data, energy, virial);
    } else {
      scalar_sum_kernel<float, float2, true, false, false> <<< nblock, nthread, shmem_size, stream >>> (
        nfft1, nfft2, nfft3, size1, size2, size3,
        recip1x, recip1y, recip1z,
        recip2x, recip2y, recip2z,
        recip3x, recip3y, recip3z,
        prefac1, prefac2, prefac3,
        fac, piv_inv, k2_00, k3_00, data, energy, virial);
    }
  } else {
    if (orderXYZ) {
      scalar_sum_kernel<float, float2, false, true, false> <<< nblock, nthread, shmem_size, stream >>> (
        nfft1, nfft2, nfft3, size1, size2, size3,
        recip1x, recip1y, recip1z,
        recip2x, recip2y, recip2z,
        recip3x, recip3y, recip3z,
        prefac1, prefac2, prefac3,
        fac, piv_inv, k2_00, k3_00, data, NULL, NULL);
    } else {
      scalar_sum_kernel<float, float2, false, false, false> <<< nblock, nthread, shmem_size, stream >>> (
        nfft1, nfft2, nfft3, size1, size2, size3,
        recip1x, recip1y, recip1z,
        recip2x, recip2y, recip2z,
        recip3x, recip3y, recip3z,
        prefac1, prefac2, prefac3,
        fac, piv_inv, k2_00, k3_00, data, NULL, NULL);
    }
  }
  cudaCheck(cudaGetLastError());

}

void spread_charge	(	const float4 *	atoms,
		const int	numAtoms,
		const int	nfftx,
		const int	nffty,
		const int	nfftz,
		const int	xsize,
		const int	ysize,
		const int	zsize,
		const int	xdim,
		const int	y00,
		const int	z00,
		const bool	periodicY,
		const bool	periodicZ,
		float *	data,
		const int	order,
		cudaStream_t	stream
	)

Definition at line 1072 of file CudaPmeSolverUtilKernel.cu.

References atoms, cudaCheck, cudaNAMD_bug(), if(), and WARPSIZE.

Referenced by CudaPmeRealSpaceCompute::spreadCharge().

                                                     {

  dim3 nthread, nblock;
  const int order3 = ((order*order*order-1)/WARPSIZE + 1)*WARPSIZE;
  switch(order) {
  case 4:
    nthread.x = WARPSIZE;
#if defined(NAMD_CUDA)
    nthread.y = 4;
#else
    nthread.y = order3 / WARPSIZE;
#endif
    nthread.z = 1;
    nblock.x = (numAtoms - 1)/nthread.x + 1;
    nblock.y = 1;
    nblock.z = 1;
    if (periodicY && periodicZ)
      spread_charge_kernel<float, 4, true, true> <<< nblock, nthread, 0, stream >>> (atoms, numAtoms,
         nfftx, nffty, nfftz, xsize, ysize, zsize, xdim, y00, z00, data);
    else if (periodicY)
      spread_charge_kernel<float, 4, true, false> <<< nblock, nthread, 0, stream >>> (atoms, numAtoms,
         nfftx, nffty, nfftz, xsize, ysize, zsize, xdim, y00, z00, data);
    else if (periodicZ)
      spread_charge_kernel<float, 4, false, true> <<< nblock, nthread, 0, stream >>> (atoms, numAtoms,
         nfftx, nffty, nfftz, xsize, ysize, zsize, xdim, y00, z00, data);
    else
      spread_charge_kernel<float, 4, false, false> <<< nblock, nthread, 0, stream >>> (atoms, numAtoms,
         nfftx, nffty, nfftz, xsize, ysize, zsize, xdim, y00, z00, data);
    break;

  case 6:
    nthread.x = WARPSIZE;
    nthread.y = order3/WARPSIZE;
    nthread.z = 1;
    nblock.x = (numAtoms - 1)/nthread.x + 1;
    nblock.y = 1;
    nblock.z = 1;
    if (periodicY && periodicZ)
      spread_charge_kernel<float, 6, true, true> <<< nblock, nthread, 0, stream >>> (atoms, numAtoms,
         nfftx, nffty, nfftz, xsize, ysize, zsize, xdim, y00, z00, data);
    else if (periodicY)
      spread_charge_kernel<float, 6, true, false> <<< nblock, nthread, 0, stream >>> (atoms, numAtoms,
         nfftx, nffty, nfftz, xsize, ysize, zsize, xdim, y00, z00, data);
    else if (periodicZ)
      spread_charge_kernel<float, 6, false, true> <<< nblock, nthread, 0, stream >>> (atoms, numAtoms,
         nfftx, nffty, nfftz, xsize, ysize, zsize, xdim, y00, z00, data);
    else
      spread_charge_kernel<float, 6, false, false> <<< nblock, nthread, 0, stream >>> (atoms, numAtoms,
         nfftx, nffty, nfftz, xsize, ysize, zsize, xdim, y00, z00, data);
    break;

  case 8:
    nthread.x = WARPSIZE;
    nthread.y = order3/WARPSIZE;
    nthread.z = 1;
    nblock.x = (numAtoms - 1)/nthread.x + 1;
    nblock.y = 1;
    nblock.z = 1;
    if (periodicY && periodicZ)
      spread_charge_kernel<float, 8, true, true> <<< nblock, nthread, 0, stream >>> (atoms, numAtoms,
         nfftx, nffty, nfftz, xsize, ysize, zsize, xdim, y00, z00, data);
    else if (periodicY)
      spread_charge_kernel<float, 8, true, false> <<< nblock, nthread, 0, stream >>> (atoms, numAtoms,
         nfftx, nffty, nfftz, xsize, ysize, zsize, xdim, y00, z00, data);
    else if (periodicZ)
      spread_charge_kernel<float, 8, false, true> <<< nblock, nthread, 0, stream >>> (atoms, numAtoms,
         nfftx, nffty, nfftz, xsize, ysize, zsize, xdim, y00, z00, data);
    else
      spread_charge_kernel<float, 8, false, false> <<< nblock, nthread, 0, stream >>> (atoms, numAtoms,
         nfftx, nffty, nfftz, xsize, ysize, zsize, xdim, y00, z00, data);
    break;

  default:
    char str[128];
    sprintf(str, "spread_charge, order %d not implemented",order);
    cudaNAMD_bug(str);
  }
  cudaCheck(cudaGetLastError());

}

void transpose_xyz_yzx	(	const int	nx,
		const int	ny,
		const int	nz,
		const int	xsize_in,
		const int	ysize_in,
		const int	ysize_out,
		const int	zsize_out,
		const float2 *	data_in,
		float2 *	data_out,
		cudaStream_t	stream
	)

Definition at line 1371 of file CudaPmeSolverUtilKernel.cu.

References cudaCheck, TILEDIM, and TILEROWS.

                                                                {

  dim3 numthread(TILEDIM, TILEROWS, 1);
  dim3 numblock((nx-1)/TILEDIM+1, (ny-1)/TILEDIM+1, nz);
  transpose_xyz_yzx_kernel<float2> <<< numblock, numthread, 0, stream >>> (
    nx, ny, nz, xsize_in, ysize_in,
    ysize_out, zsize_out,
    data_in, data_out);

  cudaCheck(cudaGetLastError());
}

void transpose_xyz_zxy	(	const int	nx,
		const int	ny,
		const int	nz,
		const int	xsize_in,
		const int	ysize_in,
		const int	zsize_out,
		const int	xsize_out,
		const float2 *	data_in,
		float2 *	data_out,
		cudaStream_t	stream
	)

Definition at line 1406 of file CudaPmeSolverUtilKernel.cu.

References cudaCheck, TILEDIM, and TILEROWS.

                                                                {

  dim3 numthread(TILEDIM, TILEROWS, 1);
  dim3 numblock((nx-1)/TILEDIM+1, (nz-1)/TILEDIM+1, ny);
  transpose_xyz_zxy_kernel<float2> <<< numblock, numthread, 0, stream >>> (
    nx, ny, nz, xsize_in, ysize_in,
    zsize_out, xsize_out,
    data_in, data_out);

  cudaCheck(cudaGetLastError());
}