<commit_before>#!/usr/bin/env python
""" A simple CUDA kernel tuner in Python

The goal of this project is to provide a - as simple as possible - tool
for tuning CUDA kernels. This implies that any CUDA kernel can be tuned
without requiring extensive changes to the original kernel code.

A very common problem in CUDA programming is that some combination of
thread block dimensions and other kernel parameters, like tiling or
unrolling factors, results in dramatically better performance than other
kernel configurations. The goal of auto-tuning is to automate the
process of finding the best performing configuration for a given device.

This kernel tuner aims that you can directly use the tuned kernels
without introducing any new dependencies. The tuned kernels can
afterwards be used independently of the programming environment, whether
that is using C/C++/Java/Fortran or Python doesn't matter.

This module currently only contains one function which is called
tune_kernel() to which you pass at least the kernel name, a string
containing the kernel code, the problem size, a list of kernel function
arguments, and a dictionary of tunable parameters. There are also a lot
of optional parameters, for a full list see the documentation of
tune_kernel().

Example usage:
    See the bottom of this file.

Author:
    <NAME> <<EMAIL>>

Copyright and License:
    Copyright 2014 Netherlands eScience Center

    Licensed under the Apache License, Version 2.0 (the "License");
    you may not use this file except in compliance with the License.
    You may obtain a copy of the License at

        http://www.apache.org/licenses/LICENSE-2.0

    Unless required by applicable law or agreed to in writing, software
    distributed under the License is distributed on an "AS IS" BASIS,
    WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
    See the License for the specific language governing permissions and
    limitations under the License.
"""

import pycuda.driver as drv
from pycuda.autoinit import context
from pycuda.compiler import SourceModule
import numpy
import itertools


def tune_kernel(kernel_name, kernel_string, problem_size, arguments,
        tune_params, cc=52, grid_div_x=["block_size_x"], grid_div_y=None):
    """ Tune a CUDA kernel given a set of tunable parameters

    Args:
        kernel_name: A string containing the kernel name
        kernel_string: A string containing the CUDA kernel code
        problem_size: A tuple containing the size from which the grid
            dimensions of the kernel will be computed
        arguments: A list of kernel arguments, use numpy arrays for arrays
        tune_params: A dictionary containing the parameter names as keys
            and lists of possible parameter settings as values.
            Currently the kernel tuner uses the convention that the following
            list of tunable parameters are used as compile-time constants
            in the code:
                block_size_x    thread block size x-dimension
                block_size_y    thread block size y-dimension
                block_size_z    thread block size z-dimension
            Options for changing these defaults will be added later.

        cc: compute capability of the CUDA device, 52 by default.
            Could be changed to detect this at runtime.
        grid_div_x: A list of names of the parameters whose values divide
            the grid dimensions in the x-direction, ["block_size_x"] by default
        grid_div_y: A list of names of the parameters whose values divide
            the grid dimensions in the y-direction, empty by default

    Returns:
        nothing for the moment it just prints a lot of stuff

    """

    original_kernel = kernel_string
    results = dict()

    #move data to GPU
    gpu_args = _create_gpu_args(arguments)

    #compute cartesian product of all tunable parameters
    for element in itertools.product(*tune_params.values()):
        params = dict(zip(tune_params.keys(), element))
        instance_string = "_".join([str(i) for i in params.values()])

        #compute thread block and grid dimensions for this kernel
        threads = _get_thread_block_dimensions(params)
        if numpy.prod(threads) > 1024:
            print "skipping config", instance_string, "reason: too many threads per block"
            continue
        grid = _get_grid_dimensions(problem_size, params,
                       grid_div_y, grid_div_x)

        #replace occurrences of the tuning parameters with their current value
        kernel_string = original_kernel
        for k, v in params.iteritems():
            kernel_string = kernel_string.replace(k, str(v))

        #rename the kernel to guarantee that PyCuda compiles a new kernel
        name = kernel_name + "_" + instance_string
        kernel_string = kernel_string.replace(kernel_name, name)

        #compile kernel func
        try:
            func = SourceModule(kernel_string, options=['-Xcompiler=-Wall'],
                    arch='compute_' + str(cc), code='sm_' + str(cc),
                    cache_dir=False).get_function(name)
        except drv.CompileError, e:
            #compiles may fail because certain kernel configurations use too
            #much shared memory for example, the desired behavior is to simply
            #skip over this configuration and try the next one
            if "uses too much shared data" in e.stderr:
                print "skipping config", instance_string, "reason: too much shared memory used"
                continue
            else:
                raise e

        #test kernel
        start = drv.Event()
        end = drv.Event()

        context.synchronize()
        start.record()
        try:
            func( *gpu_args, block=threads, grid=grid)
        except drv.LaunchError, e:
            #some launches may fail because too many registers are required
            #to run the kernel given the current thread block size
            #the desired behavior is to simply skip over this configuration
            #and proceed to try the next one
            if "too many resources requested for launch" in str(e):
                print "skipping config", instance_string, "reason: too many resources requested for launch"
                continue
            else:
                raise e
        end.record()

        context.synchronize()
        time = end.time_since(start)

        #print the result
        #later we'll save the results and return nice statistics
        print params, kernel_name, "took:", time, " ms."
        results[instance_string] = time


    #finished iterating over search space<commit_msg>added reporting of the best kernel configuration<commit_after>#!/usr/bin/env python
