def env():
kernel_string = """__global__ void vector_add(float *c, float *a, float *b, int n) {
int i = blockIdx.x * block_size_x + threadIdx.x;
if (i<n) {
c[i] = a[i] + b[i];
}
} """
args = get_vector_add_args()
params = {"block_size_x": 128}
lang = "CUDA"
kernel_source = core.KernelSource(kernel_string, lang)
verbose = True
kernel_options = Options(kernel_name="vector_add",
kernel_string=kernel_string,
problem_size=args[-1],
arguments=args,
lang=lang,
grid_div_x=None,
grid_div_y=None,
grid_div_z=None,
cmem_args=None,
texmem_args=None,
block_size_names=None)
device_options = Options(device=0,
platform=0,
lang=lang,
quiet=False,
compiler=None,
compiler_options=None)
dev = core.DeviceInterface(kernel_source, iterations=7, **device_options)
instance = dev.create_kernel_instance(kernel_source, kernel_options,
params, verbose)
return dev, instance
def __init__(self, kernel_name, kernel_string, problem_size, arguments, params, inputs=None, outputs=None, device=0, platform=0,
block_size_names=None, grid_div_x=None, grid_div_y=None, grid_div_z=None, verbose=True, lang=None):
""" Construct Python helper object to compile and call the kernel from Python
This object compiles a GPU kernel parameterized using the parameters in params.
GPU memory is allocated for each argument using its size and type as listed in arguments.
The object can be called directly as a function with the kernel arguments as function arguments.
Kernel arguments marked as inputs will be copied to the GPU on every kernel launch.
Only the kernel arguments marked as outputs will be returned, note that the result is always
returned in a list, even when there is only one output.
Most of the arguments to this function are the same as with tune_kernel or run_kernel in Kernel Tuner,
and are therefore not duplicated here. The two new arguments are:
:param inputs: a boolean list of length arguments to signal whether an argument is input to the kernel
:type inputs: list(bool)
:param outputs: a boolean list of length arguments to signal whether an argument is output of the kernel
:type outputs: list(bool)
"""
#construct device interface
kernel_source = core.KernelSource(kernel_string, lang)
self.dev = core.DeviceInterface(kernel_source, device=device)
#construct kernel_options to hold information about the kernel
opts = locals()
kernel_options = Options([(k, opts[k]) for k in _kernel_options.keys() if k in opts.keys()])
#instantiate the kernel given the parameters in params
self.kernel_instance = self.dev.create_kernel_instance(kernel_source, kernel_options, params, verbose)
#compile the kernel
self.func = self.dev.compile_kernel(self.kernel_instance, verbose)
#setup GPU memory
self.gpu_args = self.dev.ready_argument_list(arguments)
if inputs:
self.inputs = inputs
else:
self.inputs = [True for _ in arguments]
if outputs:
self.outputs = outputs
else:
self.outputs = [True for _ in arguments]
def test_check_kernel_output(dev_func_interface):
dev_func_interface.configure_mock(**mock_config)
dev = core.DeviceInterface(core.KernelSource("", lang="CUDA"))
dfi = dev.dev
answer = [np.zeros(4).astype(np.float32)]
instance = core.KernelInstance("name", None, "kernel_string", "temp_files",
(256, 1, 1), (1, 1, 1), {}, answer)
wrong = [np.array([1, 2, 3, 4]).astype(np.float32)]
atol = 1e-6
test = dev.check_kernel_output('func', answer, instance, answer, atol,
None, True)
dfi.memcpy_htod.assert_called_once_with(answer[0], answer[0])
dfi.run_kernel.assert_called_once_with('func', answer, (256, 1, 1),
(1, 1, 1))
print(dfi.mock_calls)
assert dfi.memcpy_dtoh.called == 1
for name, args, _ in dfi.mock_calls:
if name == 'memcpy_dtoh':
assert all(args[0] == answer[0])
assert all(args[1] == answer[0])
assert test
#the following call to check_kernel_output is expected to fail because
#the answer is non-zero, while the memcpy_dtoh function on the Mocked object
#obviously does not result in the result_host array containing anything
#non-zero
try:
dev.check_kernel_output('func', wrong, instance, wrong, atol, None,
True)
print("check_kernel_output failed to throw an exception")
assert False
except Exception:
assert True
def run_kernel(kernel_name, kernel_string, problem_size, arguments,
params, grid_div_x=None, grid_div_y=None, grid_div_z=None,
lang=None, device=0, platform=0, cmem_args=None, texmem_args=None, compiler=None, compiler_options=None,
block_size_names=None, quiet=False, log=None):
if log:
logging.basicConfig(filename=kernel_name + datetime.now().strftime('%Y%m%d-%H:%M:%S') + '.log', level=log)
kernel_source = core.KernelSource(kernel_string, lang)
_check_user_input(kernel_name, kernel_source, arguments, block_size_names)
#sort options into separate dicts
opts = locals()
kernel_options = Options([(k, opts[k]) for k in _kernel_options.keys()])
device_options = Options([(k, opts[k]) for k in _device_options.keys()])
#detect language and create the right device function interface
dev = core.DeviceInterface(kernel_source, iterations=1, **device_options)
#move data to the GPU
gpu_args = dev.ready_argument_list(arguments)
instance = None
try:
#create kernel instance
instance = dev.create_kernel_instance(kernel_source, kernel_options, params, False)
if instance is None:
raise Exception("cannot create kernel instance, too many threads per block")
# see if the kernel arguments have correct type
util.check_argument_list(instance.name, instance.kernel_string, arguments)
#compile the kernel
func = dev.compile_kernel(instance, False)
if func is None:
raise Exception("cannot compile kernel, too much shared memory used")
#add constant memory arguments to compiled module
if cmem_args is not None:
dev.copy_constant_memory_args(cmem_args)
#add texture memory arguments to compiled module
if texmem_args is not None:
dev.copy_texture_memory_args(texmem_args)
finally:
#delete temp files
if instance is not None:
instance.delete_temp_files()
#run the kernel
if not dev.run_kernel(func, gpu_args, instance):
raise Exception("runtime error occured, too many resources requested")
#copy data in GPU memory back to the host
results = []
for i, arg in enumerate(arguments):
if numpy.isscalar(arg):
results.append(arg)
else:
results.append(numpy.zeros_like(arg))
dev.memcpy_dtoh(results[-1], gpu_args[i])
#trying to make run_kernel work nicely with the Nvidia Visual Profiler
del dev
return results
def tune_kernel(kernel_name, kernel_string, problem_size, arguments,
tune_params, grid_div_x=None, grid_div_y=None, grid_div_z=None,
restrictions=None, answer=None, atol=1e-6, verify=None, verbose=False,
lang=None, device=0, platform=0, cmem_args=None, texmem_args=None,
compiler=None, compiler_options=None, log=None,
iterations=7, block_size_names=None, quiet=False, strategy=None, strategy_options=None,
cache=None):
if log:
logging.basicConfig(filename=kernel_name + datetime.now().strftime('%Y%m%d-%H:%M:%S') + '.log', level=log)
kernel_source = core.KernelSource(kernel_string, lang)
_check_user_input(kernel_name, kernel_source, arguments, block_size_names)
# check for forbidden names in tune parameters
util.check_tune_params_list(tune_params)
# check whether block_size_names are used as expected
util.check_block_size_params_names_list(block_size_names, tune_params)
if iterations < 1:
raise ValueError("Iterations should be at least one!")
#sort all the options into separate dicts
opts = locals()
kernel_options = Options([(k, opts[k]) for k in _kernel_options.keys()])
tuning_options = Options([(k, opts[k]) for k in _tuning_options.keys()])
device_options = Options([(k, opts[k]) for k in _device_options.keys()])
logging.debug('tune_kernel called')
logging.debug('kernel_options: %s', util.get_config_string(kernel_options))
logging.debug('tuning_options: %s', util.get_config_string(tuning_options))
logging.debug('device_options: %s', util.get_config_string(device_options))
if strategy:
if strategy in strategy_map:
strategy = strategy_map[strategy]
else:
raise ValueError("Strategy %s not recognized" % strategy)
#make strategy_options into an Options object
if tuning_options.strategy_options:
if not isinstance(strategy_options, Options):
tuning_options.strategy_options = Options(strategy_options)
#select strategy based on user options
if "fraction" in tuning_options.strategy_options and not tuning_options.strategy == 'random_sample':
raise ValueError('It is not possible to use fraction in combination with strategies other than "random_sample". ' \
'Please set strategy="random_sample", when using "fraction" in strategy_options')
#check if method is supported by the selected strategy
if "method" in tuning_options.strategy_options:
method = tuning_options.strategy_options.method
if not method in strategy.supported_methods:
raise ValueError('Method %s is not supported for strategy %s' % (method, tuning_options.strategy))
#if no strategy_options dict has been passed, create empty dictionary
else:
tuning_options.strategy_options = Options({})
#if no strategy selected
else:
strategy = brute_force
runner = SequentialRunner(kernel_source, kernel_options, device_options, iterations)
#the user-specified function may or may not have an optional atol argument;
#we normalize it so that it always accepts atol.
tuning_options.verify = util.normalize_verify_function(tuning_options.verify)
#process cache
if cache:
if cache[-5:] != ".json":
cache += ".json"
util.process_cache(cache, kernel_options, tuning_options, runner)
else:
tuning_options.cache = {}
tuning_options.cachefile = None
#call the strategy to execute the tuning process
results, env = strategy.tune(runner, kernel_options, device_options, tuning_options)
#finished iterating over search space
if not device_options.quiet:
if results: #checks if results is not empty
best_config = min(results, key=lambda x: x['time'])
units = getattr(runner, "units", None)
print("best performing configuration:", util.get_config_string(best_config, list(tune_params.keys()) + ['time'], units=units))
else:
print("no results to report")
if cache:
util.close_cache(cache)
del runner.dev
return results, env
def test_get_device_interface2():
lang = "OpenCL"
dev = core.DeviceInterface(core.KernelSource("", lang=lang))
assert isinstance(dev, core.DeviceInterface)
assert isinstance(dev.dev, opencl.OpenCLFunctions)
def test_get_device_interface1():
lang = "CUDA"
dev = core.DeviceInterface(core.KernelSource("", lang=lang))
assert isinstance(dev, core.DeviceInterface)
assert isinstance(dev.dev, cuda.CudaFunctions)
def test_get_device_interface1():
lang = "CUDA"
with core.DeviceInterface(core.KernelSource("", "", lang=lang)) as dev:
assert isinstance(dev, core.DeviceInterface)
assert isinstance(dev.dev, pycuda.PyCudaFunctions)
