<div dir="ltr">FYI, CUDA is running and here is some preliminary data on up to 1/8 of SUMMIT. This run with 4 cores/processes per GPU, so the GPU is virtualized into 4 GPUs.</div><br><div class="gmail_quote"><div dir="ltr" class="gmail_attr">On Sun, Jul 28, 2019 at 2:34 PM Karl Rupp <<a href="mailto:rupp@iue.tuwien.ac.at">rupp@iue.tuwien.ac.at</a>> wrote:<br></div><blockquote class="gmail_quote" style="margin:0px 0px 0px 0.8ex;border-left:1px solid rgb(204,204,204);padding-left:1ex">Hi Mark,<br>
<br>
feel free to submit a fresh pull request now. I looked at your latest <br>
commit in the repository in order to cherry-pick it, but it looked like <br>
it had a few other bits in it as well.<br>
<br>
Best regards,<br>
Karli<br>
<br>
<br>
On 7/28/19 6:27 PM, Mark Adams via petsc-dev wrote:<br>
> This is looking good. I'm not seeing the numerical problems, but I've <br>
> just hid them by avoiding the GPU on coarse grids.<br>
> <br>
> Should I submit a pull request now or test more or wait for Karl?<br>
> <br>
> On Sat, Jul 27, 2019 at 7:37 PM Mark Adams <<a href="mailto:mfadams@lbl.gov" target="_blank">mfadams@lbl.gov</a> <br>
> <mailto:<a href="mailto:mfadams@lbl.gov" target="_blank">mfadams@lbl.gov</a>>> wrote:<br>
> <br>
> Barry, I fixed CUDA to pin to CPUs correctly for GAMG at least.<br>
> There are some hacks here that we can work on.<br>
> <br>
> I will start testing it tomorrow, but I am pretty sure that I have<br>
> not regressed. I am hoping that this will fix the numerical<br>
> problems, which seem to be associated with empty processors.<br>
> <br>
> I did need to touch code outside of GAMG and CUDA. It might be nice<br>
> to test this in a next.<br>
> <br>
> GAMG now puts all reduced processorg grids on the CPU. This could be<br>
> looked at in the future.<br>
> <br>
> <br>
> On Sat, Jul 27, 2019 at 1:00 PM Smith, Barry F. <<a href="mailto:bsmith@mcs.anl.gov" target="_blank">bsmith@mcs.anl.gov</a><br>
> <mailto:<a href="mailto:bsmith@mcs.anl.gov" target="_blank">bsmith@mcs.anl.gov</a>>> wrote:<br>
> <br>
> <br>
> <br>
> > On Jul 27, 2019, at 11:53 AM, Mark Adams <<a href="mailto:mfadams@lbl.gov" target="_blank">mfadams@lbl.gov</a><br>
> <mailto:<a href="mailto:mfadams@lbl.gov" target="_blank">mfadams@lbl.gov</a>>> wrote:<br>
> ><br>
> ><br>
> > On Sat, Jul 27, 2019 at 11:39 AM Smith, Barry F.<br>
> <<a href="mailto:bsmith@mcs.anl.gov" target="_blank">bsmith@mcs.anl.gov</a> <mailto:<a href="mailto:bsmith@mcs.anl.gov" target="_blank">bsmith@mcs.anl.gov</a>>> wrote:<br>
> ><br>
> > Good catch. Thanks. Maybe the SeqCUDA has the same problem?<br>
> ><br>
> > THis is done (I may have done it).<br>
> ><br>
> > Now it seems to me that when you call VecPinToCPU you are<br>
> setting up and don't have data, so this copy does not seem<br>
> necessary. Maybe remove the copy here:<br>
> ><br>
> > PetscErrorCode VecPinToCPU_MPICUDA(Vec V,PetscBool pin)<br>
> > {<br>
> > PetscErrorCode ierr;<br>
> ><br>
> > PetscFunctionBegin;<br>
> > V->pinnedtocpu = pin;<br>
> > if (pin) {<br>
> > ierr = VecCUDACopyFromGPU(V);CHKERRQ(ierr); ????<br>
> <br>
> The copy from GPU should actually only do anything if the<br>
> GPU already has data and PETSC_OFFLOAD_GPU. If the GPU does not<br>
> have data<br>
> the copy doesn't do anything. When one calls VecPinToCPU() one<br>
> doesn't know where the data is so the call must be made, but it<br>
> may do nothing<br>
> <br>
> Note that VecCUDACopyFromGPU() calls<br>
> VecCUDAAllocateCheckHost() not VecCUDAAllocateCheck() so the GPU<br>
> will not allocate space,<br>
> VecCUDAAllocateCheck() is called from VecCUDACopyToGPU().<br>
> <br>
> Yes, perhaps the naming could be more consistent:<br>
> <br>
> 1) in one place it is Host in an other place it is nothing<br>
> 2) some places it is Host, Device, some places GPU,CPU<br>
> <br>
> Perhaps Karl can make these all consistent and simpler in<br>
> his refactorization<br>
> <br>
> <br>
> Barry<br>
> <br>
> <br>
> ><br>
> > or<br>
> ><br>
> > Not allocate the GPU if it is pinned by added in a check here:<br>
> ><br>
> > PetscErrorCode VecCUDAAllocateCheck(Vec v)<br>
> > {<br>
> > PetscErrorCode ierr;<br>
> > cudaError_t err;<br>
> > cudaStream_t stream;<br>
> > Vec_CUDA *veccuda;<br>
> ><br>
> > PetscFunctionBegin;<br>
> > if (!v->spptr) {<br>
> > ierr = PetscMalloc(sizeof(Vec_CUDA),&v->spptr);CHKERRQ(ierr);<br>
> > veccuda = (Vec_CUDA*)v->spptr;<br>
> > if (v->valid_GPU_array != PETSC_OFFLOAD_CPU) {<br>
> > err =<br>
> cudaMalloc((void**)&veccuda->GPUarray_allocated,sizeof(PetscScalar)*((PetscBLASInt)v->map->n));CHKERRCUDA(err);<br>
> > veccuda->GPUarray = veccuda->GPUarray_allocated;<br>
> > err = cudaStreamCreate(&stream);CHKERRCUDA(err);<br>
> > veccuda->stream = stream;<br>
> > veccuda->hostDataRegisteredAsPageLocked = PETSC_FALSE;<br>
> > if (v->valid_GPU_array == PETSC_OFFLOAD_UNALLOCATED) {<br>
> > if (v->data && ((Vec_Seq*)v->data)->array) {<br>
> > v->valid_GPU_array = PETSC_OFFLOAD_CPU;<br>
> > } else {<br>
> > v->valid_GPU_array = PETSC_OFFLOAD_GPU;<br>
> > }<br>
> > }<br>
> > }<br>
> > }<br>
> > PetscFunctionReturn(0);<br>
> > }<br>
> ><br>
> ><br>
> ><br>
> ><br>
> ><br>
> > > On Jul 27, 2019, at 10:40 AM, Mark Adams <<a href="mailto:mfadams@lbl.gov" target="_blank">mfadams@lbl.gov</a><br>
> <mailto:<a href="mailto:mfadams@lbl.gov" target="_blank">mfadams@lbl.gov</a>>> wrote:<br>
> > ><br>
> > > Yea, I just figured out the problem. VecDuplicate_MPICUDA<br>
> did not call PinToCPU or even copy pinnedtocpu. It just copied<br>
> ops, so I added and am testing:<br>
> > ><br>
> > > ierr =<br>
> VecCreate_MPICUDA_Private(*v,PETSC_TRUE,w->nghost,0);CHKERRQ(ierr);<br>
> > > vw = (Vec_MPI*)(*v)->data;<br>
> > > ierr = PetscMemcpy((*v)->ops,win->ops,sizeof(struct<br>
> _VecOps));CHKERRQ(ierr);<br>
> > > ierr = VecPinToCPU(*v,win->pinnedtocpu);CHKERRQ(ierr);<br>
> > ><br>
> > > Thanks,<br>
> > ><br>
> > > On Sat, Jul 27, 2019 at 11:33 AM Smith, Barry F.<br>
> <<a href="mailto:bsmith@mcs.anl.gov" target="_blank">bsmith@mcs.anl.gov</a> <mailto:<a href="mailto:bsmith@mcs.anl.gov" target="_blank">bsmith@mcs.anl.gov</a>>> wrote:<br>
> > ><br>
> > > I don't understand the context. Once a vector is pinned<br>
> to the CPU the flag should be PETSC_OFFLOAD_CPU permanently<br>
> until the pin to cpu is turned off. Do you have a pinned vector<br>
> that has the value PETSC_OFFLOAD_GPU? For example here it is<br>
> set to PETSC_OFFLOAD_CPU<br>
> > ><br>
> > > PetscErrorCode VecPinToCPU_MPICUDA(Vec V,PetscBool pin)<br>
> > > {<br>
> > > ....<br>
> > > if (pin) {<br>
> > > ierr = VecCUDACopyFromGPU(V);CHKERRQ(ierr);<br>
> > > V->valid_GPU_array = PETSC_OFFLOAD_CPU; /* since the<br>
> CPU code will likely change values in the vector */<br>
> > ><br>
> > ><br>
> > > Is there any way to reproduce the problem?<br>
> > ><br>
> > > Barry<br>
> > ><br>
> > ><br>
> > ><br>
> > ><br>
> > > > On Jul 27, 2019, at 10:28 AM, Mark Adams <<a href="mailto:mfadams@lbl.gov" target="_blank">mfadams@lbl.gov</a><br>
> <mailto:<a href="mailto:mfadams@lbl.gov" target="_blank">mfadams@lbl.gov</a>>> wrote:<br>
> > > ><br>
> > > > I'm not sure what to do here. The problem is that<br>
> pinned-to-cpu vectors are calling VecCUDACopyFromGPU here.<br>
> > > ><br>
> > > > Should I set x->valid_GPU_array to something else, like<br>
> PETSC_OFFLOAD_CPU, in PinToCPU so this block of code i s not<br>
> executed?<br>
> > > ><br>
> > > > PetscErrorCode VecGetArray(Vec x,PetscScalar **a)<br>
> > > > {<br>
> > > > PetscErrorCode ierr;<br>
> > > > #if defined(PETSC_HAVE_VIENNACL)<br>
> > > > PetscBool is_viennacltype = PETSC_FALSE;<br>
> > > > #endif<br>
> > > ><br>
> > > > PetscFunctionBegin;<br>
> > > > PetscValidHeaderSpecific(x,VEC_CLASSID,1);<br>
> > > > ierr = VecSetErrorIfLocked(x,1);CHKERRQ(ierr);<br>
> > > > if (x->petscnative) {<br>
> > > > #if defined(PETSC_HAVE_VIENNACL) || defined(PETSC_HAVE_CUDA)<br>
> > > > if (x->valid_GPU_array == PETSC_OFFLOAD_GPU) {<br>
> > > > #if defined(PETSC_HAVE_VIENNACL)<br>
> > > > ierr =<br>
> PetscObjectTypeCompareAny((PetscObject)x,&is_viennacltype,VECSEQVIENNACL,VECMPIVIENNACL,VECVIENNACL,"");CHKERRQ(ierr);<br>
> > > > if (is_viennacltype) {<br>
> > > > ierr = VecViennaCLCopyFromGPU(x);CHKERRQ(ierr);<br>
> > > > } else<br>
> > > > #endif<br>
> > > > {<br>
> > > > #if defined(PETSC_HAVE_CUDA)<br>
> > > > ierr = VecCUDACopyFromGPU(x);CHKERRQ(ierr);<br>
> > > > #endif<br>
> > > > }<br>
> > > > } else if (x->valid_GPU_array ==<br>
> PETSC_OFFLOAD_UNALLOCATED) {<br>
> > > > #if defined(PETSC_HAVE_VIENNACL)<br>
> > > > ierr =<br>
> PetscObjectTypeCompareAny((PetscObject)x,&is_viennacltype,VECSEQVIENNACL,VECMPIVIENNACL,VECVIENNACL,"");CHKERRQ(ierr);<br>
> > > > if (is_viennacltype) {<br>
> > > > ierr = VecViennaCLAllocateCheckHost(x);CHKERRQ(ierr);<br>
> > > > } else<br>
> > > > #endif<br>
> > > > {<br>
> > > > #if defined(PETSC_HAVE_CUDA)<br>
> > > > ierr = VecCUDAAllocateCheckHost(x);CHKERRQ(ierr);<br>
> > > > #endif<br>
> > > > }<br>
> > > > }<br>
> > > > #endif<br>
> > > > *a = *((PetscScalar**)x->data);<br>
> > > > } else {<br>
> > > ><br>
> > > ><br>
> > > > On Tue, Jul 23, 2019 at 9:18 PM Smith, Barry F.<br>
> <<a href="mailto:bsmith@mcs.anl.gov" target="_blank">bsmith@mcs.anl.gov</a> <mailto:<a href="mailto:bsmith@mcs.anl.gov" target="_blank">bsmith@mcs.anl.gov</a>>> wrote:<br>
> > > ><br>
> > > > Yes, it needs to be able to switch back and forth<br>
> between the CPU and GPU methods so you need to move into it the<br>
> setting of the methods that is currently directly in the create<br>
> method. See how MatConvert_SeqAIJ_SeqAIJViennaCL() calls ierr =<br>
> MatPinToCPU_SeqAIJViennaCL(A,PETSC_FALSE);CHKERRQ(ierr); to set<br>
> the methods for the GPU initially.<br>
> > > ><br>
> > > > Barry<br>
> > > ><br>
> > > ><br>
> > > > > On Jul 23, 2019, at 7:32 PM, Mark Adams<br>
> <<a href="mailto:mfadams@lbl.gov" target="_blank">mfadams@lbl.gov</a> <mailto:<a href="mailto:mfadams@lbl.gov" target="_blank">mfadams@lbl.gov</a>>> wrote:<br>
> > > > ><br>
> > > > ><br>
> > > > > What are the symptoms of it not working? Does it<br>
> appear to be still copying the matrices to the GPU? then running<br>
> the functions on the GPU?<br>
> > > > ><br>
> > > > ><br>
> > > > > The object is dispatching the CUDA mat-vec etc.<br>
> > > > ><br>
> > > > > I suspect the pinning is incompletely done for CUDA<br>
> (and MPIOpenCL) matrices.<br>
> > > > ><br>
> > > > ><br>
> > > > > Yes, git grep MatPinToCPU shows stuff for ViennaCL but<br>
> not CUDA.<br>
> > > > ><br>
> > > > > I guess I can add something like this below. Do we need<br>
> to set the device methods? They are already set when this method<br>
> is set, right?<br>
> > > > ><br>
> > > > > We need the equivalent of<br>
> > > > ><br>
> > > > > static PetscErrorCode MatPinToCPU_SeqAIJViennaCL(Mat<br>
> A,PetscBool flg)<br>
> > > > > {<br>
> > > > > PetscFunctionBegin;<br>
> > > > > A->pinnedtocpu = flg;<br>
> > > > > if (flg) {<br>
> > > > > A->ops->mult = MatMult_SeqAIJ;<br>
> > > > > A->ops->multadd = MatMultAdd_SeqAIJ;<br>
> > > > > A->ops->assemblyend = MatAssemblyEnd_SeqAIJ;<br>
> > > > > A->ops->duplicate = MatDuplicate_SeqAIJ;<br>
> > > > > } else {<br>
> > > > > A->ops->mult = MatMult_SeqAIJViennaCL;<br>
> > > > > A->ops->multadd = MatMultAdd_SeqAIJViennaCL;<br>
> > > > > A->ops->assemblyend = MatAssemblyEnd_SeqAIJViennaCL;<br>
> > > > > A->ops->destroy = MatDestroy_SeqAIJViennaCL;<br>
> > > > > A->ops->duplicate = MatDuplicate_SeqAIJViennaCL;<br>
> > > > > }<br>
> > > > > PetscFunctionReturn(0);<br>
> > > > > }<br>
> > > > ><br>
> > > > > for MPIViennaCL and MPISeqAIJ Cusparse but it doesn't<br>
> look like it has been written yet.<br>
> > > > ><br>
> > > > ><br>
> > > > > ><br>
> > > > > > It does not seem to work. It does not look like CUDA<br>
> has an MatCreateVecs. Should I add one and copy this flag over?<br>
> > > > ><br>
> > > > > We do need this function. But I don't see how it<br>
> relates to pinning. When the matrix is pinned to the CPU we want<br>
> it to create CPU vectors which I assume it does.<br>
> > > > ><br>
> > > > ><br>
> > > > > ><br>
> > > > > > Mark<br>
> > > > ><br>
> > > ><br>
> > ><br>
> ><br>
> <br>
</blockquote></div>