So I recently discovered the very cool
pycuda.driver.make_multichannel_2d_array() and am having some trouble
with it. As always, I appreciate any inputs :) I ordinarily would poke
around on my own but I'm hoping to do a final class project as much in
PyCUDA as possible, so what the heck :D
Here's a code snippet that describes what I'm doing:
### Start code
import pycuda.driver as cuda
import pycuda.autoinit
import numpy
realrow = 2.0 + numpy.array([1.0, 2.0, 3.0, 4.0, 5.0]).reshape(1,5)
row_quad = numpy.asarray(numpy.dstack((realrow, realrow*10,
realrow*100, realrow*1000)), dtype=numpy.float32) # Stack realrow 4
times.
# Note: row_quad.shape returns (1, 5, 4)
mod_copy_texture=cuda.SourceModule("""
texture<float4, 1> tex;
__global__ void four_to_one(float * data) {
float4 texel = tex1D(tex, 0.5f + (float)(threadIdx.x+threadIdx.y) * 0.5f);
data[ty] = texel.x + texel.y + texel.z + texel.w;
}
""")
four_to_one_kernel = mod_copy_texture.get_function("four_to_one")
texref = mod_copy_texture.get_texref("tex")
ary = cuda.make_multichannel_2d_array(row_quad, "C")
texref.set_array(ary)
texref.set_filter_mode(cuda.filter_mode.LINEAR)
gpu_output = numpy.zeros((1,5), dtype=numpy.float32)
four_to_one_kernel(cuda.Out(gpu_output), block=(1,5,1), texrefs=[texref])
print gpu_output
### End code snippet
This produces the following output:
[[ 3. 1.5 0. 0. 0. ]]
Indicating that at least the first value of row_quad, 3, is getting
read into the texture's first channel but all other elements/channels
are 0s (the 1.5 is because I'm indexing every 0.5 elements).
Any suggestions why only the very first float is getting picked up by
this 1d float4 texture? Thanks! I'm having a great time with PyCUDA,
though it is a bit mind-bending to have Vim open to a python and a C
file, and editing both o_O.
On Fri, Mar 13, 2009 at 5:20 PM, Ahmed Fasih <[email protected]> wrote:
> Greetings. I see that float2 textures in CUDA are seamlessly filled by
> complex64 Numpy arrays, but what Python representation should an array
> have if I want to load it as a float4 texture? (I want a 1d texture.)
>
> Would I create a 4Nx1 float32 array in Python in C (not Fortran)
> ordering and use the ArrayDescriptor class with num_channels=4 and
> then the Memcpy2D class ... ? doesn't sound right (I'll try it
> though), so thanks for any advice!
>
> Ahmed
>
_______________________________________________
PyCuda mailing list
[email protected]
http://tiker.net/mailman/listinfo/pycuda_tiker.net
