def setUp(self): self.dev = cupy.cuda.runtime.getDevice() assert self.dev != 1 global _test_cache_dir _test_cache_dir = tempfile.mkdtemp() os.environ['CUPY_CACHE_DIR'] = _test_cache_dir self.kern = cupy.RawKernel(_test_source1, 'test_sum', backend=self.backend) self.mod2 = cupy.RawModule(code=_test_source2, backend=self.backend) self.mod3 = cupy.RawModule(code=_test_source3, options=('-DPRECISION=2', ), backend=self.backend)
def _lombscargle(x, y, freqs, pgram, y_dot): if (pgram.dtype == 'float32'): c_type = "float" elif (pgram.dtype == 'float64'): c_type = "double" device_id = cp.cuda.Device() numSM = device_id.attributes["MultiProcessorCount"] threadsperblock = (128, ) blockspergrid = (numSM * 20, ) src = _cupy_lombscargle_src.substitute(datatype=c_type) module = cp.RawModule(code=src, options=("-std=c++11", )) kernel = module.get_function("_cupy_lombscargle") # print("Registers", kernel.num_regs) kernel_args = ( x.shape[0], freqs.shape[0], x, y, freqs, pgram, y_dot, ) kernel(blockspergrid, threadsperblock, kernel_args) cp.cuda.runtime.deviceSynchronize()
def GetModule(source, cuoptions): """Returns a cupy raw module""" if _cupy_has_RawModule(): return cp.RawModule(source, options=cuoptions) else: return cp.core.core.compile_with_cache(source, options=cuoptions, prepend_cupy_headers=False)
def compile(self, kernel_instance): """call the CUDA compiler to compile the kernel, return the device function :param kernel_name: The name of the kernel to be compiled, used to lookup the function after compilation. :type kernel_name: string :param kernel_string: The CUDA kernel code that contains the function `kernel_name` :type kernel_string: string :returns: An CUDA kernel that can be called directly. :rtype: cupy.RawKernel """ kernel_string = kernel_instance.kernel_string kernel_name = kernel_instance.name compiler_options = self.compiler_options if not any(['--std=' in opt for opt in self.compiler_options]): compiler_options = ['--std=c++11'] + self.compiler_options options = tuple(compiler_options) self.current_module = cp.RawModule(code=kernel_string, options=options, name_expressions=[kernel_name]) self.func = self.current_module.get_function(kernel_name) return self.func
def _populate_kernel_cache(): for numba_type, c_type in _SUPPORTED_TYPES.items(): # JIT compile the numba kernels, flip = 0/1 (correlate/convolve) sig = _numba_convolve_2d_signature(numba_type) _numba_kernel_cache[(0, str(numba_type))] = cuda.jit( sig, fastmath=True )(_numba_correlate_2d) _numba_kernel_cache[(1, str(numba_type))] = cuda.jit( sig, fastmath=True )(_numba_convolve_2d) # Instantiate the cupy kernel for this type and compile if isinstance(numba_type, Complex): header = "#include <cupy/complex.cuh>" else: header = "" src = loaded_from_source.substitute(datatype=c_type, header=header) module2 = cp.RawModule( code=src, options=("-std=c++11", "-use_fast_math") ) _cupy_kernel_cache[(0, str(numba_type))] = module2.get_function( "_cupy_correlate_2d" ) _cupy_kernel_cache[(1, str(numba_type))] = module2.get_function( "_cupy_convolve_2d" )
def test_invalid_compiler_flag(self): with pytest.raises(cupy.cuda.compiler.CompileException) as ex: mod = cupy.RawModule(code=_test_source3, options=('-DPRECISION=3', ), backend=self.backend) mod.get_function('test_sum') # enforce compilation assert 'precision not supported' in str(ex.value)
def main(): N = 8 module = cupy.RawModule(code=code, options=('-std=c++11', ), name_expressions=('kernel<float>', 'kernel<double>')) # The kernel computes out = A*B+C where A, B and C are 4x4 matrices. # A and B are arrays of N such matrices and C is a matrix kernel parameter. for (ctype, dtype) in zip(('float', 'double'), (numpy.float32, numpy.float64)): A = cupy.random.rand(16 * N, dtype=dtype).reshape(N, 4, 4) B = cupy.random.rand(16 * N, dtype=dtype).reshape(N, 4, 4) C = numpy.random.rand(16).astype(dtype).reshape(4, 4) out = cupy.empty_like(A) Matrix = numpy.dtype({ 'names': ['value'], 'formats': [(dtype, (4, 4))] }) kernel = module.get_function('kernel<{}>'.format(ctype)) args = (A, B, C.ravel().view(Matrix), out) kernel((1, ), (N, ), args) expected = cupy.matmul(A, B) + cupy.asarray(C[None, :, :]) cupy.testing.assert_array_almost_equal(expected, out) print("Kernel output matches expected value for " "type '{}'.".format(ctype))
def test_context_switch_RawModule6(self): # run test_template_specialization() on another device # in this test, re-compiling happens at kernel launch if self.backend == 'nvcc': self.skipTest('nvcc does not support template specialization') # compile code name_expressions = ['my_sqrt<unsigned int>'] name = name_expressions[0] with cupy.cuda.Device(0): mod = cupy.RawModule(code=test_cxx_template, options=('--std=c++11', ), name_expressions=name_expressions) # get specialized kernels ker = mod.get_function(name) # switch device with cupy.cuda.Device(1): # prepare inputs & expected outputs in_arr = cupy.testing.shaped_random((10, ), dtype=cupy.uint32) out_arr = in_arr**2 # run ker((1, ), (10, ), (in_arr, 10)) # check results assert cupy.allclose(in_arr, out_arr)
def _lombscargle(x, y, freqs, pgram, y_dot): if (str(pgram.dtype)) in _kernel_cache: kernel = _kernel_cache[(str(pgram.dtype))] else: module = cp.RawModule(code=_cupy_lombscargle_src, options=("-std=c++11", )) kernel = _kernel_cache[(str( pgram.dtype))] = module.get_function("_cupy_lombscargle_" + str(pgram.dtype)) print("Registers", kernel.num_regs) device_id = cp.cuda.Device() numSM = device_id.attributes["MultiProcessorCount"] threadsperblock = (128, ) blockspergrid = (numSM * 20, ) kernel_args = ( x.shape[0], freqs.shape[0], x, y, freqs, pgram, y_dot, ) kernel(blockspergrid, threadsperblock, kernel_args) cp.cuda.runtime.deviceSynchronize()
def _populate_kernel_cache(np_type, blocks, dim_x, dim_z, dim_u, max_tpb): # Check in np_type is a supported option if np_type not in _SUPPORTED_TYPES: raise ValueError( "Datatype {} not found for Kalman Filter".format(np_type)) if np_type == "float32": c_type = "float" else: c_type = "double" # Check CuPy version # Update to only check for v8.X in cuSignal 0.16 # Instantiate the cupy kernel for this type and compile specializations = ( "_cupy_predict<{}, {}, {}, {}, {}>".format(c_type, blocks, dim_x, dim_u, max_tpb), "_cupy_update<{}, {}, {}, {}, {}>".format(c_type, blocks, dim_x, dim_z, max_tpb), ) module = cp.RawModule( code=cuda_code_kalman, options=( "-std=c++11", "-fmad=true", ), name_expressions=specializations, ) _cupy_kernel_cache[(str(np_type), "predict")] = module.get_function(specializations[0]) _cupy_kernel_cache[(str(np_type), "update")] = module.get_function(specializations[1])
def test_template_specialization(self): if self.backend == 'nvcc': self.skipTest('nvcc does not support template specialization') # compile code name_expressions = [ 'my_sqrt<int>', 'my_sqrt<float>', 'my_sqrt<complex<double>>', 'my_func' ] mod = cupy.RawModule(code=test_cxx_template, options=('--std=c++11', ), name_expressions=name_expressions) dtypes = (cupy.int32, cupy.float32, cupy.complex128, cupy.float64) for ker_T, dtype in zip(name_expressions, dtypes): # get specialized kernels ker = mod.get_function(ker_T) # prepare inputs & expected outputs in_arr = cupy.testing.shaped_random((10, ), dtype=dtype) out_arr = in_arr**2 # run ker((1, ), (10, ), (in_arr, 10)) # check results assert cupy.allclose(in_arr, out_arr)
def test_cuFloatComplex(self): N = 100 block = 32 grid = (N + block - 1) // block dtype = cupy.complex64 mod = cupy.RawModule(code=_test_cuComplex, translate_cucomplex=True) a = cupy.random.random((N, )) + 1j * cupy.random.random((N, )) a = a.astype(dtype) b = cupy.random.random((N, )) + 1j * cupy.random.random((N, )) b = b.astype(dtype) c = cupy.random.random((N, )) + 1j * cupy.random.random((N, )) c = c.astype(dtype) out = cupy.zeros((N, ), dtype=dtype) out_float = cupy.zeros((N, ), dtype=cupy.float32) out_up = cupy.zeros((N, ), dtype=cupy.complex128) ker = mod.get_function('test_addf') ker((grid, ), (block, ), (a, b, out)) assert (out == a + b).all() ker = mod.get_function('test_subf') ker((grid, ), (block, ), (a, b, out)) assert (out == a - b).all() ker = mod.get_function('test_mulf') ker((grid, ), (block, ), (a, b, out)) assert cupy.allclose(out, a * b) ker = mod.get_function('test_divf') ker((grid, ), (block, ), (a, b, out)) assert (out == a / b).all() ker = mod.get_function('test_conjf') ker((grid, ), (block, ), (a, out)) assert (out == cupy.conj(a)).all() ker = mod.get_function('test_absf') ker((grid, ), (block, ), (a, out_float)) assert (out_float == cupy.abs(a)).all() ker = mod.get_function('test_fmaf') ker((grid, ), (block, ), (a, b, c, out)) assert cupy.allclose(out, a * b + c) ker = mod.get_function('test_makef') ker((grid, ), (block, ), (out, )) # because of precision issue, the (A==B).all() semantics would fail assert cupy.allclose(out, 1.8 - 1j * 8.7) ker = mod.get_function('test_upcast') ker((grid, ), (block, ), (a, out_up)) assert (out_up == a.astype(cupy.complex128)).all() # NumPy scalars. b = cupy.complex64(2 + 3j) ker = mod.get_function('test_addf_scalar') ker((grid, ), (block, ), (a, b, out)) assert (out == a + b).all()
def _slic_cupy( image, n_features, n_centers, max_iter, template, center_block, center_grid, image_block, image_grid, ): import cupy as cp template = 'extern "C" { ' + template + " }" module = cp.RawModule(code=template, options=("-std=c++11", )) gpu_slic_init = module.get_function("init_clusters") gpu_slic_expectation = module.get_function("expectation") gpu_slic_maximization = module.get_function("maximization") data_gpu = cp.asarray(image) centers_gpu = cp.zeros((n_centers, n_features + 3), dtype=cp.float32) labels_gpu = cp.zeros(image.shape[:3], dtype=cp.uint32) gpu_slic_init( center_grid, center_block, ( data_gpu, centers_gpu, ), ) cp.cuda.runtime.deviceSynchronize() for _ in range(max_iter): gpu_slic_expectation( image_grid, image_block, ( data_gpu, centers_gpu, labels_gpu, ), ) cp.cuda.runtime.deviceSynchronize() gpu_slic_maximization( center_grid, center_block, ( data_gpu, labels_gpu, centers_gpu, ), ) cp.cuda.runtime.deviceSynchronize() labels = np.asarray(labels_gpu.get(), dtype=np.intp) return labels
def _get_function(fatbin, func): dir = os.path.dirname(Path(__file__).parent) module = cp.RawModule( path=dir + fatbin, ) return module.get_function(func)
def load_cupy_module(fname, **kwargs): try: fname = str((Path(__file__).parent / fname).resolve()) except: pass with open(fname) as f: code = f.read() macros = ['#define {!s} {!s}'.format(k, v) for k,v in kwargs.items()] code = '\n'.join(macros + [code]) return cp.RawModule(code=code)
def test_module_load_failure(self): # in principle this test is better done in test_driver.py, but # this error is more likely to appear when using RawModule, so # let us do it here with pytest.raises(cupy.cuda.driver.CUDADriverError) as ex: cupy.RawModule(os.path.expanduser('~/this_does_not_exist.cubin'), backend=self.backend) assert 'CUDA_ERROR_FILE_NOT_FOUND' in str(ex.value)
def SART2D(p, sp, order, x0): # x0 = xinit.copy() block1D = (8, 1) grid1D = ((sp['nBins'] + block1D[0] - 1) // block1D[0], 1) block2D = (8, 8) grid2D = ((sp['nWidth'] + block2D[0] - 1) // block2D[0], (sp['nHeight'] + block2D[1] - 1) // block2D[1]) mod = cupy.RawModule(code=source_texref) fGetResiduals = mod.get_function('fGetResiduals') AssignResidualError = mod.get_function('AssignResidualError_kernel') channelDescImg = ChannelFormatDescriptor( 32, 0, 0, 0, runtime.cudaChannelFormatKindFloat) cuArrayImg = CUDAarray(channelDescImg, sp['nWidth'], sp['nHeight']) resourceDescImg = ResourceDescriptor(runtime.cudaResourceTypeArray, cuArrayImg) address_modeImg = (runtime.cudaAddressModeClamp, runtime.cudaAddressModeClamp) texDescImg = TextureDescriptor(address_modeImg, runtime.cudaFilterModePoint, runtime.cudaReadModeElementType) # 1D texture channelDesc1D = ChannelFormatDescriptor(32, 0, 0, 0, runtime.cudaChannelFormatKindFloat) cuArray1D = CUDAarray(channelDesc1D, sp['nBins']) resourceDesc1D = ResourceDescriptor(runtime.cudaResourceTypeArray, cuArray1D) address_mode1D = (runtime.cudaAddressModeClamp, runtime.cudaAddressModeClamp) texDesc1D = TextureDescriptor(address_mode1D, runtime.cudaFilterModePoint, runtime.cudaReadModeElementType) d_fResidualsData = cupy.zeros(sp['nBins'], cupy.float32) for v in range(sp['nViews']): # print('{}\n'.format(v)) nView = order[v] fLambda = sp['fRotateDir'] * 2.0 * np.pi / float( sp['nNumAngle']) * float(nView + sp['nStartAngle']) fCosLambda = np.cos(fLambda) fSinLambda = np.sin(fLambda) cuArrayImg.copy_from(x0) TextureReference(mod.get_texref('texImage'), resourceDescImg, texDescImg) getErrArgs = (d_fResidualsData, p, sp['nBins'], sp['fSod'], sp['fOdd'], sp['fCellSize'], sp['fPixelSize'], sp['fFovRadius'], fCosLambda, fSinLambda, nView, sp['fOffSet'], sp['fAngleOfSlope']) fGetResiduals(grid1D, block1D, getErrArgs) cuArray1D.copy_from(d_fResidualsData) TextureReference(mod.get_texref('texFP'), resourceDesc1D, texDesc1D) AssignResidualErrorArgs = (x0, sp['nWidth'], sp['nHeight'], sp['nBins'], sp['fSod'], sp['fOdd'], sp['fCellSize'], sp['fPixelSize'], fCosLambda, fSinLambda, sp['fOffSet'], sp['fAngleOfSlope'], sp['relax_factor']) AssignResidualError(grid2D, block2D, AssignResidualErrorArgs) return x0
def SART2DBackWard(grad, order, sp): grad_ = grad.copy() block1D = (8, 1) grid1D = ((sp['nBins'] + block1D[0] - 1) // block1D[0], 1) block2D = (8, 8) grid2D = ((sp['nWidth'] + block2D[0] - 1) // block2D[0], (sp['nHeight'] + block2D[1] - 1) // block2D[1]) mod = cupy.RawModule(code=source_texref) AssignResidualError = mod.get_function('AssignResidualError_kernel') FpKernel = mod.get_function('fGetFp_kernel') # 2D texture channelDescImg = ChannelFormatDescriptor( 32, 0, 0, 0, runtime.cudaChannelFormatKindFloat) cuArrayImg = CUDAarray(channelDescImg, sp['nWidth'], sp['nHeight']) resourceDescImg = ResourceDescriptor(runtime.cudaResourceTypeArray, cuArrayImg) address_modeImg = (runtime.cudaAddressModeClamp, runtime.cudaAddressModeClamp) texDescImg = TextureDescriptor(address_modeImg, runtime.cudaFilterModePoint, runtime.cudaReadModeElementType) # 1D texture channelDesc1D = ChannelFormatDescriptor(32, 0, 0, 0, runtime.cudaChannelFormatKindFloat) cuArray1D = CUDAarray(channelDesc1D, sp['nBins']) resourceDesc1D = ResourceDescriptor(runtime.cudaResourceTypeArray, cuArray1D) address_mode1D = (runtime.cudaAddressModeClamp, runtime.cudaAddressModeClamp) texDesc1D = TextureDescriptor(address_mode1D, runtime.cudaFilterModePoint, runtime.cudaReadModeElementType) d_fOneProj = cupy.zeros(sp['nBins'], cupy.float32) for v in range(sp['nViews']): nView = order[sp['nViews'] - 1 - v] fLambda = sp['fRotateDir'] * 2.0 * np.pi / float( sp['nNumAngle']) * float(nView + sp['nStartAngle']) fCosLambda = np.cos(fLambda) fSinLambda = np.sin(fLambda) # A*x cuArrayImg.copy_from(grad) TextureReference(mod.get_texref('texImage'), resourceDescImg, texDescImg) args = (d_fOneProj, sp['nBins'], sp['fSod'], sp['fOdd'], sp['fCellSize'], sp['fPixelSize'], sp['fFovRadius'], fCosLambda, fSinLambda, nView, sp['fOffSet'], sp['fAngleOfSlope']) FpKernel(grid1D, block1D, args) # AT*A*x cuArray1D.copy_from(d_fOneProj) TextureReference(mod.get_texref('texFP'), resourceDesc1D, texDesc1D) AssignResidualErrorArgs = (grad, sp['nWidth'], sp['nHeight'], sp['nBins'], sp['fSod'], sp['fOdd'], sp['fCellSize'], sp['fPixelSize'], fCosLambda, fSinLambda, sp['fOffSet'], sp['fAngleOfSlope'], sp['relax_factor']) AssignResidualError(grid2D, block2D, AssignResidualErrorArgs) grad = grad_ - sp['relax_factor'] * grad return grad
def test_load_ptx(self): # generate ptx in the temp dir file_path = self._generate_file('ptx') # load ptx and test the kernel mod = cupy.RawModule(path=file_path, backend=self.backend) ker = mod.get_function('test_div') x1, x2, y = self._helper(ker, cupy.float32) assert cupy.allclose(y, x1 / (x2 + 1.0))
def __init__(self, lower, upper, radius): self.lower = lower self.upper = upper self.cuda_kernel = self.compile_kernel(lower, upper) self.cuda_module = cp.RawModule( code=open("./opsi/util/cv/cuda/boxblur.cu").read()) self.apply_filter = self.cuda_module.get_function("applyFilter") self.update_radius(radius)
def __init__(self, dtype): # Load the CUDA kernel through cupy loaded_from_source = _load_cuda_kernel(dtype) median_filter_module = cp.RawModule(code=loaded_from_source) self.single_image_median_filter = median_filter_module.get_function("two_dimensional_median_filter") # Warm up the CUDA functions self._warm_up(dtype)
def setUp(self): if hasattr(self, 'clean_up'): util.clear_memo() self.dev = cupy.cuda.runtime.getDevice() assert self.dev != 1 self.temporary_cache_dir_context = use_temporary_cache_dir() self.in_memory_context = compile_in_memory(self.in_memory) self.cache_dir = self.temporary_cache_dir_context.__enter__() self.in_memory_context.__enter__() self.kern = cupy.RawKernel(_test_source1, 'test_sum', backend=self.backend) self.mod2 = cupy.RawModule(code=_test_source2, backend=self.backend) self.mod3 = cupy.RawModule(code=_test_source3, options=('-DPRECISION=2', ), backend=self.backend)
def test_invalid_compiler_flag(self): if cupy.cuda.runtime.is_hip and self.backend == 'nvrtc': self.skipTest('hiprtc does not handle #error macro properly') with pytest.raises(cupy.cuda.compiler.CompileException) as ex: mod = cupy.RawModule(code=_test_source3, options=('-DPRECISION=3',), backend=self.backend) mod.get_function('test_sum') # enforce compilation assert 'precision not supported' in str(ex.value)
def test_const_memory(self): mod = cupy.RawModule(code=test_const_mem, backend=self.backend) ker = mod.get_function('multiply_by_const') mem_ptr = mod.get_global('some_array') const_arr = cupy.ndarray((100, ), cupy.float32, mem_ptr) data = cupy.arange(100, dtype=cupy.float32) const_arr[...] = data output_arr = cupy.ones(100, dtype=cupy.float32) ker((1, ), (100, ), (output_arr, cupy.int32(100))) assert (data == output_arr).all()
def test_module_load_failure(self): # in principle this test is better done in test_driver.py, but # this error is more likely to appear when using RawModule, so # let us do it here with pytest.raises(cupy.cuda.driver.CUDADriverError) as ex: mod = cupy.RawModule( path=os.path.expanduser('~/this_does_not_exist.cubin'), backend=self.backend) mod.get_function('nonexisting_kernel') # enforce loading assert ('CUDA_ERROR_FILE_NOT_FOUND' in str(ex.value) # CUDA or 'hipErrorFileNotFound' in str(ex.value)) # HIP
def test_cuDoubleComplex(self): N = 100 block = 32 grid = (N + block - 1) // block dtype = cupy.complex128 mod = cupy.RawModule( code=_test_cuComplex, translate_cucomplex=True) a = cupy.random.random((N,)) + 1j*cupy.random.random((N,)) a = a.astype(dtype) b = cupy.random.random((N,)) + 1j*cupy.random.random((N,)) b = b.astype(dtype) c = cupy.random.random((N,)) + 1j*cupy.random.random((N,)) c = c.astype(dtype) out = cupy.zeros((N,), dtype=dtype) out_float = cupy.zeros((N,), dtype=cupy.float64) out_down = cupy.zeros((N,), dtype=cupy.complex64) ker = mod.get_function('test_add') ker((grid,), (block,), (a, b, out)) assert (out == a + b).all() ker = mod.get_function('test_sub') ker((grid,), (block,), (a, b, out)) assert (out == a - b).all() ker = mod.get_function('test_mul') ker((grid,), (block,), (a, b, out)) assert (out == a * b).all() ker = mod.get_function('test_div') ker((grid,), (block,), (a, b, out)) assert (out == a / b).all() ker = mod.get_function('test_conj') ker((grid,), (block,), (a, out)) assert (out == cupy.conj(a)).all() ker = mod.get_function('test_abs') ker((grid,), (block,), (a, out_float)) assert (out_float == cupy.abs(a)).all() ker = mod.get_function('test_fma') ker((grid,), (block,), (a, b, c, out)) assert (out == a * b + c).all() ker = mod.get_function('test_make') ker((grid,), (block,), (out,)) assert (out == 1.8 - 1j * 8.7).all() ker = mod.get_function('test_downcast') ker((grid,), (block,), (a, out_down)) assert (out_down == a.astype(cupy.complex64)).all()
def test_context_switch_RawModule4(self): # run test_load_cubin() on another device # generate cubin in the temp dir and load it on device 0 file_path = self._generate_file('cubin') mod = cupy.RawModule(path=file_path, backend=self.backend) ker = mod.get_function('test_div') # in this test, reloading happens at kernel launch cupy.cuda.runtime.setDevice(1) x1, x2, y = self._helper(ker, cupy.float32) assert cupy.allclose(y, x1 / (x2 + 1.0))
def setUp(self): global _test_cache_dir _test_cache_dir = tempfile.mkdtemp() os.environ['CUPY_CACHE_DIR'] = _test_cache_dir self.kern_grid_sync = cupy.RawKernel( _test_grid_sync, 'test_grid_sync', backend='nvcc', enable_cooperative_groups=True) self.mod_grid_sync = cupy.RawModule( code=_test_grid_sync, backend='nvcc', enable_cooperative_groups=True)
def test_compile_module(self): module = cupy.RawModule(code=_test_compile_src, backend=self.backend, options=('-DOP=+', )) log = io.StringIO() with use_temporary_cache_dir(): module.compile(log_stream=log) assert 'warning' in log.getvalue() kern = module.get_function('test_op') x1, x2, y = self._helper(kern, cupy.float32) assert cupy.allclose(y, x1 + x2)
def _helper2(self, type_str): mod2 = cupy.RawModule(code=std_code, jitify=self.jitify, name_expressions=['shift<%s>' % type_str, ], options=('--std=c++11',)) ker = mod2.get_function('shift<%s>' % type_str) N = 256 a = cupy.random.random_integers(0, 7, N).astype(cupy.int32) b = a.copy() ker((1,), (N,), (a, N)) assert cupy.allclose(a, b+100)