{"id":82467,"date":"2024-11-03T19:38:51","date_gmt":"2024-11-03T16:08:51","guid":{"rendered":"https:\/\/nabfollower.com\/blog\/on-the-programmability-of-aws-trainium-and-inferentia-4ick\/"},"modified":"2024-11-03T19:38:51","modified_gmt":"2024-11-03T16:08:51","slug":"on-the-programmability-of-aws-trainium-and-inferentia-4ick","status":"publish","type":"post","link":"https:\/\/nabfollower.com\/blog\/on-the-programmability-of-aws-trainium-and-inferentia-4ick\/","title":{"rendered":"\u062f\u0631 \u0645\u0648\u0631\u062f \u0628\u0631\u0646\u0627\u0645\u0647 \u0631\u06cc\u0632\u06cc AWS Trainium \u0648 Inferentia"},"content":{"rendered":"

Summarize this content to 400 words in Persian Lang <\/p>\n

\u062a\u0633\u0631\u06cc\u0639 \u0622\u0645\u0648\u0632\u0634 \u0645\u062f\u0644 AI\/ML \u0628\u0627 \u0627\u067e\u0631\u0627\u062a\u0648\u0631\u0647\u0627\u06cc \u0633\u0641\u0627\u0631\u0634\u06cc – \u0642\u0633\u0645\u062a 4<\/p>\n

\u0639\u06a9\u0633 \u0622\u06af\u0627\u062a\u0627 \u0628\u0631\u0633 \u062f\u0631 Unsplash<\/p>\n

\u062f\u0631 \u0627\u06cc\u0646 \u067e\u0633\u062a \u0645\u0627 \u0628\u0647 \u0628\u0631\u0631\u0633\u06cc \u0641\u0631\u0635\u062a\u200c\u0647\u0627\u06cc \u0628\u0647\u06cc\u0646\u0647\u200c\u0633\u0627\u0632\u06cc \u0632\u0645\u0627\u0646 \u0627\u062c\u0631\u0627 \u0628\u0627\u0631\u0647\u0627\u06cc \u06a9\u0627\u0631\u06cc \u06cc\u0627\u062f\u06af\u06cc\u0631\u06cc \u0645\u0627\u0634\u06cc\u0646 (ML) \u0627\u0632 \u0637\u0631\u06cc\u0642 \u062a\u0648\u0633\u0639\u0647 \u0627\u067e\u0631\u0627\u062a\u0648\u0631 \u0633\u0641\u0627\u0631\u0634\u06cc \u0627\u062f\u0627\u0645\u0647 \u0645\u06cc\u200c\u062f\u0647\u06cc\u0645. \u0627\u06cc\u0646 \u0628\u0627\u0631\u060c \u0645\u0627 \u0628\u0631 \u0627\u0628\u0632\u0627\u0631\u0647\u0627\u06cc \u0627\u0631\u0627\u0626\u0647 \u0634\u062f\u0647 \u062a\u0648\u0633\u0637 AWS Neuron SDK \u0628\u0631\u0627\u06cc \u062a\u0648\u0633\u0639\u0647 \u0648 \u0627\u062c\u0631\u0627\u06cc \u0647\u0633\u062a\u0647 \u0647\u0627\u06cc \u062c\u062f\u06cc\u062f \u062f\u0631 AWS Trainium \u0648 AWS Inferentia \u062a\u0645\u0631\u06a9\u0632 \u0645\u06cc \u06a9\u0646\u06cc\u0645. \u0628\u0627 \u062a\u0648\u0633\u0639\u0647 \u0633\u0631\u06cc\u0639 \u0627\u062c\u0632\u0627\u06cc \u0645\u062f\u0644 \u0633\u0637\u062d \u067e\u0627\u06cc\u06cc\u0646 (\u0628\u0647 \u0639\u0646\u0648\u0627\u0646 \u0645\u062b\u0627\u0644\u060c \u0644\u0627\u06cc\u0647 \u0647\u0627\u06cc \u062a\u0648\u062c\u0647) \u06a9\u0647 \u0627\u0646\u0642\u0644\u0627\u0628 \u0647\u0648\u0634 \u0645\u0635\u0646\u0648\u0639\u06cc \u0631\u0627 \u0647\u062f\u0627\u06cc\u062a \u0645\u06cc \u06a9\u0646\u062f\u060c \u0628\u0631\u0646\u0627\u0645\u0647 \u0631\u06cc\u0632\u06cc \u0634\u062a\u0627\u0628 \u062f\u0647\u0646\u062f\u0647 \u0647\u0627\u06cc \u0645\u0648\u0631\u062f \u0627\u0633\u062a\u0641\u0627\u062f\u0647 \u0628\u0631\u0627\u06cc \u0622\u0645\u0648\u0632\u0634 \u0648 \u0627\u062c\u0631\u0627\u06cc \u0645\u062f\u0644 \u0647\u0627\u06cc ML \u0628\u0633\u06cc\u0627\u0631 \u0645\u0647\u0645 \u0627\u0633\u062a. \u062a\u0631\u0627\u0634\u0647\u200c\u0647\u0627\u06cc \u0627\u062e\u062a\u0635\u0627\u0635\u06cc \u0647\u0648\u0634 \u0645\u0635\u0646\u0648\u0639\u06cc\u060c \u0628\u0647\u200c\u0648\u06cc\u0698\u0647\u060c \u0628\u0627\u06cc\u062f \u062c\u0627\u06cc\u06af\u0632\u06cc\u0646 \u0645\u0646\u0627\u0633\u0628\u06cc \u0628\u0631\u0627\u06cc \u0686\u0627\u0631\u0686\u0648\u0628\u200c\u0647\u0627\u06cc \u062a\u0648\u0633\u0639\u0647 \u067e\u0631\u06a9\u0627\u0631\u0628\u0631\u062f \u0648 \u0628\u0633\u06cc\u0627\u0631 \u062a\u0627\u062b\u06cc\u0631\u06af\u0630\u0627\u0631 GPU (GPGPU) \u0645\u0627\u0646\u0646\u062f CUDA \u0648 Triton \u0628\u0627\u0634\u0646\u062f.<\/p>\n

\u062f\u0631 \u067e\u0633\u062a\u200c\u0647\u0627\u06cc \u0642\u0628\u0644\u06cc (\u0628\u0647 \u0639\u0646\u0648\u0627\u0646 \u0645\u062b\u0627\u0644\u060c \u0627\u06cc\u0646\u062c\u0627 \u0648 \u0627\u06cc\u0646\u062c\u0627) \u0641\u0631\u0635\u062a \u0633\u0627\u062e\u062a \u0648 \u0627\u062c\u0631\u0627\u06cc \u0645\u062f\u0644\u200c\u0647\u0627\u06cc ML \u0628\u0631 \u0631\u0648\u06cc \u062a\u0631\u0627\u0634\u0647\u200c\u0647\u0627\u06cc AI \u0633\u0641\u0627\u0631\u0634\u06cc \u0633\u0627\u062e\u062a AWS \u0628\u0627 \u0627\u0633\u062a\u0641\u0627\u062f\u0647 \u0627\u0632 AWS Neuron SDK \u0631\u0627 \u0628\u0631\u0631\u0633\u06cc \u06a9\u0631\u062f\u06cc\u0645. \u062f\u0631 \u062c\u062f\u06cc\u062f\u062a\u0631\u06cc\u0646 \u0646\u0633\u062e\u0647 SDK (\u0646\u0633\u062e\u0647 2.20.0)\u060c AWS \u0631\u0627\u0628\u0637 \u0647\u0633\u062a\u0647 \u0639\u0635\u0628\u06cc (NKI) \u0631\u0627 \u0628\u0631\u0627\u06cc \u062a\u0648\u0633\u0639\u0647 \u0647\u0633\u062a\u0647\u200c\u0647\u0627\u06cc \u0633\u0641\u0627\u0631\u0634\u06cc \u0628\u0631\u0627\u06cc NeuronCore-v2\u060c \u0634\u062a\u0627\u0628\u200c\u062f\u0647\u0646\u062f\u0647 \u0632\u06cc\u0631\u0628\u0646\u0627\u06cc\u06cc \u06a9\u0647 \u0647\u0645 Trainium \u0648 \u0647\u0645 Inferentia2 \u0631\u0627 \u062a\u0627\u0645\u06cc\u0646 \u0645\u06cc\u200c\u06a9\u0646\u062f\u060c \u0645\u0639\u0631\u0641\u06cc \u06a9\u0631\u062f. \u0631\u0627\u0628\u0637 NKI \u0628\u0647 API \u062f\u06cc\u06af\u0631\u06cc \u0645\u0644\u062d\u0642 \u0645\u06cc \u0634\u0648\u062f \u06a9\u0647 \u0628\u0631\u0646\u0627\u0645\u0647 \u0646\u0648\u06cc\u0633\u06cc NeuronCore-v2 \u0631\u0627 \u0641\u0639\u0627\u0644 \u0645\u06cc \u06a9\u0646\u062f\u060c Neuron Custom C++ Operators. \u062f\u0631 \u0627\u06cc\u0646 \u067e\u0633\u062a \u0647\u0631 \u062f\u0648 \u0641\u0631\u0635\u062a \u0631\u0627 \u0628\u0631\u0631\u0633\u06cc \u06a9\u0631\u062f\u0647 \u0648 \u0622\u0646\u0647\u0627 \u0631\u0627 \u062f\u0631 \u0639\u0645\u0644 \u0646\u0634\u0627\u0646 \u062e\u0648\u0627\u0647\u06cc\u0645 \u062f\u0627\u062f.<\/p>\n

\u0633\u0644\u0628 \u0645\u0633\u0626\u0648\u0644\u06cc\u062a<\/p>\n

\u0646\u06a9\u062a\u0647 \u0645\u0647\u0645 \u0627\u06cc\u0646 \u0627\u0633\u062a \u06a9\u0647 \u0627\u06cc\u0646 \u067e\u0633\u062a \u0646\u0628\u0627\u06cc\u062f \u0628\u0647 \u0639\u0646\u0648\u0627\u0646 \u062c\u0627\u06cc\u06af\u0632\u06cc\u0646\u06cc \u0628\u0631\u0627\u06cc \u0627\u0633\u0646\u0627\u062f \u0631\u0633\u0645\u06cc AWS Neuron SDK \u062f\u0631 \u0646\u0638\u0631 \u06af\u0631\u0641\u062a\u0647 \u0634\u0648\u062f. \u062f\u0631 \u0632\u0645\u0627\u0646 \u0646\u0648\u0634\u062a\u0646 \u0627\u06cc\u0646 \u0645\u0642\u0627\u0644\u0647\u060c Neuron SDK APIs \u0628\u0631\u0627\u06cc \u062a\u0648\u0633\u0639\u0647 \u0647\u0633\u062a\u0647 \u0633\u0641\u0627\u0631\u0634\u06cc \u062f\u0631 \u0628\u062a\u0627 \u0627\u0633\u062a \u0648 \u0645\u0645\u06a9\u0646 \u0627\u0633\u062a \u0628\u0627 \u062e\u0648\u0627\u0646\u062f\u0646 \u0627\u06cc\u0646 \u0645\u0637\u0644\u0628 \u062a\u063a\u06cc\u06cc\u0631 \u06a9\u0646\u062f. \u0646\u0645\u0648\u0646\u0647 \u0647\u0627\u06cc\u06cc \u06a9\u0647 \u0645\u0627 \u0628\u0647 \u0627\u0634\u062a\u0631\u0627\u06a9 \u0645\u06cc \u06af\u0630\u0627\u0631\u06cc\u0645\u060c \u0641\u0642\u0637 \u0628\u0631\u0627\u06cc \u0627\u0647\u062f\u0627\u0641 \u0646\u0645\u0627\u06cc\u0634\u06cc \u062f\u0631 \u0646\u0638\u0631 \u06af\u0631\u0641\u062a\u0647 \u0634\u062f\u0647 \u0627\u0646\u062f. \u0645\u0627 \u0647\u06cc\u0686 \u0627\u062f\u0639\u0627\u06cc\u06cc \u062f\u0631 \u0645\u0648\u0631\u062f \u0628\u0647\u06cc\u0646\u0647 \u0628\u0648\u062f\u0646\u060c \u0627\u0633\u062a\u062d\u06a9\u0627\u0645\u060c \u062f\u0648\u0627\u0645 \u06cc\u0627 \u062f\u0642\u062a \u0622\u0646\u0647\u0627 \u0646\u062f\u0627\u0631\u06cc\u0645. \u0644\u0637\u0641\u0627\u064b \u0627\u0634\u0627\u0631\u0647 \u0645\u0627 \u0628\u0647 \u067e\u0644\u062a\u0641\u0631\u0645\u200c\u0647\u0627\u060c \u0627\u0628\u0632\u0627\u0631\u0647\u0627\u060c API\u0647\u0627 \u0648 \u063a\u06cc\u0631\u0647 \u0631\u0627 \u0628\u0647 \u0639\u0646\u0648\u0627\u0646 \u062a\u0623\u06cc\u06cc\u062f\u06cc \u0628\u0631\u0627\u06cc \u0627\u0633\u062a\u0641\u0627\u062f\u0647 \u0627\u0632 \u0622\u0646\u200c\u0647\u0627 \u062a\u0644\u0642\u06cc \u0646\u06a9\u0646\u06cc\u062f. \u0628\u0647\u062a\u0631\u06cc\u0646 \u0627\u0646\u062a\u062e\u0627\u0628 \u0628\u0631\u0627\u06cc \u0647\u0631 \u067e\u0631\u0648\u0698\u0647 \u0628\u0647 \u0648\u06cc\u0698\u06af\u06cc \u0647\u0627\u06cc \u0645\u0648\u0631\u062f \u0627\u0633\u062a\u0641\u0627\u062f\u0647 \u062f\u0631 \u062f\u0633\u062a \u0628\u0633\u062a\u06af\u06cc \u062f\u0627\u0631\u062f \u0648 \u062a\u062d\u0642\u06cc\u0642\u0627\u062a \u0648 \u062a\u062c\u0632\u06cc\u0647 \u0648 \u062a\u062d\u0644\u06cc\u0644 \u0645\u0646\u0627\u0633\u0628 \u0631\u0627 \u062a\u0636\u0645\u06cc\u0646 \u0645\u06cc \u06a9\u0646\u062f.<\/p>\n

\u062a\u0648\u0633\u0639\u0647 \u0647\u0633\u062a\u0647 \u0647\u0627\u06cc \u0633\u0641\u0627\u0631\u0634\u06cc \u0628\u0631\u0627\u06cc \u0647\u0633\u062a\u0647 \u0647\u0627\u06cc \u0639\u0635\u0628\u06cc<\/p>\n

\u0627\u06af\u0631\u0686\u0647 \u0644\u06cc\u0633\u062a \u0645\u062f\u0644 \u0647\u0627\u06cc ML \u067e\u0634\u062a\u06cc\u0628\u0627\u0646\u06cc \u0634\u062f\u0647 \u062a\u0648\u0633\u0637 Neuron SDK \u0628\u0647 \u0637\u0648\u0631 \u0645\u062f\u0627\u0648\u0645 \u062f\u0631 \u062d\u0627\u0644 \u0627\u0641\u0632\u0627\u06cc\u0634 \u0627\u0633\u062a\u060c \u0628\u0631\u062e\u06cc \u0627\u0632 \u0639\u0645\u0644\u06cc\u0627\u062a \u0647\u0627 \u06cc\u0627 \u067e\u0634\u062a\u06cc\u0628\u0627\u0646\u06cc \u0646\u0645\u06cc \u0634\u0648\u0646\u062f \u06cc\u0627 \u0628\u0647 \u0635\u0648\u0631\u062a \u063a\u06cc\u0631\u0628\u0647\u06cc\u0646\u0647 \u0627\u062c\u0631\u0627 \u0645\u06cc \u0634\u0648\u0646\u062f. \u0628\u0627 \u0627\u0641\u0634\u0627\u06cc API\u0647\u0627 \u0628\u0631\u0627\u06cc \u0633\u0641\u0627\u0631\u0634\u06cc\u200c\u0633\u0627\u0632\u06cc \u0647\u0633\u062a\u0647 \u0646\u0648\u0631\u0648\u0646\u060c SDK \u0628\u0647 \u062a\u0648\u0633\u0639\u0647\u200c\u062f\u0647\u0646\u062f\u06af\u0627\u0646 \u0627\u062c\u0627\u0632\u0647 \u0645\u06cc\u200c\u062f\u0647\u062f \u062a\u0627 \u0639\u0645\u0644\u06cc\u0627\u062a \u0633\u0637\u062d \u067e\u0627\u06cc\u06cc\u0646 \u0645\u0648\u0631\u062f \u0646\u06cc\u0627\u0632 \u062e\u0648\u062f \u0631\u0627 \u0627\u06cc\u062c\u0627\u062f \u0648\/\u06cc\u0627 \u0628\u0647\u06cc\u0646\u0647\u200c\u0633\u0627\u0632\u06cc \u06a9\u0646\u0646\u062f \u0648 \u0641\u0631\u0635\u062a \u0627\u062c\u0631\u0627\u06cc \u0628\u0627\u0631\u0647\u0627\u06cc \u06a9\u0627\u0631\u06cc ML \u062f\u0631 Trainium \u0648 Inferentia \u0631\u0627 \u0628\u0633\u06cc\u0627\u0631 \u0627\u0641\u0632\u0627\u06cc\u0634 \u0645\u06cc\u200c\u062f\u0647\u062f.<\/p>\n

\u0647\u0645\u0627\u0646\u0637\u0648\u0631 \u06a9\u0647 \u062f\u0631 \u067e\u0633\u062a \u0647\u0627\u06cc \u0642\u0628\u0644\u06cc \u0645\u0627 \u062f\u0631 \u0627\u06cc\u0646 \u0633\u0631\u06cc \u0628\u062d\u062b \u0634\u062f\u060c \u0627\u0633\u062a\u0641\u0627\u062f\u0647 \u06a9\u0627\u0645\u0644 \u0627\u0632 \u0642\u062f\u0631\u062a \u0627\u06cc\u0646 \u062a\u0631\u0627\u0634\u0647 \u0647\u0627\u06cc \u0647\u0648\u0634 \u0645\u0635\u0646\u0648\u0639\u06cc \u0645\u0633\u062a\u0644\u0632\u0645 \u062f\u0631\u06a9 \u062f\u0642\u06cc\u0642 \u0645\u0639\u0645\u0627\u0631\u06cc \u0633\u0637\u062d \u067e\u0627\u06cc\u06cc\u0646 \u0622\u0646\u0647\u0627 \u0627\u0633\u062a.<\/p>\n

\u0645\u0639\u0645\u0627\u0631\u06cc \u0647\u0633\u062a\u0647 \u0646\u0648\u0631\u0648\u0646<\/p>\n

\u0627\u0633\u0646\u0627\u062f NKI \u0634\u0627\u0645\u0644 \u0628\u062e\u0634 \u0627\u062e\u062a\u0635\u0627\u0635\u06cc \u062f\u0631 \u0645\u0648\u0631\u062f \u0637\u0631\u0627\u062d\u06cc \u0645\u0639\u0645\u0627\u0631\u06cc NeuronCore-v2 \u0648 \u067e\u06cc\u0627\u0645\u062f\u0647\u0627\u06cc \u0622\u0646 \u0628\u0631 \u062a\u0648\u0633\u0639\u0647 \u0627\u067e\u0631\u0627\u062a\u0648\u0631 \u0633\u0641\u0627\u0631\u0634\u06cc \u0627\u0633\u062a. \u0646\u06a9\u062a\u0647 \u0645\u0647\u0645 \u0627\u06cc\u0646 \u0627\u0633\u062a \u06a9\u0647 \u062a\u0641\u0627\u0648\u062a\u200c\u0647\u0627\u06cc \u0632\u06cc\u0627\u062f\u06cc \u0628\u06cc\u0646 \u0647\u0633\u062a\u0647\u200c\u0647\u0627\u06cc \u0646\u0648\u0631\u0648\u0646 \u0648 \u0647\u0645\u062a\u0627\u06cc\u0627\u0646 \u0634\u062a\u0627\u0628\u200c\u062f\u0647\u0646\u062f\u0647 \u0647\u0648\u0634 \u0645\u0635\u0646\u0648\u0639\u06cc \u0622\u0646\u200c\u0647\u0627 (\u0645\u0627\u0646\u0646\u062f GPU \u0648 TPU) \u0648\u062c\u0648\u062f \u062f\u0627\u0631\u062f. \u0628\u0647\u06cc\u0646\u0647\u200c\u0633\u0627\u0632\u06cc \u0647\u0633\u062a\u0647\u200c\u0647\u0627\u06cc \u0646\u0648\u0631\u0648\u0646 \u0628\u0647 \u0645\u062c\u0645\u0648\u0639\u0647\u200c\u0627\u06cc \u0627\u0632 \u0627\u0633\u062a\u0631\u0627\u062a\u0698\u06cc\u200c\u0647\u0627 \u0648 \u0645\u0647\u0627\u0631\u062a\u200c\u0647\u0627 \u0646\u06cc\u0627\u0632 \u062f\u0627\u0631\u062f.<\/p>\n

\u0645\u0634\u0627\u0628\u0647 \u0633\u0627\u06cc\u0631 \u062a\u0631\u0627\u0634\u0647 \u0647\u0627\u06cc \u0627\u062e\u062a\u0635\u0627\u0635\u06cc \u0647\u0648\u0634 \u0645\u0635\u0646\u0648\u0639\u06cc\u060c NeuronCore-v2 \u0634\u0627\u0645\u0644 \u0686\u0646\u062f\u06cc\u0646 \u0645\u0648\u062a\u0648\u0631 \u0634\u062a\u0627\u0628 \u062f\u0647\u0646\u062f\u0647 \u062f\u0627\u062e\u0644\u06cc \u0627\u0633\u062a \u06a9\u0647 \u0647\u0631 \u06a9\u062f\u0627\u0645 \u062f\u0631 \u0627\u0646\u062c\u0627\u0645 \u0627\u0646\u0648\u0627\u0639 \u062e\u0627\u0635\u06cc \u0627\u0632 \u0645\u062d\u0627\u0633\u0628\u0627\u062a \u062a\u062e\u0635\u0635 \u062f\u0627\u0631\u0646\u062f. \u0645\u0648\u062a\u0648\u0631\u0647\u0627 \u0645\u06cc \u062a\u0648\u0627\u0646\u0646\u062f \u0628\u0647 \u0635\u0648\u0631\u062a \u0646\u0627\u0647\u0645\u0632\u0645\u0627\u0646 \u0648 \u0645\u0648\u0627\u0632\u06cc \u06a9\u0627\u0631 \u06a9\u0646\u0646\u062f. \u06a9\u0627\u0645\u067e\u0627\u06cc\u0644\u0631 \u0646\u0648\u0631\u0648\u0646 \u0645\u0633\u0626\u0648\u0644 \u062a\u0628\u062f\u06cc\u0644 \u0645\u062f\u0644 \u0647\u0627\u06cc ML \u0628\u0647 \u0639\u0645\u0644\u06cc\u0627\u062a \u0633\u0637\u062d \u067e\u0627\u06cc\u06cc\u0646 \u0648 \u0628\u0647\u06cc\u0646\u0647 \u0633\u0627\u0632\u06cc \u0627\u0646\u062a\u062e\u0627\u0628 \u0645\u0648\u062a\u0648\u0631 \u0645\u062d\u0627\u0633\u0628\u0627\u062a\u06cc \u0628\u0631\u0627\u06cc \u0647\u0631 \u06cc\u06a9 \u0627\u0633\u062a.<\/p>\n

\u0645\u0648\u062a\u0648\u0631 Tensor \u062f\u0631 \u0636\u0631\u0628 \u0645\u0627\u062a\u0631\u06cc\u0633 \u062a\u062e\u0635\u0635 \u062f\u0627\u0631\u062f. \u0645\u0648\u062a\u0648\u0631\u0647\u0627\u06cc Vector \u0648 Scalar \u0647\u0631 \u062f\u0648 \u0628\u0631 \u0631\u0648\u06cc \u062a\u0627\u0646\u0633\u0648\u0631\u0647\u0627 \u0628\u0627 \u0645\u0648\u062a\u0648\u0631 Vector \u0645\u062a\u062e\u0635\u0635 \u062f\u0631 \u0639\u0645\u0644\u06cc\u0627\u062a \u06a9\u0627\u0647\u0634 \u0648 \u0645\u0648\u062a\u0648\u0631 Scalar \u062f\u0631 \u062a\u0648\u0627\u0628\u0639 \u063a\u06cc\u0631 \u062e\u0637\u06cc \u06a9\u0627\u0631 \u0645\u06cc \u06a9\u0646\u0646\u062f. GpSimd \u06cc\u06a9 \u0645\u0648\u062a\u0648\u0631 \u0647\u0645\u0647 \u0645\u0646\u0638\u0648\u0631\u0647 \u0627\u0633\u062a \u06a9\u0647 \u0642\u0627\u062f\u0631 \u0628\u0647 \u0627\u062c\u0631\u0627\u06cc \u0628\u0631\u0646\u0627\u0645\u0647 \u0647\u0627\u06cc \u062f\u0644\u062e\u0648\u0627\u0647 C\/C++ \u0627\u0633\u062a. \u062a\u0648\u062c\u0647 \u062f\u0627\u0634\u062a\u0647 \u0628\u0627\u0634\u06cc\u062f \u06a9\u0647 \u062f\u0631 \u062d\u0627\u0644\u06cc \u06a9\u0647 \u0631\u0627\u0628\u0637 NKI \u062f\u0633\u062a\u0631\u0633\u06cc \u0628\u0647 \u0647\u0631 \u0686\u0647\u0627\u0631 \u0645\u0648\u062a\u0648\u0631 \u0645\u062d\u0627\u0633\u0628\u0627\u062a\u06cc \u0631\u0627 \u0646\u0634\u0627\u0646 \u0645\u06cc \u062f\u0647\u062f\u060c \u0627\u067e\u0631\u0627\u062a\u0648\u0631\u0647\u0627\u06cc C++ \u0633\u0641\u0627\u0631\u0634\u06cc \u0628\u0647 \u0637\u0648\u0631 \u062e\u0627\u0635 \u0628\u0631\u0627\u06cc GpSimd \u0637\u0631\u0627\u062d\u06cc \u0634\u062f\u0647 \u0627\u0646\u062f.<\/p>\n

\u062c\u0632\u0626\u06cc\u0627\u062a \u0628\u06cc\u0634\u062a\u0631 \u062f\u0631 \u0645\u0648\u0631\u062f \u0642\u0627\u0628\u0644\u06cc\u062a \u0647\u0627\u06cc \u0647\u0631 \u0645\u0648\u062a\u0648\u0631 \u0631\u0627 \u0645\u06cc \u062a\u0648\u0627\u0646 \u062f\u0631 \u0645\u0633\u062a\u0646\u062f\u0627\u062a \u0645\u0639\u0645\u0627\u0631\u06cc \u06cc\u0627\u0641\u062a. \u0639\u0644\u0627\u0648\u0647 \u0628\u0631 \u0627\u06cc\u0646\u060c \u0645\u0633\u062a\u0646\u062f\u0627\u062a NKI Instruction Set Architecture (ISA) \u062c\u0632\u0626\u06cc\u0627\u062a\u06cc \u0631\u0627 \u062f\u0631 \u0645\u0648\u0631\u062f \u0645\u0648\u062a\u0648\u0631\u0647\u0627\u06cc\u06cc \u06a9\u0647 \u0639\u0645\u0644\u06cc\u0627\u062a \u0647\u0627\u06cc \u0645\u062e\u062a\u0644\u0641 \u0633\u0637\u062d \u067e\u0627\u06cc\u06cc\u0646 \u0628\u0631 \u0631\u0648\u06cc \u0622\u0646\u0647\u0627 \u0627\u062c\u0631\u0627 \u0645\u06cc \u0634\u0648\u062f\u060c \u0627\u0631\u0627\u0626\u0647 \u0645\u06cc \u062f\u0647\u062f.<\/p>\n

\u06cc\u06a9\u06cc \u062f\u06cc\u06af\u0631 \u0627\u0632 \u062c\u0646\u0628\u0647 \u0647\u0627\u06cc \u0645\u0647\u0645 \u062a\u0631\u0627\u0634\u0647 \u0646\u0648\u0631\u0648\u0646\u060c \u0645\u0639\u0645\u0627\u0631\u06cc \u062d\u0627\u0641\u0638\u0647 \u0622\u0646 \u0627\u0633\u062a. \u062f\u0633\u062a\u06af\u0627\u0647 Neuron \u0634\u0627\u0645\u0644 \u0633\u0647 \u0646\u0648\u0639 \u062d\u0627\u0641\u0638\u0647 HBM\u060c SBUF \u0648 PSUM \u0627\u0633\u062a. \u062f\u0631\u06a9 \u0646\u0632\u062f\u06cc\u06a9 \u0627\u0632 \u0638\u0631\u0641\u06cc\u062a \u0647\u0627 \u0648 \u0642\u0627\u0628\u0644\u06cc\u062a \u0647\u0627\u06cc \u0647\u0631 \u06cc\u06a9 \u0628\u0631\u0627\u06cc \u062a\u0648\u0633\u0639\u0647 \u0628\u0647\u06cc\u0646\u0647 \u0647\u0633\u062a\u0647 \u0628\u0633\u06cc\u0627\u0631 \u0645\u0647\u0645 \u0627\u0633\u062a.<\/p>\n

\u0628\u0627 \u062a\u0648\u062c\u0647 \u0628\u0647 \u0646\u0645\u0627\u06cc \u06a9\u0644\u06cc \u0645\u0639\u0645\u0627\u0631\u06cc\u060c \u0645\u0645\u06a9\u0646 \u0627\u0633\u062a \u0646\u062a\u06cc\u062c\u0647 \u0628\u06af\u06cc\u0631\u06cc\u062f \u06a9\u0647 \u062a\u0648\u0633\u0639\u0647 \u0647\u0633\u062a\u0647 \u0646\u0648\u0631\u0648\u0646 \u0628\u0647 \u062a\u062e\u0635\u0635 \u0628\u0627\u0644\u0627\u06cc\u06cc \u0646\u06cc\u0627\u0632 \u062f\u0627\u0631\u062f. \u0627\u06af\u0631\u0686\u0647 \u0627\u06cc\u0646 \u0645\u0645\u06a9\u0646 \u0627\u0633\u062a \u0628\u0631\u0627\u06cc \u0627\u06cc\u062c\u0627\u062f \u0647\u0633\u062a\u0647\u200c\u0647\u0627\u06cc \u06a9\u0627\u0645\u0644\u0627\u064b \u0628\u0647\u06cc\u0646\u0647\u200c\u0633\u0627\u0632\u06cc \u0634\u062f\u0647 \u06a9\u0647 \u0627\u0632 \u062a\u0645\u0627\u0645 \u0642\u0627\u0628\u0644\u06cc\u062a\u200c\u0647\u0627\u06cc \u0647\u0633\u062a\u0647 \u0646\u0648\u0631\u0648\u0646 \u0627\u0633\u062a\u0641\u0627\u062f\u0647 \u0645\u06cc\u200c\u06a9\u0646\u0646\u062f \u0635\u0627\u062f\u0642 \u0628\u0627\u0634\u062f\u060c \u0647\u062f\u0641 \u0645\u0627 \u0646\u0634\u0627\u0646 \u062f\u0627\u062f\u0646 \u0642\u0627\u0628\u0644\u06cc\u062a \u062f\u0633\u062a\u0631\u0633\u06cc\u060c \u0627\u0631\u0632\u0634 \u0648 \u067e\u062a\u0627\u0646\u0633\u06cc\u0644 API\u0647\u0627\u06cc \u0647\u0633\u062a\u0647 \u0633\u0641\u0627\u0631\u0634\u06cc Neuron – \u062d\u062a\u06cc \u0628\u0631\u0627\u06cc \u062a\u0648\u0633\u0639\u0647\u200c\u062f\u0647\u0646\u062f\u06af\u0627\u0646 \u063a\u06cc\u0631\u0645\u062a\u062e\u0635\u0635 \u0627\u0633\u062a.<\/p>\n

\u0647\u0633\u062a\u0647 \u0647\u0627\u06cc \u0633\u0641\u0627\u0631\u0634\u06cc NKI<\/p>\n

\u0631\u0627\u0628\u0637 NKI \u06cc\u06a9 API \u062f\u0631 \u0633\u0637\u062d \u067e\u0627\u06cc\u062a\u0648\u0646 \u0627\u0633\u062a \u06a9\u0647 \u0627\u0633\u062a\u0641\u0627\u062f\u0647 \u0627\u0632 \u0645\u0648\u062a\u0648\u0631\u0647\u0627\u06cc \u0645\u062d\u0627\u0633\u0628\u0627\u062a\u06cc \u0647\u0633\u062a\u0647 \u0646\u0648\u0631\u0648\u0646 \u0648 \u0645\u0646\u0627\u0628\u0639 \u062d\u0627\u0641\u0638\u0647 \u0631\u0627 \u062f\u0631 \u0627\u062e\u062a\u06cc\u0627\u0631 \u062a\u0648\u0633\u0639\u0647 \u062f\u0647\u0646\u062f\u06af\u0627\u0646 ML \u0642\u0631\u0627\u0631 \u0645\u06cc \u062f\u0647\u062f. \u0631\u0627\u0647\u0646\u0645\u0627\u06cc \u0634\u0631\u0648\u0639 NKI \u062f\u0633\u062a\u0648\u0631\u0627\u0644\u0639\u0645\u0644\u200c\u0647\u0627\u06cc \u0631\u0627\u0647\u200c\u0627\u0646\u062f\u0627\u0632\u06cc \u0631\u0627 \u0628\u0647 \u062a\u0641\u0635\u06cc\u0644 \u0634\u0631\u062d \u0645\u06cc\u200c\u062f\u0647\u062f \u0648 \u06cc\u06a9 \u0641\u0631\u0648\u062f \u0646\u0631\u0645 \u0631\u0627 \u0628\u0627 \u0647\u0633\u062a\u0647 \u0633\u0627\u062f\u0647 \u0648 “\u0633\u0644\u0627\u0645 \u062c\u0647\u0627\u0646” \u0627\u0631\u0627\u0626\u0647 \u0645\u06cc\u200c\u062f\u0647\u062f. \u0631\u0627\u0647\u0646\u0645\u0627\u06cc \u0645\u062f\u0644 \u0628\u0631\u0646\u0627\u0645\u0647\u200c\u0646\u0648\u06cc\u0633\u06cc NKI \u0633\u0647 \u0645\u0631\u062d\u0644\u0647 \u06cc\u06a9 \u0647\u0633\u062a\u0647 \u0645\u0639\u0645\u0648\u0644\u06cc NKI (\u0628\u0627\u0631\u06af\u06cc\u0631\u06cc \u0648\u0631\u0648\u062f\u06cc\u200c\u0647\u0627\u060c \u0627\u062c\u0631\u0627\u06cc \u0639\u0645\u0644\u06cc\u0627\u062a \u0631\u0648\u06cc \u0645\u0648\u062a\u0648\u0631\u0647\u0627\u06cc \u0645\u062d\u0627\u0633\u0628\u0627\u062a\u06cc \u0648 \u0630\u062e\u06cc\u0631\u0647 \u062e\u0631\u0648\u062c\u06cc\u200c\u0647\u0627) \u0631\u0627 \u0634\u0631\u062d \u0645\u06cc\u200c\u062f\u0647\u062f \u0648 \u0639\u0645\u0644\u06cc\u0627\u062a NKI Tile \u0648 Tile-based \u0631\u0627 \u0645\u0639\u0631\u0641\u06cc \u0645\u06cc\u200c\u06a9\u0646\u062f. \u0622\u0645\u0648\u0632\u0634\u200c\u0647\u0627\u06cc NKI \u0627\u0646\u0648\u0627\u0639 \u0628\u0631\u0646\u0627\u0645\u0647\u200c\u0647\u0627\u06cc \u0646\u0645\u0648\u0646\u0647 \u0647\u0633\u062a\u0647 NKI \u0631\u0627 \u0646\u0634\u0627\u0646 \u0645\u06cc\u200c\u062f\u0647\u0646\u062f \u06a9\u0647 \u0647\u0631 \u06a9\u062f\u0627\u0645 API\u0647\u0627 \u0648 \u0642\u0627\u0628\u0644\u06cc\u062a\u200c\u0647\u0627\u06cc \u0627\u0635\u0644\u06cc NKI \u0631\u0627 \u0645\u0639\u0631\u0641\u06cc \u0645\u06cc\u200c\u06a9\u0646\u0646\u062f. \u0628\u0627 \u062a\u0648\u062c\u0647 \u0628\u0647 \u0628\u0647\u06cc\u0646\u0647 \u0641\u0631\u0636\u06cc \u0647\u0633\u062a\u0647\u200c\u0647\u0627\u06cc \u0646\u0645\u0648\u0646\u0647\u060c \u06cc\u06a9 \u0627\u0633\u062a\u0631\u0627\u062a\u0698\u06cc \u0645\u0645\u06a9\u0646 \u0628\u0631\u0627\u06cc \u062a\u0648\u0633\u0639\u0647 \u0647\u0633\u062a\u0647\u200c\u0647\u0627\u06cc \u062c\u062f\u06cc\u062f \u0645\u06cc\u200c\u062a\u0648\u0627\u0646\u062f \u0627\u06cc\u0646 \u0628\u0627\u0634\u062f \u06a9\u0647 1) \u0646\u0645\u0648\u0646\u0647\u200c\u0627\u06cc \u0631\u0627 \u0634\u0646\u0627\u0633\u0627\u06cc\u06cc \u06a9\u0646\u06cc\u062f \u06a9\u0647 \u0645\u0634\u0627\u0628\u0647 \u0639\u0645\u0644\u06cc\u0627\u062a\u06cc \u0627\u0633\u062a \u06a9\u0647 \u0645\u06cc\u200c\u062e\u0648\u0627\u0647\u06cc\u062f \u0627\u062c\u0631\u0627 \u06a9\u0646\u06cc\u062f \u0648 \u0633\u067e\u0633 2) \u0627\u0632 \u0622\u0646 \u0628\u0647 \u0639\u0646\u0648\u0627\u0646 \u062e\u0637 \u067e\u0627\u06cc\u0647 \u0627\u0633\u062a\u0641\u0627\u062f\u0647 \u06a9\u0646\u06cc\u062f \u0648 \u0628\u0647 \u0637\u0648\u0631 \u0645\u06a9\u0631\u0631 \u0622\u0646 \u0631\u0627 \u0627\u0635\u0644\u0627\u062d \u0648 \u062a\u0646\u0638\u06cc\u0645 \u06a9\u0646\u06cc\u062f. \u0628\u0647 \u0639\u0645\u0644\u06a9\u0631\u062f \u062e\u0627\u0635\u06cc \u06a9\u0647 \u0646\u06cc\u0627\u0632 \u062f\u0627\u0631\u06cc\u062f \u062f\u0633\u062a \u067e\u06cc\u062f\u0627 \u06a9\u0646\u06cc\u062f.<\/p>\n

