{"id":416011,"date":"2024-06-30T00:41:40","date_gmt":"2024-06-30T00:41:40","guid":{"rendered":"http:\/\/savepearlharbor.com\/?p=416011"},"modified":"-0001-11-30T00:00:00","modified_gmt":"-0001-11-29T21:00:00","slug":"","status":"publish","type":"post","link":"https:\/\/savepearlharbor.com\/?p=416011","title":{"rendered":"<span>Triton: Open Source \u044f\u0437\u044b\u043a \u0434\u043b\u044f \u044f\u0434\u0435\u0440 Deep Learning<\/span>"},"content":{"rendered":"<div><!--[--><!--]--><\/div>\n<div id=\"post-content-body\">\n<div>\n<div class=\"article-formatted-body article-formatted-body article-formatted-body_version-2\">\n<div xmlns=\"http:\/\/www.w3.org\/1999\/xhtml\">\n<figure class=\"full-width\"><img loading=\"lazy\" decoding=\"async\" src=\"https:\/\/habrastorage.org\/r\/w780q1\/getpro\/habr\/upload_files\/2e4\/474\/ff0\/2e4474ff0adcac1f693093a2e60f9b4a.jpg\" width=\"780\" height=\"440\" data-src=\"https:\/\/habrastorage.org\/getpro\/habr\/upload_files\/2e4\/474\/ff0\/2e4474ff0adcac1f693093a2e60f9b4a.jpg\" data-blurred=\"true\"\/><figcaption><\/figcaption><\/figure>\n<p>\u041f\u043e\u0434\u043e\u0431\u043d\u044b\u0439 Python Triton \u0443\u0436\u0435 \u0440\u0430\u0431\u043e\u0442\u0430\u0435\u0442 \u0432 \u044f\u0434\u0440\u0430\u0445, \u043a\u043e\u0442\u043e\u0440\u044b\u0435 \u0432 2 \u0440\u0430\u0437\u0430 \u044d\u0444\u0444\u0435\u043a\u0442\u0438\u0432\u043d\u0435\u0435 \u044d\u043a\u0432\u0438\u0432\u0430\u043b\u0435\u043d\u0442\u043d\u044b\u0445 \u0440\u0435\u0430\u043b\u0438\u0437\u0430\u0446\u0438\u0439 Torch. \u0410 \u044f\u0434\u0440\u0430 \u043c\u0430\u0442\u0440\u0438\u0447\u043d\u043e\u0433\u043e \u0443\u043c\u043d\u043e\u0436\u0435\u043d\u0438\u044f FP16, \u0441\u0440\u0430\u0432\u043d\u0438\u043c\u044b\u0435 \u043f\u0440\u043e\u0438\u0437\u0432\u043e\u0434\u0438\u0442\u0435\u043b\u044c\u043d\u043e\u0441\u0442\u044c\u044e \u0441 cuBLAS, \u043d\u0430 Triton \u0437\u0430\u0439\u043c\u0443\u0442 \u043c\u0435\u043d\u0435\u0435 25 \u0441\u0442\u0440\u043e\u043a. \u041a\u0430\u043a \u0443\u0442\u0432\u0435\u0440\u0436\u0434\u0430\u0435\u0442 \u0430\u0432\u0442\u043e\u0440, \u043c\u043d\u043e\u0433\u0438\u0435 \u043f\u0440\u043e\u0433\u0440\u0430\u043c\u043c\u0438\u0441\u0442\u044b \u043d\u0435 \u043c\u043e\u0433\u0443\u0442 \u043d\u0430\u043f\u0438\u0441\u0430\u0442\u044c \u0442\u0430\u043a\u0438\u0435 \u044f\u0434\u0440\u0430. \u041f\u043e\u0434\u0440\u043e\u0431\u043d\u043e\u0441\u0442\u044f\u043c\u0438 \u043e Triton \u0434\u0435\u043b\u0438\u043c\u0441\u044f \u043a \u0441\u0442\u0430\u0440\u0442\u0443 <a href=\"https:\/\/skillfactory.ru\/machine-learning-i-deep-learning?utm_source=habr&amp;utm_medium=habr&amp;utm_campaign=article&amp;utm_content=data-science_mldl_230921&amp;utm_term=lead\">\u043a\u0443\u0440\u0441\u0430 \u043f\u043e ML \u0438 DL<\/a>.<\/p>\n<hr\/>\n<p>\u041d\u043e\u0432\u044b\u0435 \u0438\u0441\u0441\u043b\u0435\u0434\u043e\u0432\u0430\u0442\u0435\u043b\u044c\u0441\u043a\u0438\u0435 \u0438\u0434\u0435\u0438 \u0432 \u043e\u0431\u043b\u0430\u0441\u0442\u0438 Deep Learning \u043e\u0431\u044b\u0447\u043d\u043e \u0440\u0435\u0430\u043b\u0438\u0437\u0443\u044e\u0442\u0441\u044f \u0441 \u043f\u043e\u043c\u043e\u0449\u044c\u044e \u043a\u043e\u043c\u0431\u0438\u043d\u0430\u0446\u0438\u0438 \u043e\u043f\u0435\u0440\u0430\u0442\u043e\u0440\u043e\u0432 \u043d\u0430\u0442\u0438\u0432\u043d\u043e\u0433\u043e \u0444\u0440\u0435\u0439\u043c\u0432\u043e\u0440\u043a\u0430. \u042d\u0442\u043e \u0443\u0434\u043e\u0431\u043d\u043e, \u043d\u043e \u0447\u0430\u0441\u0442\u043e \u0442\u0440\u0435\u0431\u0443\u0435\u0442 \u0441\u043e\u0437\u0434\u0430\u043d\u0438\u044f (\u0438\/\u0438\u043b\u0438 \u043f\u0435\u0440\u0435\u043c\u0435\u0449\u0435\u043d\u0438\u044f) \u043c\u043d\u043e\u0436\u0435\u0441\u0442\u0432\u0430 \u0432\u0440\u0435\u043c\u0435\u043d\u043d\u044b\u0445 \u0442\u0435\u043d\u0437\u043e\u0440\u043e\u0432, \u0447\u0442\u043e \u043c\u043e\u0436\u0435\u0442 \u0441\u043d\u0438\u0437\u0438\u0442\u044c \u043f\u0440\u043e\u0438\u0437\u0432\u043e\u0434\u0438\u0442\u0435\u043b\u044c\u043d\u043e\u0441\u0442\u044c \u043d\u0435\u0439\u0440\u043e\u0441\u0435\u0442\u0435\u0439 \u0432 \u0431\u043e\u043b\u044c\u0448\u043e\u043c \u043c\u0430\u0441\u0448\u0442\u0430\u0431\u0435. \u041f\u0440\u043e\u0431\u043b\u0435\u043c\u044b \u0440\u0435\u0448\u0430\u044e\u0442\u0441\u044f \u0441\u043f\u0435\u0446\u0438\u0430\u043b\u0438\u0437\u0438\u0440\u043e\u0432\u0430\u043d\u043d\u044b\u043c\u0438 \u044f\u0434\u0440\u0430\u043c\u0438 \u0434\u043b\u044f GPU, \u043d\u043e \u043d\u0430\u043f\u0438\u0441\u0430\u0442\u044c \u0438\u0445 \u043c\u043e\u0436\u0435\u0442 \u0431\u044b\u0442\u044c \u0443\u0434\u0438\u0432\u0438\u0442\u0435\u043b\u044c\u043d\u043e \u0441\u043b\u043e\u0436\u043d\u043e \u0438\u0437-\u0437\u0430 \u0442\u043e\u043d\u043a\u043e\u0441\u0442\u0435\u0439 \u043f\u0440\u043e\u0433\u0440\u0430\u043c\u043c\u0438\u0440\u043e\u0432\u0430\u043d\u0438\u044f \u043d\u0430 GPU <a href=\"https:\/\/openai.com\/blog\/triton\/#rf1\"><sup>1<\/sup><\/a><sup>, <\/sup><a href=\"https:\/\/openai.com\/blog\/triton\/#rf2\"><sup>2<\/sup><\/a><sup>, <\/sup><a href=\"https:\/\/openai.com\/blog\/triton\/#rf3\"><sup>3<\/sup><\/a>.<\/p>\n<p>\u0418, \u0445\u043e\u0442\u044f \u0432 \u043f\u043e\u0441\u043b\u0435\u0434\u043d\u0435\u0435 \u0432\u0440\u0435\u043c\u044f \u043f\u043e\u044f\u0432\u0438\u043b\u043e\u0441\u044c \u043c\u043d\u043e\u0436\u0435\u0441\u0442\u0432\u043e \u043e\u0431\u043b\u0435\u0433\u0447\u0430\u044e\u0449\u0438\u0445 \u0442\u0440\u0443\u0434 \u0441\u0438\u0441\u0442\u0435\u043c<a href=\"https:\/\/openai.com\/blog\/triton\/#rf4\"><sup>4<\/sup><\/a><sup>, <\/sup><a href=\"https:\/\/openai.com\/blog\/triton\/#rf5\"><sup>5<\/sup><\/a>, \u043c\u044b \u043e\u0431\u043d\u0430\u0440\u0443\u0436\u0438\u043b\u0438, \u0447\u0442\u043e \u043e\u043d\u0438 \u043b\u0438\u0431\u043e \u0441\u043b\u0438\u0448\u043a\u043e\u043c \u043c\u043d\u043e\u0433\u043e\u0441\u043b\u043e\u0432\u043d\u044b, \u043b\u0438\u0431\u043e \u043d\u0435\u0434\u043e\u0441\u0442\u0430\u0442\u043e\u0447\u043d\u043e \u0433\u0438\u0431\u043a\u0438\u0435, \u043b\u0438\u0431\u043e \u0433\u0435\u043d\u0435\u0440\u0438\u0440\u0443\u044e\u0442 \u043a\u043e\u0434 \u0437\u0430\u043c\u0435\u0442\u043d\u043e \u043c\u0435\u0434\u043b\u0435\u043d\u043d\u0435\u0435, \u0447\u0435\u043c \u043d\u0430\u0441\u0442\u0440\u043e\u0435\u043d\u043d\u044b\u0439 \u0432\u0440\u0443\u0447\u043d\u0443\u044e \u0431\u0430\u0437\u043e\u0432\u044b\u0439. \u042d\u0442\u043e \u043f\u0440\u0438\u0432\u0435\u043b\u043e \u043a \u0440\u0430\u0441\u0448\u0438\u0440\u0435\u043d\u0438\u044e \u0438 \u0443\u043b\u0443\u0447\u0448\u0435\u043d\u0438\u044e Triton<a href=\"https:\/\/openai.com\/blog\/triton\/#rf6\"><sup>6<\/sup><\/a>, \u2014 \u043d\u043e\u0432\u043e\u0433\u043e \u044f\u0437\u044b\u043a\u0430 \u0438 \u043a\u043e\u043c\u043f\u0438\u043b\u044f\u0442\u043e\u0440\u0430, \u0441\u043e\u0437\u0434\u0430\u0442\u0435\u043b\u044c \u043a\u043e\u0442\u043e\u0440\u043e\u0433\u043e \u0441\u0435\u0439\u0447\u0430\u0441 \u0440\u0430\u0431\u043e\u0442\u0430\u0435\u0442 \u0432 OpenAI.<\/p>\n<h2>\u041f\u0440\u043e\u0431\u043b\u0435\u043c\u044b \u043f\u0440\u043e\u0433\u0440\u0430\u043c\u043c\u0438\u0440\u043e\u0432\u0430\u043d\u0438\u044f GPU<\/h2>\n<p>\u0410\u0440\u0445\u0438\u0442\u0435\u043a\u0442\u0443\u0440\u0430 \u0441\u043e\u0432\u0440\u0435\u043c\u0435\u043d\u043d\u044b\u0445 GPU \u0443\u0441\u043b\u043e\u0432\u043d\u043e \u0434\u0435\u043b\u0438\u0442\u0441\u044f \u043d\u0430 \u0442\u0440\u0438 \u043e\u0441\u043d\u043e\u0432\u043d\u044b\u0445 \u043a\u043e\u043c\u043f\u043e\u043d\u0435\u043d\u0442\u0430 \u2014 DRAM, SRAM \u0438 ALU. \u041e\u043f\u0442\u0438\u043c\u0438\u0437\u0438\u0440\u0443\u044f \u043a\u043e\u0434 \u043d\u0430 CUDA, \u043d\u0443\u0436\u043d\u043e \u0443\u0447\u0438\u0442\u044b\u0432\u0430\u0442\u044c \u043a\u0430\u0436\u0434\u044b\u0439.<\/p>\n<ul>\n<li>\n<p>\u041f\u0435\u0440\u0435\u0434\u0430\u0447\u0438 \u043f\u0430\u043c\u044f\u0442\u0438 \u0438\u0437 DRAM \u0434\u043e\u043b\u0436\u043d\u044b \u0431\u044b\u0442\u044c \u043e\u0431\u044a\u0435\u0434\u0438\u043d\u0435\u043d\u044b \u0432 \u0431\u043e\u043b\u044c\u0448\u0438\u0435 \u0442\u0440\u0430\u043d\u0437\u0430\u043a\u0446\u0438\u0438, \u0447\u0442\u043e\u0431\u044b \u0438\u0441\u043f\u043e\u043b\u044c\u0437\u043e\u0432\u0430\u0442\u044c \u0431\u043e\u043b\u044c\u0448\u0443\u044e \u0448\u0438\u0440\u0438\u043d\u0443 \u0448\u0438\u043d\u044b \u0441\u043e\u0432\u0440\u0435\u043c\u0435\u043d\u043d\u044b\u0445 \u0438\u043d\u0442\u0435\u0440\u0444\u0435\u0439\u0441\u043e\u0432 \u043f\u0430\u043c\u044f\u0442\u0438.<\/p>\n<\/li>\n<li>\n<p>\u0414\u0430\u043d\u043d\u044b\u0435 \u0434\u043e\u043b\u0436\u043d\u044b \u0431\u044b\u0442\u044c \u0432\u0440\u0443\u0447\u043d\u0443\u044e \u0441\u043e\u0445\u0440\u0430\u043d\u0435\u043d\u044b \u0432 SRAM \u043f\u0435\u0440\u0435\u0434 \u043f\u043e\u0432\u0442\u043e\u0440\u043d\u044b\u043c \u0438\u0441\u043f\u043e\u043b\u044c\u0437\u043e\u0432\u0430\u043d\u0438\u0435\u043c \u0438 \u0443\u043f\u0440\u0430\u0432\u043b\u044f\u0442\u044c\u0441\u044f \u0442\u0430\u043a\u0438\u043c \u043e\u0431\u0440\u0430\u0437\u043e\u043c, \u0447\u0442\u043e\u0431\u044b \u043c\u0438\u043d\u0438\u043c\u0438\u0437\u0438\u0440\u043e\u0432\u0430\u0442\u044c \u043a\u043e\u043d\u0444\u043b\u0438\u043a\u0442\u044b \u043c\u0435\u0436\u0434\u0443 \u0431\u0430\u043d\u043a\u0430\u043c\u0438 \u043e\u0431\u0449\u0435\u0439 \u043f\u0430\u043c\u044f\u0442\u0438 \u043f\u0440\u0438 \u0438\u0437\u0432\u043b\u0435\u0447\u0435\u043d\u0438\u0438.<\/p>\n<\/li>\n<li>\n<p>\u0412\u044b\u0447\u0438\u0441\u043b\u0435\u043d\u0438\u044f \u0434\u043e\u043b\u0436\u043d\u044b \u0431\u044b\u0442\u044c \u0440\u0430\u0437\u0434\u0435\u043b\u0435\u043d\u044b \u0438 \u0442\u0449\u0430\u0442\u0435\u043b\u044c\u043d\u043e \u0441\u043f\u043b\u0430\u043d\u0438\u0440\u043e\u0432\u0430\u043d\u044b \u043a\u0430\u043a \u043c\u0435\u0436\u0434\u0443 \u043f\u043e\u0442\u043e\u043a\u043e\u0432\u044b\u043c\u0438 \u043c\u0443\u043b\u044c\u0442\u0438\u043f\u0440\u043e\u0446\u0435\u0441\u0441\u043e\u0440\u0430\u043c\u0438 (SM), \u0442\u0430\u043a \u0438 \u0432\u043d\u0443\u0442\u0440\u0438 \u043d\u0438\u0445, \u0447\u0442\u043e\u0431\u044b \u0441\u043f\u043e\u0441\u043e\u0431\u0441\u0442\u0432\u043e\u0432\u0430\u0442\u044c \u043f\u0430\u0440\u0430\u043b\u043b\u0435\u043b\u0438\u0437\u043c\u0443 \u043d\u0430 \u0443\u0440\u043e\u0432\u043d\u0435 \u0438\u043d\u0441\u0442\u0440\u0443\u043a\u0446\u0438\u0439\/\u043f\u043e\u0442\u043e\u043a\u043e\u0432 \u0438 \u0438\u0441\u043f\u043e\u043b\u044c\u0437\u043e\u0432\u0430\u0442\u044c \u0441\u043f\u0435\u0446\u0438\u0430\u043b\u0438\u0437\u0438\u0440\u043e\u0432\u0430\u043d\u043d\u044b\u0435 ALU (\u043d\u0430\u043f\u0440\u0438\u043c\u0435\u0440, \u0442\u0435\u043d\u0437\u043e\u0440\u043d\u044b\u0435 \u044f\u0434\u0440\u0430).<\/p>\n<\/li>\n<\/ul>\n<figure class=\"full-width\"><img loading=\"lazy\" decoding=\"async\" src=\"https:\/\/habrastorage.org\/r\/w1560\/getpro\/habr\/upload_files\/85c\/a5f\/877\/85ca5f8773675eac0b4ab4ab8d6c6657.png\" alt=\"\u0411\u0430\u0437\u043e\u0432\u0430\u044f \u0430\u0440\u0445\u0438\u0442\u0435\u043a\u0442\u0443\u0440\u0430 GPU.\" title=\"\u0411\u0430\u0437\u043e\u0432\u0430\u044f \u0430\u0440\u0445\u0438\u0442\u0435\u043a\u0442\u0443\u0440\u0430 GPU.\" width=\"796\" height=\"254\" data-src=\"https:\/\/habrastorage.org\/getpro\/habr\/upload_files\/85c\/a5f\/877\/85ca5f8773675eac0b4ab4ab8d6c6657.png\"\/><figcaption>\u0411\u0430\u0437\u043e\u0432\u0430\u044f \u0430\u0440\u0445\u0438\u0442\u0435\u043a\u0442\u0443\u0440\u0430 GPU.<\/figcaption><\/figure>\n<p>\u0423\u0447\u0438\u0442\u044b\u0432\u0430\u0442\u044c \u044d\u0442\u0438 \u0444\u0430\u043a\u0442\u043e\u0440\u044b \u043c\u043e\u0436\u0435\u0442 \u0431\u044b\u0442\u044c \u0441\u043b\u043e\u0436\u043d\u043e \u0434\u0430\u0436\u0435 \u043e\u043f\u044b\u0442\u043d\u044b\u043c \u043f\u0440\u043e\u0433\u0440\u0430\u043c\u043c\u0438\u0441\u0442\u0430\u043c CUDA. \u0426\u0435\u043b\u044c Triton \u2014 \u043f\u043e\u043b\u043d\u043e\u0441\u0442\u044c\u044e \u0430\u0432\u0442\u043e\u043c\u0430\u0442\u0438\u0437\u0438\u0440\u043e\u0432\u0430\u0442\u044c \u044d\u0442\u0438 \u043e\u043f\u0442\u0438\u043c\u0438\u0437\u0430\u0446\u0438\u0438, \u0447\u0442\u043e\u0431\u044b \u0440\u0430\u0437\u0440\u0430\u0431\u043e\u0442\u0447\u0438\u043a\u0438 \u0441\u043e\u0441\u0440\u0435\u0434\u043e\u0442\u043e\u0447\u0438\u043b\u0438\u0441\u044c \u043d\u0430 \u0432\u044b\u0441\u043e\u043a\u043e\u0443\u0440\u043e\u0432\u043d\u0435\u0432\u043e\u0439 \u043b\u043e\u0433\u0438\u043a\u0435 \u043f\u0430\u0440\u0430\u043b\u043b\u0435\u043b\u044c\u043d\u043e\u0433\u043e \u043a\u043e\u0434\u0430.\u00a0<\/p>\n<p>Triton \u0441\u0442\u0440\u0435\u043c\u0438\u0442\u0441\u044f \u043a \u0448\u0438\u0440\u043e\u043a\u043e\u043c\u0443 \u043f\u0440\u0438\u043c\u0435\u043d\u0435\u043d\u0438\u044e, \u0438 \u043f\u043e\u044d\u0442\u043e\u043c\u0443 \u043d\u0435 \u0441\u043e\u0441\u0442\u0430\u0432\u043b\u044f\u0435\u0442 \u0430\u0432\u0442\u043e\u043c\u0430\u0442\u0438\u0447\u0435\u0441\u043a\u043e\u0433\u043e \u0440\u0430\u0441\u043f\u0438\u0441\u0430\u043d\u0438\u044f \u0440\u0430\u0431\u043e\u0442\u044b \u043c\u0435\u0436\u0434\u0443 SM, \u043e\u0441\u0442\u0430\u0432\u043b\u044f\u044f \u043d\u0435\u043a\u043e\u0442\u043e\u0440\u044b\u0435 \u0432\u0430\u0436\u043d\u044b\u0435 \u0430\u043b\u0433\u043e\u0440\u0438\u0442\u043c\u0438\u0447\u0435\u0441\u043a\u0438\u0435 \u0441\u043e\u043e\u0431\u0440\u0430\u0436\u0435\u043d\u0438\u044f (\u043d\u0430\u043f\u0440\u0438\u043c\u0435\u0440, \u0442\u0430\u0439\u043b\u0438\u043d\u0433, \u0441\u0438\u043d\u0445\u0440\u043e\u043d\u0438\u0437\u0430\u0446\u0438\u044f \u043c\u0435\u0436\u0434\u0443 SM) \u043d\u0430 \u0443\u0441\u043c\u043e\u0442\u0440\u0435\u043d\u0438\u0435 \u0440\u0430\u0437\u0440\u0430\u0431\u043e\u0442\u0447\u0438\u043a\u043e\u0432.<\/p>\n<p>\u041e\u043f\u0442\u0438\u043c\u0438\u0437\u0430\u0446\u0438\u044f \u043a\u043e\u043c\u043f\u0438\u043b\u044f\u0442\u043e\u0440\u0430 \u0432 CUDA \u043f\u043e \u0441\u0440\u0430\u0432\u043d\u0435\u043d\u0438\u044e \u0441 Triton.<\/p>\n<div>\n<div class=\"table\">\n<table>\n<tbody>\n<tr>\n<th>\n<p><strong><br \/><\/strong><\/p>\n<\/th>\n<th>\n<p><strong>CUDA<\/strong><\/p>\n<\/th>\n<td>\n<p><strong>TRITON<\/strong><\/p>\n<\/td>\n<\/tr>\n<tr>\n<td>\n<p>\u041e\u0431\u044a\u0435\u0434\u0438\u043d\u0435\u043d\u0438\u0435 \u043f\u0430\u043c\u044f\u0442\u0438<\/p>\n<\/td>\n<td>\n<p>\u0412\u0440\u0443\u0447\u043d\u0443\u044e<\/p>\n<\/td>\n<td>\n<p>\u0410\u0432\u0442\u043e\u043c\u0430\u0442\u0438\u0447\u0435\u0441\u043a\u0438<\/p>\n<\/td>\n<\/tr>\n<tr>\n<td>\n<p>\u0423\u043f\u0440\u0430\u0432\u043b\u0435\u043d\u0438\u0435 \u043e\u0431\u0449\u0435\u0439 \u043f\u0430\u043c\u044f\u0442\u044c\u044e<\/p>\n<\/td>\n<td>\n<p>\u0412\u0440\u0443\u0447\u043d\u0443\u044e<\/p>\n<\/td>\n<td>\n<p>\u0410\u0432\u0442\u043e\u043c\u0430\u0442\u0438\u0447\u0435\u0441\u043a\u0438<\/p>\n<\/td>\n<\/tr>\n<tr>\n<td>\n<p>\u041f\u043b\u0430\u043d\u0438\u0440\u043e\u0432\u0430\u043d\u0438\u0435 (\u0432\u043d\u0443\u0442\u0440\u0438 SM)<\/p>\n<\/td>\n<td>\n<p>\u0412\u0440\u0443\u0447\u043d\u0443\u044e<\/p>\n<\/td>\n<td>\n<p>\u0410\u0432\u0442\u043e\u043c\u0430\u0442\u0438\u0447\u0435\u0441\u043a\u0438<\/p>\n<\/td>\n<\/tr>\n<tr>\n<td>\n<p>\u041f\u043b\u0430\u043d\u0438\u0440\u043e\u0432\u0430\u043d\u0438\u0435 (\u043f\u043e \u0432\u0441\u0435\u043c SM)<\/p>\n<\/td>\n<td>\n<p>\u0412\u0440\u0443\u0447\u043d\u0443\u044e<\/p>\n<\/td>\n<td>\n<p>\u0412\u0440\u0443\u0447\u043d\u0443\u044e<\/p>\n<\/td>\n<\/tr>\n<\/tbody>\n<\/table>\n<\/div>\n<\/div>\n<h3>\u041c\u043e\u0434\u0435\u043b\u044c \u043f\u0440\u043e\u0433\u0440\u0430\u043c\u043c\u0438\u0440\u043e\u0432\u0430\u043d\u0438\u044f<\/h3>\n<p>\u0418\u0437 \u0432\u0441\u0435\u0445 \u0434\u043e\u0441\u0442\u0443\u043f\u043d\u044b\u0445 \u0434\u043e\u043c\u0435\u043d\u043d\u043e-\u0441\u043f\u0435\u0446\u0438\u0444\u0438\u0447\u0435\u0441\u043a\u0438\u0445 \u044f\u0437\u044b\u043a\u043e\u0432 \u0438 JIT-\u043a\u043e\u043c\u043f\u0438\u043b\u044f\u0442\u043e\u0440\u043e\u0432 Triton, \u043f\u043e\u0436\u0430\u043b\u0443\u0439, \u043d\u0430\u0438\u0431\u043e\u043b\u0435\u0435 \u043f\u043e\u0445\u043e\u0436 \u043d\u0430 Numba: \u044f\u0434\u0440\u0430 \u043e\u043f\u0440\u0435\u0434\u0435\u043b\u044f\u044e\u0442\u0441\u044f \u043a\u0430\u043a \u043e\u0444\u043e\u0440\u043c\u043b\u0435\u043d\u043d\u044b\u0435 \u0444\u0443\u043d\u043a\u0446\u0438\u0438 Python \u0438 \u0437\u0430\u043f\u0443\u0441\u043a\u0430\u044e\u0442\u0441\u044f \u043f\u0430\u0440\u0430\u043b\u043b\u0435\u043b\u044c\u043d\u043e \u0441 \u0440\u0430\u0437\u043d\u044b\u043c\u0438 program_id \u043d\u0430 \u0441\u0435\u0442\u043a\u0435 \u0442\u0430\u043a \u043d\u0430\u0437\u044b\u0432\u0430\u0435\u043c\u044b\u0445 <em>instances<\/em> (\u0438\u043d\u0441\u0442\u0430\u043d\u0441\u043e\u0432, \u044d\u043a\u0437\u0435\u043c\u043f\u043b\u044f\u0440\u043e\u0432).\u00a0<\/p>\n<p>\u041e\u0434\u043d\u0430\u043a\u043e, \u043a\u0430\u043a \u043f\u043e\u043a\u0430\u0437\u0430\u043d\u043e \u0432 \u043f\u0440\u0438\u0432\u0435\u0434\u0451\u043d\u043d\u043e\u043c \u043d\u0438\u0436\u0435 \u0444\u0440\u0430\u0433\u043c\u0435\u043d\u0442\u0435 \u043a\u043e\u0434\u0430, \u043d\u0430 \u044d\u0442\u043e\u043c \u0441\u0445\u043e\u0434\u0441\u0442\u0432\u043e \u0437\u0430\u043a\u0430\u043d\u0447\u0438\u0432\u0430\u0435\u0442\u0441\u044f: Triton \u0440\u0430\u0441\u043a\u0440\u044b\u0432\u0430\u0435\u0442 \u0432\u043d\u0443\u0442\u0440\u0438\u044d\u043a\u0437\u0435\u043c\u043f\u043b\u044f\u0440\u043d\u044b\u0439 \u043f\u0430\u0440\u0430\u043b\u043b\u0435\u043b\u0438\u0437\u043c \u0447\u0435\u0440\u0435\u0437 \u043e\u043f\u0435\u0440\u0430\u0446\u0438\u0438 \u043d\u0430\u0434 <em>\u0431\u043b\u043e\u043a\u0430\u043c\u0438<\/em> \u2014 \u043d\u0435\u0431\u043e\u043b\u044c\u0448\u0438\u043c\u0438 \u043c\u0430\u0441\u0441\u0438\u0432\u0430\u043c\u0438, \u0440\u0430\u0437\u043c\u0435\u0440\u044b \u043a\u043e\u0442\u043e\u0440\u044b\u0445 \u0440\u0430\u0432\u043d\u044b \u0441\u0442\u0435\u043f\u0435\u043d\u0438 \u0434\u0432\u043e\u0439\u043a\u0438, \u2014 \u0432\u043c\u0435\u0441\u0442\u043e \u043c\u043e\u0434\u0435\u043b\u0438 \u0432\u044b\u043f\u043e\u043b\u043d\u0435\u043d\u0438\u044f Single Instruction, Multiple Thread (SIMT, \u043e\u0434\u043d\u0430 \u0438\u043d\u0441\u0442\u0440\u0443\u043a\u0446\u0438\u044f \u2014 \u043c\u043d\u043e\u0433\u043e \u043f\u043e\u0442\u043e\u043a\u043e\u0432)<a href=\"https:\/\/openai.com\/blog\/triton\/#rf7\"><sup>7<\/sup><\/a>.\u00a0<\/p>\n<p>\u041f\u0440\u0438 \u044d\u0442\u043e\u043c Triton \u044d\u0444\u0444\u0435\u043a\u0442\u0438\u0432\u043d\u043e \u0430\u0431\u0441\u0442\u0440\u0430\u0433\u0438\u0440\u0443\u0435\u0442\u0441\u044f \u043e\u0442 \u0432\u0441\u0435\u0445 \u043f\u0440\u043e\u0431\u043b\u0435\u043c, \u0441\u0432\u044f\u0437\u0430\u043d\u043d\u044b\u0445 \u0441 \u043f\u0430\u0440\u0430\u043b\u043b\u0435\u043b\u0438\u0437\u043c\u043e\u043c <em>\u0432\u043d\u0443\u0442\u0440\u0438<\/em> \u0431\u043b\u043e\u043a\u043e\u0432 \u043f\u043e\u0442\u043e\u043a\u043e\u0432 CUDA (\u043d\u0430\u043f\u0440\u0438\u043c\u0435\u0440, \u043e\u0431\u044a\u0435\u0434\u0438\u043d\u0435\u043d\u0438\u0435 \u043f\u0430\u043c\u044f\u0442\u0438, \u0441\u0438\u043d\u0445\u0440\u043e\u043d\u0438\u0437\u0430\u0446\u0438\u044f\/\u043a\u043e\u043d\u0444\u043b\u0438\u043a\u0442\u044b \u043e\u0431\u0449\u0435\u0439 \u043f\u0430\u043c\u044f\u0442\u0438, \u043f\u043b\u0430\u043d\u0438\u0440\u043e\u0432\u0430\u043d\u0438\u0435 \u0440\u0430\u0431\u043e\u0442\u044b \u0442\u0435\u043d\u0437\u043e\u0440\u043d\u044b\u0445 \u044f\u0434\u0435\u0440).<\/p>\n<p>\u0414\u043e\u0431\u0430\u0432\u043b\u0435\u043d\u0438\u0435 \u0432\u0435\u043a\u0442\u043e\u0440\u0430:<\/p>\n<pre><code class=\"python\">BLOCK = 512  # This is a GPU kernel in Numba. # Different instances of this # function may run in parallel. @jit def add(X, Y, Z, N):    # In Numba\/CUDA, each kernel     # instance itself uses an SIMT execution    # model, where instructions are executed in    # parallel for different values of threadIdx    tid = threadIdx.x    bid = blockIdx.x    # scalar index    idx = bid * BLOCK + tid    if id &lt; N:      # There is no pointer in Numba.      # Z,X,Y are dense tensors      Z[idx] = X[idx] + Y[idx]   ... grid = (ceil_div(N, BLOCK),) block = (BLOCK,) add[grid, block](x, y, z, x.shape[0])<\/code><\/pre>\n<pre><code class=\"python\">BLOCK = 512  # This is a GPU kernel in Triton. # Different instances of this # function may run in parallel. @jit def add(X, Y, Z, N):    # In Triton, each kernel instance    # executes block operations on a    # single thread: there is no construct    # analogous to threadIdx    pid = program_id(0)    # block of indices    idx = pid * BLOCK + arange(BLOCK)    mask = idx &lt; N    # Triton uses pointer arithmetics      # rather than indexing operators    x = load(X + idx, mask=mask)    y = load(Y + idx, mask=mask)    store(Z + idx, x + y, mask=mask)   ... grid = (ceil_div(N, BLOCK),) # no thread-block add[grid](x, y, z, x.shape[0])<\/code><\/pre>\n<p>\u0425\u043e\u0442\u044f \u044d\u0442\u043e \u043c\u043e\u0436\u0435\u0442 \u0431\u044b\u0442\u044c \u043d\u0435 \u043e\u0441\u043e\u0431\u0435\u043d\u043d\u043e \u043f\u043e\u043b\u0435\u0437\u043d\u043e \u0434\u043b\u044f \u043f\u043e\u044d\u043b\u0435\u043c\u0435\u043d\u0442\u043d\u044b\u0445 \u0432\u044b\u0447\u0438\u0441\u043b\u0435\u043d\u0438\u0439, \u043f\u043e\u0434\u0445\u043e\u0434 \u043c\u043e\u0436\u0435\u0442 \u0437\u043d\u0430\u0447\u0438\u0442\u0435\u043b\u044c\u043d\u043e \u0443\u043f\u0440\u043e\u0441\u0442\u0438\u0442\u044c \u0440\u0430\u0437\u0440\u0430\u0431\u043e\u0442\u043a\u0443 \u0431\u043e\u043b\u0435\u0435 \u0441\u043b\u043e\u0436\u043d\u044b\u0445 \u043f\u0440\u043e\u0433\u0440\u0430\u043c\u043c \u043d\u0430 GPU. \u0420\u0430\u0441\u0441\u043c\u043e\u0442\u0440\u0438\u043c, \u043d\u0430\u043f\u0440\u0438\u043c\u0435\u0440, \u0441\u043b\u0443\u0447\u0430\u0439 \u044f\u0434\u0440\u0430 fused softmax, \u0433\u0434\u0435 \u043a\u0430\u0436\u0434\u044b\u0439 \u044d\u043a\u0437\u0435\u043c\u043f\u043b\u044f\u0440 \u043d\u043e\u0440\u043c\u0438\u0440\u0443\u0435\u0442 \u043e\u0442\u0434\u0435\u043b\u044c\u043d\u0443\u044e \u0441\u0442\u0440\u043e\u043a\u0443 \u0437\u0430\u0434\u0430\u043d\u043d\u043e\u0433\u043e \u0432\u0445\u043e\u0434\u043d\u043e\u0433\u043e \u0442\u0435\u043d\u0437\u043e\u0440\u0430 <em>X<\/em>\u2208R<em><sup>M<\/sup><\/em><sup>\u00d7<\/sup><em><sup>N<\/sup><\/em>.\u00a0<\/p>\n<p>\u0421\u0442\u0430\u043d\u0434\u0430\u0440\u0442\u043d\u044b\u0435 CUDA-\u0440\u0435\u0430\u043b\u0438\u0437\u0430\u0446\u0438\u0438 \u044d\u0442\u043e\u0439 \u0441\u0442\u0440\u0430\u0442\u0435\u0433\u0438\u0438 \u0440\u0430\u0441\u043f\u0430\u0440\u0430\u043b\u043b\u0435\u043b\u0438\u0432\u0430\u043d\u0438\u044f \u043c\u043e\u0433\u0443\u0442 \u0431\u044b\u0442\u044c \u0441\u043b\u043e\u0436\u043d\u044b\u043c\u0438 \u0432 \u043d\u0430\u043f\u0438\u0441\u0430\u043d\u0438\u0438, \u0442\u0440\u0435\u0431\u0443\u044f \u044f\u0432\u043d\u043e\u0439 \u0441\u0438\u043d\u0445\u0440\u043e\u043d\u0438\u0437\u0430\u0446\u0438\u0438 \u043c\u0435\u0436\u0434\u0443 \u043f\u043e\u0442\u043e\u043a\u0430\u043c\u0438, \u043f\u043e\u0441\u043a\u043e\u043b\u044c\u043a\u0443 \u043e\u043d\u0438 \u043e\u0434\u043d\u043e\u0432\u0440\u0435\u043c\u0435\u043d\u043d\u043e \u0443\u043c\u0435\u043d\u044c\u0448\u0430\u044e\u0442 \u043e\u0434\u043d\u0443 \u0438 \u0442\u0443 \u0436\u0435 \u0441\u0442\u0440\u043e\u043a\u0443 X<em>X<\/em>. \u0411\u043e\u043b\u044c\u0448\u0430\u044f \u0447\u0430\u0441\u0442\u044c \u044d\u0442\u043e\u0439 \u0441\u043b\u043e\u0436\u043d\u043e\u0441\u0442\u0438 \u0438\u0441\u0447\u0435\u0437\u0430\u0435\u0442 \u0432 Triton, \u0433\u0434\u0435 \u043a\u0430\u0436\u0434\u044b\u0439 \u044d\u043a\u0437\u0435\u043c\u043f\u043b\u044f\u0440 \u044f\u0434\u0440\u0430 \u0437\u0430\u0433\u0440\u0443\u0436\u0430\u0435\u0442 \u0438\u043d\u0442\u0435\u0440\u0435\u0441\u0443\u044e\u0449\u0443\u044e \u0441\u0442\u0440\u043e\u043a\u0443 \u0438 \u043f\u043e\u0441\u043b\u0435\u0434\u043e\u0432\u0430\u0442\u0435\u043b\u044c\u043d\u043e \u043d\u043e\u0440\u043c\u0430\u043b\u0438\u0437\u0443\u0435\u0442 \u0435\u0451, \u0438\u0441\u043f\u043e\u043b\u044c\u0437\u0443\u044f NumPy-\u043f\u043e\u0434\u043e\u0431\u043d\u044b\u0435 \u043f\u0440\u0438\u043c\u0438\u0442\u0438\u0432\u044b.<\/p>\n<p>Fused softmax \u0432 Triton:<\/p>\n<pre><code class=\"python\">import triton import triton.language as tl  @triton.jit def softmax(Y, stride_ym, stride_yn, X, stride_xm, stride_xn, M, N):     # row index     m = tl.program_id(0)     # col indices     # this specific kernel only works for matrices that      # have less than BLOCK_SIZE columns     BLOCK_SIZE = 1024     n = tl.arange(0, BLOCK_SIZE)     # the memory address of all the elements     # that we want to load can be computed as follows     X = X + m * stride_xm + n * stride_xn     # load input data; pad out-of-bounds elements with 0      x = tl.load(X, mask=n &lt; N, other=-float('inf'))     # compute numerically-stable softmax     z = x - tl.max(x, axis=0)     num = tl.exp(z)     denom = tl.sum(num, axis=0)     y = num \/ denom     # write back to Y     Y = Y + m * stride_ym + n * stride_yn     tl.store(Y, y, mask=n &lt; N)  import torch # Allocate input\/output tensors X = torch.normal(0, 1, size=(583, 931), device='cuda') Y = torch.empty_like(X) # SPMD launch grid grid = (X.shape[0], ) # enqueue GPU kernel softmax[grid](Y, Y.stride(0), Y.stride(1),                X, X.stride(0), X.stride(1),               X.shape[0]    , X.shape[1])<\/code><\/pre>\n<p>\u041e\u0431\u0440\u0430\u0442\u0438\u0442\u0435 \u0432\u043d\u0438\u043c\u0430\u043d\u0438\u0435, \u0447\u0442\u043e Triton JIT \u0440\u0430\u0441\u0441\u043c\u0430\u0442\u0440\u0438\u0432\u0430\u0435\u0442 X \u0438 Y \u043a\u0430\u043a <em>\u0443\u043a\u0430\u0437\u0430\u0442\u0435\u043b\u0438<\/em>, \u0430 \u043d\u0435 \u0442\u0435\u043d\u0437\u043e\u0440\u044b; \u043c\u044b \u0447\u0443\u0432\u0441\u0442\u0432\u043e\u0432\u0430\u043b\u0438, \u0447\u0442\u043e \u0441\u043e\u0445\u0440\u0430\u043d\u0435\u043d\u0438\u0435 \u043d\u0438\u0437\u043a\u043e\u0443\u0440\u043e\u0432\u043d\u0435\u0432\u043e\u0433\u043e \u043a\u043e\u043d\u0442\u0440\u043e\u043b\u044f \u0434\u043e\u0441\u0442\u0443\u043f\u0430 \u043a \u043f\u0430\u043c\u044f\u0442\u0438 \u0432\u0430\u0436\u043d\u043e \u0432 \u043e\u0431\u0440\u0430\u0449\u0435\u043d\u0438\u0438 \u0441 \u0431\u043e\u043b\u0435\u0435 \u0441\u043b\u043e\u0436\u043d\u044b\u043c\u0438 \u0441\u0442\u0440\u0443\u043a\u0442\u0443\u0440\u0430\u043c\u0438 \u0434\u0430\u043d\u043d\u044b\u0445 (\u043d\u0430\u043f\u0440\u0438\u043c\u0435\u0440, \u0431\u043b\u043e\u0447\u043d\u043e-\u0440\u0430\u0441\u0441\u0435\u044f\u043d\u043d\u044b\u0445 (block-sparse) \u0442\u0435\u043d\u0437\u043e\u0440\u043e\u0432). \u0412\u0430\u0436\u043d\u043e \u043e\u0442\u043c\u0435\u0442\u0438\u0442\u044c, \u0447\u0442\u043e \u044d\u0442\u0430 \u043a\u043e\u043d\u043a\u0440\u0435\u0442\u043d\u0430\u044f \u0440\u0435\u0430\u043b\u0438\u0437\u0430\u0446\u0438\u044f softmax \u0441\u043e\u0445\u0440\u0430\u043d\u044f\u0435\u0442 \u0441\u0442\u0440\u043e\u043a\u0438 X<em>X<\/em> \u0432 SRAM \u043d\u0430 \u043f\u0440\u043e\u0442\u044f\u0436\u0435\u043d\u0438\u0438 \u0432\u0441\u0435\u0433\u043e \u043f\u0440\u043e\u0446\u0435\u0441\u0441\u0430 \u043d\u043e\u0440\u043c\u0430\u043b\u0438\u0437\u0430\u0446\u0438\u0438, \u0447\u0442\u043e \u043f\u043e\u0437\u0432\u043e\u043b\u044f\u0435\u0442 \u043c\u0430\u043a\u0441\u0438\u043c\u0430\u043b\u044c\u043d\u043e \u043f\u0435\u0440\u0435\u0438\u0441\u043f\u043e\u043b\u044c\u0437\u043e\u0432\u0430\u0442\u044c \u0434\u0430\u043d\u043d\u044b\u0435, \u043a\u043e\u0433\u0434\u0430 \u044d\u0442\u043e \u0432\u043e\u0437\u043c\u043e\u0436\u043d\u043e (~&lt;32K \u0441\u0442\u043e\u043b\u0431\u0446\u043e\u0432).\u00a0<\/p>\n<p>\u0418\u043d\u0430\u0447\u0435 \u0434\u0435\u043b\u043e \u043e\u0431\u0441\u0442\u043e\u0438\u0442 \u0432\u043e \u0432\u043d\u0443\u0442\u0440\u0435\u043d\u043d\u0435\u043c \u043a\u043e\u0434\u0435 CUDA PyTorch, \u0447\u044c\u0451 \u0438\u0441\u043f\u043e\u043b\u044c\u0437\u043e\u0432\u0430\u043d\u0438\u0435 \u0432\u0440\u0435\u043c\u0435\u043d\u043d\u043e\u0439 \u043f\u0430\u043c\u044f\u0442\u0438 \u0434\u0435\u043b\u0430\u0435\u0442 \u0435\u0433\u043e \u0431\u043e\u043b\u0435\u0435 \u043e\u0431\u0449\u0438\u043c, \u043d\u043e \u0437\u043d\u0430\u0447\u0438\u0442\u0435\u043b\u044c\u043d\u043e \u0431\u043e\u043b\u0435\u0435 \u043c\u0435\u0434\u043b\u0435\u043d\u043d\u044b\u043c (\u0441\u043c. \u043d\u0438\u0436\u0435).\u00a0<\/p>\n<blockquote>\n<p>Triton \u0443\u043f\u0440\u043e\u0449\u0430\u0435\u0442 \u0440\u0430\u0437\u0440\u0430\u0431\u043e\u0442\u043a\u0443 \u0441\u043f\u0435\u0446\u0438\u0430\u043b\u0438\u0437\u0438\u0440\u043e\u0432\u0430\u043d\u043d\u044b\u0445 \u044f\u0434\u0435\u0440, \u043a\u043e\u0442\u043e\u0440\u044b\u0435 \u043c\u043e\u0433\u0443\u0442 \u0431\u044b\u0442\u044c \u043d\u0430\u043c\u043d\u043e\u0433\u043e \u0431\u044b\u0441\u0442\u0440\u0435\u0435 \u044f\u0434\u0435\u0440 \u0432 \u0431\u0438\u0431\u043b\u0438\u043e\u0442\u0435\u043a\u0430\u0445 \u043e\u0431\u0449\u0435\u0433\u043e \u043d\u0430\u0437\u043d\u0430\u0447\u0435\u043d\u0438\u044f.<\/p>\n<\/blockquote>\n<figure class=\"full-width\"><img loading=\"lazy\" decoding=\"async\" src=\"https:\/\/habrastorage.org\/r\/w1560\/getpro\/habr\/upload_files\/b28\/278\/e8b\/b28278e8b29322e9c37f5d4ffb06f3e4.png\" alt=\"\u041f\u0440\u043e\u0438\u0437\u0432\u043e\u0434\u0438\u0442\u0435\u043b\u044c\u043d\u043e\u0441\u0442\u044c A100 fused softmax \u043f\u0440\u0438 M=4096\" title=\"\u041f\u0440\u043e\u0438\u0437\u0432\u043e\u0434\u0438\u0442\u0435\u043b\u044c\u043d\u043e\u0441\u0442\u044c A100 fused softmax \u043f\u0440\u0438 M=4096\" width=\"1300\" height=\"600\" data-src=\"https:\/\/habrastorage.org\/getpro\/habr\/upload_files\/b28\/278\/e8b\/b28278e8b29322e9c37f5d4ffb06f3e4.png\"\/><figcaption>\u041f\u0440\u043e\u0438\u0437\u0432\u043e\u0434\u0438\u0442\u0435\u043b\u044c\u043d\u043e\u0441\u0442\u044c A100 fused softmax \u043f\u0440\u0438 M=4096<\/figcaption><\/figure>\n<p>\u0411\u043e\u043b\u0435\u0435 \u043d\u0438\u0437\u043a\u0430\u044f \u043f\u0440\u043e\u0438\u0437\u0432\u043e\u0434\u0438\u0442\u0435\u043b\u044c\u043d\u043e\u0441\u0442\u044c Torch (v1.9) JIT \u043f\u043e\u0434\u0447\u0451\u0440\u043a\u0438\u0432\u0430\u0435\u0442 \u0441\u043b\u043e\u0436\u043d\u043e\u0441\u0442\u044c \u0430\u0432\u0442\u043e\u043c\u0430\u0442\u0438\u0447\u0435\u0441\u043a\u043e\u0439 \u0433\u0435\u043d\u0435\u0440\u0430\u0446\u0438\u0438 \u043a\u043e\u0434\u0430 CUDA \u0438\u0437 \u043f\u043e\u0441\u043b\u0435\u0434\u043e\u0432\u0430\u0442\u0435\u043b\u044c\u043d\u043e\u0441\u0442\u0435\u0439 \u0432\u044b\u0441\u043e\u043a\u043e\u0443\u0440\u043e\u0432\u043d\u0435\u0432\u044b\u0445 \u0442\u0435\u043d\u0437\u043e\u0440\u043d\u044b\u0445 \u043e\u043f\u0435\u0440\u0430\u0446\u0438\u0439. Fused softmax \u043d\u0430 Torch JIT:<\/p>\n<pre><code class=\"python\">@torch.jit.script def softmax(x):     x_max = x.max(dim=1)[0]     z = x - x_max[:, None]     numerator = torch.exp(x)     denominator = numerator.sum(dim=1)     return numerator \/ denominator[:, None]<\/code><\/pre>\n<h2>\u041c\u0430\u0442\u0440\u0438\u0447\u043d\u043e\u0435 \u0443\u043c\u043d\u043e\u0436\u0435\u043d\u0438\u0435<\/h2>\n<p>\u0412\u043e\u0437\u043c\u043e\u0436\u043d\u043e\u0441\u0442\u044c \u043f\u0438\u0441\u0430\u0442\u044c &#171;\u0441\u043b\u0438\u0442\u044b\u0435&#187; (fused) \u044f\u0434\u0440\u0430 (\u044f\u0434\u0440\u0430 \u0441 \u043e\u0431\u043e\u0431\u0449\u0435\u043d\u0438\u0435\u043c \u044d\u0442\u0430\u043f\u043e\u0432 \u0432\u044b\u0447\u0438\u0441\u043b\u0435\u043d\u0438\u0439) \u0434\u043b\u044f \u043e\u043f\u0435\u0440\u0430\u0446\u0438\u0439 \u0438 \u0441\u043e\u043a\u0440\u0430\u0449\u0435\u043d\u0438\u0439 \u043f\u043e \u044d\u043b\u0435\u043c\u0435\u043d\u0442\u0430\u043c \u0432\u0430\u0436\u043d\u0430, \u043d\u043e \u043d\u0435\u0434\u043e\u0441\u0442\u0430\u0442\u043e\u0447\u043d\u0430, \u0443\u0447\u0438\u0442\u044b\u0432\u0430\u044f, \u0447\u0442\u043e \u0432 \u043d\u0435\u0439\u0440\u043e\u043d\u043d\u044b\u0445 \u0441\u0435\u0442\u044f\u0445 \u0437\u0430\u0434\u0430\u0447\u0438 \u0443\u043c\u043d\u043e\u0436\u0435\u043d\u0438\u044f \u043c\u0430\u0442\u0440\u0438\u0446 \u0437\u0430\u043d\u0438\u043c\u0430\u044e\u0442 \u0432\u0430\u0436\u043d\u043e\u0435 \u043c\u0435\u0441\u0442\u043e.\u00a0\u041a\u0430\u043a \u043e\u043a\u0430\u0437\u0430\u043b\u043e\u0441\u044c, Triton \u043e\u0442\u043b\u0438\u0447\u043d\u043e \u043f\u043e\u0434\u0445\u043e\u0434\u0438\u0442 \u0438 \u0434\u043b\u044f \u043d\u0438\u0445, \u0434\u043e\u0441\u0442\u0438\u0433\u0430\u044f \u043f\u0438\u043a\u043e\u0432\u043e\u0439 \u043f\u0440\u043e\u0438\u0437\u0432\u043e\u0434\u0438\u0442\u0435\u043b\u044c\u043d\u043e\u0441\u0442\u0438 \u043c\u0435\u043d\u0435\u0435 \u0447\u0435\u043c \u0432 25 \u0441\u0442\u0440\u043e\u043a Python. \u0421 \u0434\u0440\u0443\u0433\u043e\u0439 \u0441\u0442\u043e\u0440\u043e\u043d\u044b, \u0440\u0435\u0430\u043b\u0438\u0437\u0430\u0446\u0438\u044f \u0447\u0435\u0433\u043e-\u0442\u043e \u043f\u043e\u0434\u043e\u0431\u043d\u043e\u0433\u043e \u0432 CUDA \u043f\u043e\u0442\u0440\u0435\u0431\u0443\u0435\u0442 <a href=\"https:\/\/github.com\/openai\/blocksparse\/blob\/master\/src\/matmul_op_gpu.cu\">\u0433\u043e\u0440\u0430\u0437\u0434\u043e \u0431\u043e\u043b\u044c\u0448\u0435 \u0443\u0441\u0438\u043b\u0438\u0439<\/a> \u0438 \u0434\u0430\u0436\u0435, \u0441\u043a\u043e\u0440\u0435\u0435 \u0432\u0441\u0435\u0433\u043e, \u043f\u0440\u0438\u0432\u0435\u0434\u0451\u0442 \u043a \u0441\u043d\u0438\u0436\u0435\u043d\u0438\u044e \u043f\u0440\u043e\u0438\u0437\u0432\u043e\u0434\u0438\u0442\u0435\u043b\u044c\u043d\u043e\u0441\u0442\u0438.<\/p>\n<p>\u0423\u043c\u043d\u043e\u0436\u0435\u043d\u0438\u0435 \u043c\u0430\u0442\u0440\u0438\u0446 \u043f\u0440\u0438 \u043f\u043e\u043c\u043e\u0449\u0438 Triton:<\/p>\n<pre><code class=\"python\">@triton.jit def matmul(A, B, C, M, N, K, stride_am, stride_ak,              stride_bk, stride_bn, stride_cm, stride_cn,             **META):     # extract metaparameters     BLOCK_M, GROUP_M = META['BLOCK_M'], META['GROUP_M']     BLOCK_N = META['BLOCK_N']     BLOCK_K = META['BLOCK_K']     # programs are grouped together to improve L2 hit rate     _pid_m = tl.program_id(0)     _pid_n = tl.program_id(1)     pid_m = _pid_m \/\/ GROUP_M     pid_n = (_pid_n * GROUP_M) + (_pid_m % GROUP_M)     # rm (resp. rn) denotes a range of indices     # for rows (resp. col) of C     rm = pid_m * BLOCK_M + tl.arange(0, BLOCK_M)     rn = pid_n * BLOCK_N + tl.arange(0, BLOCK_N)     # rk denotes a range of indices for columns      # (resp. rows) of A (resp. B)     rk = tl.arange(0, BLOCK_K)     # the memory addresses of elements in the first block of     # A and B can be computed using numpy-style broadcasting     A = A + (rm[:, None] * stride_am + rk[None, :] * stride_ak)     B = B + (rk [:, None] * stride_bk  + rn[None, :] * stride_bn)     # initialize and iteratively update accumulator     acc = tl.zeros((BLOCK_M, BLOCK_N), dtype=tl.float32)     for k in range(K, 0, -BLOCK_K):         a = tl.load(A)         b = tl.load(B)         # block level matrix multiplication         acc += tl.dot(a, b)         # increment pointers so that the next blocks of A and B         # are loaded during the next iteration         A += BLOCK_K * stride_ak         B += BLOCK_K * stride_bk     # fuse leaky ReLU if desired     # acc = tl.where(acc >= 0, acc, alpha * acc)     # write back result     C = C + (rm[:, None] * stride_cm + rn[None, :] * stride_cn)     mask = (rm[:, None] &lt; M) &amp; (rn[None, :] &lt; N)     tl.store(C, acc, mask=mask)<\/code><\/pre>\n<p>\u041e\u0434\u043d\u0438\u043c \u0438\u0437 \u0432\u0430\u0436\u043d\u044b\u0445 \u043f\u0440\u0435\u0438\u043c\u0443\u0449\u0435\u0441\u0442\u0432 \u0440\u0443\u043a\u043e\u043f\u0438\u0441\u043d\u044b\u0445 \u044f\u0434\u0435\u0440 \u043c\u0430\u0442\u0440\u0438\u0447\u043d\u043e\u0433\u043e \u0443\u043c\u043d\u043e\u0436\u0435\u043d\u0438\u044f \u044f\u0432\u043b\u044f\u0435\u0442\u0441\u044f \u0442\u043e, \u0447\u0442\u043e \u043e\u043d\u0438 \u043c\u043e\u0433\u0443\u0442 \u0431\u044b\u0442\u044c \u043d\u0430\u0441\u0442\u0440\u043e\u0435\u043d\u044b \u043f\u043e \u0436\u0435\u043b\u0430\u043d\u0438\u044e, \u0447\u0442\u043e\u0431\u044b \u0440\u0430\u0437\u043c\u0435\u0441\u0442\u0438\u0442\u044c \u0441\u043b\u0438\u0442\u044b\u0435 \u043f\u0440\u0435\u043e\u0431\u0440\u0430\u0437\u043e\u0432\u0430\u043d\u0438\u044f \u0438\u0445 \u0432\u0445\u043e\u0434\u043d\u044b\u0445 \u0434\u0430\u043d\u043d\u044b\u0445 (\u043d\u0430\u043f\u0440\u0438\u043c\u0435\u0440, \u0441\u043b\u0430\u0439\u0441\u0438\u043d\u0433\u0430) \u0438 \u0432\u044b\u0445\u043e\u0434\u043d\u044b\u0445 (\u043d\u0430\u043f\u0440\u0438\u043c\u0435\u0440, Leaky ReLU). \u0411\u0435\u0437 Triton \u043d\u0435\u0442\u0440\u0438\u0432\u0438\u0430\u043b\u044c\u043d\u044b\u0435 \u043c\u043e\u0434\u0438\u0444\u0438\u043a\u0430\u0446\u0438\u0438 \u044f\u0434\u0435\u0440 \u043c\u0430\u0442\u0440\u0438\u0447\u043d\u043e\u0433\u043e \u0443\u043c\u043d\u043e\u0436\u0435\u043d\u0438\u044f \u043d\u0435\u0434\u043e\u0441\u0442\u0443\u043f\u043d\u044b \u0440\u0430\u0437\u0440\u0430\u0431\u043e\u0442\u0447\u0438\u043a\u0430\u043c \u0431\u0435\u0437 \u0438\u0441\u043a\u043b\u044e\u0447\u0438\u0442\u0435\u043b\u044c\u043d\u043e\u0433\u043e \u043e\u043f\u044b\u0442\u0430 \u043f\u0440\u043e\u0433\u0440\u0430\u043c\u043c\u0438\u0440\u043e\u0432\u0430\u043d\u0438\u044f GPU.<\/p>\n<figure class=\"full-width\"><img loading=\"lazy\" decoding=\"async\" src=\"https:\/\/habrastorage.org\/r\/w1560\/getpro\/habr\/upload_files\/399\/527\/bdd\/399527bdd43a36eff87b53f6d0773769.png\" alt=\"V100 \u043f\u0440\u043e\u0438\u0437\u0432\u043e\u0434\u0438\u0442\u0435\u043b\u044c\u043d\u043e\u0441\u0442\u044c \u0442\u0435\u043d\u0437\u043e\u0440\u043d\u044b\u0445 \u044f\u0434\u0435\u0440 \u043f\u0440\u0438 \u0443\u043c\u043d\u043e\u0436\u0435\u043d\u0438\u0438 \u043c\u0430\u0442\u0440\u0438\u0446 \u0441 \u0441\u043e\u043e\u0442\u0432\u0435\u0442\u0441\u0442\u0432\u0443\u044e\u0449\u0438\u043c \u043e\u0431\u0440\u0430\u0437\u043e\u043c \u043d\u0430\u0441\u0442\u0440\u043e\u0435\u043d\u043d\u044b\u043c\u0438 \u0437\u043d\u0430\u0447\u0435\u043d\u0438\u044f\u043c\u0438 \u0434\u043b\u044f BLOCK_MM, BLOCK_NN, BLOCK_KK, GROUP_MM.\" title=\"V100 \u043f\u0440\u043e\u0438\u0437\u0432\u043e\u0434\u0438\u0442\u0435\u043b\u044c\u043d\u043e\u0441\u0442\u044c \u0442\u0435\u043d\u0437\u043e\u0440\u043d\u044b\u0445 \u044f\u0434\u0435\u0440 \u043f\u0440\u0438 \u0443\u043c\u043d\u043e\u0436\u0435\u043d\u0438\u0438 \u043c\u0430\u0442\u0440\u0438\u0446 \u0441 \u0441\u043e\u043e\u0442\u0432\u0435\u0442\u0441\u0442\u0432\u0443\u044e\u0449\u0438\u043c \u043e\u0431\u0440\u0430\u0437\u043e\u043c \u043d\u0430\u0441\u0442\u0440\u043e\u0435\u043d\u043d\u044b\u043c\u0438 \u0437\u043d\u0430\u0447\u0435\u043d\u0438\u044f\u043c\u0438 \u0434\u043b\u044f BLOCK_MM, BLOCK_NN, BLOCK_KK, GROUP_MM.\" width=\"1386\" height=\"820\" data-src=\"https:\/\/habrastorage.org\/getpro\/habr\/upload_files\/399\/527\/bdd\/399527bdd43a36eff87b53f6d0773769.png\"\/><figcaption>V100 \u043f\u0440\u043e\u0438\u0437\u0432\u043e\u0434\u0438\u0442\u0435\u043b\u044c\u043d\u043e\u0441\u0442\u044c \u0442\u0435\u043d\u0437\u043e\u0440\u043d\u044b\u0445 \u044f\u0434\u0435\u0440 \u043f\u0440\u0438 \u0443\u043c\u043d\u043e\u0436\u0435\u043d\u0438\u0438 \u043c\u0430\u0442\u0440\u0438\u0446 \u0441 \u0441\u043e\u043e\u0442\u0432\u0435\u0442\u0441\u0442\u0432\u0443\u044e\u0449\u0438\u043c \u043e\u0431\u0440\u0430\u0437\u043e\u043c \u043d\u0430\u0441\u0442\u0440\u043e\u0435\u043d\u043d\u044b\u043c\u0438 \u0437\u043d\u0430\u0447\u0435\u043d\u0438\u044f\u043c\u0438 \u0434\u043b\u044f BLOCK_MM, BLOCK_NN, BLOCK_KK, GROUP_MM.<\/figcaption><\/figure>\n<h2>\u041e\u0431\u0437\u043e\u0440 \u0430\u0440\u0445\u0438\u0442\u0435\u043a\u0442\u0443\u0440\u044b<\/h2>\n<p>\u0425\u043e\u0440\u043e\u0448\u0430\u044f \u043f\u0440\u043e\u0438\u0437\u0432\u043e\u0434\u0438\u0442\u0435\u043b\u044c\u043d\u043e\u0441\u0442\u044c Triton \u043e\u0431\u0443\u0441\u043b\u043e\u0432\u043b\u0435\u043d\u0430 \u043c\u043e\u0434\u0443\u043b\u044c\u043d\u043e\u0439 \u0430\u0440\u0445\u0438\u0442\u0435\u043a\u0442\u0443\u0440\u043e\u0439 \u0441\u0438\u0441\u0442\u0435\u043c\u044b, \u0432 \u0446\u0435\u043d\u0442\u0440\u0435 \u043a\u043e\u0442\u043e\u0440\u043e\u0439 \u2014 Triton-IR, \u0442\u043e \u0435\u0441\u0442\u044c \u043f\u0440\u043e\u043c\u0435\u0436\u0443\u0442\u043e\u0447\u043d\u043e\u0435 \u043f\u0440\u0435\u0434\u0441\u0442\u0430\u0432\u043b\u0435\u043d\u0438\u0435 \u043d\u0430 \u043e\u0441\u043d\u043e\u0432\u0435 LLVM, \u0433\u0434\u0435 \u043c\u043d\u043e\u0433\u043e\u043c\u0435\u0440\u043d\u044b\u0435 \u0431\u043b\u043e\u043a\u0438 \u0437\u043d\u0430\u0447\u0435\u043d\u0438\u0439 \u2014 \u044d\u0442\u043e \u0441\u0430\u043c\u043e\u0435 \u0432\u0430\u0436\u043d\u043e\u0435.<\/p>\n<figure class=\"full-width\"><img loading=\"lazy\" decoding=\"async\" src=\"https:\/\/habrastorage.org\/r\/w1560\/getpro\/habr\/upload_files\/660\/26e\/e5c\/66026ee5cbd9d9a3f4f39ad6d122b012.png\" alt=\"\u0412\u044b\u0441\u043e\u043a\u043e\u0443\u0440\u043e\u0432\u043d\u0435\u0432\u0430\u044f \u0430\u0440\u0445\u0438\u0442\u0435\u043a\u0442\u0443\u0440\u0430 Triton\" title=\"\u0412\u044b\u0441\u043e\u043a\u043e\u0443\u0440\u043e\u0432\u043d\u0435\u0432\u0430\u044f \u0430\u0440\u0445\u0438\u0442\u0435\u043a\u0442\u0443\u0440\u0430 Triton\" width=\"2648\" height=\"1392\" data-src=\"https:\/\/habrastorage.org\/getpro\/habr\/upload_files\/660\/26e\/e5c\/66026ee5cbd9d9a3f4f39ad6d122b012.png\"\/><figcaption>\u0412\u044b\u0441\u043e\u043a\u043e\u0443\u0440\u043e\u0432\u043d\u0435\u0432\u0430\u044f \u0430\u0440\u0445\u0438\u0442\u0435\u043a\u0442\u0443\u0440\u0430 Triton<\/figcaption><\/figure>\n<details class=\"spoiler\">\n<summary>\u041a\u043e\u0434 \u043d\u0430 \u0438\u0437\u043e\u0431\u0440\u0430\u0436\u0435\u043d\u0438\u044f\u0445 \u0432\u044b\u0448\u0435<\/summary>\n<div class=\"spoiler__content\">\n<pre><code class=\"python\">@jit def add(X, Y, Z, N):    pid = program_id(0)    idx= pid * 512 + arange(512)    mask = idx &lt; N    x = load(X + idx, mask=mask)    y = load(Y + idx, mask=mask)    store(Z + idx, x + y, mask=mask)<\/code><\/pre>\n<pre><code class=\"python\">def void add(i32* X .aligned(16) , i32* Y .aligned(16) , i32* Z .aligned(16) , i32 N .multipleof(2) ) { entry:   %0 = get_program_id[0] i32;   %1 = mul i32 %0, 512;   %3 = make_range[0 : 512] i32&lt;512>;   %4 = splat i32&lt;512> %1;   %6 = add i32&lt;512> %4, %3;   %9 = splat i32&lt;512> N;   %11 = icmp_slt i1&lt;512> %6, %9;   %14 = splat i32*&lt;512> X;   %16 = getelementptr i32*&lt;512> %14, %6;   %19 = broadcast i1&lt;512> %11;   %21 = splat i32&lt;512> undef;   %22 = masked_load i32&lt;512> %16, %19, %21;   %26 = splat i32*&lt;512> Y;   %28 = getelementptr i32*&lt;512> %26, %6;   %31 = broadcast i1&lt;512> %11;   %33 = splat i32&lt;512> undef;   %34 = masked_load i32&lt;512> %28, %31, %33;   %38 = splat i32*&lt;512> Z;   %40 = getelementptr i32*&lt;512> %38, %6;   %43 = add i32&lt;512> %22, %34;   %46 = broadcast i32&lt;512> %43;   %48 = broadcast i1&lt;512> %11;   masked_store void %40, %46, %48;   ret void; }<\/code><\/pre>\n<pre><code class=\"python\">.visible .entry add(     .param .u64 add_param_0, .param .u64 add_param_1,     .param .u64 add_param_2, .param .u32 add_param_3 ) .maxntid 128, 1, 1 {     .reg .pred     %p&lt;4>;     .reg .b32     %r&lt;18>;     .reg .b64     %rd&lt;8>;     ld.param.u64     %rd4, [add_param_0];     ld.param.u64     %rd5, [add_param_1];     mov.u32     %r13, %tid.x;     ld.param.u32     %r14, [add_param_3];     shl.b32     %r15, %r13, 2;     mov.u32     %r16, %ctaid.x;     mad.lo.s32     %r17, %r16, 512, %r15;     setp.ge.s32     %p3, %r17, %r14;     setp.lt.s32     %p1, %r17, %r14;     mul.wide.s32     %rd7, %r17, 4;     add.s64     %rd2, %rd4, %rd7;     @%p1 ld.global.cg.v4.b32 {%r5,%r6,%r7,%r8}, [ %rd2 + 0];     add.s64     %rd3, %rd5, %rd7;     @%p1 ld.global.cg.v4.b32 {%r9,%r10,%r11,%r12}, [ %rd3 + 0];     @%p3 bra     LBB0_2;     ld.param.u64     %rd6, [add_param_2];     add.s64     %rd1, %rd6, %rd7;     add.s32     %r1, %r5, %r9;     add.s32     %r2, %r6, %r10;     add.s32     %r3, %r7, %r11;     add.s32     %r4, %r8, %r12;     st.global.v4.u32     [%rd1], {%r1, %r2, %r3, %r4}; LBB0_2:     ret; }<\/code><\/pre>\n<\/div>\n<\/details>\n<p>\u0414\u0435\u043a\u043e\u0440\u0430\u0442\u043e\u0440 @triton.jit \u0440\u0430\u0431\u043e\u0442\u0430\u0435\u0442, \u043f\u0440\u043e\u0441\u043c\u0430\u0442\u0440\u0438\u0432\u0430\u044f AST \u043f\u0440\u0435\u0434\u043e\u0441\u0442\u0430\u0432\u043b\u0435\u043d\u043d\u043e\u0439 \u0444\u0443\u043d\u043a\u0446\u0438\u0438 Python, \u0447\u0442\u043e\u0431\u044b \u043d\u0430 \u043b\u0435\u0442\u0443 \u0441\u0433\u0435\u043d\u0435\u0440\u0438\u0440\u043e\u0432\u0430\u0442\u044c Triton-IR, \u0438\u0441\u043f\u043e\u043b\u044c\u0437\u0443\u044f \u043e\u0431\u044b\u0447\u043d\u044b\u0439 \u0430\u043b\u0433\u043e\u0440\u0438\u0442\u043c \u043f\u043e\u0441\u0442\u0440\u043e\u0435\u043d\u0438\u044f SSA<a href=\"https:\/\/openai.com\/blog\/triton\/#rf8\"><sup>8<\/sup><\/a>.<\/p>\n<blockquote>\n<p>\u041f\u043e\u043b\u0443\u0447\u0435\u043d\u043d\u044b\u0439 IR-\u043a\u043e\u0434 \u0437\u0430\u0442\u0435\u043c \u0443\u043f\u0440\u043e\u0449\u0430\u0435\u0442\u0441\u044f, \u043e\u043f\u0442\u0438\u043c\u0438\u0437\u0438\u0440\u0443\u0435\u0442\u0441\u044f \u0438 \u0430\u0432\u0442\u043e\u043c\u0430\u0442\u0438\u0447\u0435\u0441\u043a\u0438 \u0440\u0430\u0441\u043f\u0430\u0440\u0430\u043b\u043b\u0435\u043b\u0438\u0432\u0430\u0435\u0442\u0441\u044f \u043d\u0430\u0448\u0438\u043c \u0431\u044d\u043a\u0435\u043d\u0434\u043e\u043c \u043a\u043e\u043c\u043f\u0438\u043b\u044f\u0442\u043e\u0440\u0430, \u043f\u043e\u0441\u043b\u0435 \u0447\u0435\u0433\u043e \u043f\u0440\u0435\u043e\u0431\u0440\u0430\u0437\u0443\u0435\u0442\u0441\u044f \u0432 \u0432\u044b\u0441\u043e\u043a\u043e\u043a\u0430\u0447\u0435\u0441\u0442\u0432\u0435\u043d\u043d\u044b\u0439 LLVM-IR \u0438, \u0432 \u043a\u043e\u043d\u0435\u0447\u043d\u043e\u043c \u0441\u0447\u0451\u0442\u0435, \u0432 PTX \u0434\u043b\u044f \u0432\u044b\u043f\u043e\u043b\u043d\u0435\u043d\u0438\u044f \u043d\u0430 \u043d\u043e\u0432\u0435\u0439\u0448\u0438\u0445 \u0433\u0440\u0430\u0444\u0438\u0447\u0435\u0441\u043a\u0438\u0445 \u043f\u0440\u043e\u0446\u0435\u0441\u0441\u043e\u0440\u0430\u0445 NVIDIA.\u00a0<\/p>\n<\/blockquote>\n<p>CPU \u0438 GPU AMD \u0441\u0435\u0439\u0447\u0430\u0441 \u043d\u0435 \u043f\u043e\u0434\u0434\u0435\u0440\u0436\u0438\u0432\u0430\u044e\u0442\u0441\u044f, \u043d\u043e \u043c\u044b \u043f\u0440\u0438\u0432\u0435\u0442\u0441\u0442\u0432\u0443\u0435\u043c \u0443\u0441\u0442\u0440\u0430\u043d\u044f\u044e\u0449\u0438\u0439 \u044d\u0442\u043e \u043e\u0433\u0440\u0430\u043d\u0438\u0447\u0435\u043d\u0438\u0435 \u043a\u043e\u0434.<\/p>\n<h2>\u0411\u044d\u043a\u0435\u043d\u0434 \u043a\u043e\u043c\u043f\u0438\u043b\u044f\u0442\u043e\u0440\u0430<\/h2>\n<p>\u041c\u044b \u043e\u0431\u043d\u0430\u0440\u0443\u0436\u0438\u043b\u0438, \u0447\u0442\u043e \u0438\u0441\u043f\u043e\u043b\u044c\u0437\u043e\u0432\u0430\u043d\u0438\u0435 \u0431\u043b\u043e\u0447\u043d\u044b\u0445 \u043f\u0440\u0435\u0434\u0441\u0442\u0430\u0432\u043b\u0435\u043d\u0438\u0439 \u043f\u0440\u043e\u0433\u0440\u0430\u043c\u043c Triton-IR \u043f\u043e\u0437\u0432\u043e\u043b\u044f\u0435\u0442 \u043d\u0430\u0448\u0435\u043c\u0443 \u043a\u043e\u043c\u043f\u0438\u043b\u044f\u0442\u043e\u0440\u0443 \u0430\u0432\u0442\u043e\u043c\u0430\u0442\u0438\u0447\u0435\u0441\u043a\u0438 \u0432\u044b\u043f\u043e\u043b\u043d\u044f\u0442\u044c \u0448\u0438\u0440\u043e\u043a\u0438\u0439 \u0441\u043f\u0435\u043a\u0442\u0440 \u0432\u0430\u0436\u043d\u044b\u0445 \u043e\u043f\u0442\u0438\u043c\u0438\u0437\u0430\u0446\u0438\u0439.\u00a0<\/p>\n<p>\u0414\u0430\u043d\u043d\u044b\u0435 \u043c\u043e\u0433\u0443\u0442 \u0430\u0432\u0442\u043e\u043c\u0430\u0442\u0438\u0447\u0435\u0441\u043a\u0438 \u0443\u043a\u043b\u0430\u0434\u044b\u0432\u0430\u0442\u044c\u0441\u044f \u0432 \u043e\u0431\u0449\u0443\u044e \u043f\u0430\u043c\u044f\u0442\u044c, \u043f\u0440\u043e\u0441\u043c\u0430\u0442\u0440\u0438\u0432\u0430\u044f \u043e\u043f\u0435\u0440\u0430\u043d\u0434\u044b \u0432\u044b\u0447\u0438\u0441\u043b\u0438\u0442\u0435\u043b\u044c\u043d\u043e \u0438\u043d\u0442\u0435\u043d\u0441\u0438\u0432\u043d\u044b\u0445 \u043e\u043f\u0435\u0440\u0430\u0446\u0438\u0439 \u043d\u0430 \u0443\u0440\u043e\u0432\u043d\u0435 \u0431\u043b\u043e\u043a\u043e\u0432 (\u043d\u0430\u043f\u0440\u0438\u043c\u0435\u0440, tl.dot), \u0438 \u0432\u044b\u0434\u0435\u043b\u044f\u0442\u044c\u0441\u044f\/\u0441\u0438\u043d\u0445\u0440\u043e\u043d\u0438\u0437\u0438\u0440\u043e\u0432\u0430\u0442\u044c\u0441\u044f \u0441 \u043f\u043e\u043c\u043e\u0449\u044c\u044e \u0441\u0442\u0430\u043d\u0434\u0430\u0440\u0442\u043d\u044b\u0445 \u043c\u0435\u0442\u043e\u0434\u043e\u0432 \u0430\u043d\u0430\u043b\u0438\u0437\u0430 \u0436\u0438\u0432\u0443\u0447\u0435\u0441\u0442\u0438.<\/p>\n<figure class=\"full-width\"><img loading=\"lazy\" decoding=\"async\" src=\"https:\/\/habrastorage.org\/r\/w1560\/getpro\/habr\/upload_files\/324\/08e\/3d1\/32408e3d1234322b38749ee91bee2032.png\" alt=\"\u041a\u043e\u043c\u043f\u0438\u043b\u044f\u0442\u043e\u0440 Triton \u0432\u044b\u0434\u0435\u043b\u044f\u0435\u0442 \u043e\u0431\u0449\u0443\u044e \u043f\u0430\u043c\u044f\u0442\u044c, \u0430\u043d\u0430\u043b\u0438\u0437\u0438\u0440\u0443\u044f \u0436\u0438\u0432\u043e\u0439 \u0434\u0438\u0430\u043f\u0430\u0437\u043e\u043d \u0431\u043b\u043e\u0447\u043d\u044b\u0445 \u043f\u0435\u0440\u0435\u043c\u0435\u043d\u043d\u044b\u0445, \u0437\u0430\u0434\u0435\u0439\u0441\u0442\u0432\u043e\u0432\u0430\u043d\u043d\u044b\u0445 \u0432 \u0432\u044b\u0447\u0438\u0441\u043b\u0438\u0442\u0435\u043b\u044c\u043d\u043e \u0438\u043d\u0442\u0435\u043d\u0441\u0438\u0432\u043d\u044b\u0445 \u043e\u043f\u0435\u0440\u0430\u0446\u0438\u044f\u0445.\" title=\"\u041a\u043e\u043c\u043f\u0438\u043b\u044f\u0442\u043e\u0440 Triton \u0432\u044b\u0434\u0435\u043b\u044f\u0435\u0442 \u043e\u0431\u0449\u0443\u044e \u043f\u0430\u043c\u044f\u0442\u044c, \u0430\u043d\u0430\u043b\u0438\u0437\u0438\u0440\u0443\u044f \u0436\u0438\u0432\u043e\u0439 \u0434\u0438\u0430\u043f\u0430\u0437\u043e\u043d \u0431\u043b\u043e\u0447\u043d\u044b\u0445 \u043f\u0435\u0440\u0435\u043c\u0435\u043d\u043d\u044b\u0445, \u0437\u0430\u0434\u0435\u0439\u0441\u0442\u0432\u043e\u0432\u0430\u043d\u043d\u044b\u0445 \u0432 \u0432\u044b\u0447\u0438\u0441\u043b\u0438\u0442\u0435\u043b\u044c\u043d\u043e \u0438\u043d\u0442\u0435\u043d\u0441\u0438\u0432\u043d\u044b\u0445 \u043e\u043f\u0435\u0440\u0430\u0446\u0438\u044f\u0445.\" width=\"1416\" height=\"498\" data-src=\"https:\/\/habrastorage.org\/getpro\/habr\/upload_files\/324\/08e\/3d1\/32408e3d1234322b38749ee91bee2032.png\"\/><figcaption>\u041a\u043e\u043c\u043f\u0438\u043b\u044f\u0442\u043e\u0440 Triton \u0432\u044b\u0434\u0435\u043b\u044f\u0435\u0442 \u043e\u0431\u0449\u0443\u044e \u043f\u0430\u043c\u044f\u0442\u044c, \u0430\u043d\u0430\u043b\u0438\u0437\u0438\u0440\u0443\u044f \u0436\u0438\u0432\u043e\u0439 \u0434\u0438\u0430\u043f\u0430\u0437\u043e\u043d \u0431\u043b\u043e\u0447\u043d\u044b\u0445 \u043f\u0435\u0440\u0435\u043c\u0435\u043d\u043d\u044b\u0445, \u0437\u0430\u0434\u0435\u0439\u0441\u0442\u0432\u043e\u0432\u0430\u043d\u043d\u044b\u0445 \u0432 \u0432\u044b\u0447\u0438\u0441\u043b\u0438\u0442\u0435\u043b\u044c\u043d\u043e \u0438\u043d\u0442\u0435\u043d\u0441\u0438\u0432\u043d\u044b\u0445 \u043e\u043f\u0435\u0440\u0430\u0446\u0438\u044f\u0445.<\/figcaption><\/figure>\n<p>\u0421 \u0434\u0440\u0443\u0433\u043e\u0439 \u0441\u0442\u043e\u0440\u043e\u043d\u044b, \u043f\u0440\u043e\u0433\u0440\u0430\u043c\u043c\u044b Triton \u043c\u043e\u0433\u0443\u0442 \u0431\u044b\u0442\u044c \u044d\u0444\u0444\u0435\u043a\u0442\u0438\u0432\u043d\u043e \u0438 \u0430\u0432\u0442\u043e\u043c\u0430\u0442\u0438\u0447\u0435\u0441\u043a\u0438 \u0440\u0430\u0441\u043f\u0430\u0440\u0430\u043b\u043b\u0435\u043b\u0435\u043d\u044b \u043a\u0430\u043a (1) \u043c\u0435\u0436\u0434\u0443 SM \u2014 \u043f\u0443\u0442\u0451\u043c \u043e\u0434\u043d\u043e\u0432\u0440\u0435\u043c\u0435\u043d\u043d\u043e\u0433\u043e \u0432\u044b\u043f\u043e\u043b\u043d\u0435\u043d\u0438\u044f \u0440\u0430\u0437\u043b\u0438\u0447\u043d\u044b\u0445 \u044d\u043a\u0437\u0435\u043c\u043f\u043b\u044f\u0440\u043e\u0432 \u044f\u0434\u0440\u0430, \u0442\u0430\u043a \u0438 (2) \u0432\u043d\u0443\u0442\u0440\u0438 SM \u2014 \u043f\u0443\u0442\u0451\u043c \u0430\u043d\u0430\u043b\u0438\u0437\u0430 \u043f\u0440\u043e\u0441\u0442\u0440\u0430\u043d\u0441\u0442\u0432\u0430 \u0438\u0442\u0435\u0440\u0430\u0446\u0438\u0439 \u043a\u0430\u0436\u0434\u043e\u0439 \u043e\u043f\u0435\u0440\u0430\u0446\u0438\u0438 \u043d\u0430 \u0443\u0440\u043e\u0432\u043d\u0435 \u0431\u043b\u043e\u043a\u043e\u0432 \u0438 \u0430\u0434\u0435\u043a\u0432\u0430\u0442\u043d\u043e\u0433\u043e \u0440\u0430\u0437\u0431\u0438\u0435\u043d\u0438\u044f \u0435\u0433\u043e \u043d\u0430 \u0440\u0430\u0437\u043b\u0438\u0447\u043d\u044b\u0435 SIMD-\u0431\u043b\u043e\u043a\u0438, \u043a\u0430\u043a \u043f\u043e\u043a\u0430\u0437\u0430\u043d\u043e \u043d\u0438\u0436\u0435.<\/p>\n<figure class=\"full-width\"><img loading=\"lazy\" decoding=\"async\" src=\"https:\/\/habrastorage.org\/r\/w1560\/getpro\/habr\/upload_files\/6f6\/94a\/9f2\/6f694a9f2eada04d308a604f8b0a6dae.png\" alt=\"\u0410\u0432\u0442\u043e\u043c\u0430\u0442\u0438\u0447\u0435\u0441\u043a\u043e\u0435 \u0440\u0430\u0441\u043f\u0430\u0440\u0430\u043b\u043b\u0435\u043b\u0438\u0432\u0430\u043d\u0438\u0435 \u0432 Triton. \u041a\u0430\u0436\u0434\u0430\u044f \u043e\u043f\u0435\u0440\u0430\u0446\u0438\u044f \u043d\u0430 \u0443\u0440\u043e\u0432\u043d\u0435 \u0431\u043b\u043e\u043a\u043e\u0432 \u043e\u043f\u0440\u0435\u0434\u0435\u043b\u044f\u0435\u0442 \u0431\u043b\u043e\u0447\u043d\u043e\u0435 \u043f\u0440\u043e\u0441\u0442\u0440\u0430\u043d\u0441\u0442\u0432\u043e \u0438\u0442\u0435\u0440\u0430\u0446\u0438\u0439, \u043a\u043e\u0442\u043e\u0440\u043e\u0435 \u0430\u0432\u0442\u043e\u043c\u0430\u0442\u0438\u0447\u0435\u0441\u043a\u0438 \u0440\u0430\u0441\u043f\u0430\u0440\u0430\u043b\u043b\u0435\u043b\u0438\u0432\u0430\u0435\u0442\u0441\u044f \u0434\u043b\u044f \u0438\u0441\u043f\u043e\u043b\u044c\u0437\u043e\u0432\u0430\u043d\u0438\u044f \u0440\u0435\u0441\u0443\u0440\u0441\u043e\u0432, \u0434\u043e\u0441\u0442\u0443\u043f\u043d\u044b\u0445 \u043d\u0430 \u043f\u043e\u0442\u043e\u043a\u043e\u0432\u043e\u043c \u043c\u0443\u043b\u044c\u0442\u0438\u043f\u0440\u043e\u0446\u0435\u0441\u0441\u043e\u0440\u0435 (SM).\" title=\"\u0410\u0432\u0442\u043e\u043c\u0430\u0442\u0438\u0447\u0435\u0441\u043a\u043e\u0435 \u0440\u0430\u0441\u043f\u0430\u0440\u0430\u043b\u043b\u0435\u043b\u0438\u0432\u0430\u043d\u0438\u0435 \u0432 Triton. \u041a\u0430\u0436\u0434\u0430\u044f \u043e\u043f\u0435\u0440\u0430\u0446\u0438\u044f \u043d\u0430 \u0443\u0440\u043e\u0432\u043d\u0435 \u0431\u043b\u043e\u043a\u043e\u0432 \u043e\u043f\u0440\u0435\u0434\u0435\u043b\u044f\u0435\u0442 \u0431\u043b\u043e\u0447\u043d\u043e\u0435 \u043f\u0440\u043e\u0441\u0442\u0440\u0430\u043d\u0441\u0442\u0432\u043e \u0438\u0442\u0435\u0440\u0430\u0446\u0438\u0439, \u043a\u043e\u0442\u043e\u0440\u043e\u0435 \u0430\u0432\u0442\u043e\u043c\u0430\u0442\u0438\u0447\u0435\u0441\u043a\u0438 \u0440\u0430\u0441\u043f\u0430\u0440\u0430\u043b\u043b\u0435\u043b\u0438\u0432\u0430\u0435\u0442\u0441\u044f \u0434\u043b\u044f \u0438\u0441\u043f\u043e\u043b\u044c\u0437\u043e\u0432\u0430\u043d\u0438\u044f \u0440\u0435\u0441\u0443\u0440\u0441\u043e\u0432, \u0434\u043e\u0441\u0442\u0443\u043f\u043d\u044b\u0445 \u043d\u0430 \u043f\u043e\u0442\u043e\u043a\u043e\u0432\u043e\u043c \u043c\u0443\u043b\u044c\u0442\u0438\u043f\u0440\u043e\u0446\u0435\u0441\u0441\u043e\u0440\u0435 (SM).\" width=\"2194\" height=\"856\" data-src=\"https:\/\/habrastorage.org\/getpro\/habr\/upload_files\/6f6\/94a\/9f2\/6f694a9f2eada04d308a604f8b0a6dae.png\"\/><figcaption>\u0410\u0432\u0442\u043e\u043c\u0430\u0442\u0438\u0447\u0435\u0441\u043a\u043e\u0435 \u0440\u0430\u0441\u043f\u0430\u0440\u0430\u043b\u043b\u0435\u043b\u0438\u0432\u0430\u043d\u0438\u0435 \u0432 Triton. \u041a\u0430\u0436\u0434\u0430\u044f \u043e\u043f\u0435\u0440\u0430\u0446\u0438\u044f \u043d\u0430 \u0443\u0440\u043e\u0432\u043d\u0435 \u0431\u043b\u043e\u043a\u043e\u0432 \u043e\u043f\u0440\u0435\u0434\u0435\u043b\u044f\u0435\u0442 \u0431\u043b\u043e\u0447\u043d\u043e\u0435 \u043f\u0440\u043e\u0441\u0442\u0440\u0430\u043d\u0441\u0442\u0432\u043e \u0438\u0442\u0435\u0440\u0430\u0446\u0438\u0439, \u043a\u043e\u0442\u043e\u0440\u043e\u0435 \u0430\u0432\u0442\u043e\u043c\u0430\u0442\u0438\u0447\u0435\u0441\u043a\u0438 \u0440\u0430\u0441\u043f\u0430\u0440\u0430\u043b\u043b\u0435\u043b\u0438\u0432\u0430\u0435\u0442\u0441\u044f \u0434\u043b\u044f \u0438\u0441\u043f\u043e\u043b\u044c\u0437\u043e\u0432\u0430\u043d\u0438\u044f \u0440\u0435\u0441\u0443\u0440\u0441\u043e\u0432, \u0434\u043e\u0441\u0442\u0443\u043f\u043d\u044b\u0445 \u043d\u0430 \u043f\u043e\u0442\u043e\u043a\u043e\u0432\u043e\u043c \u043c\u0443\u043b\u044c\u0442\u0438\u043f\u0440\u043e\u0446\u0435\u0441\u0441\u043e\u0440\u0435 (SM).<\/figcaption><\/figure>\n<details class=\"spoiler\">\n<summary>\u041a\u043e\u0434 \u0438 \u043a\u0440\u0443\u043f\u043d\u044b\u0435 \u0438\u043b\u043b\u044e\u0441\u0442\u0440\u0430\u0446\u0438\u0438<\/summary>\n<div class=\"spoiler__content\">\n<pre><code class=\"python\">S1 float A[4,4] = ... S2 float B[4,4] = ... S3 float C[4,4] = A + B<\/code><\/pre>\n<pre><code class=\"python\">S1 half A[4,2] = ... S2 half B[2,2] = ... S3 float C[4,2] = dot(A,B)<\/code><\/pre>\n<p>\u0412\u044b\u0448\u0435 \u2014 \u043e\u043f\u0440\u0435\u0434\u0435\u043b\u0435\u043d\u0438\u0435 \u043f\u0440\u043e\u0433\u0440\u0430\u043c\u043c\u044b Triton <em>P<\/em>, \u0441\u043e\u0441\u0442\u043e\u044f\u0449\u0435\u0439 \u0438\u0437 \u0442\u0440\u0451\u0445 \u0443\u0442\u0432\u0435\u0440\u0436\u0434\u0435\u043d\u0438\u0439 S1, S2, S3.<\/p>\n<figure class=\"\"><img loading=\"lazy\" decoding=\"async\" src=\"https:\/\/habrastorage.org\/r\/w1560\/getpro\/habr\/upload_files\/f7b\/a66\/717\/f7ba66717d87d8122029cae0b91a702c.png\" width=\"240\" height=\"120\" data-src=\"https:\/\/habrastorage.org\/getpro\/habr\/upload_files\/f7b\/a66\/717\/f7ba66717d87d8122029cae0b91a702c.png\"\/><figcaption><\/figcaption><\/figure>\n<figure class=\"\"><img loading=\"lazy\" decoding=\"async\" src=\"https:\/\/habrastorage.org\/r\/w1560\/getpro\/habr\/upload_files\/174\/1cd\/f07\/1741cdf07e80091233484d493826ee1f.png\" width=\"240\" height=\"120\" data-src=\"https:\/\/habrastorage.org\/getpro\/habr\/upload_files\/174\/1cd\/f07\/1741cdf07e80091233484d493826ee1f.png\"\/><figcaption><\/figcaption><\/figure>\n<p>\u041f\u0440\u043e\u0441\u0442\u0440\u0430\u043d\u0441\u0442\u0432\u043e \u0438\u0442\u0435\u0440\u0430\u0446\u0438\u0438 S3.<\/p>\n<figure class=\"\"><img loading=\"lazy\" decoding=\"async\" src=\"https:\/\/habrastorage.org\/r\/w1560\/getpro\/habr\/upload_files\/94d\/8ca\/bab\/94d8cababd377b5b20637eddc023e759.png\" width=\"241\" height=\"300\" data-src=\"https:\/\/habrastorage.org\/getpro\/habr\/upload_files\/94d\/8ca\/bab\/94d8cababd377b5b20637eddc023e759.png\"\/><figcaption><\/figcaption><\/figure>\n<p>\u041e\u0442\u043e\u0431\u0440\u0430\u0436\u0435\u043d\u0438\u0435 S3 \u043d\u0430 \u043f\u043e\u0442\u043e\u043a\u043e\u0432\u044b\u0439 \u043c\u0443\u043b\u044c\u0442\u0438\u043f\u0440\u043e\u0446\u0435\u0441\u0441\u043e\u0440 (SM).<\/p>\n<figure class=\"\"><img loading=\"lazy\" decoding=\"async\" src=\"https:\/\/habrastorage.org\/r\/w1560\/getpro\/habr\/upload_files\/bdb\/75c\/474\/bdb75c47433f5523c10b1e450f46df09.png\" width=\"241\" height=\"300\" data-src=\"https:\/\/habrastorage.org\/getpro\/habr\/upload_files\/bdb\/75c\/474\/bdb75c47433f5523c10b1e450f46df09.png\"\/><figcaption><\/figcaption><\/figure>\n<p>\u041e\u0442\u043e\u0431\u0440\u0430\u0436\u0435\u043d\u0438\u0435 <em>P<\/em> \u043d\u0430 GPU.<\/p>\n<\/div>\n<\/details>\n<h2>\u0417\u0430\u043a\u043b\u044e\u0447\u0435\u043d\u0438\u0435<\/h2>\n<p>\u0412 OpenAI \u043d\u0430\u043c\u0435\u0440\u0435\u043d\u044b, \u0447\u0442\u043e\u0431\u044b \u043f\u0440\u043e\u0435\u043a\u0442 Triton \u0443\u043f\u0440\u0430\u0432\u043b\u044f\u043b\u0441\u044f \u0441\u043e\u043e\u0431\u0449\u0435\u0441\u0442\u0432\u043e\u043c. \u0421\u0432\u043e\u0431\u043e\u0434\u043d\u043e \u0444\u043e\u0440\u043a\u0430\u0439\u0442\u0435 \u0440\u0435\u043f\u043e\u0437\u0438\u0442\u043e\u0440\u0438\u0439 \u043d\u0430 <a href=\"https:\/\/github.com\/openai\/triton\">GitHub<\/a>. \u0414\u043e\u043a\u0443\u043c\u0435\u043d\u0442\u0430\u0446\u0438\u044f Triton \u043d\u0430\u0445\u043e\u0434\u0438\u0442\u0441\u044f <a href=\"https:\/\/triton-lang.org\/\">\u0437\u0434\u0435\u0441\u044c<\/a>. \u041f\u043e\u043f\u0440\u043e\u0431\u043e\u0432\u0430\u0442\u044c Triton \u0432\u044b \u0441\u043c\u043e\u0436\u0435\u0442\u0435 \u043d\u0430 \u043d\u0430\u0448\u0438\u0445 \u043a\u0443\u0440\u0441\u0430\u0445:<\/p>\n<ul>\n<li>\n<p><a href=\"https:\/\/skillfactory.ru\/machine-learning-i-deep-learning?utm_source=habr&amp;utm_medium=habr&amp;utm_campaign=article&amp;utm_content=data-science_mldl_230921&amp;utm_term=conc\">\u041a\u0443\u0440\u0441 \u043f\u043e Machine Learning \u0438 Deep Learning (6 \u043c\u0435\u0441\u044f\u0446\u0435\u0432)<\/a><\/p>\n<\/li>\n<li>\n<p><a href=\"https:\/\/skillfactory.ru\/data-scientist-pro?utm_source=habr&amp;utm_medium=habr&amp;utm_campaign=article&amp;utm_content=data-science_dspr_230921&amp;utm_term=conc\">\u041f\u0440\u043e\u0444\u0435\u0441\u0441\u0438\u044f Data Scientist (24 \u043c\u0435\u0441\u044f\u0446\u0430)<\/a><\/p>\n<\/li>\n<\/ul>\n<p>\u0410 \u0447\u0442\u043e\u0431\u044b \u0443\u0432\u0438\u0434\u0435\u0442\u044c, \u043a\u0430\u043a \u043c\u044b \u0433\u043e\u0442\u043e\u0432\u0438\u043c \u0441\u043f\u0435\u0446\u0438\u0430\u043b\u0438\u0441\u0442\u043e\u0432 \u0432 \u0434\u0440\u0443\u0433\u0438\u0445 \u043d\u0430\u043f\u0440\u0430\u0432\u043b\u0435\u043d\u0438\u044f\u0445 IT, \u0432\u044b \u043c\u043e\u0436\u0435\u0442\u0435 \u043f\u0435\u0440\u0435\u0439\u0442\u0438 \u043d\u0430 \u0441\u0442\u0440\u0430\u043d\u0438\u0446\u044b <a href=\"https:\/\/skillfactory.ru\/catalogue?utm_source=habr&amp;utm_medium=habr&amp;utm_campaign=article&amp;utm_content=sf_allcourses_230921&amp;utm_term=conc\">\u0438\u0437 \u043a\u0430\u0442\u0430\u043b\u043e\u0433\u0430<\/a>.<\/p>\n<details class=\"spoiler\">\n<summary>\u0421\u0441\u044b\u043b\u043a\u0438<\/summary>\n<div class=\"spoiler__content\">\n<ul>\n<li>\n<p>Gray, S. (2017). <a href=\"https:\/\/github.com\/NervanaSystems\/maxas\/wiki\/SGEMM\">SGEMM Walkthrough<\/a>. <em>URL<\/em>. <a href=\"#rfref1\">\u21a9\ufe0e<\/a><\/p>\n<\/li>\n<li>\n<p>Kerr, A. (2020). <a href=\"https:\/\/developer.nvidia.com\/gtc\/2020\/video\/s21745-vid\">Developing CUDA kernels to push Tensor Cores to the Absolute Limit on NVIDIA A100<\/a>. <a href=\"#rfref2\">\u21a9\ufe0e<\/a><\/p>\n<\/li>\n<li>\n<p>Yan, D., Wang, W., &amp; Chu, X. (2020, May). <a href=\"https:\/\/ieeexplore.ieee.org\/abstract\/document\/9139835\/\">Demystifying tensor cores to optimize half-precision matrix multiply<\/a>. In <em>2020 IEEE International Parallel and Distributed Processing Symposium (IPDPS)<\/em>.\u00a0IEEE. <a href=\"#rfref3\">\u21a9\ufe0e<\/a><\/p>\n<\/li>\n<li>\n<p><a href=\"https:\/\/github.com\/NVIDIA\/cutlass\">NVIDIA\u00a0CUTLASS<\/a> <a href=\"#rfref4\">\u21a9\ufe0e<\/a><\/p>\n<\/li>\n<li>\n<p><a href=\"https:\/\/github.com\/apache\/tvm\">Apache\u00a0TVM<\/a> <a href=\"#rfref5\">\u21a9\ufe0e<\/a><\/p>\n<\/li>\n<li>\n<p>Tillet, P., Kung, H. T., &amp; Cox, D. (2019, June). <a href=\"https:\/\/dl.acm.org\/doi\/abs\/10.1145\/3315508.3329973\">Triton: an intermediate language and compiler for tiled neural network computations<\/a>. In <em>Proceedings of the 3rd ACM SIGPLAN International Workshop on Machine Learning and Programming Languages<\/em> (pp.\u00a010-19). <a href=\"#rfref6\">\u21a9\ufe0e<\/a><\/p>\n<\/li>\n<li>\n<p>Lin, Y. &amp; Grover, V. (2018). <a href=\"https:\/\/developer.nvidia.com\/blog\/using-cuda-warp-level-primitives\/\">Using CUDA Warp-Level Primitives<\/a>. <em>URL <\/em><a href=\"https:\/\/developer.nvidia.com\/blog\/using-cuda-warp-level-primitives\/\"><em>https:\/\/developer.nvidia.com\/blog\/using-cuda-warp-level-primitives\/<\/em><\/a>. <a href=\"#rfref7\">\u21a9\ufe0e<\/a><\/p>\n<\/li>\n<li>\n<p>Braun, M., Buchwald, S., Hack, S., Lei\u00dfa, R., Mallon, C., &amp; Zwinkau, A. (2013, March). <a href=\"https:\/\/link.springer.com\/chapter\/10.1007\/978-3-642-37051-9_6\">Simple and efficient construction of static single assignment form<\/a>. In <em>International Conference on Compiler Construction<\/em> (pp. 102-122). Springer, Berlin,\u00a0Heidelberg. <a href=\"#rfref8\">\u21a9\ufe0e<\/a><\/p>\n<\/li>\n<\/ul>\n<\/div>\n<\/details>\n<figure class=\"full-width\"><img loading=\"lazy\" decoding=\"async\" src=\"https:\/\/habrastorage.org\/r\/w1560\/getpro\/habr\/upload_files\/f18\/c7e\/c56\/f18c7ec567c5d0e6b3c137a013aca42f.png\" width=\"1000\" height=\"200\" data-src=\"https:\/\/habrastorage.org\/getpro\/habr\/upload_files\/f18\/c7e\/c56\/f18c7ec567c5d0e6b3c137a013aca42f.png\"\/><figcaption><\/figcaption><\/figure>\n<details class=\"spoiler\">\n<summary>\u041f\u0440\u043e\u0444\u0435\u0441\u0441\u0438\u0438 \u0438 \u043a\u0443\u0440\u0441\u044b<\/summary>\n<div class=\"spoiler__content\">\n<p><strong>Data Science \u0438 Machine Learning<\/strong><\/p>\n<ul>\n<li>\n<p><a href=\"https:\/\/skillfactory.ru\/data-scientist-pro?utm_source=habr&amp;utm_medium=habr&amp;utm_campaign=article&amp;utm_content=data-science_dspr_230921&amp;utm_term=cat\">\u041f\u0440\u043e\u0444\u0435\u0441\u0441\u0438\u044f Data Scientist<\/a><\/p>\n<\/li>\n<li>\n<p><a href=\"https:\/\/skillfactory.ru\/data-analyst-pro?utm_source=habr&amp;utm_medium=habr&amp;utm_campaign=article&amp;utm_content=analytics_dapr_230921&amp;utm_term=cat\">\u041f\u0440\u043e\u0444\u0435\u0441\u0441\u0438\u044f Data Analyst<\/a><\/p>\n<\/li>\n<li>\n<p><a href=\"https:\/\/skillfactory.ru\/matematika-dlya-data-science#syllabus?utm_source=habr&amp;utm_medium=habr&amp;utm_campaign=article&amp;utm_content=data-science_mat_230921&amp;utm_term=cat\">\u041a\u0443\u0440\u0441 \u00ab\u041c\u0430\u0442\u0435\u043c\u0430\u0442\u0438\u043a\u0430 \u0434\u043b\u044f Data Science\u00bb<\/a><\/p>\n<\/li>\n<li>\n<p><a href=\"https:\/\/skillfactory.ru\/matematika-i-machine-learning-dlya-data-science?utm_source=habr&amp;utm_medium=habr&amp;utm_campaign=article&amp;utm_content=data-science_matml_230921&amp;utm_term=cat\">\u041a\u0443\u0440\u0441 \u00ab\u041c\u0430\u0442\u0435\u043c\u0430\u0442\u0438\u043a\u0430 \u0438 Machine Learning \u0434\u043b\u044f Data Science\u00bb<\/a><\/p>\n<\/li>\n<li>\n<p><a href=\"https:\/\/skillfactory.ru\/data-engineer?utm_source=habr&amp;utm_medium=habr&amp;utm_campaign=article&amp;utm_content=data-science_dea_230921&amp;utm_term=cat\">\u041a\u0443\u0440\u0441 \u043f\u043e Data Engineering<\/a><\/p>\n<\/li>\n<li>\n<p><a href=\"https:\/\/skillfactory.ru\/machine-learning-i-deep-learning?utm_source=habr&amp;utm_medium=habr&amp;utm_campaign=article&amp;utm_content=data-science_mldl_230921&amp;utm_term=cat\">\u041a\u0443\u0440\u0441 \u00abMachine Learning \u0438 Deep Learning\u00bb<\/a><\/p>\n<\/li>\n<li>\n<p><a href=\"https:\/\/skillfactory.ru\/machine-learning?utm_source=habr&amp;utm_medium=habr&amp;utm_campaign=article&amp;utm_content=data-science_ml_230921&amp;utm_term=cat\">\u041a\u0443\u0440\u0441 \u043f\u043e Machine Learning<\/a><\/p>\n<\/li>\n<\/ul>\n<p><strong>Python, \u0432\u0435\u0431-\u0440\u0430\u0437\u0440\u0430\u0431\u043e\u0442\u043a\u0430<\/strong><\/p>\n<ul>\n<li>\n<p><a href=\"https:\/\/skillfactory.ru\/python-fullstack-web-developer?utm_source=habr&amp;utm_medium=habr&amp;utm_campaign=article&amp;utm_content=coding_fpw_230921&amp;utm_term=cat\">\u041f\u0440\u043e\u0444\u0435\u0441\u0441\u0438\u044f Fullstack-\u0440\u0430\u0437\u0440\u0430\u0431\u043e\u0442\u0447\u0438\u043a \u043d\u0430 Python<\/a><\/p>\n<\/li>\n<li>\n<p><a href=\"https:\/\/skillfactory.ru\/python-for-web-developers?utm_source=habr&amp;utm_medium=habr&amp;utm_campaign=article&amp;utm_content=coding_pws_230921&amp;utm_term=cat\">\u041a\u0443\u0440\u0441 \u00abPython \u0434\u043b\u044f \u0432\u0435\u0431-\u0440\u0430\u0437\u0440\u0430\u0431\u043e\u0442\u043a\u0438\u00bb<\/a><\/p>\n<\/li>\n<li>\n<p><a href=\"https:\/\/skillfactory.ru\/frontend-razrabotchik?utm_source=habr&amp;utm_medium=habr&amp;utm_campaign=article&amp;utm_content=coding_fr_230921&amp;utm_term=cat\">\u041f\u0440\u043e\u0444\u0435\u0441\u0441\u0438\u044f Frontend-\u0440\u0430\u0437\u0440\u0430\u0431\u043e\u0442\u0447\u0438\u043a<\/a><\/p>\n<\/li>\n<li>\n<p><a href=\"https:\/\/skillfactory.ru\/webdev?utm_source=habr&amp;utm_medium=habr&amp;utm_campaign=article&amp;utm_content=coding_webdev_230921&amp;utm_term=cat\">\u041f\u0440\u043e\u0444\u0435\u0441\u0441\u0438\u044f \u0412\u0435\u0431-\u0440\u0430\u0437\u0440\u0430\u0431\u043e\u0442\u0447\u0438\u043a<\/a><\/p>\n<\/li>\n<\/ul>\n<p><strong>\u041c\u043e\u0431\u0438\u043b\u044c\u043d\u0430\u044f \u0440\u0430\u0437\u0440\u0430\u0431\u043e\u0442\u043a\u0430<\/strong><\/p>\n<ul>\n<li>\n<p><a href=\"https:\/\/skillfactory.ru\/ios-razrabotchik-s-nulya?utm_source=habr&amp;utm_medium=habr&amp;utm_campaign=article&amp;utm_content=coding_ios_230921&amp;utm_term=cat\">\u041f\u0440\u043e\u0444\u0435\u0441\u0441\u0438\u044f iOS-\u0440\u0430\u0437\u0440\u0430\u0431\u043e\u0442\u0447\u0438\u043a<\/a><\/p>\n<\/li>\n<li>\n<p><a href=\"https:\/\/skillfactory.ru\/android-razrabotchik?utm_source=habr&amp;utm_medium=habr&amp;utm_campaign=article&amp;utm_content=coding_andr_230921&amp;utm_term=cat\">\u041f\u0440\u043e\u0444\u0435\u0441\u0441\u0438\u044f Android-\u0440\u0430\u0437\u0440\u0430\u0431\u043e\u0442\u0447\u0438\u043a<\/a><\/p>\n<\/li>\n<\/ul>\n<p><strong>Java \u0438 C#<\/strong><\/p>\n<ul>\n<li>\n<p><a href=\"https:\/\/skillfactory.ru\/java-razrabotchik?utm_source=habr&amp;utm_medium=habr&amp;utm_campaign=article&amp;utm_content=coding_java_230921&amp;utm_term=cat\">\u041f\u0440\u043e\u0444\u0435\u0441\u0441\u0438\u044f Java-\u0440\u0430\u0437\u0440\u0430\u0431\u043e\u0442\u0447\u0438\u043a<\/a><\/p>\n<\/li>\n<li>\n<p><a href=\"https:\/\/skillfactory.ru\/java-qa-engineer-testirovshik-po?utm_source=habr&amp;utm_medium=habr&amp;utm_campaign=article&amp;utm_content=coding_qaja_230921&amp;utm_term=cat\">\u041f\u0440\u043e\u0444\u0435\u0441\u0441\u0438\u044f QA-\u0438\u043d\u0436\u0435\u043d\u0435\u0440 \u043d\u0430 JAVA<\/a><\/p>\n<\/li>\n<li>\n<p><a href=\"https:\/\/skillfactory.ru\/c-sharp-razrabotchik?utm_source=habr&amp;utm_medium=habr&amp;utm_campaign=article&amp;utm_content=coding_cdev_230921&amp;utm_term=cat\">\u041f\u0440\u043e\u0444\u0435\u0441\u0441\u0438\u044f C#-\u0440\u0430\u0437\u0440\u0430\u0431\u043e\u0442\u0447\u0438\u043a<\/a><\/p>\n<\/li>\n<li>\n<p><a href=\"https:\/\/skillfactory.ru\/game-razrabotchik-na-unity-i-c-sharp?utm_source=habr&amp;utm_medium=habr&amp;utm_campaign=article&amp;utm_content=coding_gamedev_230921&amp;utm_term=cat\">\u041f\u0440\u043e\u0444\u0435\u0441\u0441\u0438\u044f \u0420\u0430\u0437\u0440\u0430\u0431\u043e\u0442\u0447\u0438\u043a \u0438\u0433\u0440 \u043d\u0430 Unity<\/a><\/p>\n<\/li>\n<\/ul>\n<p><strong>\u041e\u0442 \u043e\u0441\u043d\u043e\u0432 \u2014 \u0432 \u0433\u043b\u0443\u0431\u0438\u043d\u0443<\/strong><\/p>\n<ul>\n<li>\n<p><a href=\"https:\/\/skillfactory.ru\/algoritmy-i-struktury-dannyh?utm_source=habr&amp;utm_medium=habr&amp;utm_campaign=article&amp;utm_content=coding_algo_230921&amp;utm_term=cat\">\u041a\u0443\u0440\u0441 \u00ab\u0410\u043b\u0433\u043e\u0440\u0438\u0442\u043c\u044b \u0438 \u0441\u0442\u0440\u0443\u043a\u0442\u0443\u0440\u044b \u0434\u0430\u043d\u043d\u044b\u0445\u00bb<\/a><\/p>\n<\/li>\n<li>\n<p><a href=\"https:\/\/skillfactory.ru\/c-plus-plus-razrabotchik?utm_source=habr&amp;utm_medium=habr&amp;utm_campaign=article&amp;utm_content=coding_cplus_230921&amp;utm_term=cat\">\u041f\u0440\u043e\u0444\u0435\u0441\u0441\u0438\u044f C++ \u0440\u0430\u0437\u0440\u0430\u0431\u043e\u0442\u0447\u0438\u043a<\/a><\/p>\n<\/li>\n<li>\n<p><a href=\"https:\/\/skillfactory.ru\/cyber-security-etichnij-haker?utm_source=habr&amp;utm_medium=habr&amp;utm_campaign=article&amp;utm_content=coding_hacker_230921&amp;utm_term=cat\">\u041f\u0440\u043e\u0444\u0435\u0441\u0441\u0438\u044f \u042d\u0442\u0438\u0447\u043d\u044b\u0439 \u0445\u0430\u043a\u0435\u0440<\/a><\/p>\n<\/li>\n<\/ul>\n<p><strong>\u0410 \u0442\u0430\u043a\u0436\u0435:<\/strong><\/p>\n<ul>\n<li>\n<p><a href=\"https:\/\/skillfactory.ru\/devops-ingineer?utm_source=habr&amp;utm_medium=habr&amp;utm_campaign=article&amp;utm_content=coding_devops_230921&amp;utm_term=cat\">\u041a\u0443\u0440\u0441 \u043f\u043e DevOps<\/a><\/p>\n<\/li>\n<\/ul>\n<\/div>\n<\/details>\n<\/div>\n<\/div>\n<\/div>\n<p><!----><!----><\/div>\n<p><!----><!----><br \/> \u0441\u0441\u044b\u043b\u043a\u0430 \u043d\u0430 \u043e\u0440\u0438\u0433\u0438\u043d\u0430\u043b \u0441\u0442\u0430\u0442\u044c\u0438 <a href=\"https:\/\/habr.com\/ru\/articles\/579380\/\"> https:\/\/habr.com\/ru\/articles\/579380\/<\/a><\/p>\n","protected":false},"excerpt":{"rendered":"<div><!--[--><!--]--><\/div>\n<div id=\"post-content-body\">\n<div>\n<div class=\"article-formatted-body article-formatted-body article-formatted-body_version-2\">\n<div xmlns=\"http:\/\/www.w3.org\/1999\/xhtml\">\n<figure class=\"full-width\"><figcaption><\/figcaption><\/figure>\n<p>\u041f\u043e\u0434\u043e\u0431\u043d\u044b\u0439 Python Triton \u0443\u0436\u0435 \u0440\u0430\u0431\u043e\u0442\u0430\u0435\u0442 \u0432 \u044f\u0434\u0440\u0430\u0445, \u043a\u043e\u0442\u043e\u0440\u044b\u0435 \u0432 2 \u0440\u0430\u0437\u0430 \u044d\u0444\u0444\u0435\u043a\u0442\u0438\u0432\u043d\u0435\u0435 \u044d\u043a\u0432\u0438\u0432\u0430\u043b\u0435\u043d\u0442\u043d\u044b\u0445 \u0440\u0435\u0430\u043b\u0438\u0437\u0430\u0446\u0438\u0439 Torch. \u0410 \u044f\u0434\u0440\u0430 \u043c\u0430\u0442\u0440\u0438\u0447\u043d\u043e\u0433\u043e \u0443\u043c\u043d\u043e\u0436\u0435\u043d\u0438\u044f FP16, \u0441\u0440\u0430\u0432\u043d\u0438\u043c\u044b\u0435 \u043f\u0440\u043e\u0438\u0437\u0432\u043e\u0434\u0438\u0442\u0435\u043b\u044c\u043d\u043e\u0441\u0442\u044c\u044e \u0441 cuBLAS, \u043d\u0430 Triton \u0437\u0430\u0439\u043c\u0443\u0442 \u043c\u0435\u043d\u0435\u0435 25 \u0441\u0442\u0440\u043e\u043a. \u041a\u0430\u043a \u0443\u0442\u0432\u0435\u0440\u0436\u0434\u0430\u0435\u0442 \u0430\u0432\u0442\u043e\u0440, \u043c\u043d\u043e\u0433\u0438\u0435 \u043f\u0440\u043e\u0433\u0440\u0430\u043c\u043c\u0438\u0441\u0442\u044b \u043d\u0435 \u043c\u043e\u0433\u0443\u0442 \u043d\u0430\u043f\u0438\u0441\u0430\u0442\u044c \u0442\u0430\u043a\u0438\u0435 \u044f\u0434\u0440\u0430. \u041f\u043e\u0434\u0440\u043e\u0431\u043d\u043e\u0441\u0442\u044f\u043c\u0438 \u043e Triton \u0434\u0435\u043b\u0438\u043c\u0441\u044f \u043a \u0441\u0442\u0430\u0440\u0442\u0443 <a href=\"https:\/\/skillfactory.ru\/machine-learning-i-deep-learning?utm_source=habr&amp;utm_medium=habr&amp;utm_campaign=article&amp;utm_content=data-science_mldl_230921&amp;utm_term=lead\">\u043a\u0443\u0440\u0441\u0430 \u043f\u043e ML \u0438 DL<\/a>.<\/p>\n<hr\/>\n<p>\u041d\u043e\u0432\u044b\u0435 \u0438\u0441\u0441\u043b\u0435\u0434\u043e\u0432\u0430\u0442\u0435\u043b\u044c\u0441\u043a\u0438\u0435 \u0438\u0434\u0435\u0438 \u0432 \u043e\u0431\u043b\u0430\u0441\u0442\u0438 Deep Learning \u043e\u0431\u044b\u0447\u043d\u043e \u0440\u0435\u0430\u043b\u0438\u0437\u0443\u044e\u0442\u0441\u044f \u0441 \u043f\u043e\u043c\u043e\u0449\u044c\u044e \u043a\u043e\u043c\u0431\u0438\u043d\u0430\u0446\u0438\u0438 \u043e\u043f\u0435\u0440\u0430\u0442\u043e\u0440\u043e\u0432 \u043d\u0430\u0442\u0438\u0432\u043d\u043e\u0433\u043e \u0444\u0440\u0435\u0439\u043c\u0432\u043e\u0440\u043a\u0430. \u042d\u0442\u043e \u0443\u0434\u043e\u0431\u043d\u043e, \u043d\u043e \u0447\u0430\u0441\u0442\u043e \u0442\u0440\u0435\u0431\u0443\u0435\u0442 \u0441\u043e\u0437\u0434\u0430\u043d\u0438\u044f (\u0438\/\u0438\u043b\u0438 \u043f\u0435\u0440\u0435\u043c\u0435\u0449\u0435\u043d\u0438\u044f) \u043c\u043d\u043e\u0436\u0435\u0441\u0442\u0432\u0430 \u0432\u0440\u0435\u043c\u0435\u043d\u043d\u044b\u0445 \u0442\u0435\u043d\u0437\u043e\u0440\u043e\u0432, \u0447\u0442\u043e \u043c\u043e\u0436\u0435\u0442 \u0441\u043d\u0438\u0437\u0438\u0442\u044c \u043f\u0440\u043e\u0438\u0437\u0432\u043e\u0434\u0438\u0442\u0435\u043b\u044c\u043d\u043e\u0441\u0442\u044c \u043d\u0435\u0439\u0440\u043e\u0441\u0435\u0442\u0435\u0439 \u0432 \u0431\u043e\u043b\u044c\u0448\u043e\u043c \u043c\u0430\u0441\u0448\u0442\u0430\u0431\u0435. \u041f\u0440\u043e\u0431\u043b\u0435\u043c\u044b \u0440\u0435\u0448\u0430\u044e\u0442\u0441\u044f \u0441\u043f\u0435\u0446\u0438\u0430\u043b\u0438\u0437\u0438\u0440\u043e\u0432\u0430\u043d\u043d\u044b\u043c\u0438 \u044f\u0434\u0440\u0430\u043c\u0438 \u0434\u043b\u044f GPU, \u043d\u043e \u043d\u0430\u043f\u0438\u0441\u0430\u0442\u044c \u0438\u0445 \u043c\u043e\u0436\u0435\u0442 \u0431\u044b\u0442\u044c \u0443\u0434\u0438\u0432\u0438\u0442\u0435\u043b\u044c\u043d\u043e \u0441\u043b\u043e\u0436\u043d\u043e \u0438\u0437-\u0437\u0430 \u0442\u043e\u043d\u043a\u043e\u0441\u0442\u0435\u0439 \u043f\u0440\u043e\u0433\u0440\u0430\u043c\u043c\u0438\u0440\u043e\u0432\u0430\u043d\u0438\u044f \u043d\u0430 GPU <a href=\"https:\/\/openai.com\/blog\/triton\/#rf1\"><sup>1<\/sup><\/a><sup>, <\/sup><a href=\"https:\/\/openai.com\/blog\/triton\/#rf2\"><sup>2<\/sup><\/a><sup>, <\/sup><a href=\"https:\/\/openai.com\/blog\/triton\/#rf3\"><sup>3<\/sup><\/a>.<\/p>\n<p>\u0418, \u0445\u043e\u0442\u044f \u0432 \u043f\u043e\u0441\u043b\u0435\u0434\u043d\u0435\u0435 \u0432\u0440\u0435\u043c\u044f \u043f\u043e\u044f\u0432\u0438\u043b\u043e\u0441\u044c \u043c\u043d\u043e\u0436\u0435\u0441\u0442\u0432\u043e \u043e\u0431\u043b\u0435\u0433\u0447\u0430\u044e\u0449\u0438\u0445 \u0442\u0440\u0443\u0434 \u0441\u0438\u0441\u0442\u0435\u043c<a href=\"https:\/\/openai.com\/blog\/triton\/#rf4\"><sup>4<\/sup><\/a><sup>, <\/sup><a href=\"https:\/\/openai.com\/blog\/triton\/#rf5\"><sup>5<\/sup><\/a>, \u043c\u044b \u043e\u0431\u043d\u0430\u0440\u0443\u0436\u0438\u043b\u0438, \u0447\u0442\u043e \u043e\u043d\u0438 \u043b\u0438\u0431\u043e \u0441\u043b\u0438\u0448\u043a\u043e\u043c \u043c\u043d\u043e\u0433\u043e\u0441\u043b\u043e\u0432\u043d\u044b, \u043b\u0438\u0431\u043e \u043d\u0435\u0434\u043e\u0441\u0442\u0430\u0442\u043e\u0447\u043d\u043e \u0433\u0438\u0431\u043a\u0438\u0435, \u043b\u0438\u0431\u043e \u0433\u0435\u043d\u0435\u0440\u0438\u0440\u0443\u044e\u0442 \u043a\u043e\u0434 \u0437\u0430\u043c\u0435\u0442\u043d\u043e \u043c\u0435\u0434\u043b\u0435\u043d\u043d\u0435\u0435, \u0447\u0435\u043c \u043d\u0430\u0441\u0442\u0440\u043e\u0435\u043d\u043d\u044b\u0439 \u0432\u0440\u0443\u0447\u043d\u0443\u044e \u0431\u0430\u0437\u043e\u0432\u044b\u0439. \u042d\u0442\u043e \u043f\u0440\u0438\u0432\u0435\u043b\u043e \u043a \u0440\u0430\u0441\u0448\u0438\u0440\u0435\u043d\u0438\u044e \u0438 \u0443\u043b\u0443\u0447\u0448\u0435\u043d\u0438\u044e Triton<a href=\"https:\/\/openai.com\/blog\/triton\/#rf6\"><sup>6<\/sup><\/a>, \u2014 \u043d\u043e\u0432\u043e\u0433\u043e \u044f\u0437\u044b\u043a\u0430 \u0438 \u043a\u043e\u043c\u043f\u0438\u043b\u044f\u0442\u043e\u0440\u0430, \u0441\u043e\u0437\u0434\u0430\u0442\u0435\u043b\u044c \u043a\u043e\u0442\u043e\u0440\u043e\u0433\u043e \u0441\u0435\u0439\u0447\u0430\u0441 \u0440\u0430\u0431\u043e\u0442\u0430\u0435\u0442 \u0432 OpenAI.<\/p>\n<h2>\u041f\u0440\u043e\u0431\u043b\u0435\u043c\u044b \u043f\u0440\u043e\u0433\u0440\u0430\u043c\u043c\u0438\u0440\u043e\u0432\u0430\u043d\u0438\u044f GPU<\/h2>\n<p>\u0410\u0440\u0445\u0438\u0442\u0435\u043a\u0442\u0443\u0440\u0430 \u0441\u043e\u0432\u0440\u0435\u043c\u0435\u043d\u043d\u044b\u0445 GPU \u0443\u0441\u043b\u043e\u0432\u043d\u043e \u0434\u0435\u043b\u0438\u0442\u0441\u044f \u043d\u0430 \u0442\u0440\u0438 \u043e\u0441\u043d\u043e\u0432\u043d\u044b\u0445 \u043a\u043e\u043c\u043f\u043e\u043d\u0435\u043d\u0442\u0430 \u2014 DRAM, SRAM \u0438 ALU. \u041e\u043f\u0442\u0438\u043c\u0438\u0437\u0438\u0440\u0443\u044f \u043a\u043e\u0434 \u043d\u0430 CUDA, \u043d\u0443\u0436\u043d\u043e \u0443\u0447\u0438\u0442\u044b\u0432\u0430\u0442\u044c \u043a\u0430\u0436\u0434\u044b\u0439.<\/p>\n<ul>\n<li>\n<p>\u041f\u0435\u0440\u0435\u0434\u0430\u0447\u0438 \u043f\u0430\u043c\u044f\u0442\u0438 \u0438\u0437 DRAM \u0434\u043e\u043b\u0436\u043d\u044b \u0431\u044b\u0442\u044c \u043e\u0431\u044a\u0435\u0434\u0438\u043d\u0435\u043d\u044b \u0432 \u0431\u043e\u043b\u044c\u0448\u0438\u0435 \u0442\u0440\u0430\u043d\u0437\u0430\u043a\u0446\u0438\u0438, \u0447\u0442\u043e\u0431\u044b \u0438\u0441\u043f\u043e\u043b\u044c\u0437\u043e\u0432\u0430\u0442\u044c \u0431\u043e\u043b\u044c\u0448\u0443\u044e \u0448\u0438\u0440\u0438\u043d\u0443 \u0448\u0438\u043d\u044b \u0441\u043e\u0432\u0440\u0435\u043c\u0435\u043d\u043d\u044b\u0445 \u0438\u043d\u0442\u0435\u0440\u0444\u0435\u0439\u0441\u043e\u0432 \u043f\u0430\u043c\u044f\u0442\u0438.<\/p>\n<\/li>\n<li>\n<p>\u0414\u0430\u043d\u043d\u044b\u0435 \u0434\u043e\u043b\u0436\u043d\u044b \u0431\u044b\u0442\u044c \u0432\u0440\u0443\u0447\u043d\u0443\u044e \u0441\u043e\u0445\u0440\u0430\u043d\u0435\u043d\u044b \u0432 SRAM \u043f\u0435\u0440\u0435\u0434 \u043f\u043e\u0432\u0442\u043e\u0440\u043d\u044b\u043c \u0438\u0441\u043f\u043e\u043b\u044c\u0437\u043e\u0432\u0430\u043d\u0438\u0435\u043c \u0438 \u0443\u043f\u0440\u0430\u0432\u043b\u044f\u0442\u044c\u0441\u044f \u0442\u0430\u043a\u0438\u043c \u043e\u0431\u0440\u0430\u0437\u043e\u043c, \u0447\u0442\u043e\u0431\u044b \u043c\u0438\u043d\u0438\u043c\u0438\u0437\u0438\u0440\u043e\u0432\u0430\u0442\u044c \u043a\u043e\u043d\u0444\u043b\u0438\u043a\u0442\u044b \u043c\u0435\u0436\u0434\u0443 \u0431\u0430\u043d\u043a\u0430\u043c\u0438 \u043e\u0431\u0449\u0435\u0439 \u043f\u0430\u043c\u044f\u0442\u0438 \u043f\u0440\u0438 \u0438\u0437\u0432\u043b\u0435\u0447\u0435\u043d\u0438\u0438.<\/p>\n<\/li>\n<li>\n<p>\u0412\u044b\u0447\u0438\u0441\u043b\u0435\u043d\u0438\u044f \u0434\u043e\u043b\u0436\u043d\u044b \u0431\u044b\u0442\u044c \u0440\u0430\u0437\u0434\u0435\u043b\u0435\u043d\u044b \u0438 \u0442\u0449\u0430\u0442\u0435\u043b\u044c\u043d\u043e \u0441\u043f\u043b\u0430\u043d\u0438\u0440\u043e\u0432\u0430\u043d\u044b \u043a\u0430\u043a \u043c\u0435\u0436\u0434\u0443 \u043f\u043e\u0442\u043e\u043a\u043e\u0432\u044b\u043c\u0438 \u043c\u0443\u043b\u044c\u0442\u0438\u043f\u0440\u043e\u0446\u0435\u0441\u0441\u043e\u0440\u0430\u043c\u0438 (SM), \u0442\u0430\u043a \u0438 \u0432\u043d\u0443\u0442\u0440\u0438 \u043d\u0438\u0445, \u0447\u0442\u043e\u0431\u044b \u0441\u043f\u043e\u0441\u043e\u0431\u0441\u0442\u0432\u043e\u0432\u0430\u0442\u044c \u043f\u0430\u0440\u0430\u043b\u043b\u0435\u043b\u0438\u0437\u043c\u0443 \u043d\u0430 \u0443\u0440\u043e\u0432\u043d\u0435 \u0438\u043d\u0441\u0442\u0440\u0443\u043a\u0446\u0438\u0439\/\u043f\u043e\u0442\u043e\u043a\u043e\u0432 \u0438 \u0438\u0441\u043f\u043e\u043b\u044c\u0437\u043e\u0432\u0430\u0442\u044c \u0441\u043f\u0435\u0446\u0438\u0430\u043b\u0438\u0437\u0438\u0440\u043e\u0432\u0430\u043d\u043d\u044b\u0435 ALU (\u043d\u0430\u043f\u0440\u0438\u043c\u0435\u0440, \u0442\u0435\u043d\u0437\u043e\u0440\u043d\u044b\u0435 \u044f\u0434\u0440\u0430).<\/p>\n<\/li>\n<\/ul>\n<figure class=\"full-width\"><figcaption>\u0411\u0430\u0437\u043e\u0432\u0430\u044f \u0430\u0440\u0445\u0438\u0442\u0435\u043a\u0442\u0443\u0440\u0430 GPU.<\/figcaption><\/figure>\n<p>\u0423\u0447\u0438\u0442\u044b\u0432\u0430\u0442\u044c \u044d\u0442\u0438 \u0444\u0430\u043a\u0442\u043e\u0440\u044b \u043c\u043e\u0436\u0435\u0442 \u0431\u044b\u0442\u044c \u0441\u043b\u043e\u0436\u043d\u043e \u0434\u0430\u0436\u0435 \u043e\u043f\u044b\u0442\u043d\u044b\u043c \u043f\u0440\u043e\u0433\u0440\u0430\u043c\u043c\u0438\u0441\u0442\u0430\u043c CUDA. \u0426\u0435\u043b\u044c Triton \u2014 \u043f\u043e\u043b\u043d\u043e\u0441\u0442\u044c\u044e \u0430\u0432\u0442\u043e\u043c\u0430\u0442\u0438\u0437\u0438\u0440\u043e\u0432\u0430\u0442\u044c \u044d\u0442\u0438 \u043e\u043f\u0442\u0438\u043c\u0438\u0437\u0430\u0446\u0438\u0438, \u0447\u0442\u043e\u0431\u044b \u0440\u0430\u0437\u0440\u0430\u0431\u043e\u0442\u0447\u0438\u043a\u0438 \u0441\u043e\u0441\u0440\u0435\u0434\u043e\u0442\u043e\u0447\u0438\u043b\u0438\u0441\u044c \u043d\u0430 \u0432\u044b\u0441\u043e\u043a\u043e\u0443\u0440\u043e\u0432\u043d\u0435\u0432\u043e\u0439 \u043b\u043e\u0433\u0438\u043a\u0435 \u043f\u0430\u0440\u0430\u043b\u043b\u0435\u043b\u044c\u043d\u043e\u0433\u043e \u043a\u043e\u0434\u0430.\u00a0<\/p>\n<p>Triton \u0441\u0442\u0440\u0435\u043c\u0438\u0442\u0441\u044f \u043a \u0448\u0438\u0440\u043e\u043a\u043e\u043c\u0443 \u043f\u0440\u0438\u043c\u0435\u043d\u0435\u043d\u0438\u044e, \u0438 \u043f\u043e\u044d\u0442\u043e\u043c\u0443 \u043d\u0435 \u0441\u043e\u0441\u0442\u0430\u0432\u043b\u044f\u0435\u0442 \u0430\u0432\u0442\u043e\u043c\u0430\u0442\u0438\u0447\u0435\u0441\u043a\u043e\u0433\u043e \u0440\u0430\u0441\u043f\u0438\u0441\u0430\u043d\u0438\u044f \u0440\u0430\u0431\u043e\u0442\u044b \u043c\u0435\u0436\u0434\u0443 SM, \u043e\u0441\u0442\u0430\u0432\u043b\u044f\u044f \u043d\u0435\u043a\u043e\u0442\u043e\u0440\u044b\u0435 \u0432\u0430\u0436\u043d\u044b\u0435 \u0430\u043b\u0433\u043e\u0440\u0438\u0442\u043c\u0438\u0447\u0435\u0441\u043a\u0438\u0435 \u0441\u043e\u043e\u0431\u0440\u0430\u0436\u0435\u043d\u0438\u044f (\u043d\u0430\u043f\u0440\u0438\u043c\u0435\u0440, \u0442\u0430\u0439\u043b\u0438\u043d\u0433, \u0441\u0438\u043d\u0445\u0440\u043e\u043d\u0438\u0437\u0430\u0446\u0438\u044f \u043c\u0435\u0436\u0434\u0443 SM) \u043d\u0430 \u0443\u0441\u043c\u043e\u0442\u0440\u0435\u043d\u0438\u0435 \u0440\u0430\u0437\u0440\u0430\u0431\u043e\u0442\u0447\u0438\u043a\u043e\u0432.<\/p>\n<p>\u041e\u043f\u0442\u0438\u043c\u0438\u0437\u0430\u0446\u0438\u044f \u043a\u043e\u043c\u043f\u0438\u043b\u044f\u0442\u043e\u0440\u0430 \u0432 CUDA \u043f\u043e \u0441\u0440\u0430\u0432\u043d\u0435\u043d\u0438\u044e \u0441 Triton.<\/p>\n<div>\n<div class=\"table\">\n<table>\n<tbody>\n<tr>\n<th>\n<p><strong><br \/><\/strong><\/p>\n<\/th>\n<th>\n<p><strong>CUDA<\/strong><\/p>\n<\/th>\n<td>\n<p><strong>TRITON<\/strong><\/p>\n<\/td>\n<\/tr>\n<tr>\n<td>\n<p>\u041e\u0431\u044a\u0435\u0434\u0438\u043d\u0435\u043d\u0438\u0435 \u043f\u0430\u043c\u044f\u0442\u0438<\/p>\n<\/td>\n<td>\n<p>\u0412\u0440\u0443\u0447\u043d\u0443\u044e<\/p>\n<\/td>\n<td>\n<p>\u0410\u0432\u0442\u043e\u043c\u0430\u0442\u0438\u0447\u0435\u0441\u043a\u0438<\/p>\n<\/td>\n<\/tr>\n<tr>\n<td>\n<p>\u0423\u043f\u0440\u0430\u0432\u043b\u0435\u043d\u0438\u0435 \u043e\u0431\u0449\u0435\u0439 \u043f\u0430\u043c\u044f\u0442\u044c\u044e<\/p>\n<\/td>\n<td>\n<p>\u0412\u0440\u0443\u0447\u043d\u0443\u044e<\/p>\n<\/td>\n<td>\n<p>\u0410\u0432\u0442\u043e\u043c\u0430\u0442\u0438\u0447\u0435\u0441\u043a\u0438<\/p>\n<\/td>\n<\/tr>\n<tr>\n<td>\n<p>\u041f\u043b\u0430\u043d\u0438\u0440\u043e\u0432\u0430\u043d\u0438\u0435 (\u0432\u043d\u0443\u0442\u0440\u0438 SM)<\/p>\n<\/td>\n<td>\n<p>\u0412\u0440\u0443\u0447\u043d\u0443\u044e<\/p>\n<\/td>\n<td>\n<p>\u0410\u0432\u0442\u043e\u043c\u0430\u0442\u0438\u0447\u0435\u0441\u043a\u0438<\/p>\n<\/td>\n<\/tr>\n<tr>\n<td>\n<p>\u041f\u043b\u0430\u043d\u0438\u0440\u043e\u0432\u0430\u043d\u0438\u0435 (\u043f\u043e \u0432\u0441\u0435\u043c SM)<\/p>\n<\/td>\n<td>\n<p>\u0412\u0440\u0443\u0447\u043d\u0443\u044e<\/p>\n<\/td>\n<td>\n<p>\u0412\u0440\u0443\u0447\u043d\u0443\u044e<\/p>\n<\/td>\n<\/tr>\n<\/tbody>\n<\/table>\n<\/div>\n<\/div>\n<h3>\u041c\u043e\u0434\u0435\u043b\u044c \u043f\u0440\u043e\u0433\u0440\u0430\u043c\u043c\u0438\u0440\u043e\u0432\u0430\u043d\u0438\u044f<\/h3>\n<p>\u0418\u0437 \u0432\u0441\u0435\u0445 \u0434\u043e\u0441\u0442\u0443\u043f\u043d\u044b\u0445 \u0434\u043e\u043c\u0435\u043d\u043d\u043e-\u0441\u043f\u0435\u0446\u0438\u0444\u0438\u0447\u0435\u0441\u043a\u0438\u0445 \u044f\u0437\u044b\u043a\u043e\u0432 \u0438 JIT-\u043a\u043e\u043c\u043f\u0438\u043b\u044f\u0442\u043e\u0440\u043e\u0432 Triton, \u043f\u043e\u0436\u0430\u043b\u0443\u0439, \u043d\u0430\u0438\u0431\u043e\u043b\u0435\u0435 \u043f\u043e\u0445\u043e\u0436 \u043d\u0430 Numba: \u044f\u0434\u0440\u0430 \u043e\u043f\u0440\u0435\u0434\u0435\u043b\u044f\u044e\u0442\u0441\u044f \u043a\u0430\u043a \u043e\u0444\u043e\u0440\u043c\u043b\u0435\u043d\u043d\u044b\u0435 \u0444\u0443\u043d\u043a\u0446\u0438\u0438 Python \u0438 \u0437\u0430\u043f\u0443\u0441\u043a\u0430\u044e\u0442\u0441\u044f \u043f\u0430\u0440\u0430\u043b\u043b\u0435\u043b\u044c\u043d\u043e \u0441 \u0440\u0430\u0437\u043d\u044b\u043c\u0438 program_id \u043d\u0430 \u0441\u0435\u0442\u043a\u0435 \u0442\u0430\u043a \u043d\u0430\u0437\u044b\u0432\u0430\u0435\u043c\u044b\u0445 <em>instances<\/em> (\u0438\u043d\u0441\u0442\u0430\u043d\u0441\u043e\u0432, \u044d\u043a\u0437\u0435\u043c\u043f\u043b\u044f\u0440\u043e\u0432).\u00a0<\/p>\n<p>\u041e\u0434\u043d\u0430\u043a\u043e, \u043a\u0430\u043a \u043f\u043e\u043a\u0430\u0437\u0430\u043d\u043e \u0432 \u043f\u0440\u0438\u0432\u0435\u0434\u0451\u043d\u043d\u043e\u043c \u043d\u0438\u0436\u0435 \u0444\u0440\u0430\u0433\u043c\u0435\u043d\u0442\u0435 \u043a\u043e\u0434\u0430, \u043d\u0430 \u044d\u0442\u043e\u043c \u0441\u0445\u043e\u0434\u0441\u0442\u0432\u043e \u0437\u0430\u043a\u0430\u043d\u0447\u0438\u0432\u0430\u0435\u0442\u0441\u044f: Triton \u0440\u0430\u0441\u043a\u0440\u044b\u0432\u0430\u0435\u0442 \u0432\u043d\u0443\u0442\u0440\u0438\u044d\u043a\u0437\u0435\u043c\u043f\u043b\u044f\u0440\u043d\u044b\u0439 \u043f\u0430\u0440\u0430\u043b\u043b\u0435\u043b\u0438\u0437\u043c \u0447\u0435\u0440\u0435\u0437 \u043e\u043f\u0435\u0440\u0430\u0446\u0438\u0438 \u043d\u0430\u0434 <em>\u0431\u043b\u043e\u043a\u0430\u043c\u0438<\/em> \u2014 \u043d\u0435\u0431\u043e\u043b\u044c\u0448\u0438\u043c\u0438 \u043c\u0430\u0441\u0441\u0438\u0432\u0430\u043c\u0438, \u0440\u0430\u0437\u043c\u0435\u0440\u044b \u043a\u043e\u0442\u043e\u0440\u044b\u0445 \u0440\u0430\u0432\u043d\u044b \u0441\u0442\u0435\u043f\u0435\u043d\u0438 \u0434\u0432\u043e\u0439\u043a\u0438, \u2014 \u0432\u043c\u0435\u0441\u0442\u043e \u043c\u043e\u0434\u0435\u043b\u0438 \u0432\u044b\u043f\u043e\u043b\u043d\u0435\u043d\u0438\u044f Single Instruction, Multiple Thread (SIMT, \u043e\u0434\u043d\u0430 \u0438\u043d\u0441\u0442\u0440\u0443\u043a\u0446\u0438\u044f \u2014 \u043c\u043d\u043e\u0433\u043e \u043f\u043e\u0442\u043e\u043a\u043e\u0432)<a href=\"https:\/\/openai.com\/blog\/triton\/#rf7\"><sup>7<\/sup><\/a>.\u00a0<\/p>\n<p>\u041f\u0440\u0438 \u044d\u0442\u043e\u043c Triton \u044d\u0444\u0444\u0435\u043a\u0442\u0438\u0432\u043d\u043e \u0430\u0431\u0441\u0442\u0440\u0430\u0433\u0438\u0440\u0443\u0435\u0442\u0441\u044f \u043e\u0442 \u0432\u0441\u0435\u0445 \u043f\u0440\u043e\u0431\u043b\u0435\u043c, \u0441\u0432\u044f\u0437\u0430\u043d\u043d\u044b\u0445 \u0441 \u043f\u0430\u0440\u0430\u043b\u043b\u0435\u043b\u0438\u0437\u043c\u043e\u043c <em>\u0432\u043d\u0443\u0442\u0440\u0438<\/em> \u0431\u043b\u043e\u043a\u043e\u0432 \u043f\u043e\u0442\u043e\u043a\u043e\u0432 CUDA (\u043d\u0430\u043f\u0440\u0438\u043c\u0435\u0440, \u043e\u0431\u044a\u0435\u0434\u0438\u043d\u0435\u043d\u0438\u0435 \u043f\u0430\u043c\u044f\u0442\u0438, \u0441\u0438\u043d\u0445\u0440\u043e\u043d\u0438\u0437\u0430\u0446\u0438\u044f\/\u043a\u043e\u043d\u0444\u043b\u0438\u043a\u0442\u044b \u043e\u0431\u0449\u0435\u0439 \u043f\u0430\u043c\u044f\u0442\u0438, \u043f\u043b\u0430\u043d\u0438\u0440\u043e\u0432\u0430\u043d\u0438\u0435 \u0440\u0430\u0431\u043e\u0442\u044b \u0442\u0435\u043d\u0437\u043e\u0440\u043d\u044b\u0445 \u044f\u0434\u0435\u0440).<\/p>\n<p>\u0414\u043e\u0431\u0430\u0432\u043b\u0435\u043d\u0438\u0435 \u0432\u0435\u043a\u0442\u043e\u0440\u0430:<\/p>\n<pre><code class=\"python\">BLOCK = 512  # This is a GPU kernel in Numba. # Different instances of this # function may run in parallel. @jit def add(X, Y, Z, N):    # In Numba\/CUDA, each kernel     # instance itself uses an SIMT execution    # model, where instructions are executed in    # parallel for different values of threadIdx    tid = threadIdx.x    bid = blockIdx.x    # scalar index    idx = bid * BLOCK + tid    if id &lt; N:      # There is no pointer in Numba.      # Z,X,Y are dense tensors      Z[idx] = X[idx] + Y[idx]   ... grid = (ceil_div(N, BLOCK),) block = (BLOCK,) add[grid, block](x, y, z, x.shape[0])<\/code><\/pre>\n<pre><code class=\"python\">BLOCK = 512  # This is a GPU kernel in Triton. # Different instances of this # function may run in parallel. @jit def add(X, Y, Z, N):    # In Triton, each kernel instance    # executes block operations on a    # single thread: there is no construct    # analogous to threadIdx    pid = program_id(0)    # block of indices    idx = pid * BLOCK + arange(BLOCK)    mask = idx &lt; N    # Triton uses pointer arithmetics      # rather than indexing operators    x = load(X + idx, mask=mask)    y = load(Y + idx, mask=mask)    store(Z + idx, x + y, mask=mask)   ... grid = (ceil_div(N, BLOCK),) # no thread-block add[grid](x, y, z, x.shape[0])<\/code><\/pre>\n<p>\u0425\u043e\u0442\u044f \u044d\u0442\u043e \u043c\u043e\u0436\u0435\u0442 \u0431\u044b\u0442\u044c \u043d\u0435 \u043e\u0441\u043e\u0431\u0435\u043d\u043d\u043e \u043f\u043e\u043b\u0435\u0437\u043d\u043e \u0434\u043b\u044f \u043f\u043e\u044d\u043b\u0435\u043c\u0435\u043d\u0442\u043d\u044b\u0445 \u0432\u044b\u0447\u0438\u0441\u043b\u0435\u043d\u0438\u0439, \u043f\u043e\u0434\u0445\u043e\u0434 \u043c\u043e\u0436\u0435\u0442 \u0437\u043d\u0430\u0447\u0438\u0442\u0435\u043b\u044c\u043d\u043e \u0443\u043f\u0440\u043e\u0441\u0442\u0438\u0442\u044c \u0440\u0430\u0437\u0440\u0430\u0431\u043e\u0442\u043a\u0443 \u0431\u043e\u043b\u0435\u0435 \u0441\u043b\u043e\u0436\u043d\u044b\u0445 \u043f\u0440\u043e\u0433\u0440\u0430\u043c\u043c \u043d\u0430 GPU. \u0420\u0430\u0441\u0441\u043c\u043e\u0442\u0440\u0438\u043c, \u043d\u0430\u043f\u0440\u0438\u043c\u0435\u0440, \u0441\u043b\u0443\u0447\u0430\u0439 \u044f\u0434\u0440\u0430 fused softmax, \u0433\u0434\u0435 \u043a\u0430\u0436\u0434\u044b\u0439 \u044d\u043a\u0437\u0435\u043c\u043f\u043b\u044f\u0440 \u043d\u043e\u0440\u043c\u0438\u0440\u0443\u0435\u0442 \u043e\u0442\u0434\u0435\u043b\u044c\u043d\u0443\u044e \u0441\u0442\u0440\u043e\u043a\u0443 \u0437\u0430\u0434\u0430\u043d\u043d\u043e\u0433\u043e \u0432\u0445\u043e\u0434\u043d\u043e\u0433\u043e \u0442\u0435\u043d\u0437\u043e\u0440\u0430 <em>X<\/em>\u2208R<em><sup>M<\/sup><\/em><sup>\u00d7<\/sup><em><sup>N<\/sup><\/em>.\u00a0<\/p>\n<p>\u0421\u0442\u0430\u043d\u0434\u0430\u0440\u0442\u043d\u044b\u0435 CUDA-\u0440\u0435\u0430\u043b\u0438\u0437\u0430\u0446\u0438\u0438 \u044d\u0442\u043e\u0439 \u0441\u0442\u0440\u0430\u0442\u0435\u0433\u0438\u0438 \u0440\u0430\u0441\u043f\u0430\u0440\u0430\u043b\u043b\u0435\u043b\u0438\u0432\u0430\u043d\u0438\u044f \u043c\u043e\u0433\u0443\u0442 \u0431\u044b\u0442\u044c \u0441\u043b\u043e\u0436\u043d\u044b\u043c\u0438 \u0432 \u043d\u0430\u043f\u0438\u0441\u0430\u043d\u0438\u0438, \u0442\u0440\u0435\u0431\u0443\u044f \u044f\u0432\u043d\u043e\u0439 \u0441\u0438\u043d\u0445\u0440\u043e\u043d\u0438\u0437\u0430\u0446\u0438\u0438 \u043c\u0435\u0436\u0434\u0443 \u043f\u043e\u0442\u043e\u043a\u0430\u043c\u0438, \u043f\u043e\u0441\u043a\u043e\u043b\u044c\u043a\u0443 \u043e\u043d\u0438 \u043e\u0434\u043d\u043e\u0432\u0440\u0435\u043c\u0435\u043d\u043d\u043e \u0443\u043c\u0435\u043d\u044c\u0448\u0430\u044e\u0442 \u043e\u0434\u043d\u0443 \u0438 \u0442\u0443 \u0436\u0435 \u0441\u0442\u0440\u043e\u043a\u0443 X<em>X<\/em>. \u0411\u043e\u043b\u044c\u0448\u0430\u044f \u0447\u0430\u0441\u0442\u044c \u044d\u0442\u043e\u0439 \u0441\u043b\u043e\u0436\u043d\u043e\u0441\u0442\u0438 \u0438\u0441\u0447\u0435\u0437\u0430\u0435\u0442 \u0432 Triton, \u0433\u0434\u0435 \u043a\u0430\u0436\u0434\u044b\u0439 \u044d\u043a\u0437\u0435\u043c\u043f\u043b\u044f\u0440 \u044f\u0434\u0440\u0430 \u0437\u0430\u0433\u0440\u0443\u0436\u0430\u0435\u0442 \u0438\u043d\u0442\u0435\u0440\u0435\u0441\u0443\u044e\u0449\u0443\u044e \u0441\u0442\u0440\u043e\u043a\u0443 \u0438 \u043f\u043e\u0441\u043b\u0435\u0434\u043e\u0432\u0430\u0442\u0435\u043b\u044c\u043d\u043e \u043d\u043e\u0440\u043c\u0430\u043b\u0438\u0437\u0443\u0435\u0442 \u0435\u0451, \u0438\u0441\u043f\u043e\u043b\u044c\u0437\u0443\u044f NumPy-\u043f\u043e\u0434\u043e\u0431\u043d\u044b\u0435 \u043f\u0440\u0438\u043c\u0438\u0442\u0438\u0432\u044b.<\/p>\n<p>Fused softmax \u0432 Triton:<\/p>\n<pre><code class=\"python\">import triton import triton.language as tl  @triton.jit def softmax(Y, stride_ym, stride_yn, X, stride_xm, stride_xn, M, N):     # row index     m = tl.program_id(0)     # col indices     # this specific kernel only works for matrices that      # have less than BLOCK_SIZE columns     BLOCK_SIZE = 1024     n = tl.arange(0, BLOCK_SIZE)     # the memory address of all the elements     # that we want to load can be computed as follows     X = X + m * stride_xm + n * stride_xn     # load input data; pad out-of-bounds elements with 0      x = tl.load(X, mask=n &lt; N, other=-float('inf'))     # compute numerically-stable softmax     z = x - tl.max(x, axis=0)     num = tl.exp(z)     denom = tl.sum(num, axis=0)     y = num \/ denom     # write back to Y     Y = Y + m * stride_ym + n * stride_yn     tl.store(Y, y, mask=n &lt; N)  import torch # Allocate input\/output tensors X = torch.normal(0, 1, size=(583, 931), device='cuda') Y = torch.empty_like(X) # SPMD launch grid grid = (X.shape[0], ) # enqueue GPU kernel softmax[grid](Y, Y.stride(0), Y.stride(1),                X, X.stride(0), X.stride(1),               X.shape[0]    , X.shape[1])<\/code><\/pre>\n<p>\u041e\u0431\u0440\u0430\u0442\u0438\u0442\u0435 \u0432\u043d\u0438\u043c\u0430\u043d\u0438\u0435, \u0447\u0442\u043e Triton JIT \u0440\u0430\u0441\u0441\u043c\u0430\u0442\u0440\u0438\u0432\u0430\u0435\u0442 X \u0438 Y \u043a\u0430\u043a <em>\u0443\u043a\u0430\u0437\u0430\u0442\u0435\u043b\u0438<\/em>, \u0430 \u043d\u0435 \u0442\u0435\u043d\u0437\u043e\u0440\u044b; \u043c\u044b \u0447\u0443\u0432\u0441\u0442\u0432\u043e\u0432\u0430\u043b\u0438, \u0447\u0442\u043e \u0441\u043e\u0445\u0440\u0430\u043d\u0435\u043d\u0438\u0435 \u043d\u0438\u0437\u043a\u043e\u0443\u0440\u043e\u0432\u043d\u0435\u0432\u043e\u0433\u043e \u043a\u043e\u043d\u0442\u0440\u043e\u043b\u044f \u0434\u043e\u0441\u0442\u0443\u043f\u0430 \u043a \u043f\u0430\u043c\u044f\u0442\u0438 \u0432\u0430\u0436\u043d\u043e \u0432 \u043e\u0431\u0440\u0430\u0449\u0435\u043d\u0438\u0438 \u0441 \u0431\u043e\u043b\u0435\u0435 \u0441\u043b\u043e\u0436\u043d\u044b\u043c\u0438 \u0441\u0442\u0440\u0443\u043a\u0442\u0443\u0440\u0430\u043c\u0438 \u0434\u0430\u043d\u043d\u044b\u0445 (\u043d\u0430\u043f\u0440\u0438\u043c\u0435\u0440, \u0431\u043b\u043e\u0447\u043d\u043e-\u0440\u0430\u0441\u0441\u0435\u044f\u043d\u043d\u044b\u0445 (block-sparse) \u0442\u0435\u043d\u0437\u043e\u0440\u043e\u0432). \u0412\u0430\u0436\u043d\u043e \u043e\u0442\u043c\u0435\u0442\u0438\u0442\u044c, \u0447\u0442\u043e \u044d\u0442\u0430 \u043a\u043e\u043d\u043a\u0440\u0435\u0442\u043d\u0430\u044f \u0440\u0435\u0430\u043b\u0438\u0437\u0430\u0446\u0438\u044f softmax \u0441\u043e\u0445\u0440\u0430\u043d\u044f\u0435\u0442 \u0441\u0442\u0440\u043e\u043a\u0438 X<em>X<\/em> \u0432 SRAM \u043d\u0430 \u043f\u0440\u043e\u0442\u044f\u0436\u0435\u043d\u0438\u0438 \u0432\u0441\u0435\u0433\u043e \u043f\u0440\u043e\u0446\u0435\u0441\u0441\u0430 \u043d\u043e\u0440\u043c\u0430\u043b\u0438\u0437\u0430\u0446\u0438\u0438, \u0447\u0442\u043e \u043f\u043e\u0437\u0432\u043e\u043b\u044f\u0435\u0442 \u043c\u0430\u043a\u0441\u0438\u043c\u0430\u043b\u044c\u043d\u043e \u043f\u0435\u0440\u0435\u0438\u0441\u043f\u043e\u043b\u044c\u0437\u043e\u0432\u0430\u0442\u044c \u0434\u0430\u043d\u043d\u044b\u0435, \u043a\u043e\u0433\u0434\u0430 \u044d\u0442\u043e \u0432\u043e\u0437\u043c\u043e\u0436\u043d\u043e (~&lt;32K \u0441\u0442\u043e\u043b\u0431\u0446\u043e\u0432).\u00a0<\/p>\n<p>\u0418\u043d\u0430\u0447\u0435 \u0434\u0435\u043b\u043e \u043e\u0431\u0441\u0442\u043e\u0438\u0442 \u0432\u043e \u0432\u043d\u0443\u0442\u0440\u0435\u043d\u043d\u0435\u043c \u043a\u043e\u0434\u0435 CUDA PyTorch, \u0447\u044c\u0451 \u0438\u0441\u043f\u043e\u043b\u044c\u0437\u043e\u0432\u0430\u043d\u0438\u0435 \u0432\u0440\u0435\u043c\u0435\u043d\u043d\u043e\u0439 \u043f\u0430\u043c\u044f\u0442\u0438 \u0434\u0435\u043b\u0430\u0435\u0442 \u0435\u0433\u043e \u0431\u043e\u043b\u0435\u0435 \u043e\u0431\u0449\u0438\u043c, \u043d\u043e \u0437\u043d\u0430\u0447\u0438\u0442\u0435\u043b\u044c\u043d\u043e \u0431\u043e\u043b\u0435\u0435 \u043c\u0435\u0434\u043b\u0435\u043d\u043d\u044b\u043c (\u0441\u043c. \u043d\u0438\u0436\u0435).\u00a0<\/p>\n<blockquote>\n<p>Triton \u0443\u043f\u0440\u043e\u0449\u0430\u0435\u0442 \u0440\u0430\u0437\u0440\u0430\u0431\u043e\u0442\u043a\u0443 \u0441\u043f\u0435\u0446\u0438\u0430\u043b\u0438\u0437\u0438\u0440\u043e\u0432\u0430\u043d\u043d\u044b\u0445 \u044f\u0434\u0435\u0440, \u043a\u043e\u0442\u043e\u0440\u044b\u0435 \u043c\u043e\u0433\u0443\u0442 \u0431\u044b\u0442\u044c \u043d\u0430\u043c\u043d\u043e\u0433\u043e \u0431\u044b\u0441\u0442\u0440\u0435\u0435 \u044f\u0434\u0435\u0440 \u0432 \u0431\u0438\u0431\u043b\u0438\u043e\u0442\u0435\u043a\u0430\u0445 \u043e\u0431\u0449\u0435\u0433\u043e \u043d\u0430\u0437\u043d\u0430\u0447\u0435\u043d\u0438\u044f.<\/p>\n<\/blockquote>\n<figure class=\"full-width\"><figcaption>\u041f\u0440\u043e\u0438\u0437\u0432\u043e\u0434\u0438\u0442\u0435\u043b\u044c\u043d\u043e\u0441\u0442\u044c A100 fused softmax \u043f\u0440\u0438 M=4096<\/figcaption><\/figure>\n<p>\u0411\u043e\u043b\u0435\u0435 \u043d\u0438\u0437\u043a\u0430\u044f \u043f\u0440\u043e\u0438\u0437\u0432\u043e\u0434\u0438\u0442\u0435\u043b\u044c\u043d\u043e\u0441\u0442\u044c Torch (v1.9) JIT \u043f\u043e\u0434\u0447\u0451\u0440\u043a\u0438\u0432\u0430\u0435\u0442 \u0441\u043b\u043e\u0436\u043d\u043e\u0441\u0442\u044c \u0430\u0432\u0442\u043e\u043c\u0430\u0442\u0438\u0447\u0435\u0441\u043a\u043e\u0439 \u0433\u0435\u043d\u0435\u0440\u0430\u0446\u0438\u0438 \u043a\u043e\u0434\u0430 CUDA \u0438\u0437 \u043f\u043e\u0441\u043b\u0435\u0434\u043e\u0432\u0430\u0442\u0435\u043b\u044c\u043d\u043e\u0441\u0442\u0435\u0439 \u0432\u044b\u0441\u043e\u043a\u043e\u0443\u0440\u043e\u0432\u043d\u0435\u0432\u044b\u0445 \u0442\u0435\u043d\u0437\u043e\u0440\u043d\u044b\u0445 \u043e\u043f\u0435\u0440\u0430\u0446\u0438\u0439. Fused softmax \u043d\u0430 Torch JIT:<\/p>\n<pre><code class=\"python\">@torch.jit.script def softmax(x):     x_max = x.max(dim=1)[0]     z = x - x_max[:, None]     numerator = torch.exp(x)     denominator = numerator.sum(dim=1)     return numerator \/ denominator[:, None]<\/code><\/pre>\n<h2>\u041c\u0430\u0442\u0440\u0438\u0447\u043d\u043e\u0435 \u0443\u043c\u043d\u043e\u0436\u0435\u043d\u0438\u0435<\/h2>\n<p>\u0412\u043e\u0437\u043c\u043e\u0436\u043d\u043e\u0441\u0442\u044c \u043f\u0438\u0441\u0430\u0442\u044c &#171;\u0441\u043b\u0438\u0442\u044b\u0435&#187; (fused) \u044f\u0434\u0440\u0430 (\u044f\u0434\u0440\u0430 \u0441 \u043e\u0431\u043e\u0431\u0449\u0435\u043d\u0438\u0435\u043c \u044d\u0442\u0430\u043f\u043e\u0432 \u0432\u044b\u0447\u0438\u0441\u043b\u0435\u043d\u0438\u0439) \u0434\u043b\u044f \u043e\u043f\u0435\u0440\u0430\u0446\u0438\u0439 \u0438 \u0441\u043e\u043a\u0440\u0430\u0449\u0435\u043d\u0438\u0439 \u043f\u043e \u044d\u043b\u0435\u043c\u0435\u043d\u0442\u0430\u043c \u0432\u0430\u0436\u043d\u0430, \u043d\u043e \u043d\u0435\u0434\u043e\u0441\u0442\u0430\u0442\u043e\u0447\u043d\u0430, \u0443\u0447\u0438\u0442\u044b\u0432\u0430\u044f, \u0447\u0442\u043e \u0432 \u043d\u0435\u0439\u0440\u043e\u043d\u043d\u044b\u0445 \u0441\u0435\u0442\u044f\u0445 \u0437\u0430\u0434\u0430\u0447\u0438 \u0443\u043c\u043d\u043e\u0436\u0435\u043d\u0438\u044f \u043c\u0430\u0442\u0440\u0438\u0446 \u0437\u0430\u043d\u0438\u043c\u0430\u044e\u0442 \u0432\u0430\u0436\u043d\u043e\u0435 \u043c\u0435\u0441\u0442\u043e.\u00a0\u041a\u0430\u043a \u043e\u043a\u0430\u0437\u0430\u043b\u043e\u0441\u044c, Triton \u043e\u0442\u043b\u0438\u0447\u043d\u043e \u043f\u043e\u0434\u0445\u043e\u0434\u0438\u0442 \u0438 \u0434\u043b\u044f \u043d\u0438\u0445, \u0434\u043e\u0441\u0442\u0438\u0433\u0430\u044f \u043f\u0438\u043a\u043e\u0432\u043e\u0439 \u043f\u0440\u043e\u0438\u0437\u0432\u043e\u0434\u0438\u0442\u0435\u043b\u044c\u043d\u043e\u0441\u0442\u0438 \u043c\u0435\u043d\u0435\u0435 \u0447\u0435\u043c \u0432 25 \u0441\u0442\u0440\u043e\u043a Python. \u0421 \u0434\u0440\u0443\u0433\u043e\u0439 \u0441\u0442\u043e\u0440\u043e\u043d\u044b, \u0440\u0435\u0430\u043b\u0438\u0437\u0430\u0446\u0438\u044f \u0447\u0435\u0433\u043e-\u0442\u043e \u043f\u043e\u0434\u043e\u0431\u043d\u043e\u0433\u043e \u0432 CUDA \u043f\u043e\u0442\u0440\u0435\u0431\u0443\u0435\u0442 <a href=\"https:\/\/github.com\/openai\/blocksparse\/blob\/master\/src\/matmul_op_gpu.cu\">\u0433\u043e\u0440\u0430\u0437\u0434\u043e \u0431\u043e\u043b\u044c\u0448\u0435 \u0443\u0441\u0438\u043b\u0438\u0439<\/a> \u0438 \u0434\u0430\u0436\u0435, \u0441\u043a\u043e\u0440\u0435\u0435 \u0432\u0441\u0435\u0433\u043e, \u043f\u0440\u0438\u0432\u0435\u0434\u0451\u0442 \u043a \u0441\u043d\u0438\u0436\u0435\u043d\u0438\u044e \u043f\u0440\u043e\u0438\u0437\u0432\u043e\u0434\u0438\u0442\u0435\u043b\u044c\u043d\u043e\u0441\u0442\u0438.<\/p>\n<p>\u0423\u043c\u043d\u043e\u0436\u0435\u043d\u0438\u0435 \u043c\u0430\u0442\u0440\u0438\u0446 \u043f\u0440\u0438 \u043f\u043e\u043c\u043e\u0449\u0438 Triton:<\/p>\n<pre><code class=\"python\">@triton.jit def matmul(A, B, C, M, N, K, stride_am, stride_ak,              stride_bk, stride_bn, stride_cm, stride_cn,             **META):     # extract metaparameters     BLOCK_M, GROUP_M = META['BLOCK_M'], META['GROUP_M']     BLOCK_N = META['BLOCK_N']     BLOCK_K = META['BLOCK_K']     # programs are grouped together to improve L2 hit rate     _pid_m = tl.program_id(0)     _pid_n = tl.program_id(1)     pid_m = _pid_m \/\/ GROUP_M     pid_n = (_pid_n * GROUP_M) + (_pid_m % GROUP_M)     # rm (resp. rn) denotes a range of indices     # for rows (resp. col) of C     rm = pid_m * BLOCK_M + tl.arange(0, BLOCK_M)     rn = pid_n * BLOCK_N + tl.arange(0, BLOCK_N)     # rk denotes a range of indices for columns      # (resp. rows) of A (resp. B)     rk = tl.arange(0, BLOCK_K)     # the memory addresses of elements in the first block of     # A and B can be computed using numpy-style broadcasting     A = A + (rm[:, None] * stride_am + rk[None, :] * stride_ak)     B = B + (rk [:, None] * stride_bk  + rn[None, :] * stride_bn)     # initialize and iteratively update accumulator     acc = tl.zeros((BLOCK_M, BLOCK_N), dtype=tl.float32)     for k in range(K, 0, -BLOCK_K):         a = tl.load(A)         b = tl.load(B)         # block level matrix multiplication         acc += tl.dot(a, b)         # increment pointers so that the next blocks of A and B         # are loaded during the next iteration         A += BLOCK_K * stride_ak         B += BLOCK_K * stride_bk     # fuse leaky ReLU if desired     # acc = tl.where(acc >= 0, acc, alpha * acc)     # write back result     C = C + (rm[:, None] * stride_cm + rn[None, :] * stride_cn)     mask = (rm[:, None] &lt; M) &amp; (rn[None, :] &lt; N)     tl.store(C, acc, mask=mask)<\/code><\/pre>\n<p>\u041e\u0434\u043d\u0438\u043c \u0438\u0437 \u0432\u0430\u0436\u043d\u044b\u0445 \u043f\u0440\u0435\u0438\u043c\u0443\u0449\u0435\u0441\u0442\u0432 \u0440\u0443\u043a\u043e\u043f\u0438\u0441\u043d\u044b\u0445 \u044f\u0434\u0435\u0440<\/p>\n<\/div>\n<\/div>\n<\/div>\n<\/div>\n","protected":false},"author":1,"featured_media":0,"comment_status":"open","ping_status":"open","sticky":false,"template":"","format":"standard","meta":{"footnotes":""},"categories":[],"tags":[],"class_list":["post-416011","post","type-post","status-publish","format-standard","hentry"],"_links":{"self":[{"href":"https:\/\/savepearlharbor.com\/index.php?rest_route=\/wp\/v2\/posts\/416011","targetHints":{"allow":["GET"]}}],"collection":[{"href":"https:\/\/savepearlharbor.com\/index.php?rest_route=\/wp\/v2\/posts"}],"about":[{"href":"https:\/\/savepearlharbor.com\/index.php?rest_route=\/wp\/v2\/types\/post"}],"author":[{"embeddable":true,"href":"https:\/\/savepearlharbor.com\/index.php?rest_route=\/wp\/v2\/users\/1"}],"replies":[{"embeddable":true,"href":"https:\/\/savepearlharbor.com\/index.php?rest_route=%2Fwp%2Fv2%2Fcomments&post=416011"}],"version-history":[{"count":0,"href":"https:\/\/savepearlharbor.com\/index.php?rest_route=\/wp\/v2\/posts\/416011\/revisions"}],"wp:attachment":[{"href":"https:\/\/savepearlharbor.com\/index.php?rest_route=%2Fwp%2Fv2%2Fmedia&parent=416011"}],"wp:term":[{"taxonomy":"category","embeddable":true,"href":"https:\/\/savepearlharbor.com\/index.php?rest_route=%2Fwp%2Fv2%2Fcategories&post=416011"},{"taxonomy":"post_tag","embeddable":true,"href":"https:\/\/savepearlharbor.com\/index.php?rest_route=%2Fwp%2Fv2%2Ftags&post=416011"}],"curies":[{"name":"wp","href":"https:\/\/api.w.org\/{rel}","templated":true}]}}