ICODE-ADC/lammps
4
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

git-svn-id: svn://svn.icms.temple.edu/lammps-ro/trunk@6262 f3b2605a-c512-4ea7-a41b-209d697bcdaa

Browse Source
This commit is contained in:
sjplimp
2011-05-31 21:09:51 +00:00
parent 2ea09945bf
commit 7af10d8e0f
167 changed files with 19369 additions and 0 deletions
Download Patch File Download Diff File Expand all files Collapse all files

156
lib/cuda/cuda_data_kernel.cu Normal file
View File

@ -0,0 +1,156 @@
__global__ void CudaData_Upload_Kernel_DoubleFloat(double* buffer,float* dev_data,
unsigned nx,unsigned ny,unsigned nz,copy_mode mode)
{
if(mode==x) mode=xx;
unsigned length=nx;
if(ny>0) length*=ny;
if(nz>0) length*=nz;
unsigned i=(blockIdx.x*gridDim.y+blockIdx.y)*blockDim.x+threadIdx.x,j,k,l;
if(i>=length) return;
switch(mode)
{
case xx:
{
dev_data[i]=buffer[i];
}
case xy:
{
dev_data[i]=buffer[i];
}
case yx:
{
j=i/ny;
k=i%ny;
dev_data[k*nx+j]=buffer[j*ny+k];
}
case xyz:
{
dev_data[i]=buffer[i];
}
case xzy:
{
j=i/(ny*nz);
k=(i%(ny*nz))/nz;
l=i%nz;
dev_data[j*ny*nz+l*ny+k]=buffer[j*ny*nz+k*nz+l];
}
}
}
__global__ void CudaData_Upload_Kernel_DoubleDouble(double* buffer,double* dev_data,
unsigned nx,unsigned ny,unsigned nz,copy_mode mode)
{
if(mode==x) mode=xx;
unsigned length=nx;
if(ny>0) length*=ny;
if(nz>0) length*=nz;
unsigned i=(blockIdx.x*gridDim.y+blockIdx.y)*blockDim.x+threadIdx.x,j,k,l;
if(i>=length) return;
switch(mode)
{
case xx:
dev_data[i]=buffer[i];
case xy:
dev_data[i]=buffer[i];
case yx:
j=i/ny;
k=i%ny;
dev_data[k*nx+j]=buffer[j*ny+k];
case xyz:
dev_data[i]=buffer[i];
case xzy:
j=i/(ny*nz);
k=(i%(ny*nz))/nz;
l=i%nz;
dev_data[j*ny*nz+l*ny+k]=buffer[j*ny*nz+k*nz+l];
}
}
__global__ void CudaData_Upload_Kernel_FloatDouble(float* buffer,double* dev_data,
unsigned nx,unsigned ny,unsigned nz,copy_mode mode)
{
if(mode==x) mode=xx;
unsigned length=nx;
if(ny>0) length*=ny;
if(nz>0) length*=nz;
unsigned i=(blockIdx.x*gridDim.y+blockIdx.y)*blockDim.x+threadIdx.x,j,k,l;
if(i>=length) return;
switch(mode)
{
case xx:
dev_data[i]=buffer[i];
case xy:
dev_data[i]=buffer[i];
case yx:
j=i/ny;
k=i%ny;
dev_data[k*nx+j]=buffer[j*ny+k];
case xyz:
dev_data[i]=buffer[i];
case xzy:
j=i/(ny*nz);
k=(i%(ny*nz))/nz;
l=i%nz;
dev_data[j*ny*nz+l*ny+k]=buffer[j*ny*nz+k*nz+l];
}
}
__global__ void CudaData_Upload_Kernel_FloatFloat(float* buffer,float* dev_data,
unsigned nx,unsigned ny,unsigned nz,copy_mode mode)
{
if(mode==x) mode=xx;
unsigned length=nx;
if(ny>0) length*=ny;
if(nz>0) length*=nz;
unsigned i=(blockIdx.x*gridDim.y+blockIdx.y)*blockDim.x+threadIdx.x,j,k,l;
if(i>=length) return;
switch(mode)
{
case xx:
dev_data[i]=buffer[i];
case xy:
dev_data[i]=buffer[i];
case yx:
j=i/ny;
k=i%ny;
dev_data[k*nx+j]=buffer[j*ny+k];
case xyz:
dev_data[i]=buffer[i];
case xzy:
j=i/(ny*nz);
k=(i%(ny*nz))/nz;
l=i%nz;
dev_data[j*ny*nz+l*ny+k]=buffer[j*ny*nz+k*nz+l];
}
}
__global__ void CudaData_Upload_Kernel_IntInt(int* buffer,int* dev_data,
unsigned nx,unsigned ny,unsigned nz,copy_mode mode)
{
if(mode==x) mode=xx;
unsigned length=nx;
if(ny>0) length*=ny;
if(nz>0) length*=nz;
unsigned i=(blockIdx.x*gridDim.y+blockIdx.y)*blockDim.x+threadIdx.x,j,k,l;
if(i>=length) return;
switch(mode)
{
case xx:
dev_data[i]=buffer[i];
case xy:
dev_data[i]=buffer[i];
case yx:
j=i/ny;
k=i%ny;
dev_data[k*nx+j]=buffer[j*ny+k];
case xyz:
dev_data[i]=buffer[i];
case xzy:
j=i/(ny*nz);
k=(i%(ny*nz))/nz;
l=i%nz;
dev_data[j*ny*nz+l*ny+k]=buffer[j*ny*nz+k*nz+l];
}
}