{"id":341925,"date":"2022-11-30T09:00:11","date_gmt":"2022-11-30T09:00:11","guid":{"rendered":"http:\/\/savepearlharbor.com\/?p=341925"},"modified":"-0001-11-30T00:00:00","modified_gmt":"-0001-11-29T21:00:00","slug":"","status":"publish","type":"post","link":"https:\/\/savepearlharbor.com\/?p=341925","title":{"rendered":"<span>\u041f\u0438\u0448\u0435\u043c \u043a\u0430\u0441\u0442\u043e\u043c\u043d\u044b\u0435 CUDA-\u044f\u0434\u0440\u0430 \u043d\u0430 Triton<\/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-1\">\n<div xmlns=\"http:\/\/www.w3.org\/1999\/xhtml\"><img decoding=\"async\" src=\"https:\/\/habrastorage.org\/r\/w1560\/webt\/gt\/au\/sk\/gtausku8oyxsy-a60nmlxenbypq.png\" alt=\"image\" data-src=\"https:\/\/habrastorage.org\/webt\/gt\/au\/sk\/gtausku8oyxsy-a60nmlxenbypq.png\"\/><\/p>\n<p>  Triton \u2013 \u044d\u0442\u043e \u044f\u0437\u044b\u043a\u043e\u0432\u043e\u0439 \u043a\u043e\u043c\u043f\u0438\u043b\u044f\u0442\u043e\u0440 \u0434\u043b\u044f \u0441\u043e\u0437\u0434\u0430\u043d\u0438\u044f \u0441\u0438\u043b\u044c\u043d\u043e \u043e\u043f\u0442\u0438\u043c\u0438\u0437\u0438\u0440\u043e\u0432\u0430\u043d\u043d\u044b\u0445 \u044f\u0434\u0435\u0440 CUDA. \u0417\u0434\u0435\u0441\u044c \u0431\u0443\u0434\u0443\u0442 \u0438\u0437\u043b\u043e\u0436\u0435\u043d\u044b \u043e\u0441\u043d\u043e\u0432\u044b \u043f\u0440\u043e\u0433\u0440\u0430\u043c\u043c\u0438\u0440\u043e\u0432\u0430\u043d\u0438\u044f \u0434\u043b\u044f GPU \u0438 \u0440\u0430\u0441\u0441\u043a\u0430\u0437\u0430\u043d\u043e, \u043a\u0430\u043a \u0434\u043b\u044f \u044d\u0442\u043e\u0439 \u0446\u0435\u043b\u0438 \u0438\u0441\u043f\u043e\u043b\u044c\u0437\u0443\u0435\u0442\u0441\u044f Triton.<\/p>\n<p>  \u0423\u0447\u0438\u0442\u044b\u0432\u0430\u044f \u043d\u044b\u043d\u0435\u0448\u043d\u0438\u0439 \u0443\u0441\u043f\u0435\u0445 \u0433\u043b\u0443\u0431\u043e\u043a\u043e\u0433\u043e \u043e\u0431\u0443\u0447\u0435\u043d\u0438\u044f \u0438 \u0432\u0430\u043b \u0438\u0441\u0441\u043b\u0435\u0434\u043e\u0432\u0430\u0442\u0435\u043b\u044c\u0441\u043a\u0438\u0445 \u0441\u0442\u0430\u0442\u0435\u0439 \u043d\u0430 \u044d\u0442\u0443 \u0442\u0435\u043c\u0443, \u0447\u0430\u0441\u0442\u043e \u0432\u043e\u0437\u043d\u0438\u043a\u0430\u0435\u0442 \u0442\u0430\u043a\u0430\u044f \u0441\u0438\u0442\u0443\u0430\u0446\u0438\u044f: \u0440\u043e\u0436\u0434\u0430\u0435\u0442\u0441\u044f \u043a\u0430\u043a\u0430\u044f-\u043d\u0438\u0431\u0443\u0434\u044c \u043d\u043e\u0432\u0430\u044f \u0438\u0434\u0435\u044f, \u0438 \u0432\u044b\u044f\u0441\u043d\u044f\u0435\u0442\u0441\u044f, \u0447\u0442\u043e \u0434\u043b\u044f \u043d\u0435\u0435 \u043d\u0435 \u043f\u043e\u0434\u0434\u0435\u0440\u0436\u0438\u0432\u0430\u0435\u0442\u0441\u044f \u0430\u043f\u043f\u0430\u0440\u0430\u0442\u043d\u043e\u0435 \u0443\u0441\u043a\u043e\u0440\u0435\u043d\u0438\u0435. \u0422\u043e\u0447\u043d\u0435\u0435, \u0441\u0442\u043e\u0438\u0442 \u0432\u0430\u043c \u0438\u0437\u043e\u0431\u0440\u0435\u0441\u0442\u0438 \u043d\u043e\u0432\u0443\u044e \u0444\u0443\u043d\u043a\u0446\u0438\u044e \u0430\u043a\u0442\u0438\u0432\u0430\u0446\u0438\u0438 \u0438\u043b\u0438 \u043c\u0435\u0445\u0430\u043d\u0438\u0437\u043c \u0441\u0430\u043c\u043e\u0432\u043d\u0438\u043c\u0430\u043d\u0438\u044f \u2013 \u043d\u0430\u043c \u0441\u0440\u0430\u0437\u0443 \u043f\u0440\u0438\u0445\u043e\u0434\u0438\u0442\u0441\u044f \u043f\u0440\u0438\u0431\u0435\u0433\u0430\u0442\u044c \u043a \u0432\u043e\u0437\u043c\u043e\u0436\u043d\u043e\u0441\u0442\u044f\u043c PyTorch\/Tensorflow \u0434\u043b\u044f \u043e\u0431\u0440\u0430\u0431\u043e\u0442\u043a\u0438 \u043f\u0440\u044f\u043c\u043e\u0433\u043e \u0438 \u043e\u0431\u0440\u0430\u0442\u043d\u043e\u0433\u043e \u043f\u0440\u043e\u0445\u043e\u0434\u0430 \u0447\u0435\u0440\u0435\u0437 \u043c\u043e\u0434\u0443\u043b\u044c.<\/p>\n<p>  \u0412 \u0442\u0430\u043a\u0438\u0445 \u0441\u043b\u0443\u0447\u0430\u044f\u0445 \u043f\u0440\u0438\u043c\u0435\u043d\u0438\u043c, \u043d\u0430\u043f\u0440\u0438\u043c\u0435\u0440, PyTorch JIT. \u041d\u043e PyTorch JIT \u2013 \u044d\u0442\u043e \u0432\u044b\u0441\u043e\u043a\u043e\u0443\u0440\u043e\u0432\u043d\u0435\u0432\u044b\u0439 \u043a\u043e\u043c\u043f\u0438\u043b\u044f\u0442\u043e\u0440, \u0441\u043f\u043e\u0441\u043e\u0431\u043d\u044b\u0439 \u043e\u043f\u0442\u0438\u043c\u0438\u0437\u0438\u0440\u043e\u0432\u0430\u0442\u044c \u043b\u0438\u0448\u044c \u043d\u0435\u043a\u043e\u0442\u043e\u0440\u044b\u0435 \u0447\u0430\u0441\u0442\u0438 \u043a\u043e\u0434\u0430, \u043d\u043e \u043d\u0435\u043f\u0440\u0438\u0433\u043e\u0434\u043d\u044b\u0439 \u0434\u043b\u044f \u043d\u0430\u043f\u0438\u0441\u0430\u043d\u0438\u044f \u0441\u043f\u0435\u0446\u0438\u0430\u043b\u0438\u0437\u0438\u0440\u043e\u0432\u0430\u043d\u043d\u044b\u0445 \u044f\u0434\u0435\u0440 CUDA.<br \/>  <a name=\"habracut\"><\/a><br \/>  \u0421 \u043d\u0430\u043f\u0438\u0441\u0430\u043d\u0438\u0435\u043c \u044f\u0434\u0435\u0440 CUDA \u0435\u0441\u0442\u044c \u0438 \u0435\u0449\u0435 \u043e\u0434\u043d\u0430 \u043f\u0440\u043e\u0431\u043b\u0435\u043c\u0430 \u2013 \u0441\u043e\u0437\u0434\u0430\u0442\u044c \u0442\u0430\u043a\u043e\u0435 \u044f\u0434\u0440\u043e \u043d\u0435\u0432\u0435\u0440\u043e\u044f\u0442\u043d\u043e \u0441\u043b\u043e\u0436\u043d\u043e. \u041e\u043f\u0442\u0438\u043c\u0438\u0437\u0430\u0446\u0438\u044f \u0432\u044b\u0447\u0438\u0441\u043b\u0435\u043d\u0438\u0439 \u0441 \u0440\u0430\u0441\u0447\u0435\u0442\u043e\u043c \u043d\u0430 \u0438\u0441\u043f\u043e\u043b\u044c\u0437\u043e\u0432\u0430\u043d\u0438\u0435 \u043b\u043e\u043a\u0430\u043b\u044c\u043d\u043e\u0441\u0442\u0438 \u0438 \u043f\u0430\u0440\u0430\u043b\u043b\u0435\u043b\u0438\u0437\u043c\u0430 \u0442\u0440\u0435\u0431\u0443\u0435\u0442 \u043c\u043d\u043e\u0433\u043e \u0432\u0440\u0435\u043c\u0435\u043d\u0438 \u0438 \u0447\u0440\u0435\u0432\u0430\u0442\u0430 \u043e\u0448\u0438\u0431\u043a\u0430\u043c\u0438, \u0438 \u044d\u043a\u0441\u043f\u0435\u0440\u0442\u0430\u043c \u0432 \u044d\u0442\u0438\u0445 \u043e\u0431\u043b\u0430\u0441\u0442\u044f\u0445 \u0447\u0430\u0441\u0442\u043e \u0442\u0440\u0435\u0431\u0443\u0435\u0442\u0441\u044f \u0445\u043e\u0440\u043e\u0448\u043e \u043f\u043e\u0442\u0440\u0443\u0434\u0438\u0442\u044c\u0441\u044f, \u0447\u0442\u043e\u0431\u044b \u043d\u0430\u0443\u0447\u0438\u0442\u044c\u0441\u044f \u043f\u0438\u0441\u0430\u0442\u044c \u043a\u043e\u0434 CUDA. \u041a\u0440\u043e\u043c\u0435 \u0442\u043e\u0433\u043e, \u0430\u0440\u0445\u0438\u0442\u0435\u043a\u0442\u0443\u0440\u044b GPU \u0441\u0442\u0440\u0435\u043c\u0438\u0442\u0435\u043b\u044c\u043d\u043e \u0440\u0430\u0437\u0432\u0438\u0432\u0430\u044e\u0442\u0441\u044f; \u043d\u0430\u043f\u0440\u0438\u043c\u0435\u0440, \u0443\u0436\u0435 \u043f\u043e\u044f\u0432\u0438\u043b\u0438\u0441\u044c \u043d\u043e\u0432\u0435\u0439\u0448\u0438\u0435 \u0442\u0435\u043d\u0437\u043e\u0440\u043d\u044b\u0435 \u044f\u0434\u0440\u0430. \u041f\u043e\u044d\u0442\u043e\u043c\u0443 \u043e\u043a\u0430\u0437\u044b\u0432\u0430\u0435\u0442\u0441\u044f \u0435\u0449\u0435 \u0441\u043b\u043e\u0436\u043d\u0435\u0435 \u043f\u0438\u0441\u0430\u0442\u044c \u0442\u0430\u043a\u043e\u0439 \u043a\u043e\u0434, \u043a\u043e\u0442\u043e\u0440\u044b\u0439 \u0432\u044b\u0436\u0438\u043c\u0430\u043b \u0431\u044b \u0438\u0437 \u0438\u043c\u0435\u044e\u0449\u0435\u0433\u043e\u0441\u044f \u043e\u0431\u043e\u0440\u0443\u0434\u043e\u0432\u0430\u043d\u0438\u044f \u043c\u0430\u043a\u0441\u0438\u043c\u0443\u043c \u043f\u0440\u043e\u0438\u0437\u0432\u043e\u0434\u0438\u0442\u0435\u043b\u044c\u043d\u043e\u0441\u0442\u0438.\u0418\u043c\u0435\u043d\u043d\u043e \u0437\u0434\u0435\u0441\u044c \u043d\u0430 \u0441\u0446\u0435\u043d\u0443 \u0432\u044b\u0445\u043e\u0434\u0438\u0442 <a href=\"https:\/\/github.com\/openai\/triton\">Triton<\/a> \u043e\u0442 OpenAI. \u0412 \u0441\u043e\u0441\u0442\u0430\u0432 Triton \u0432\u0445\u043e\u0434\u0438\u0442 \u0442\u0440\u0438 \u043e\u0441\u043d\u043e\u0432\u043d\u044b\u0445 \u043a\u043e\u043c\u043f\u043e\u043d\u0435\u043d\u0442\u0430.<\/p>\n<p>  <img decoding=\"async\" src=\"https:\/\/habrastorage.org\/r\/w1560\/webt\/4i\/hp\/zy\/4ihpzyfhauvratbbpsceickvjss.png\" alt=\"image\" data-src=\"https:\/\/habrastorage.org\/webt\/4i\/hp\/zy\/4ihpzyfhauvratbbpsceickvjss.png\"\/><br \/>  <b><i><font color=\"#1c22be\">\u0420\u0438\u0441\u0443\u043d\u043e\u043a 1: \u041e\u0431\u0437\u043e\u0440 \u043e\u0441\u043d\u043e\u0432\u043d\u044b\u0445 \u043a\u043e\u043c\u043f\u043e\u043d\u0435\u043d\u0442\u043e\u0432 Triton.<\/font><\/i><\/b>  <\/p>\n<ol>\n<li>Triton-C: \u044d\u0442\u043e C-\u043f\u043e\u0434\u043e\u0431\u043d\u044b\u0439 \u044f\u0437\u044b\u043a, \u043e\u0440\u0438\u0435\u043d\u0442\u0438\u0440\u043e\u0432\u0430\u043d\u043d\u044b\u0439 \u0432 \u043e\u0441\u043d\u043e\u0432\u043d\u043e\u043c \u043d\u0430 \u043f\u0440\u043e\u0433\u0440\u0430\u043c\u043c\u0438\u0441\u0442\u043e\u0432, \u0443\u0436\u0435 \u0437\u043d\u0430\u043a\u043e\u043c\u044b\u0445 \u0441 CUDA.<\/li>\n<li>Triton-IR: \u044d\u0442\u043e \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. \u041f\u0440\u043e\u0433\u0440\u0430\u043c\u043c\u044b \u043d\u0430 Triton-IR \u0441\u043e\u0431\u0438\u0440\u0430\u044e\u0442\u0441\u044f \u043d\u0435\u043f\u043e\u0441\u0440\u0435\u0434\u0441\u0442\u0432\u0435\u043d\u043d\u043e \u0438\u0437 Triton-C. \u0415\u0441\u043b\u0438 \u043a\u043e\u0440\u043e\u0442\u043a\u043e, \u0432 LLVM \u043f\u0440\u0435\u0434\u043e\u0441\u0442\u0430\u0432\u043b\u044f\u0435\u0442\u0441\u044f \u043c\u043d\u043e\u0433\u043e \u0441\u043f\u0435\u0446\u0438\u0444\u0438\u0447\u043d\u044b\u0445 \u0430\u043f\u043f\u0430\u0440\u0430\u0442\u043d\u044b\u0445 \u043e\u043f\u0442\u0438\u043c\u0438\u0437\u0430\u0446\u0438\u0439, \u0447\u0442\u043e \u043f\u043e\u0437\u0432\u043e\u043b\u044f\u0435\u0442 \u043d\u0430\u043c \u043d\u0430\u043f\u0440\u044f\u043c\u0443\u044e \u0438\u0441\u043f\u043e\u043b\u044c\u0437\u043e\u0432\u0430\u0442\u044c \u043a\u043e\u043c\u043f\u0438\u043b\u044f\u0442\u043e\u0440 CUDA \u043e\u0442 Nvidia (NVCC), \u0447\u0442\u043e\u0431\u044b \u043f\u0440\u0438\u0441\u043f\u043e\u0441\u043e\u0431\u0438\u0442\u044c \u043d\u0430\u0448 \u043a\u043e\u0434 \u043f\u043e\u0434 \u043a\u043e\u043d\u043a\u0440\u0435\u0442\u043d\u043e\u0435 \u0430\u043f\u043f\u0430\u0440\u0430\u0442\u043d\u043e\u0435 \u043e\u0431\u0435\u0441\u043f\u0435\u0447\u0435\u043d\u0438\u0435.<\/li>\n<li>Triton-JIT: \u044d\u0442\u043e \u0431\u0435\u043a\u0435\u043d\u0434 \u0441 \u0434\u0438\u043d\u0430\u043c\u0438\u0447\u0435\u0441\u043a\u0438\u043c (JIT) \u043a\u043e\u043c\u043f\u0438\u043b\u044f\u0442\u043e\u0440\u043e\u043c \u0438 \u0438\u043d\u0441\u0442\u0440\u0443\u043c\u0435\u043d\u0442\u043e\u043c \u0433\u0435\u043d\u0435\u0440\u0430\u0446\u0438\u0438 \u043a\u043e\u0434\u0430. \u041f\u0440\u0435\u0434\u043d\u0430\u0437\u043d\u0430\u0447\u0435\u043d \u0434\u043b\u044f \u043a\u043e\u043c\u043f\u0438\u043b\u044f\u0446\u0438\u0438 \u043f\u0440\u043e\u0433\u0440\u0430\u043c\u043c, \u043d\u0430\u043f\u0438\u0441\u0430\u043d\u043d\u044b\u0445 \u043d\u0430 Triton-IR, \u0432 \u044d\u0444\u0444\u0435\u043a\u0442\u0438\u0432\u043d\u044b\u0439 \u0431\u0438\u0442\u043a\u043e\u0434 \u0434\u043b\u044f LLVM. \u0422\u0430\u043a\u0436\u0435 \u043e\u043d \u0441\u043e\u0434\u0435\u0440\u0436\u0438\u0442 \u043c\u043d\u043e\u0436\u0435\u0441\u0442\u0432\u043e \u043c\u0430\u0448\u0438\u043d\u043d\u043e-\u043d\u0435\u0437\u0430\u0432\u0438\u0441\u0438\u043c\u044b\u0445 \u043e\u043f\u0442\u0438\u043c\u0438\u0437\u0430\u0446\u0438\u0439, \u0447\u0442\u043e, \u043e\u043f\u044f\u0442\u044c \u0436\u0435, \u0441\u043e\u043a\u0440\u0430\u0449\u0430\u0435\u0442 \u0434\u043b\u044f \u0432\u0430\u0441 \u043e\u0431\u044a\u0435\u043c \u0440\u0430\u0431\u043e\u0442\u044b.<\/li>\n<\/ol>\n<p>  \u0421\u0430\u043c\u0430\u044f \u0438\u043d\u0442\u0435\u0440\u0435\u0441\u043d\u0430\u044f \u0447\u0430\u0441\u0442\u044c \u043f\u0440\u043e\u0435\u043a\u0442\u0430 Triton \u2013 \u044d\u0442\u043e \u0438\u043c\u0435\u043d\u043d\u043e Triton-JIT. \u0411\u043b\u0430\u0433\u043e\u0434\u0430\u0440\u044f \u044d\u0442\u043e\u043c\u0443 \u043a\u043e\u043c\u043f\u0438\u043b\u044f\u0442\u043e\u0440\u0443, \u043f\u0440\u043e\u0433\u0440\u0430\u043c\u043c\u0438\u0441\u0442\u044b \u043f\u043e\u0447\u0442\u0438 \u0431\u0435\u0437 \u043e\u043f\u044b\u0442\u0430 \u0440\u0430\u0431\u043e\u0442\u044b \u0441 CUDA \u043c\u043e\u0433\u0443\u0442 \u043f\u0438\u0441\u0430\u0442\u044c \u043d\u0430 Python \u0441\u0438\u043b\u044c\u043d\u043e \u043e\u043f\u0442\u0438\u043c\u0438\u0437\u0438\u0440\u043e\u0432\u0430\u043d\u043d\u044b\u0435 \u044f\u0434\u0440\u0430 CUDA. \u041f\u0440\u0435\u0436\u0434\u0435 \u0447\u0435\u043c \u043e\u0431\u0441\u0443\u0436\u0434\u0430\u0442\u044c Triton, \u0434\u0430\u0432\u0430\u0439\u0442\u0435 \u0440\u0430\u0437\u0431\u0435\u0440\u0435\u043c\u0441\u044f, \u043a\u0430\u043a \u0438\u043c\u0435\u043d\u043d\u043e \u043f\u0440\u043e\u0433\u0440\u0430\u043c\u043c\u044b CUDA \u0440\u0430\u0431\u043e\u0442\u0430\u044e\u0442 \u043d\u0430 GPU.<\/p>\n<p>  \u041f\u043e\u043b\u0435\u0437\u043d\u044b\u0435 \u0441\u0441\u044b\u043b\u043a\u0438:  <\/p>\n<ul>\n<li><a href=\"http:\/\/www.eecs.harvard.edu\/~htk\/publication\/2019-mapl-tillet-kung-cox.pdf\">Triton: An Intermediate Language and Compiler for Tiled Neural Network Computations<\/a><\/li>\n<li><a href=\"https:\/\/openai.com\/blog\/triton\/\">Introducing Triton: Open-Source GPU Programming for Neural Networks<\/a><\/li>\n<li><a href=\"https:\/\/github.com\/openai\/triton\">triton github<\/a><\/li>\n<li><a href=\"https:\/\/triton-lang.org\/\">\u0414\u043e\u043a\u0443\u043c\u0435\u043d\u0442\u0430\u0446\u0438\u044f \u043f\u043e triton<\/a><\/li>\n<\/ul>\n<p>  <\/p>\n<h2><font color=\"#1c22be\">\u276f<\/font> <font color=\"#454cee\">\u041e\u0441\u043d\u043e\u0432\u044b \u043f\u0440\u043e\u0433\u0440\u0430\u043c\u043c\u0438\u0440\u043e\u0432\u0430\u043d\u0438\u044f \u0434\u043b\u044f GPU<\/font><\/h2>\n<p>  \u041d\u0430\u0447\u043d\u0435\u043c \u0441 \u0426\u041f (\u0445\u043e\u0441\u0442\u0430). \u0426\u041f \u043e\u0431\u043b\u0430\u0434\u0430\u0435\u0442 \u0434\u043e\u0441\u0442\u0443\u043f\u043e\u043c \u043a \u043e\u043f\u0435\u0440\u0430\u0442\u0438\u0432\u043d\u043e\u0439 \u043f\u0430\u043c\u044f\u0442\u0438 (RAM), \u0434\u0438\u0441\u043a\u0430\u043c \u0434\u043b\u044f \u0445\u0440\u0430\u043d\u0435\u043d\u0438\u044f \u0434\u0430\u043d\u043d\u044b\u0445 \u0438 \u043a \u043b\u044e\u0431\u043e\u0439 \u043f\u043e\u0434\u043a\u043b\u044e\u0447\u0435\u043d\u043d\u043e\u0439 \u043f\u0435\u0440\u0438\u0444\u0435\u0440\u0438\u0438. \u0412 \u0441\u0432\u043e\u044e \u043e\u0447\u0435\u0440\u0435\u0434\u044c, GPU (\u0443\u0441\u0442\u0440\u043e\u0439\u0441\u0442\u0432\u043e) \u043d\u0435 \u0438\u043c\u0435\u0435\u0442 \u0434\u043e\u0441\u0442\u0443\u043f\u0430 \u043d\u0438 \u043a RAM, \u043d\u0438 \u043a \u0447\u0435\u043c\u0443 \u0438\u0437 \u0432\u044b\u0448\u0435\u043f\u0435\u0440\u0435\u0447\u0438\u0441\u043b\u0435\u043d\u043d\u043e\u0433\u043e. \u0423 GPU \u0435\u0441\u0442\u044c \u0441\u0432\u043e\u044f \u0441\u043e\u0431\u0441\u0442\u0432\u0435\u043d\u043d\u0430\u044f \u043f\u0430\u043c\u044f\u0442\u044c, \u043a\u043e\u0442\u043e\u0440\u0430\u044f \u043d\u0430\u0437\u044b\u0432\u0430\u0435\u0442\u0441\u044f VRAM. \u0427\u0442\u043e\u0431\u044b GPU \u043c\u043e\u0433 \u0440\u0430\u0431\u043e\u0442\u0430\u0442\u044c, \u0434\u0430\u043d\u043d\u044b\u0435 \u043d\u0443\u0436\u043d\u043e \u0441\u043a\u043e\u043f\u0438\u0440\u043e\u0432\u0430\u0442\u044c \u0441 \u0426\u041f \u043d\u0430 GPU, \u0430 \u0437\u0430\u0442\u0435\u043c \u043e\u0431\u0440\u0430\u0431\u043e\u0442\u0430\u043d\u043d\u044b\u0435 \u0434\u0430\u043d\u043d\u044b\u0435 \u0434\u043e\u043b\u0436\u043d\u044b \u0431\u044b\u0442\u044c \u0441\u043a\u043e\u043f\u0438\u0440\u043e\u0432\u0430\u043d\u044b \u043e\u0431\u0440\u0430\u0442\u043d\u043e \u0441 GPU \u043d\u0430 \u0426\u041f, \u0447\u0442\u043e\u0431\u044b \u0426\u041f \u043c\u043e\u0433 \u0438\u0445 \u0433\u0434\u0435-\u043d\u0438\u0431\u0443\u0434\u044c \u0441\u043e\u0445\u0440\u0430\u043d\u0438\u0442\u044c \u0438\u043b\u0438 \u043f\u043e\u0434\u0435\u043b\u0438\u0442\u044c\u0441\u044f \u0438\u043c\u0438 \u0441 \u043f\u043e\u0434\u043a\u043b\u044e\u0447\u0435\u043d\u043d\u043e\u0439 \u043f\u0435\u0440\u0438\u0444\u0435\u0440\u0438\u0435\u0439.<\/p>\n<p>  <i><b>\u041f\u0440\u0438\u043c\u0435\u0447\u0430\u043d\u0438\u0435:<\/b> \u0438\u043c\u0435\u043d\u043d\u043e \u043f\u043e\u044d\u0442\u043e\u043c\u0443 \u0441\u043b\u0435\u0434\u0443\u0435\u0442 \u043c\u0438\u043d\u0438\u043c\u0438\u0437\u0438\u0440\u043e\u0432\u0430\u0442\u044c \u043a\u0443\u0440\u0441\u0438\u0440\u043e\u0432\u0430\u043d\u0438\u0435 \u0434\u0430\u043d\u043d\u044b\u0445 \u043c\u0435\u0436\u0434\u0443 \u0426\u041f \u0438 GPU, \u043d\u0430\u0441\u043a\u043e\u043b\u044c\u043a\u043e \u044d\u0442\u043e \u0432\u043e\u0437\u043c\u043e\u0436\u043d\u043e. \u041f\u0440\u043e\u0434\u0443\u043c\u0430\u0439\u0442\u0435, \u043a\u0430\u043a \u043c\u043e\u0436\u043d\u043e \u0437\u0430\u0433\u0440\u0443\u0436\u0430\u0442\u044c \u0434\u0430\u043d\u043d\u044b\u0435 \u043d\u0435\u0431\u043e\u043b\u044c\u0448\u0438\u043c\u0438 \u0444\u0440\u0430\u0433\u043c\u0435\u043d\u0442\u0430\u043c\u0438 \u0434\u043b\u044f \u043f\u0430\u0440\u0430\u043b\u043b\u0435\u043b\u044c\u043d\u043e\u0439 \u043e\u0431\u0440\u0430\u0431\u043e\u0442\u043a\u0438 \u0438\u043b\u0438 \u0438\u043c\u043f\u043e\u0440\u0442\u0438\u0440\u043e\u0432\u0430\u0442\u044c \u0434\u0430\u043d\u043d\u044b\u0435 \u0442\u0430\u043a, \u0447\u0442\u043e\u0431\u044b \u043f\u0440\u0438\u0441\u043f\u043e\u0441\u043e\u0431\u0438\u0442\u044c \u0438\u0445 \u0434\u043b\u044f \u043c\u043d\u043e\u0433\u043e\u043a\u0440\u0430\u0442\u043d\u043e\u0433\u043e \u0438\u0441\u043f\u043e\u043b\u044c\u0437\u043e\u0432\u0430\u043d\u0438\u044f, \u043f\u0440\u0435\u0436\u0434\u0435, \u0447\u0435\u043c \u0438\u043c\u043f\u043e\u0440\u0442\u0438\u0440\u043e\u0432\u0430\u0442\u044c \u0441\u043b\u0435\u0434\u0443\u044e\u0449\u0443\u044e \u043f\u043e\u0440\u0446\u0438\u044e \u0434\u0430\u043d\u043d\u044b\u0445.<\/i><\/p>\n<p>  \u0412 CUDA \u0437\u0430\u043f\u0443\u0441\u043a\u0430\u0435\u0442\u0441\u044f \u043c\u043d\u043e\u0436\u0435\u0441\u0442\u0432\u043e <b>\u043f\u043e\u0442\u043e\u043a\u043e\u0432<\/b>, \u0441\u0433\u0440\u0443\u043f\u043f\u0438\u0440\u043e\u0432\u0430\u043d\u043d\u044b\u0445 \u0432 \u0431\u043b\u043e\u043a\u0438, \u043a\u043e\u0442\u043e\u0440\u044b\u0435 \u043e\u0431\u0440\u0430\u0437\u0443\u044e\u0442 <b>\u0433\u0440\u0438\u0434<\/b>. \u0412\u0441\u0435 <i>\u043f\u043e\u0442\u043e\u043a\u0438<\/i> \u0432 \u043e\u0434\u043d\u043e\u043c \u0431\u043b\u043e\u043a\u0435 \u043c\u043e\u0433\u0443\u0442 \u043e\u0431\u043c\u0435\u043d\u0438\u0432\u0430\u0442\u044c\u0441\u044f \u0438\u043d\u0444\u043e\u0440\u043c\u0430\u0446\u0438\u0435\u0439 \u0434\u0440\u0443\u0433 \u0441 \u0434\u0440\u0443\u0433\u043e\u043c. \u0412 \u043a\u0430\u0436\u0434\u043e\u043c \u0431\u043b\u043e\u043a\u0435 \u043c\u043e\u0436\u043d\u043e \u0437\u0430\u043f\u0443\u0441\u0442\u0438\u0442\u044c \u0434\u043e 1024 \u043f\u043e\u0442\u043e\u043a\u043e\u0432, \u0430 \u0437\u0430 \u043e\u0434\u0438\u043d \u0437\u0430\u043f\u0443\u0441\u043a \u043c\u043e\u0436\u0435\u0442 \u0431\u044b\u0442\u044c \u0441\u0434\u0435\u043b\u0430\u043d\u043e \u0434\u043e 2^32\u20141 \u0431\u043b\u043e\u043a\u043e\u0432. \u0422\u0430\u043a\u043e\u0439 \u043f\u0440\u0438\u043c\u0435\u0440 \u043f\u0440\u0438\u0432\u0435\u0434\u0435\u043d \u043d\u0430 \u0440\u0438\u0441. 2.<\/p>\n<p>  <img decoding=\"async\" src=\"https:\/\/habrastorage.org\/r\/w1560\/webt\/sg\/9s\/eg\/sg9segav4gvim5hmgzuua-4upps.png\" alt=\"image\" data-src=\"https:\/\/habrastorage.org\/webt\/sg\/9s\/eg\/sg9segav4gvim5hmgzuua-4upps.png\"\/><br \/>  <b><i><font color=\"#1c22be\">\u0420\u0438\u0441\u0443\u043d\u043e\u043a 2: \u0410\u0440\u0445\u0438\u0442\u0435\u043a\u0442\u0443\u0440\u0430 \u043f\u0440\u043e\u0433\u0440\u0430\u043c\u043c CUDA.<\/font><\/i><\/b><\/p>\n<p>  \u0418\u0434\u0435\u044f, \u043d\u0430 \u043a\u043e\u0442\u043e\u0440\u043e\u0439 \u043e\u0441\u043d\u043e\u0432\u0430\u043d\u043e \u0438\u0441\u043f\u043e\u043b\u044c\u0437\u043e\u0432\u0430\u043d\u0438\u0435 \u0431\u043b\u043e\u043a\u043e\u0432 \u0442\u0430\u043a\u043e\u0432\u0430, \u0447\u0442\u043e \u0432\u0430\u043c \u043d\u0435 \u043f\u043e\u0442\u0440\u0435\u0431\u0443\u0435\u0442\u0441\u044f \u043f\u0435\u0440\u0435\u043f\u0438\u0441\u044b\u0432\u0430\u0442\u044c \u0432\u0430\u0448 \u043a\u043e\u0434, \u0435\u0441\u043b\u0438 \u0432 \u0431\u0443\u0434\u0443\u0449\u0435\u043c \u0434\u043e\u0432\u0435\u0434\u0435\u0442\u0441\u044f \u0438\u043c\u0435\u0442\u044c \u0434\u0435\u043b\u043e \u0441 \u043d\u043e\u0432\u044b\u043c GPU. \u0422\u0430\u043a \u0447\u0442\u043e \u043d\u043e\u0432\u044b\u0439 GPU \u043f\u0440\u043e\u0441\u0442\u043e \u0441\u043c\u043e\u0436\u0435\u0442 \u043a\u043e\u043d\u043a\u0443\u0440\u0435\u043d\u0442\u043d\u043e \u0432\u044b\u043f\u043e\u043b\u043d\u044f\u0442\u044c \u0431\u043e\u043b\u044c\u0448\u0435\u0435 \u043a\u043e\u043b\u0438\u0447\u0435\u0441\u0442\u0432\u043e \u0431\u043b\u043e\u043a\u043e\u0432, \u0431\u0435\u0437 \u043a\u0430\u043a\u0438\u0445-\u043b\u0438\u0431\u043e \u0438\u0437\u043c\u0435\u043d\u0435\u043d\u0438\u0439 \u0432 \u043a\u043e\u0434\u0435.<\/p>\n<h2><font color=\"#1c22be\">\u276f<\/font> <font color=\"#454cee\">\u0421\u0440\u0430\u0432\u043d\u0438\u0432\u0430\u0435\u043c \u043f\u0440\u043e\u0433\u0440\u0430\u043c\u043c\u0443 \u0434\u043b\u044f \u0426\u041f \u0438 \u043f\u0440\u043e\u0433\u0440\u0430\u043c\u043c\u0443 \u0434\u043b\u044f CUDA<\/font><\/h2>\n<p>  \u041d\u0435 \u0432\u0434\u0430\u0432\u0430\u044f\u0441\u044c \u0432 \u0442\u0435\u0445\u043d\u0438\u0447\u0435\u0441\u043a\u0438\u0435 \u0442\u043e\u043d\u043a\u043e\u0441\u0442\u0438, \u0434\u0430\u0432\u0430\u0439\u0442\u0435 \u0440\u0430\u0441\u0441\u043c\u043e\u0442\u0440\u0438\u043c \u043f\u0440\u043e\u0441\u0442\u043e\u0439 \u043f\u0440\u0438\u043c\u0435\u0440: \u043a\u0430\u043a \u0441\u043b\u043e\u0436\u0438\u0442\u044c \u0434\u0432\u0430 \u043c\u0430\u0441\u0441\u0438\u0432\u0430, \u0438\u043c\u0435\u044e\u0449\u0438\u0445 \u0434\u043b\u0438\u043d\u0443 3.<\/p>\n<pre><code class=\"python\">arr1 = [1, 2, 3]     arr2 = [10, 11, 12]<\/code><\/pre>\n<p>  \u0415\u0441\u043b\u0438 \u0431\u044b \u043c\u044b \u0445\u043e\u0442\u0435\u043b\u0438 \u0441\u043b\u043e\u0436\u0438\u0442\u044c \u044d\u0442\u0438 \u043c\u0430\u0441\u0441\u0438\u0432\u044b \u0432 C++, \u0442\u043e \u0441\u043e\u0437\u0434\u0430\u043b\u0438 \u0431\u044b \u0446\u0438\u043a\u043b for, \u043a\u043e\u0442\u043e\u0440\u044b\u0439 \u0441\u043e\u0432\u0435\u0440\u0448\u0438\u043b \u0431\u044b \u0442\u0440\u0438 \u0438\u0442\u0435\u0440\u0430\u0446\u0438\u0438 (\u043f\u0440\u0435\u0434\u043f\u043e\u043b\u0430\u0433\u0430\u0435\u0442\u0441\u044f, \u0447\u0442\u043e \u0443 \u043d\u0430\u0441 \u043e\u0434\u043d\u043e\u043f\u043e\u0442\u043e\u0447\u043d\u0430\u044f \u043f\u0440\u043e\u0433\u0440\u0430\u043c\u043c\u0430).<br \/>  \u041d\u043e \u0432 CUDA \u043c\u044b \u0437\u0430\u043f\u0443\u0441\u0442\u0438\u043c 3 \u043f\u043e\u0442\u043e\u043a\u0430, \u0438 \u043a\u0430\u0436\u0434\u044b\u0439 \u0438\u0437 \u043d\u0438\u0445 \u0432\u044b\u043f\u043e\u043b\u043d\u0438\u0442 \u0441\u043b\u043e\u0436\u0435\u043d\u0438\u0435 \u0432 \u0441\u0432\u043e\u0435\u043c \u0438\u043d\u0434\u0435\u043a\u0441\u0435, \u0430 \u0446\u0438\u043a\u043b for \u0441\u0440\u0430\u0431\u043e\u0442\u0430\u0435\u0442 \u0437\u0430 \u043e\u0434\u0438\u043d \u0448\u0430\u0433. \u0424\u0430\u043a\u0442\u0438\u0447\u0435\u0441\u043a\u0438, \u0437\u0434\u0435\u0441\u044c \u043f\u0440\u043e\u0438\u0437\u043e\u0439\u0434\u0443\u0442 \u0441\u043b\u0435\u0434\u0443\u044e\u0449\u0438\u0435 \u043e\u043f\u0435\u0440\u0430\u0446\u0438\u0438  <\/p>\n<ol>\n<li>\u041a\u043e\u043f\u0438\u0440\u043e\u0432\u0430\u043d\u0438\u0435 <code>arr1<\/code>, <code>arr2<\/code> \u0441 \u0426\u041f \u043d\u0430 GPU.<\/li>\n<li>\u0421\u043e\u0437\u0434\u0430\u043d\u0438\u0435 \u043d\u043e\u0432\u043e\u0433\u043e \u043c\u0430\u0441\u0441\u0438\u0432\u0430 \u0440\u0430\u0437\u043c\u0435\u0440\u043e\u043c 3 (\u0438\u043b\u0438 \u0441\u043e\u0445\u0440\u0430\u043d\u0435\u043d\u0438\u0435 \u0440\u0435\u0437\u0443\u043b\u044c\u0442\u0438\u0440\u0443\u044e\u0449\u0435\u0439 \u0441\u0443\u043c\u043c\u044b \u0432 <code>arr1<\/code>).<\/li>\n<li>\u0417\u0430\u043f\u0443\u0441\u043a 3 \u043f\u043e\u0442\u043e\u043a\u043e\u0432 \u0434\u043b\u044f \u0432\u044b\u043f\u043e\u043b\u043d\u0435\u043d\u0438\u044f \u0441\u043b\u043e\u0436\u0435\u043d\u0438\u044f \u0438 \u0441\u043e\u0445\u0440\u0430\u043d\u0435\u043d\u0438\u044f \u0440\u0435\u0437\u0443\u043b\u044c\u0442\u0430\u0442\u0430 \u0432 \u043d\u043e\u0432\u043e\u043c \u043c\u0430\u0441\u0441\u0438\u0432\u0435.<\/li>\n<li>\u041a\u043e\u043f\u0438\u0440\u043e\u0432\u0430\u043d\u0438\u0435 \u0440\u0435\u0437\u0443\u043b\u044c\u0442\u0430\u0442\u0430 \u0441 GPU \u043d\u0430 \u0426\u041f.<\/li>\n<\/ol>\n<p>  \u041f\u043e\u0441\u043a\u043e\u043b\u044c\u043a\u0443 \u0443 GPU \u0442\u044b\u0441\u044f\u0447\u0438 \u044f\u0434\u0435\u0440, \u043d\u0430 \u043d\u0438\u0445 \u0433\u043e\u0440\u0430\u0437\u0434\u043e \u0431\u044b\u0441\u0442\u0440\u0435\u0435, \u0447\u0435\u043c \u043d\u0430 \u043e\u0431\u044b\u0447\u043d\u044b\u0445 \u0426\u041f, \u0432\u044b\u043f\u043e\u043b\u043d\u044f\u044e\u0442\u0441\u044f \u043f\u0440\u043e\u0441\u0442\u044b\u0435 \u0432\u0435\u0449\u0438, \u0432 \u0447\u0430\u0441\u0442\u043d\u043e\u0441\u0442\u0438, \u0441\u043b\u043e\u0436\u0435\u043d\u0438\u0435, \u043f\u0435\u0440\u0435\u043c\u043d\u043e\u0436\u0435\u043d\u0438\u0435 \u043c\u0430\u0442\u0440\u0438\u0446 (\u043f\u0440\u0438 \u0443\u0441\u043b\u043e\u0432\u0438\u0438, \u0447\u0442\u043e \u0432\u044b\u0438\u0433\u0440\u044b\u0448 \u0432 \u0441\u043a\u043e\u0440\u043e\u0441\u0442\u0438 \u043a\u043e\u043c\u043f\u0435\u043d\u0441\u0438\u0440\u0443\u0435\u0442 \u0432\u0440\u0435\u043c\u044f, \u0437\u0430\u0442\u0440\u0430\u0447\u0438\u0432\u0430\u0435\u043c\u043e\u0435 \u043d\u0430 \u043f\u0435\u0440\u0435\u043d\u043e\u0441 \u0434\u0430\u043d\u043d\u044b\u0445 \u043c\u0435\u0436\u0434\u0443 \u0426\u041f \u0438 GPU).<\/p>\n<h2><font color=\"#1c22be\">\u276f<\/font> <font color=\"#454cee\">\u0421\u0440\u0430\u0432\u043d\u0435\u043d\u0438\u0435 CUDA \u0438 Triton<\/font><\/h2>\n<p>  \u0412\u044b\u0448\u0435 \u0431\u044b\u043b\u0430 \u043f\u043e\u043a\u0430\u0437\u0430\u043d\u0430 \u043c\u043e\u0434\u0435\u043b\u044c \u0432\u044b\u043f\u043e\u043b\u043d\u0435\u043d\u0438\u044f CUDA. \u0422\u0435\u043f\u0435\u0440\u044c \u0434\u0430\u0432\u0430\u0439\u0442\u0435 \u043f\u043e\u0441\u043c\u043e\u0442\u0440\u0438\u043c, \u0447\u0435\u043c \u043e\u0442 \u043d\u0435\u0435 \u043e\u0442\u043b\u0438\u0447\u0430\u0435\u0442\u0441\u044f \u043c\u043e\u0434\u0435\u043b\u044c \u0432\u044b\u043f\u043e\u043b\u043d\u0435\u043d\u0438\u044f \u0432 Triton.<\/p>\n<p>  \u0412 CUDA \u043a\u0430\u0436\u0434\u043e\u0435 \u044f\u0434\u0440\u043e \u0430\u0441\u0441\u043e\u0446\u0438\u0438\u0440\u043e\u0432\u0430\u043d\u043e \u0441 <i>\u0431\u043b\u043e\u043a\u043e\u043c \u043f\u043e\u0442\u043e\u043a\u043e\u0432<\/i> (\u0442. e. \u0441 \u043a\u043e\u043b\u043b\u0435\u043a\u0446\u0438\u0435\u0439 \u043f\u043e\u0442\u043e\u043a\u043e\u0432). \u0412 Triton \u043a\u0430\u0436\u0434\u043e\u0435 \u044f\u0434\u0440\u043e \u0430\u0441\u0441\u043e\u0446\u0438\u0438\u0440\u043e\u0432\u0430\u043d\u043e \u0441 \u0435\u0434\u0438\u043d\u0441\u0442\u0432\u0435\u043d\u043d\u044b\u043c \u043f\u043e\u0442\u043e\u043a\u043e\u043c. \u041f\u0440\u0438 \u0442\u0430\u043a\u043e\u0439 \u043f\u0430\u0440\u0430\u0434\u0438\u0433\u043c\u0435 \u0432\u044b\u043f\u043e\u043b\u043d\u0435\u043d\u0438\u044f \u0440\u0435\u0448\u0430\u0435\u0442\u0441\u044f \u043f\u0440\u043e\u0431\u043b\u0435\u043c\u0430 \u0441 \u043c\u0435\u0436\u043f\u043e\u0442\u043e\u0447\u043d\u043e\u0439 \u0441\u0438\u043d\u0445\u0440\u043e\u043d\u0438\u0437\u0430\u0446\u0438\u0435\u0439 \u043f\u0430\u043c\u044f\u0442\u0438, \u0432\u043d\u0443\u0442\u0440\u0438\u043f\u043e\u0442\u043e\u0447\u043d\u043e\u0439 \u043a\u043e\u043c\u043c\u0443\u043d\u0438\u043a\u0430\u0446\u0438\u0435\u0439, \u0438 \u0432 \u0442\u043e \u0436\u0435 \u0432\u0440\u0435\u043c\u044f \u0440\u0430\u0437\u0440\u0435\u0448\u0430\u0435\u0442\u0441\u044f \u0430\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.<\/p>\n<p>  \u0422\u0435\u043f\u0435\u0440\u044c \u043c\u044b \u0443\u0436\u0435 \u043d\u0435 \u0445\u0440\u0430\u043d\u0438\u043c \u043f\u043e\u0442\u043e\u043a\u0438 \u0432 \u0441\u043f\u0435\u0446\u0438\u0430\u043b\u044c\u043d\u043e\u043c \u0431\u043b\u043e\u043a\u0435, \u0430 \u0441\u0430\u043c \u0431\u043b\u043e\u043a \u043f\u0440\u0435\u0434\u0441\u0442\u0430\u0432\u043b\u044f\u0435\u0442 \u0441\u043e\u0431\u043e\u0439 <i>\u0434\u0438\u0430\u043f\u0430\u0437\u043e\u043d<\/i> (range) \u043c\u043e\u0436\u043d\u043e \u0441\u043a\u0430\u0437\u0430\u0442\u044c, \u0437\u0430\u043c\u043e\u0449\u0435\u043d\u043d\u044b\u0439 \u0443\u043a\u0430\u0437\u0430\u0442\u0435\u043b\u044f\u043c\u0438 \u043d\u0430 \u043f\u043e\u0442\u043e\u043a\u0438. \u0412 \u0434\u0430\u043d\u043d\u043e\u043c \u0441\u043b\u0443\u0447\u0430\u0435 \u043d\u0430\u0438\u0431\u043e\u043b\u0435\u0435 \u0438\u043d\u0442\u0435\u0440\u0435\u0441\u043d\u043e, \u0447\u0442\u043e \u0443 \u0432\u0430\u0441 \u043c\u043e\u0436\u0435\u0442 \u0431\u044b\u0442\u044c \u0441\u0442\u043e\u043b\u044c\u043a\u043e \u0434\u0438\u0430\u043f\u0430\u0437\u043e\u043d\u043e\u0432, \u0441\u043a\u043e\u043b\u044c\u043a\u043e \u0437\u0430\u0445\u043e\u0442\u0438\u0442\u0435. \u0422\u0430\u043a, \u0435\u0441\u043b\u0438 \u0432\u0432\u043e\u0434 \u0443 \u0432\u0430\u0441 \u0432 2D, \u043c\u043e\u0436\u043d\u043e \u0443\u043a\u0430\u0437\u0430\u0442\u044c Range(10) \u0434\u043b\u044f \u043e\u0441\u0438 x \u0438 Range(5) \u0434\u043b\u044f \u043e\u0441\u0438 y \u0432\u0441\u0435\u0433\u043e \u0434\u043b\u044f 50 \u043f\u043e\u0442\u043e\u043a\u043e\u0432. \u0410\u043d\u0430\u043b\u043e\u0433\u0438\u0447\u043d\u043e, \u0434\u0438\u0430\u043f\u0430\u0437\u043e\u043d\u044b \u043c\u043e\u0436\u043d\u043e \u043e\u043f\u0440\u0435\u0434\u0435\u043b\u044f\u0442\u044c \u0432 \u0442\u0430\u043a\u043e\u043c \u043a\u043e\u043b\u0438\u0447\u0435\u0441\u0442\u0432\u0435 \u0438\u0437\u043c\u0435\u0440\u0435\u043d\u0438\u0439, \u0432 \u043a\u0430\u043a\u043e\u043c \u0445\u043e\u0442\u0438\u0442\u0435.<\/p>\n<p>  <img decoding=\"async\" src=\"https:\/\/habrastorage.org\/r\/w1560\/webt\/ya\/tg\/vj\/yatgvjxlaljjrjrvgw1dwgkfucq.png\" alt=\"image\" data-src=\"https:\/\/habrastorage.org\/webt\/ya\/tg\/vj\/yatgvjxlaljjrjrvgw1dwgkfucq.png\"\/><br \/>  <b><i><font color=\"#1c22be\">\u0420\u0438\u0441\u0443\u043d\u043e\u043a 3: \u0421\u0440\u0430\u0432\u043d\u0435\u043d\u0438\u0435 \u043c\u043e\u0434\u0435\u043b\u0438 \u0432\u044b\u043f\u043e\u043b\u043d\u0435\u043d\u0438\u044f CUDA \u0438 \u043c\u043e\u0434\u0435\u043b\u0438 \u0432\u044b\u043f\u043e\u043b\u043d\u0435\u043d\u0438\u044f Triton.<\/font><\/i><\/b><\/p>\n<h2><font color=\"#1c22be\">\u276f<\/font> <font color=\"#454cee\">\u0421\u043b\u043e\u0436\u0435\u043d\u0438\u0435 \u0434\u0432\u0443\u0445 \u043c\u0430\u0441\u0441\u0438\u0432\u043e\u0432 \u043f\u0440\u0438 \u043f\u043e\u043c\u043e\u0449\u0438 Triton<\/font><\/h2>\n<p>  \u0418\u0442\u0430\u043a, \u043c\u044b \u0441\u043e\u0441\u0442\u0430\u0432\u0438\u043b\u0438 \u0432\u043f\u0435\u0447\u0430\u0442\u043b\u0435\u043d\u0438\u0435 \u043e \u0442\u043e\u043c, \u043a\u0430\u043a \u0440\u0430\u0431\u043e\u0442\u0430\u044e\u0442 CUDA \u0438 Triton, \u0438 \u043c\u043e\u0436\u0435\u043c \u043f\u0438\u0441\u0430\u0442\u044c \u043f\u0440\u043e\u0433\u0440\u0430\u043c\u043c\u044b \u043d\u0430 Triton. Triton \u0443\u0441\u0442\u0430\u043d\u0430\u0432\u043b\u0438\u0432\u0430\u0435\u0442\u0441\u044f \u0441\u043b\u0435\u0434\u0443\u044e\u0449\u0435\u0439 \u043a\u043e\u043c\u0430\u043d\u0434\u043e\u0439 <code>pip install triton<\/code>.<\/p>\n<p>  \u0412\u043a\u0440\u0430\u0442\u0446\u0435 \u044d\u0442\u0430\u043f\u044b \u0440\u0430\u0431\u043e\u0442\u044b \u0442\u0430\u043a\u043e\u0432\u044b:<\/p>\n<p>  <b>1.<\/b> \u041e\u043f\u0440\u0435\u0434\u0435\u043b\u0438\u0442\u044c <i>\u0431\u043b\u043e\u043a<\/i>. \u041a\u0430\u043a \u0438\u0437\u0432\u0435\u0441\u0442\u043d\u043e, <i>\u0431\u043b\u043e\u043a\u0438<\/i> \u043e\u043f\u0440\u0435\u0434\u0435\u043b\u044f\u044e\u0442\u0441\u044f \u043f\u0443\u0442\u0435\u043c \u0443\u043a\u0430\u0437\u0430\u043d\u0438\u044f \u0434\u0438\u0430\u043f\u0430\u0437\u043e\u043d\u0430. \u041f\u043e\u044d\u0442\u043e\u043c\u0443 \u043f\u0440\u0438 \u0441\u043b\u043e\u0436\u0435\u043d\u0438\u0438 \u043c\u044b \u0434\u043e\u043b\u0436\u043d\u044b \u043f\u0440\u043e\u0441\u0442\u043e \u043e\u043f\u0440\u0435\u0434\u0435\u043b\u0438\u0442\u044c \u0434\u0438\u0430\u043f\u0430\u0437\u043e\u043d \u0432 \u043e\u0434\u043d\u043e\u043c \u0438\u0437\u043c\u0435\u0440\u0435\u043d\u0438\u0438. \u041f\u0443\u0441\u0442\u044c \u043e\u043d \u0431\u0443\u0434\u0435\u0442 \u0440\u0430\u0432\u0435\u043d 512, \u0438 \u043c\u044b \u043e\u043f\u0440\u0435\u0434\u0435\u043b\u0438\u043c \u0435\u0433\u043e \u043a\u0430\u043a \u0433\u043b\u043e\u0431\u0430\u043b\u044c\u043d\u044b\u0439 <code>BLOCK_SIZE = 512<\/code>.<\/p>\n<p>  <b>2.<\/b> \u0424\u0430\u043a\u0442\u0438\u0447\u0435\u0441\u043a\u0438, \u0434\u0438\u0430\u043f\u0430\u0437\u043e\u043d 512 \u043e\u0437\u043d\u0430\u0447\u0430\u0435\u0442, \u0447\u0442\u043e \u0434\u043b\u044f \u0432\u044b\u043f\u043e\u043b\u043d\u0435\u043d\u0438\u044f \u0432\u044b\u0447\u0438\u0441\u043b\u0435\u043d\u0438\u0439 \u043c\u044b \u0437\u0430\u043f\u0443\u0441\u043a\u0430\u0435\u043c 512 \u043f\u043e\u0442\u043e\u043a\u043e\u0432.<\/p>\n<p>  <b>3.<\/b> \u0414\u0430\u043b\u0435\u0435 \u043f\u043e\u043b\u0443\u0447\u0430\u0435\u043c \u0438\u043d\u0434\u0435\u043a\u0441 \u0432\u0445\u043e\u0434\u043d\u044b\u0445 \u0434\u0430\u043d\u043d\u044b\u0445. \u0414\u043e\u043f\u0443\u0441\u0442\u0438\u043c, \u0440\u0430\u0437\u043c\u0435\u0440 \u0432\u0445\u043e\u0434\u043d\u043e\u0433\u043e \u043c\u0430\u0441\u0441\u0438\u0432\u0430 \u0440\u0430\u0432\u0435\u043d 1000. \u041f\u043e\u0441\u043a\u043e\u043b\u044c\u043a\u0443 \u043c\u044b \u043e\u043f\u0440\u0435\u0434\u0435\u043b\u0438\u043b\u0438 \u0431\u043b\u043e\u043a \u0440\u0430\u0437\u043c\u0435\u0440\u043e\u043c 512, \u0432\u0445\u043e\u0434\u043d\u043e\u0439 \u043c\u0430\u0441\u0441\u0438\u0432 \u043c\u044b \u0431\u0443\u0434\u0435\u043c \u043e\u0431\u0440\u0430\u0431\u0430\u0442\u044b\u0432\u0430\u0442\u044c \u0444\u0440\u0430\u0433\u043c\u0435\u043d\u0442\u0430\u043c\u0438 \u043f\u043e 512 \u043a\u0430\u0436\u0434\u044b\u0439. \u0422\u0430\u043a\u0438\u043c \u043e\u0431\u0440\u0430\u0437\u043e\u043c, \u043f\u0435\u0440\u0432\u044b\u0439 \u0444\u0440\u0430\u0433\u043c\u0435\u043d\u0442 \u043d\u0430\u0447\u043d\u0435\u0442\u0441\u044f \u0441 <code>0:512<\/code>, \u0430 \u0432\u0442\u043e\u0440\u043e\u0439 \u0441 <code>512:1024<\/code>. \u042d\u0442\u043e \u0434\u0435\u043b\u0430\u0435\u0442\u0441\u044f \u043f\u0440\u0438 \u043f\u043e\u043c\u043e\u0449\u0438 \u043a\u043e\u0434\u0430, \u043f\u043e\u043a\u0430\u0437\u0430\u043d\u043d\u043e\u0433\u043e \u043d\u0438\u0436\u0435.<\/p>\n<p>  <b>4.<\/b> <\/p>\n<pre><code class=\"python\"># \u0421\u043b\u043e\u0436\u0435\u043d\u0438\u0435 \u043e\u0434\u043d\u043e\u043c\u0435\u0440\u043d\u043e, \u043f\u043e\u044d\u0442\u043e\u043c\u0443 \u043d\u0430\u043c \u043d\u0443\u0436\u043d\u043e \u0442\u043e\u043b\u044c\u043a\u043e \u043f\u043e\u043b\u0443\u0447\u0438\u0442\u044c \u0438\u043d\u0434\u0435\u043a\u0441 \u043f\u043e axis=0 pid = triton.language.program_id(axis=0)<\/code><\/pre>\n<p>   <\/p>\n<pre><code class=\"python\"># \u041f\u043e\u043a\u0430\u0437\u0430\u043d\u043d\u044b\u0435 \u043d\u0438\u0436\u0435 \u0441\u043c\u0435\u0449\u0435\u043d\u0438\u044f \u043e\u0442\u043d\u043e\u0441\u044f\u0442\u0441\u044f \u043a \u0441\u043f\u0438\u0441\u043a\u0430\u043c \u0443\u043a\u0430\u0437\u0430\u0442\u0435\u043b\u0435\u0439 block_start = pid * BLOCK_SIZE offsets = block_start + triton.language.arange(0, BLOCK_SIZE)<\/code><\/pre>\n<p>  <b>5.<\/b> \u041c\u0430\u0441\u043a\u0438\u0440\u043e\u0432\u043a\u0430 \u0434\u043b\u044f \u0437\u0430\u0449\u0438\u0442\u044b \u043e\u043f\u0435\u0440\u0430\u0446\u0438\u0439 \u0432 \u043f\u0430\u043c\u044f\u0442\u0438. \u0412 \u0432\u044b\u0448\u0435\u043f\u0440\u0438\u0432\u0435\u0434\u0435\u043d\u043d\u043e\u043c \u043f\u0440\u0438\u043c\u0435\u0440\u0435 \u0440\u0430\u0437\u043c\u0435\u0440 \u0432\u0445\u043e\u0434\u043d\u043e\u0433\u043e \u043c\u0430\u0441\u0441\u0438\u0432\u0430 <code>N=1000<\/code>, \u043d\u043e, \u0431\u043b\u0430\u0433\u043e\u0434\u0430\u0440\u044f \u0441\u043c\u0435\u0449\u0435\u043d\u0438\u044e, \u043e\u043d \u043d\u0430\u0447\u0438\u043d\u0430\u0435\u0442\u0441\u044f \u0441 <code>512:1024<\/code>. \u041f\u043e\u044d\u0442\u043e\u043c\u0443 \u043c\u044b \u0434\u043e\u043b\u0436\u043d\u044b \u0443\u043a\u0430\u0437\u0430\u0442\u044c \u043c\u0430\u0441\u043a\u0443, \u043a\u043e\u0442\u043e\u0440\u0430\u044f \u0437\u0430\u0449\u0438\u0442\u0438\u0442 \u043d\u0430\u0441 \u043e\u0442 \u0432\u044b\u0445\u043e\u0434\u0430 \u0437\u0430 \u0433\u0440\u0430\u043d\u0438\u0446\u044b \u043c\u0430\u0441\u0441\u0438\u0432\u0430 \u043f\u0440\u0438 \u0437\u0430\u043f\u0440\u043e\u0441\u0430\u0445. \u0422\u0430\u043a\u0443\u044e \u043c\u0430\u0441\u043a\u0443 \u043d\u0443\u0436\u043d\u043e \u0443\u043a\u0430\u0437\u044b\u0432\u0430\u0442\u044c \u0434\u043b\u044f \u043a\u0430\u0436\u0434\u043e\u0439 \u043e\u0441\u0438.<\/p>\n<pre><code class=\"python\">mask = offsets &lt; N<\/code><\/pre>\n<p>  <b>6.<\/b> \u0417\u0430\u0433\u0440\u0443\u0436\u0430\u0435\u043c \u0434\u0430\u043d\u043d\u044b\u0435. \u0422\u0435\u043f\u0435\u0440\u044c, \u043a\u043e\u0433\u0434\u0430 \u0443 \u043d\u0430\u0441 \u043e\u043f\u0440\u0435\u0434\u0435\u043b\u0435\u043d\u044b \u0441\u043c\u0435\u0449\u0435\u043d\u0438\u044f \u0438 \u043c\u0430\u0441\u043a\u0430, \u043c\u043e\u0436\u043d\u043e \u0437\u0430\u0433\u0440\u0443\u0437\u0438\u0442\u044c \u0434\u0430\u043d\u043d\u044b\u0435 \u0438\u0437 RAM \u0438 \u0437\u0430\u043a\u0440\u044b\u0442\u044c \u043c\u0430\u0441\u043a\u0430\u043c\u0438 \u0432\u0441\u0435 \u043b\u0438\u0448\u043d\u0438\u0435 \u044d\u043b\u0435\u043c\u0435\u043d\u0442\u044b.<\/p>\n<pre><code class=\"python\">def add_kernel(arr1_ptr, arr2_ptr, output_ptr, ...): ... arr1 = triton.language.load(arr1_ptr + offsets, mask=mask) arr2 = triton.language.load(arr2_ptr + offsets, mask=mask)<\/code><\/pre>\n<p>  <b>7.<\/b> \u0412\u044b\u043f\u043e\u043b\u043d\u044f\u0435\u043c \u043d\u0443\u0436\u043d\u0443\u044e \u043e\u043f\u0435\u0440\u0430\u0446\u0438\u044e. \u0412 \u0434\u0430\u043d\u043d\u043e\u043c \u0441\u043b\u0443\u0447\u0430\u0435 \u0442\u0440\u0435\u0431\u0443\u0435\u0442\u0441\u044f \u0441\u0434\u0435\u043b\u0430\u0442\u044c \u0442\u043e\u043b\u044c\u043a\u043e \u0441\u043b\u043e\u0436\u0435\u043d\u0438\u0435.<\/p>\n<pre><code class=\"python\">output = arr1 + arr2<\/code><\/pre>\n<p>  <b>8.<\/b> \u041f\u0440\u043e\u0438\u0437\u0432\u0435\u0434\u044f \u0432\u044b\u0447\u0438\u0441\u043b\u0435\u043d\u0438\u044f, \u0441\u043e\u0445\u0440\u0430\u043d\u044f\u0435\u043c \u0440\u0435\u0437\u0443\u043b\u044c\u0442\u0430\u0442 \u0432 RAM. GPU \u043d\u0435 \u0438\u043c\u0435\u0435\u0442 \u0434\u043e\u0441\u0442\u0443\u043f\u0430 \u043a \u0445\u0440\u0430\u043d\u0438\u043b\u0438\u0449\u0443 \u0434\u0430\u043d\u043d\u044b\u0445, \u043f\u043e\u044d\u0442\u043e\u043c\u0443 \u0441\u043d\u0430\u0447\u0430\u043b\u0430 \u043d\u0443\u0436\u043d\u043e \u043f\u0435\u0440\u0435\u043d\u0435\u0441\u0442\u0438 \u0434\u0430\u043d\u043d\u044b\u0435 \u0432 RAM, \u0430 \u0437\u0430\u0442\u0435\u043c \u043c\u044b \u0441\u043c\u043e\u0436\u0435\u043c \u0441\u043e\u0445\u0440\u0430\u043d\u0438\u0442\u044c \u0438\u0445 \u043d\u0430 \u0434\u0438\u0441\u043a\u0435, \u0435\u0441\u043b\u0438 \u0437\u0430\u0445\u043e\u0442\u0438\u043c.<\/p>\n<pre><code class=\"python\">triton.language.store(output_ptr + offsets, output, mask=mask)<\/code><\/pre>\n<p>  \u041d\u0438\u0436\u0435 \u043f\u0440\u0438\u0432\u0435\u0434\u0435\u043d \u0432\u0435\u0441\u044c \u043a\u043e\u0434 \u044f\u0434\u0440\u0430.<\/p>\n<pre><code class=\"python\">import triton import triton.language as tl     BLOCK_SIZE = 512     @triton.jit def add_kernel(arr1_ptr, arr2_ptr, output_ptr, N): # \u0428\u0430\u0433 1: \u043f\u043e\u043b\u0443\u0447\u0430\u0435\u043c \u0432\u0435\u0441\u044c \u0434\u0438\u0430\u043f\u0430\u0437\u043e\u043d \u043e\u0441\u0438 pid = tl.program_id(axis=0)     # \u0428\u0430\u0433 2: \u043e\u043f\u0440\u0435\u0434\u0435\u043b\u044f\u0435\u043c \u0441\u043c\u0435\u0449\u0435\u043d\u0438\u044f \u0438 \u043c\u0430\u0441\u043a\u0443 block_start = pid * BLOCK_SIZE offsets = block_start + tl.arange(0, BLOCK_SIZE) mask = offsets &lt; N     # \u0428\u0430\u0433 3: \u0437\u0430\u0433\u0440\u0443\u0436\u0430\u0435\u043c \u0434\u0430\u043d\u043d\u044b\u0435 \u0438\u0437 RAM arr1 = tl.load(arr1_ptr + offsets, mask=mask) arr2 = tl.load(arr2_ptr + offsets, mask=mask)     # \u0428\u0430\u0433 4: \u0432\u044b\u043f\u043e\u043b\u043d\u044f\u0435\u043c \u0432\u044b\u0447\u0438\u0441\u043b\u0435\u043d\u0438\u044f output = arr1 + arr2     # \u0428\u0430\u0433 5: \u0421\u043e\u0445\u0440\u0430\u043d\u044f\u0435\u043c \u0440\u0435\u0437\u0443\u043b\u044c\u0442\u0430\u0442 \u0432 RAM tl.store(output_ptr + offsets, output, mask=mask)<\/code><\/pre>\n<p>  \u0427\u0442\u043e\u0431\u044b \u0438\u0441\u043f\u043e\u043b\u044c\u0437\u043e\u0432\u0430\u0442\u044c \u044f\u0434\u0440\u043e, \u043e\u043f\u0440\u0435\u0434\u0435\u043b\u044f\u0435\u043c \u0432\u0441\u043f\u043e\u043c\u043e\u0433\u0430\u0442\u0435\u043b\u044c\u043d\u0443\u044e \u0444\u0443\u043d\u043a\u0446\u0438\u044e, \u043a\u0430\u043a \u043f\u043e\u043a\u0430\u0437\u0430\u043d\u043e \u043d\u0438\u0436\u0435<\/p>\n<pre><code class=\"python\">def add(arr1: torch.Tensor, arr2: torch.Tensor): output = torch.empty_like(arr1) N = output.numel()   grid = lambda meta: (triton.cdiv(N, BLOCK_SIZE),)     add_kernel[grid](arr1, arr2, output, N) return output<\/code><\/pre>\n<p>  \u0412 \u043f\u0440\u0438\u043d\u0446\u0438\u043f\u0435, <code>grid<\/code> \u0443\u043a\u0430\u0437\u044b\u0432\u0430\u0435\u0442 \u0442\u043e \u043f\u0440\u043e\u0441\u0442\u0440\u0430\u043d\u0441\u0442\u0432\u043e, \u043d\u0430\u0434 \u043a\u043e\u0442\u043e\u0440\u044b\u043c \u043c\u044b \u0431\u0443\u0434\u0435\u043c \u0440\u0430\u0431\u043e\u0442\u0430\u0442\u044c. \u0412 \u043d\u0430\u0448\u0435\u043c \u0441\u043b\u0443\u0447\u0430\u0435 \u0433\u0440\u0438\u0434 \u043e\u0434\u043d\u043e\u043c\u0435\u0440\u043d\u044b\u0439, \u0438 \u043c\u044b \u0443\u043a\u0430\u0437\u044b\u0432\u0430\u0435\u043c, \u043a\u0430\u043a \u0434\u0430\u043d\u043d\u044b\u0435 \u0431\u0443\u0434\u0443\u0442 \u0440\u0430\u0441\u043f\u0440\u0435\u0434\u0435\u043b\u0435\u043d\u044b \u0432 \u0433\u0440\u0438\u0434\u0435. \u041f\u043e\u044d\u0442\u043e\u043c\u0443, \u0435\u0441\u043b\u0438 \u0432\u0445\u043e\u0434\u043d\u044b\u0435 \u043c\u0430\u0441\u0441\u0438\u0432\u044b \u0438\u043c\u0435\u044e\u0442 \u043d\u0443\u0436\u043d\u044b\u0439 \u0440\u0430\u0437\u043c\u0435\u0440, \u0442\u043e \u043e\u043f\u0440\u0435\u0434\u0435\u043b\u0438\u043c \u0433\u0440\u0438\u0434 \u0441\u043b\u0435\u0434\u0443\u044e\u0449\u0438\u043c \u043e\u0431\u0440\u0430\u0437\u043e\u043c: <code>[0:512], [512:1024]<\/code>. \u041d\u0430 \u0434\u0430\u043d\u043d\u043e\u043c \u0448\u0430\u0433\u0435 \u043c\u044b \u0443\u043a\u0430\u0437\u044b\u0432\u0430\u0435\u043c, \u043a\u0430\u043a \u0440\u0430\u0437\u0434\u0435\u043b\u0438\u0442\u044c \u0432\u0445\u043e\u0434\u043d\u044b\u0435 \u0434\u0430\u043d\u043d\u044b\u0435 \u0438 \u043f\u0435\u0440\u0435\u0434\u0430\u0442\u044c \u0438\u0445 \u044f\u0434\u0440\u0443.<br \/>  \u041f\u043e \u0443\u043c\u043e\u043b\u0447\u0430\u043d\u0438\u044e grid \u043f\u0440\u0438\u043d\u0438\u043c\u0430\u0435\u0442 \u043e\u0434\u0438\u043d \u0430\u0440\u0433\u0443\u043c\u0435\u043d\u0442, \u043e\u0431\u043e\u0437\u043d\u0430\u0447\u0430\u044e\u0449\u0438\u0439 \u043f\u043e\u0437\u0438\u0446\u0438\u044e \u2013 \u043d\u0430\u0437\u043e\u0432\u0435\u043c \u0435\u0433\u043e <code>meta<\/code>. \u0410\u0440\u0433\u0443\u043c\u0435\u043d\u0442 <code>meta<\/code> \u043d\u0443\u0436\u0435\u043d \u0434\u043b\u044f \u043f\u0440\u0435\u0434\u043e\u0441\u0442\u0430\u0432\u043b\u0435\u043d\u0438\u044f \u0442\u0430\u043a\u043e\u0439 \u0438\u043d\u0444\u043e\u0440\u043c\u0430\u0446\u0438\u0438 \u043a\u0430\u043a <code>BLOCK_SIZE<\/code>, \u043d\u043e \u043c\u044b \u043e\u043f\u0440\u0435\u0434\u0435\u043b\u0438\u043b\u0438 \u0440\u0430\u0437\u043c\u0435\u0440 \u0431\u043b\u043e\u043a\u0430 \u0432 \u0433\u043b\u043e\u0431\u0430\u043b\u044c\u043d\u043e\u0439 \u043f\u0435\u0440\u0435\u043c\u0435\u043d\u043d\u043e\u0439.<\/p>\n<p>  \u0422\u0435\u043f\u0435\u0440\u044c \u043c\u044b \u0432\u044b\u0437\u044b\u0432\u0430\u0435\u043c \u0444\u0443\u043d\u043a\u0446\u0438\u044e <code>add<\/code> \u043a\u0430\u043a \u043e\u0431\u044b\u0447\u043d\u0443\u044e Python-\u0444\u0443\u043d\u043a\u0446\u0438\u044e, \u0441\u043c. \u043d\u0438\u0436\u0435. \u041f\u0435\u0440\u0435\u0434 \u044d\u0442\u0438\u043c \u043d\u0443\u0436\u043d\u043e \u0443\u0431\u0435\u0434\u0438\u0442\u044c\u0441\u044f, \u0447\u0442\u043e \u0434\u0430\u043d\u043d\u044b\u0435, \u043f\u0435\u0440\u0435\u0434\u0430\u0432\u0430\u0435\u043c\u044b\u0435 \u0444\u0443\u043d\u043a\u0446\u0438\u0438 \u043d\u0430 \u0432\u0445\u043e\u0434, \u0443\u0436\u0435 \u043d\u0430\u0445\u043e\u0434\u044f\u0442\u0441\u044f \u043d\u0430 GPU.<\/p>\n<pre><code class=\"python\">arr_size = 100_000 arr1 = torch.rand(arr_size, device='cuda') arr2 = torch.rand(arr_size, device='cuda')     pytorch_out = arr1 + arr2 triton_out = add(arr1, arr2)     print(torch.sum(torch.abs(pytorch_out - triton_out)))<\/code><\/pre>\n<p>  \u0412\u044b\u0432\u043e\u0434<\/p>\n<p>  <code>\u276f python main.py<br \/>  tensor(0., device='cuda:0')<\/code><\/p>\n<h2><font color=\"#1c22be\">\u276f<\/font> <font color=\"#454cee\">\u0421\u043b\u043e\u0436\u0435\u043d\u0438\u0435 \u0434\u043b\u044f \u0442\u0435\u043d\u0437\u043e\u0440\u043e\u0432 \u0431\u043e\u043b\u0435\u0435 \u0432\u044b\u0441\u043e\u043a\u0438\u0445 \u0438\u0437\u043c\u0435\u0440\u0435\u043d\u0438\u0439<\/font><\/h2>\n<p>  \u0422\u0435\u043c \u0436\u0435 \u0441\u0430\u043c\u044b\u043c \u044f\u0434\u0440\u043e\u043c \u043c\u043e\u0436\u043d\u043e \u0432\u043e\u0441\u043f\u043e\u043b\u044c\u0437\u043e\u0432\u0430\u0442\u044c\u0441\u044f \u0438 \u0434\u043b\u044f \u0440\u0430\u0431\u043e\u0442\u044b \u0441 N-\u043c\u0435\u0440\u043d\u044b\u043c\u0438 \u0442\u0435\u043d\u0437\u043e\u0440\u0430\u043c\u0438. \u042d\u0442\u043e \u0433\u0438\u0431\u043a\u0438\u0439 \u043f\u043e\u0434\u0445\u043e\u0434, \u043f\u043e\u0437\u0432\u043e\u043b\u044f\u044e\u0449\u0438\u0439 \u043e\u0431\u043e\u0439\u0442\u0438\u0441\u044c \u0431\u0435\u0437 \u043d\u0430\u043f\u0438\u0441\u0430\u043d\u0438\u044f \u043c\u043d\u043e\u0436\u0435\u0441\u0442\u0432\u0430 \u044f\u0434\u0435\u0440, \u043a\u0430\u0436\u0434\u043e\u0435 \u0438\u0437 \u043a\u043e\u0442\u043e\u0440\u044b\u0445 \u043e\u0442\u0432\u0435\u0447\u0430\u043b\u043e \u0431\u044b \u0437\u0430 \u0432\u0432\u043e\u0434 \u0441 \u043a\u043e\u043d\u043a\u0440\u0435\u0442\u043d\u044b\u043c \u043a\u043e\u043b\u0438\u0447\u0435\u0441\u0442\u0432\u043e\u043c \u0438\u0437\u043c\u0435\u0440\u0435\u043d\u0438\u0439. \u0418\u0434\u0435\u044f \u043f\u0440\u043e\u0441\u0442\u0430: \u0432\u043e \u0432\u0441\u043f\u043e\u043c\u043e\u0433\u0430\u0442\u0435\u043b\u044c\u043d\u043e\u0439 \u0444\u0443\u043d\u043a\u0446\u0438\u0438 \u043f\u0435\u0440\u0435\u0444\u043e\u0440\u043c\u0438\u0440\u0443\u0435\u043c \u0438\u043c\u0435\u044e\u0449\u0438\u0439\u0441\u044f \u0442\u0435\u043d\u0437\u043e\u0440 \u0432 \u043e\u0434\u043d\u043e\u043c\u0435\u0440\u043d\u044b\u0439, \u0430 \u0437\u0430\u0442\u0435\u043c \u043f\u0435\u0440\u0435\u0444\u043e\u0440\u043c\u0438\u0440\u0443\u0435\u043c \u0432\u044b\u0445\u043e\u0434\u043d\u043e\u0439 \u0442\u0435\u043d\u0437\u043e\u0440.<\/p>\n<p>  \u0422\u0430\u043a\u0430\u044f \u043e\u043f\u0435\u0440\u0430\u0446\u0438\u044f \u043f\u043e \u043f\u0435\u0440\u0435\u0444\u043e\u0440\u043c\u0438\u0440\u043e\u0432\u0430\u043d\u0438\u044e \u043d\u0435 \u043f\u043e\u0442\u0440\u0435\u0431\u0443\u0435\u0442 \u043c\u043d\u043e\u0433\u043e \u0432\u0440\u0435\u043c\u0435\u043d\u0438, \u0442\u0430\u043a \u043a\u0430\u043a \u043c\u044b \u0432\u0441\u0435\u0433\u043e \u043b\u0438\u0448\u044c \u0438\u0437\u043c\u0435\u043d\u044f\u0435\u043c \u0437\u043d\u0430\u0447\u0435\u043d\u0438\u044f \u0448\u0430\u0433\u043e\u0432 \u0432 \u043a\u043b\u0430\u0441\u0441\u0435 tensor. \u041d\u0438\u0436\u0435 \u043f\u043e\u043a\u0430\u0437\u0430\u043d\u0430 \u043c\u043e\u0434\u0438\u0444\u0438\u0446\u0438\u0440\u043e\u0432\u0430\u043d\u043d\u0430\u044f \u0432\u0441\u043f\u043e\u043c\u043e\u0433\u0430\u0442\u0435\u043b\u044c\u043d\u0430\u044f \u0444\u0443\u043d\u043a\u0446\u0438\u044f.<\/p>\n<pre><code class=\"python\">def add(arr1: torch.Tensor, arr2: torch.Tensor): input_shape = arr1.shape arr1 = arr1.view(-1) arr2 = arr2.view(-1)     output = torch.empty_like(arr1) N = output.numel()     grid = lambda meta: (triton.cdiv(N, BLOCK_SIZE),)     add_kernel[grid](arr1, arr2, output, N)     output = output.view(input_shape) return output<\/code><\/pre>\n<p>  \u0410 \u0437\u0430\u0442\u0435\u043c \u0432\u044b\u0437\u044b\u0432\u0430\u0435\u043c \u0444\u0443\u043d\u043a\u0446\u0438\u044e \u0442\u043e\u0447\u043d\u043e \u0442\u0430\u043a \u0436\u0435, \u043a\u0430\u043a \u0438 \u0440\u0430\u043d\u044c\u0448\u0435.<\/p>\n<pre><code class=\"python\">arr_size = (100, 100, 100) arr1 = torch.rand(arr_size, device='cuda') arr2 = torch.rand(arr_size, device='cuda')     pytorch_out = arr1 + arr2 triton_out = add(arr1, arr2)     print(torch.sum(torch.abs(pytorch_out - triton_out)))<\/code><\/pre>\n<p>  \u0412\u044b\u0432\u043e\u0434<\/p>\n<pre><code class=\"python\">\u276f python main.py tensor(0., device='cuda:0')<\/code><\/pre>\n<p>  \u041c\u044b \u0440\u0430\u0437\u043e\u0431\u0440\u0430\u043b\u0438 \u043f\u0440\u043e\u0441\u0442\u043e\u0439 \u043f\u0440\u0438\u043c\u0435\u0440. \u041d\u043e \u0432 \u0431\u043e\u043b\u0435\u0435 \u0441\u043b\u043e\u0436\u043d\u044b\u0445 \u0441\u043b\u0443\u0447\u0430\u044f\u0445, \u0441\u043a\u0430\u0436\u0435\u043c, \u043f\u0440\u0438 \u043f\u0435\u0440\u0435\u043c\u043d\u043e\u0436\u0435\u043d\u0438\u0438 \u043c\u0430\u0442\u0440\u0438\u0446, \u043f\u0440\u043e\u0438\u0437\u0432\u043e\u0434\u0438\u0442\u0435\u043b\u044c\u043d\u043e\u0441\u0442\u044c \u0441\u0438\u043b\u044c\u043d\u0435\u0439\u0448\u0438\u043c \u043e\u0431\u0440\u0430\u0437\u043e\u043c \u0437\u0430\u0432\u0438\u0441\u0438\u0442 \u043e\u0442 \u0442\u043e\u0433\u043e, \u043a\u0430\u043a \u0438\u043c\u0435\u043d\u043d\u043e \u0432\u044b \u0440\u0430\u0437\u0434\u0435\u043b\u0438\u0442\u0435 \u0434\u0430\u043d\u043d\u044b\u0435. \u0412 \u0434\u043e\u043a\u0443\u043c\u0435\u043d\u0442\u0430\u0446\u0438\u0438 \u043f\u043e Triton \u0435\u0441\u0442\u044c <a href=\"https:\/\/triton-lang.org\/getting-started\/tutorials\/03-matrix-multiplication.html\">\u0440\u0443\u043a\u043e\u0432\u043e\u0434\u0441\u0442\u0432\u043e<\/a>, \u043f\u043e\u044f\u0441\u043d\u044f\u044e\u0449\u0435\u0435, \u043a\u0430\u043a \u043d\u0430\u043f\u0438\u0441\u0430\u0442\u044c \u044f\u0434\u0440\u043e \u0434\u043b\u044f \u044d\u0444\u0444\u0435\u043a\u0442\u0438\u0432\u043d\u043e\u0433\u043e \u043f\u0435\u0440\u0435\u043c\u043d\u043e\u0436\u0435\u043d\u0438\u044f \u043c\u0430\u0442\u0440\u0438\u0446 \u2013 \u0442\u0430\u043c \u0432\u0441\u0435 \u0440\u0430\u0441\u0441\u043a\u0430\u0437\u0430\u043d\u043e \u0432 \u043f\u043e\u0434\u0440\u043e\u0431\u043d\u043e\u0441\u0442\u044f\u0445.<\/p>\n<p>  \u042d\u0442\u043e \u0431\u044b\u043b\u043e \u043a\u0440\u0430\u0442\u043a\u043e\u0435 \u0440\u0443\u043a\u043e\u0432\u043e\u0434\u0441\u0442\u0432\u043e, \u0431\u043b\u0430\u0433\u043e\u0434\u0430\u0440\u044f \u043a\u043e\u0442\u043e\u0440\u043e\u043c\u0443 \u0432\u044b \u043f\u043e\u0437\u043d\u0430\u043a\u043e\u043c\u0438\u043b\u0438\u0441\u044c \u0441 \u043e\u0441\u043d\u043e\u0432\u0430\u043c\u0438 \u043f\u0440\u043e\u0433\u0440\u0430\u043c\u043c\u0438\u0440\u043e\u0432\u0430\u043d\u0438\u044f \u0434\u043b\u044f GPU \u0438 \u0441 Triton. \u041f\u043e\u0434\u0440\u043e\u0431\u043d\u0435\u0435 \u043f\u043e\u0447\u0438\u0442\u0430\u0442\u044c \u043e \u043f\u0440\u043e\u0435\u043a\u0442\u0435 Triton \u043c\u043e\u0436\u043d\u043e \u0432 <a href=\"https:\/\/github.com\/openai\/triton\">openai\/github<\/a>.<\/p>\n<p>  <a href=\"https:\/\/timeweb.cloud\/\/vds-vps?utm_source=habr&amp;utm_medium=banner&amp;utm_campaign=vds-promo-6-rub\"><img decoding=\"async\" src=\"https:\/\/habrastorage.org\/r\/w1560\/webt\/p-\/u9\/l2\/p-u9l27ynelxi92bcmdxhu76ma8.png\" data-src=\"https:\/\/habrastorage.org\/webt\/p-\/u9\/l2\/p-u9l27ynelxi92bcmdxhu76ma8.png\"\/><\/a><\/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\/company\/timeweb\/blog\/702298\/\"> https:\/\/habr.com\/ru\/company\/timeweb\/blog\/702298\/<\/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-1\">\n<div xmlns=\"http:\/\/www.w3.org\/1999\/xhtml\"><img decoding=\"async\" src=\"https:\/\/habrastorage.org\/r\/w1560\/webt\/gt\/au\/sk\/gtausku8oyxsy-a60nmlxenbypq.png\" alt=\"image\" data-src=\"https:\/\/habrastorage.org\/webt\/gt\/au\/sk\/gtausku8oyxsy-a60nmlxenbypq.png\"\/><\/p>\n<p>  Triton \u2013 \u044d\u0442\u043e \u044f\u0437\u044b\u043a\u043e\u0432\u043e\u0439 \u043a\u043e\u043c\u043f\u0438\u043b\u044f\u0442\u043e\u0440 \u0434\u043b\u044f \u0441\u043e\u0437\u0434\u0430\u043d\u0438\u044f \u0441\u0438\u043b\u044c\u043d\u043e \u043e\u043f\u0442\u0438\u043c\u0438\u0437\u0438\u0440\u043e\u0432\u0430\u043d\u043d\u044b\u0445 \u044f\u0434\u0435\u0440 CUDA. \u0417\u0434\u0435\u0441\u044c \u0431\u0443\u0434\u0443\u0442 \u0438\u0437\u043b\u043e\u0436\u0435\u043d\u044b \u043e\u0441\u043d\u043e\u0432\u044b \u043f\u0440\u043e\u0433\u0440\u0430\u043c\u043c\u0438\u0440\u043e\u0432\u0430\u043d\u0438\u044f \u0434\u043b\u044f GPU \u0438 \u0440\u0430\u0441\u0441\u043a\u0430\u0437\u0430\u043d\u043e, \u043a\u0430\u043a \u0434\u043b\u044f \u044d\u0442\u043e\u0439 \u0446\u0435\u043b\u0438 \u0438\u0441\u043f\u043e\u043b\u044c\u0437\u0443\u0435\u0442\u0441\u044f Triton.<\/p>\n<p>  \u0423\u0447\u0438\u0442\u044b\u0432\u0430\u044f \u043d\u044b\u043d\u0435\u0448\u043d\u0438\u0439 \u0443\u0441\u043f\u0435\u0445 \u0433\u043b\u0443\u0431\u043e\u043a\u043e\u0433\u043e \u043e\u0431\u0443\u0447\u0435\u043d\u0438\u044f \u0438 \u0432\u0430\u043b \u0438\u0441\u0441\u043b\u0435\u0434\u043e\u0432\u0430\u0442\u0435\u043b\u044c\u0441\u043a\u0438\u0445 \u0441\u0442\u0430\u0442\u0435\u0439 \u043d\u0430 \u044d\u0442\u0443 \u0442\u0435\u043c\u0443, \u0447\u0430\u0441\u0442\u043e \u0432\u043e\u0437\u043d\u0438\u043a\u0430\u0435\u0442 \u0442\u0430\u043a\u0430\u044f \u0441\u0438\u0442\u0443\u0430\u0446\u0438\u044f: \u0440\u043e\u0436\u0434\u0430\u0435\u0442\u0441\u044f \u043a\u0430\u043a\u0430\u044f-\u043d\u0438\u0431\u0443\u0434\u044c \u043d\u043e\u0432\u0430\u044f \u0438\u0434\u0435\u044f, \u0438 \u0432\u044b\u044f\u0441\u043d\u044f\u0435\u0442\u0441\u044f, \u0447\u0442\u043e \u0434\u043b\u044f \u043d\u0435\u0435 \u043d\u0435 \u043f\u043e\u0434\u0434\u0435\u0440\u0436\u0438\u0432\u0430\u0435\u0442\u0441\u044f \u0430\u043f\u043f\u0430\u0440\u0430\u0442\u043d\u043e\u0435 \u0443\u0441\u043a\u043e\u0440\u0435\u043d\u0438\u0435. \u0422\u043e\u0447\u043d\u0435\u0435, \u0441\u0442\u043e\u0438\u0442 \u0432\u0430\u043c \u0438\u0437\u043e\u0431\u0440\u0435\u0441\u0442\u0438 \u043d\u043e\u0432\u0443\u044e \u0444\u0443\u043d\u043a\u0446\u0438\u044e \u0430\u043a\u0442\u0438\u0432\u0430\u0446\u0438\u0438 \u0438\u043b\u0438 \u043c\u0435\u0445\u0430\u043d\u0438\u0437\u043c \u0441\u0430\u043c\u043e\u0432\u043d\u0438\u043c\u0430\u043d\u0438\u044f \u2013 \u043d\u0430\u043c \u0441\u0440\u0430\u0437\u0443 \u043f\u0440\u0438\u0445\u043e\u0434\u0438\u0442\u0441\u044f \u043f\u0440\u0438\u0431\u0435\u0433\u0430\u0442\u044c \u043a \u0432\u043e\u0437\u043c\u043e\u0436\u043d\u043e\u0441\u0442\u044f\u043c PyTorch\/Tensorflow \u0434\u043b\u044f \u043e\u0431\u0440\u0430\u0431\u043e\u0442\u043a\u0438 \u043f\u0440\u044f\u043c\u043e\u0433\u043e \u0438 \u043e\u0431\u0440\u0430\u0442\u043d\u043e\u0433\u043e \u043f\u0440\u043e\u0445\u043e\u0434\u0430 \u0447\u0435\u0440\u0435\u0437 \u043c\u043e\u0434\u0443\u043b\u044c.<\/p>\n<p>  \u0412 \u0442\u0430\u043a\u0438\u0445 \u0441\u043b\u0443\u0447\u0430\u044f\u0445 \u043f\u0440\u0438\u043c\u0435\u043d\u0438\u043c, \u043d\u0430\u043f\u0440\u0438\u043c\u0435\u0440, PyTorch JIT. \u041d\u043e PyTorch JIT \u2013 \u044d\u0442\u043e \u0432\u044b\u0441\u043e\u043a\u043e\u0443\u0440\u043e\u0432\u043d\u0435\u0432\u044b\u0439 \u043a\u043e\u043c\u043f\u0438\u043b\u044f\u0442\u043e\u0440, \u0441\u043f\u043e\u0441\u043e\u0431\u043d\u044b\u0439 \u043e\u043f\u0442\u0438\u043c\u0438\u0437\u0438\u0440\u043e\u0432\u0430\u0442\u044c \u043b\u0438\u0448\u044c \u043d\u0435\u043a\u043e\u0442\u043e\u0440\u044b\u0435 \u0447\u0430\u0441\u0442\u0438 \u043a\u043e\u0434\u0430, \u043d\u043e \u043d\u0435\u043f\u0440\u0438\u0433\u043e\u0434\u043d\u044b\u0439 \u0434\u043b\u044f \u043d\u0430\u043f\u0438\u0441\u0430\u043d\u0438\u044f \u0441\u043f\u0435\u0446\u0438\u0430\u043b\u0438\u0437\u0438\u0440\u043e\u0432\u0430\u043d\u043d\u044b\u0445 \u044f\u0434\u0435\u0440 CUDA.  <\/p>\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-341925","post","type-post","status-publish","format-standard","hentry"],"_links":{"self":[{"href":"https:\/\/savepearlharbor.com\/index.php?rest_route=\/wp\/v2\/posts\/341925","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=341925"}],"version-history":[{"count":0,"href":"https:\/\/savepearlharbor.com\/index.php?rest_route=\/wp\/v2\/posts\/341925\/revisions"}],"wp:attachment":[{"href":"https:\/\/savepearlharbor.com\/index.php?rest_route=%2Fwp%2Fv2%2Fmedia&parent=341925"}],"wp:term":[{"taxonomy":"category","embeddable":true,"href":"https:\/\/savepearlharbor.com\/index.php?rest_route=%2Fwp%2Fv2%2Fcategories&post=341925"},{"taxonomy":"post_tag","embeddable":true,"href":"https:\/\/savepearlharbor.com\/index.php?rest_route=%2Fwp%2Fv2%2Ftags&post=341925"}],"curies":[{"name":"wp","href":"https:\/\/api.w.org\/{rel}","templated":true}]}}