<br><br><div class="gmail_quote">On Wed, Mar 23, 2011 at 4:40 PM, Hans-Christoph Steiner <span dir="ltr"><<a href="mailto:hans@at.or.at" target="_blank">hans@at.or.at</a>></span> wrote:<br><blockquote class="gmail_quote" style="border-left: 1px solid rgb(204, 204, 204); margin: 0pt 0pt 0pt 0.8ex; padding-left: 1ex;">
<div style="word-wrap: break-word;"><div><br></div><div>I think its going to be quite difficult to have a single object running in the CUDA/GPU while the rest of the patch runs on the CPU in regular Pd. My guess is that the best first step would be to implement a basic Pd in CUDA, then work from there about integrating it. Perhaps then you could use the [pd~] model.</div>
</div></blockquote><div><br>That's given me something to think about. How it might be done: functions need to be declared __device__ to run as serial processes on the GPU. There's no loss of what you can do in C on the GPU--it's just a lot slower clock and threading is different. Then, all the dynamic memory allocation routines would need to be changed. <br>
<br>The functions to be called from the host would be declared __global__.<br><br>It sounds cumbersome, but possible. Perhaps, it could be done for libpd--but then I know there's still ~10,000s of lines of code in libpd. I'm not a libpd user, so I don't know how I'd test it without the gui/editor. That's just conjecture, but I can't see a good way to simplify it yet.<br>
<br>I think it's better to only apply CUDA to the major vectorizable operations, which is just signal operations, but also--I couldn't just apply it to *all* signal operations. So, I'd try to add the basic compatibility stuff for managing signals allocated on GPU... and then go back and try to incrementally add new signal externals from the Pd math tilde objects.<br>
<br></div>
<blockquote class="gmail_quote" style="border-left: 1px solid rgb(204, 204, 204); margin: 0pt 0pt 0pt 0.8ex; padding-left: 1ex;"><div style="word-wrap: break-word;"><div><br></div><div>For examples of reimplmentations of Pd, check out ZenGarden (C++) and Webpd (Javascript) <a href="http://mccormick.cx/projects/WebPd/" target="_blank">http://mccormick.cx/projects/WebPd/</a></div>
<div><br></div><div>.hc</div></div></blockquote><div><br>I've been having fun playing with WebPd--but I'm not needing that tool just yet. Soon... soon...<br><br> </div><blockquote class="gmail_quote" style="border-left: 1px solid rgb(204, 204, 204); margin: 0pt 0pt 0pt 0.8ex; padding-left: 1ex;">
<div style="word-wrap: break-word;"><div><div></div><div><br><div><div>On Mar 23, 2011, at 5:20 PM, Charles Henry wrote:</div><br><blockquote type="cite">Hi, hc<br><br>Let me explain a little further here. The end goal is to have an external library that allows one to create externals that use memory on GPUs. Big idea here is that once you've got a system for handling the memory allocation and dsp sorting in *exactly* the same way as Pd, then you can write externals for CUDA or CL in a way that's consistent with existing externals. <br>
<br>With Pd handling the signal memory allocation inside of d_ugen.c and called from canvas_dodsp, I wanted for my external library to have its own canvas class and different methods for handling the memory allocation differently. In fact, I think of that as being the key class to create the library.<br>
<br>I worked through it for a while, and I think it's just plain impossible to have another canvas class in an external library, unless there's something good I don't understand. And I really want to understand :)<br>
<br>So, since then, I've been thinking that I'd have to modify the pd-vanilla source code. I've got something that loads a CUDA device and gets me to run code during canvas_new, canvas_free, and canvas_dsp. I'm still trying to organize around the idea of having an external library to load and make as few changes directly in the vanilla source code.<br>
<br>Chuck<br><br><div class="gmail_quote">On Wed, Mar 23, 2011 at 1:51 PM, Hans-Christoph Steiner <span dir="ltr"><<a href="mailto:hans@at.or.at" target="_blank">hans@at.or.at</a>></span> wrote:<br><blockquote class="gmail_quote" style="border-left: 1px solid rgb(204, 204, 204); margin: 0pt 0pt 0pt 0.8ex; padding-left: 1ex;">
<br> What's the end goal here? You want an object that acts like a t_canvas/t_glist?<br> <br> .hc<div><div></div><div><br> <br> On Mar 13, 2011, at 3:31 PM, Charles Henry wrote:<br> <br> </div></div><blockquote class="gmail_quote" style="border-left: 1px solid rgb(204, 204, 204); margin: 0pt 0pt 0pt 0.8ex; padding-left: 1ex;">
<div><div></div><div> I've been working through my CUDA Pd project, and I ran into the problem of making externals that copy the canvas class.<br> <br> My first idea was that I wanted a completely separate class with different methods using glist. Calls from Pd looking for t_canvas work just fine, but functions like pd_findbyclass that look for canvas_class fail. I started mucking around in the pd src, and I think it's just too difficult and would make onerous changes that I don't like.<br>
<br> Is there something I'm not getting about canvas classes and externals?<br> <br> My second approach to creating an external library is to modify glist by adding an "unsigned int gl_hascuda" variable. I'd still prefer solutions that make use of entirely external libraries over modifying src, but this small change gets me half the way there. Then, I just need to write the creator functions and the class works.<br>
<br> Chuck<br></div></div> _______________________________________________<br> Pd-dev mailing list<br> <a href="mailto:Pd-dev@iem.at" target="_blank">Pd-dev@iem.at</a><br> <a href="http://lists.puredata.info/listinfo/pd-dev" target="_blank">http://lists.puredata.info/listinfo/pd-dev</a><br>
</blockquote> <br> <br> <br> <br> <br> ----------------------------------------------------------------------------<br> <br> "[T]he greatest purveyor of violence in the world today [is] my own government." - Martin Luther King, Jr.<br>
<br> <br> <br> </blockquote></div><br></blockquote></div><br></div></div><div> <span style="border-collapse: separate; color: rgb(0, 0, 0); font-family: Helvetica; font-size: medium; font-style: normal; font-variant: normal; font-weight: normal; letter-spacing: normal; line-height: normal; text-indent: 0px; text-transform: none; white-space: normal; word-spacing: 0px;"><span style="font-size: 12px;"><div>
<br>----------------------------------------------------------------------------<br></div></span><div><br></div><div>"A cellphone to me is just an opportunity to be irritated wherever you are." - Linus Torvalds</div>
</span> </div><br></div></blockquote></div><br>