class OrderedDictLevel(PriceLevel):
    def is_empty(self):
        return len(self.orders) == 0

    def __init__(self):
        PriceLevel.__init__(self)
        self.orders = OrderedDict()

    def _add(self, order):
        self.orders[order[O_ID]] = order

    def _remove(self, order: tuple) -> list:
        return self.orders.pop(order[O_ID])

    def is_not_empty(self) -> bool:
        return len(self.orders) > 0

    def __len__(self):
        return len(self.orders)

    def get_first(self) -> list:
        for order in self.orders.values():
            return order

    def _remove_first(self):
        self.orders.popitem(last=False)

    def get_last(self):
        for order in self.orders.values().__reversed__():
            return order

    def _remove_last(self):
        self.orders.popitem(last=True)
Esempio n. 2
0
class _ProcedureWorker(Worker):
    def __init__(self, cl_environment, compile_flags, cl_function, kernel_data,
                 double_precision, use_local_reduction):
        super().__init__(cl_environment)
        self._cl_function = cl_function
        self._kernel_data = OrderedDict(sorted(kernel_data.items()))
        self._double_precision = double_precision
        self._use_local_reduction = use_local_reduction

        self._mot_float_dtype = np.float32
        if double_precision:
            self._mot_float_dtype = np.float64

        for data in self._kernel_data.values():
            data.set_mot_float_dtype(self._mot_float_dtype)

        self._kernel = self._build_kernel(self._get_kernel_source(),
                                          compile_flags)

        self._workgroup_size = self._kernel.run_procedure.get_work_group_info(
            cl.kernel_work_group_info.PREFERRED_WORK_GROUP_SIZE_MULTIPLE,
            self._cl_environment.device)
        if not self._use_local_reduction:
            self._workgroup_size = 1

        self._kernel_inputs = {
            name: data.get_kernel_inputs(self._cl_context,
                                         self._workgroup_size)
            for name, data in self._kernel_data.items()
        }

    def calculate(self, range_start, range_end):
        nmr_problems = range_end - range_start

        func = self._kernel.run_procedure
        func.set_scalar_arg_dtypes(self.get_scalar_arg_dtypes())

        kernel_inputs_list = []
        for inputs in [
                self._kernel_inputs[name] for name in self._kernel_data
        ]:
            kernel_inputs_list.extend(inputs)

        func(self._cl_queue, (int(nmr_problems * self._workgroup_size), ),
             (int(self._workgroup_size), ),
             *kernel_inputs_list,
             global_offset=(int(range_start * self._workgroup_size), ))

        for name, data in self._kernel_data.items():
            data.enqueue_readouts(self._cl_queue, self._kernel_inputs[name],
                                  range_start, range_end)

    def _build_kernel(self, kernel_source, compile_flags=()):
        """Convenience function for building the kernel for this worker.

        Args:
            kernel_source (str): the kernel source to use for building the kernel

        Returns:
            cl.Program: a compiled CL kernel
        """
        from mot import configuration
        if configuration.should_ignore_kernel_compile_warnings():
            warnings.simplefilter("ignore")
        return cl.Program(self._cl_context,
                          kernel_source).build(' '.join(compile_flags))

    def _get_kernel_source(self):
        assignment = ''
        if self._cl_function.get_return_type() != 'void':
            assignment = '__results[gid] = '

        variable_inits = []
        function_call_inputs = []
        post_function_callbacks = []
        for parameter in self._cl_function.get_parameters():
            data = self._kernel_data[parameter.name]
            call_args = (parameter.name, '_' + parameter.name, 'gid',
                         parameter.data_type.address_space)

            variable_inits.append(data.initialize_variable(*call_args))
            function_call_inputs.append(
                data.get_function_call_input(*call_args))
            post_function_callbacks.append(
                data.post_function_callback(*call_args))

        kernel_source = ''
        kernel_source += get_float_type_def(self._double_precision)
        kernel_source += '\n'.join(data.get_type_definitions()
                                   for data in self._kernel_data.values())
        kernel_source += self._cl_function.get_cl_code()
        kernel_source += '''
            __kernel void run_procedure(''' + ",\n".join(self._get_kernel_arguments()) + '''){
                ulong gid = (ulong)(get_global_id(0) / get_local_size(0));
                
                ''' + '\n'.join(variable_inits) + '''     
                
                ''' + assignment + ' ' + self._cl_function.get_cl_function_name() + '(' + \
                         ', '.join(function_call_inputs) + ''');
                
                ''' + '\n'.join(post_function_callbacks) + '''
            }
        '''
        return kernel_source

    def _get_kernel_arguments(self):
        """Get the list of kernel arguments for loading the kernel data elements into the kernel.

        This will use the sorted keys for looping through the kernel input items.

        Returns:
            list of str: the list of parameter definitions
        """
        declarations = []
        for name, data in self._kernel_data.items():
            declarations.extend(data.get_kernel_parameters('_' + name))
        return declarations

    def get_scalar_arg_dtypes(self):
        """Get the location and types of the input scalars.

        Returns:
            list: for every kernel input element either None if the data is a buffer or the numpy data type if
                if is a scalar.
        """
        dtypes = []
        for name, data in self._kernel_data.items():
            dtypes.extend(data.get_scalar_arg_dtypes())
        return dtypes