Hi Andreas,
thanks for the quick answer! It seems that my question is somehow related to
scikits.cuda. Executing the code attached to this email sketches the issue: By
uncommenting the line "#model.fit(a_gpu, b_gpu)", the profiling output
"compiler.py:185(compile)" vanishes. The manual PyCUDA kernel seems be compiled
beforehand (during the initializing the model). The call "culinalg.dot",
however, seems to cause the compiler output after having initialized the model
...
For this toy example, there seems to be no big time difference for the second
call of "fit" (i.e., whether the first "fit" line is uncommented or not; on my
machine, it takes about 0.41 seconds in both cases). However, for the project I
am working on (which is too big to share), it makes a difference (1.35 seconds
instead of 1.662 seconds, as mentioned in my previous email).
Cheers
Fabian
Andreas Kloeckner <li...@informa.tiker.net> schrieb am 17:08 Donnerstag,
19.Februar 2015:
<fas...@yahoo.de> writes:
> Hi,
>
> is it possible to "precompile" the invoked kernels beforehand? My code makes
> use of several CUDA kernels, which are basically called within a "fit"
> function. Profiling the code with cProfile yields:
>
> 42272 function calls (42228 primitive calls) in 1.662 seconds
> ...
>
> 11 0.000 0.000 0.344 0.031 compiler.py:185(compile)
> 11 0.002 0.000 0.346 0.031 compiler.py:245(__init__)
> 4 0.000 0.000 0.317 0.079 compiler.py:33(preprocess_source)
> 11 0.000 0.000 0.342 0.031 compiler.py:66(compile_plain)
> ...
>
> Thus, about 0.344 of the 1.662 seconds are spent on compiling the
> code. When executing the function "fit" twice, the code is not
> compiled again (hence, saving these 0.344 seconds for the second call
> of "fit"). I would like to somehow precompile all involved kernels as
> soon as the object the "fit" function belongs to is initialized...
>
>
> Can one invoke the overall compilation process beforehand?
Sure! That's what the SourceModule constructor does. Just keep the
instance around.
Andreas
import sys
import math
import copy
import time
import numpy
import logging
import pycuda
import pycuda.gpuarray as gpuarray
import scikits.cuda.linalg as culinalg
from pycuda.compiler import SourceModule
import pycuda.cumath as cumath
import cProfile
class MyModel(object):
def __init__(self, device_id=0):
self.device_id = device_id
self._init_device()
def __del__(self):
self.ctx.pop()
def _init_device(self):
""" Initializes the device.
"""
# sanity check for device id
if (self.device_id < 0) or (self.device_id > pycuda.driver.Device.count() - 1):
raise Exception("Invalid device id. Maximum device id is " + str(pycuda.driver.Device.count() - 1))
# init driver and context
pycuda.driver.init()
self.device = pycuda.driver.Device(self.device_id)
self.ctx = self.device.make_context()
# init linalg for cuda
culinalg.init()
# init kernels
self._init_kernels()
def _init_kernels(self):
kernel_elementwise_multiply_mod = SourceModule("""
__global__ void elementwise_multiply(float *dest, float *a, int n, int m)
{
int tidx = threadIdx.x + blockDim.x * blockIdx.x;
int tidy = threadIdx.y + blockDim.y * blockIdx.y;
if (tidx >= n || tidy >= m){
return;
}
dest[tidx*m + tidy] = a[tidx*m + tidy]*a[tidx*m + tidy];
}
""")
self.kernel_elementwise_multiply = kernel_elementwise_multiply_mod.get_function("elementwise_multiply")
def _elementwise_mult(self, a):
n = a.shape[0]
m = a.shape[1]
dest = gpuarray.empty((n, m), numpy.float32)
blocksize = 16
block = (blocksize, blocksize, 1)
grid = (int(math.ceil(float(n) / blocksize)), int(math.ceil(float(m) / blocksize)))
self.kernel_elementwise_multiply(dest, a, numpy.int32(n), numpy.int32(m), block=block, grid=grid)
return dest
def _mult(self, a, b, transa=False, transb=False):
if transa:
transa = "T"
else:
transa = "N"
if transb:
transb = "T"
else:
transb = "N"
return culinalg.dot(a, b, transa=transa, transb=transb)
def fit(self, a_gpu, b_gpu):
test = self._mult(a_gpu, b_gpu)
test2 = self._elementwise_mult(a_gpu)
return test2 - test
# initialize model
model = MyModel(device_id=0)
# test arrays
a = numpy.random.randn(8000*8000).astype(numpy.float32).reshape((8000,8000))
b = numpy.random.randn(8000*8000).astype(numpy.float32).reshape((8000,8000))
a_gpu = gpuarray.to_gpu(a)
b_gpu = gpuarray.to_gpu(b)
# NOTE: By uncommenting this line, the profiling output "compiler.py:185(compile)" vanishes
#model.fit(a_gpu, b_gpu)
# this call is profiled
def run():
res = model.fit(a_gpu, b_gpu)
print res.get()
start = time.time()
cProfile.run("run()")
end = time.time()
print("Elapsed time: " + str(end-start))
_______________________________________________
PyCUDA mailing list
PyCUDA@tiker.net
http://lists.tiker.net/listinfo/pycuda