NKI API Reference Manual API Python \u0628\u0631\u0627\u06cc \u062a\u0648\u0633\u0639\u0647 \u0647\u0633\u062a\u0647 \u0631\u0627 \u0634\u0631\u062d \u0645\u06cc \u062f\u0647\u062f. \u0628\u0627 \u0646\u062d\u0648 \u0648 \u0645\u0639\u0646\u0627\u06cc\u06cc \u06a9\u0647 \u0645\u0634\u0627\u0628\u0647 Triton \u0648 NumPy \u0627\u0633\u062a\u060c \u062a\u0639\u0631\u06cc\u0641 \u0632\u0628\u0627\u0646 NKI \u0628\u0647 \u062d\u062f\u0627\u06a9\u062b\u0631 \u0631\u0633\u0627\u0646\u062f\u0646 \u062f\u0633\u062a\u0631\u0633\u06cc \u0648 \u0633\u0647\u0648\u0644\u062a \u0627\u0633\u062a\u0641\u0627\u062f\u0647 \u0627\u0633\u062a. \u0628\u0627 \u0627\u06cc\u0646 \u062d\u0627\u0644\u060c \u062a\u0648\u062c\u0647 \u0628\u0647 \u0627\u06cc\u0646 \u0646\u06a9\u062a\u0647 \u0645\u0647\u0645 \u0627\u0633\u062a \u06a9\u0647 \u062a\u0648\u0633\u0639\u0647 \u0647\u0633\u062a\u0647 NKI \u0628\u0647 \u0639\u0645\u0644\u06cc\u0627\u062a \u062a\u0639\u0631\u06cc\u0641 \u0634\u062f\u0647 \u062f\u0631 \u06a9\u062a\u0627\u0628\u062e\u0627\u0646\u0647 NKI \u0645\u062d\u062f\u0648\u062f \u0645\u06cc \u0634\u0648\u062f\u060c \u06a9\u0647 (\u062a\u0627 \u0632\u0645\u0627\u0646 \u0646\u06af\u0627\u0631\u0634 \u0627\u06cc\u0646 \u0645\u0642\u0627\u0644\u0647) \u0646\u0633\u0628\u062a \u0628\u0647 \u06a9\u062a\u0627\u0628\u062e\u0627\u0646\u0647 \u0647\u0627\u06cc\u06cc \u0645\u0627\u0646\u0646\u062f Triton \u0648 NumPy \u06a9\u0645\u062a\u0631 \u0648 \u0645\u062d\u062f\u0648\u062f\u062a\u0631 \u0647\u0633\u062a\u0646\u062f.<\/p>\n

