Commits

Karl Rupp committed 451b912 Merge

Merge branch 'cj/fix-petsccusp'

  • Participants
  • Parent commits ad0a213, 3c0b9f6

Comments (0)

Files changed (4)

include/petsccusp.h

 #if !defined(__PETSCCUSP_H)
 #define __PETSCCUSP_H
-/*
-    This should only be included in user code that uses CUSP directly and hence the file name ends with .cu
-*/
-#include <../src/vec/vec/impls/dvecimpl.h>
-#include <../src/vec/vec/impls/seq/seqcusp/cuspvecimpl.h>
+
+#include <petscsys.h>
+#include <petscmath.h>
+#include <petscvec.h>
+#include <cusp/array1d.h>
+
+PETSC_EXTERN PetscErrorCode VecCUSPGetArrayReadWrite(Vec v, cusp::array1d<PetscScalar,cusp::device_memory> **a);
+PETSC_EXTERN PetscErrorCode VecCUSPRestoreArrayReadWrite(Vec v, cusp::array1d<PetscScalar,cusp::device_memory> **a);
+
+PETSC_EXTERN PetscErrorCode VecCUSPGetArrayRead(Vec v, cusp::array1d<PetscScalar,cusp::device_memory> **a);
+PETSC_EXTERN PetscErrorCode VecCUSPRestoreArrayRead(Vec v, cusp::array1d<PetscScalar,cusp::device_memory> **a);
+
+PETSC_EXTERN PetscErrorCode VecCUSPGetArrayWrite(Vec v, cusp::array1d<PetscScalar,cusp::device_memory> **a);
+PETSC_EXTERN PetscErrorCode VecCUSPRestoreArrayWrite(Vec v, cusp::array1d<PetscScalar,cusp::device_memory> **a);
+
 #endif

src/snes/examples/tutorials/ex47cu.cu

 #include <petscsnes.h>
 #include <petsccusp.h>
 
+#include <thrust/for_each.h>
+#include <thrust/tuple.h>
+#include <thrust/iterator/constant_iterator.h>
+#include <thrust/iterator/counting_iterator.h>
+#include <thrust/iterator/zip_iterator.h>
+
 extern PetscErrorCode ComputeFunction(SNES,Vec,Vec,void*), ComputeJacobian(SNES,Vec,Mat*,Mat*,MatStructure*,void*);
 PetscBool useCUSP = PETSC_FALSE;
 
 
 PetscErrorCode ComputeFunction(SNES snes,Vec x,Vec f,void *ctx)
 {
-  PetscInt       i,Mx,xs,xm,xstartshift,xendshift,fstart;
+  PetscInt       i,Mx,xs,xm,xstartshift,xendshift,fstart,lsize;
   PetscScalar    *xx,*ff,hx;
   DM             da = (DM) ctx;
   Vec            xlocal;
   PetscErrorCode ierr;
   PetscMPIInt    rank,size;
   MPI_Comm       comm;
-  CUSPARRAY      *xarray,*farray;
+  cusp::array1d<PetscScalar,cusp::device_memory> *xarray,*farray;
 
   ierr = DMDAGetInfo(da,PETSC_IGNORE,&Mx,PETSC_IGNORE,PETSC_IGNORE,PETSC_IGNORE,PETSC_IGNORE,PETSC_IGNORE,PETSC_IGNORE,PETSC_IGNORE,PETSC_IGNORE,PETSC_IGNORE,PETSC_IGNORE,PETSC_IGNORE);CHKERRQ(ierr);
   hx   = 1.0/(PetscReal)(Mx-1);
     if (rank != size-1) xendshift = 1;
     else xendshift = 0;
     ierr = VecGetOwnershipRange(f,&fstart,NULL);CHKERRQ(ierr);
+    ierr = VecGetLocalSize(x,&lsize);CHKERRQ(ierr);
     try {
       thrust::for_each(
         thrust::make_zip_iterator(
             xarray->end()-xendshift,
             xarray->end()-xendshift + 1,
             xarray->end()-xendshift - 1,
-            thrust::counting_iterator<int>(fstart) + x->map->n,
+            thrust::counting_iterator<int>(fstart) + lsize,
             thrust::constant_iterator<int>(Mx),
             thrust::constant_iterator<PetscScalar>(hx))),
         ApplyStencil());

src/vec/vec/impls/seq/seqcusp/cuspvecimpl.h

 #if !defined(__CUSPVECIMPL)
 #define __CUSPVECIMPL
 
+#include <petsccusp.h>
 #include <petsc-private/vecimpl.h>
 
 #include <algorithm>
 #endif
 };
 
