Does numba.cuda has tex2D and tex3D as cuda for using texture memory?

[Pei Li]

Hi, Is there anyone knowing whether numba.cuda or accelerate has functions for using and managing texture memory of GPU? The texture memory is always used to finish the streaming process of Lattice Boltzmann Method, but I do not find how to use texture memory with numba so far....... 

[Siu Kwan Lam]

Hi Pei Li,

We don't have texture support yet.  I don't have first-hand experience for texture but I have heard mixed opinion on texture performance (esp. in new compute capabilities).  Because of that, we have avoided the extra complexity of texture memory so far.  

Can you tell us about your experience with texture memory?  Are you observing significant speedup with texture memory when comparing to global memory in your application?  

If I were to implement texture memory support, I would probably go for the bindless texture object introduced for Kepler.

On Tue, Aug 2, 2016 at 3:40 PM Pei Li <lipeihyd...@gmail.com> wrote:

Hi, Is there anyone knowing whether numba.cuda or accelerate has functions for using and managing texture memory of GPU? The texture memory is always used to finish the streaming process of Lattice Boltzmann Method, but I do not find how to use texture memory with numba so far....... 

--
You received this message because you are subscribed to the Google Groups "Numba Public Discussion - Public" group.
To unsubscribe from this group and stop receiving emails from it, send an email to numba-users...@continuum.io.
To post to this group, send email to numba...@continuum.io.
To view this discussion on the web visit https://groups.google.com/a/continuum.io/d/msgid/numba-users/eec9795e-abca-40e9-8ebc-f09f3622c900%40continuum.io.
For more options, visit https://groups.google.com/a/continuum.io/d/optout.

--

Siu Kwan Lam

Software Engineer

Continuum Analytics

[Pei Li]

Hi, Siu,

Thanks so much for your explanation. Sorry for missing the mail reminder of your reply. Indeed, I am a freshman for CUDA and numba, so I did not really implement any codes of using texture memory. So far, I used another way of using shared memory to solve my problem partially. With the shared memory, the grids can exchange distribution value much faster than using global memory (at least 10 times.) However, it limits the number of grids in each direction. It means that the number of grids must be 2^n; otherwise, the distribution value cannot be transferred correctly in the last column of blocks. After reading several papers, I found that using texture memory can fix this limitation. This is the reason why I wondered whether numba.cuda can control texture memory. If there is no function for using texture memory of GPU, probably, I may write a part of CUDA C to use texture memory and keep you updated.

Thank you very much,

Pei

[Siu Kwan Lam]

Hi Pei,

I read a little on Lattice Boltzmann method.  Don't you need writable memory?  Beware that texture memory is read-only.

If you are writing in C, I strongly suggest doing a quick benchmark of texture vs shared memory on your GPUs first.  Writing in C may take much longer and you want to be sure it is worth it.  

To view this discussion on the web visit https://groups.google.com/a/continuum.io/d/msgid/numba-users/5ef59765-ea09-48cd-93ec-93d81c52fb9f%40continuum.io.

For more options, visit https://groups.google.com/a/continuum.io/d/optout.

[Pei Li]

Hi, Siu,

Thanks very much for your advice. For LBM, there is one process called 'streaming' and in this process each distribution function of specific direction on the lattice/grid just needs to update values from neighbors. The texture memory has the advantage of reading values from adjcent address and fetching the values on the boundary of each block. The latter one is greatly important to finish 'streaming process'. This is why using texture memory is very common and popular in LBM.  I already found one example code in CUDA C and, probably, it can help me save time to learn the principle of using texture memory.

To unsubscribe from this group and stop receiving emails from it, send an email to numba-users+unsubscribe@continuum.io.

To post to this group, send email to numba...@continuum.io.
To view this discussion on the web visit https://groups.google.com/a/continuum.io/d/msgid/numba-users/5ef59765-ea09-48cd-93ec-93d81c52fb9f%40continuum.io.
For more options, visit https://groups.google.com/a/continuum.io/d/optout.

--

Siu Kwan Lam

Software Engineer

Continuum Analytics

--

You received this message because you are subscribed to the Google Groups "Numba Public Discussion - Public" group.

To unsubscribe from this group and stop receiving emails from it, send an email to numba-users+unsubscribe@continuum.io.

To post to this group, send email to numba...@continuum.io.

To view this discussion on the web visit https://groups.google.com/a/continuum.io/d/msgid/numba-users/CAHjo7%3DXSLNKeRcPEprVM%3DaS8-DLR7L5eMX%3DDkFB0R3QhUXZtSg%40mail.gmail.com.

[Siu Kwan Lam]

Pei,

