<html><head><meta http-equiv="Content-Type" content="text/html; charset=us-ascii"></head><body style="word-wrap: break-word; -webkit-nbsp-mode: space; line-break: after-white-space;" class=""><br class=""><div><br class=""><blockquote type="cite" class=""><div class="">On Dec 13, 2020, at 8:39 AM, Mark Adams <<a href="mailto:mfadams@lbl.gov" class="">mfadams@lbl.gov</a>> wrote:</div><br class="Apple-interchange-newline"><div class=""><div dir="ltr" class="">So how do I declare <div class=""><br class=""></div><div class="">struct myMat {<div class=""> PetscScalar *a;</div><div class=""> PetscInt *i, *j, nrows;</div><div class="">}</div><div class=""><br class=""></div><div class="">so that I can use, or give me an alternative idiom, </div><div class=""><br class=""></div><div class="">struct myMat h_mat, *d_mat;<br class=""></div><div class="">....</div><div class="">cudaMemcpy( d_mat, &h_mat, sizeof(myMat), cudaMemcpyHostToDevice))<br class=""></div></div></div></div></blockquote><div><br class=""></div> This will copy the pointers (not the entries in the arrays) which means in both d_mat and h_mat a and i,j refer to Kokkos (or CUDA etc) memory so </div><div><br class=""></div><div><blockquote type="cite" class=""><div dir="ltr" class=""><div class="">struct myMat {<div class=""> PetscKokkosScalar *a;</div><div class=""> PetscKokkosInt *i, *j;</div></div></div></blockquote> PetscInt nrows;<blockquote type="cite" class=""><div dir="ltr" class=""><div class=""><div class="">}</div></div></div></blockquote><br class=""></div><div> the nrows is annoying because in d_mat it refers to device memory while in h_mat it refers to host memory so we cannot label it clearly.</div><div><br class=""></div><div> Sorry, by code reader I meant someone reading or maintaining the code. </div><div><br class=""></div><div> To me the above struct is clearer, I know immediately that *a and *i, *j point to Kokkos space. While labeled as PetscScalar *a I have to figure out somehow by examining other code where it points (this is a waste of time I think). </div><div><br class=""><blockquote type="cite" class=""><div class=""><div dir="ltr" class=""><div class=""><div class=""><br class=""></div><div class="">And what is PetscKokkosReal*? I guess this is in the default execution space in Kokkos. I guess that is fine.</div><div class=""><br class=""></div><div class="">But I guess I don't understand what a "code reader" is. If it's a person or a compiler you could do what I do is have a variable naming convention like h_ and d_. </div><div class=""><br class=""></div><div class="">Junchao used _d and _h in aijkokkosimpl.hpp, and these are typed for device and host. He had to type them to get it into the Kokkos API, which does type checking so I think it would catch having KokkosReal be a double and PetscReal be a float.</div><div class=""><br class=""></div><div class=""><br class=""></div></div></div><br class=""><div class="gmail_quote"><div dir="ltr" class="gmail_attr">On Sat, Dec 12, 2020 at 5:42 PM Barry Smith <<a href="mailto:bsmith@petsc.dev" class="">bsmith@petsc.dev</a>> wrote:<br class=""></div><blockquote class="gmail_quote" style="margin:0px 0px 0px 0.8ex;border-left:1px solid rgb(204,204,204);padding-left:1ex"><div style="overflow-wrap: break-word;" class=""><br class=""><div class=""><br class=""><blockquote type="cite" class=""><div class="">On Dec 12, 2020, at 4:23 PM, Mark Adams <<a href="mailto:mfadams@lbl.gov" target="_blank" class="">mfadams@lbl.gov</a>> wrote:</div><br class=""><div class=""><div dir="ltr" class=""><div dir="ltr" class=""><br class=""></div><br class=""><div class="gmail_quote"><div dir="ltr" class="gmail_attr">On Sat, Dec 12, 2020 at 10:44 AM Barry Smith <<a href="mailto:bsmith@petsc.dev" target="_blank" class="">bsmith@petsc.dev</a>> wrote:<br class=""></div><blockquote class="gmail_quote" style="margin:0px 0px 0px 0.8ex;border-left:1px solid rgb(204,204,204);padding-left:1ex"><br class="">
Currently we use PetscScalar and PetscScalar * to refer to variables that could be in any memory space. On the CPU, on the GPU, in Kokkos, etc.<br class="">
<br class="">
Would it make sense to use typedef to indicate at each location the true type of the memory location when possible? <br class=""></blockquote><div class=""><br class=""></div><div class="">No. Absolutely not.</div><div class=""><br class=""></div><div class="">Because Cuda is simple C code (eg, printf is provided but few standard libs are provided and you can't call non-device functions from the device), you can put kernels in a header file and include it in the .cu file to get Cuda code. You need to #define things like the device declaration syntax (eg __device__) and things like atomicAdd. This is how MatSetValuesDevice works.</div><div class=""><br class=""></div><div class="">The way I do deep copies in Cuda I declare a host a device struct, like:</div><div class=""><br class=""></div><div class="">Mat h_mat, *d_mat.</div><div class=""><br class=""></div><div class="">Then do cuda mallocs into pointers in h_mat, then a cuda malloc on to get d_mat. Then a cuda copy-to-device to put any data (cuda malloced) or metadata (eg, array size, dim, etc) from h_mat into d_mat. I don't know how I could do this if h_mat and d_mat are not the same without even more gymnastics.</div></div></div></div></blockquote><div class=""><br class=""></div>Oh, you would not need to change your code at all, the only difference is in certain places you would have variables declared as PetscKokkosReal that point to Kokkos memory instead of declared as PetscReal. It is, as Jed notes, just for readers/maintainers of the code ease.</div><div class=""><br class=""><blockquote type="cite" class=""><div class=""><div dir="ltr" class=""><div class="gmail_quote"><div class=""><br class=""></div><div class="">The Kokkos people have been working with this for a long time and I think they have probably learned the hard way what works. I would look at what they do. If they or SYCL does it I would take a look.</div></div></div></div></blockquote><div class=""><br class=""></div> CUDA has always had the same syntax for pointers to GPU memory and CPU memory; this seemed odd to me because it means each code reader has to find out some other way which pointers actually refer to host memory and which point to GPU memory instead of just directly saying it. Kokkos seems to have just inherited this approach. Of course with unified memory there is no distinguishing so maybe no reason to have different names.<br class=""><blockquote type="cite" class=""><div class=""><div dir="ltr" class=""><div class="gmail_quote"><div class=""><br class=""></div><div class=""> </div><blockquote class="gmail_quote" style="margin:0px 0px 0px 0.8ex;border-left:1px solid rgb(204,204,204);padding-left:1ex">
<br class="">
typedef PetscReal PetscKokkosReal means the variable is in the Kokkos memory space<br class=""></blockquote><div class=""><br class=""></div><div class="">There is no such thing. THere is the default execution space, default host space, Cuda memory space, etc.</div><div class=""> </div><blockquote class="gmail_quote" style="margin:0px 0px 0px 0.8ex;border-left:1px solid rgb(204,204,204);padding-left:1ex">
typedef PetscReal PetscCUDAReal<br class="">
typedef PetscReal PetscNVSHEMReal<br class="">
<br class="">
etc. <br class="">
<br class="">
Then one could write things like <br class="">
<br class="">
struct {<br class="">
...<br class="">
PetscNVSHEMReal *values;<br class="">
}<br class="">
<br class="">
Similarly inside kernels one would use the type type associated with the kernel, cuda with cuda etc. <br class="">
<br class="">
I think the resulting code will be much clearer and easier to dive into, then having to first figure out where each variable lives.<br class="">
<br class="">
I find the current code confusing because one cannot immediately see a variable declaration and know where it lives, even though it does live somewhere in particular..<br class="">
<br class="">
Barry<br class="">
<br class="">
<br class="">
<br class="">
<br class="">
<br class="">
</blockquote></div></div>
</div></blockquote></div><br class=""></div></blockquote></div>
</div></blockquote></div><br class=""></body></html>