\u0646\u0645\u0648\u0646\u0647 \u0627\u0633\u0628\u0627\u0628 \u0628\u0627\u0632\u06cc – \u0647\u0633\u062a\u0647 GIOU<\/p>\n

\u0647\u0645\u0627\u0646\u0637\u0648\u0631 \u06a9\u0647 \u062f\u0631 \u067e\u0633\u062a\u200c\u0647\u0627\u06cc \u0642\u0628\u0644\u06cc\u200c\u0645\u0627\u0646\u060c \u0627\u0633\u062a\u0641\u0627\u062f\u0647 \u0627\u0632 NKI \u0631\u0627 \u0628\u0627 \u0627\u06cc\u062c\u0627\u062f \u06cc\u06a9 \u067e\u06cc\u0627\u062f\u0647\u200c\u0633\u0627\u0632\u06cc \u0633\u0641\u0627\u0631\u0634\u06cc \u0627\u0632 \u0639\u0645\u0644\u06cc\u0627\u062a \u062a\u0642\u0627\u0637\u0639 \u0639\u0645\u0648\u0645\u06cc \u0628\u0631 \u0631\u0648\u06cc \u0627\u062a\u062d\u0627\u062f\u06cc\u0647 (GIOU) \u0631\u0648\u06cc \u06cc\u06a9 \u062c\u0641\u062a \u062f\u0633\u062a\u0647 \u0627\u0632 \u062c\u0639\u0628\u0647\u200c\u0647\u0627\u06cc \u0648\u0631\u0648\u062f\u06cc \u0627\u0631\u0632\u06cc\u0627\u0628\u06cc \u0645\u06cc\u200c\u06a9\u0646\u06cc\u0645. \u0627\u0632 \u0622\u0646\u062c\u0627\u06cc\u06cc \u06a9\u0647 GIOU \u0634\u0627\u0645\u0644 \u0639\u0645\u0644\u06cc\u0627\u062a \u0645\u0628\u062a\u0646\u06cc \u0628\u0631 \u067e\u06cc\u06a9\u0633\u0644 \u0627\u0633\u062a\u060c \u0645\u0627 \u0627\u0632 \u0647\u0633\u062a\u0647 *exp* \u0627\u0632 \u0631\u0627\u0647\u0646\u0645\u0627\u06cc \u0628\u0631\u0646\u0627\u0645\u0647 \u0646\u0648\u06cc\u0633\u06cc NKI \u0628\u0647 \u0639\u0646\u0648\u0627\u0646 \u0646\u0642\u0637\u0647 \u0645\u0631\u062c\u0639 \u0627\u0633\u062a\u0641\u0627\u062f\u0647 \u06a9\u0631\u062f\u06cc\u0645 \u0648 \u0627\u0633\u062a\u0641\u0627\u062f\u0647 \u0627\u0632 \u0646\u0645\u0627\u06cc\u0647 \u0633\u0627\u0632\u06cc \u062a\u0627\u0646\u0633\u0648\u0631 \u067e\u06cc\u0634\u0631\u0641\u062a\u0647 NKI \u0631\u0627 \u062f\u0631 \u067e\u06cc\u0627\u062f\u0647 \u0633\u0627\u0632\u06cc \u062e\u0648\u062f \u0628\u0647 \u06a9\u0627\u0631 \u0628\u0631\u062f\u06cc\u0645. \u0628\u0631\u0627\u06cc \u062a\u0633\u0647\u06cc\u0644 \u0627\u0634\u06a9\u0627\u0644\u200c\u0632\u062f\u0627\u06cc\u06cc \u062f\u0631 \u06cc\u06a9 \u0645\u062d\u06cc\u0637 CPU\u060c \u0645\u0627 \u0647\u0645\u0686\u0646\u06cc\u0646 \u06af\u0632\u06cc\u0646\u0647\u200c\u0647\u0627\u06cc\u06cc \u0628\u0631\u0627\u06cc \u0627\u062c\u0631\u0627\u06cc \u06a9\u062f \u0628\u0627 \u0627\u0633\u062a\u0641\u0627\u062f\u0647 \u0627\u0632 API\u0647\u0627\u06cc nki.simulate_kernel \u0648 nki.language.device_print.html \u0627\u0636\u0627\u0641\u0647 \u06a9\u0631\u062f\u06cc\u0645.<\/p>\n