-
-#undef __FUNCT__
-#define __FUNCT__ "VecCUSPGetArrayReadWrite"
-PETSC_STATIC_INLINE PetscErrorCode VecCUSPGetArrayReadWrite(Vec v, CUSPARRAY **a)
-{
-  PetscErrorCode ierr;
-
-  PetscFunctionBegin;
-  *a   = 0;
-  ierr = VecCUSPCopyToGPU(v);CHKERRQ(ierr);
-  *a   = ((Vec_CUSP*)v->spptr)->GPUarray;
-  PetscFunctionReturn(0);
-}
-
-#undef __FUNCT__
-#define __FUNCT__ "VecCUSPRestoreArrayReadWrite"
-PETSC_STATIC_INLINE PetscErrorCode VecCUSPRestoreArrayReadWrite(Vec v, CUSPARRAY **a)
-{
-  PetscErrorCode ierr;
-
-  PetscFunctionBegin;
-  v->valid_GPU_array = PETSC_CUSP_GPU;
-
-  ierr = PetscObjectStateIncrease((PetscObject)v);CHKERRQ(ierr);
-  PetscFunctionReturn(0);
-}
-
-#undef __FUNCT__
-#define __FUNCT__ "VecCUSPGetArrayRead"
-PETSC_STATIC_INLINE PetscErrorCode VecCUSPGetArrayRead(Vec v, CUSPARRAY **a)
-{
-  PetscErrorCode ierr;
-
-  PetscFunctionBegin;
-  *a   = 0;
-  ierr = VecCUSPCopyToGPU(v);CHKERRQ(ierr);
-  *a   = ((Vec_CUSP*)v->spptr)->GPUarray;
-  PetscFunctionReturn(0);
-}
-
-#undef __FUNCT__
-#define __FUNCT__ "VecCUSPRestoreArrayRead"
-PETSC_STATIC_INLINE PetscErrorCode VecCUSPRestoreArrayRead(Vec v, CUSPARRAY **a)
-{
-  PetscFunctionBegin;
-  PetscFunctionReturn(0);
-}
-
-#undef __FUNCT__
-#define __FUNCT__ "VecCUSPGetArrayWrite"
-PETSC_STATIC_INLINE PetscErrorCode VecCUSPGetArrayWrite(Vec v, CUSPARRAY **a)
-{
-  PetscErrorCode ierr;
-
-  PetscFunctionBegin;
-  *a   = 0;
-  ierr = VecCUSPAllocateCheck(v);CHKERRQ(ierr);
-  *a   = ((Vec_CUSP*)v->spptr)->GPUarray;
-  PetscFunctionReturn(0);
-}
-
-#undef __FUNCT__
-#define __FUNCT__ "VecCUSPRestoreArrayWrite"
-PETSC_STATIC_INLINE PetscErrorCode VecCUSPRestoreArrayWrite(Vec v, CUSPARRAY **a)
-{
-  PetscErrorCode ierr;
-
-  PetscFunctionBegin;
-  v->valid_GPU_array = PETSC_CUSP_GPU;
-
-  ierr = PetscObjectStateIncrease((PetscObject)v);CHKERRQ(ierr);
-  PetscFunctionReturn(0);
-}
 #endif

src/vec/vec/impls/seq/seqcusp/veccusp.cu

   ierr = VecSet(V,0.0);CHKERRQ(ierr);
   PetscFunctionReturn(0);
 }
+
+#undef __FUNCT__
+#define __FUNCT__ "VecCUSPGetArrayReadWrite"
+PETSC_EXTERN PetscErrorCode VecCUSPGetArrayReadWrite(Vec v, CUSPARRAY **a)
+{
+  PetscErrorCode ierr;
+
+  PetscFunctionBegin;
+  *a   = 0;
+  ierr = VecCUSPCopyToGPU(v);CHKERRQ(ierr);
+  *a   = ((Vec_CUSP*)v->spptr)->GPUarray;
+  PetscFunctionReturn(0);
+}
+
+#undef __FUNCT__
+#define __FUNCT__ "VecCUSPRestoreArrayReadWrite"
+PETSC_EXTERN PetscErrorCode VecCUSPRestoreArrayReadWrite(Vec v, CUSPARRAY **a)
+{
+  PetscErrorCode ierr;
+
+  PetscFunctionBegin;
+  v->valid_GPU_array = PETSC_CUSP_GPU;
+
+  ierr = PetscObjectStateIncrease((PetscObject)v);CHKERRQ(ierr);
+  PetscFunctionReturn(0);
+}
+
+#undef __FUNCT__
+#define __FUNCT__ "VecCUSPGetArrayRead"
+PETSC_EXTERN PetscErrorCode VecCUSPGetArrayRead(Vec v, CUSPARRAY **a)
+{
+  PetscErrorCode ierr;
+
+  PetscFunctionBegin;
+  *a   = 0;
+  ierr = VecCUSPCopyToGPU(v);CHKERRQ(ierr);
+  *a   = ((Vec_CUSP*)v->spptr)->GPUarray;
+  PetscFunctionReturn(0);
+}
+
+#undef __FUNCT__
+#define __FUNCT__ "VecCUSPRestoreArrayRead"
+PETSC_EXTERN PetscErrorCode VecCUSPRestoreArrayRead(Vec v, CUSPARRAY **a)
+{
+  PetscFunctionBegin;
+  PetscFunctionReturn(0);
+}
+
+#undef __FUNCT__
+#define __FUNCT__ "VecCUSPGetArrayWrite"
+PETSC_EXTERN PetscErrorCode VecCUSPGetArrayWrite(Vec v, CUSPARRAY **a)
+{
+  PetscErrorCode ierr;
+
+  PetscFunctionBegin;
+  *a   = 0;
+  ierr = VecCUSPAllocateCheck(v);CHKERRQ(ierr);
+  *a   = ((Vec_CUSP*)v->spptr)->GPUarray;
+  PetscFunctionReturn(0);
+}
+
+#undef __FUNCT__
+#define __FUNCT__ "VecCUSPRestoreArrayWrite"
+PETSC_EXTERN PetscErrorCode VecCUSPRestoreArrayWrite(Vec v, CUSPARRAY **a)
+{
+  PetscErrorCode ierr;
+
+  PetscFunctionBegin;
+  v->valid_GPU_array = PETSC_CUSP_GPU;
+
+  ierr = PetscObjectStateIncrease((PetscObject)v);CHKERRQ(ierr);
+  PetscFunctionReturn(0);
+}