""" A simple CUDA kernel tuner in Python

The goal of this project is to provide a - as simple as possible - tool
for tuning CUDA kernels. This implies that any CUDA kernel can be tuned
without requiring extensive changes to the original kernel code.

A very common problem in CUDA programming is that some combination of
thread block dimensions and other kernel parameters, like tiling or
unrolling factors, results in dramatically better performance than other
kernel configurations. The goal of auto-tuning is to automate the
process of finding the best performing configuration for a given device.

This kernel tuner aims that you can directly use the tuned kernels
without introducing any new dependencies. The tuned kernels can
afterwards be used independently of the programming environment, whether
that is using C/C++/Java/Fortran or Python doesn't matter.

This module currently only contains one function which is called
tune_kernel() to which you pass at least the kernel name, a string
containing the kernel code, the problem size, a list of kernel function
arguments, and a dictionary of tunable parameters. There are also a lot
of optional parameters, for a full list see the documentation of
tune_kernel().

Example usage:
    See the bottom of this file.

Author:
    <NAME> <<EMAIL>>

Copyright and License:
    Copyright 2014 Netherlands eScience Center

    Licensed under the Apache License, Version 2.0 (the "License");
    you may not use this file except in compliance with the License.
    You may obtain a copy of the License at

        http://www.apache.org/licenses/LICENSE-2.0

    Unless required by applicable law or agreed to in writing, software
    distributed under the License is distributed on an "AS IS" BASIS,
    WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
    See the License for the specific language governing permissions and
    limitations under the License.
"""

import pycuda.driver as drv
from pycuda.autoinit import context
from pycuda.compiler import SourceModule
import numpy
import itertools


def tune_kernel(kernel_name, kernel_string, problem_size, arguments,
        tune_params, cc=52, grid_div_x=["block_size_x"], grid_div_y=None):
    """ Tune a CUDA kernel given a set of tunable parameters

    Args:
        kernel_name: A string containing the kernel name
        kernel_string: A string containing the CUDA kernel code
        problem_size: A tuple containing the size from which the grid
            dimensions of the kernel will be computed
        arguments: A list of kernel arguments, use numpy arrays for arrays
        tune_params: A dictionary containing the parameter names as keys
            and lists of possible parameter settings as values.
            Currently the kernel tuner uses the convention that the following
            list of tunable parameters are used as compile-time constants
            in the code:
                block_size_x    thread block size x-dimension
                block_size_y    thread block size y-dimension
                block_size_z    thread block size z-dimension
            Options for changing these defaults will be added later.

        cc: compute capability of the CUDA device, 52 by default.
            Could be changed to detect this at runtime.
        grid_div_x: A list of names of the parameters whose values divide
            the grid dimensions in the x-direction, ["block_size_x"] by default
        grid_div_y: A list of names of the parameters whose values divide
            the grid dimensions in the y-direction, empty by default

    Returns:
        nothing for the moment it just prints a lot of stuff

    """

    original_kernel = kernel_string
    results = dict()

    #move data to GPU
    gpu_args = _create_gpu_args(arguments)

    #compute cartesian product of all tunable parameters
    for element in itertools.product(*tune_params.values()):
        params = dict(zip(tune_params.keys(), element))
        instance_string = "_".join([str(i) for i in params.values()])

        #compute thread block and grid dimensions for this kernel
        threads = _get_thread_block_dimensions(params)
        if numpy.prod(threads) > 1024:
            print "skipping config", instance_string, "reason: too many threads per block"
            continue
        grid = _get_grid_dimensions(problem_size, params,
                       grid_div_y, grid_div_x)

        #replace occurrences of the tuning parameters with their current value
        kernel_string = original_kernel
        for k, v in params.iteritems():
            kernel_string = kernel_string.replace(k, str(v))

        #rename the kernel to guarantee that PyCuda compiles a new kernel
        name = kernel_name + "_" + instance_string
        kernel_string = kernel_string.replace(kernel_name, name)

        #compile kernel func
        try:
            func = SourceModule(kernel_string, options=['-Xcompiler=-Wall'],
                    arch='compute_' + str(cc), code='sm_' + str(cc),
                    cache_dir=False).get_function(name)
        except drv.CompileError, e:
            #compiles may fail because certain kernel configurations use too
            #much shared memory for example, the desired behavior is to simply
            #skip over this configuration and try the next one
            if "uses too much shared data" in e.stderr:
                print "skipping config", instance_string, "reason: too much shared memory used"
                continue
            else:
                raise e

        #test kernel
        start = drv.Event()
        end = drv.Event()

        context.synchronize()
        start.record()
        try:
            func( *gpu_args, block=threads, grid=grid)
        except drv.LaunchError, e:
            #some launches may fail because too many registers are required
            #to run the kernel given the current thread block size
            #the desired behavior is to simply skip over this configuration
            #and proceed to try the next one
            if "too many resources requested for launch" in str(e):
                print "skipping config", instance_string, "reason: too many resources requested for launch"
                continue
            else:
                raise e
        end.record()

        context.synchronize()
        time = end.time_since(start)

        #print the result
        #later we'll save the results and return nice statistics
        print params, kernel_name, "took:", time, " ms."
        results[instance_string] = time

    #finished iterating over search space
    best_config = min(results, key=results.get)
    print "best performing configuration: ", best_config, "took:", results[best_config], "ms."
{'fraction_non_alphanumeric': 0.05183760999778156, 'fraction_numerical': 0.002366338830141241, 'mean_word_length': 2.6880283610580857, 'pattern_counts': {'":': 0, '<': 9, '<?xml version=': 0, '>': 11, 'https://': 0, 'lorem ipsum': 0, 'www.': 2, 'xml': 0}, 'pii_count': 0, 'substrings_counts': 0, 'word_list_counts': {'cursed_substrings.json': 0, 'profanity_word_list.json': 0, 'sexual_word_list.json': 0, 'zh_pornsignals.json': 0}}
{'dir': 'git-commits-cleaned', 'id': '1016505', 'n_tokens_mistral': 3342, 'n_tokens_neox': 3195, 'n_words': 1558}