Ah, indeed after some more debugging it turns out that the error was being caused by my `grid' argument -- the entries were floats (should have been cast to int), and were causing the exception.  For anyone interested, here's a simple working example (using only pycuda) of passing in an int:

import pycuda.driver as cuda
import pycuda.autoinit
from pycuda.compiler import SourceModule

import numpy as np

a = np.random.randn(4,4)
a = a.astype(np.float32)
a_gpu = cuda.mem_alloc(a.nbytes)
# cast the DeviceAllocation to int, and wrap as numpy type
intptr = int(a_gpu)
intptr = np.uint64(intptr)

print intptr
print hex(intptr)

cuda.memcpy_htod(a_gpu, a)

mod = SourceModule("""
  __global__ void doublify(float *a)
  {
    int idx = threadIdx.x + threadIdx.y*4;
    a[idx] *= 2;
  }
  """)

block = (4,4,1)
grid = (1,1)
func = mod.get_function("doublify")
# func(a_gpu, block=block, grid=grid)
# pass in the numpy integer
func(intptr, block=block, grid=grid)

a_doubled = np.empty_like(a)
cuda.memcpy_dtoh(a_doubled, a_gpu)

print a_doubled
print a
print a_doubled/a



On Fri, Sep 27, 2013 at 3:34 PM, Andreas Kloeckner <lists@informa.tiker.net> wrote:
Sam Preston <jsam@sci.utah.edu> writes:

> Hi all,
>
> I would like to use pycuda to write a few kernels to interoperate with a
> larger cuda/python library I'm already using.  I can get the raw device
> memory address as a python int, and from searching past threads on the
> mailing list it sounds like I should be able to pass this as an argument to
> my kernel.  However everything I try seems to generate an error -- I'm
> currently doing something like this:
>
> .
> .
> .
> myfunc = mod.get_function("myfunc")
> # myob.get() returns memory address as int
> addr = np.uint32(myob.get())
> myfunc(addr, block=block, grid=grid)

Can you post a full reproducing example? Just passing a (sized) integer
(to a kernel function) *should* work. For the rest of the PyCUDA API,
you should probably use just a plain Python int.

Andreas



--
--------------------
J. Samuel Preston
Research Assistant
Scientific Computing and Imaging Institute
University of Utah