Thanks Ahmed and Andreas!

That was mighty useful! Got me going :)

@Andreas: Don't you miss the funny BAaaddisch people speak around here? I sure would ;)

Holger

Am 10.03.2009 um 20:24 schrieb Andreas Klöckner:

Hey Holger,

as a quick addition to what Ahmed said, there are texture examples in
test/test_driver.py. Those might help, too.

Andreas

PS: Hallo Karlsruhe! :)

On Dienstag 10 März 2009, Ahmed Fasih wrote:
Hi, I just started a few days ago on a similar problem (though linear
interpolation is fine for me). See the below code snippet to get a
feel for what's going on with textures, though I strongly recommend
reading the 4-5 sections on the subject scattered throughout the CUDA
Programming Guide.

This example gets more complicated if like me you want to use
complex-valued data, but I'm guessing you don't. It just demonstrates
copying a Numpy array to a texture and some linear interpolation with
normalized indexing in the kernel that writes to a standard
non-textured array in global memory.

Note that set_filter_mode() isn't mentioned in the docs.

### Start. Sorry, I just today learned about PEP-8 :(

import pycuda.driver as cuda
import pycuda.autoinit
import numpy

realrow = numpy.array([1.0, 2.0, 3.0, 4.0, 5.0],
               dtype=numpy.float32).reshape(1,5)

mod_copy_texture=cuda.SourceModule("""
texture<float, 1> tex;

__global__ void copy_texture_kernel(float * data) {
   int ty=threadIdx.y;
   data[ty] = tex1D(tex, (float)(ty)/2.0f);
}
""")

copy_texture_func = mod_copy_texture.get_function("copy_texture_kernel")
texref = mod_copy_texture.get_texref("tex")

cuda.matrix_to_texref(realrow, texref, order="C")

texref.set_flags(cuda.TRSF_NORMALIZED_COORDINATES)
texref.set_filter_mode(cuda.filter_mode.LINEAR)

gpu_output = numpy.zeros_like(realrow)
copy_texture_func(cuda.Out(gpu_output), block=(1,5,1), texrefs=[texref])

print "Output:"
print gpu_output

### End

Thanks to Andreas for sorting me out so far :), I appreciate it.

Ahmed

On Tue, Mar 10, 2009 at 12:39 PM, Holger Rapp <[email protected]> wrote:
Hey,

I discovered pycuda today (after deciding it is time to finally put my cuda card to some use) and was very thrilled to find a library so much in
my style: write code in a nice language, make sure the inner loop is
fast! I was also quite pleased how fast I got my toy example (1d cubic
spline interpolation) to run on the gpu. I then turned to my real
problem.

I need cubic interpolation of numpy arrays, so I can sample my pixels at pixel positions (x=2.345,y=pi). I used scipy.interpolate before, but now
I'm looking for a cuda implementation. I found
http://www.dannyruijters.nl/cubicinterpolation/ which seems like exactly what i want, but i was unable to get it to work with pycuda: The kernels rely on texture<> to write their results to and I have not understood how I can feed the memory of my numpy array as a texture to the kernels. I have some vague understanding of textures in cuda in general, so I think some preprocessing is needed (feeding alignement informations, how the data should be adressed and it must be transferred to the cuda device).
Does somebody have sample code using pycuda? Something simple like a
rotation kernel would be a perfect example!

Hope somebody can help!
Greetings and thanks for the great work with pycuda. I will follow it
closely.

Holger



----------------------------------------
Dipl.-Phys. Holger Rapp

Institut für Mess- und Regelungstechnik
Universität Karlsruhe (TH)
Engler-Bunte-Ring 21
76131 Karlsruhe, Germany
Geb. 40.32, Zi. 232, zweite Etage

Tel:  +49 (0)721 / 608-2341
Fax:  +49 (0)721 / 661874
Mail: [email protected]
Web:  www.mrt.uni-karlsruhe.de
----------------------------------------




_______________________________________________
PyCuda mailing list
[email protected]
http://tiker.net/mailman/listinfo/pycuda_tiker.net

_______________________________________________
PyCuda mailing list
[email protected]
http://tiker.net/mailman/listinfo/pycuda_tiker.net




_______________________________________________
PyCuda mailing list
[email protected]
http://tiker.net/mailman/listinfo/pycuda_tiker.net

Reply via email to