import torch
\nimport neuronxcc.nki as nki
\nimport neuronxcc.nki.language as nl
\nimport numpy as np<\/p>\n

simulate = False<\/p>\n

try:
\n # if torch libraries are installed assume that we are running on Neuron
\n import torch_xla.core.xla_model as xm
\n import torch_neuronx
\n from torch_neuronx import nki_jit<\/p>\n

device = xm.xla_device()<\/p>\n

# empty implementation
\n def debug_print(*args, **kwargs):
\n pass
\nexcept:
\n # if torch libraries are not installed assume that we are running on CPU
\n # and program script to use nki simulation
\n simulate = True
\n nki_jit = nki.trace
\n debug_print = nl.device_print
\n device=”cpu”<\/p>\n

@nki_jit
\ndef giou_kernel(preds_ptr,
\n targets_ptr,
\n output_ptr):
\n epsilon = 1e-5
\n TILE_M = nl.tile_size.pmax # 128
\n TILE_N = nl.tile_size.psum_fmax # 512
\n TILE_N_OUT = TILE_N \/\/ 4<\/p>\n

p_1, p_2 = preds_ptr.shape
\n t_1, t_2 = targets_ptr.shape
\n o_1, o_2 = output_ptr.shape<\/p>\n

# verify input
\n # batch size must be multiple of 128
\n assert p_1 % TILE_M == 0
\n assert p_1 == t_1
\n assert p_1 == o_1
\n # num boxes box *4 must be multiple of 512
\n assert p_2 % TILE_N == 0
\n assert p_2 == t_2
\n assert p_2 \/\/ 4 == o_2<\/p>\n

