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
signature.asc
Description: This is a digitally signed message part.
_______________________________________________ PyCuda mailing list [email protected] http://tiker.net/mailman/listinfo/pycuda_tiker.net
