code: def test_streamed_kernel(self):
      # this differs from the "simple_kernel" case in that *all* computation
      # and data copying is asynchronous. Observe how this necessitates the
      # use of page-locked memory.
  
      mod = SourceModule(
          """
          __global__ void multiply_them(float *dest, float *a, float *b)
          {
            const int i = threadIdx.x*blockDim.y + threadIdx.y;
            dest[i] = a[i] * b[i];
          }
          """
      )
  
      multiply_them = mod.get_function("multiply_them")
  
      shape = (32, 8)
      a = drv.pagelocked_zeros(shape, dtype=np.float32)
      b = drv.pagelocked_zeros(shape, dtype=np.float32)
      a[:] = np.random.randn(*shape)
      b[:] = np.random.randn(*shape)
  
      a_gpu = drv.mem_alloc(a.nbytes)
      b_gpu = drv.mem_alloc(b.nbytes)
  
      strm = drv.Stream()
      drv.memcpy_htod_async(a_gpu, a, strm)
      drv.memcpy_htod_async(b_gpu, b, strm)
      strm.synchronize()
  
      dest = drv.pagelocked_empty_like(a)
      multiply_them(
          drv.Out(dest), a_gpu, b_gpu, block=shape + (1,), stream=strm
      )
      strm.synchronize()
  
      drv.memcpy_dtoh_async(a, a_gpu, strm)
      drv.memcpy_dtoh_async(b, b_gpu, strm)
      strm.synchronize()
  
      assert la.norm(dest - a * b) == 0
  
imports: import numpy
  import numpy.linalg
  from pycuda.tools import mark_cuda_test, dtype_to_ctype
  import pytest
  import pycuda.gpuarray
  import pycuda.driver
  from pycuda.compiler import SourceModule
label_desc: Use image node_id 5 for calling the TestDriver obj's underlying member method code with example usage: obj.test_streamed_kernel() without return types