num_tiles_m = p_1 \/\/ TILE_M
\n num_tiles_n = p_2 \/\/ TILE_N<\/p>\n

# Generate tensors for advanced indexing
\n i_p = nl.arange(TILE_M)[:, None]\n i_f = nl.arange(TILE_N \/\/ 4)[None, :]\n i_f_0 = (4 * i_f)
\n i_f_1 = (4 * i_f + 1)
\n i_f_2 = (4 * i_f + 2)
\n i_f_3 = (4 * i_f + 3)<\/p>\n

# Use affine_range to loop over tiles
\n for m in nl.affine_range(num_tiles_m):
\n for n in nl.affine_range(num_tiles_n):
\n # Load input data from HBM
\n preds = nl.load(preds_ptr[m * TILE_M:(m + 1) * TILE_M,
\n n * TILE_N:(n + 1) * TILE_N])
\n targets = nl.load(targets_ptr[m * TILE_M:(m + 1) * TILE_M,
\n n * TILE_N:(n + 1) * TILE_N])
\n debug_print(‘preds’, preds)
\n preds_left = preds[i_p, i_f_0]\n preds_top = preds[i_p, i_f_1]\n preds_right = preds[i_p, i_f_2]\n preds_bottom = preds[i_p, i_f_3]\n

gt_left = targets[i_p, i_f_0]\n gt_top = targets[i_p, i_f_1]\n gt_right = targets[i_p, i_f_2]\n gt_bottom = targets[i_p, i_f_3]\n

# Compute the area of each box
\n area1 = (preds_right – preds_left) * (preds_bottom – preds_top)
\n area2 = (gt_right – gt_left) * (gt_bottom – gt_top)<\/p>\n

# Compute the intersection
\n left = nl.maximum(preds_left, gt_left)
\n top = nl.maximum(preds_top, gt_top)
\n right = nl.minimum(preds_right, gt_right)
\n bottom = nl.minimum(preds_bottom, gt_bottom)<\/p>\n

inter_w = nl.maximum(right – left, 0)
\n inter_h = nl.maximum(bottom – top, 0)
\n inter_area = inter_w * inter_h<\/p>\n

union_area = area1 + area2 – inter_area<\/p>\n

iou_val = inter_area \/ nl.maximum(union_area, epsilon)<\/p>\n

# Compute the smallest enclosing box
\n enclose_left = nl.minimum(preds_left, gt_left)
\n enclose_top = nl.minimum(preds_top, gt_top)
\n enclose_right = nl.maximum(preds_right, gt_right)
\n enclose_bottom = nl.maximum(preds_bottom, gt_bottom)<\/p>\n

enclose_w = nl.maximum(enclose_right – enclose_left, 0)
\n enclose_h = nl.maximum(enclose_bottom – enclose_top, 0)
\n enclose_area = enclose_w * enclose_h<\/p>\n

# Compute GIOU
\n delta_area = (enclose_area – union_area)
\n enclose_area = nl.maximum(enclose_area, epsilon)
\n giou = iou_val – delta_area \/ enclose_area<\/p>\n

# Store results
\n nl.store(output_ptr[m * TILE_M:(m + 1) * TILE_M,
\n n * TILE_N_OUT:(n + 1) * TILE_N_OUT],
\n giou)<\/p>\n

\u0628\u0631\u0627\u06cc \u0627\u062c\u0631\u0627\u06cc \u0647\u0633\u062a\u0647 GIOU \u062e\u0648\u062f\u060c \u062f\u0648 \u062f\u0633\u062a\u0647 \u0627\u0632 \u062c\u0639\u0628\u0647 \u0647\u0627\u06cc \u062a\u0635\u0627\u062f\u0641\u06cc \u062a\u0648\u0644\u06cc\u062f \u0645\u06cc \u06a9\u0646\u06cc\u0645 \u0648 \u0622\u0646\u0647\u0627 \u0631\u0627 \u0628\u0647 \u062a\u0627\u0628\u0639 \u062e\u0648\u062f \u062a\u063a\u0630\u06cc\u0647 \u0645\u06cc \u06a9\u0646\u06cc\u0645:<\/p>\n

# generate random data in np
\nnp.random.seed(0)
\nbatch_size = 1024
\nn_boxes = 256
\nimg_size = 256
\nboxes = []\n

for i in range(2):
\n # Randomly generate box sizes and positions
\n box_sizes = np.random.randint(1, img_size, size=(batch_size,n_boxes,2))
\n top_left = np.random.randint(0, img_size-1, size=(batch_size,n_boxes,2))
\n bottom_right = np.clip(top_left + box_sizes, 0, img_size – 1)<\/p>\n

# Concatenate top-left and bottom-right coordinates
\n rand_boxes = np.concatenate((top_left, bottom_right), axis=2)<\/p>\n

boxes.append(rand_boxes.astype(np.float32))<\/p>\n

out = np.empty((batch_size, n_boxes), np.float32)<\/p>\n

# convert tensors to PyTorch
\nt_boxes_0 = torch.tensor(boxes[0]).to(device)
\nt_boxes_1 = torch.tensor(boxes[1]).to(device)
\nt_out = torch.tensor(out).to(device)<\/p>\n

if simulate:
\n # the simulation API requires numpy input
\n nki.simulate_kernel(giou_kernel,
\n boxes[0].reshape((batch_size, -1)),
\n boxes[1].reshape((batch_size, -1)),
\n out)
\nelse:
\n giou_kernel(t_boxes_0.view((batch_size, -1)),
\n t_boxes_1.view((batch_size, -1)),
\n t_out)<\/p>\n

\u0628\u0631\u0627\u06cc \u0627\u0631\u0632\u06cc\u0627\u0628\u06cc \u0639\u0645\u0644\u06a9\u0631\u062f \u0647\u0633\u062a\u0647 NKI \u062e\u0648\u062f\u060c \u0622\u0646 \u0631\u0627 \u0628\u0627 \u0627\u062c\u0631\u0627\u06cc \u0633\u0627\u062f\u0647 GIOU \u0632\u06cc\u0631 \u062f\u0631 PyTorch \u0645\u0642\u0627\u06cc\u0633\u0647 \u0645\u06cc \u06a9\u0646\u06cc\u0645:<\/p>\n

def torch_giou(boxes1, boxes2):
\n # loosely based on torchvision generalized_box_iou_loss code
\n epsilon = 1e-5<\/p>\n

# Compute areas of both sets of boxes
\n area1 = (boxes1[…,2]-boxes1[…,0])*(boxes1[…,3]-boxes1[…,1])
\n area2 = (boxes2[…,2]-boxes2[…,0])*(boxes2[…,3]-boxes2[…,1])<\/p>\n

# Corners of intersection
\n lt = torch.max(boxes1[…, :2], boxes2[…, :2])
\n rb = torch.min(boxes1[…, 2:], boxes2[…, 2:])<\/p>\n

# Width and height of intersection
\n wh = (rb – lt).clamp(min=0)<\/p>\n

# Area of the intersection
\n inter = wh[…, 0] * wh[…, 1]\n

# Union of the two boxes
\n union = area1 + area2 – inter
\n iou = inter \/ union.clamp(epsilon)<\/p>\n

# Corners of enclosing box
\n lti = torch.min(boxes1[…, :2], boxes2[…, :2])
\n rbi = torch.max(boxes1[…, 2:], boxes2[…, 2:])<\/p>\n

# Width and height of the enclosing box
\n whi = (rbi – lti).clamp(min=0)<\/p>\n

# Area of the enclosing box
\n areai = (whi[…, 0] * whi[…, 1]).clamp(epsilon)<\/p>\n

return iou – (areai – union) \/ areai<\/p>\n

\u0645\u0627 \u0627\u0632 \u0627\u0628\u0632\u0627\u0631 \u0633\u0646\u062c\u0634 \u0632\u06cc\u0631 \u0628\u0631\u0627\u06cc \u0645\u0642\u0627\u06cc\u0633\u0647 \u0639\u0645\u0644\u06a9\u0631\u062f \u0632\u0645\u0627\u0646 \u0627\u062c\u0631\u0627 \u062f\u0648 \u0639\u0645\u0644\u06a9\u0631\u062f \u062e\u0648\u062f \u0627\u0633\u062a\u0641\u0627\u062f\u0647 \u0645\u06cc \u06a9\u0646\u06cc\u0645:<\/p>\n

import time
\ndef benchmark(f, warmup_iters=20, ntrials: int = 100):
\n def run(*args, **kwargs):
\n # warmup
\n for _ in range(warmup_iters):
\n f(*args, **kwargs)
\n start_time = time.time()
\n for _ in range(ntrials):
\n f(*args, **kwargs)
\n end_time = time.time()
\n # Calculate average time per iteration
\n avg_time = (end_time – start_time) \/ ntrials
\n return avg_time<\/p>\n

return run<\/p>\n

avg_time = benchmark(torch_giou)(t_boxes_0, t_boxes_1)
\nprint(f’torch_giou: {avg_time}’)<\/p>\n

avg_time = benchmark(giou_kernel)(t_boxes_0.view((batch_size, -1)),
\n t_boxes_1.view((batch_size, -1)),
\n t_out)
\nprint(f’giou_kernel: {avg_time}’)<\/p>\n

\u0645\u062d\u06cc\u0637 \u0632\u0645\u0627\u0646 \u0627\u062c\u0631\u0627<\/p>\n

\u0645\u0627 \u0627\u0633\u06a9\u0631\u06cc\u067e\u062a \u062e\u0648\u062f \u0631\u0627 \u0631\u0648\u06cc \u06cc\u06a9 \u0646\u0645\u0648\u0646\u0647 Amazon EC2 inf2.xlarge (\u0634\u0627\u0645\u0644 \u062f\u0648 \u0647\u0633\u062a\u0647 Neuron \u0648 \u0686\u0647\u0627\u0631 vCPU) \u0627\u062c\u0631\u0627 \u06a9\u0631\u062f\u06cc\u0645. \u0645\u0627 \u0627\u0632 \u062c\u062f\u06cc\u062f\u062a\u0631\u06cc\u0646 \u0646\u0633\u062e\u0647 Deep Learning AMI \u0628\u0631\u0627\u06cc Neuron \u06a9\u0647 \u062f\u0631 \u0632\u0645\u0627\u0646 \u0646\u06af\u0627\u0631\u0634 \u0627\u06cc\u0646 \u0645\u0642\u0627\u0644\u0647 \u0645\u0648\u062c\u0648\u062f \u0628\u0648\u062f\u060c “Deep Learning AMI Neuron (Ubuntu 22.04) 20241027” \u0628\u0627 AWS Neuron 2.20.1 \u0648 PyTorch 2.1 \u0627\u0633\u062a\u0641\u0627\u062f\u0647 \u06a9\u0631\u062f\u06cc\u0645.<\/p>\n

\u0646\u062a\u0627\u06cc\u062c<\/p>\n

