OilfieldToolsNet/LBPM
4
0
Fork 0
Code Issues Pull Requests Packages Projects Releases Wiki Activity

Refactored cpu branch

Browse Source
This commit is contained in:
James E McClure
2014-03-19 09:22:50 -04:00
parent d9e42aa064
commit 05bbbd3e5d
8 changed files with 69 additions and 69 deletions
Download Patch File Download Diff File Expand all files Collapse all files

46
cpu/Color.cpp
View File

@@ -1,6 +1,6 @@
#include <math.h>
extern "C" void dvc_InitDenColor(char *ID, double *Den, double *Phi, double das, double dbs, int Nx, int Ny, int Nz)
extern "C" void InitDenColor(char *ID, double *Den, double *Phi, double das, double dbs, int Nx, int Ny, int Nz)
{
int i,j,k,n,N;
@@ -34,7 +34,7 @@ extern "C" void dvc_InitDenColor(char *ID, double *Den, double *Phi, double das,
}
}
}
extern "C" void dvc_InitDenColorDistancePacked(char *ID, double *Den, double *Phi, double *Distance,
extern "C" void InitDenColorDistancePacked(char *ID, double *Den, double *Phi, double *Distance,
double das, double dbs, double beta, double xp, int Nx, int Ny, int Nz)
{
int i,j,k,n,N;
@@ -78,7 +78,7 @@ extern "C" void dvc_InitDenColorDistancePacked(char *ID, double *Den, double *Ph
}
}
extern "C" void dvc_InitDenColorDistance(char *ID, double *Den, double *Phi, double *Distance,
extern "C" void InitDenColorDistance(char *ID, double *Den, double *Phi, double *Distance,
double das, double dbs, double beta, double xp, int Nx, int Ny, int Nz)
{
int i,j,k,n,N;
@@ -114,7 +114,7 @@ extern "C" void dvc_InitDenColorDistance(char *ID, double *Den, double *Phi, dou
}
extern "C" void dvc_Compute_VELOCITY(char *ID, double *disteven, double *distodd, double *vel, int Nx, int Ny, int Nz, int S)
extern "C" void Compute_VELOCITY(char *ID, double *disteven, double *distodd, double *vel, int Nx, int Ny, int Nz, int S)
{
int n,N;
// distributions
@@ -162,7 +162,7 @@ extern "C" void dvc_Compute_VELOCITY(char *ID, double *disteven, double *distodd
}
}
extern "C" void dvc_ComputePressureD3Q19(char *ID, double *disteven, double *distodd, double *Pressure,
extern "C" void ComputePressureD3Q19(char *ID, double *disteven, double *distodd, double *Pressure,
int Nx, int Ny, int Nz)
{
int n,N;
@@ -207,7 +207,7 @@ extern "C" void dvc_ComputePressureD3Q19(char *ID, double *disteven, double *dis
}
//*************************************************************************
extern "C" void dvc_ColorBC_inlet(double *Phi, double *Den, double *A_even, double *A_odd,
extern "C" void ColorBC_inlet(double *Phi, double *Den, double *A_even, double *A_odd,
double *B_even, double *B_odd, int Nx, int Ny, int Nz)
{
int i,j,k,n,N;
@@ -250,7 +250,7 @@ extern "C" void dvc_ColorBC_inlet(double *Phi, double *Den, double *A_even, doub
}
}
//*************************************************************************
extern "C" void dvc_ColorBC_outlet(double *Phi, double *Den, double *A_even, double *A_odd,
extern "C" void ColorBC_outlet(double *Phi, double *Den, double *A_even, double *A_odd,
double *B_even, double *B_odd, int Nx, int Ny, int Nz)
{
int i,j,k,n,N;
@@ -294,7 +294,7 @@ extern "C" void dvc_ColorBC_outlet(double *Phi, double *Den, double *A_even, dou
}
}
//*************************************************************************
extern "C" void dvc_PressureBC_inlet(double *disteven, double *distodd, double din,
extern "C" void PressureBC_inlet(double *disteven, double *distodd, double din,
int Nx, int Ny, int Nz)
{
int n,N;
@@ -363,7 +363,7 @@ extern "C" void dvc_PressureBC_inlet(double *disteven, double *distodd, double d
}
}
extern "C" void dvc_PressureBC_outlet(double *disteven, double *distodd, double dout,
extern "C" void PressureBC_outlet(double *disteven, double *distodd, double dout,
int Nx, int Ny, int Nz, int outlet)
{
int n,N;
@@ -430,7 +430,7 @@ extern "C" void dvc_PressureBC_outlet(double *disteven, double *distodd, double
}
}
//*************************************************************************
extern "C" void dvc_ComputeColorGradient(char *ID, double *phi, double *ColorGrad, int Nx, int Ny, int Nz)
extern "C" void ComputeColorGradient(char *ID, double *phi, double *ColorGrad, int Nx, int Ny, int Nz)
{
int n,N,i,j,k,nn;
// distributions
@@ -557,7 +557,7 @@ extern "C" void dvc_ComputeColorGradient(char *ID, double *phi, double *ColorGra
}
}
//*************************************************************************
extern "C" void dvc_ColorCollide( char *ID, double *disteven, double *distodd, double *ColorGrad,
extern "C" void ColorCollide( char *ID, double *disteven, double *distodd, double *ColorGrad,
double *Velocity, int Nx, int Ny, int Nz, double rlx_setA, double rlx_setB,
double alpha, double beta, double Fx, double Fy, double Fz, bool pBC)
{
@@ -772,7 +772,7 @@ extern "C" void dvc_ColorCollide( char *ID, double *disteven, double *distodd, d
} // loop over n
}
extern "C" void dvc_ColorCollideOpt( char *ID, double *disteven, double *distodd, double *phi, double *ColorGrad,
extern "C" void ColorCollideOpt( char *ID, double *disteven, double *distodd, double *phi, double *ColorGrad,
double *Velocity, int Nx, int Ny, int Nz, double rlx_setA, double rlx_setB,
double alpha, double beta, double Fx, double Fy, double Fz)
{
@@ -1173,7 +1173,7 @@ extern "C" void dvc_ColorCollideOpt( char *ID, double *disteven, double *distodd
} // loop over n
}
extern "C" void dvc_MassColorCollideD3Q7(char *ID, double *A_even, double *A_odd, double *B_even, double *B_odd,
extern "C" void MassColorCollideD3Q7(char *ID, double *A_even, double *A_odd, double *B_even, double *B_odd,
double *Den, double *Phi, double *ColorGrad, double *Velocity, double beta, int N, bool pBC)
{
char id;
@@ -1279,7 +1279,7 @@ extern "C" void dvc_MassColorCollideD3Q7(char *ID, double *A_even, double *A_odd
}
//*************************************************************************
extern "C" void dvc_DensityStreamD3Q7(char *ID, double *Den, double *Copy, double *Phi, double *ColorGrad, double *Velocity,
extern "C" void DensityStreamD3Q7(char *ID, double *Den, double *Copy, double *Phi, double *ColorGrad, double *Velocity,
double beta, int Nx, int Ny, int Nz, bool pBC, int S)
{
char id;
@@ -1431,7 +1431,7 @@ extern "C" void dvc_DensityStreamD3Q7(char *ID, double *Den, double *Copy, doubl
}
}
/*
extern "C" void dvc_ComputePhi(char *ID, double *Phi, double *Copy, double *Den, int N, int S)
extern "C" void ComputePhi(char *ID, double *Phi, double *Copy, double *Den, int N, int S)
{
int n;
double Na,Nb;
@@ -1455,7 +1455,7 @@ extern "C" void dvc_ComputePhi(char *ID, double *Phi, double *Copy, double *Den,
//...................................................................
}
*/
extern "C" void dvc_ComputePhi(char *ID, double *Phi, double *Den, int N)
extern "C" void ComputePhi(char *ID, double *Phi, double *Den, int N)
{
int n;
double Na,Nb;
@@ -1475,19 +1475,19 @@ extern "C" void dvc_ComputePhi(char *ID, double *Phi, double *Den, int N)
/*
//*************************************************************************
extern "C" void dvc_InitDenColor( int nblocks, int nthreads, int S,
extern "C" void InitDenColor( int nblocks, int nthreads, int S,
char *ID, double *Den, double *Phi, double das, double dbs, int Nx, int Ny, int Nz)
{
InitDenColor <<<nblocks, nthreads>>> (ID, Den, Phi, das, dbs, Nx, Ny, Nz, S);
}
//*************************************************************************
extern "C" void dvc_ComputeColorGradient(int nBlocks, int nthreads, int S,
extern "C" void ComputeColorGradient(int nBlocks, int nthreads, int S,
char *ID, double *Phi, double *ColorGrad, int Nx, int Ny, int Nz)
{
ComputeColorGradient<<<nBlocks,nthreads>>>(ID, Phi, ColorGrad, Nx, Ny, Nz, S);
}
//*************************************************************************
extern "C" void dvc_ColorCollide(int nBlocks, int nthreads, int S,
extern "C" void ColorCollide(int nBlocks, int nthreads, int S,
char *ID, double *f_even, double *f_odd, double *ColorGrad, double *Velocity,
double rlxA, double rlxB,double alpha, double beta, double Fx, double Fy, double Fz,
int Nx, int Ny, int Nz, bool pBC)
@@ -1496,7 +1496,7 @@ extern "C" void dvc_ColorCollide(int nBlocks, int nthreads, int S,
rlxA, rlxB, alpha, beta, Fx, Fy, Fz, pBC);
}
//*************************************************************************
extern "C" void dvc_ColorCollideOpt(int nBlocks, int nthreads, int S,
extern "C" void ColorCollideOpt(int nBlocks, int nthreads, int S,
char *ID, double *f_even, double *f_odd, double *Phi, double *ColorGrad,
double *Velocity, int Nx, int Ny, int Nz,double rlxA, double rlxB,
double alpha, double beta, double Fx, double Fy, double Fz)
@@ -1509,20 +1509,20 @@ extern "C" void dvc_ColorCollideOpt(int nBlocks, int nthreads, int S,
}
//*************************************************************************
extern "C" void dvc_DensityStreamD3Q7(int nBlocks, int nthreads, int S,
extern "C" void DensityStreamD3Q7(int nBlocks, int nthreads, int S,
char *ID, double *Den, double *Copy, double *Phi, double *ColorGrad, double *Velocity,
double beta, int Nx, int Ny, int Nz, bool pBC)
{
DensityStreamD3Q7<<<nBlocks, nthreads>>>(ID,Den,Copy,Phi,ColorGrad,Velocity,beta,Nx,Ny,Nz,pBC,S);
}
//*************************************************************************
extern "C" void dvc_ComputePhi(int nBlocks, int nthreads, int S,
extern "C" void ComputePhi(int nBlocks, int nthreads, int S,
char *ID, double *Phi, double *Copy, double *Den, int N)
{
ComputePhi<<<nBlocks, nthreads>>>(ID,Phi,Copy,Den,N,S);
}
//*************************************************************************
extern "C" void dvc_ComputePressure(int nBlocks, int nthreads, int S,
extern "C" void ComputePressure(int nBlocks, int nthreads, int S,
char *ID, double *disteven, double *distodd,
double *Pressure, int Nx, int Ny, int Nz)
{

30
cpu/Color.h
View File

@@ -1,28 +1,28 @@
extern "C" void dvc_InitDenColor(char *ID, double *Den, double *Phi, double das, double dbs, int Nx, int Ny, int Nz);
extern "C" void dvc_InitDenColorDistance(char *ID, double *Den, double *Phi, double *Distance,
extern "C" void InitDenColor(char *ID, double *Den, double *Phi, double das, double dbs, int Nx, int Ny, int Nz);
extern "C" void InitDenColorDistance(char *ID, double *Den, double *Phi, double *Distance,
double das, double dbs, double beta, double xp, int Nx, int Ny, int Nz);
extern "C" void dvc_Compute_VELOCITY(char *ID, double *disteven, double *distodd, double *vel, int Nx, int Ny, int Nz);
extern "C" void dvc_ComputePressureD3Q19(char *ID, double *disteven, double *distodd, double *Pressure,
extern "C" void Compute_VELOCITY(char *ID, double *disteven, double *distodd, double *vel, int Nx, int Ny, int Nz);
extern "C" void ComputePressureD3Q19(char *ID, double *disteven, double *distodd, double *Pressure,
int Nx, int Ny, int Nz);
extern "C" void dvc_PressureBC_inlet(double *disteven, double *distodd, double din,
extern "C" void PressureBC_inlet(double *disteven, double *distodd, double din,
int Nx, int Ny, int Nz);
extern "C" void dvc_PressureBC_outlet(double *disteven, double *distodd, double dout,
extern "C" void PressureBC_outlet(double *disteven, double *distodd, double dout,
int Nx, int Ny, int Nz, int outlet);
extern "C" void dvc_ComputeColorGradient(char *ID, double *phi, double *ColorGrad, int Nx, int Ny, int Nz);
extern "C" void dvc_ColorCollide( char *ID, double *disteven, double *distodd, double *ColorGrad,
extern "C" void ComputeColorGradient(char *ID, double *phi, double *ColorGrad, int Nx, int Ny, int Nz);
extern "C" void ColorCollide( char *ID, double *disteven, double *distodd, double *ColorGrad,
double *Velocity, int Nx, int Ny, int Nz,double rlx_setA, double rlx_setB,
double alpha, double beta, double Fx, double Fy, double Fz, bool pBC);
extern "C" void dvc_ColorCollideOpt( char *ID, double *disteven, double *distodd, double *phi, double *ColorGrad,
extern "C" void ColorCollideOpt( char *ID, double *disteven, double *distodd, double *phi, double *ColorGrad,
double *Velocity, int Nx, int Ny, int Nz,double rlx_setA, double rlx_setB,
double alpha, double beta, double Fx, double Fy, double Fz);
extern "C" void dvc_DensityStreamD3Q7(char *ID, double *Den, double *Copy, double *Phi, double *ColorGrad, double *Velocity,
extern "C" void DensityStreamD3Q7(char *ID, double *Den, double *Copy, double *Phi, double *ColorGrad, double *Velocity,
double beta, int Nx, int Ny, int Nz, bool pBC);
//extern "C" void dvc_ComputePhi(char *ID, double *Phi, double *Copy, double *Den, int N);
//extern "C" void ComputePhi(char *ID, double *Phi, double *Copy, double *Den, int N);
// Added by J.E. McClure 2-12-2014
extern "C" void dvc_ComputePhi(char *ID, double *Phi, double *Den, int N);
extern "C" void dvc_MassColorCollideD3Q7(char *ID, double *A_even, double *A_odd, double *B_even, double *B_odd,
extern "C" void ComputePhi(char *ID, double *Phi, double *Den, int N);
extern "C" void MassColorCollideD3Q7(char *ID, double *A_even, double *A_odd, double *B_even, double *B_odd,
double *Den, double *Phi, double *ColorGrad, double *Velocity, double beta, int N, bool pBC);
extern "C" void dvc_ColorBC_inlet(double *Phi, double *Den, double *A_even, double *A_odd,
extern "C" void ColorBC_inlet(double *Phi, double *Den, double *A_even, double *A_odd,
double *B_even, double *B_odd, int Nx, int Ny, int Nz);
extern "C" void dvc_ColorBC_outlet(double *Phi, double *Den, double *A_even, double *A_odd,
extern "C" void ColorBC_outlet(double *Phi, double *Den, double *A_even, double *A_odd,
double *B_even, double *B_odd, int Nx, int Ny, int Nz);

8
cpu/D3Q19.cpp
View File

@@ -1,4 +1,4 @@
extern "C" void dvc_PackDist(int q, int *list, int start, int count, double *sendbuf, double *dist, int N){
extern "C" void PackDist(int q, int *list, int start, int count, double *sendbuf, double *dist, int N){
//....................................................................................
// Pack distribution q into the send buffer for the listed lattice sites
// dist may be even or odd distributions stored by stream layout
@@ -10,7 +10,7 @@ extern "C" void dvc_PackDist(int q, int *list, int start, int count, double *sen
}
}
extern "C" void dvc_UnpackDist(int q, int Cqx, int Cqy, int Cqz, int *list, int start, int count,
extern "C" void UnpackDist(int q, int Cqx, int Cqy, int Cqz, int *list, int start, int count,
double *recvbuf, double *dist, int Nx, int Ny, int Nz){
//....................................................................................
// Unack distribution from the recv buffer
@@ -52,7 +52,7 @@ extern "C" void dvc_UnpackDist(int q, int Cqx, int Cqy, int Cqz, int *list, int
}
}
extern "C" void dvc_InitD3Q19(char *ID, double *f_even, double *f_odd, int Nx, int Ny, int Nz)
extern "C" void InitD3Q19(char *ID, double *f_even, double *f_odd, int Nx, int Ny, int Nz)
{
int n,N;
N = Nx*Ny*Nz;
@@ -91,7 +91,7 @@ extern "C" void dvc_InitD3Q19(char *ID, double *f_even, double *f_odd, int Nx, i
}
//*************************************************************************
extern "C" void dvc_SwapD3Q19(char *ID, double *disteven, double *distodd, int Nx, int Ny, int Nz)
extern "C" void SwapD3Q19(char *ID, double *disteven, double *distodd, int Nx, int Ny, int Nz)
{
int i,j,k,n,nn,N;
// distributions

8
cpu/D3Q19.h
View File

@@ -1,6 +1,6 @@
extern "C" void dvc_PackDist(int q, int *list, int start, int count, double *sendbuf, double *dist, int N);
extern "C" void dvc_UnpackDist(int q, int Cqx, int Cqy, int Cqz, int *list, int start, int count,
extern "C" void PackDist(int q, int *list, int start, int count, double *sendbuf, double *dist, int N);
extern "C" void UnpackDist(int q, int Cqx, int Cqy, int Cqz, int *list, int start, int count,
double *recvbuf, double *dist, int Nx, int Ny, int Nz);
//*************************************************************************
extern "C" void dvc_InitD3Q19(char *ID, double *f_even, double *f_odd, int Nx, int Ny, int Nz);
extern "C" void dvc_SwapD3Q19(char *ID, double *disteven, double *distodd, int Nx, int Ny, int Nz);
extern "C" void InitD3Q19(char *ID, double *f_even, double *f_odd, int Nx, int Ny, int Nz);
extern "C" void SwapD3Q19(char *ID, double *disteven, double *distodd, int Nx, int Ny, int Nz);

14
cpu/D3Q7.cpp
View File

@@ -1,6 +1,6 @@
// CPU Functions for D3Q7 Lattice Boltzmann Methods
extern "C" void dvc_PackValues(int *list, int count, double *sendbuf, double *Data, int N){
extern "C" void PackValues(int *list, int count, double *sendbuf, double *Data, int N){
//....................................................................................
// Pack distribution q into the send buffer for the listed lattice sites
// dist may be even or odd distributions stored by stream layout
@@ -11,7 +11,7 @@ extern "C" void dvc_PackValues(int *list, int count, double *sendbuf, double *Da
sendbuf[idx] = Data[n];
}
}
extern "C" void dvc_UnpackValues(int *list, int count, double *recvbuf, double *Data, int N){
extern "C" void UnpackValues(int *list, int count, double *recvbuf, double *Data, int N){
//....................................................................................
// Pack distribution q into the send buffer for the listed lattice sites
// dist may be even or odd distributions stored by stream layout
@@ -23,7 +23,7 @@ extern "C" void dvc_UnpackValues(int *list, int count, double *recvbuf, double *
}
}
extern "C" void dvc_PackDenD3Q7(int *list, int count, double *sendbuf, int number, double *Data, int N){
extern "C" void PackDenD3Q7(int *list, int count, double *sendbuf, int number, double *Data, int N){
//....................................................................................
// Pack distribution into the send buffer for the listed lattice sites
//....................................................................................
@@ -38,7 +38,7 @@ extern "C" void dvc_PackDenD3Q7(int *list, int count, double *sendbuf, int numbe
}
extern "C" void dvc_UnpackDenD3Q7(int *list, int count, double *recvbuf, int number, double *Data, int N){
extern "C" void UnpackDenD3Q7(int *list, int count, double *recvbuf, int number, double *Data, int N){
//....................................................................................
// Unack distribution from the recv buffer
// Sum to the existing density value
@@ -52,7 +52,7 @@ extern "C" void dvc_UnpackDenD3Q7(int *list, int count, double *recvbuf, int nu
}
}
extern "C" void dvc_InitD3Q7(char *ID, double *f_even, double *f_odd, double *Den, int Nx, int Ny, int Nz)
extern "C" void InitD3Q7(char *ID, double *f_even, double *f_odd, double *Den, int Nx, int Ny, int Nz)
{
int n,N;
N = Nx*Ny*Nz;
@@ -81,7 +81,7 @@ extern "C" void dvc_InitD3Q7(char *ID, double *f_even, double *f_odd, double *De
}
//*************************************************************************
extern "C" void dvc_SwapD3Q7(char *ID, double *disteven, double *distodd, int Nx, int Ny, int Nz)
extern "C" void SwapD3Q7(char *ID, double *disteven, double *distodd, int Nx, int Ny, int Nz)
{
int i,j,k,n,nn,N;
// distributions
@@ -142,7 +142,7 @@ extern "C" void dvc_SwapD3Q7(char *ID, double *disteven, double *distodd, int Nx
}
//*************************************************************************
extern "C" void dvc_ComputeDensityD3Q7(char *ID, double *disteven, double *distodd, double *Den,
extern "C" void ComputeDensityD3Q7(char *ID, double *disteven, double *distodd, double *Den,
int Nx, int Ny, int Nz)
{
char id;

14
cpu/D3Q7.h
View File

@@ -1,16 +1,16 @@
// CPU Functions for D3Q7 Lattice Boltzmann Methods
extern "C" void dvc_PackValues(int *list, int count, double *sendbuf, double *Data, int N);
extern "C" void PackValues(int *list, int count, double *sendbuf, double *Data, int N);
extern "C" void dvc_UnpackValues(int *list, int count, double *recvbuf, double *Data, int N);
extern "C" void UnpackValues(int *list, int count, double *recvbuf, double *Data, int N);
extern "C" void dvc_PackDenD3Q7(int *list, int count, double *sendbuf, int number, double *Data, int N);
extern "C" void PackDenD3Q7(int *list, int count, double *sendbuf, int number, double *Data, int N);
extern "C" void dvc_UnpackDenD3Q7(int *list, int count, double *recvbuf, int number, double *Data, int N);
extern "C" void UnpackDenD3Q7(int *list, int count, double *recvbuf, int number, double *Data, int N);
extern "C" void dvc_InitD3Q7(char *ID, double *f_even, double *f_odd, double *Den, int Nx, int Ny, int Nz);
extern "C" void InitD3Q7(char *ID, double *f_even, double *f_odd, double *Den, int Nx, int Ny, int Nz);
extern "C" void dvc_SwapD3Q7(char *ID, double *disteven, double *distodd, int Nx, int Ny, int Nz);
extern "C" void SwapD3Q7(char *ID, double *disteven, double *distodd, int Nx, int Ny, int Nz);
extern "C" void dvc_ComputeDensityD3Q7(char *ID, double *disteven, double *distodd, double *Den,
extern "C" void ComputeDensityD3Q7(char *ID, double *disteven, double *distodd, double *Den,
int Nx, int Ny, int Nz);

8
cpu/Extras.cpp
View File

@@ -3,7 +3,7 @@
#include <stdio.h>
#include <string.h>
extern "C" void dvc_AllocateDeviceMemory(void** address, size_t size){
extern "C" void AllocateDeviceMemory(void** address, size_t size){
//cudaMalloc(address,size);
(*address) = malloc(size);
@@ -12,17 +12,17 @@ extern "C" void dvc_AllocateDeviceMemory(void** address, size_t size){
}
}
extern "C" void dvc_CopyToDevice(void* dest, void* source, size_t size){
extern "C" void CopyToDevice(void* dest, void* source, size_t size){
// cudaMemcpy(dest,source,size,cudaMemcpyHostToDevice);
memcpy(dest, source, size);
}
extern "C" void dvc_CopyToHost(void* dest, void* source, size_t size){
extern "C" void CopyToHost(void* dest, void* source, size_t size){
// cudaMemcpy(dest,source,size,cudaMemcpyDeviceToHost);
memcpy(dest, source, size);
}
extern "C" void dvc_Barrier(){
extern "C" void Barrier(){
// cudaDeviceSynchronize();
}

8
cpu/Extras.h
View File

@@ -1,7 +1,7 @@
extern "C" void dvc_AllocateDeviceMemory(void** address, size_t size);
extern "C" void AllocateDeviceMemory(void** address, size_t size);
extern "C" void dvc_CopyToDevice(void* dest, void* source, size_t size);
extern "C" void CopyToDevice(void* dest, void* source, size_t size);
extern "C" void dvc_CopyToHost(void* dest, void* source, size_t size);
extern "C" void CopyToHost(void* dest, void* source, size_t size);
extern "C" void dvc_Barrier();
extern "C" void DeviceBarrier();