Can you paste the link to the example code CUDA C code?  I can use it to learn about this usecase.

Thanks,

To view this discussion on the web visit https://groups.google.com/a/continuum.io/d/msgid/numba-users/5ef59765-ea09-48cd-93ec-93d81c52fb9f%40continuum.io.
For more options, visit https://groups.google.com/a/continuum.io/d/optout.

--

Siu Kwan Lam

Software Engineer

Continuum Analytics

--
You received this message because you are subscribed to the Google Groups "Numba Public Discussion - Public" group.

To unsubscribe from this group and stop receiving emails from it, send an email to numba-users...@continuum.io.
To post to this group, send email to numba...@continuum.io.

To view this discussion on the web visit https://groups.google.com/a/continuum.io/d/msgid/numba-users/CAHjo7%3DXSLNKeRcPEprVM%3DaS8-DLR7L5eMX%3DDkFB0R3QhUXZtSg%40mail.gmail.com.

For more options, visit https://groups.google.com/a/continuum.io/d/optout.

--
You received this message because you are subscribed to the Google Groups "Numba Public Discussion - Public" group.
To unsubscribe from this group and stop receiving emails from it, send an email to numba-users...@continuum.io.
To post to this group, send email to numba...@continuum.io.

To view this discussion on the web visit https://groups.google.com/a/continuum.io/d/msgid/numba-users/CA%2BuhXPeqJ2mDuD2Uyv1-_kPnEwM%2BWDhi_GvxFtN%3DQ2-u8ovOTw%40mail.gmail.com.

For more options, visit https://groups.google.com/a/continuum.io/d/optout.

[Siu Kwan Lam]

Actually, I found a sample C code at: http://www.many-core.group.cam.ac.uk/projects/LBdemo.shtml and the GEM book chapter: http://http.developer.nvidia.com/GPUGems2/gpugems2_chapter47.html.   However, all these are quite old.  The cache hierarchy is very different in old CUDA arch.  I wonder how the performance is in modern hardware.

[Pei Li]

Hi, Siu. They are exactly as the same as what I found. You're right that both of them are old and, so far, I did not really run the code on my own computer. Probably, this project called Sailfish is also helpful: 

http://sailfish.us.edu.pl/#

It is python package for general LBM but written with PyCUDA. 

To unsubscribe from this group and stop receiving emails from it, send an email to numba-users+unsubscribe@continuum.io.

To post to this group, send email to numba...@continuum.io.
To view this discussion on the web visit https://groups.google.com/a/continuum.io/d/msgid/numba-users/5ef59765-ea09-48cd-93ec-93d81c52fb9f%40continuum.io.
For more options, visit https://groups.google.com/a/continuum.io/d/optout.

--

Siu Kwan Lam

Software Engineer

Continuum Analytics

--
You received this message because you are subscribed to the Google Groups "Numba Public Discussion - Public" group.

To unsubscribe from this group and stop receiving emails from it, send an email to numba-users+unsubscribe@continuum.io.

To post to this group, send email to numba...@continuum.io.

To view this discussion on the web visit https://groups.google.com/a/continuum.io/d/msgid/numba-users/CAHjo7%3DXSLNKeRcPEprVM%3DaS8-DLR7L5eMX%3DDkFB0R3QhUXZtSg%40mail.gmail.com.

For more options, visit https://groups.google.com/a/continuum.io/d/optout.

--
You received this message because you are subscribed to the Google Groups "Numba Public Discussion - Public" group.

To unsubscribe from this group and stop receiving emails from it, send an email to numba-users+unsubscribe@continuum.io.

To post to this group, send email to numba...@continuum.io.
To view this discussion on the web visit https://groups.google.com/a/continuum.io/d/msgid/numba-users/CA%2BuhXPeqJ2mDuD2Uyv1-_kPnEwM%2BWDhi_GvxFtN%3DQ2-u8ovOTw%40mail.gmail.com.
For more options, visit https://groups.google.com/a/continuum.io/d/optout.

--

Siu Kwan Lam

Software Engineer

Continuum Analytics

--

Siu Kwan Lam

Software Engineer

Continuum Analytics

--

You received this message because you are subscribed to the Google Groups "Numba Public Discussion - Public" group.

To unsubscribe from this group and stop receiving emails from it, send an email to numba-users+unsubscribe@continuum.io.

To post to this group, send email to numba...@continuum.io.

To view this discussion on the web visit https://groups.google.com/a/continuum.io/d/msgid/numba-users/CAHjo7%3DXLfGMj1z5n1hAnzsLZ6-yA1%2BoHqYrZJ_JuzoHKhgkqcQ%40mail.gmail.com.