\u0647\u0633\u062a\u0647 GIOU \u0633\u0641\u0627\u0631\u0634\u06cc \u0645\u0627 \u0645\u06cc\u0627\u0646\u06af\u06cc\u0646 \u0632\u0645\u0627\u0646 \u0627\u062c\u0631\u0627 0.211 \u0645\u06cc\u0644\u06cc \u062b\u0627\u0646\u06cc\u0647 \u0631\u0627 \u062f\u0631 \u0645\u0642\u0627\u06cc\u0633\u0647 \u0628\u0627 0.293 \u0646\u0634\u0627\u0646 \u0645\u06cc \u062f\u0647\u062f \u06a9\u0647 \u0628\u0647 \u0645\u06cc\u0632\u0627\u0646 39\u066a \u0627\u0641\u0632\u0627\u06cc\u0634 \u0639\u0645\u0644\u06a9\u0631\u062f \u0631\u0627 \u0646\u0634\u0627\u0646 \u0645\u06cc \u062f\u0647\u062f. \u0628\u0647 \u062e\u0627\u0637\u0631 \u062f\u0627\u0634\u062a\u0647 \u0628\u0627\u0634\u06cc\u062f \u06a9\u0647 \u0627\u06cc\u0646 \u0646\u062a\u0627\u06cc\u062c \u0628\u0631\u0627\u06cc \u0646\u0645\u0648\u0646\u0647 \u0627\u0633\u0628\u0627\u0628 \u0628\u0627\u0632\u06cc \u0645\u0627 \u0645\u0646\u062d\u0635\u0631 \u0628\u0647 \u0641\u0631\u062f \u0627\u0633\u062a. \u0633\u0627\u06cc\u0631 \u0639\u0645\u0644\u06af\u0631\u0647\u0627\u060c \u0628\u0647 \u0648\u06cc\u0698\u0647 \u0622\u0646\u0647\u0627\u06cc\u06cc \u06a9\u0647 \u0634\u0627\u0645\u0644 \u0636\u0631\u0628 \u0645\u0627\u062a\u0631\u06cc\u0633 \u0647\u0633\u062a\u0646\u062f (\u0648 \u0627\u0632 \u0645\u0648\u062a\u0648\u0631 Tensor \u0627\u0633\u062a\u0641\u0627\u062f\u0647 \u0645\u06cc \u06a9\u0646\u0646\u062f) \u0627\u062d\u062a\u0645\u0627\u0644\u0627\u064b \u0646\u062a\u0627\u06cc\u062c \u0645\u0642\u0627\u06cc\u0633\u0647 \u0627\u06cc \u0645\u062a\u0641\u0627\u0648\u062a\u06cc \u0631\u0627 \u0646\u0634\u0627\u0646 \u0645\u06cc \u062f\u0647\u0646\u062f.<\/p>\n

\u0628\u0647\u06cc\u0646\u0647 \u0633\u0627\u0632\u06cc \u0639\u0645\u0644\u06a9\u0631\u062f \u0647\u0633\u062a\u0647 NKI<\/p>\n

\u06af\u0627\u0645 \u0628\u0639\u062f\u06cc \u062f\u0631 \u062a\u0648\u0633\u0639\u0647 \u0647\u0633\u062a\u0647 \u0645\u0627 – \u0641\u0631\u0627\u062a\u0631 \u0627\u0632 \u0645\u062d\u062f\u0648\u062f\u0647 \u0627\u06cc\u0646 \u067e\u0633\u062a – \u062a\u062c\u0632\u06cc\u0647 \u0648 \u062a\u062d\u0644\u06cc\u0644 \u0639\u0645\u0644\u06a9\u0631\u062f \u0647\u0633\u062a\u0647 GIOU \u0628\u0627 \u0627\u0633\u062a\u0641\u0627\u062f\u0647 \u0627\u0632 Neuron Profiler \u0627\u062e\u062a\u0635\u0627\u0635\u06cc \u0628\u0647 \u0645\u0646\u0638\u0648\u0631 \u0634\u0646\u0627\u0633\u0627\u06cc\u06cc \u062a\u0646\u06af\u0646\u0627\u0647\u0627 \u0648 \u0628\u0647\u06cc\u0646\u0647 \u0633\u0627\u0632\u06cc \u067e\u06cc\u0627\u062f\u0647 \u0633\u0627\u0632\u06cc \u0645\u0627 \u062e\u0648\u0627\u0647\u062f \u0628\u0648\u062f. \u0644\u0637\u0641\u0627\u064b \u0631\u0627\u0647\u0646\u0645\u0627\u06cc \u0639\u0645\u0644\u06a9\u0631\u062f NKI \u0631\u0627 \u0628\u0631\u0627\u06cc \u062c\u0632\u0626\u06cc\u0627\u062a \u0628\u06cc\u0634\u062a\u0631 \u0628\u0628\u06cc\u0646\u06cc\u062f.<\/p>\n

Neuron Custom C++ Operators<\/p>\n

\u0631\u0648\u0634 \u062f\u0648\u0645 \u0628\u0631\u0627\u06cc \u0627\u06cc\u062c\u0627\u062f \u06cc\u06a9 \u0647\u0633\u062a\u0647 Neuron \u0633\u0641\u0627\u0631\u0634\u06cc\u060c \u0633\u0627\u062e\u062a \u06cc\u06a9 \u0627\u067e\u0631\u0627\u062a\u0648\u0631 C++ \u0628\u0631\u0627\u06cc \u0645\u0648\u062a\u0648\u0631 GpSimd \u0627\u0633\u062a. \u0627\u06cc\u0646 \u0631\u0648\u0634 \u062f\u0631 Neuron Custom C++ Operators Developer Guide \u0648 \u062f\u0631 Neuron Custom C++ Operators \u062f\u0631 MLP \u0648 Neuron Custom C++ Operators Performance Optimization \u0646\u0645\u0627\u06cc\u0634 \u062f\u0627\u062f\u0647 \u0634\u062f\u0647 \u0627\u0633\u062a.<\/p>\n

Neuron Custom C++ Operators \u0628\u0627 \u062a\u0633\u0647\u06cc\u0644 \u062a\u0631\u06a9\u06cc\u0628 \u0686\u0646\u062f\u06cc\u0646 \u0639\u0645\u0644\u06cc\u0627\u062a \u0633\u0637\u062d \u067e\u0627\u06cc\u06cc\u0646 \u062f\u0631 \u06cc\u06a9 \u0647\u0633\u062a\u0647 \u0648\u0627\u062d\u062f\u060c \u0641\u0631\u0635\u062a\u06cc \u0631\u0627 \u0628\u0631\u0627\u06cc “\u062a\u0644\u0641\u06cc\u0642 \u0647\u0633\u062a\u0647” \u062f\u0631 \u0645\u0648\u062a\u0648\u0631 GpSimd \u0627\u0631\u0627\u0626\u0647 \u0645\u06cc \u062f\u0647\u062f. \u0627\u06cc\u0646 \u0631\u0648\u06cc\u06a9\u0631\u062f \u0645\u06cc\u200c\u062a\u0648\u0627\u0646\u062f \u0647\u0632\u06cc\u0646\u0647\u200c\u0647\u0627\u06cc \u0633\u0631\u0628\u0627\u0631 \u0645\u0631\u062a\u0628\u0637 \u0628\u0627 \u0645\u0648\u0627\u0631\u062f \u0632\u06cc\u0631 \u0631\u0627 \u0628\u0647 \u0645\u06cc\u0632\u0627\u0646 \u0642\u0627\u0628\u0644 \u062a\u0648\u062c\u0647\u06cc \u06a9\u0627\u0647\u0634 \u062f\u0647\u062f: 1) \u0628\u0627\u0631\u06af\u06cc\u0631\u06cc \u0686\u0646\u062f\u06cc\u0646 \u0647\u0633\u062a\u0647 \u062c\u062f\u0627\u06af\u0627\u0646\u0647\u060c \u0648 2) \u0627\u0646\u062a\u0642\u0627\u0644 \u062f\u0627\u062f\u0647 \u0628\u06cc\u0646 \u0645\u0646\u0627\u0637\u0642 \u0645\u062e\u062a\u0644\u0641 \u062d\u0627\u0641\u0638\u0647.<\/p>\n

\u0645\u062b\u0627\u0644 \u0627\u0633\u0628\u0627\u0628 \u0628\u0627\u0632\u06cc – \u0647\u0633\u062a\u0647 GIOU C++<\/p>\n

\u062f\u0631 \u0628\u0644\u0648\u06a9 \u06a9\u062f \u0632\u06cc\u0631 \u06cc\u06a9 \u0639\u0645\u0644\u06af\u0631 C++ GIOU \u0631\u0627 \u0628\u0631\u0627\u06cc Neuron \u067e\u06cc\u0627\u062f\u0647 \u0633\u0627\u0632\u06cc \u0645\u06cc \u06a9\u0646\u06cc\u0645 \u0648 \u0622\u0646 \u0631\u0627 \u062f\u0631 \u0641\u0627\u06cc\u0644\u06cc \u0628\u0647 \u0646\u0627\u0645 \u0630\u062e\u06cc\u0631\u0647 \u0645\u06cc \u06a9\u0646\u06cc\u0645 giou.cpp. \u0647\u0633\u062a\u0647 \u0645\u0627 \u0627\u0632 \u062f\u0633\u062a\u0631\u0633\u06cc TCM \u0628\u0631\u0627\u06cc \u0628\u0647\u06cc\u0646\u0647 \u0633\u0627\u0632\u06cc \u0639\u0645\u0644\u06a9\u0631\u062f \u062e\u0648\u0627\u0646\u062f\u0646 \u0648 \u0646\u0648\u0634\u062a\u0646 \u062d\u0627\u0641\u0638\u0647 \u0627\u0633\u062a\u0641\u0627\u062f\u0647 \u0645\u06cc \u06a9\u0646\u062f \u0648 \u062a\u0646\u0638\u06cc\u0645 *\u0686\u0646\u062f \u0647\u0633\u062a\u0647 \u0627\u06cc * \u0631\u0627 \u0628\u0631\u0627\u06cc \u0627\u0633\u062a\u0641\u0627\u062f\u0647 \u0627\u0632 \u0647\u0631 \u0647\u0634\u062a \u067e\u0631\u062f\u0627\u0632\u0646\u062f\u0647 \u062f\u0627\u062e\u0644\u06cc GpSimd \u0627\u0639\u0645\u0627\u0644 \u0645\u06cc \u06a9\u0646\u062f.<\/p>\n

#include <stdint.h>
\n#include <stdlib.h>
\n#include <torch\/torch.h>
\n#include <neuron\/neuron-utils.hpp>
\n#include <algorithm><\/p>\n

\/\/ input boxes of shape 1024x256x4
\n\/\/ output scores of shape 1024×256
\ntorch::Tensor giou(const torch::Tensor& t_pred,
\n const torch::Tensor& t_target) {
\n size_t num_samples = t_pred.sizes()[0];
\n size_t num_boxes = t_pred.sizes()[1];
\n torch::Tensor t_out = get_dst_tensor();<\/p>\n

\/\/ get the number of GpSimd processors (8 in NeuronCoreV2)
\n uint32_t cpu_count = get_cpu_count();
\n \/\/ get index of current processor
\n uint32_t cpu_id = get_cpu_id();<\/p>\n

\/\/ divide the batch size into 8 partitions
\n uint32_t partition = num_samples \/ cpu_count;<\/p>\n

\/\/ use tcm buffers to load and write data
\n size_t tcm_in_size = num_boxes*4;
\n size_t tcm_out_size = num_boxes;
\n float *tcm_pred = (float*)torch::neuron::tcm_malloc(
\n sizeof(float)*tcm_in_size);
\n float *tcm_target = (float*)torch::neuron::tcm_malloc(
\n sizeof(float)*tcm_in_size);
\n float *tcm_output = (float*)torch::neuron::tcm_malloc(
\n sizeof(float)*tcm_in_size);
\n auto t_pred_tcm_acc = t_pred.tcm_accessor();
\n auto t_target_tcm_acc = t_target.tcm_accessor();
\n auto t_out_tcm_acc = t_out.tcm_accessor();<\/p>\n

\/\/ iterate over each of the entries in the partition
\n for (size_t i = 0; i < partition; i++) {
\n \/\/ load the pred and target boxes into local memory
\n t_pred_tcm_acc.tensor_to_tcm<float>(tcm_pred,
\n partition*cpu_id + i*tcm_in_size,
\n tcm_in_size);
\n t_target_tcm_acc.tensor_to_tcm<float>(tcm_target,
\n partition*cpu_id + i*tcm_in_size,
\n tcm_in_size);<\/p>\n

\/\/ iterate over each of the boxes in the entry
\n for (size_t j = 0; j < num_boxes; j++) {
\n const float epsilon = 1e-5;
\n const float* box1 = &tcm_pred[j * 4];
\n const float* box2 = &tcm_target[j * 4];
\n \/\/ Compute area of each box
\n float area1 = (box1[2] – box1[0]) * (box1[3] – box1[1]);
\n float area2 = (box2[2] – box2[0]) * (box2[3] – box2[1]);<\/p>\n

\/\/ Compute the intersection
\n float left = std::max(box1[0], box2[0]);
\n float top = std::max(box1[1], box2[1]);
\n float right = std::min(box1[2], box2[2]);
\n float bottom = std::min(box1[3], box2[3]);<\/p>\n

float inter_w = std::max(right – left, 0.f);
\n float inter_h = std::max(bottom – top, 0.f);
\n float inter_area = inter_w * inter_h;<\/p>\n

\/\/ Compute the union area
\n float union_area = area1 + area2 – inter_area;<\/p>\n

\/\/ IoU
\n float iou_val = inter_area \/ std::max(union_area, epsilon);<\/p>\n

\/\/ Compute the smallest enclosing box
\n float enclose_left = std::min(box1[0], box2[0]);
\n float enclose_top = std::min(box1[1], box2[1]);
\n float enclose_right = std::max(box1[2], box2[2]);
\n float enclose_bottom = std::max(box1[3], box2[3]);<\/p>\n

float enclose_w = std::max(enclose_right – enclose_left, 0.f);
\n float enclose_h = std::max(enclose_bottom – enclose_top, 0.f);
\n float enclose_area = std::max(enclose_w * enclose_h, epsilon);<\/p>\n

float result = iou_val – (enclose_area-union_area)\/enclose_area;
\n tcm_output[j] = result;
\n }<\/p>\n

\/\/ write the giou scores of all boxes in the current entry
\n t_out_tcm_acc.tcm_to_tensor<float>(tcm_output,
\n partition*cpu_id + i*tcm_out_size,
\n tcm_out_size);
\n }<\/p>\n

torch::neuron::tcm_free(tcm_pred);
\n torch::neuron::tcm_free(tcm_target);
\n return t_out;
\n}<\/p>\n

