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
