Hi Andreas,
aha, so I changed
...
a_gpu = ga.to_gpu(a)
a_gpu.bind_to_texref_ext(my_tex, allow_double_hack=True)
...
to
...
a_gpu = drv.matrix_to_array(a, order="C", allow_double_hack=True)
my_tex.set_array(a_gpu)
...
and now it seems to work. :-)
Thanks a lot,
Markus
Am Samstag, den 12.03.2011, 00:13 -0500 schrieb Andreas Kloeckner:
> Hi Markus,
>
> you should only use fp_tex2D if your texref is bound to a 2D array--but
> you are binding to linear memory. In this case, fp_tex1Dfetch() is the
> right thing.
>
> HTH,
> Andreas
>
> On Fri, 11 Mar 2011 17:39:19 +0100, Markus Wollgarten
> <wollgar...@helmholtz-berlin.de> wrote:
> > Dear Mailing-Listeners!
> >
> > Executing the following code (adapted from test_driver.py):
> >
> > -------------------
> > import pycuda.driver as drv
> > import pycuda.gpuarray as ga
> > from pycuda.compiler import SourceModule
> >
> > print "pycuda.VERSION "+str(pycuda.VERSION)
> > print "Compute Capapility
> > "+str(drv.Context.get_device().compute_capability())
> >
> > for tp in [numpy.float32, numpy.float64]:
> > from pycuda.tools import dtype_to_ctype
> >
> > tp_cstr = dtype_to_ctype(tp)
> > mod = SourceModule("""
> > #include <pycuda-helpers.hpp>
> >
> > texture<fp_tex_%(tp)s, 2>my_tex;
> >
> > __global__ void copy_texture(%(tp)s *dest)
> > {
> >
> > dest[threadIdx.xthreadIdx.y*8]=fp_tex2D(my_tex,threadIdx.y,threadIdx.x);
> > }
> > """ % {"tp": tp_cstr})
> >
> > copy_texture = mod.get_function("copy_texture")
> > my_tex = mod.get_texref("my_tex")
> >
> > shape = (8,2,)
> > a = numpy.random.randn(*shape).astype(tp)
> >
> > a_gpu = ga.to_gpu(a)
> > a_gpu.bind_to_texref_ext(my_tex, allow_double_hack=True)
> > blck=shape+(1,)
> >
> > dest = numpy.zeros(shape, dtype=tp)
> > g_dest=drv.to_device(dest)
> > copy_texture.prepare("P",blck,texrefs=[my_tex])
> > time=copy_texture.prepared_timed_call((1,1),g_dest)
> > dest=drv.from_device(g_dest, dest.shape, dest.dtype, order='C')
> > print a
> > print dest
> >
> > ---------
> >
> > returns:
> >
> >
> > pycuda.VERSION (0, 94, 2)
> > Compute Capapility (1, 3)
> > [[-0.92633218 0.20489018]
> > [ 1.14500916 0.23236905]
> > [ 0.43516356 0.4719891 ]
> > [-0.8008799 0.81867486]
> > [-0.20814744 -0.55152911]
> > [ 0.81224 -1.37392473]
> > [ 1.99982738 0.11174646]
> > [ 0.11471771 -1.01642931]]
> > [[-0.92633218 -0.92633218]
> > [-0.92633218 -0.92633218]
> > [-0.92633218 -0.92633218]
> > [-0.92633218 -0.92633218]
> > [ 0. 0. ]
> > [ 0. 0. ]
> > [ 0. 0. ]
> > [ 0. 0. ]]
> > [[ -6.39043527e-01 -1.95960158e-01]
> > [ 1.69915072e+00 1.16279297e+00]
> > [ -4.03001846e-01 -1.23898467e+00]
> > [ 3.62089701e-01 1.84103824e-01]
> > [ 5.73958324e-01 -3.26678644e-04]
> > [ -2.28391102e-01 1.59704601e+00]
> > [ 1.43664545e+00 -1.15527274e-01]
> > [ 4.74887599e-01 7.09184358e-01]]
> > [[-0.63904353 -0.63904353]
> > [-0.63904353 -0.63904353]
> > [-0.63904353 -0.63904353]
> > [-0.63904353 -0.63904353]
> > [ 0. 0. ]
> > [ 0. 0. ]
> > [ 0. 0. ]
> > [ 0. 0. ]]
> >
> > which is not what I expected, i.e. getting the same array back. The 1D
> > test from test_driver.py works but for 2D I probably did something
> > wrong. However, I have no clue what and would appreciate your help very
> > much.
> >
> > Best wishes,
> >
> > Markus
> >
> >
> > _______________________________________________
> > PyCUDA mailing list
> > PyCUDA@tiker.net
> > http://lists.tiker.net/listinfo/pycuda
> >
>
_______________________________________________
PyCUDA mailing list
PyCUDA@tiker.net
http://lists.tiker.net/listinfo/pycuda