\u0645\u0627 \u0646\u06cc\u0627\u0632 \u0628\u0647 \u06cc\u06a9 \u062c\u062f\u0627\u06af\u0627\u0646\u0647 \u062f\u0627\u0631\u06cc\u0645 shape.cpp \u0641\u0627\u06cc\u0644\u06cc \u06a9\u0647 \u0634\u06a9\u0644 \u062e\u0631\u0648\u062c\u06cc \u062a\u0627\u0628\u0639 GIOU \u0645\u0627 \u0631\u0627 \u062a\u0639\u0631\u06cc\u0641 \u0645\u06cc \u06a9\u0646\u062f \u0648 \u0639\u0645\u0644\u06af\u0631 \u0633\u0641\u0627\u0631\u0634\u06cc \u0645\u0627 \u0631\u0627 \u0628\u0627 \u06a9\u062a\u0627\u0628\u062e\u0627\u0646\u0647 Neuron \u062b\u0628\u062a \u0645\u06cc \u06a9\u0646\u062f:<\/p>\n

#include <stdint.h>
\n#include <stdlib.h>
\n#include <torch\/torch.h>
\n#include “torchneuron\/register.h”<\/p>\n

torch::Tensor giou_shape(torch::Tensor boxes1, torch::Tensor boxes2) {
\n torch::Tensor t_out = torch::zeros({boxes1.sizes()[0],
\n boxes1.sizes()[1]},
\n torch::kFloat);
\n return t_out;
\n}<\/p>\n

NEURON_LIBRARY(my_ops, m) {
\n m.def(“giou”, &giou_shape, “giou”);
\n}<\/p>\n

\u0631\u0627 build.py \u0627\u0633\u06a9\u0631\u06cc\u067e\u062a \u0639\u0645\u0644\u06af\u0631 C++ \u0631\u0627 \u06a9\u0627\u0645\u067e\u0627\u06cc\u0644 \u0645\u06cc \u06a9\u0646\u062f \u0648 \u0622\u0646 \u0631\u0627 \u0628\u0647 \u0639\u0646\u0648\u0627\u0646 \u06cc\u06a9 API \u067e\u0627\u06cc\u062a\u0648\u0646 \u0646\u0645\u0627\u06cc\u0634 \u0645\u06cc \u062f\u0647\u062f:<\/p>\n

import os
\nimport torch_neuronx
\nfrom torch_neuronx.xla_impl import custom_op<\/p>\n

custom_op.load(
\n name=”giou”,
\n compute_srcs=[‘giou.cpp’],
\n shape_srcs=[‘shape.cpp’],
\n build_directory=os.getcwd(),
\n multicore=True,
\n verbose=True
\n)<\/p>\n

\u0627\u0633\u06a9\u0631\u06cc\u067e\u062a \u06a9\u0627\u0645\u067e\u0627\u06cc\u0644 \u06cc\u06a9 \u06a9\u062a\u0627\u0628\u062e\u0627\u0646\u0647 *libgiou.so * \u0627\u06cc\u062c\u0627\u062f \u0645\u06cc \u06a9\u0646\u062f \u06a9\u0647 \u0634\u0627\u0645\u0644 \u0627\u062c\u0631\u0627\u06cc \u0639\u0645\u0644\u06af\u0631 C++ GIOU \u0645\u0627 \u0627\u0633\u062a. \u062f\u0631 \u0628\u0644\u0648\u06a9 \u06a9\u062f \u0632\u06cc\u0631\u060c \u06a9\u062a\u0627\u0628\u062e\u0627\u0646\u0647 \u0631\u0627 \u0628\u0627\u0631\u06af\u0630\u0627\u0631\u06cc \u0645\u06cc \u06a9\u0646\u06cc\u0645 \u0648 \u0639\u0645\u0644\u06a9\u0631\u062f \u0647\u0633\u062a\u0647 \u0633\u0641\u0627\u0631\u0634\u06cc \u062e\u0648\u062f \u0631\u0627 \u0628\u0627 \u0627\u0633\u062a\u0641\u0627\u062f\u0647 \u0627\u0632 \u0627\u0628\u0632\u0627\u0631 \u0633\u0646\u062c\u0634 \u062a\u0639\u0631\u06cc\u0641 \u0634\u062f\u0647 \u062f\u0631 \u0628\u0627\u0644\u0627 \u0627\u0646\u062f\u0627\u0632\u0647 \u06af\u06cc\u0631\u06cc \u0645\u06cc \u06a9\u0646\u06cc\u0645:<\/p>\n

from torch_neuronx.xla_impl import custom_op
\ncustom_op.load_library(‘libgiou.so’)<\/p>\n

avg_time = benchmark(torch.ops.my_ops.giou)(t_boxes_0, t_boxes_1)
\nprint(f’C++ giou: {avg_time}’)<\/p>\n

\u0645\u062d\u06cc\u0637 \u0632\u0645\u0627\u0646 \u0627\u062c\u0631\u0627<\/p>\n

\u0645\u0627 \u0627\u0632 \u0647\u0645\u0627\u0646 \u0645\u062d\u06cc\u0637 Neuron \u0627\u0632 \u0622\u0632\u0645\u0627\u06cc\u0634\u200c\u0647\u0627\u06cc NKI \u062e\u0648\u062f \u0628\u0631\u0627\u06cc \u06a9\u0627\u0645\u067e\u0627\u06cc\u0644 \u0648 \u0622\u0632\u0645\u0627\u06cc\u0634 \u0647\u0633\u062a\u0647 C ++ \u062e\u0648\u062f \u0627\u0633\u062a\u0641\u0627\u062f\u0647 \u06a9\u0631\u062f\u06cc\u0645. \u0644\u0637\u0641\u0627\u064b \u0628\u0647 \u0645\u0631\u0627\u062d\u0644 \u0646\u0635\u0628\u06cc \u06a9\u0647 \u0628\u0631\u0627\u06cc \u062a\u0648\u0633\u0639\u0647 \u0627\u067e\u0631\u0627\u062a\u0648\u0631 C++ \u0633\u0641\u0627\u0631\u0634\u06cc \u0644\u0627\u0632\u0645 \u0627\u0633\u062a \u062a\u0648\u062c\u0647 \u06a9\u0646\u06cc\u062f.<\/p>\n

\u0646\u062a\u0627\u06cc\u062c<\/p>\n

\u0647\u0633\u062a\u0647 C++ GIOU \u0645\u0627 \u0645\u06cc\u0627\u0646\u06af\u06cc\u0646 \u0632\u0645\u0627\u0646 \u0627\u062c\u0631\u0627 0.061 \u0645\u06cc\u0644\u06cc \u062b\u0627\u0646\u06cc\u0647 \u0631\u0627 \u0646\u0634\u0627\u0646 \u062f\u0627\u062f – \u062a\u0642\u0631\u06cc\u0628\u0627\u064b \u067e\u0646\u062c \u0628\u0631\u0627\u0628\u0631 \u0633\u0631\u06cc\u0639\u062a\u0631 \u0627\u0632 \u0627\u062c\u0631\u0627\u06cc \u062e\u0637 \u067e\u0627\u06cc\u0647 \u0645\u0627. \u0627\u06cc\u0646 \u0627\u062d\u062a\u0645\u0627\u0644\u0627\u064b \u0646\u062a\u06cc\u062c\u0647 “\u0647\u0645\u062c\u0648\u0634\u06cc \u0647\u0633\u062a\u0647” \u0627\u0633\u062a\u060c \u0647\u0645\u0627\u0646\u0637\u0648\u0631 \u06a9\u0647 \u062f\u0631 \u0628\u0627\u0644\u0627 \u0645\u0648\u0631\u062f \u0628\u062d\u062b \u0642\u0631\u0627\u0631 \u06af\u0631\u0641\u062a.<\/p>\n

\u0646\u062a\u06cc\u062c\u0647 \u06af\u06cc\u0631\u06cc<\/p>\n

\u062c\u062f\u0648\u0644 \u0632\u06cc\u0631 \u0646\u062a\u0627\u06cc\u062c \u0632\u0645\u0627\u0646 \u0627\u062c\u0631\u0627 \u0622\u0632\u0645\u0627\u06cc\u0634 \u0647\u0627\u06cc \u0645\u0627 \u0631\u0627 \u062e\u0644\u0627\u0635\u0647 \u0645\u06cc \u06a9\u0646\u062f.
\n\u0645\u06cc\u0627\u0646\u06af\u06cc\u0646 \u0632\u0645\u0627\u0646 \u067e\u06cc\u0627\u062f\u0647 \u0633\u0627\u0632\u06cc \u0647\u0627\u06cc \u0645\u062e\u062a\u0644\u0641 GIOU (\u06a9\u0645\u062a\u0631 \u0628\u0647\u062a\u0631 \u0627\u0633\u062a) – \u062a\u0648\u0633\u0637 \u0646\u0648\u06cc\u0633\u0646\u062f\u0647<\/p>\n

\u0644\u0637\u0641\u0627\u064b \u0628\u0647 \u062e\u0627\u0637\u0631 \u062f\u0627\u0634\u062a\u0647 \u0628\u0627\u0634\u06cc\u062f \u06a9\u0647 \u0627\u06cc\u0646 \u0646\u062a\u0627\u06cc\u062c \u0645\u062e\u062a\u0635 \u0646\u0645\u0648\u0646\u0647 \u0627\u0633\u0628\u0627\u0628 \u0628\u0627\u0632\u06cc \u0648 \u0645\u062d\u06cc\u0637 \u0632\u0645\u0627\u0646 \u0627\u062c\u0631\u0627 \u0645\u0648\u0631\u062f \u0627\u0633\u062a\u0641\u0627\u062f\u0647 \u062f\u0631 \u0627\u06cc\u0646 \u0645\u0637\u0627\u0644\u0639\u0647 \u0627\u0633\u062a. \u0646\u062a\u0627\u06cc\u062c \u0645\u0642\u0627\u06cc\u0633\u0647 \u0633\u0627\u06cc\u0631 \u0647\u0633\u062a\u0647 \u0647\u0627 \u0645\u0645\u06a9\u0646 \u0627\u0633\u062a \u0628\u0633\u06cc\u0627\u0631 \u0645\u062a\u0641\u0627\u0648\u062a \u0628\u0627\u0634\u062f – \u0628\u0633\u062a\u0647 \u0628\u0647 \u062f\u0631\u062c\u0647 \u0627\u06cc \u06a9\u0647 \u0622\u0646\u0647\u0627 \u0645\u06cc \u062a\u0648\u0627\u0646\u0646\u062f \u0627\u0632 \u0645\u0648\u062a\u0648\u0631\u0647\u0627\u06cc \u0645\u062d\u0627\u0633\u0628\u0627\u062a\u06cc \u062f\u0627\u062e\u0644\u06cc \u0647\u0633\u062a\u0647 \u0646\u0648\u0631\u0648\u0646 \u0627\u0633\u062a\u0641\u0627\u062f\u0647 \u06a9\u0646\u0646\u062f.<\/p>\n

