o io@sddlZddlZddlZddlmZmZmZddlZddl m Z ddl m Z e r.ddl m Z ddlmZmZmZmZmZmZddlmZmZmZmZmZdd l mZGd d d ejZGd d d ejZeZeddeiZ iZ!GdddZ"Gdddej#Z$de%de%fddZ&ed$de%de'defddZ(ede%defddZ)ede%defd d!Z*ede%defd"d#Z+dS)%N)CallableTupleList)_TemplateJITPlugin)IS_AOT_ENABLED)_TemplateAOTPlugin)_parse_register_inputs_parse_register_return_validate_autotune_validate_impl_validate_aot_impl_validate_name_and_namespace)_built_in_to_plugin_field_type _join_with_numpy_to_plugin_field_type_is_numpy_array_infer_numpy_type) public_apics4eZdZfddZddZddZddZZS) _PluginNamespacecstd|||_dS)Nztensorrt.plugin.op.)super__init__ _namespace)self namespace __class__Q/home/ubuntu/.local/lib/python3.10/site-packages/tensorrt_bindings/plugin/_lib.pyr0s z_PluginNamespace.__init__cCst||rJt|||dSN)hasattrsetattr)rname plugin_defrrrdefine4sz_PluginNamespace.definecCs"td|jjd|jd|d)N'z ' object 'z' has no attribute ')AttributeErrorr__name__rrr"rrr __getattr__8sz_PluginNamespace.__getattr__cCsd|jdS)Nz_PluginNamespace(namespace="z"))rrrrr__repr__=sz_PluginNamespace.__repr__)r' __module__ __qualname__rr$r)r+ __classcell__rrrrr/s  rcs,eZdZfddZddZddZZS)_OpcstddS)Nztensorrt.plugin.op)rrr*rrrrCsz _Op.__init__cCs,t||r t||St|}t||||Sr)r getattrrr!)rrnsrrr define_or_getFs   z_Op.define_or_getcCstd|jjd|d)Nr%z' object has no attribute ')r&rr'r(rrrr)Osz_Op.__getattr__)r'r,r-rr2r)r.rrrrr/Bs  r/op)symbolc@s:eZdZddZdeeejeejejffddZ dS) PluginDefcCsRd|_d|_d|_d|_d|_d|_d|_d|_d|_d|_ d|_ d|_ d|_ dSr) plugin_id register_func impl_func aot_impl_func autotune_funcautotune_attr_namesinput_tensor_names input_attrsimpl_attr_namesaot_impl_attr_names num_outputsinput_arg_schemaexpects_tacticr*rrrr]s zPluginDef.__init__returnc s |jd\gg}|D]}t|tjstdt||d|qi}|D]s\}}||j vrGtd|d|j d|||<|j |}t|t j rt |t j krt t |dd} t | t |jkrtd t |jd |d | d n|t|urtd t|d |d |d ||q0dtdgt|jt|j dd} dd|d} | | krtd| d| d|jtvrtdn'i} |D]\}}t|t j rd|jf| |<qdt|f| |<qt| g|D]r\}}t|t j rDt |j} | t jkr6t||tjjqt||t| qt|trZt|| tjj!qt|t"rnt||tjjqt|t #|gt$t|qddfdd }|S)N::z+Expected trt.ITensor but got input of type ITensorUnexpected attribute z provided. Expected one of .rrzUnexpected dtype 'z' for attribute 'z '. Expected 'z'.zUnexpected type '(T)z, zUnexpected schema z received. Expected 1Fquick_plugin_creation_requesttrt.QuickPluginCreationRequestcsJ|durttjj}nttjj|}g|fSr) create_plugintrtPluginFieldCollection TensorRTPhaseBUILD)rKplgfields input_tensorsr"r plg_creatorrrcreate_plugin_instances z2PluginDef.__call__..create_plugin_instancer)rKrL)%r6split isinstancerNrE ValueErrortypeappenditemsr=keysnpndarraytyping get_originget_argsdtyperlenr<join QDP_CREATORSget_plugin_registry get_creator_register_plugin_creatorfloat16 PluginFieldtobytesPluginFieldTypeUNKNOWNrstrencodeCHARbytesarrayr)rargskwargs schema_chunkstattrskeyvalueattr_annotationnp_dtypeexpected_schemaschema attrs_typesnp_typerWrrSr__call__ls                   zPluginDef.__call__N) r'r,r-rrrrNrE IPluginV3rrrrrr5\s r5c@s eZdZddZdd ddZdS) _TemplatePluginCreatorc Cstj|||_||_d|_g}|D]L\}\}}|rO|tur.|t |dtj j q|t ur?|t |dtj j q|t |tgt|q|t |tgtt|qt||_dS)NrJ)rNIPluginCreatorV3Quickrr"plugin_namespaceplugin_versionr]rpr\rlrnrrrsror_rtrrrdrO field_names)rr"rryrbuiltintype_rrrrs4 z_TemplatePluginCreator.__init__NqpcrrLc Csft|d|}|}|}i}|D]\}|j|jvr(td|jd|jd|j|j} t| rTt| } | tjkrJtj |j tjd||j<q|j | ||j<qt | tre|j d||j<q| |j ||j<qd} |durt|||j} | |j||j|j|j|j|j| S|tjjkr|jdurtd|jdd } nS|tjjkr|jdurtd |jd d } n<|jdu} |jdu}| r|r|tjj krd } n%|tjj!krd } ntd |jd| rd } n|rd } n td |jd| dusJ| rt|||j} | |j||j|j|j|j|j| St"|||j} | |j||j#|j|j|j| S)NrDrFz, provided to create_plugin. Expected one of rG)rdzutf-8z3AOT implementation requested, but not defined for 'z$'. Was @trt.plugin.aot_impl defined?Fz3JIT implementation requested, but not defined for 'z '. Was @trt.plugin.impl defined?TzPlugin 'z' has both AOT and JIT implementations. NetworkDefinitionCreationFlag.PREFER_AOT_PYTHON_PLUGINS or NetworkDefinitionCreationFlag.PREFER_JIT_PYTHON_PLUGINS should be specified.z3' does not have either a AOT or JIT implementation.)$ QDP_REGISTRYr"r=AssertionErrorr^rrr_rk frombufferdatarmastype issubclassrpdecoderr@initr7r>r8r;r:rBrNQuickPluginCreationRequest STRICT_AOTr9rZr6 STRICT_JIT PREFER_AOT PREFER_JITrr?)rr"rfcphaserdescryfattr_type_annotr jit_or_aotrR aot_defined jit_definedrrrrMs              z$_TemplatePluginCreator.create_pluginr)rrL)r'r,r-rrMrrrrrsrr"rcCsDt}t|||}|||||d|}|t|d|<|S)NrJrD)rNrhrregister_creatorrirg)r"rr plg_registryrVrrrrjgs   rjFr6 lazy_registerrCcsdtffdd }|S)a( Wraps a function to register and describe a TensorRT plugin's IO characteristics. In addition, a complete plugin at least needs an `trt.plugin.impl` function to be registered. This API is only intended to be used as a decorator. The decorated function must have type hints for all inputs as well as return value. .. code-block:: text (inp0: TensorDesc, inp1: TensorDesc, ..., attr0: SupportedAttrType, attr1: SupportedAttrType, ...) -> Union[TensorDesc, Tuple[TensorDesc]] * Input tensors are declared first, each described by a tensor descriptor TensorDesc. * Plugin attributes are declared next. "SupportedAttrType" must be one of: * Supported built-in types: int, float, str, bool, bytes (Note: Lists/tuples of these types are not supported) * 1-D Numpy arrays of the following types: int8, int16, int32, int64, float16, float32, float64, bool. These must be annotated with 'numpy.typing.NDArray[dtype]', where 'dtype' is the expected numpy dtype. * If the plugin has only one output, the return annotation could be TensorDesc. Tuple[TensorDesc] could be used for any number of outputs. By default, the plugin will be immediately registered in the TRT plugin registry. Use the lazy_register argument to change this. Args: plugin_id: An ID for the plugin in the form "{namespace}::{name}", e.g. "my_project::add_plugin". The namespace is used to avoid collisions so using your product/project name is recommended. lazy_register: During plugin development/when building engine, lazy registration may be used to delay plugin registration until the plugin is explicitly instantiated using `trt.plugin.op.ns.plugin_name(...)` .. code-block:: python :linenos: :caption: Registration of an elementwise plugin (output has same characteristics as the input) import tensorrt.plugin as trtp @trtp.register("my::add_plugin") def add_plugin_desc(inp0: trtp.TensorDesc, block_size: int) -> Tuple[trtp.TensorDesc]: return inp0.like() r7c sd\}}t||t|}t||r#tdtjjd|dt|\}}}}t }|_ ||_ ||_ ||_ ||_t|} | |_|t<sQt|||||||S)NrDr%z' already has a defintion for ')rXrr3r2r rZrr'r r5r6r7r<r=rAr r@rrjr$) r7 plugin_ns plugin_name op_namespace tensor_namesr=rArr#r@rr6rr decorators4      zregister..decoratorr)r6rrrrrregisterss&&rcdtffdd }|S)a Wraps a function to define an implementation for a plugin already registered through `trt.plugin.register`. This API is only intended to be used as a decorator. The decorated function is not required to have type hints for input arguments or return value; however, any type hints specified will be validated against the `trt.plugin.register` signature for consistency. The schema for the function is as follows: .. code-block:: text (inp0: Tensor, inp1: Tensor, ..., attr0: SupportedAttrType, attr1: SupportedAttrType, outputs: Tuple[Tensor], stream: int, tactic: Optional[int]) -> None * Input tensors are passed first, each described by a `Tensor`. * Plugin attributes are declared next. * Not all attributes included in `trt.plugin.register` must be specified here -- they could be a subset. * Included attributes will be serialized to the TRT engine. Therefore, only attributes the plugin actually needs to perform inference (within the body of `trt.plugin.impl`) should be included. * `tactic` is an optional argument. If the plugin is using custom tactics, it must be specified to receive the tactic value to use for the current execution of the plugin. Args: plugin_id: The ID for the plugin in the form "{namespace}::{name}", which must match that used during `trt.plugin.register` .. code-block:: python :linenos: :caption: Implementation of an elementwise plugin with an OpenAI Triton kernel import tensorrt.plugin as trtp import triton import triton.language as tl @triton.jit def add_kernel(x_ptr, y_ptr, n_elements, BLOCK_SIZE: tl.constexpr): pid = tl.program_id(0) offsets = pid * BLOCK_SIZE + tl.arange(0, BLOCK_SIZE) mask = offsets < n_elements x = tl.load(x_ptr + offsets, mask=mask) tl.store(y_ptr + offsets, x + 1, mask=mask) @trtp.register("my::add_plugin") def add_plugin_desc(inp0: trtp.TensorDesc, block_size: int) -> Tuple[trtp.TensorDesc]: return inp0.like() @trtp.impl("my::add_plugin") def add_plugin_impl(inp0: trtp.Tensor, block_size: int, outputs: Tuple[trtp.Tensor], stream: int) -> None: n = inp0.numel() inp0_t = torch.as_tensor(inp0, device="cuda") out_t = torch.as_tensor(outputs[0], device="cuda") add_kernel[(triton.cdiv(n, block_size),)](inp0_t, out_t, n, BLOCK_SIZE = block_size) r8csDtvr tddt}t||\}}||_||_||_|SNzPlugin zJ is not registered. Did you register it with tensorrt.plugin.register API?)rrZr r8r>rB)r8r#r> found_tacticr6rrrs zimpl..decoratorrr6rrrrimpls5rcr)aW Wraps a function to define an Ahead-of-Time (AOT) implementation for a plugin already registered through `trt.plugin.register`. This API is only intended to be used as a decorator. The decorated function is not required to have type hints for input arguments or return value; however, any type hints specified will be validated against the `trt.plugin.register` signature for consistency. The schema for the function is as follows: .. code-block:: text (inp0: TensorDesc, inp1: TensorDesc, ..., attr0: SupportedAttrType, attr1: SupportedAttrType, outputs: Tuple[TensorDesc], tactic: Optional[int]) -> Tuple[str, str, KernelLaunchParams, SymExprs] * Input tensors are passed first, each described by a `TensorDesc`. * Plugin attributes are declared next. * Not all attributes included in `trt.plugin.register` must be specified here -- they could be a subset. * NOTE: Plugin attributes are not serialized into the engine when using an AOT implementation. * `tactic` is an optional argument. If the plugin is using custom tactics, it must be specified to receive the tactic value to use for the current execution of the plugin. Args: plugin_id: The ID for the plugin in the form "{namespace}::{name}", which must match that used during `trt.plugin.register` :returns: - kernel_name: The name of the kernel. - compiled_kernel: Compiled form of the kernel. Presently, only PTX is supported. - launch_params: The launch parameters for the kernel - extra_args: Symbolic expressions for scalar inputs to the kernel, located after the tensor inputs and before the tensor outputs .. code-block:: python :linenos: :caption: Implementation of an elementwise plugin with an OpenAI Triton kernel import tensorrt.plugin as trtp import triton import triton.language as tl @triton.jit def add_kernel(x_ptr, n_elements, y_ptr, BLOCK_SIZE: tl.constexpr): pid = tl.program_id(0) offsets = pid * BLOCK_SIZE + tl.arange(0, BLOCK_SIZE) mask = offsets < n_elements x = tl.load(x_ptr + offsets, mask=mask) tl.store(y_ptr + offsets, x + 1, mask=mask) @trtp.register("my::add_plugin") def add_plugin_desc(inp0: trtp.TensorDesc, block_size: int) -> Tuple[trtp.TensorDesc]: return inp0.like() @trtp.aot_impl("my::elemwise_add_plugin") def add_plugin_aot_impl( inp0: trtp.TensorDesc, block_size: int, single_tactic: bool, outputs: Tuple[trtp.TensorDesc], tactic: int ) -> Tuple[Union[str, bytes], Union[str, bytes], trtp.KernelLaunchParams, trtp.SymExprs]: type_str = "fp32" if inp0.dtype == trt.float32 else "fp16" src = triton.compiler.ASTSource( fn=add_kernel, signature=f"*{type_str},i32,*{type_str}", constants={ "BLOCK_SIZE": block_size, }, ) compiled_kernel = triton.compile(src) N = inp0.shape_expr.numel() launch_params = trtp.KernelLaunchParams() # grid dims launch_params.grid_x = trtp.cdiv(N, block_size) # block dims launch_params.block_x = compiled_kernel.metadata.num_warps * 32 # shared memory launch_params.shared_mem = compiled_kernel.metadata.shared extra_args = trtp.SymIntExprs(1) extra_args[0] = trtp.SymInt32(N) return compiled_kernel.metadata.name, compiled_kernel.asm["ptx"], launch_params, extra_args r9c:tvr tddt}t||}||_||_|Sr)rrZr r9r?)r9r#r?rrrrYs  zaot_impl..decoratorrrrrraot_impl sP rcr)a Wraps a function to define autotune logic for a plugin already registered through `trt.plugin.register`. Autotuning is the process by which TensorRT executes the plugin over IO type/format combinations, and any custom tactics advertised as being supported by the plugin. The (type, format, tactic) combination with the lowest latency is used to execute the plugin once the engine is built. .. note:: An autotune function is optional. If not specified, TensorRT will assume the plugin only supports input types specified at network creation, output types specifeid through `trt.plugin.register`, and linear formats for all I/O. This API is only intended to be used as a decorator. The decorated function is not required to have type hints for input arguments or return value; however, any type hints specified will be validated against the `trt.plugin.register` signature for consistency. The schema for the function is as follows: .. code-block:: text (inp0: TensorDesc, inp1: TensorDesc, ..., attr0: SupportedAttrType, attr1: SupportedAttrType, outputs: Tuple[TensorDesc]) -> List[AutoTuneCombination] * Input tensors are passed first, each described by a :class:`TensorDesc`. * Plugin attributes are declared next. Not all attributes included in `trt.plugin.register` must be specified here -- they could be a subset. * The function should return a list of :class:`AutoTuneCombination`\s. Args: plugin_id: The ID for the plugin in the form "{namespace}::{name}", which must match that used during `trt.plugin.register` .. code-block:: python :linenos: :caption: An elementwise add plugin which supports both FP32 and FP16 linear I/O and wants to be tuned over 2 custom tactics. import tensorrt.plugin as trtp @trtp.register("my::add_plugin") def add_plugin_desc(inp0: trtp.TensorDesc, block_size: int) -> Tuple[trtp.TensorDesc]: return inp0.like() @trtp.autotune("my::add_plugin") def add_plugin_autotune(inp0: trtp.TensorDesc, block_size: int, outputs: Tuple[trtp.TensorDesc]) -> List[trtp.AutoTuneCombination]: return [trtp.AutoTuneCombination("FP32|FP16, FP32|FP16", "LINEAR", [1, 2])] .. code-block:: python :linenos: :caption: Same as above example but using index-by-index construction of an `AutoTuneCombination` import tensorrt.plugin as trtp @trtp.register("my::add_plugin") def add_plugin_desc(inp0: trtp.TensorDesc, block_size: int) -> Tuple[trtp.TensorDesc]: return inp0.like() @trtp.autotune("my::add_plugin") def add_plugin_autotune(inp0: trtp.TensorDesc, block_size: int, outputs: Tuple[trtp.TensorDesc]) -> List[trtp.AutoTuneCombination]: c = trtp.AutoTuneCombination() c.pos(0, "FP32|FP16", "LINEAR") c.pos(1, "FP32|FP16") # index 1 is the output. Omitting format is the same as declaring it to be LINEAR. c.tactics([1, 2]) return [c] r:crr)rrZr r:r;)r:r#r;rrrrs  zautotune..decoratorrrrrrautotunejs;r)F),tensorrtrNtypesrarrrnumpyr_ _plugin_classr_exportrr _validater r r r r r_utilsrrrrrr ModuleTyperr/r3rgrr5rrrprjboolrrrrrrrrs>      OE`