Newly created instances of cupy.ElementwiseKernel always synchronize

[niboshi]

#725 moved the kernel cache from the global (@util.memoize) to each instance of ElementwiseKernel (ElementwiseKernel._kernel_memo). With this change, if ElementwiseKernel instance is created in each repeated call with the identical CUDA code, the loaded module is not reused and always loaded from disk, causing host-device synchronization at this line.

Example:

import cupy

h = cupy.ones([1000, 1000, 1000], 'float32')
a = cupy.array([1, 2], 'float32')
b = cupy.array([1, 2], 'float32')

for i in range(10):
    h.sum()  # dummy load on GPU

    # This call synchronizes
    cupy.ElementwiseKernel(
        'T a, T b',
        'T out',
        'out = a + b')(a, b, a)

print(a)

[niboshi]

Several functions in CuPy are affected by this (like tri, percentile, random.shuffle, random.permutation and those indirectly using them).
I consider this a bug.