\u062c\u062f\u0648\u0644 \u0632\u06cc\u0631 \u0628\u0631\u062e\u06cc \u0627\u0632 \u062a\u0641\u0627\u0648\u062a\u200c\u0647\u0627\u06cc\u06cc \u0631\u0627 \u06a9\u0647 \u0628\u06cc\u0646 \u062f\u0648 \u0631\u0648\u0634 \u0633\u0641\u0627\u0631\u0634\u06cc\u200c\u0633\u0627\u0632\u06cc \u0647\u0633\u062a\u0647 \u0646\u0648\u0631\u0648\u0646 AWS \u0645\u0634\u0627\u0647\u062f\u0647 \u06a9\u0631\u062f\u06cc\u0645\u060c \u062e\u0644\u0627\u0635\u0647 \u0645\u06cc\u200c\u06a9\u0646\u062f.<\/p>\n

\u0645\u0642\u0627\u06cc\u0633\u0647 \u0628\u06cc\u0646 \u0627\u0628\u0632\u0627\u0631 \u0633\u0641\u0627\u0631\u0634\u06cc \u0633\u0627\u0632\u06cc \u0647\u0633\u062a\u0647 (\u062a\u0648\u0633\u0637 \u0646\u0648\u06cc\u0633\u0646\u062f\u0647)<\/p>\n

\u0627\u0632 \u0637\u0631\u06cc\u0642 \u0631\u0627\u0628\u0637 \u0633\u0637\u062d \u0628\u0627\u0644\u0627\u06cc \u067e\u0627\u06cc\u062a\u0648\u0646\u060c API \u0647\u0627\u06cc NKI\u060c \u0642\u062f\u0631\u062a \u0645\u0648\u062a\u0648\u0631\u0647\u0627\u06cc \u0634\u062a\u0627\u0628 \u062f\u0647\u0646\u062f\u0647 \u0646\u0648\u0631\u0648\u0646 \u0631\u0627 \u0628\u0647 \u0634\u06a9\u0644\u06cc \u062f\u0631 \u062f\u0633\u062a\u0631\u0633 \u0648 \u06a9\u0627\u0631\u0628\u0631\u067e\u0633\u0646\u062f \u062f\u0631 \u0627\u062e\u062a\u06cc\u0627\u0631 \u062a\u0648\u0633\u0639\u0647 \u062f\u0647\u0646\u062f\u06af\u0627\u0646 ML \u0642\u0631\u0627\u0631 \u0645\u06cc \u062f\u0647\u0646\u062f. \u06a9\u062a\u0627\u0628\u062e\u0627\u0646\u0647 \u0633\u0637\u062d \u067e\u0627\u06cc\u06cc\u0646 C++ Custom Operators \u0628\u0631\u0646\u0627\u0645\u0647 \u0631\u06cc\u0632\u06cc \u062d\u062a\u06cc \u0628\u06cc\u0634\u062a\u0631 \u0631\u0627 \u0627\u0645\u06a9\u0627\u0646 \u067e\u0630\u06cc\u0631 \u0645\u06cc \u06a9\u0646\u062f\u060c \u0627\u0645\u0627 \u0628\u0647 \u0645\u0648\u062a\u0648\u0631 GpSimd \u0645\u062d\u062f\u0648\u062f \u0645\u06cc \u0634\u0648\u062f. \u0628\u0627 \u062a\u0631\u06a9\u06cc\u0628 \u0645\u0648\u062b\u0631 \u0647\u0631 \u062f\u0648 \u0627\u0628\u0632\u0627\u0631\u060c \u062a\u0648\u0633\u0639\u0647 \u062f\u0647\u0646\u062f\u06af\u0627\u0646 \u0645\u06cc \u062a\u0648\u0627\u0646\u0646\u062f \u0628\u0647 \u0637\u0648\u0631 \u06a9\u0627\u0645\u0644 \u0627\u0632 \u0642\u0627\u0628\u0644\u06cc\u062a \u0647\u0627\u06cc \u0645\u0639\u0645\u0627\u0631\u06cc AWS Neuron \u0628\u0647\u0631\u0647 \u0628\u0628\u0631\u0646\u062f.<\/p>\n

\u062e\u0644\u0627\u0635\u0647<\/p>\n

\u0628\u0627 \u0631\u0634\u062f \u06a9\u0627\u0645\u0644 \u0627\u0646\u0642\u0644\u0627\u0628 \u0647\u0648\u0634 \u0645\u0635\u0646\u0648\u0639\u06cc\u060c \u0628\u0633\u06cc\u0627\u0631\u06cc \u0627\u0632 \u0634\u0631\u06a9\u062a\u200c\u0647\u0627 \u062f\u0631 \u062d\u0627\u0644 \u062a\u0648\u0633\u0639\u0647 \u062a\u0631\u0627\u0634\u0647\u200c\u0647\u0627\u06cc \u067e\u06cc\u0634\u0631\u0641\u062a\u0647 \u0647\u0648\u0634 \u0645\u0635\u0646\u0648\u0639\u06cc \u0628\u0631\u0627\u06cc \u067e\u0627\u0633\u062e\u06af\u0648\u06cc\u06cc \u0628\u0647 \u062a\u0642\u0627\u0636\u0627\u06cc \u0631\u0648 \u0628\u0647 \u0631\u0634\u062f \u0628\u0631\u0627\u06cc \u0645\u062d\u0627\u0633\u0628\u0627\u062a \u0647\u0633\u062a\u0646\u062f. \u062f\u0631 \u062d\u0627\u0644\u06cc \u06a9\u0647 \u0627\u0637\u0644\u0627\u0639\u06cc\u0647 \u0647\u0627\u06cc \u0639\u0645\u0648\u0645\u06cc \u0627\u063a\u0644\u0628 \u0639\u0645\u0644\u06a9\u0631\u062f \u0627\u06cc\u0646 \u062a\u0631\u0627\u0634\u0647 \u0647\u0627\u060c \u0635\u0631\u0641\u0647 \u062c\u0648\u06cc\u06cc \u062f\u0631 \u0647\u0632\u06cc\u0646\u0647 \u0648 \u0628\u0647\u0631\u0647 \u0648\u0631\u06cc \u0627\u0646\u0631\u0698\u06cc \u0631\u0627 \u0628\u0631\u062c\u0633\u062a\u0647 \u0645\u06cc \u06a9\u0646\u062f\u060c \u0686\u0646\u062f\u06cc\u0646 \u0642\u0627\u0628\u0644\u06cc\u062a \u0627\u0635\u0644\u06cc \u0636\u0631\u0648\u0631\u06cc \u0627\u0633\u062a \u062a\u0627 \u0627\u06cc\u0646 \u062a\u0631\u0627\u0634\u0647 \u0647\u0627 \u0648 \u067e\u0634\u062a\u0647 \u0647\u0627\u06cc \u0646\u0631\u0645 \u0627\u0641\u0632\u0627\u0631\u06cc \u0622\u0646\u0647\u0627 \u0648\u0627\u0642\u0639\u0627\u064b \u0628\u0631\u0627\u06cc \u062a\u0648\u0633\u0639\u0647 ML \u0642\u0627\u0628\u0644 \u062f\u0648\u0627\u0645 \u0628\u0627\u0634\u0646\u062f. \u0627\u06cc\u0646 \u0642\u0627\u0628\u0644\u06cc\u062a \u0647\u0627 \u0634\u0627\u0645\u0644 \u0627\u0628\u0632\u0627\u0631\u0647\u0627\u06cc \u0627\u0634\u06a9\u0627\u0644 \u0632\u062f\u0627\u06cc\u06cc \u0642\u0648\u06cc\u060c \u0627\u0628\u0632\u0627\u0631\u0647\u0627\u06cc \u062a\u062c\u0632\u06cc\u0647 \u0648 \u062a\u062d\u0644\u06cc\u0644 \u0639\u0645\u0644\u06a9\u0631\u062f \u0648 \u0628\u0647\u06cc\u0646\u0647 \u0633\u0627\u0632\u06cc\u060c \u0642\u0627\u0628\u0644\u06cc\u062a \u0628\u0631\u0646\u0627\u0645\u0647 \u0631\u06cc\u0632\u06cc \u0648 \u063a\u06cc\u0631\u0647 \u0627\u0633\u062a.<\/p>\n

\u062f\u0631 \u0627\u06cc\u0646 \u067e\u0633\u062a\u060c \u0645\u0627 \u0628\u0631 \u0627\u0628\u0632\u0627\u0631\u0647\u0627\u06cc \u0645\u0648\u062c\u0648\u062f \u0628\u0631\u0627\u06cc \u0628\u0631\u0646\u0627\u0645\u0647\u200c\u0646\u0648\u06cc\u0633\u06cc \u0634\u062a\u0627\u0628\u200c\u062f\u0647\u0646\u062f\u0647\u200c\u0647\u0627\u06cc \u0647\u0648\u0634 \u0645\u0635\u0646\u0648\u0639\u06cc \u0628\u0648\u0645\u06cc AWS\u060c Trainium \u0648 Inferentia \u062a\u0645\u0631\u06a9\u0632 \u06a9\u0631\u062f\u06cc\u0645 \u0648 \u0627\u0633\u062a\u0641\u0627\u062f\u0647 \u0627\u0632 \u0622\u0646\u0647\u0627 \u0631\u0627 \u062f\u0631 \u0633\u0627\u062e\u062a \u0639\u0645\u0644\u06cc\u0627\u062a ML \u0633\u0641\u0627\u0631\u0634\u06cc \u0646\u0634\u0627\u0646 \u062f\u0627\u062f\u06cc\u0645. \u0627\u06cc\u0646 \u0627\u0628\u0632\u0627\u0631\u0647\u0627 \u062a\u0648\u0633\u0639\u0647 \u062f\u0647\u0646\u062f\u06af\u0627\u0646 \u0631\u0627 \u0642\u0627\u062f\u0631 \u0645\u06cc \u0633\u0627\u0632\u062f \u062a\u0627 \u0639\u0645\u0644\u06a9\u0631\u062f \u0645\u062f\u0644 \u0647\u0627\u06cc ML \u062e\u0648\u062f \u0631\u0627 \u0628\u0631 \u0631\u0648\u06cc \u062a\u0631\u0627\u0634\u0647 \u0647\u0627\u06cc \u0647\u0648\u0634 \u0645\u0635\u0646\u0648\u0639\u06cc AWS \u0628\u0647\u06cc\u0646\u0647 \u06a9\u0646\u0646\u062f \u0648 \u0641\u0631\u0635\u062a \u0647\u0627\u06cc \u062c\u062f\u06cc\u062f\u06cc \u0631\u0627 \u0628\u0631\u0627\u06cc \u0646\u0648\u0622\u0648\u0631\u06cc \u0648 \u062e\u0644\u0627\u0642\u06cc\u062a \u0628\u0627\u0632 \u06a9\u0646\u0646\u062f.<\/p>\n

\n
\n

\u0641\u0647\u0631\u0633\u062a \u0645\u0637\u0627\u0644\u0628<\/p>\nToggle<\/span><\/path><\/svg><\/svg><\/span><\/span><\/span><\/a><\/span><\/div>\n