On Thu, Nov 1, 2012 at 12:07 AM, Ahmed Fasih <[email protected]> wrote:
> Hi PyCUDA community and author, thanks for your continued fantastic
> support. Today I am running into an unexpected problem when I try to bind a
> big chunk of memory to a texture. I've reduced the problem to the following
> stand-alone illustrative example, but I apologize, because the example
> contains some specifics about my usecase (float2 data, 2D texture---the end
> use of all this is to do matrix filtering):
>
> ### code starts
> import pycuda.driver as cuda
> import pycuda.autoinit
> from pycuda.compiler import SourceModule
>
> # borrowed from PyCUDA test
> mod = SourceModule("""
> #define CHANNELS 2
> texture<float2, 2, cudaReadModeElementType> mtx_tex;
> __global__ void copy_texture(float *dest)
> {
> int row = threadIdx.x;
> int col = threadIdx.y;
> int w = blockDim.y;
> float2 texval = tex2D(mtx_tex, row, col);
> dest[(row*w+col)*CHANNELS + 0] = texval.x;
> dest[(row*w+col)*CHANNELS + 1] = texval.y;
> }
> """)
> mtx_tex = mod.get_texref("mtx_tex")
>
> width = 32*1024
> height = 4330
> nbytes = width * height * 2 * 4 # 2 floats per float2, 4 bytes per float
> gpu = cuda.mem_alloc(nbytes)
> HACK = 3
> mtx_tex.set_address(gpu, nbytes / HACK)
>
> ### code ends
>
> The above code runs on my Telsa C2050, for values of "HACK" >= 3, that is,
> asking set_address() to bind only one-third (or less) of the device memory
> needed into the texture unit. I need it to work for HACK=1, but I get an
> error:
>
> ### error begins
> Traceback (most recent call last):
> File "texbind.py", line 26, in <module>
> mtx_tex.set_address(gpu, nbytes / HACK)
> pycuda._driver.LogicError: cuTexRefSetAddress failed: invalid value
> ### error ends
>
>
Sorry for the extra email. I see on [1] that there is a limit to the width
of a 1D linear memory texture reference, of 2**27, which is just a bit less
than my width*height = 2**27.08. I'll try resolving this issue by
using set_address_2d() instead of set_address(), and will post an update
tomorrow.
Again, apologies for additionally burdening your inboxes,
Ahmed
[1] http://en.wikipedia.org/wiki/CUDA#Version_features_and_specifications
_______________________________________________
PyCUDA mailing list
[email protected]
http://lists.tiker.net/listinfo/pycuda
