"""Native LiteRT lowerings for shape-dependent reductions and shape arithmetic.""" import operator import torch from litert_torch import fx_infra from litert_torch.backend.experimental import torch_tfl from litert_torch.backend.experimental.torch_tfl._lowerings import _ir_operation from litert_torch.backend.lowerings import registry,utils from litert_converter.mlir import ir fx_infra.decomp.update_pre_lower_decomp(torch_tfl.decomps) def dynamic_layer_norm(x,normalized_shape,weight,bias,eps): axes=list(range(x.ndim-len(normalized_shape),x.ndim)) mean=x.mean(axes,keepdim=True) variance=((x-mean)*(x-mean)).mean(axes,keepdim=True) invstd=torch.rsqrt(variance+eps) output=(x-mean)*invstd if weight is not None:output=output*weight if bias is not None:output=output+bias return output,mean,invstd fx_infra.decomp.update_pre_lower_decomp({torch.ops.aten.native_layer_norm.default:dynamic_layer_norm}) @registry.lower(operator.floordiv) def lower_floor_div(lctx,lhs,rhs): if isinstance(lhs,int) and isinstance(rhs,int):return lhs//rhs if isinstance(lhs,int):lhs=utils.splat(lhs,rhs.type.element_type,rhs.type.shape) if isinstance(rhs,int):rhs=utils.splat(rhs,lhs.type.element_type,lhs.type.shape) if lhs.type!=rhs.type:raise TypeError('Shape floor division requires matching scalar integer types') return _ir_operation('tfl.floor_div',results=[lhs.type],operands=[lhs,rhs]) @registry.lower(torch.ops.aten.relu.default) def lower_relu(lctx,x): return _ir_operation('tfl.relu',results=utils.node_meta_to_ir_types(lctx.node),operands=[x]) @registry.lower(torch.ops.aten.cumsum.default) def lower_cumsum(lctx,x,dim,dtype=None): result=utils.node_meta_to_ir_types(lctx.node) if x.type.element_type!=result[0].element_type: typ=ir.RankedTensorType.get(x.type.shape,result[0].element_type) x=_ir_operation('tfl.cast',results=[typ],operands=[x]) axis=utils.splat(dim,ir.IntegerType.get_signless(32),[]) return _ir_operation('tfl.cumsum',results=result,operands=[x,axis],attributes={'exclusive':ir.BoolAttr.get(False),'reverse':ir.BoolAttr.get(False)}) @registry.lower(torch.ops.aten._is_all_true.default) def lower_all_true(lctx,x): if len(x.type.shape)==0:return x import numpy as np axes=utils.numpy_array_constant(np.arange(len(x.type.shape),dtype=np.int32)) return _ir_operation('tfl.reduce_all',results=utils.node_meta_to_ir_types(lctx.node),operands=[x,axes],attributes={'keep_dims':ir.BoolAttr.get(False)}) # Preserve symbolic dimensions in the operator schema. The upstream int[] # schema can specialize reshape dimensions during decomposition. @torch.library.custom_op('vds::reshape',mutates_args=(),schema='(Tensor x, SymInt[] shape) -> Tensor') def dynamic_reshape(x,shape): return x.reshape(shape).clone() @dynamic_reshape.register_fake def dynamic_reshape_fake(x,shape): return x.reshape(shape).clone() @registry.lower(torch.ops.vds.reshape.default) def lower_reshape(lctx,x,shape): return _ir_operation('tfl.reshape',results=utils.node_meta_to_ir_types(lctx.node),operands=[x,utils.convert_shape_to_ir_value(shape)]) def dynamic_view(x,shape):return torch.ops.vds.reshape(x,shape) fx_infra.decomp.update_pre_lower_decomp({torch.ops.aten.view.default:dynamic_view,torch.ops.aten._unsafe_view.default:dynamic_view}) def dynamic_arange(start,end,step=1,*,dtype=None,layout=None,device=None,pin_memory=None): result=torch.ops.tfl.range(start,end,step) return result.to(dtype=dtype) if dtype is not None else result fx_infra.decomp.update_pre_lower_decomp({torch.ops.aten.arange.start_step:dynamic_arange}) @torch.library.custom_op('vds::broadcast_to',mutates_args=(),schema='(Tensor x, SymInt[] shape) -> Tensor') def dynamic_broadcast(x,shape):return x.expand(shape).clone() @dynamic_broadcast.register_fake def dynamic_broadcast_fake(x,shape):return x.expand(shape).clone() @registry.lower(torch.ops.vds.broadcast_to.default) def lower_broadcast(lctx,x,shape): return _ir_operation('tfl.broadcast_to',results=utils.node_meta_to_ir_types(lctx.node),operands=[x,utils.convert_shape_to_ir_value(shape)]) fx_infra.decomp.update_pre_lower_decomp({torch.ops.aten.expand.default:lambda x,shape,implicit=False:torch.ops.vds.broadcast_to(x,shape)}) @torch.library.custom_op('vds::linear',mutates_args=()) def dynamic_linear(x:torch.Tensor,weight:torch.Tensor,bias:torch.Tensor|None)->torch.Tensor: return torch.nn.functional.linear(x,weight,bias) @dynamic_linear.register_fake def dynamic_linear_fake(x,weight,bias): return x.new_empty(x.shape[:-1] if weight.ndim==1 else (*x.shape[:-1],weight.shape[0])) @registry.lower(torch.ops.vds.linear.default) def lower_linear(lctx,x,weight,bias): if len(weight.type.shape)==1: from litert_converter.mlir.dialects import stablehlo from litert_torch.backend.experimental.torch_tfl._lowerings import _tfl_squeeze_lowering weight=stablehlo.reshape(ir.RankedTensorType.get([1,weight.type.shape[0]],weight.type.element_type),weight) if bias is None:bias=utils.splat(0.0,weight.type.element_type,[1]) typ=utils.node_meta_to_ir_types(lctx.node)[0] expanded=ir.RankedTensorType.get([*typ.shape,1],typ.element_type) y=_ir_operation('tfl.fully_connected',results=[expanded],operands=[x,weight,bias],attributes={'fused_activation_function':ir.StringAttr.get('NONE'),'weights_format':ir.StringAttr.get('DEFAULT'),'keep_num_dims':ir.BoolAttr.get(True),'asymmetric_quantize_inputs':ir.BoolAttr.get(False)}) return _tfl_squeeze_lowering(lctx,y,[-1]) if bias is None:bias=utils.splat(0.0,weight.type.element_type,[weight.type.shape[0]]) return _ir_operation('tfl.fully_connected',results=utils.node_meta_to_ir_types(lctx.node),operands=[x,weight,bias],attributes={'fused_activation_function':ir.StringAttr.get('NONE'),'weights_format':ir.StringAttr.get('DEFAULT'),'keep_num_dims':ir.BoolAttr.get(True),'asymmetric_quantize_inputs':ir.BoolAttr.get(False)}) fx_infra.decomp.update_pre_lower_decomp({torch.ops.aten.linear.default:lambda x,weight,bias=None:torch.ops.vds.linear(x,weight,bias)}) fx_infra.decomp.update_pre_convert_decomp({ torch.ops.aten.linear.default:lambda x,weight,bias=None:torch.ops.vds.linear(x,weight,bias), torch.ops.aten.view.default:dynamic_view, torch.ops.aten._unsafe_view.default:dynamic_view, torch.ops.aten.native_layer_norm.default:dynamic_layer_norm, torch.ops.aten.arange.start_step:dynamic_arange, torch.ops.aten.expand.default:lambda x,shape,implicit=False:torch.ops.vds.broadcast_to(x,shape), }) @torch.library.custom_op('vds::resize2d',mutates_args=(),schema='(Tensor x, SymInt[] size, bool align_corners, bool nearest) -> Tensor') def dynamic_resize2d(x,size,align_corners,nearest): if nearest:return torch.nn.functional.interpolate(x,size=size,mode='nearest') return torch.nn.functional.interpolate(x,size=size,mode='bilinear',align_corners=align_corners) @dynamic_resize2d.register_fake def dynamic_resize2d_fake(x,size,align_corners,nearest):return x.new_empty((x.shape[0],x.shape[1],*size)) @registry.lower(torch.ops.vds.resize2d.default) def lower_resize2d(lctx,x,size,align_corners,nearest): from litert_converter.mlir.dialects import stablehlo nhwc=stablehlo.transpose(x,[0,2,3,1]) result=utils.node_meta_to_ir_types(lctx.node)[0] typ=ir.RankedTensorType.get([result.shape[i] for i in [0,2,3,1]],result.element_type) resized=_ir_operation('tfl.resize_nearest_neighbor' if nearest else 'tfl.resize_bilinear',results=[typ],operands=[nhwc,utils.convert_shape_to_ir_value(size)],attributes={'align_corners':ir.BoolAttr.get(align_corners),'half_pixel_centers':ir.BoolAttr.get(not align_corners and not nearest)}) return stablehlo.transpose(resized,[0,3,1,2]) def bilinear_vec(x,output_size,align_corners,scale_factors=None): size=output_size if output_size is not None else [int(x.shape[-2+i]*scale_factors[i]) for i in range(2)] return torch.ops.vds.resize2d(x,size,align_corners,False) def bilinear_default(x,output_size,align_corners,scales_h=None,scales_w=None): return torch.ops.vds.resize2d(x,output_size,align_corners,False) fx_infra.decomp.update_pre_convert_decomp({torch.ops.aten.upsample_bilinear2d.vec:bilinear_vec,torch.ops.aten.upsample_bilinear2d.default:bilinear_default}) from litert_torch.backend.experimental.torch_tfl._lowerings import _tfl_mul_lowering @registry.lower(torch.ops.tfl.mul.default) def lower_bool_mul(lctx,lhs,rhs,fused_activation_function='NONE'): if isinstance(lhs,ir.Value) and isinstance(rhs,ir.Value) and str(lhs.type.element_type)=='i1' and str(rhs.type.element_type)=='i1': return _ir_operation('tfl.logical_and',results=utils.node_meta_to_ir_types(lctx.node),operands=[lhs,rhs]) return _tfl_mul_lowering(lctx,lhs,rhs,fused_activation_function) @torch.library.custom_op('vds::strided_slice',mutates_args=(),schema='(Tensor x, SymInt[] begin, SymInt[] end, int[] strides) -> Tensor') def dynamic_strided_slice(x,begin,end,strides): return x[tuple(slice(b,e,s) for b,e,s in zip(begin,end,strides))].clone() @dynamic_strided_slice.register_fake def dynamic_strided_slice_fake(x,begin,end,strides): return x[tuple(slice(b,e,s) for b,e,s in zip(begin,end,strides))].clone() @registry.lower(torch.ops.vds.strided_slice.default) def lower_strided_slice(lctx,x,begin,end,strides): from litert_torch.backend.experimental.torch_tfl._lowerings import _tfl_strided_slice_lowering return _tfl_strided_slice_lowering(lctx,x,begin,end,strides) def dynamic_slice(x,dim=0,start=None,end=None,step=1): n=x.shape[dim];b=0 if start is None else start;e=n if end is None else end if b<0:b+=n if e<0:e+=n b=max(0,min(b,n));e=max(b,min(e,n)) begin=[0]*x.ndim;ends=list(x.shape);strides=[1]*x.ndim begin[dim]=b;ends[dim]=e;strides[dim]=step return torch.ops.vds.strided_slice(x,begin,ends,strides) def dynamic_embedding(weight,indices,padding_idx=-1,scale_grad_by_freq=False,sparse=False): flat=torch.ops.vds.reshape(indices,[-1]).to(torch.int32) result=torch.ops.tfl.embedding_lookup(flat,weight) return torch.ops.vds.reshape(result,[*indices.shape,weight.shape[-1]]) fx_infra.decomp.update_pre_convert_decomp({torch.ops.aten.slice.Tensor:dynamic_slice,torch.ops.aten.embedding.default:dynamic_embedding})