On Fri, Oct 7, 2011 at 4:00 PM, Andreas Kloeckner <li...@informa.tiker.net> wrote: > On Fri, 7 Oct 2011 15:38:59 +0200, Nick Rayrider > <nick.rayri...@googlemail.com> wrote: >> On Thu, Oct 6, 2011 at 2:47 PM, Andreas Kloeckner >> <li...@informa.tiker.net> wrote: >> > On Thu, 6 Oct 2011 12:25:12 +0200, Nick Rayrider >> > <nick.rayri...@googlemail.com> wrote: >> >> Hi, >> >> >> >> first thanks for this fine piece of software. >> >> >> >> Optimizing my kernels, the nvidia's visual profiler recommended, that >> >> I should use more pinned memory. I read the PyCUDA documentation [1] >> >> and tried to understand the sparse solve example [3], but I could not >> >> make out how to turn an existing numpy array into pinned memory. I did >> >> not find further examples of PageLockedMemoryPool [2]. >> > >> > pool = PageLockedMemoryPool() >> > empty_pinned_array = pool.allocate((300, 300), np.float64) >> > >> > Now empty_pinned_array is backed by pinned storage, and when you memcpy >> > to/from it, it'll go faster. >> >> Thanks for the fast answer. I read that you shouldn't use in() as it >> performs a copy [1], so I tried following version to no avail. >> What am I missing? Problably something about how numnpy and pycuda >> handle pointers... >> >> from pycuda.tools import PageLockedMemoryPool >> pool = PageLockedMemoryPool() >> empty_pinned_array = pool.allocate(data.shape,np.float32) >> empty_pinned_array = gpuarray.to_gpu(data) >> my_kernel(empty_pinned_array,... ) >> >> [1] http://lists.tiker.net/pipermail/pycuda/2009-August/001784.html > > If you'd like to pass a GPUArray to a kernel, you need to pass > empty_pinned_array.gpudata.
I think I am totally on the wrong path. I attached what I have so far. The example compiles and 'calculates' the correct results, but the visual profiler still says that the "host mem transfer" is still pageable. #!/usr/bin/env python # -*- coding: utf-8 -*- import numpy as np from pycuda import driver, compiler, gpuarray from pycuda.tools import PageLockedMemoryPool import pycuda.autoinit pool = PageLockedMemoryPool() data = np.random.randn(5).astype(np.float32) odata = gpuarray.zeros(5, dtype=np.float32) print data empty_pinned_array = pool.allocate(data.shape,np.float32) empty_pinned_array = gpuarray.to_gpu(data) kernel_code = """ // the kernel definition __global__ void kernel(float *data, float *odata) { int index = blockIdx.x * blockDim.x + threadIdx.x ; odata[index] = data[index]+1.0f; } """ mod = compiler.SourceModule(kernel_code) kernel = mod.get_function("kernel") kernel( empty_pinned_array.gpudata,odata, grid = (20,1), block = (1, 1, 1), ) print odata.get() pool.free_held() _______________________________________________ PyCUDA mailing list PyCUDA@tiker.net http://lists.tiker.net/listinfo/pycuda