(self, operation)
| 491 | """ |
| 492 | |
| 493 | def emit(self, operation): |
| 494 | |
| 495 | warp_shape = [ |
| 496 | operation.tile_description.threadblock_shape[idx] |
| 497 | // operation.tile_description.warp_count[idx] |
| 498 | for idx in range(3) |
| 499 | ] |
| 500 | |
| 501 | epilogue_vector_length = int( |
| 502 | min(operation.C.alignment * DataTypeSize[operation.C.element], 128) |
| 503 | / DataTypeSize[operation.C.element] |
| 504 | ) |
| 505 | |
| 506 | residual = "" |
| 507 | |
| 508 | values = { |
| 509 | "operation_name": operation.procedural_name(), |
| 510 | "element_a": DataTypeTag[operation.A.element], |
| 511 | "layout_a": LayoutTag[operation.A.layout], |
| 512 | "element_b": DataTypeTag[operation.B.element], |
| 513 | "layout_b": LayoutTag[operation.B.layout], |
| 514 | "element_c": DataTypeTag[operation.C.element], |
| 515 | "layout_c": LayoutTag[operation.C.layout], |
| 516 | "element_accumulator": DataTypeTag[operation.accumulator_type()], |
| 517 | "opcode_class": OpcodeClassTag[ |
| 518 | operation.tile_description.math_instruction.opcode_class |
| 519 | ], |
| 520 | "arch": "cutlass::arch::Sm%d" % operation.arch, |
| 521 | "threadblock_shape_m": str(operation.tile_description.threadblock_shape[0]), |
| 522 | "threadblock_shape_n": str(operation.tile_description.threadblock_shape[1]), |
| 523 | "threadblock_shape_k": str(operation.tile_description.threadblock_shape[2]), |
| 524 | "warp_shape_m": str(warp_shape[0]), |
| 525 | "warp_shape_n": str(warp_shape[1]), |
| 526 | "warp_shape_k": str(warp_shape[2]), |
| 527 | "instruction_shape_m": str( |
| 528 | operation.tile_description.math_instruction.instruction_shape[0] |
| 529 | ), |
| 530 | "instruction_shape_n": str( |
| 531 | operation.tile_description.math_instruction.instruction_shape[1] |
| 532 | ), |
| 533 | "instruction_shape_k": str( |
| 534 | operation.tile_description.math_instruction.instruction_shape[2] |
| 535 | ), |
| 536 | "epilogue_vector_length": str(epilogue_vector_length), |
| 537 | "element_epilogue": str(DataTypeTag[operation.element_epilogue]), |
| 538 | "epilogue_functor": EpilogueFunctorTag[operation.epilogue_functor], |
| 539 | "swizzling_functor": SwizzlingFunctorTag[operation.swizzling_functor], |
| 540 | "stages": str(operation.tile_description.stages), |
| 541 | "align_a": str(operation.A.alignment), |
| 542 | "align_b": str(operation.B.alignment), |
| 543 | "transform_a": ComplexTransformTag[operation.A.complex_transform], |
| 544 | "transform_b": ComplexTransformTag[operation.B.complex_transform], |
| 545 | "math_operation": MathOperationTag[ |
| 546 | operation.tile_description.math_instruction.math_operation |
| 547 | ], |
| 548 | "residual": residual, |
| 549 | } |
| 550 |
no test coverage detected