Re: [PyCUDA] Problem binding large DeviceAllocations to textures

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