{"id":373625,"date":"2024-05-21T05:40:33","date_gmt":"2024-05-21T05:40:33","guid":{"rendered":"http:\/\/savepearlharbor.com\/?p=373625"},"modified":"-0001-11-30T00:00:00","modified_gmt":"-0001-11-29T21:00:00","slug":"","status":"publish","type":"post","link":"https:\/\/savepearlharbor.com\/?p=373625","title":{"rendered":"<span>\u041a\u0440\u0430\u0442\u043a\u043e \u043f\u0440\u043e \u044f\u0437\u044b\u043a \u043f\u0440\u043e\u0433\u0440\u0430\u043c\u043c\u0438\u0440\u043e\u0432\u0430\u043d\u0438\u044f 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-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\/w1560\/getpro\/habr\/upload_files\/587\/60d\/e7b\/58760de7b81309b0fa392a42a0a99349.png\" width=\"1920\" height=\"1080\" data-src=\"https:\/\/habrastorage.org\/getpro\/habr\/upload_files\/587\/60d\/e7b\/58760de7b81309b0fa392a42a0a99349.png\"\/><\/figure>\n<p>\u041f\u0440\u0438\u0432\u0435\u0442, \u0425\u0430\u0431\u0440!<\/p>\n<p><strong>Triton<\/strong> \u0431\u044b\u043b \u0440\u0430\u0437\u0440\u0430\u0431\u043e\u0442\u0430\u043d \u0441\u043f\u0435\u0446\u0438\u0430\u043b\u044c\u043d\u043e \u0434\u043b\u044f \u0432\u044b\u043f\u043e\u043b\u043d\u0435\u043d\u0438\u044f \u043d\u0430 GPU \u0438 \u043f\u0440\u0435\u0434\u043e\u0441\u0442\u0430\u0432\u043b\u044f\u0435\u0442 \u0443\u0434\u043e\u0431\u043d\u0443\u044e Python-\u043e\u0440\u0438\u0435\u043d\u0442\u0438\u0440\u043e\u0432\u0430\u043d\u043d\u0443\u044e \u0441\u0440\u0435\u0434\u0443.<\/p>\n<p>Triton \u043f\u043e\u0437\u0432\u043e\u043b\u044f\u0435\u0442 \u0438\u0441\u043f\u043e\u043b\u044c\u0437\u043e\u0432\u0430\u0442\u044c \u043c\u043e\u0434\u0435\u043b\u044c \u043f\u0440\u043e\u0433\u0440\u0430\u043c\u043c\u0438\u0440\u043e\u0432\u0430\u043d\u0438\u044f, <strong>\u043e\u0441\u043d\u043e\u0432\u0430\u043d\u043d\u0443\u044e \u043d\u0430 \u0431\u043b\u043e\u043a\u0430\u0445,<\/strong> \u043a\u043e\u0442\u043e\u0440\u0430\u044f \u0437\u043d\u0430\u0447\u0438\u0442\u0435\u043b\u044c\u043d\u043e \u043e\u0442\u043b\u0438\u0447\u0430\u0435\u0442\u0441\u044f \u043e\u0442 \u0442\u0440\u0430\u0434\u0438\u0446\u0438\u043e\u043d\u043d\u043e\u0439 \u043c\u043e\u0434\u0435\u043b\u0438 CUDA. \u0412\u043c\u0435\u0441\u0442\u043e \u0443\u043f\u0440\u0430\u0432\u043b\u0435\u043d\u0438\u044f \u043f\u043e\u0442\u043e\u043a\u0430\u043c\u0438 \u043d\u0430 \u0443\u0440\u043e\u0432\u043d\u0435 \u0441\u043a\u0430\u043b\u044f\u0440\u043d\u044b\u0445 \u0438\u043d\u0441\u0442\u0440\u0443\u043a\u0446\u0438\u0439, <strong>Triton \u043e\u043f\u0435\u0440\u0438\u0440\u0443\u0435\u0442 \u0431\u043b\u043e\u043a\u0430\u043c\u0438 \u0434\u0430\u043d\u043d\u044b\u0445, \u0447\u0442\u043e \u0432 \u0446\u0435\u043b\u043e\u043c \u0434\u0430\u0435\u0442 \u0431\u043e\u043b\u0435\u0435 \u043b\u0443\u0447\u0448\u0443\u044e \u043f\u0440\u043e\u0438\u0437\u0432\u043e\u0434\u0438\u0442\u0435\u043b\u044c\u043d\u043e\u0441\u0442\u044c.<\/strong><\/p>\n<p>\u0412 \u043e\u0442\u043b\u0438\u0447\u0438\u0435 \u043e\u0442 \u0441\u0442\u0430\u043d\u0434\u0430\u0440\u0442\u043d\u043e\u0433\u043e \u043f\u043e\u0434\u0445\u043e\u0434\u0430 CUDA, \u0433\u0434\u0435 \u0438\u0441\u043f\u043e\u043b\u043d\u0435\u043d\u0438\u0435 \u043a\u043e\u0434\u0430 \u043e\u0440\u0433\u0430\u043d\u0438\u0437\u0443\u0435\u0442\u0441\u044f \u0447\u0435\u0440\u0435\u0437 \u0432\u0437\u0430\u0438\u043c\u043e\u0434\u0435\u0439\u0441\u0442\u0432\u0438\u0435 \u043c\u043d\u043e\u0436\u0435\u0441\u0442\u0432\u0430 \u043f\u043e\u0442\u043e\u043a\u043e\u0432, Triton \u0441\u0442\u0440\u0443\u043a\u0442\u0443\u0440\u0438\u0440\u0443\u0435\u0442 \u0432\u044b\u043f\u043e\u043b\u043d\u0435\u043d\u0438\u0435 \u043d\u0430 \u0443\u0440\u043e\u0432\u043d\u0435 \u043f\u0440\u043e\u0433\u0440\u0430\u043c\u043c. <strong>\u0422.\u0435 \u043a\u0430\u0436\u0434\u044b\u0439 \u0431\u043b\u043e\u043a \u043f\u0440\u043e\u0433\u0440\u0430\u043c\u043c\u044b \u043c\u043e\u0436\u0435\u0442 \u0431\u044b\u0442\u044c \u0438\u0441\u043f\u043e\u043b\u043d\u0435\u043d \u043d\u0435\u0437\u0430\u0432\u0438\u0441\u0438\u043c\u043e, \u0441 \u0432\u043e\u0437\u043c\u043e\u0436\u043d\u043e\u0441\u0442\u044c\u044e \u043e\u0431\u0440\u0430\u0449\u0435\u043d\u0438\u044f \u043a \u0433\u043b\u043e\u0431\u0430\u043b\u044c\u043d\u043e\u0439 \u043f\u0430\u043c\u044f\u0442\u0438 GPU<\/strong> \u0438 \u0432\u044b\u043f\u043e\u043b\u043d\u0435\u043d\u0438\u044f \u0430\u0441\u0438\u043d\u0445\u0440\u043e\u043d\u043d\u044b\u0445 \u043e\u043f\u0435\u0440\u0430\u0446\u0438\u0439 \u0431\u0435\u0437 \u044f\u0432\u043d\u043e\u0433\u043e \u0443\u043f\u0440\u0430\u0432\u043b\u0435\u043d\u0438\u044f \u0441\u0438\u043d\u0445\u0440\u043e\u043d\u0438\u0437\u0430\u0446\u0438\u0435\u0439 \u043f\u043e\u0442\u043e\u043a\u043e\u0432.<\/p>\n<p>\u041a\u043e\u043c\u043f\u0438\u043b\u044f\u0442\u043e\u0440 Triton \u043f\u0440\u0438\u043c\u0435\u043d\u044f\u0435\u0442 \u0441\u043b\u043e\u0436\u043d\u044b\u0435 \u0441\u0442\u0440\u0430\u0442\u0435\u0433\u0438\u0438 \u043e\u043f\u0442\u0438\u043c\u0438\u0437\u0430\u0446\u0438\u0438, \u043a \u043f\u0440\u0438\u043c\u0435\u0440\u0443 \u0442\u0430\u043a\u0438\u0435 \u043a\u0430\u043a \u0430\u043d\u0430\u043b\u0438\u0437 \u043f\u043e\u0442\u043e\u043a\u0430 \u0434\u0430\u043d\u043d\u044b\u0445 \u0438 \u0443\u043f\u0440\u0430\u0432\u043b\u0435\u043d\u0438\u0435 \u043f\u0430\u043c\u044f\u0442\u044c\u044e \u043d\u0430 \u0443\u0440\u043e\u0432\u043d\u0435 \u0431\u043b\u043e\u043a\u043e\u0432. \u042d\u0442\u043e \u0432\u043a\u043b\u044e\u0447\u0430\u0435\u0442 \u0432 \u0441\u0435\u0431\u044f: <em>\u0430\u0432\u0442\u043e\u043c\u0430\u0442\u0438\u0447\u0435\u0441\u043a\u0430\u044f \u0432\u0435\u043a\u0442\u043e\u0440\u0438\u0437\u0430\u0446\u0438\u044f<\/em>, <em>\u043f\u0440\u0435\u0434\u0432\u0430\u0440\u0438\u0442\u0435\u043b\u044c\u043d\u0430\u044f \u0432\u044b\u0431\u043e\u0440\u043a\u0430 \u0434\u0430\u043d\u043d\u044b\u0445,<\/em> \u0438 <em>\u0438\u0441\u043f\u043e\u043b\u044c\u0437\u043e\u0432\u0430\u043d\u0438\u0435 \u0442\u0435\u043d\u0437\u043e\u0440\u043d\u044b\u0445 \u044f\u0434\u0435\u0440<\/em>, \u0433\u0434\u0435 \u044d\u0442\u043e \u0432\u043e\u0437\u043c\u043e\u0436\u043d\u043e. \u0422\u0430\u043a\u043e\u0439 \u043f\u043e\u0434\u0445\u043e\u0434 \u043f\u043e\u0437\u0432\u043e\u043b\u044f\u0435\u0442 \u043c\u0430\u043a\u0441\u0438\u043c\u0430\u043b\u044c\u043d\u043e \u0438\u0441\u043f\u043e\u043b\u044c\u0437\u043e\u0432\u0430\u0442\u044c \u0432\u043e\u0437\u043c\u043e\u0436\u043d\u043e\u0441\u0442\u0438 GPU.<\/p>\n<h4>\u0423\u0441\u0442\u0430\u043d\u043e\u0432\u0438\u043c<\/h4>\n<p>\u0421\u0430\u043c\u044b\u0439 \u043f\u0440\u043e\u0441\u0442\u043e\u0439 \u0438 \u0434\u043e\u0441\u0442\u0443\u043f\u043d\u044b\u0439 \u0441\u043f\u043e\u0441\u043e\u0431 \u0443\u0441\u0442\u0430\u043d\u043e\u0432\u0438\u0442\u044c Triton \u2014 <strong>\u044d\u0442\u043e \u0447\u0435\u0440\u0435\u0437 pip:<\/strong><\/p>\n<pre><code>pip install triton<\/code><\/pre>\n<p>Triton <strong>\u043c\u043e\u0436\u043d\u043e \u0442\u0430\u043a\u0436\u0435 \u0441\u043e\u0431\u0440\u0430\u0442\u044c \u0438\u0437 \u0438\u0441\u0445\u043e\u0434\u043d\u0438\u043a\u043e\u0432.<\/strong> \u0414\u043b\u044f \u044d\u0442\u043e\u0433\u043e \u043a\u043b\u043e\u043d\u0438\u0440\u0443\u0435\u043c \u0440\u0435\u043f\u043e\u0437\u0438\u0442\u043e\u0440\u0438\u0439 \u0438 \u0443\u0441\u0442\u0430\u043d\u0430\u0432\u043b\u0438\u0432\u0430\u0435\u043c \u043d\u0435\u043e\u0431\u0445\u043e\u0434\u0438\u043c\u044b\u0435 \u0437\u0430\u0432\u0438\u0441\u0438\u043c\u043e\u0441\u0442\u0438:<\/p>\n<pre><code>git clone https:\/\/github.com\/openai\/triton.git cd triton\/python pip install ninja cmake wheel pip install -e .<\/code><\/pre>\n<p>\u0423\u0447\u0442\u0438\u0442\u0435, \u0447\u0442\u043e \u0435\u0441\u043b\u0438 \u043d\u0430 \u0441\u0438\u0441\u0442\u0435\u043c\u0435 \u043d\u0435 \u0443\u0441\u0442\u0430\u043d\u043e\u0432\u043b\u0435\u043d llvm, \u0441\u043a\u0440\u0438\u043f\u0442 <a href=\"http:\/\/setup.py\">setup.py<\/a> \u0441\u043a\u0430\u0447\u0430\u0435\u0442 \u043e\u0444\u0438\u0446\u0438\u0430\u043b\u044c\u043d\u044b\u0435 \u0441\u0442\u0430\u0442\u0438\u0447\u0435\u0441\u043a\u0438\u0435 \u0431\u0438\u0431\u043b\u0438\u043e\u0442\u0435\u043a\u0438 LLVM \u0438 \u0441\u0432\u044f\u0436\u0435\u0442 \u0438\u0445.<\/p>\n<h4>\u0420\u0430\u0441\u0441\u043c\u043e\u0442\u0440\u0438\u043c \u0432\u0435\u0441\u044c \u043e\u0441\u043d\u043e\u0432\u043d\u043e\u0439 \u0441\u0438\u043d\u0442\u0430\u043a\u0441\u0438\u0441<\/h4>\n<p>Triton \u0438\u0441\u043f\u043e\u043b\u044c\u0437\u0443\u0435\u0442 \u0434\u0435\u043a\u043e\u0440\u0430\u0442\u043e\u0440 <code>@triton.jit<\/code> \u0434\u043b\u044f \u043a\u043e\u043c\u043f\u0438\u043b\u044f\u0446\u0438\u0438 Python-\u0444\u0443\u043d\u043a\u0446\u0438\u0439 \u0432 GPU \u044f\u0434\u0440\u0430. \u041f\u0440\u0438\u043c\u0435\u0440 \u043e\u043f\u0440\u0435\u0434\u0435\u043b\u0435\u043d\u0438\u044f \u0444\u0443\u043d\u043a\u0446\u0438\u0438:<\/p>\n<pre><code class=\"python\">@triton.jit def my_kernel(x_ptr, y_ptr, z_ptr, N):     # \u0442\u0435\u043b\u043e \u0444\u0443\u043d\u043a\u0446\u0438\u0438<\/code><\/pre>\n<p>\u041c\u043e\u0436\u043d\u043e \u0438\u0441\u043f\u043e\u043b\u044c\u0437\u043e\u0432\u0430\u0442\u044c \u043e\u043f\u0435\u0440\u0430\u0446\u0438\u0438 \u0441 \u0442\u0435\u043d\u0437\u043e\u0440\u0430\u043c\u0438, \u043d\u0430\u043f\u043e\u0434\u043e\u0431\u0438\u0435 \u0442\u0435\u0445, \u0447\u0442\u043e \u0435\u0441\u0442\u044c \u0432 NumPy. \u041f\u0440\u0438\u043c\u0435\u0440 \u0441\u043e\u0437\u0434\u0430\u043d\u0438\u044f \u0442\u0435\u043d\u0437\u043e\u0440\u0430 \u0438 \u0432\u044b\u043f\u043e\u043b\u043d\u0435\u043d\u0438\u044f \u043e\u043f\u0435\u0440\u0430\u0446\u0438\u0438:<\/p>\n<pre><code class=\"python\">x = triton.testing.randn((N,), dtype=torch.float32, device=device) y = triton.testing.randn((N,), dtype=torch.float32, device=device)<\/code><\/pre>\n<p>\u0422\u0430\u043a\u0436\u0435 \u0435\u0441\u0442\u044c \u043f\u043e\u0434\u0434\u0435\u0440\u0436\u043a\u0430 \u043e\u043f\u0435\u0440\u0430\u0446\u0438\u0439 \u0438\u043d\u0434\u0435\u043a\u0441\u0430\u0446\u0438\u0438 \u0438 \u0441\u0440\u0435\u0437\u043e\u0432, \u0430\u043d\u0430\u043b\u043e\u0433\u0438\u0447\u043d\u044b\u0435 Python:<\/p>\n<pre><code class=\"python\">x = x_ptr + tl.arange(0, N, dtype=tl.int64)<\/code><\/pre>\n<p>\u0415\u0441\u0442\u044c \u0440\u0430\u0437\u043b\u0438\u0447\u043d\u044b\u0435 \u0444\u0443\u043d\u043a\u0446\u0438\u0438 \u0434\u043b\u044f \u0443\u043f\u0440\u0430\u0432\u043b\u0435\u043d\u0438\u044f \u043f\u0430\u043c\u044f\u0442\u044c\u044e, \u0432\u043a\u043b\u044e\u0447\u0430\u044f \u0437\u0430\u0433\u0440\u0443\u0437\u043a\u0443 <code>tl.load<\/code> \u0438 \u0441\u043e\u0445\u0440\u0430\u043d\u0435\u043d\u0438\u0435 <code>tl.store<\/code> \u0434\u0430\u043d\u043d\u044b\u0445:  <\/p>\n<pre><code class=\"python\">x = tl.load(x_ptr + idx) tl.store(z_ptr + idx, x + 5)<\/code><\/pre>\n<p>\u041c\u043e\u0436\u043d\u043e \u043a\u043e\u043d\u0442\u0440\u043e\u043b\u0438\u0440\u043e\u0432\u0430\u0442\u044c \u0440\u0430\u0441\u043f\u0440\u0435\u0434\u0435\u043b\u0435\u043d\u0438\u0435 \u043f\u0430\u043c\u044f\u0442\u0438 \u0438 \u0432\u044b\u043f\u043e\u043b\u043d\u0435\u043d\u0438\u0435 \u043f\u043e\u0442\u043e\u043a\u043e\u0432 \u043d\u0430 GPU \u0442\u0430\u043a\u0438\u043c \u043e\u0431\u0440\u0430\u0437\u043e\u043c:<\/p>\n<pre><code class=\"python\">BLOCK_SIZE = 128 grid = lambda meta: (triton.cdiv(N, BLOCK_SIZE),)<\/code><\/pre>\n<p>\u0415\u0441\u0442\u044c \u0442\u0430\u043a\u0436\u0435 \u043c\u043d\u043e\u0436\u0435\u0441\u0442\u0432\u043e \u0432\u0441\u0442\u0440\u043e\u0435\u043d\u043d\u044b\u0445 \u043c\u0430\u0442\u0435\u043c\u0430\u0442\u0438\u0447\u0435\u0441\u043a\u0438\u0445 \u0444\u0443\u043d\u043a\u0446\u0438\u0439:<\/p>\n<ol>\n<li>\n<p><strong>abs<\/strong> &#8212; \u0432\u044b\u0447\u0438\u0441\u043b\u044f\u0435\u0442 \u043f\u043e\u044d\u043b\u0435\u043c\u0435\u043d\u0442\u043d\u043e\u0435 \u0430\u0431\u0441\u043e\u043b\u044e\u0442\u043d\u043e\u0435 \u0437\u043d\u0430\u0447\u0435\u043d\u0438\u0435 <code>x<\/code>.<\/p>\n<\/li>\n<li>\n<p><strong>cdiv<\/strong> &#8212; \u0432\u044b\u0447\u0438\u0441\u043b\u044f\u0435\u0442 \u043f\u043e\u0442\u043e\u043b\u043e\u043a \u043e\u0442 \u0434\u0435\u043b\u0435\u043d\u0438\u044f <code>x<\/code> \u043d\u0430 <code>div<\/code>.<\/p>\n<\/li>\n<li>\n<p><strong>clamp<\/strong> &#8212; \u043e\u0433\u0440\u0430\u043d\u0438\u0447\u0438\u0432\u0430\u0435\u0442 \u0442\u0435\u043d\u0437\u043e\u0440 <code>x<\/code> \u0432 \u043f\u0440\u0435\u0434\u0435\u043b\u0430\u0445 \u0443\u043a\u0430\u0437\u0430\u043d\u043d\u043e\u0433\u043e \u0434\u0438\u0430\u043f\u0430\u0437\u043e\u043d\u0430 <code>[min, max]<\/code>.<\/p>\n<\/li>\n<li>\n<p><strong>cos<\/strong> &#8212; \u0432\u044b\u0447\u0438\u0441\u043b\u044f\u0435\u0442 \u043f\u043e\u044d\u043b\u0435\u043c\u0435\u043d\u0442\u043d\u044b\u0439 \u043a\u043e\u0441\u0438\u043d\u0443\u0441 <code>x<\/code>.<\/p>\n<\/li>\n<li>\n<p><strong>div_rn<\/strong> &#8212; \u0432\u044b\u0447\u0438\u0441\u043b\u044f\u0435\u0442 \u043f\u043e\u044d\u043b\u0435\u043c\u0435\u043d\u0442\u043d\u043e\u0435 \u0442\u043e\u0447\u043d\u043e\u0435 \u0434\u0435\u043b\u0435\u043d\u0438\u0435 <code>x<\/code> \u043d\u0430 <code>y<\/code> \u0441 \u043e\u043a\u0440\u0443\u0433\u043b\u0435\u043d\u0438\u0435\u043c \u043a \u0431\u043b\u0438\u0436\u0430\u0439\u0448\u0435\u043c\u0443 \u0446\u0435\u043b\u043e\u043c\u0443.<\/p>\n<\/li>\n<li>\n<p><strong>erf<\/strong> &#8212; \u0432\u044b\u0447\u0438\u0441\u043b\u044f\u0435\u0442 \u043f\u043e\u044d\u043b\u0435\u043c\u0435\u043d\u0442\u043d\u0443\u044e \u0444\u0443\u043d\u043a\u0446\u0438\u044e \u043e\u0448\u0438\u0431\u043e\u043a <code>x<\/code>.<\/p>\n<\/li>\n<li>\n<p><strong>exp<\/strong> &#8212; \u0432\u044b\u0447\u0438\u0441\u043b\u044f\u0435\u0442 \u043f\u043e\u044d\u043b\u0435\u043c\u0435\u043d\u0442\u043d\u0443\u044e \u044d\u043a\u0441\u043f\u043e\u043d\u0435\u043d\u0442\u0443 <code>x<\/code>.<\/p>\n<\/li>\n<li>\n<p><strong>exp2<\/strong> &#8212; \u0432\u044b\u0447\u0438\u0441\u043b\u044f\u0435\u0442 \u043f\u043e\u044d\u043b\u0435\u043c\u0435\u043d\u0442\u043d\u0443\u044e \u044d\u043a\u0441\u043f\u043e\u043d\u0435\u043d\u0442\u0443 <code>x<\/code> \u043f\u043e \u043e\u0441\u043d\u043e\u0432\u0430\u043d\u0438\u044e 2.<\/p>\n<\/li>\n<li>\n<p><strong>fma<\/strong> &#8212; \u0432\u044b\u0447\u0438\u0441\u043b\u044f\u0435\u0442 \u043f\u043e\u044d\u043b\u0435\u043c\u0435\u043d\u0442\u043d\u043e \u0441\u043b\u0438\u044f\u043d\u0438\u0435 \u0443\u043c\u043d\u043e\u0436\u0435\u043d\u0438\u044f \u0438 \u0441\u043b\u043e\u0436\u0435\u043d\u0438\u044f \u0434\u043b\u044f <code>x<\/code>, <code>y<\/code> \u0438 <code>z<\/code> (x * y + z).<\/p>\n<\/li>\n<li>\n<p><strong>fdiv<\/strong> &#8212; \u0432\u044b\u0447\u0438\u0441\u043b\u044f\u0435\u0442 \u043f\u043e\u044d\u043b\u0435\u043c\u0435\u043d\u0442\u043d\u043e\u0435 \u0431\u044b\u0441\u0442\u0440\u043e\u0435 \u0434\u0435\u043b\u0435\u043d\u0438\u0435 <code>x<\/code> \u043d\u0430 <code>y<\/code>.<\/p>\n<\/li>\n<li>\n<p><strong>floor<\/strong> &#8212; \u0432\u044b\u0447\u0438\u0441\u043b\u044f\u0435\u0442 \u043f\u043e\u044d\u043b\u0435\u043c\u0435\u043d\u0442\u043d\u043e\u0435 \u043e\u043a\u0440\u0443\u0433\u043b\u0435\u043d\u0438\u0435 <code>x<\/code> \u0432\u043d\u0438\u0437.<\/p>\n<\/li>\n<li>\n<p><strong>log<\/strong> &#8212; \u0432\u044b\u0447\u0438\u0441\u043b\u044f\u0435\u0442 \u043f\u043e\u044d\u043b\u0435\u043c\u0435\u043d\u0442\u043d\u044b\u0439 \u043d\u0430\u0442\u0443\u0440\u0430\u043b\u044c\u043d\u044b\u0439 \u043b\u043e\u0433\u0430\u0440\u0438\u0444\u043c <code>x<\/code>.<\/p>\n<\/li>\n<li>\n<p><strong>log2<\/strong> &#8212; \u0432\u044b\u0447\u0438\u0441\u043b\u044f\u0435\u0442 \u043f\u043e\u044d\u043b\u0435\u043c\u0435\u043d\u0442\u043d\u044b\u0439 \u043b\u043e\u0433\u0430\u0440\u0438\u0444\u043c <code>x<\/code> \u043f\u043e \u043e\u0441\u043d\u043e\u0432\u0430\u043d\u0438\u044e 2.<\/p>\n<\/li>\n<li>\n<p><strong>maximum<\/strong> &#8212; \u0432\u044b\u0447\u0438\u0441\u043b\u044f\u0435\u0442 \u043f\u043e\u044d\u043b\u0435\u043c\u0435\u043d\u0442\u043d\u044b\u0439 \u043c\u0430\u043a\u0441\u0438\u043c\u0443\u043c \u0438\u0437 <code>x<\/code> \u0438 <code>y<\/code>.<\/p>\n<\/li>\n<li>\n<p><strong>minimum<\/strong> &#8212; \u0432\u044b\u0447\u0438\u0441\u043b\u044f\u0435\u0442 \u043f\u043e\u044d\u043b\u0435\u043c\u0435\u043d\u0442\u043d\u044b\u0439 \u043c\u0438\u043d\u0438\u043c\u0443\u043c \u0438\u0437 <code>x<\/code> \u0438 <code>y<\/code>.<\/p>\n<\/li>\n<li>\n<p><strong>sigmoid<\/strong> &#8212; \u0432\u044b\u0447\u0438\u0441\u043b\u044f\u0435\u0442 \u043f\u043e\u044d\u043b\u0435\u043c\u0435\u043d\u0442\u043d\u0443\u044e \u0441\u0438\u0433\u043c\u043e\u0438\u0434\u043d\u0443\u044e \u0444\u0443\u043d\u043a\u0446\u0438\u044e <code>x<\/code>.<\/p>\n<\/li>\n<li>\n<p><strong>sin<\/strong> &#8212; \u0432\u044b\u0447\u0438\u0441\u043b\u044f\u0435\u0442 \u043f\u043e\u044d\u043b\u0435\u043c\u0435\u043d\u0442\u043d\u044b\u0439 \u0441\u0438\u043d\u0443\u0441 <code>x<\/code>.<\/p>\n<\/li>\n<li>\n<p><strong>softmax<\/strong> &#8212; \u0432\u044b\u0447\u0438\u0441\u043b\u044f\u0435\u0442 \u043f\u043e\u044d\u043b\u0435\u043c\u0435\u043d\u0442\u043d\u0443\u044e softmax \u0444\u0443\u043d\u043a\u0446\u0438\u044e <code>x<\/code>.<\/p>\n<\/li>\n<li>\n<p><strong>sqrt<\/strong> &#8212; \u0432\u044b\u0447\u0438\u0441\u043b\u044f\u0435\u0442 \u043f\u043e\u044d\u043b\u0435\u043c\u0435\u043d\u0442\u043d\u044b\u0439 \u0431\u044b\u0441\u0442\u0440\u044b\u0439 \u043a\u0432\u0430\u0434\u0440\u0430\u0442\u043d\u044b\u0439 \u043a\u043e\u0440\u0435\u043d\u044c <code>x<\/code>.<\/p>\n<\/li>\n<li>\n<p><strong>sqrt_rn<\/strong> &#8212; \u0432\u044b\u0447\u0438\u0441\u043b\u044f\u0435\u0442 \u043f\u043e\u044d\u043b\u0435\u043c\u0435\u043d\u0442\u043d\u044b\u0439 \u0442\u043e\u0447\u043d\u044b\u0439 \u043a\u0432\u0430\u0434\u0440\u0430\u0442\u043d\u044b\u0439 \u043a\u043e\u0440\u0435\u043d\u044c <code>x<\/code> \u0441 \u043e\u043a\u0440\u0443\u0433\u043b\u0435\u043d\u0438\u0435\u043c \u043a \u0431\u043b\u0438\u0436\u0430\u0439\u0448\u0435\u043c\u0443.<\/p>\n<\/li>\n<li>\n<p><strong>umulhi<\/strong> &#8212; \u0432\u044b\u0447\u0438\u0441\u043b\u044f\u0435\u0442 \u043f\u043e\u044d\u043b\u0435\u043c\u0435\u043d\u0442\u043d\u043e \u0441\u0442\u0430\u0440\u0448\u0438\u0435 N \u0431\u0438\u0442 \u0438\u0437 2N-\u0431\u0438\u0442\u043d\u043e\u0433\u043e \u043f\u0440\u043e\u0438\u0437\u0432\u0435\u0434\u0435\u043d\u0438\u044f <code>x<\/code> \u0438 <code>y<\/code>.<\/p>\n<\/li>\n<\/ol>\n<p>\u042e\u0437\u0430\u0442\u044c \u0438\u0445 \u0434\u043e\u0441\u0442\u0430\u0442\u043e\u0447\u043d\u043e \u043f\u0440\u043e\u0441\u0442\u043e, \u043a \u043f\u0440\u0438\u043c\u0435\u0440\u0443:<\/p>\n<pre><code class=\"python\">z = tl.maximum(x, y)<\/code><\/pre>\n<p>\u041f\u043e\u0441\u043b\u0435 \u043e\u043f\u0440\u0435\u0434\u0435\u043b\u0435\u043d\u0438\u044f \u0444\u0443\u043d\u043a\u0446\u0438\u0438 \u0435\u0451 \u043c\u043e\u0436\u043d\u043e \u0441\u043a\u043e\u043c\u043f\u0438\u043b\u0438\u0440\u043e\u0432\u0430\u0442\u044c \u0438 \u0432\u044b\u043f\u043e\u043b\u043d\u0438\u0442\u044c, \u043f\u0435\u0440\u0435\u0434\u0430\u0432 \u0441\u043e\u043e\u0442\u0432\u0435\u0442\u0441\u0442\u0432\u0443\u044e\u0449\u0438\u0435 \u043f\u0430\u0440\u0430\u043c\u0435\u0442\u0440\u044b:<\/p>\n<pre><code class=\"python\">my_kernel[grid](x_ptr, y_ptr, z_ptr, N)<\/code><\/pre>\n<p>\u0415\u0441\u0442\u044c \u043f\u043e\u0434\u0434\u0435\u0440\u0436\u043a\u0430 \u0432\u0435\u043a\u0442\u043e\u0440\u0438\u0437\u0430\u0446\u0438\u0438:<\/p>\n<pre><code class=\"python\">x = tl.load(x_ptr + tl.arange(0, BLOCK_SIZE))<\/code><\/pre>\n<p>\u0415\u0441\u0442\u044c \u0441\u0442\u0430\u043d\u0434\u0430\u0440\u0442\u043d\u044b\u0435 \u0443\u0441\u043b\u043e\u0432\u043d\u044b\u0435 \u043e\u043f\u0435\u0440\u0430\u0442\u043e\u0440\u044b, \u0442\u0430\u043a\u0438\u0435 \u043a\u0430\u043a <code>if<\/code>, <code>else<\/code> :<\/p>\n<pre><code class=\"python\">if idx &lt; N:     x = x_ptr[idx]     y = y_ptr[idx]     z_ptr[idx] = x + y<\/code><\/pre>\n<p>\u041a\u043e\u0433\u0434\u0430 \u043d\u0443\u0436\u043d\u043e, \u0447\u0442\u043e\u0431\u044b \u043d\u0435\u0441\u043a\u043e\u043b\u044c\u043a\u043e \u043f\u043e\u0442\u043e\u043a\u043e\u0432 \u0434\u043e\u043b\u0436\u043d\u044b \u0431\u0435\u0437\u043e\u043f\u0430\u0441\u043d\u043e \u043e\u0431\u043d\u043e\u0432\u043b\u044f\u0442\u044c \u043e\u0434\u043d\u0438 \u0438 \u0442\u0435 \u0436\u0435 \u0434\u0430\u043d\u043d\u044b\u0435, \u043c\u043e\u0436\u043d\u043e \u044e\u0437\u0430\u0442\u044c \u0430\u0442\u043e\u043c\u0430\u0440\u043d\u044b\u0435 \u043e\u043f\u0435\u0440\u0430\u0446\u0438\u0438, \u0442\u0430\u043a\u0438\u0435 \u043a\u0430\u043a <code>tl.atomic_add<\/code>:  <\/p>\n<pre><code class=\"python\">tl.atomic_add(z_ptr[idx], x + y)<\/code><\/pre>\n<p>\u0417\u0434\u0435\u0441\u044c \u043c\u044b \u0432\u044b\u043f\u043e\u043b\u043d\u0438\u043b\u0438 \u0430\u0442\u043e\u043c\u0430\u0440\u043d\u043e\u0435 \u0441\u043b\u043e\u0436\u0435\u043d\u0438\u0435 \u043f\u043e \u0443\u043a\u0430\u0437\u0430\u043d\u043d\u043e\u043c\u0443 \u0430\u0434\u0440\u0435\u0441\u0443 \u043f\u0430\u043c\u044f\u0442\u0438. <\/p>\n<p>\u041f\u0440\u043e\u0447\u0438\u0435 \u0430\u0442\u043e\u043c\u0430\u0440\u043d\u044b\u0435 \u043e\u043f\u0435\u0440\u0430\u0446\u0438\u0438:<\/p>\n<ol>\n<li>\n<p><strong>atomic_cas<\/strong><\/p>\n<ul>\n<li>\n<p>\u0412\u044b\u043f\u043e\u043b\u043d\u044f\u0435\u0442 \u0430\u0442\u043e\u043c\u0430\u0440\u043d\u0443\u044e \u043e\u043f\u0435\u0440\u0430\u0446\u0438\u044e \u0441\u0440\u0430\u0432\u043d\u0435\u043d\u0438\u044f \u0438 \u0437\u0430\u043c\u0435\u043d\u044b \u043f\u043e \u0443\u043a\u0430\u0437\u0430\u043d\u043d\u043e\u043c\u0443 \u0430\u0434\u0440\u0435\u0441\u0443 \u043f\u0430\u043c\u044f\u0442\u0438.<\/p>\n<\/li>\n<\/ul>\n<\/li>\n<li>\n<p><strong>atomic_max<\/strong><\/p>\n<ul>\n<li>\n<p>\u0412\u044b\u043f\u043e\u043b\u043d\u044f\u0435\u0442 \u0430\u0442\u043e\u043c\u0430\u0440\u043d\u043e\u0435 \u043d\u0430\u0445\u043e\u0436\u0434\u0435\u043d\u0438\u0435 \u043c\u0430\u043a\u0441\u0438\u043c\u0443\u043c\u0430 \u043f\u043e \u0443\u043a\u0430\u0437\u0430\u043d\u043d\u043e\u043c\u0443 \u0430\u0434\u0440\u0435\u0441\u0443 \u043f\u0430\u043c\u044f\u0442\u0438.<\/p>\n<\/li>\n<\/ul>\n<\/li>\n<li>\n<p><strong>atomic_min<\/strong><\/p>\n<ul>\n<li>\n<p>\u0412\u044b\u043f\u043e\u043b\u043d\u044f\u0435\u0442 \u0430\u0442\u043e\u043c\u0430\u0440\u043d\u043e\u0435 \u043d\u0430\u0445\u043e\u0436\u0434\u0435\u043d\u0438\u0435 \u043c\u0438\u043d\u0438\u043c\u0443\u043c\u0430 \u043f\u043e \u0443\u043a\u0430\u0437\u0430\u043d\u043d\u043e\u043c\u0443 \u0430\u0434\u0440\u0435\u0441\u0443 \u043f\u0430\u043c\u044f\u0442\u0438.<\/p>\n<\/li>\n<\/ul>\n<\/li>\n<li>\n<p><strong>atomic_or<\/strong><\/p>\n<ul>\n<li>\n<p>\u0412\u044b\u043f\u043e\u043b\u043d\u044f\u0435\u0442 \u0430\u0442\u043e\u043c\u0430\u0440\u043d\u0443\u044e \u043b\u043e\u0433\u0438\u0447\u0435\u0441\u043a\u0443\u044e \u043e\u043f\u0435\u0440\u0430\u0446\u0438\u044e \u0418\u041b\u0418 \u043f\u043e \u0443\u043a\u0430\u0437\u0430\u043d\u043d\u043e\u043c\u0443 \u0430\u0434\u0440\u0435\u0441\u0443 \u043f\u0430\u043c\u044f\u0442\u0438.<\/p>\n<\/li>\n<\/ul>\n<\/li>\n<li>\n<p><strong>atomic_xchg<\/strong><\/p>\n<ul>\n<li>\n<p>\u0412\u044b\u043f\u043e\u043b\u043d\u044f\u0435\u0442 \u0430\u0442\u043e\u043c\u0430\u0440\u043d\u044b\u0439 \u043e\u0431\u043c\u0435\u043d \u0437\u043d\u0430\u0447\u0435\u043d\u0438\u044f\u043c\u0438 \u043f\u043e \u0443\u043a\u0430\u0437\u0430\u043d\u043d\u043e\u043c\u0443 \u0430\u0434\u0440\u0435\u0441\u0443 \u043f\u0430\u043c\u044f\u0442\u0438.<\/p>\n<\/li>\n<\/ul>\n<\/li>\n<li>\n<p><strong>atomic_xor<\/strong><\/p>\n<ul>\n<li>\n<p>\u0412\u044b\u043f\u043e\u043b\u043d\u044f\u0435\u0442 \u0430\u0442\u043e\u043c\u0430\u0440\u043d\u0443\u044e \u043b\u043e\u0433\u0438\u0447\u0435\u0441\u043a\u0443\u044e \u043e\u043f\u0435\u0440\u0430\u0446\u0438\u044e \u0438\u0441\u043a\u043b\u044e\u0447\u0430\u044e\u0449\u0435\u0435 \u0418\u041b\u0418 \u043f\u043e \u0443\u043a\u0430\u0437\u0430\u043d\u043d\u043e\u043c\u0443 \u0430\u0434\u0440\u0435\u0441\u0443 \u043f\u0430\u043c\u044f\u0442\u0438.<\/p>\n<\/li>\n<\/ul>\n<\/li>\n<\/ol>\n<p>\u0422\u0430\u043a\u0436\u0435 \u0435\u0441\u0442\u044c \u043e\u043f\u0435\u0440\u0430\u0446\u0438\u0438 \u0441\u043a\u0430\u043d\u0438\u0440\u043e\u0432\u0430\u043d\u0438\u044f \u0438 \u0441\u043e\u0440\u0442\u0438\u0440\u043e\u0432\u043a\u0438:<\/p>\n<ol>\n<li>\n<p><strong>associative_scan<\/strong><\/p>\n<ul>\n<li>\n<p>\u041f\u0440\u0438\u043c\u0435\u043d\u044f\u0435\u0442 \u0444\u0443\u043d\u043a\u0446\u0438\u044e <code>combine_fn<\/code> \u043a \u043a\u0430\u0436\u0434\u043e\u043c\u0443 \u044d\u043b\u0435\u043c\u0435\u043d\u0442\u0443 \u0441 \u0441\u043e\u0445\u0440\u0430\u043d\u0435\u043d\u0438\u0435\u043c \u043f\u0440\u043e\u043c\u0435\u0436\u0443\u0442\u043e\u0447\u043d\u043e\u0433\u043e \u0437\u043d\u0430\u0447\u0435\u043d\u0438\u044f <code>carry<\/code> \u0432 \u0442\u0435\u043d\u0437\u043e\u0440\u0430\u0445 \u0432\u0434\u043e\u043b\u044c \u0443\u043a\u0430\u0437\u0430\u043d\u043d\u043e\u0439 \u043e\u0441\u0438 \u0438 \u043e\u0431\u043d\u043e\u0432\u043b\u044f\u0435\u0442 <code>carry<\/code>.<\/p>\n<\/li>\n<\/ul>\n<\/li>\n<li>\n<p><strong>cumprod<\/strong><\/p>\n<ul>\n<li>\n<p>\u0412\u043e\u0437\u0432\u0440\u0430\u0449\u0430\u0435\u0442 \u043a\u0443\u043c\u0443\u043b\u044f\u0442\u0438\u0432\u043d\u043e\u0435 \u043f\u0440\u043e\u0438\u0437\u0432\u0435\u0434\u0435\u043d\u0438\u0435 \u0432\u0441\u0435\u0445 \u044d\u043b\u0435\u043c\u0435\u043d\u0442\u043e\u0432 \u0432 \u0442\u0435\u043d\u0437\u043e\u0440\u0435 \u0432\u0434\u043e\u043b\u044c \u0443\u043a\u0430\u0437\u0430\u043d\u043d\u043e\u0439 \u043e\u0441\u0438.<\/p>\n<\/li>\n<\/ul>\n<\/li>\n<li>\n<p><strong>cumsum<\/strong><\/p>\n<ul>\n<li>\n<p>\u0412\u043e\u0437\u0432\u0440\u0430\u0449\u0430\u0435\u0442 \u043a\u0443\u043c\u0443\u043b\u044f\u0442\u0438\u0432\u043d\u0443\u044e \u0441\u0443\u043c\u043c\u0443 \u0432\u0441\u0435\u0445 \u044d\u043b\u0435\u043c\u0435\u043d\u0442\u043e\u0432 \u0432 \u0442\u0435\u043d\u0437\u043e\u0440\u0435 \u0432\u0434\u043e\u043b\u044c \u0443\u043a\u0430\u0437\u0430\u043d\u043d\u043e\u0439 \u043e\u0441\u0438.<\/p>\n<\/li>\n<\/ul>\n<\/li>\n<li>\n<p><strong>histogram<\/strong><\/p>\n<ul>\n<li>\n<p>\u0412\u044b\u0447\u0438\u0441\u043b\u044f\u0435\u0442 \u0433\u0438\u0441\u0442\u043e\u0433\u0440\u0430\u043c\u043c\u0443 \u043d\u0430 \u043e\u0441\u043d\u043e\u0432\u0435 \u0432\u0445\u043e\u0434\u043d\u043e\u0433\u043e \u0442\u0435\u043d\u0437\u043e\u0440\u0430 \u0441 \u0437\u0430\u0434\u0430\u043d\u043d\u044b\u043c \u0447\u0438\u0441\u043b\u043e\u043c \u043a\u043e\u0440\u0437\u0438\u043d <code>num_bins<\/code>, \u043a\u043e\u0440\u0437\u0438\u043d\u044b \u0438\u043c\u0435\u044e\u0442 \u0448\u0438\u0440\u0438\u043d\u0443 1 \u0438 \u043d\u0430\u0447\u0438\u043d\u0430\u044e\u0442\u0441\u044f \u0441 0.<\/p>\n<\/li>\n<\/ul>\n<\/li>\n<li>\n<p><strong>sort<\/strong><\/p>\n<ul>\n<li>\n<p>\u0421\u043e\u0440\u0442\u0438\u0440\u0443\u0435\u0442 \u044d\u043b\u0435\u043c\u0435\u043d\u0442\u044b \u0442\u0435\u043d\u0437\u043e\u0440\u0430. \u041c\u043e\u0436\u0435\u0442 \u0431\u044b\u0442\u044c \u0432\u044b\u0437\u0432\u0430\u043d\u0430 \u043a\u0430\u043a \u043c\u0435\u0442\u043e\u0434 \u0442\u0435\u043d\u0437\u043e\u0440\u0430 <code>x.sort(...)<\/code>.<\/p>\n<\/li>\n<\/ul>\n<\/li>\n<\/ol>\n<p>\u0418 \u0441\u0442\u0430\u043d\u0434\u0430\u0440\u0442\u043d\u044b\u0435 \u043e\u043f\u0435\u0440\u0430\u0446\u0438\u0438 \u0440\u0435\u0434\u0443\u043a\u0446\u0438\u0438:<\/p>\n<ol>\n<li>\n<p><strong>argmax<\/strong><\/p>\n<ul>\n<li>\n<p>\u0412\u043e\u0437\u0432\u0440\u0430\u0449\u0430\u0435\u0442 \u0438\u043d\u0434\u0435\u043a\u0441 \u043c\u0430\u043a\u0441\u0438\u043c\u0430\u043b\u044c\u043d\u043e\u0433\u043e \u044d\u043b\u0435\u043c\u0435\u043d\u0442\u0430 \u0432 \u0442\u0435\u043d\u0437\u043e\u0440\u0435 \u0432\u0434\u043e\u043b\u044c \u0443\u043a\u0430\u0437\u0430\u043d\u043d\u043e\u0439 \u043e\u0441\u0438.<\/p>\n<\/li>\n<\/ul>\n<\/li>\n<li>\n<p><strong>argmin<\/strong><\/p>\n<ul>\n<li>\n<p>\u0412\u043e\u0437\u0432\u0440\u0430\u0449\u0430\u0435\u0442 \u0438\u043d\u0434\u0435\u043a\u0441 \u043c\u0438\u043d\u0438\u043c\u0430\u043b\u044c\u043d\u043e\u0433\u043e \u044d\u043b\u0435\u043c\u0435\u043d\u0442\u0430 \u0432 \u0442\u0435\u043d\u0437\u043e\u0440\u0435 \u0432\u0434\u043e\u043b\u044c \u0443\u043a\u0430\u0437\u0430\u043d\u043d\u043e\u0439 \u043e\u0441\u0438.<\/p>\n<\/li>\n<\/ul>\n<\/li>\n<li>\n<p><strong>max<\/strong><\/p>\n<ul>\n<li>\n<p>\u0412\u043e\u0437\u0432\u0440\u0430\u0449\u0430\u0435\u0442 \u043c\u0430\u043a\u0441\u0438\u043c\u0430\u043b\u044c\u043d\u043e\u0435 \u0437\u043d\u0430\u0447\u0435\u043d\u0438\u0435 \u0441\u0440\u0435\u0434\u0438 \u0432\u0441\u0435\u0445 \u044d\u043b\u0435\u043c\u0435\u043d\u0442\u043e\u0432 \u0442\u0435\u043d\u0437\u043e\u0440\u0430 \u0432\u0434\u043e\u043b\u044c \u0443\u043a\u0430\u0437\u0430\u043d\u043d\u043e\u0439 \u043e\u0441\u0438.<\/p>\n<\/li>\n<\/ul>\n<\/li>\n<li>\n<p><strong>min<\/strong><\/p>\n<ul>\n<li>\n<p>\u0412\u043e\u0437\u0432\u0440\u0430\u0449\u0430\u0435\u0442 \u043c\u0438\u043d\u0438\u043c\u0430\u043b\u044c\u043d\u043e\u0435 \u0437\u043d\u0430\u0447\u0435\u043d\u0438\u0435 \u0441\u0440\u0435\u0434\u0438 \u0432\u0441\u0435\u0445 \u044d\u043b\u0435\u043c\u0435\u043d\u0442\u043e\u0432 \u0442\u0435\u043d\u0437\u043e\u0440\u0430 \u0432\u0434\u043e\u043b\u044c \u0443\u043a\u0430\u0437\u0430\u043d\u043d\u043e\u0439 \u043e\u0441\u0438.<\/p>\n<\/li>\n<\/ul>\n<\/li>\n<li>\n<p><strong>reduce<\/strong><\/p>\n<ul>\n<li>\n<p>\u041f\u0440\u0438\u043c\u0435\u043d\u044f\u0435\u0442 \u0444\u0443\u043d\u043a\u0446\u0438\u044e \u043a\u043e\u043c\u0431\u0438\u043d\u0438\u0440\u043e\u0432\u0430\u043d\u0438\u044f <code>combine_fn<\/code> \u043a\u043e \u0432\u0441\u0435\u043c \u044d\u043b\u0435\u043c\u0435\u043d\u0442\u0430\u043c \u0432\u0445\u043e\u0434\u043d\u044b\u0445 \u0442\u0435\u043d\u0437\u043e\u0440\u043e\u0432 \u0432\u0434\u043e\u043b\u044c \u0443\u043a\u0430\u0437\u0430\u043d\u043d\u043e\u0439 \u043e\u0441\u0438.<\/p>\n<\/li>\n<\/ul>\n<\/li>\n<\/ol>\n<h4>\u0418\u0442\u0430\u043a, \u043a\u0430\u043a \u0432\u0441\u0435 \u044d\u0442\u043e \u043f\u0440\u0438\u043c\u0435\u043d\u044f\u0442\u044c?<\/h4>\n<p>\u041c\u043e\u0436\u043d\u043e \u044e\u0437\u0430\u0442\u044c <code>atomic_add<\/code> \u0434\u043b\u044f \u044d\u0444\u0444\u0435\u043a\u0442\u0438\u0432\u043d\u043e\u0439 \u0440\u0435\u0430\u043b\u0438\u0437\u0430\u0446\u0438\u0438 \u0433\u0440\u0430\u0434\u0438\u0435\u043d\u0442\u043d\u043e\u0433\u043e \u0441\u043f\u0443\u0441\u043a\u0430 \u0432 \u0441\u0440\u0435\u0434\u0435 \u0441 \u0432\u044b\u0441\u043e\u043a\u043e\u0439 \u0441\u0442\u0435\u043f\u0435\u043d\u044c\u044e \u043f\u0430\u0440\u0430\u043b\u043b\u0435\u043b\u0438\u0437\u043c\u0430. \u041d\u0430\u043f\u0440\u0438\u043c\u0435\u0440, \u043f\u0440\u0438 \u043e\u0431\u0443\u0447\u0435\u043d\u0438\u0438 \u0431\u043e\u043b\u044c\u0448\u0438\u0445 \u043d\u0435\u0439\u0440\u043e\u043d\u043d\u044b\u0445 \u0441\u0435\u0442\u0435\u0439 \u043d\u0430 GPU, \u0430\u0442\u043e\u043c\u0430\u0440\u043d\u044b\u0435 \u043e\u043f\u0435\u0440\u0430\u0446\u0438\u0438 \u043c\u043e\u0433\u0443\u0442 \u043e\u0431\u0435\u0441\u043f\u0435\u0447\u0438\u0442\u044c \u043e\u0431\u043d\u043e\u0432\u043b\u0435\u043d\u0438\u0435 \u0432\u0435\u0441\u043e\u0432 \u0431\u0435\u0437 \u043a\u043e\u043d\u0444\u043b\u0438\u043a\u0442\u043e\u0432 \u043c\u0435\u0436\u0434\u0443 \u043f\u043e\u0442\u043e\u043a\u0430\u043c\u0438:<\/p>\n<pre><code class=\"python\">@triton.jit def update_weights(grads, weights, learning_rate):     pid = tl.program_id(0)     for i in range(pid, grads.shape[0], tl.num_programs()):         tl.atomic_add(weights, i, -learning_rate * grads[i])<\/code><\/pre>\n<p>\u041c\u043e\u0436\u043d\u043e \u0442\u0430\u043a\u0436\u0435 \u044e\u0437\u0430\u0442\u044c <code>sort<\/code> \u0438 <code>reduce<\/code> \u0434\u043b\u044f \u0440\u0435\u0430\u043b\u0438\u0437\u0430\u0446\u0438\u0438 \u0430\u043b\u0433\u043e\u0440\u0438\u0442\u043c\u043e\u0432 \u043a\u043b\u0430\u0441\u0442\u0435\u0440\u0438\u0437\u0430\u0446\u0438\u0438, \u0442\u0430\u043a\u0438\u0445 \u043a\u0430\u043a K-means, \u0433\u0434\u0435 \u043d\u0435\u043e\u0431\u0445\u043e\u0434\u0438\u043c\u043e \u0441\u043e\u0440\u0442\u0438\u0440\u043e\u0432\u0430\u0442\u044c \u0434\u0430\u043d\u043d\u044b\u0435 \u0438 \u0432\u044b\u0447\u0438\u0441\u043b\u044f\u0442\u044c \u0446\u0435\u043d\u0442\u0440\u043e\u0438\u0434\u044b:<\/p>\n<pre><code class=\"python\">@triton.jit def k_means_update(x, centroids, labels, num_clusters):     # \u0432\u044b\u0447\u0438\u0441\u043b\u0435\u043d\u0438\u0435 \u0440\u0430\u0441\u0441\u0442\u043e\u044f\u043d\u0438\u0439 \u0438 \u043f\u0440\u0438\u0441\u0432\u043e\u0435\u043d\u0438\u0435 \u043c\u0435\u0442\u043e\u043a     block_start = tl.program_id(0) * tl.num_programs()     distances = tl.zeros((num_clusters,), dtype=tl.float32)     for i in range(block_start, min(block_start + BLOCK_SIZE, x.shape[0])):         for c in range(num_clusters):             distances[c] = tl.sum(tl.pow(x[i] - centroids[c], 2))         labels[i] = tl.argmin(distances)     # \u043e\u0431\u043d\u043e\u0432\u043b\u0435\u043d\u0438\u0435 \u0446\u0435\u043d\u0442\u0440\u043e\u0438\u0434\u043e\u0432     for c in range(num_clusters):         assigned_pts = x[labels == c]         centroids[c] = tl.sum(assigned_pts, axis=0) \/ len(assigned_pts)<\/code><\/pre>\n<p>\u0410 \u0432\u043e\u0442 <code>exp<\/code>, <code>max<\/code>, \u0438 <code>log<\/code> \u043c\u043e\u0436\u043d\u043e \u044e\u0437\u0430\u0442\u044c \u043d\u0430\u043f\u0440\u0438\u043c\u0435\u0440, \u0434\u043b\u044f \u0444\u0443\u043d\u0446\u0438\u0439 <em>softmax<\/em>:<\/p>\n<pre><code class=\"python\">@triton.jit def softmax(x):     max_val = tl.max(x, axis=1, keepdims=True)     exp_x = tl.exp(x - max_val)     sum_exp_x = tl.sum(exp_x, axis=1, keepdims=True)     return exp_x \/ sum_exp_x<\/code><\/pre>\n<p>\u041c\u043e\u0436\u043d\u043e \u0438\u0441\u043f\u043e\u043b\u044c\u0437\u043e\u0432\u0430\u0442\u044c <code>cumsum<\/code> \u0434\u043b\u044f \u0440\u0430\u0441\u0447\u0451\u0442\u0430 \u0441\u043a\u043e\u043b\u044c\u0437\u044f\u0449\u0435\u0433\u043e \u0441\u0440\u0435\u0434\u043d\u0435\u0433\u043e (\u0447\u0442\u043e \u043f\u0440\u0438\u043c\u0435\u043d\u0438\u043c\u043e \u0432\u043e \u0432\u0440\u0435\u043c\u0435\u043d\u043d\u044b\u0445 \u0440\u044f\u0434\u0430\u0445):<\/p>\n<pre><code class=\"python\">@triton.jit def moving_average(data, window_size):     cum_sum = tl.cumsum(data, axis=0)     return (cum_sum[window_size:] - cum_sum[:-window_size]) \/ window_size<\/code><\/pre>\n<p><code>associative_scan<\/code> \u043f\u043e\u0434\u0445\u043e\u0434\u0438\u0442 \u0434\u043b\u044f \u0440\u0435\u0430\u043b\u0438\u0437\u0430\u0446\u0438\u0438 \u0430\u043b\u0433\u043e\u0440\u0438\u0442\u043c\u0430 \u043f\u0440\u044f\u043c\u043e\u0433\u043e \u0440\u0430\u0441\u043f\u0440\u043e\u0441\u0442\u0440\u0430\u043d\u0435\u043d\u0438\u044f \u0432 RNN, \u0433\u0434\u0435 \u043f\u043e\u0441\u043b\u0435\u0434\u043e\u0432\u0430\u0442\u0435\u043b\u044c\u043d\u044b\u0435 \u0437\u0430\u0432\u0438\u0441\u0438\u043c\u043e\u0441\u0442\u0438 \u043c\u043e\u0433\u0443\u0442 \u0431\u044b\u0442\u044c \u044d\u0444\u0444\u0435\u043a\u0442\u0438\u0432\u043d\u043e \u043e\u0431\u0440\u0430\u0431\u043e\u0442\u0430\u043d\u044b \u0441 \u043f\u043e\u043c\u043e\u0449\u044c\u044e \u044d\u0442\u043e\u0439 \u043e\u043f\u0435\u0440\u0430\u0446\u0438\u0438:<\/p>\n<pre><code class=\"python\"> @triton.jit def rnn_step(hidden, input, weights):     carry = tl.zeros_like(hidden)     for t in range(input.shape[0]):         carry = tl.associative_scan(lambda h, x: tl.tanh(x @ weights + h), input[t], axis=0)         hidden[t] = carry     return hidden<\/code><\/pre>\n<p>\u0422\u0430\u043a\u0436\u0435 \u043c\u043e\u0436\u043d\u043e \u043f\u0440\u0438\u043c\u0435\u043d\u0438\u0442\u044c <code>associative_scan<\/code> \u0434\u043b\u044f \u0440\u0435\u0430\u043b\u0438\u0437\u0430\u0446\u0438\u0438 \u043f\u0430\u0440\u0430\u043b\u043b\u0435\u043b\u044c\u043d\u043e\u0433\u043e \u0430\u043b\u0433\u043e\u0440\u0438\u0442\u043c\u0430 \u0431\u044b\u0441\u0442\u0440\u043e\u0433\u043e \u043f\u0440\u0435\u043e\u0431\u0440\u0430\u0437\u043e\u0432\u0430\u043d\u0438\u044f \u0424\u0443\u0440\u044c\u0435:<\/p>\n<pre><code class=\"python\">@triton.jit def fft_step(data, step_size):     idx = tl.program_id(0) * step_size * 2 + tl.arange(0, step_size)     t = exp(-2j * pi \/ (2 * step_size) * idx)     u = data[idx]     v = data[idx + step_size] * t     data[idx] = u + v     data[idx + step_size] = u - v<\/code><\/pre>\n<hr\/>\n<p>\u0411\u043e\u043b\u0435\u0435 \u043f\u043e\u0434\u0440\u043e\u0431\u043d\u043e \u0441 Triton <a href=\"https:\/\/triton-lang.org\/\">\u043c\u043e\u0436\u043d\u043e \u043e\u0437\u043d\u0430\u043a\u043e\u043c\u0438\u0442\u044c\u0441\u044f \u0437\u0434\u0435\u0441\u044c.<\/a><\/p>\n<p>\u041c\u0430\u0442\u0435\u0440\u0438\u0430\u043b \u043f\u043e\u0434\u0433\u043e\u0442\u043e\u0432\u043b\u0435\u043d \u0432 \u043f\u0440\u0435\u0434\u0434\u0432\u0435\u0440\u0438\u0438 \u0441\u0442\u0430\u0440\u0442\u0430 <a href=\"https:\/\/otus.pw\/H4dT\/\">\u0441\u043f\u0435\u0446\u0438\u0430\u043b\u0438\u0437\u0430\u0446\u0438\u0438 &#171;Machine Learning&#187;.<\/a><\/p>\n<\/p>\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\/807655\/\"> https:\/\/habr.com\/ru\/articles\/807655\/<\/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\"><\/figure>\n<p>\u041f\u0440\u0438\u0432\u0435\u0442, \u0425\u0430\u0431\u0440!<\/p>\n<p><strong>Triton<\/strong> \u0431\u044b\u043b \u0440\u0430\u0437\u0440\u0430\u0431\u043e\u0442\u0430\u043d \u0441\u043f\u0435\u0446\u0438\u0430\u043b\u044c\u043d\u043e \u0434\u043b\u044f \u0432\u044b\u043f\u043e\u043b\u043d\u0435\u043d\u0438\u044f \u043d\u0430 GPU \u0438 \u043f\u0440\u0435\u0434\u043e\u0441\u0442\u0430\u0432\u043b\u044f\u0435\u0442 \u0443\u0434\u043e\u0431\u043d\u0443\u044e Python-\u043e\u0440\u0438\u0435\u043d\u0442\u0438\u0440\u043e\u0432\u0430\u043d\u043d\u0443\u044e \u0441\u0440\u0435\u0434\u0443.<\/p>\n<p>Triton \u043f\u043e\u0437\u0432\u043e\u043b\u044f\u0435\u0442 \u0438\u0441\u043f\u043e\u043b\u044c\u0437\u043e\u0432\u0430\u0442\u044c \u043c\u043e\u0434\u0435\u043b\u044c \u043f\u0440\u043e\u0433\u0440\u0430\u043c\u043c\u0438\u0440\u043e\u0432\u0430\u043d\u0438\u044f, <strong>\u043e\u0441\u043d\u043e\u0432\u0430\u043d\u043d\u0443\u044e \u043d\u0430 \u0431\u043b\u043e\u043a\u0430\u0445,<\/strong> \u043a\u043e\u0442\u043e\u0440\u0430\u044f \u0437\u043d\u0430\u0447\u0438\u0442\u0435\u043b\u044c\u043d\u043e \u043e\u0442\u043b\u0438\u0447\u0430\u0435\u0442\u0441\u044f \u043e\u0442 \u0442\u0440\u0430\u0434\u0438\u0446\u0438\u043e\u043d\u043d\u043e\u0439 \u043c\u043e\u0434\u0435\u043b\u0438 CUDA. \u0412\u043c\u0435\u0441\u0442\u043e \u0443\u043f\u0440\u0430\u0432\u043b\u0435\u043d\u0438\u044f \u043f\u043e\u0442\u043e\u043a\u0430\u043c\u0438 \u043d\u0430 \u0443\u0440\u043e\u0432\u043d\u0435 \u0441\u043a\u0430\u043b\u044f\u0440\u043d\u044b\u0445 \u0438\u043d\u0441\u0442\u0440\u0443\u043a\u0446\u0438\u0439, <strong>Triton \u043e\u043f\u0435\u0440\u0438\u0440\u0443\u0435\u0442 \u0431\u043b\u043e\u043a\u0430\u043c\u0438 \u0434\u0430\u043d\u043d\u044b\u0445, \u0447\u0442\u043e \u0432 \u0446\u0435\u043b\u043e\u043c \u0434\u0430\u0435\u0442 \u0431\u043e\u043b\u0435\u0435 \u043b\u0443\u0447\u0448\u0443\u044e \u043f\u0440\u043e\u0438\u0437\u0432\u043e\u0434\u0438\u0442\u0435\u043b\u044c\u043d\u043e\u0441\u0442\u044c.<\/strong><\/p>\n<p>\u0412 \u043e\u0442\u043b\u0438\u0447\u0438\u0435 \u043e\u0442 \u0441\u0442\u0430\u043d\u0434\u0430\u0440\u0442\u043d\u043e\u0433\u043e \u043f\u043e\u0434\u0445\u043e\u0434\u0430 CUDA, \u0433\u0434\u0435 \u0438\u0441\u043f\u043e\u043b\u043d\u0435\u043d\u0438\u0435 \u043a\u043e\u0434\u0430 \u043e\u0440\u0433\u0430\u043d\u0438\u0437\u0443\u0435\u0442\u0441\u044f \u0447\u0435\u0440\u0435\u0437 \u0432\u0437\u0430\u0438\u043c\u043e\u0434\u0435\u0439\u0441\u0442\u0432\u0438\u0435 \u043c\u043d\u043e\u0436\u0435\u0441\u0442\u0432\u0430 \u043f\u043e\u0442\u043e\u043a\u043e\u0432, Triton \u0441\u0442\u0440\u0443\u043a\u0442\u0443\u0440\u0438\u0440\u0443\u0435\u0442 \u0432\u044b\u043f\u043e\u043b\u043d\u0435\u043d\u0438\u0435 \u043d\u0430 \u0443\u0440\u043e\u0432\u043d\u0435 \u043f\u0440\u043e\u0433\u0440\u0430\u043c\u043c. <strong>\u0422.\u0435 \u043a\u0430\u0436\u0434\u044b\u0439 \u0431\u043b\u043e\u043a \u043f\u0440\u043e\u0433\u0440\u0430\u043c\u043c\u044b \u043c\u043e\u0436\u0435\u0442 \u0431\u044b\u0442\u044c \u0438\u0441\u043f\u043e\u043b\u043d\u0435\u043d \u043d\u0435\u0437\u0430\u0432\u0438\u0441\u0438\u043c\u043e, \u0441 \u0432\u043e\u0437\u043c\u043e\u0436\u043d\u043e\u0441\u0442\u044c\u044e \u043e\u0431\u0440\u0430\u0449\u0435\u043d\u0438\u044f \u043a \u0433\u043b\u043e\u0431\u0430\u043b\u044c\u043d\u043e\u0439 \u043f\u0430\u043c\u044f\u0442\u0438 GPU<\/strong> \u0438 \u0432\u044b\u043f\u043e\u043b\u043d\u0435\u043d\u0438\u044f \u0430\u0441\u0438\u043d\u0445\u0440\u043e\u043d\u043d\u044b\u0445 \u043e\u043f\u0435\u0440\u0430\u0446\u0438\u0439 \u0431\u0435\u0437 \u044f\u0432\u043d\u043e\u0433\u043e \u0443\u043f\u0440\u0430\u0432\u043b\u0435\u043d\u0438\u044f \u0441\u0438\u043d\u0445\u0440\u043e\u043d\u0438\u0437\u0430\u0446\u0438\u0435\u0439 \u043f\u043e\u0442\u043e\u043a\u043e\u0432.<\/p>\n<p>\u041a\u043e\u043c\u043f\u0438\u043b\u044f\u0442\u043e\u0440 Triton \u043f\u0440\u0438\u043c\u0435\u043d\u044f\u0435\u0442 \u0441\u043b\u043e\u0436\u043d\u044b\u0435 \u0441\u0442\u0440\u0430\u0442\u0435\u0433\u0438\u0438 \u043e\u043f\u0442\u0438\u043c\u0438\u0437\u0430\u0446\u0438\u0438, \u043a \u043f\u0440\u0438\u043c\u0435\u0440\u0443 \u0442\u0430\u043a\u0438\u0435 \u043a\u0430\u043a \u0430\u043d\u0430\u043b\u0438\u0437 \u043f\u043e\u0442\u043e\u043a\u0430 \u0434\u0430\u043d\u043d\u044b\u0445 \u0438 \u0443\u043f\u0440\u0430\u0432\u043b\u0435\u043d\u0438\u0435 \u043f\u0430\u043c\u044f\u0442\u044c\u044e \u043d\u0430 \u0443\u0440\u043e\u0432\u043d\u0435 \u0431\u043b\u043e\u043a\u043e\u0432. \u042d\u0442\u043e \u0432\u043a\u043b\u044e\u0447\u0430\u0435\u0442 \u0432 \u0441\u0435\u0431\u044f: <em>\u0430\u0432\u0442\u043e\u043c\u0430\u0442\u0438\u0447\u0435\u0441\u043a\u0430\u044f \u0432\u0435\u043a\u0442\u043e\u0440\u0438\u0437\u0430\u0446\u0438\u044f<\/em>, <em>\u043f\u0440\u0435\u0434\u0432\u0430\u0440\u0438\u0442\u0435\u043b\u044c\u043d\u0430\u044f \u0432\u044b\u0431\u043e\u0440\u043a\u0430 \u0434\u0430\u043d\u043d\u044b\u0445,<\/em> \u0438 <em>\u0438\u0441\u043f\u043e\u043b\u044c\u0437\u043e\u0432\u0430\u043d\u0438\u0435 \u0442\u0435\u043d\u0437\u043e\u0440\u043d\u044b\u0445 \u044f\u0434\u0435\u0440<\/em>, \u0433\u0434\u0435 \u044d\u0442\u043e \u0432\u043e\u0437\u043c\u043e\u0436\u043d\u043e. \u0422\u0430\u043a\u043e\u0439 \u043f\u043e\u0434\u0445\u043e\u0434 \u043f\u043e\u0437\u0432\u043e\u043b\u044f\u0435\u0442 \u043c\u0430\u043a\u0441\u0438\u043c\u0430\u043b\u044c\u043d\u043e \u0438\u0441\u043f\u043e\u043b\u044c\u0437\u043e\u0432\u0430\u0442\u044c \u0432\u043e\u0437\u043c\u043e\u0436\u043d\u043e\u0441\u0442\u0438 GPU.<\/p>\n<h4>\u0423\u0441\u0442\u0430\u043d\u043e\u0432\u0438\u043c<\/h4>\n<p>\u0421\u0430\u043c\u044b\u0439 \u043f\u0440\u043e\u0441\u0442\u043e\u0439 \u0438 \u0434\u043e\u0441\u0442\u0443\u043f\u043d\u044b\u0439 \u0441\u043f\u043e\u0441\u043e\u0431 \u0443\u0441\u0442\u0430\u043d\u043e\u0432\u0438\u0442\u044c Triton \u2014 <strong>\u044d\u0442\u043e \u0447\u0435\u0440\u0435\u0437 pip:<\/strong><\/p>\n<pre><code>pip install triton<\/code><\/pre>\n<p>Triton <strong>\u043c\u043e\u0436\u043d\u043e \u0442\u0430\u043a\u0436\u0435 \u0441\u043e\u0431\u0440\u0430\u0442\u044c \u0438\u0437 \u0438\u0441\u0445\u043e\u0434\u043d\u0438\u043a\u043e\u0432.<\/strong> \u0414\u043b\u044f \u044d\u0442\u043e\u0433\u043e \u043a\u043b\u043e\u043d\u0438\u0440\u0443\u0435\u043c \u0440\u0435\u043f\u043e\u0437\u0438\u0442\u043e\u0440\u0438\u0439 \u0438 \u0443\u0441\u0442\u0430\u043d\u0430\u0432\u043b\u0438\u0432\u0430\u0435\u043c \u043d\u0435\u043e\u0431\u0445\u043e\u0434\u0438\u043c\u044b\u0435 \u0437\u0430\u0432\u0438\u0441\u0438\u043c\u043e\u0441\u0442\u0438:<\/p>\n<pre><code>git clone https:\/\/github.com\/openai\/triton.git cd triton\/python pip install ninja cmake wheel pip install -e .<\/code><\/pre>\n<p>\u0423\u0447\u0442\u0438\u0442\u0435, \u0447\u0442\u043e \u0435\u0441\u043b\u0438 \u043d\u0430 \u0441\u0438\u0441\u0442\u0435\u043c\u0435 \u043d\u0435 \u0443\u0441\u0442\u0430\u043d\u043e\u0432\u043b\u0435\u043d llvm, \u0441\u043a\u0440\u0438\u043f\u0442 <a href=\"http:\/\/setup.py\">setup.py<\/a> \u0441\u043a\u0430\u0447\u0430\u0435\u0442 \u043e\u0444\u0438\u0446\u0438\u0430\u043b\u044c\u043d\u044b\u0435 \u0441\u0442\u0430\u0442\u0438\u0447\u0435\u0441\u043a\u0438\u0435 \u0431\u0438\u0431\u043b\u0438\u043e\u0442\u0435\u043a\u0438 LLVM \u0438 \u0441\u0432\u044f\u0436\u0435\u0442 \u0438\u0445.<\/p>\n<h4>\u0420\u0430\u0441\u0441\u043c\u043e\u0442\u0440\u0438\u043c \u0432\u0435\u0441\u044c \u043e\u0441\u043d\u043e\u0432\u043d\u043e\u0439 \u0441\u0438\u043d\u0442\u0430\u043a\u0441\u0438\u0441<\/h4>\n<p>Triton \u0438\u0441\u043f\u043e\u043b\u044c\u0437\u0443\u0435\u0442 \u0434\u0435\u043a\u043e\u0440\u0430\u0442\u043e\u0440 <code>@triton.jit<\/code> \u0434\u043b\u044f \u043a\u043e\u043c\u043f\u0438\u043b\u044f\u0446\u0438\u0438 Python-\u0444\u0443\u043d\u043a\u0446\u0438\u0439 \u0432 GPU \u044f\u0434\u0440\u0430. \u041f\u0440\u0438\u043c\u0435\u0440 \u043e\u043f\u0440\u0435\u0434\u0435\u043b\u0435\u043d\u0438\u044f \u0444\u0443\u043d\u043a\u0446\u0438\u0438:<\/p>\n<pre><code class=\"python\">@triton.jit def my_kernel(x_ptr, y_ptr, z_ptr, N):     # \u0442\u0435\u043b\u043e \u0444\u0443\u043d\u043a\u0446\u0438\u0438<\/code><\/pre>\n<p>\u041c\u043e\u0436\u043d\u043e \u0438\u0441\u043f\u043e\u043b\u044c\u0437\u043e\u0432\u0430\u0442\u044c \u043e\u043f\u0435\u0440\u0430\u0446\u0438\u0438 \u0441 \u0442\u0435\u043d\u0437\u043e\u0440\u0430\u043c\u0438, \u043d\u0430\u043f\u043e\u0434\u043e\u0431\u0438\u0435 \u0442\u0435\u0445, \u0447\u0442\u043e \u0435\u0441\u0442\u044c \u0432 NumPy. \u041f\u0440\u0438\u043c\u0435\u0440 \u0441\u043e\u0437\u0434\u0430\u043d\u0438\u044f \u0442\u0435\u043d\u0437\u043e\u0440\u0430 \u0438 \u0432\u044b\u043f\u043e\u043b\u043d\u0435\u043d\u0438\u044f \u043e\u043f\u0435\u0440\u0430\u0446\u0438\u0438:<\/p>\n<pre><code class=\"python\">x = triton.testing.randn((N,), dtype=torch.float32, device=device) y = triton.testing.randn((N,), dtype=torch.float32, device=device)<\/code><\/pre>\n<p>\u0422\u0430\u043a\u0436\u0435 \u0435\u0441\u0442\u044c \u043f\u043e\u0434\u0434\u0435\u0440\u0436\u043a\u0430 \u043e\u043f\u0435\u0440\u0430\u0446\u0438\u0439 \u0438\u043d\u0434\u0435\u043a\u0441\u0430\u0446\u0438\u0438 \u0438 \u0441\u0440\u0435\u0437\u043e\u0432, \u0430\u043d\u0430\u043b\u043e\u0433\u0438\u0447\u043d\u044b\u0435 Python:<\/p>\n<pre><code class=\"python\">x = x_ptr + tl.arange(0, N, dtype=tl.int64)<\/code><\/pre>\n<p>\u0415\u0441\u0442\u044c \u0440\u0430\u0437\u043b\u0438\u0447\u043d\u044b\u0435 \u0444\u0443\u043d\u043a\u0446\u0438\u0438 \u0434\u043b\u044f \u0443\u043f\u0440\u0430\u0432\u043b\u0435\u043d\u0438\u044f \u043f\u0430\u043c\u044f\u0442\u044c\u044e, \u0432\u043a\u043b\u044e\u0447\u0430\u044f \u0437\u0430\u0433\u0440\u0443\u0437\u043a\u0443 <code>tl.load<\/code> \u0438 \u0441\u043e\u0445\u0440\u0430\u043d\u0435\u043d\u0438\u0435 <code>tl.store<\/code> \u0434\u0430\u043d\u043d\u044b\u0445:  <\/p>\n<pre><code class=\"python\">x = tl.load(x_ptr + idx) tl.store(z_ptr + idx, x + 5)<\/code><\/pre>\n<p>\u041c\u043e\u0436\u043d\u043e \u043a\u043e\u043d\u0442\u0440\u043e\u043b\u0438\u0440\u043e\u0432\u0430\u0442\u044c \u0440\u0430\u0441\u043f\u0440\u0435\u0434\u0435\u043b\u0435\u043d\u0438\u0435 \u043f\u0430\u043c\u044f\u0442\u0438 \u0438 \u0432\u044b\u043f\u043e\u043b\u043d\u0435\u043d\u0438\u0435 \u043f\u043e\u0442\u043e\u043a\u043e\u0432 \u043d\u0430 GPU \u0442\u0430\u043a\u0438\u043c \u043e\u0431\u0440\u0430\u0437\u043e\u043c:<\/p>\n<pre><code class=\"python\">BLOCK_SIZE = 128 grid = lambda meta: (triton.cdiv(N, BLOCK_SIZE),)<\/code><\/pre>\n<p>\u0415\u0441\u0442\u044c \u0442\u0430\u043a\u0436\u0435 \u043c\u043d\u043e\u0436\u0435\u0441\u0442\u0432\u043e \u0432\u0441\u0442\u0440\u043e\u0435\u043d\u043d\u044b\u0445 \u043c\u0430\u0442\u0435\u043c\u0430\u0442\u0438\u0447\u0435\u0441\u043a\u0438\u0445 \u0444\u0443\u043d\u043a\u0446\u0438\u0439:<\/p>\n<ol>\n<li>\n<p><strong>abs<\/strong> &#8212; \u0432\u044b\u0447\u0438\u0441\u043b\u044f\u0435\u0442 \u043f\u043e\u044d\u043b\u0435\u043c\u0435\u043d\u0442\u043d\u043e\u0435 \u0430\u0431\u0441\u043e\u043b\u044e\u0442\u043d\u043e\u0435 \u0437\u043d\u0430\u0447\u0435\u043d\u0438\u0435 <code>x<\/code>.<\/p>\n<\/li>\n<li>\n<p><strong>cdiv<\/strong> &#8212; \u0432\u044b\u0447\u0438\u0441\u043b\u044f\u0435\u0442 \u043f\u043e\u0442\u043e\u043b\u043e\u043a \u043e\u0442 \u0434\u0435\u043b\u0435\u043d\u0438\u044f <code>x<\/code> \u043d\u0430 <code>div<\/code>.<\/p>\n<\/li>\n<li>\n<p><strong>clamp<\/strong> &#8212; \u043e\u0433\u0440\u0430\u043d\u0438\u0447\u0438\u0432\u0430\u0435\u0442 \u0442\u0435\u043d\u0437\u043e\u0440 <code>x<\/code> \u0432 \u043f\u0440\u0435\u0434\u0435\u043b\u0430\u0445 \u0443\u043a\u0430\u0437\u0430\u043d\u043d\u043e\u0433\u043e \u0434\u0438\u0430\u043f\u0430\u0437\u043e\u043d\u0430 <code>[min, max]<\/code>.<\/p>\n<\/li>\n<li>\n<p><strong>cos<\/strong> &#8212; \u0432\u044b\u0447\u0438\u0441\u043b\u044f\u0435\u0442 \u043f\u043e\u044d\u043b\u0435\u043c\u0435\u043d\u0442\u043d\u044b\u0439 \u043a\u043e\u0441\u0438\u043d\u0443\u0441 <code>x<\/code>.<\/p>\n<\/li>\n<li>\n<p><strong>div_rn<\/strong> &#8212; \u0432\u044b\u0447\u0438\u0441\u043b\u044f\u0435\u0442 \u043f\u043e\u044d\u043b\u0435\u043c\u0435\u043d\u0442\u043d\u043e\u0435 \u0442\u043e\u0447\u043d\u043e\u0435 \u0434\u0435\u043b\u0435\u043d\u0438\u0435 <code>x<\/code> \u043d\u0430 <code>y<\/code> \u0441 \u043e\u043a\u0440\u0443\u0433\u043b\u0435\u043d\u0438\u0435\u043c \u043a \u0431\u043b\u0438\u0436\u0430\u0439\u0448\u0435\u043c\u0443 \u0446\u0435\u043b\u043e\u043c\u0443.<\/p>\n<\/li>\n<li>\n<p><strong>erf<\/strong> &#8212; \u0432\u044b\u0447\u0438\u0441\u043b\u044f\u0435\u0442 \u043f\u043e\u044d\u043b\u0435\u043c\u0435\u043d\u0442\u043d\u0443\u044e \u0444\u0443\u043d\u043a\u0446\u0438\u044e \u043e\u0448\u0438\u0431\u043e\u043a <code>x<\/code>.<\/p>\n<\/li>\n<li>\n<p><strong>exp<\/strong> &#8212; \u0432\u044b\u0447\u0438\u0441\u043b\u044f\u0435\u0442 \u043f\u043e\u044d\u043b\u0435\u043c\u0435\u043d\u0442\u043d\u0443\u044e \u044d\u043a\u0441\u043f\u043e\u043d\u0435\u043d\u0442\u0443 <code>x<\/code>.<\/p>\n<\/li>\n<li>\n<p><strong>exp2<\/strong> &#8212; \u0432\u044b\u0447\u0438\u0441\u043b\u044f\u0435\u0442 \u043f\u043e\u044d\u043b\u0435\u043c\u0435\u043d\u0442\u043d\u0443\u044e \u044d\u043a\u0441\u043f\u043e\u043d\u0435\u043d\u0442\u0443 <code>x<\/code> \u043f\u043e \u043e\u0441\u043d\u043e\u0432\u0430\u043d\u0438\u044e 2.<\/p>\n<\/li>\n<li>\n<p><strong>fma<\/strong> &#8212; \u0432\u044b\u0447\u0438\u0441\u043b\u044f\u0435\u0442 \u043f\u043e\u044d\u043b\u0435\u043c\u0435\u043d\u0442\u043d\u043e \u0441\u043b\u0438\u044f\u043d\u0438\u0435 \u0443\u043c\u043d\u043e\u0436\u0435\u043d\u0438\u044f \u0438 \u0441\u043b\u043e\u0436\u0435\u043d\u0438\u044f \u0434\u043b\u044f <code>x<\/code>, <code>y<\/code> \u0438 <code>z<\/code> (x * y + z).<\/p>\n<\/li>\n<li>\n<p><strong>fdiv<\/strong> &#8212; \u0432\u044b\u0447\u0438\u0441\u043b\u044f\u0435\u0442 \u043f\u043e\u044d\u043b\u0435\u043c\u0435\u043d\u0442\u043d\u043e\u0435 \u0431\u044b\u0441\u0442\u0440\u043e\u0435 \u0434\u0435\u043b\u0435\u043d\u0438\u0435 <code>x<\/code> \u043d\u0430 <code>y<\/code>.<\/p>\n<\/li>\n<li>\n<p><strong>floor<\/strong> &#8212; \u0432\u044b\u0447\u0438\u0441\u043b\u044f\u0435\u0442 \u043f\u043e\u044d\u043b\u0435\u043c\u0435\u043d\u0442\u043d\u043e\u0435 \u043e\u043a\u0440\u0443\u0433\u043b\u0435\u043d\u0438\u0435 <code>x<\/code> \u0432\u043d\u0438\u0437.<\/p>\n<\/li>\n<li>\n<p><strong>log<\/strong> &#8212; \u0432\u044b\u0447\u0438\u0441\u043b\u044f\u0435\u0442 \u043f\u043e\u044d\u043b\u0435\u043c\u0435\u043d\u0442\u043d\u044b\u0439 \u043d\u0430\u0442\u0443\u0440\u0430\u043b\u044c\u043d\u044b\u0439 \u043b\u043e\u0433\u0430\u0440\u0438\u0444\u043c <code>x<\/code>.<\/p>\n<\/li>\n<li>\n<p><strong>log2<\/strong> &#8212; \u0432\u044b\u0447\u0438\u0441\u043b\u044f\u0435\u0442 \u043f\u043e\u044d\u043b\u0435\u043c\u0435\u043d\u0442\u043d\u044b\u0439 \u043b\u043e\u0433\u0430\u0440\u0438\u0444\u043c <code>x<\/code> \u043f\u043e \u043e\u0441\u043d\u043e\u0432\u0430\u043d\u0438\u044e 2.<\/p>\n<\/li>\n<li>\n<p><strong>maximum<\/strong> &#8212; \u0432\u044b\u0447\u0438\u0441\u043b\u044f\u0435\u0442 \u043f\u043e\u044d\u043b\u0435\u043c\u0435\u043d\u0442\u043d\u044b\u0439 \u043c\u0430\u043a\u0441\u0438\u043c\u0443\u043c \u0438\u0437 <code>x<\/code> \u0438 <code>y<\/code>.<\/p>\n<\/li>\n<li>\n<p><strong>minimum<\/strong> &#8212; \u0432\u044b\u0447\u0438\u0441\u043b\u044f\u0435\u0442 \u043f\u043e\u044d\u043b\u0435\u043c\u0435\u043d\u0442\u043d\u044b\u0439 \u043c\u0438\u043d\u0438\u043c\u0443\u043c \u0438\u0437 <code>x<\/code> \u0438 <code>y<\/code>.<\/p>\n<\/li>\n<li>\n<p><strong>sigmoid<\/strong> &#8212; \u0432\u044b\u0447\u0438\u0441\u043b\u044f\u0435\u0442 \u043f\u043e\u044d\u043b\u0435\u043c\u0435\u043d\u0442\u043d\u0443\u044e \u0441\u0438\u0433\u043c\u043e\u0438\u0434\u043d\u0443\u044e \u0444\u0443\u043d\u043a\u0446\u0438\u044e <code>x<\/code>.<\/p>\n<\/li>\n<li>\n<p><strong>sin<\/strong> &#8212; \u0432\u044b\u0447\u0438\u0441\u043b\u044f\u0435\u0442 \u043f\u043e\u044d\u043b\u0435\u043c\u0435\u043d\u0442\u043d\u044b\u0439 \u0441\u0438\u043d\u0443\u0441 <code>x<\/code>.<\/p>\n<\/li>\n<li>\n<p><strong>softmax<\/strong> &#8212; \u0432\u044b\u0447\u0438\u0441\u043b\u044f\u0435\u0442 \u043f\u043e\u044d\u043b\u0435\u043c\u0435\u043d\u0442\u043d\u0443\u044e softmax \u0444\u0443\u043d\u043a\u0446\u0438\u044e <code>x<\/code>.<\/p>\n<\/li>\n<li>\n<p><strong>sqrt<\/strong> &#8212; \u0432\u044b\u0447\u0438\u0441\u043b\u044f\u0435\u0442 \u043f\u043e\u044d\u043b\u0435\u043c\u0435\u043d\u0442\u043d\u044b\u0439 \u0431\u044b\u0441\u0442\u0440\u044b\u0439 \u043a\u0432\u0430\u0434\u0440\u0430\u0442\u043d\u044b\u0439 \u043a\u043e\u0440\u0435\u043d\u044c <code>x<\/code>.<\/p>\n<\/li>\n<li>\n<p><strong>sqrt_rn<\/strong> &#8212; \u0432\u044b\u0447\u0438\u0441\u043b\u044f\u0435\u0442 \u043f\u043e\u044d\u043b\u0435\u043c\u0435\u043d\u0442\u043d\u044b\u0439 \u0442\u043e\u0447\u043d\u044b\u0439 \u043a\u0432\u0430\u0434\u0440\u0430\u0442\u043d\u044b\u0439 \u043a\u043e\u0440\u0435\u043d\u044c <code>x<\/code> \u0441 \u043e\u043a\u0440\u0443\u0433\u043b\u0435\u043d\u0438\u0435\u043c \u043a \u0431\u043b\u0438\u0436\u0430\u0439\u0448\u0435\u043c\u0443.<\/p>\n<\/li>\n<li>\n<p><strong>umulhi<\/strong> &#8212; \u0432\u044b\u0447\u0438\u0441\u043b\u044f\u0435\u0442 \u043f\u043e\u044d\u043b\u0435\u043c\u0435\u043d\u0442\u043d\u043e \u0441\u0442\u0430\u0440\u0448\u0438\u0435 N \u0431\u0438\u0442 \u0438\u0437 2N-\u0431\u0438\u0442\u043d\u043e\u0433\u043e \u043f\u0440\u043e\u0438\u0437\u0432\u0435\u0434\u0435\u043d\u0438\u044f <code>x<\/code> \u0438 <code>y<\/code>.<\/p>\n<\/li>\n<\/ol>\n<p>\u042e\u0437\u0430\u0442\u044c \u0438\u0445 \u0434\u043e\u0441\u0442\u0430\u0442\u043e\u0447\u043d\u043e \u043f\u0440\u043e\u0441\u0442\u043e, \u043a \u043f\u0440\u0438\u043c\u0435\u0440\u0443:<\/p>\n<pre><code class=\"python\">z = tl.maximum(x, y)<\/code><\/pre>\n<p>\u041f\u043e\u0441\u043b\u0435 \u043e\u043f\u0440\u0435\u0434\u0435\u043b\u0435\u043d\u0438\u044f \u0444\u0443\u043d\u043a\u0446\u0438\u0438 \u0435\u0451 \u043c\u043e\u0436\u043d\u043e \u0441\u043a\u043e\u043c\u043f\u0438\u043b\u0438\u0440\u043e\u0432\u0430\u0442\u044c \u0438 \u0432\u044b\u043f\u043e\u043b\u043d\u0438\u0442\u044c, \u043f\u0435\u0440\u0435\u0434\u0430\u0432 \u0441\u043e\u043e\u0442\u0432\u0435\u0442\u0441\u0442\u0432\u0443\u044e\u0449\u0438\u0435 \u043f\u0430\u0440\u0430\u043c\u0435\u0442\u0440\u044b:<\/p>\n<pre><code class=\"python\">my_kernel[grid](x_ptr, y_ptr, z_ptr, N)<\/code><\/pre>\n<p>\u0415\u0441\u0442\u044c \u043f\u043e\u0434\u0434\u0435\u0440\u0436\u043a\u0430 \u0432\u0435\u043a\u0442\u043e\u0440\u0438\u0437\u0430\u0446\u0438\u0438:<\/p>\n<pre><code class=\"python\">x = tl.load(x_ptr + tl.arange(0, BLOCK_SIZE))<\/code><\/pre>\n<p>\u0415\u0441\u0442\u044c \u0441\u0442\u0430\u043d\u0434\u0430\u0440\u0442\u043d\u044b\u0435 \u0443\u0441\u043b\u043e\u0432\u043d\u044b\u0435 \u043e\u043f\u0435\u0440\u0430\u0442\u043e\u0440\u044b, \u0442\u0430\u043a\u0438\u0435 \u043a\u0430\u043a <code>if<\/code>, <code>else<\/code> :<\/p>\n<pre><code class=\"python\">if idx &lt; N:     x = x_ptr[idx]     y = y_ptr[idx]     z_ptr[idx] = x + y<\/code><\/pre>\n<p>\u041a\u043e\u0433\u0434\u0430 \u043d\u0443\u0436\u043d\u043e, \u0447\u0442\u043e\u0431\u044b \u043d\u0435\u0441\u043a\u043e\u043b\u044c\u043a\u043e \u043f\u043e\u0442\u043e\u043a\u043e\u0432 \u0434\u043e\u043b\u0436\u043d\u044b \u0431\u0435\u0437\u043e\u043f\u0430\u0441\u043d\u043e \u043e\u0431\u043d\u043e\u0432\u043b\u044f\u0442\u044c \u043e\u0434\u043d\u0438 \u0438 \u0442\u0435 \u0436\u0435 \u0434\u0430\u043d\u043d\u044b\u0435, \u043c\u043e\u0436\u043d\u043e \u044e\u0437\u0430\u0442\u044c \u0430\u0442\u043e\u043c\u0430\u0440\u043d\u044b\u0435 \u043e\u043f\u0435\u0440\u0430\u0446\u0438\u0438, \u0442\u0430\u043a\u0438\u0435 \u043a\u0430\u043a <code>tl.atomic_add<\/code>:  <\/p>\n<pre><code class=\"python\">tl.atomic_add(z_ptr[idx], x + y)<\/code><\/pre>\n<p>\u0417\u0434\u0435\u0441\u044c \u043c\u044b \u0432\u044b\u043f\u043e\u043b\u043d\u0438\u043b\u0438 \u0430\u0442\u043e\u043c\u0430\u0440\u043d\u043e\u0435 \u0441\u043b\u043e\u0436\u0435\u043d\u0438\u0435 \u043f\u043e \u0443\u043a\u0430\u0437\u0430\u043d\u043d\u043e\u043c\u0443 \u0430\u0434\u0440\u0435\u0441\u0443 \u043f\u0430\u043c\u044f\u0442\u0438. <\/p>\n<p>\u041f\u0440\u043e\u0447\u0438\u0435 \u0430\u0442\u043e\u043c\u0430\u0440\u043d\u044b\u0435 \u043e\u043f\u0435\u0440\u0430\u0446\u0438\u0438:<\/p>\n<ol>\n<li>\n<p><strong>atomic_cas<\/strong><\/p>\n<ul>\n<li>\n<p>\u0412\u044b\u043f\u043e\u043b\u043d\u044f\u0435\u0442 \u0430\u0442\u043e\u043c\u0430\u0440\u043d\u0443\u044e \u043e\u043f\u0435\u0440\u0430\u0446\u0438\u044e \u0441\u0440\u0430\u0432\u043d\u0435\u043d\u0438\u044f \u0438 \u0437\u0430\u043c\u0435\u043d\u044b \u043f\u043e \u0443\u043a\u0430\u0437\u0430\u043d\u043d\u043e\u043c\u0443 \u0430\u0434\u0440\u0435\u0441\u0443 \u043f\u0430\u043c\u044f\u0442\u0438.<\/p>\n<\/li>\n<\/ul>\n<\/li>\n<li>\n<p><strong>atomic_max<\/strong><\/p>\n<ul>\n<li>\n<p>\u0412\u044b\u043f\u043e\u043b\u043d\u044f\u0435\u0442 \u0430\u0442\u043e\u043c\u0430\u0440\u043d\u043e\u0435 \u043d\u0430\u0445\u043e\u0436\u0434\u0435\u043d\u0438\u0435 \u043c\u0430\u043a\u0441\u0438\u043c\u0443\u043c\u0430 \u043f\u043e \u0443\u043a\u0430\u0437\u0430\u043d\u043d\u043e\u043c\u0443 \u0430\u0434\u0440\u0435\u0441\u0443 \u043f\u0430\u043c\u044f\u0442\u0438.<\/p>\n<\/li>\n<\/ul>\n<\/li>\n<li>\n<p><strong>atomic_min<\/strong><\/p>\n<ul>\n<li>\n<p>\u0412\u044b\u043f\u043e\u043b\u043d\u044f\u0435\u0442 \u0430\u0442\u043e\u043c\u0430\u0440\u043d\u043e\u0435 \u043d\u0430\u0445\u043e\u0436\u0434\u0435\u043d\u0438\u0435 \u043c\u0438\u043d\u0438\u043c\u0443\u043c\u0430 \u043f\u043e \u0443\u043a\u0430\u0437\u0430\u043d\u043d\u043e\u043c\u0443 \u0430\u0434\u0440\u0435\u0441\u0443 \u043f\u0430\u043c\u044f\u0442\u0438.<\/p>\n<\/li>\n<\/ul>\n<\/li>\n<li>\n<p><strong>atomic_or<\/strong><\/p>\n<ul>\n<li>\n<p>\u0412\u044b\u043f\u043e\u043b\u043d\u044f\u0435\u0442 \u0430\u0442\u043e\u043c\u0430\u0440\u043d\u0443\u044e \u043b\u043e\u0433\u0438\u0447\u0435\u0441\u043a\u0443\u044e \u043e\u043f\u0435\u0440\u0430\u0446\u0438\u044e \u0418\u041b\u0418 \u043f\u043e \u0443\u043a\u0430\u0437\u0430\u043d\u043d\u043e\u043c\u0443 \u0430\u0434\u0440\u0435\u0441\u0443 \u043f\u0430\u043c\u044f\u0442\u0438.<\/p>\n<\/li>\n<\/ul>\n<\/li>\n<li>\n<p><strong>atomic_xchg<\/strong><\/p>\n<ul>\n<li>\n<p>\u0412\u044b\u043f\u043e\u043b\u043d\u044f\u0435\u0442 \u0430\u0442\u043e\u043c\u0430\u0440\u043d\u044b\u0439 \u043e\u0431\u043c\u0435\u043d \u0437\u043d\u0430\u0447\u0435\u043d\u0438\u044f\u043c\u0438 \u043f\u043e \u0443\u043a\u0430\u0437\u0430\u043d\u043d\u043e\u043c\u0443 \u0430\u0434\u0440\u0435\u0441\u0443 \u043f\u0430\u043c\u044f\u0442\u0438.<\/p>\n<\/li>\n<\/ul>\n<\/li>\n<li>\n<p><strong>atomic_xor<\/strong><\/p>\n<ul>\n<li>\n<p>\u0412\u044b\u043f\u043e\u043b\u043d\u044f\u0435\u0442 \u0430\u0442\u043e\u043c\u0430\u0440\u043d\u0443\u044e \u043b\u043e\u0433\u0438\u0447\u0435\u0441\u043a\u0443\u044e \u043e\u043f\u0435\u0440\u0430\u0446\u0438\u044e \u0438\u0441\u043a\u043b\u044e\u0447\u0430\u044e\u0449\u0435\u0435 \u0418\u041b\u0418 \u043f\u043e \u0443\u043a\u0430\u0437\u0430\u043d\u043d\u043e\u043c\u0443 \u0430\u0434\u0440\u0435\u0441\u0443 \u043f\u0430\u043c\u044f\u0442\u0438.<\/p>\n<\/li>\n<\/ul>\n<\/li>\n<\/ol>\n<p>\u0422\u0430\u043a\u0436\u0435 \u0435\u0441\u0442\u044c \u043e\u043f\u0435\u0440\u0430\u0446\u0438\u0438 \u0441\u043a\u0430\u043d\u0438\u0440\u043e\u0432\u0430\u043d\u0438\u044f \u0438 \u0441\u043e\u0440\u0442\u0438\u0440\u043e\u0432\u043a\u0438:<\/p>\n<ol>\n<li>\n<p><strong>associative_scan<\/strong><\/p>\n<ul>\n<li>\n<p>\u041f\u0440\u0438\u043c\u0435\u043d\u044f\u0435\u0442 \u0444\u0443\u043d\u043a\u0446\u0438\u044e <code>combine_fn<\/code> \u043a \u043a\u0430\u0436\u0434\u043e\u043c\u0443 \u044d\u043b\u0435\u043c\u0435\u043d\u0442\u0443 \u0441 \u0441\u043e\u0445\u0440\u0430\u043d\u0435\u043d\u0438\u0435\u043c \u043f\u0440\u043e\u043c\u0435\u0436\u0443\u0442\u043e\u0447\u043d\u043e\u0433\u043e \u0437\u043d\u0430\u0447\u0435\u043d\u0438\u044f <code>carry<\/code> \u0432 \u0442\u0435\u043d\u0437\u043e\u0440\u0430\u0445 \u0432\u0434\u043e\u043b\u044c \u0443\u043a\u0430\u0437\u0430\u043d\u043d\u043e\u0439 \u043e\u0441\u0438 \u0438 \u043e\u0431\u043d\u043e\u0432\u043b\u044f\u0435\u0442 <code>carry<\/code>.<\/p>\n<\/li>\n<\/ul>\n<\/li>\n<li>\n<p><strong>cumprod<\/strong><\/p>\n<ul>\n<li>\n<p>\u0412\u043e\u0437\u0432\u0440\u0430\u0449\u0430\u0435\u0442 \u043a\u0443\u043c\u0443\u043b\u044f\u0442\u0438\u0432\u043d\u043e\u0435 \u043f\u0440\u043e\u0438\u0437\u0432\u0435\u0434\u0435\u043d\u0438\u0435 \u0432\u0441\u0435\u0445 \u044d\u043b\u0435\u043c\u0435\u043d\u0442\u043e\u0432 \u0432 \u0442\u0435\u043d\u0437\u043e\u0440\u0435 \u0432\u0434\u043e\u043b\u044c \u0443\u043a\u0430\u0437\u0430\u043d\u043d\u043e\u0439 \u043e\u0441\u0438.<\/p>\n<\/li>\n<\/ul>\n<\/li>\n<li>\n<p><strong>cumsum<\/strong><\/p>\n<ul>\n<li>\n<p>\u0412\u043e\u0437\u0432\u0440\u0430\u0449\u0430\u0435\u0442 \u043a\u0443\u043c\u0443\u043b\u044f\u0442\u0438\u0432\u043d\u0443\u044e \u0441\u0443\u043c\u043c\u0443 \u0432\u0441\u0435\u0445 \u044d\u043b\u0435\u043c\u0435\u043d\u0442\u043e\u0432 \u0432 \u0442\u0435\u043d\u0437\u043e\u0440\u0435 \u0432\u0434\u043e\u043b\u044c \u0443\u043a\u0430\u0437\u0430\u043d\u043d\u043e\u0439 \u043e\u0441\u0438.<\/p>\n<\/li>\n<\/ul>\n<\/li>\n<li>\n<p><strong>histogram<\/strong><\/p>\n<ul>\n<li>\n<p>\u0412\u044b\u0447\u0438\u0441\u043b\u044f\u0435\u0442 \u0433\u0438\u0441\u0442\u043e\u0433\u0440\u0430\u043c\u043c\u0443 \u043d\u0430 \u043e\u0441\u043d\u043e\u0432\u0435 \u0432\u0445\u043e\u0434\u043d\u043e\u0433\u043e \u0442\u0435\u043d\u0437\u043e\u0440\u0430 \u0441 \u0437\u0430\u0434\u0430\u043d\u043d\u044b\u043c \u0447\u0438\u0441\u043b\u043e\u043c \u043a\u043e\u0440\u0437\u0438\u043d <code>num_bins<\/code>, \u043a\u043e\u0440\u0437\u0438\u043d\u044b \u0438\u043c\u0435\u044e\u0442 \u0448\u0438\u0440\u0438\u043d\u0443 1 \u0438 \u043d\u0430\u0447\u0438\u043d\u0430\u044e\u0442\u0441\u044f \u0441 0.<\/p>\n<\/li>\n<\/ul>\n<\/li>\n<li>\n<p><strong>sort<\/strong><\/p>\n<ul>\n<li>\n<p>\u0421\u043e\u0440\u0442\u0438\u0440\u0443\u0435\u0442 \u044d\u043b\u0435\u043c\u0435\u043d\u0442\u044b \u0442\u0435\u043d\u0437\u043e\u0440\u0430. \u041c\u043e\u0436\u0435\u0442 \u0431\u044b\u0442\u044c \u0432\u044b\u0437\u0432\u0430\u043d\u0430 \u043a\u0430\u043a \u043c\u0435\u0442\u043e\u0434 \u0442\u0435\u043d\u0437\u043e\u0440\u0430 <code>x.sort(...)<\/code>.<\/p>\n<\/li>\n<\/ul>\n<\/li>\n<\/ol>\n<p>\u0418 \u0441\u0442\u0430\u043d\u0434\u0430\u0440\u0442\u043d\u044b\u0435 \u043e\u043f\u0435\u0440\u0430\u0446\u0438\u0438 \u0440\u0435\u0434\u0443\u043a\u0446\u0438\u0438:<\/p>\n<ol>\n<li>\n<p><strong>argmax<\/strong><\/p>\n<ul>\n<li>\n<p>\u0412\u043e\u0437\u0432\u0440\u0430\u0449\u0430\u0435\u0442 \u0438\u043d\u0434\u0435\u043a\u0441 \u043c\u0430\u043a\u0441\u0438\u043c\u0430\u043b\u044c\u043d\u043e\u0433\u043e \u044d\u043b\u0435\u043c\u0435\u043d\u0442\u0430 \u0432 \u0442\u0435\u043d\u0437\u043e\u0440\u0435 \u0432\u0434\u043e\u043b\u044c \u0443\u043a\u0430\u0437\u0430\u043d\u043d\u043e\u0439 \u043e\u0441\u0438.<\/p>\n<\/li>\n<\/ul>\n<\/li>\n<li>\n<p><strong>argmin<\/strong><\/p>\n<ul>\n<li>\n<p>\u0412\u043e\u0437\u0432\u0440\u0430\u0449\u0430\u0435\u0442 \u0438\u043d\u0434\u0435\u043a\u0441 \u043c\u0438\u043d\u0438\u043c\u0430\u043b\u044c\u043d\u043e\u0433\u043e \u044d\u043b\u0435\u043c\u0435\u043d\u0442\u0430 \u0432 \u0442\u0435\u043d\u0437\u043e\u0440\u0435 \u0432\u0434\u043e\u043b\u044c \u0443\u043a\u0430\u0437\u0430\u043d\u043d\u043e\u0439 \u043e\u0441\u0438.<\/p>\n<\/li>\n<\/ul>\n<\/li>\n<li>\n<p><strong>max<\/strong><\/p>\n<ul>\n<li>\n<p>\u0412\u043e\u0437\u0432\u0440\u0430\u0449\u0430\u0435\u0442 \u043c\u0430\u043a\u0441\u0438\u043c\u0430\u043b\u044c\u043d\u043e\u0435 \u0437\u043d\u0430\u0447\u0435\u043d\u0438\u0435 \u0441\u0440\u0435\u0434\u0438 \u0432\u0441\u0435\u0445 \u044d\u043b\u0435\u043c\u0435\u043d\u0442\u043e\u0432 \u0442\u0435\u043d\u0437\u043e\u0440\u0430 \u0432\u0434\u043e\u043b\u044c \u0443\u043a\u0430\u0437\u0430\u043d\u043d\u043e\u0439 \u043e\u0441\u0438.<\/p>\n<\/li>\n<\/ul>\n<\/li>\n<li>\n<p><strong>min<\/strong><\/p>\n<ul>\n<li>\n<p>\u0412\u043e\u0437\u0432\u0440\u0430\u0449\u0430\u0435\u0442 \u043c\u0438\u043d\u0438\u043c\u0430\u043b\u044c\u043d\u043e\u0435 \u0437\u043d\u0430\u0447\u0435\u043d\u0438\u0435 \u0441\u0440\u0435\u0434\u0438 \u0432\u0441\u0435\u0445 \u044d\u043b\u0435\u043c\u0435\u043d\u0442\u043e\u0432 \u0442\u0435\u043d\u0437\u043e\u0440\u0430 \u0432\u0434\u043e\u043b\u044c \u0443\u043a\u0430\u0437\u0430\u043d\u043d\u043e\u0439 \u043e\u0441\u0438.<\/p>\n<\/li>\n<\/ul>\n<\/li>\n<li>\n<p><strong>reduce<\/strong><\/p>\n<ul>\n<li>\n<p>\u041f\u0440\u0438\u043c\u0435\u043d\u044f\u0435\u0442 \u0444\u0443\u043d\u043a\u0446\u0438\u044e \u043a\u043e\u043c\u0431\u0438\u043d\u0438\u0440\u043e\u0432\u0430\u043d\u0438\u044f <code>combine_fn<\/code> \u043a\u043e \u0432\u0441\u0435\u043c \u044d\u043b\u0435\u043c\u0435\u043d\u0442\u0430\u043c \u0432\u0445\u043e\u0434\u043d\u044b\u0445 \u0442\u0435\u043d\u0437\u043e\u0440\u043e\u0432 \u0432\u0434\u043e\u043b\u044c \u0443\u043a\u0430\u0437\u0430\u043d\u043d\u043e\u0439 \u043e\u0441\u0438.<\/p>\n<\/li>\n<\/ul>\n<\/li>\n<\/ol>\n<h4>\u0418\u0442\u0430\u043a, \u043a\u0430\u043a \u0432\u0441\u0435 \u044d\u0442\u043e \u043f\u0440\u0438\u043c\u0435\u043d\u044f\u0442\u044c?<\/h4>\n<p>\u041c\u043e\u0436\u043d\u043e \u044e\u0437\u0430\u0442\u044c <code>atomic_add<\/code> \u0434\u043b\u044f \u044d\u0444\u0444\u0435\u043a\u0442\u0438\u0432\u043d\u043e\u0439 \u0440\u0435\u0430\u043b\u0438\u0437\u0430\u0446\u0438\u0438 \u0433\u0440\u0430\u0434\u0438\u0435\u043d\u0442\u043d\u043e\u0433\u043e \u0441\u043f\u0443\u0441\u043a\u0430 \u0432 \u0441\u0440\u0435\u0434\u0435 \u0441 \u0432\u044b\u0441\u043e\u043a\u043e\u0439 \u0441\u0442\u0435\u043f\u0435\u043d\u044c\u044e \u043f\u0430\u0440\u0430\u043b\u043b\u0435\u043b\u0438\u0437\u043c\u0430. \u041d\u0430\u043f\u0440\u0438\u043c\u0435\u0440, \u043f\u0440\u0438 \u043e\u0431\u0443\u0447\u0435\u043d\u0438\u0438 \u0431\u043e\u043b\u044c\u0448\u0438\u0445 \u043d\u0435\u0439\u0440\u043e\u043d\u043d\u044b\u0445 \u0441\u0435\u0442\u0435\u0439 \u043d\u0430 GPU, \u0430\u0442\u043e\u043c\u0430\u0440\u043d\u044b\u0435 \u043e\u043f\u0435\u0440\u0430\u0446\u0438\u0438 \u043c\u043e\u0433\u0443\u0442 \u043e\u0431\u0435\u0441\u043f\u0435\u0447\u0438\u0442\u044c \u043e\u0431\u043d\u043e\u0432\u043b\u0435\u043d\u0438\u0435 \u0432\u0435\u0441\u043e\u0432 \u0431\u0435\u0437 \u043a\u043e\u043d\u0444\u043b\u0438\u043a\u0442\u043e\u0432 \u043c\u0435\u0436\u0434\u0443 \u043f\u043e\u0442\u043e\u043a\u0430\u043c\u0438:<\/p>\n<pre><code class=\"python\">@triton.jit def update_weights(grads, weights, learning_rate):     pid = tl.program_id(0)     for i in range(pid, grads.shape[0], tl.num_programs()):         tl.atomic_add(weights, i, -learning_rate * grads[i])<\/code><\/pre>\n<p>\u041c\u043e\u0436\u043d\u043e \u0442\u0430\u043a\u0436\u0435 \u044e\u0437\u0430\u0442\u044c <code>sort<\/code> \u0438 <code>reduce<\/code> \u0434\u043b\u044f \u0440\u0435\u0430\u043b\u0438\u0437\u0430\u0446\u0438\u0438 \u0430\u043b\u0433\u043e\u0440\u0438\u0442\u043c\u043e\u0432 \u043a\u043b\u0430\u0441\u0442\u0435\u0440\u0438\u0437\u0430\u0446\u0438\u0438, \u0442\u0430\u043a\u0438\u0445 \u043a\u0430\u043a K-means, \u0433\u0434\u0435 \u043d\u0435\u043e\u0431\u0445\u043e\u0434\u0438\u043c\u043e \u0441\u043e\u0440\u0442\u0438\u0440\u043e\u0432\u0430\u0442\u044c \u0434\u0430\u043d\u043d\u044b\u0435 \u0438 \u0432\u044b\u0447\u0438\u0441\u043b\u044f\u0442\u044c \u0446\u0435\u043d\u0442\u0440\u043e\u0438\u0434\u044b:<\/p>\n<pre><code class=\"python\">@triton.jit def k_means_update(x, centroids, labels, num_clusters):     # \u0432\u044b\u0447\u0438\u0441\u043b\u0435\u043d\u0438\u0435 \u0440\u0430\u0441\u0441\u0442\u043e\u044f\u043d\u0438\u0439 \u0438 \u043f\u0440\u0438\u0441\u0432\u043e\u0435\u043d\u0438\u0435 \u043c\u0435\u0442\u043e\u043a     block_start = tl.program_id(0) * tl.num_programs()     distances = tl.zeros((num_clusters,), dtype=tl.float32)     for i in range(block_start, min(block_start + BLOCK_SIZE, x.shape[0])):         for c in range(num_clusters):             distances[c] = tl.sum(tl.pow(x[i] - centroids[c], 2))         labels[i] = tl.argmin(distances)     # \u043e\u0431\u043d\u043e\u0432\u043b\u0435\u043d\u0438\u0435 \u0446\u0435\u043d\u0442\u0440\u043e\u0438\u0434\u043e\u0432     for c in range(num_clusters):         assigned_pts = x[labels == c]         centroids[c] = tl.sum(assigned_pts, axis=0) \/ len(assigned_pts)<\/code><\/pre>\n<p>\u0410 \u0432\u043e\u0442 <code>exp<\/code>, <code>max<\/code>, \u0438 <code>log<\/code> \u043c\u043e\u0436\u043d\u043e \u044e\u0437\u0430\u0442\u044c \u043d\u0430\u043f\u0440\u0438\u043c\u0435\u0440, \u0434\u043b\u044f \u0444\u0443\u043d\u0446\u0438\u0439 <em>softmax<\/em>:<\/p>\n<pre><code class=\"python\">@triton.jit def softmax(x):     max_val = tl.max(x, axis=1, keepdims=True)     exp_x = tl.exp(x - max_val)     sum_exp_x = tl.sum(exp_x, axis=1, keepdims=True)     return exp_x \/ sum_exp_x<\/code><\/pre>\n<p>\u041c\u043e\u0436\u043d\u043e \u0438\u0441\u043f\u043e\u043b\u044c\u0437\u043e\u0432\u0430\u0442\u044c <code>cumsum<\/code> \u0434\u043b\u044f \u0440\u0430\u0441\u0447\u0451\u0442\u0430 \u0441\u043a\u043e\u043b\u044c\u0437\u044f\u0449\u0435\u0433\u043e \u0441\u0440\u0435\u0434\u043d\u0435\u0433\u043e (\u0447\u0442\u043e \u043f\u0440\u0438\u043c\u0435\u043d\u0438\u043c\u043e \u0432\u043e \u0432\u0440\u0435\u043c\u0435\u043d\u043d\u044b\u0445 \u0440\u044f\u0434\u0430\u0445):<\/p>\n<pre><code class=\"python\">@triton.jit def moving_average(data, window_size):     cum_sum = tl.cumsum(data, axis=0)     return (cum_sum[window_size:] - cum_sum[:-window_size]) \/ window_size<\/code><\/pre>\n<p><code>associative_scan<\/code> \u043f\u043e\u0434\u0445\u043e\u0434\u0438\u0442 \u0434\u043b\u044f \u0440\u0435\u0430\u043b\u0438\u0437\u0430\u0446\u0438\u0438 \u0430\u043b\u0433\u043e\u0440\u0438\u0442\u043c\u0430 \u043f\u0440\u044f\u043c\u043e\u0433\u043e \u0440\u0430\u0441\u043f\u0440\u043e\u0441\u0442\u0440\u0430\u043d\u0435\u043d\u0438\u044f \u0432 RNN, \u0433\u0434\u0435 \u043f\u043e\u0441\u043b\u0435\u0434\u043e\u0432\u0430\u0442\u0435\u043b\u044c\u043d\u044b\u0435 \u0437\u0430\u0432\u0438\u0441\u0438\u043c\u043e\u0441\u0442\u0438 \u043c\u043e\u0433\u0443\u0442 \u0431\u044b\u0442\u044c \u044d\u0444\u0444\u0435\u043a\u0442\u0438\u0432\u043d\u043e \u043e\u0431\u0440\u0430\u0431\u043e\u0442\u0430\u043d\u044b \u0441 \u043f\u043e\u043c\u043e\u0449\u044c\u044e \u044d\u0442\u043e\u0439 \u043e\u043f\u0435\u0440\u0430\u0446\u0438\u0438:<\/p>\n<pre><code class=\"python\"> @triton.jit def rnn_step(hidden, input, weights):     carry = tl.zeros_like(hidden)     for t in range(input.shape[0]):         carry = tl.associative_scan(lambda h, x: tl.tanh(x @ weights + h), input[t], axis=0)         hidden[t] = carry     return hidden<\/code><\/pre>\n<p>\u0422\u0430\u043a\u0436\u0435 \u043c\u043e\u0436\u043d\u043e \u043f\u0440\u0438\u043c\u0435\u043d\u0438\u0442\u044c <code>associative_scan<\/code> \u0434\u043b\u044f \u0440\u0435\u0430\u043b\u0438\u0437\u0430\u0446\u0438\u0438 \u043f\u0430\u0440\u0430\u043b\u043b\u0435\u043b\u044c\u043d\u043e\u0433\u043e \u0430\u043b\u0433\u043e\u0440\u0438\u0442\u043c\u0430 \u0431\u044b\u0441\u0442\u0440\u043e\u0433\u043e \u043f\u0440\u0435\u043e\u0431\u0440\u0430\u0437\u043e\u0432\u0430\u043d\u0438\u044f \u0424\u0443\u0440\u044c\u0435:<\/p>\n<pre><code class=\"python\">@triton.jit def fft_step(data, step_size):     idx = tl.program_id(0) * step_size * 2 + tl.arange(0, step_size)     t = exp(-2j * pi \/ (2 * step_size) * idx)     u = data[idx]     v = data[idx + step_size] * t     data[idx] = u + v     data[idx + step_size] = u - v<\/code><\/pre>\n<hr\/>\n<p>\u0411\u043e\u043b\u0435\u0435 \u043f\u043e\u0434\u0440\u043e\u0431\u043d\u043e \u0441 Triton <a href=\"https:\/\/triton-lang.org\/\">\u043c\u043e\u0436\u043d\u043e \u043e\u0437\u043d\u0430\u043a\u043e\u043c\u0438\u0442\u044c\u0441\u044f \u0437\u0434\u0435\u0441\u044c.<\/a><\/p>\n<p>\u041c\u0430\u0442\u0435\u0440\u0438\u0430\u043b \u043f\u043e\u0434\u0433\u043e\u0442\u043e\u0432\u043b\u0435\u043d \u0432 \u043f\u0440\u0435\u0434\u0434\u0432\u0435\u0440\u0438\u0438 \u0441\u0442\u0430\u0440\u0442\u0430 <a href=\"https:\/\/otus.pw\/H4dT\/\">\u0441\u043f\u0435\u0446\u0438\u0430\u043b\u0438\u0437\u0430\u0446\u0438\u0438 &#171;Machine Learning&#187;.<\/a><\/p>\n<\/p>\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\/807655\/\"> https:\/\/habr.com\/ru\/articles\/807655\/<\/a><br \/><\/br><\/br><\/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-373625","post","type-post","status-publish","format-standard","hentry"],"_links":{"self":[{"href":"https:\/\/savepearlharbor.com\/index.php?rest_route=\/wp\/v2\/posts\/373625","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=373625"}],"version-history":[{"count":0,"href":"https:\/\/savepearlharbor.com\/index.php?rest_route=\/wp\/v2\/posts\/373625\/revisions"}],"wp:attachment":[{"href":"https:\/\/savepearlharbor.com\/index.php?rest_route=%2Fwp%2Fv2%2Fmedia&parent=373625"}],"wp:term":[{"taxonomy":"category","embeddable":true,"href":"https:\/\/savepearlharbor.com\/index.php?rest_route=%2Fwp%2Fv2%2Fcategories&post=373625"},{"taxonomy":"post_tag","embeddable":true,"href":"https:\/\/savepearlharbor.com\/index.php?rest_route=%2Fwp%2Fv2%2Ftags&post=373625"}],"curies":[{"name":"wp","href":"https:\/\/api.w.org\/{rel}","templated":true}]}}