{"id":496918,"date":"2026-10-01T12:40:48","date_gmt":"2026-10-01T12:40:48","guid":{"rendered":"https:\/\/savepearlharbor.com\/?p=496918"},"modified":"-0001-11-30T00:00:00","modified_gmt":"-0001-11-29T21:00:00","slug":"","status":"publish","type":"post","link":"https:\/\/savepearlharbor.com\/?p=496918","title":{"rendered":"ROUGE\u2011V: \u043a\u0430\u043a\u00a0\u044f \u0442\u0440\u0430\u043d\u0441\u043b\u0438\u0440\u043e\u0432\u0430\u043b CUDA\/PTX \u0432\u00a0LLVM IR \u0438 \u0437\u0430\u043f\u0443\u0441\u043a\u0430\u043b \u0435\u0433\u043e \u0431\u0435\u0437\u00a0CUDA runtime"},"content":{"rendered":"<div xmlns=\"http:\/\/www.w3.org\/1999\/xhtml\">\n<blockquote>\n<p><strong>TL;DR:<\/strong> \u041f\u0440\u043e\u0431\u043b\u0435\u043c\u0430 \u043f\u0435\u0440\u0435\u043d\u043e\u0441\u0438\u043c\u043e\u0441\u0442\u0438 CUDA \u0434\u043b\u044f\u00a0\u043c\u0435\u043d\u044f \u043e\u043a\u0430\u0437\u0430\u043b\u0430\u0441\u044c \u043d\u0435\u00a0\u0441\u0442\u043e\u043b\u044c\u043a\u043e \u0432\u00a0\u0441\u0430\u043c\u0438\u0445 \u0430\u0440\u0438\u0444\u043c\u0435\u0442\u0438\u0447\u0435\u0441\u043a\u0438\u0445 \u043e\u043f\u0435\u0440\u0430\u0446\u0438\u044f\u0445, \u0441\u043a\u043e\u043b\u044c\u043a\u043e \u0432\u00a0\u043e\u0433\u0440\u043e\u043c\u043d\u043e\u043c \u043a\u043e\u043b\u0438\u0447\u0435\u0441\u0442\u0432\u0435 \u0438\u043d\u0444\u0440\u0430\u0441\u0442\u0440\u0443\u043a\u0442\u0443\u0440\u044b \u0432\u043e\u043a\u0440\u0443\u0433 \u043d\u0438\u0445. \u041f\u043e\u044d\u0442\u043e\u043c\u0443 \u0432\u043c\u0435\u0441\u0442\u043e \u043f\u043e\u043f\u044b\u0442\u043a\u0438 \u043f\u043e\u0432\u0442\u043e\u0440\u044f\u0442\u044c CUDA runtime \u044f \u0440\u0435\u0448\u0438\u043b \u043f\u0435\u0440\u0435\u0445\u0432\u0430\u0442\u0438\u0442\u044c \u043f\u0440\u043e\u0433\u0440\u0430\u043c\u043c\u0443 \u0440\u0430\u043d\u044c\u0448\u0435\u00a0\u2014 \u043d\u0430\u00a0\u0443\u0440\u043e\u0432\u043d\u0435 PTX. \u042f \u043d\u0430\u043f\u0438\u0441\u0430\u043b ROUGE\u2011V\u00a0\u2014 \u043e\u0442\u043a\u0440\u044b\u0442\u044b\u0439 AOT\u2011\u043a\u043e\u043c\u043f\u0438\u043b\u044f\u0442\u043e\u0440, \u0440\u0430\u0437\u0440\u0430\u0431\u043e\u0442\u0430\u043d\u043d\u044b\u0439 \u043f\u043e\u00a0\u043f\u0440\u0438\u043d\u0446\u0438\u043f\u0443 clean\u2011room: \u0431\u0435\u0440\u0451\u0442 PTX, \u0441\u0433\u0435\u043d\u0435\u0440\u0438\u0440\u043e\u0432\u0430\u043d\u043d\u044b\u0439 \u043e\u0444\u0438\u0446\u0438\u0430\u043b\u044c\u043d\u044b\u043c <code>nvcc<\/code> 13.4, \u0442\u0440\u0430\u043d\u0441\u043b\u0438\u0440\u0443\u0435\u0442 \u0435\u0433\u043e \u0432\u00a0LLVM IR, \u0434\u0430\u043b\u044c\u0448\u0435 \u0441\u043e\u0431\u0438\u0440\u0430\u0435\u0442 \u043e\u0431\u044b\u0447\u043d\u044b\u0439 <code>clang<\/code>. \u0418\u0441\u043f\u043e\u043b\u043d\u044f\u0435\u0442\u0441\u044f \u043d\u0430\u00a0x86-64\u00a0\u0438 \u043d\u0430\u00a0\u0436\u0438\u0432\u043e\u0439 GeForce MX450\u00a0(\u043f\u043e\u0431\u0438\u0442\u043e\u0432\u043e \u0441\u00a0\u0445\u043e\u0441\u0442\u2011\u044d\u0442\u0430\u043b\u043e\u043d\u043e\u043c); AMD \u0438 RISC\u2011V\u00a0\u2014 \u043f\u043e\u043a\u0430 \u0442\u043e\u043b\u044c\u043a\u043e \u0441\u0431\u043e\u0440\u043a\u0430, \u0431\u0435\u0437\u00a0\u0432\u0440\u0430\u043d\u044c\u044f \u043f\u0440\u043e \u00ab\u043b\u044e\u0431\u043e\u0435\u00a0\u0436\u0435\u043b\u0435\u0437\u043e\u00bb.<\/p>\n<p>\u041c\u043e\u0439 GitHub: <a href=\"https:\/\/github.com\/MrModelOS\/ROUGE-V\" rel=\"noopener nofollow\">github.com\/MrModelOS\/ROUGE\u2011V<\/a> (\u0430\u0432\u0442\u043e\u0440: MrModel).<\/p>\n<p>\u041f\u043e\u043a\u0440\u044b\u0442\u0438\u0435: <code>ctest<\/code>\u00a0\u2014 21\/21; 8\u00a0\u044f\u0434\u0435\u0440 (7\u00a0\u043a\u0430\u043d\u043e\u043d\u0438\u0447\u0435\u0441\u043a\u0438\u0445 + 1\u00a0\u043d\u0430\u0441\u0442\u043e\u044f\u0449\u0435\u0435 \u043e\u0442 <code>nvcc<\/code>) \u0438\u0441\u043f\u043e\u043b\u043d\u0435\u043d\u044b \u043d\u0430\u00a0GeForce MX450.<\/p>\n<\/blockquote>\n<h3>1. \u0417\u0430\u0447\u0435\u043c \u044f \u0441\u0434\u0435\u043b\u0430\u043b \u0441\u0432\u043e\u0439 \u0442\u0440\u0430\u043d\u0441\u043b\u044f\u0442\u043e\u0440, \u043a\u043e\u0433\u0434\u0430 \u0435\u0441\u0442\u044c \u0434\u0440\u0443\u0433\u0438\u0435 \u043f\u0440\u043e\u0435\u043a\u0442\u044b?<\/h3>\n<p>\u041e\u0434\u0438\u043d \u0438\u0437\u00a0\u043e\u0447\u0435\u0432\u0438\u0434\u043d\u044b\u0445 \u0441\u043f\u043e\u0441\u043e\u0431\u043e\u0432 \u0431\u043e\u0440\u043e\u0442\u044c\u0441\u044f \u0441\u00a0vendor lock\u2011in NVIDIA\u00a0\u2014 \u0440\u0435\u0438\u043c\u043f\u043b\u0435\u043c\u0435\u043d\u0442\u0438\u0440\u043e\u0432\u0430\u0442\u044c runtime \u0438\u043b\u0438\u00a0\u043f\u0435\u0440\u0435\u0445\u0432\u0430\u0442\u044b\u0432\u0430\u0442\u044c API \u043d\u0430\u00a0\u043b\u0435\u0442\u0443. \u042d\u0442\u043e \u0432\u0435\u0447\u043d\u0430\u044f \u0433\u043e\u043d\u043a\u0430 \u0437\u0430\u00a0\u0447\u0443\u0436\u0438\u043c \u0438\u043d\u0442\u0435\u0440\u0444\u0435\u0439\u0441\u043e\u043c: \u0432\u0435\u043d\u0434\u043e\u0440 \u043f\u043e\u043c\u0435\u043d\u044f\u043b \u043f\u043e\u0432\u0435\u0434\u0435\u043d\u0438\u0435\u00a0\u2014 \u0442\u044b \u0434\u043e\u0433\u043e\u043d\u044f\u0435\u0448\u044c.<\/p>\n<p>\u042f \u043f\u043e\u0448\u0451\u043b \u043f\u043e\u00a0\u043f\u0443\u0442\u0438 Clean\u2011Room \u0440\u0430\u0437\u0440\u0430\u0431\u043e\u0442\u043a\u0438.<\/p>\n<p>\u041a\u043e\u043c\u043f\u0438\u043b\u044f\u0442\u043e\u0440 <code>nvcc<\/code> \u0433\u0435\u043d\u0435\u0440\u0438\u0440\u0443\u0435\u0442 \u043d\u0435\u00a0\u043c\u0430\u0448\u0438\u043d\u043d\u044b\u0439 \u043a\u043e\u0434, \u0430\u00a0PTX\u00a0\u2014 \u0442\u0435\u043a\u0441\u0442\u043e\u0432\u044b\u0439 \u043f\u0440\u043e\u043c\u0435\u0436\u0443\u0442\u043e\u0447\u043d\u044b\u0439 \u0430\u0441\u0441\u0435\u043c\u0431\u043b\u0435\u0440 \u0441\u00a0\u043f\u0443\u0431\u043b\u0438\u0447\u043d\u043e \u043e\u043f\u0443\u0431\u043b\u0438\u043a\u043e\u0432\u0430\u043d\u043d\u043e\u0439 \u0441\u043f\u0435\u0446\u0438\u0444\u0438\u043a\u0430\u0446\u0438\u0435\u0439 (\u0432\u044b\u0445\u043b\u043e\u043f \u043d\u0430\u0448\u0435\u0433\u043e <code>nvcc<\/code> 13.4\u00a0\u2014 PTX ISA 9.4). \u042f \u0432\u044b\u0431\u0440\u0430\u043b \u0435\u0433\u043e \u043a\u0430\u043a\u00a0\u0442\u043e\u0447\u043a\u0443 \u043f\u0435\u0440\u0435\u0445\u0432\u0430\u0442\u0430:<\/p>\n<ol>\n<li>\n<p>\u0411\u0435\u0440\u0443 \u0441\u0442\u0430\u043d\u0434\u0430\u0440\u0442\u043d\u044b\u0439 <code>.cu<\/code>\u2011\u043a\u043e\u0434 \u0438 \u043a\u043e\u043c\u043f\u0438\u043b\u0438\u0440\u0443\u044e \u0435\u0433\u043e \u0447\u0435\u0440\u0435\u0437 <code>nvcc<\/code> \u0432 <code>.ptx<\/code>.<\/p>\n<\/li>\n<li>\n<p>\u041c\u043e\u0439 \u0442\u0440\u0430\u043d\u0441\u043b\u044f\u0442\u043e\u0440 <code>ptx2ir<\/code> \u043f\u0430\u0440\u0441\u0438\u0442 PTX \u0438 \u0441\u0442\u0440\u043e\u0438\u0442 \u044d\u043a\u0432\u0438\u0432\u0430\u043b\u0435\u043d\u0442\u043d\u043e\u0435 \u043f\u0440\u0435\u0434\u0441\u0442\u0430\u0432\u043b\u0435\u043d\u0438\u0435 \u0432\u00a0LLVM IR.<\/p>\n<\/li>\n<li>\n<p>\u0421\u0442\u0430\u043d\u0434\u0430\u0440\u0442\u043d\u044b\u0439 <code>clang<\/code> \/ LLVM \u0441\u043e\u0431\u0438\u0440\u0430\u0435\u0442 \u043a\u043e\u043d\u0435\u0447\u043d\u044b\u0439 \u0431\u0438\u043d\u0430\u0440\u043d\u0438\u043a \u043f\u043e\u0434\u00a0\u0446\u0435\u043b\u0435\u0432\u0443\u044e \u0430\u0440\u0445\u0438\u0442\u0435\u043a\u0442\u0443\u0440\u0443.<\/p>\n<\/li>\n<\/ol>\n<p>\u041d\u0438\u043a\u0430\u043a\u0438\u0445 \u043f\u0440\u043e\u043f\u0440\u0438\u0435\u0442\u0430\u0440\u043d\u044b\u0445 \u0431\u0438\u043d\u0430\u0440\u043d\u0438\u043a\u043e\u0432 \u0438 \u0437\u0430\u0433\u043e\u043b\u043e\u0432\u043a\u043e\u0432 NVIDIA \u044f \u043d\u0435\u00a0\u043a\u043e\u043f\u0438\u0440\u0443\u044e \u0438 \u043d\u0435\u00a0\u0440\u0435\u0432\u0435\u0440\u0441\u0438\u0440\u0443\u044e\u00a0\u2014 \u043f\u0440\u0430\u0432\u0438\u043b\u043e clean\u2011room \u0437\u0430\u043a\u0440\u0435\u043f\u043b\u0435\u043d\u043e \u0432 <a href=\"https:\/\/github.com\/MrModelOS\/ROUGE-V\/blob\/v0.1\/CONTRIBUTING.md\" rel=\"noopener nofollow\"><code>CONTRIBUTING.md<\/code><\/a>.<\/p>\n<h3>2. \u0410\u0440\u0445\u0438\u0442\u0435\u043a\u0442\u0443\u0440\u0430 \u0442\u0440\u0430\u043d\u0441\u043b\u044f\u0442\u043e\u0440\u0430 \u0438 4\u00a0\u0448\u0430\u0433\u0430 \u0434\u043e\u00a0\u0437\u0430\u043f\u0443\u0441\u043a\u0430<\/h3>\n<p>\u041f\u0430\u0439\u043f\u043b\u0430\u0439\u043d \u0443\u00a0\u043c\u0435\u043d\u044f \u0432\u044b\u0433\u043b\u044f\u0434\u0438\u0442 \u0442\u0430\u043a:<\/p>\n<pre><code>CUDA \/ .cu    \u2502    \u25bcnvcc --ptx    \u2502    \u25bcPTX    \u2502    \u25bcptx2ir    \u2502    \u25bcLLVM IR    \u2502    \u25bcLLVM \/ clang    \u2502    \u25bcnative codeNVPTX \u2192 GeForce MX450 \u2192 bit-exact check<\/code><div class=\"code-explainer\"><a href=\"https:\/\/sourcecraft.dev\/\" class=\"tm-button code-explainer__link\" style=\"visibility: hidden;\"><span class=\"sc-logo-wrapper\"><\/span><\/a><\/div><\/pre>\n<p>\u0427\u0435\u0441\u0442\u043d\u0430\u044f \u043b\u0435\u0433\u0435\u043d\u0434\u0430: AMD (<code>amdgcn<\/code>) \u0438 RISC\u2011V (<code>rv64gcv<\/code>)\u00a0\u2014 \u043a\u0440\u043e\u0441\u0441\u2011\u043a\u043e\u043c\u043f\u0438\u043b\u044f\u0446\u0438\u044f, \u0438\u0441\u043f\u043e\u043b\u043d\u0435\u043d\u0438\u044f \u043d\u0430\u00a0\u0436\u0435\u043b\u0435\u0437\u0435 \u0443\u00a0\u043c\u0435\u043d\u044f \u043d\u0435\u00a0\u0431\u044b\u043b\u043e. \u0418\u0441\u043f\u043e\u043b\u043d\u044f\u0435\u0442\u0441\u044f\u00a0\u2014 \u0445\u043e\u0441\u0442 \u0438 NVIDIA.<\/p>\n<p>\u041d\u0430\u00a0\u043f\u0440\u0430\u043a\u0442\u0438\u043a\u0435 \u0432\u0435\u0441\u044c \u043f\u0440\u043e\u0446\u0435\u0441\u0441 \u0437\u0430\u043d\u0438\u043c\u0430\u0435\u0442 4\u00a0\u043a\u043e\u043c\u0430\u043d\u0434\u044b (\u043f\u0440\u043e\u0432\u0435\u0440\u0435\u043d\u043e \u043d\u0430\u00a0\u0436\u0438\u0432\u043e\u043c \u0441\u0442\u0435\u043d\u0434\u0435):<\/p>\n<pre><code class=\"bash\"># 1. \u0421\u043a\u043e\u043c\u043f\u0438\u043b\u0438\u0440\u043e\u0432\u0430\u043b .cu \u0432 PTX \u0432\u0435\u043d\u0434\u043e\u0440\u043d\u044b\u043c nvcc 13.4nvcc -arch=sm_75 -ptx sq.cu -o sq.ptx# 2. \u0422\u0440\u0430\u043d\u0441\u043b\u0438\u0440\u043e\u0432\u0430\u043b PTX \u0432 LLVM IR \u0441\u0432\u043e\u0438\u043c ptx2irptx2ir --target nvptx64-nvidia-cuda sq.ptx sq.ll# 3. \u0421\u043e\u0431\u0440\u0430\u043b \u043d\u0430\u0441\u0442\u043e\u044f\u0449\u0438\u043c \u0431\u044d\u043a\u0435\u043d\u0434\u043e\u043c clang \u0432 PTX \u044f\u0434\u0440\u0430clang --target=nvptx64-nvidia-cuda -march=sm_75 -S sq.ll -o sq_gpu.ptx# 4. \u0417\u0430\u043f\u0443\u0441\u0442\u0438\u043b \u043d\u0430 \u0432\u0438\u0434\u0435\u043e\u043a\u0430\u0440\u0442\u0435 \u0447\u0435\u0440\u0435\u0437 CUDA Driver API.\/gosq# Output: OK: 1024 \u043a\u0432\u0430\u0434\u0440\u0430\u0442\u0430 \u043f\u043e\u0441\u0447\u0438\u0442\u0430\u043d\u044b \u043d\u0430 GeForce MX450<\/code><div class=\"code-explainer\"><a href=\"https:\/\/sourcecraft.dev\/\" class=\"tm-button code-explainer__link\" style=\"visibility: hidden;\"><span class=\"sc-logo-wrapper icon-only\"><\/span><\/a><\/div><\/pre>\n<p>\u0427\u0442\u043e\u00a0\u0432\u043d\u0443\u0442\u0440\u0438 <code>sq_gpu.ptx<\/code>\u00a0\u2014 \u043d\u0430\u0441\u0442\u043e\u044f\u0449\u0435\u0435 \u044f\u0434\u0440\u043e, \u0430\u00a0\u043d\u0435\u00a0\u0437\u0430\u0433\u043b\u0443\u0448\u043a\u0430:<\/p>\n<pre><code>.visible .entry _Z2sqPKfPfj(    mul.lo.s32  %r7, %r6, %r5;    ld.global.b32   %r12, [%rd9];    mul.rn.f32  %r14, %r13, %r13;    st.global.b32   [%rd14], %r15;<\/code><div class=\"code-explainer\"><a href=\"https:\/\/sourcecraft.dev\/\" class=\"tm-button code-explainer__link\" style=\"visibility: hidden;\"><span class=\"sc-logo-wrapper icon-only\"><\/span><\/a><\/div><\/pre>\n<h3>3. \u0413\u0440\u0430\u0431\u043b\u0438: \u043a\u0430\u043a\u00a0\u044f \u043b\u043e\u0432\u0438\u043b \u0442\u0438\u0445\u0443\u044e \u043f\u043e\u0440\u0447\u0443 \u0434\u0430\u043d\u043d\u044b\u0445<\/h3>\n<pre><code>\u0440\u0443\u0447\u043d\u043e\u0439 PTX \u2192 \u0432\u0441\u0451 \u0437\u0435\u043b\u0451\u043d\u043e\u0435 \u2192 \u043d\u0430\u0441\u0442\u043e\u044f\u0449\u0438\u0439 nvcc \u2192 \u043f\u0430\u0440\u0441\u0435\u0440 \u043b\u043e\u043c\u0430\u0435\u0442\u0441\u044f    \u2192 \u0447\u0438\u043d\u044e \u0433\u0440\u0430\u043c\u043c\u0430\u0442\u0438\u043a\u0443 \u2192 \u0440\u0435\u0430\u043b\u044c\u043d\u0430\u044f GPU \u2192 \u0432\u044b\u043b\u0430\u0437\u044f\u0442 address-space \u0431\u0430\u0433\u0438    \u2192 \u0438\u0441\u043f\u0440\u0430\u0432\u043b\u044f\u044e pipeline<\/code><div class=\"code-explainer\"><a href=\"https:\/\/sourcecraft.dev\/\" class=\"tm-button code-explainer__link\" style=\"visibility: hidden;\"><span class=\"sc-logo-wrapper icon-only\"><\/span><\/a><\/div><\/pre>\n<p>\u041d\u0430\u043f\u0438\u0441\u0430\u0442\u044c \u043f\u0430\u0440\u0441\u0435\u0440 \u0434\u043b\u044f <code>add<\/code>\/<code>mul<\/code>\u00a0\u2014 \u043f\u0440\u043e\u0441\u0442\u043e. \u0421\u043b\u043e\u0436\u043d\u043e\u0441\u0442\u0438 \u043d\u0430\u0447\u0430\u043b\u0438\u0441\u044c \u043d\u0430\u00a0\u0432\u044b\u0445\u043b\u043e\u043f\u0435 \u0441\u0432\u0435\u0436\u0435\u0433\u043e <code>nvcc<\/code>.<\/p>\n<p><strong>\u041a\u0430\u043c\u0435\u043d\u044c 1: \u043e\u0434\u043d\u043e\u0441\u0442\u0440\u043e\u0447\u043d\u044b\u0435 \u0441\u0438\u0433\u043d\u0430\u0442\u0443\u0440\u044b.<\/strong> \u0421\u043e\u0432\u0440\u0435\u043c\u0435\u043d\u043d\u044b\u0439 <code>nvcc<\/code> \u043f\u0438\u0448\u0435\u0442 <code>.visible .entry _Z...(\u2026)<\/code> \u0432\u00a0\u043e\u0434\u043d\u0443 \u0441\u0442\u0440\u043e\u043a\u0443, \u0430\u00a0\u043c\u043e\u0439 \u043f\u0430\u0440\u0441\u0435\u0440 \u0436\u0434\u0430\u043b \u0441\u043a\u043e\u0431\u043a\u0443 \u043d\u0430\u00a0\u043e\u0442\u0434\u0435\u043b\u044c\u043d\u043e\u0439 \u0441\u0442\u0440\u043e\u043a\u0435\u00a0\u2014 \u0438 \u0442\u0435\u043b\u043e \u044f\u0434\u0440\u0430 \u043e\u0441\u0442\u0430\u0432\u0430\u043b\u043e\u0441\u044c \u043f\u0443\u0441\u0442\u044b\u043c. \u041f\u043b\u044e\u0441 37\u00a0\u043c\u0435\u043b\u043a\u0438\u0445 \u043f\u0440\u043e\u0431\u0435\u043b\u043e\u0432 \u0432\u00a0\u0433\u0440\u0430\u043c\u043c\u0430\u0442\u0438\u043a\u0435: <code>and<\/code>\/<code>or<\/code>\/<code>xor<\/code>, \u0441\u0434\u0432\u0438\u0433\u0438, <code>selp<\/code>, <code>div<\/code>\/<code>rem<\/code>, <code>mad.lo.s32<\/code>, <code>bfi<\/code>\/<code>bfe<\/code>, <code>f16x2<\/code>, \u0432\u0430\u0440\u043f\u2011<code>shfl<\/code>. \u0411\u0430\u0437\u043e\u0432\u0430\u044f \u0430\u0440\u0438\u0444\u043c\u0435\u0442\u0438\u043a\u0430, \u043a\u043e\u0442\u043e\u0440\u0443\u044e \u043a\u043e\u043c\u043f\u0438\u043b\u044f\u0442\u043e\u0440 \u0433\u0435\u043d\u0435\u0440\u0438\u0440\u0443\u0435\u0442 \u0432\u00a0\u043a\u0430\u0436\u0434\u043e\u043c \u044f\u0434\u0440\u0435.<\/p>\n<p><strong>\u041a\u0430\u043c\u0435\u043d\u044c 2: \u0438\u043d\u0432\u0435\u0440\u0441\u0438\u044f \u043f\u0440\u0435\u0434\u0438\u043a\u0430\u0442\u043e\u0432.<\/strong> \u0417\u0430\u043f\u0438\u0441\u044c <code>@!%p1 bra LABEL<\/code> \u043e\u0437\u043d\u0430\u0447\u0430\u0435\u0442 \u00ab\u043f\u0435\u0440\u0435\u0439\u0442\u0438, \u0435\u0441\u043b\u0438 \u043f\u0440\u0435\u0434\u0438\u043a\u0430\u0442 \u043b\u043e\u0436\u0435\u043d\u00bb\u00a0\u2014 \u0444\u043b\u0430\u0433 <code>!<\/code> \u043e\u0431\u044f\u0437\u0430\u043d \u044f\u0432\u043d\u043e \u0438\u043d\u0432\u0435\u0440\u0442\u0438\u0440\u043e\u0432\u0430\u0442\u044c\u0441\u044f \u0432\u00a0IR. \u0423\u00a0\u043c\u0435\u043d\u044f \u044d\u0442\u043e \u0437\u0430\u043a\u0440\u044b\u0442\u043e \u043d\u0430\u00a0\u0443\u0440\u043e\u0432\u043d\u0435 \u0447\u0442\u0435\u043d\u0438\u044f \u043f\u0440\u0435\u0434\u0438\u043a\u0430\u0442\u043d\u044b\u0445 \u043e\u043f\u0435\u0440\u0430\u043d\u0434\u043e\u0432: \u0432\u0435\u0434\u0443\u0449\u0438\u0439 <code>!<\/code> \u043f\u0440\u0435\u0432\u0440\u0430\u0449\u0430\u0435\u0442\u0441\u044f \u0432 <code>xor i1 \u2026, true<\/code>, \u0438 \u0442\u043e\u00a0\u0436\u0435 \u043f\u0440\u0430\u0432\u0438\u043b\u043e \u0434\u0435\u0439\u0441\u0442\u0432\u0443\u0435\u0442 \u0432\u00a0\u0438\u043d\u0442\u0435\u0440\u043f\u0440\u0435\u0442\u0430\u0442\u043e\u0440\u0435. \u0422\u0438\u0445\u043e\u0439 \u043f\u043e\u0440\u0447\u0438 \u0437\u0434\u0435\u0441\u044c \u043d\u0435\u0442 \u0438\u043c\u0435\u043d\u043d\u043e \u043f\u043e\u0442\u043e\u043c\u0443, \u0447\u0442\u043e\u00a0\u0438\u043d\u0432\u0435\u0440\u0441\u0438\u044f\u00a0\u2014 \u0447\u0430\u0441\u0442\u044c \u043a\u043e\u043d\u0442\u0440\u0430\u043a\u0442\u0430 \u0447\u0442\u0435\u043d\u0438\u044f \u043e\u043f\u0435\u0440\u0430\u043d\u0434\u0430, \u0430\u00a0\u043d\u0435 \u00ab\u043d\u0430\u0434\u0435\u0436\u0434\u0430, \u0447\u0442\u043e\u00a0\u0431\u044d\u043a\u0435\u043d\u0434 \u0443\u0433\u0430\u0434\u0430\u0435\u0442\u00bb.<\/p>\n<p><strong>\u041a\u0430\u043c\u0435\u043d\u044c 3: <\/strong><code><strong>setp<\/strong><\/code><strong> \u0438 <\/strong><code><strong>div<\/strong><\/code><strong>.<\/strong> <code>setp.lt.f32<\/code> \u0443\u00a0\u043c\u0435\u043d\u044f \u0441\u0440\u0430\u0432\u043d\u0438\u0432\u0430\u043b \u0431\u0438\u0442\u044b \u0446\u0435\u043b\u043e\u0447\u0438\u0441\u043b\u0435\u043d\u043d\u044b\u043c <code>icmp<\/code> \u0432\u043c\u0435\u0441\u0442\u043e <code>fcmp<\/code>, \u0430 <code>div.s32<\/code> \u044d\u043c\u0438\u0442\u0438\u043b <code>mul<\/code>. \u041e\u0431\u0430 \u043d\u0435\u00a0\u043f\u0430\u0434\u0430\u043b\u0438\u00a0\u2014 \u043c\u043e\u043b\u0447\u0430 \u0441\u0447\u0438\u0442\u0430\u043b\u0438 \u043d\u0435\u043f\u0440\u0430\u0432\u0438\u043b\u044c\u043d\u043e. \u041d\u0430\u0448\u043b\u0438\u0441\u044c \u0442\u043e\u043b\u044c\u043a\u043e \u0441\u0432\u0435\u0440\u043a\u043e\u0439 \u0441\u00a0\u043d\u0430\u0441\u0442\u043e\u044f\u0449\u0438\u043c \u0432\u0445\u043e\u0434\u043e\u043c \u043e\u0442 <code>nvcc<\/code>.<\/p>\n<blockquote>\n<p>\u041c\u043e\u0440\u0430\u043b\u044c, \u043a\u043e\u0442\u043e\u0440\u0443\u044e \u044f \u0432\u044b\u043d\u0435\u0441: \u0442\u0435\u0441\u0442 \u043d\u0430\u00a0\u0441\u043e\u0431\u0441\u0442\u0432\u0435\u043d\u043d\u044b\u0445 \u0434\u0430\u043d\u043d\u044b\u0445 \u0434\u043e\u043a\u0430\u0437\u044b\u0432\u0430\u0435\u0442, \u0447\u0442\u043e\u00a0\u043a\u043e\u0434 \u0440\u0430\u0431\u043e\u0442\u0430\u0435\u0442 \u043d\u0430\u00a0\u0442\u0432\u043e\u0438\u0445 \u0434\u0430\u043d\u043d\u044b\u0445. \u041d\u0443\u0436\u0435\u043d \u0447\u0443\u0436\u043e\u0439 \u0432\u0445\u043e\u0434\u00a0\u2014 \u043d\u0430\u0441\u0442\u043e\u044f\u0449\u0438\u0439 PTX \u0438\u0437\u00a0CI (<code>compiler_nvcc_ptx<\/code>).<\/p>\n<\/blockquote>\n<p>\u0410\u00a0\u043f\u043e\u0442\u043e\u043c \u0432\u0438\u0434\u0435\u043e\u043a\u0430\u0440\u0442\u0430 \u043d\u0430\u0448\u043b\u0430 \u0435\u0449\u0451 \u0434\u0432\u0430 \u0431\u0430\u0433\u0430, \u043d\u0435\u0432\u0438\u0434\u0438\u043c\u044b\u0445 \u043d\u0430\u00a0\u0445\u043e\u0441\u0442\u0435: \u0432\u044b\u0432\u043e\u0434 \u043c\u0430\u0442\u0440\u0438\u0446\u044b \u043f\u0438\u0441\u0430\u043b\u0441\u044f \u0432\u00a0shared \u0432\u043c\u0435\u0441\u0442\u043e global (\u043c\u0435\u0442\u043a\u0430 \u0440\u0435\u0433\u0438\u0441\u0442\u0440\u0430 \u043f\u0435\u0440\u0435\u0436\u0438\u0432\u0430\u043b\u0430 \u043f\u0435\u0440\u0435\u0437\u0430\u043f\u0438\u0441\u044c\u00a0\u2014 \u0442\u0435\u043f\u0435\u0440\u044c \u043f\u0440\u043e\u0441\u0442\u0440\u0430\u043d\u0441\u0442\u0432\u043e \u0431\u0435\u0440\u0451\u0442\u0441\u044f \u0438\u0437\u00a0\u0441\u0443\u0444\u0444\u0438\u043a\u0441\u0430 \u0438\u043d\u0441\u0442\u0440\u0443\u043a\u0446\u0438\u0438), \u0430\u00a0\u0430\u0434\u0440\u0435\u0441 device\u2011\u0433\u043b\u043e\u0431\u0430\u043b\u0430 \u0447\u0438\u0442\u0430\u043b\u0441\u044f \u043a\u0430\u043a\u00a0\u0437\u0430\u0433\u0440\u0443\u0437\u043a\u0430 \u0438\u0437\u00a0\u043d\u0435\u0433\u043e. \u041d\u0430\u00a0\u043f\u043b\u043e\u0441\u043a\u043e\u0439 \u0445\u043e\u0441\u0442\u2011\u043f\u0430\u043c\u044f\u0442\u0438 \u043e\u0431\u0430 \u043c\u043e\u043b\u0447\u0430\u043b\u0438.<\/p>\n<h3>4. \u041c\u043e\u0438 \u0440\u0435\u0437\u0443\u043b\u044c\u0442\u0430\u0442\u044b \u0438 \u0442\u0435\u0441\u0442\u044b<\/h3>\n<p><code>ctest<\/code>\u00a0\u2014 21\/21, \u043d\u043e\u043b\u044c \u0432\u043e\u0440\u043d\u0438\u043d\u0433\u043e\u0432. \u0418\u0437\u00a0\u043d\u0438\u0445:<\/p>\n<ul>\n<li>\n<p>7\u00a0\u043a\u0430\u043d\u043e\u043d\u0438\u0447\u0435\u0441\u043a\u0438\u0445 \u044f\u0434\u0435\u0440 (<code>vadd<\/code>, <code>block_reduce<\/code>, <code>atomic_reduce<\/code>, <code>fp16_reduce<\/code>, <code>gemm_tile<\/code>, <code>shfl_reduce<\/code>, <code>atom_cas<\/code>) \u0438\u0441\u043f\u043e\u043b\u043d\u0435\u043d\u044b \u043d\u0430\u00a0GeForce MX450\u00a0\u0438 \u0441\u043e\u0448\u043b\u0438\u0441\u044c <strong>\u043f\u043e\u0431\u0438\u0442\u043e\u0432\u043e<\/strong> \u0441\u00a0\u0445\u043e\u0441\u0442\u2011\u044d\u0442\u0430\u043b\u043e\u043d\u043e\u043c;<\/p>\n<\/li>\n<li>\n<p>\u043d\u0430\u0441\u0442\u043e\u044f\u0449\u0435\u0435 \u044f\u0434\u0440\u043e \u043e\u0442 <code>nvcc<\/code> (\u043c\u0430\u043d\u0433\u043b\u0438\u0440\u043e\u0432\u0430\u043d\u043d\u043e\u0435 \u0438\u043c\u044f, device\u2011\u0433\u043b\u043e\u0431\u0430\u043b, <code><strong>syncthreads<\/strong><\/code><strong>)\u00a0\u2014 \u043d\u0430\u0441\u043a\u0432\u043e\u0437\u044c <\/strong><code><strong>nvcc \u2192 ptx2ir \u2192 GPU<\/strong><\/code><strong>. \u0427\u0435\u0441\u0442\u043d\u0430\u044f \u043e\u0433\u043e\u0432\u043e\u0440\u043a\u0430: \u0432\u00a0\u0435\u0433\u043e \u0438\u0441\u0445\u043e\u0434\u043d\u0438\u043a\u0435 \u0433\u043e\u043d\u043a\u0430 (\u0440\u0435\u0434\u0443\u043a\u0446\u0438\u044f \u0431\u0435\u0437 <\/strong><code>syncthreads<\/code>), \u0432\u0435\u043d\u0434\u043e\u0440\u043d\u0430\u044f \u0441\u0431\u043e\u0440\u043a\u0430 \u0433\u043e\u043d\u0438\u0442\u0441\u044f \u0442\u0430\u043a\u00a0\u0436\u0435\u00a0\u2014 \u044d\u0442\u0430\u043b\u043e\u043d \u043f\u0440\u0438\u043d\u0438\u043c\u0430\u0435\u0442 \u0434\u043e\u043a\u0430\u0437\u0430\u043d\u043d\u044b\u0439 \u0438\u043d\u0442\u0435\u0440\u0432\u0430\u043b, \u0430\u00a0\u043d\u0435\u00a0\u043e\u0434\u043d\u043e \u0447\u0438\u0441\u043b\u043e;<\/p>\n<\/li>\n<li>\n<p><code>shfl.sync.bfly<\/code> \u043f\u0440\u043e\u0432\u0435\u0440\u0435\u043d \u043e\u0442\u0434\u0435\u043b\u044c\u043d\u044b\u043c NVPTX\u2011\u0442\u0435\u0441\u0442\u043e\u043c (<code>shfl_reduce<\/code>) \u043d\u0430\u00a0\u0440\u0435\u0430\u043b\u044c\u043d\u043e\u0439 GPU\u00a0\u2014 \u0447\u0435\u0440\u0435\u0437 \u043d\u0430\u0441\u0442\u043e\u044f\u0449\u0438\u0439 NVVM\u2011\u0438\u043d\u0442\u0440\u0438\u043d\u0441\u0438\u043a; \u043e\u0441\u0442\u0430\u043b\u044c\u043d\u044b\u0435 warp\u2011\u043a\u043e\u043b\u043b\u0435\u043a\u0442\u0438\u0432\u044b \u043f\u043e\u043a\u0440\u044b\u0442\u044b\u043c\u0438 \u0442\u0435\u0441\u0442\u0430\u043c\u0438 \u043d\u0435\u00a0\u0441\u0447\u0438\u0442\u0430\u044e: <code>vote<\/code> \/ <code>activemask<\/code> \/ \u043f\u0440\u0435\u0434\u0438\u043a\u0430\u0442\u043d\u044b\u0439 <code>shfl<\/code>\u00a0\u2014 \u044f\u0432\u043d\u044b\u0439 \u043e\u0442\u043a\u0430\u0437, \u0430\u00a0\u043d\u0435\u00a0\u0442\u0438\u0445\u0438\u0439 \u0444\u043e\u043b\u0431\u044d\u043a;<\/p>\n<\/li>\n<li>\n<p>\u0431\u0435\u0437\u00a0\u043a\u0430\u0440\u0442\u044b sm_75+ GPU\u2011\u0442\u0435\u0441\u0442\u044b \u0433\u043e\u0432\u043e\u0440\u044f\u0442 \u0447\u0435\u0441\u0442\u043d\u044b\u0439 SKIP, \u0430\u00a0\u043d\u0435 \u00ab\u0443\u0441\u043f\u0435\u0445\u00bb.<\/p>\n<\/li>\n<\/ul>\n<p>\u041f\u0440\u043e\u0438\u0437\u0432\u043e\u0434\u0438\u0442\u0435\u043b\u044c\u043d\u043e\u0441\u0442\u0438 \u044f \u043d\u0435\u00a0\u043c\u0435\u0440\u044f\u043b\u00a0\u2014 \u0442\u043e\u043b\u044c\u043a\u043e \u0431\u0438\u0442\u2011\u0438\u0434\u0435\u043d\u0442\u0438\u0447\u043d\u043e\u0441\u0442\u044c. \u041a\u0442\u043e \u043e\u0431\u0435\u0449\u0430\u0435\u0442 \u0441\u043a\u043e\u0440\u043e\u0441\u0442\u044c \u0431\u0435\u0437\u00a0\u0437\u0430\u043c\u0435\u0440\u043e\u0432, \u0442\u043e\u0442 \u0432\u0440\u0451\u0442; \u044f \u043d\u0435\u00a0\u043e\u0431\u0435\u0449\u0430\u044e.<\/p>\n<h3>5. \u0427\u0442\u043e\u00a0\u0434\u0430\u043b\u044c\u0448\u0435?<\/h3>\n<ul>\n<li>\n<p>SIMT \u2192 RVV \u0432\u0435\u043a\u0442\u043e\u0440\u0438\u0437\u0430\u0446\u0438\u044f: \u0441\u0432\u043e\u0440\u0430\u0447\u0438\u0432\u0430\u0442\u044c 32\u00a0\u043f\u043e\u0442\u043e\u043a\u0430 \u0432\u00a0\u0432\u0435\u043a\u0442\u043e\u0440\u043d\u044b\u0435\u00a0\u043b\u0438\u043d\u0438\u0438 (MLIR\u2011\u043a\u043e\u043d\u0442\u0443\u0440 \u0443\u0436\u0435 \u0430\u043d\u0430\u043b\u0438\u0437\u0438\u0440\u0443\u0435\u0442 \u043f\u0430\u0442\u0442\u0435\u0440\u043d\u044b, \u043f\u043e\u043b\u043d\u043e\u0446\u0435\u043d\u043d\u043e\u0439 \u0432\u0435\u043a\u0442\u043e\u0440\u0438\u0437\u0430\u0446\u0438\u0438 \u043f\u043e\u043a\u0430 \u043d\u0435\u0442);<\/p>\n<\/li>\n<li>\n<p><code>mma<\/code> \u0438 \u0442\u0435\u043d\u0437\u043e\u0440\u043d\u044b\u0435 \u0438\u043d\u0441\u0442\u0440\u0443\u043a\u0446\u0438\u0438\u00a0\u2014 \u043f\u043e\u043a\u0430 \u043e\u0441\u043e\u0437\u043d\u0430\u043d\u043d\u044b\u0439 \u043e\u0442\u043a\u0430\u0437: \u00ab\u043f\u0440\u0438\u043c\u0435\u0440\u043d\u043e \u043f\u0440\u0430\u0432\u0438\u043b\u044c\u043d\u043e\u00bb \u043e\u043f\u0430\u0441\u043d\u0435\u0435, \u0447\u0435\u043c \u043d\u0438\u043a\u0430\u043a;<\/p>\n<\/li>\n<li>\n<p>\u0431\u043e\u043b\u044c\u0448\u0435 \u043d\u0430\u0441\u0442\u043e\u044f\u0449\u0438\u0445 CUDA\u2011\u0431\u0438\u043d\u0430\u0440\u043d\u0438\u043a\u043e\u0432 \u043d\u0430\u00a0\u0432\u0445\u043e\u0434\u0435 \u0432\u043c\u0435\u0441\u0442\u043e \u0441\u0438\u043d\u0442\u0435\u0442\u0438\u043a\u0438;<\/p>\n<\/li>\n<li>\n<p>\u043f\u0443\u0431\u043b\u0438\u0447\u043d\u044b\u0435 \u0431\u0435\u043d\u0447\u043c\u0430\u0440\u043a\u0438\u00a0\u2014 \u043d\u043e\u00a0\u0441\u043d\u0430\u0447\u0430\u043b\u0430 \u043f\u043e\u0432\u0442\u043e\u0440\u044f\u0435\u043c\u044b\u0439 \u0441\u0442\u0435\u043d\u0434.<\/p>\n<\/li>\n<\/ul>\n<h3>\u0413\u0434\u0435 \u043f\u043e\u0449\u0443\u043f\u0430\u0442\u044c<\/h3>\n<p>\u0411\u0435\u0437\u00a0\u0441\u0431\u043e\u0440\u043a\u0438, \u0434\u0432\u0430 \u0431\u0438\u043d\u0430\u0440\u043d\u0438\u043a\u0430 (\u043a\u0440\u043e\u043c\u0435 \u0441\u0438\u0441\u0442\u0435\u043c\u043d\u043e\u0433\u043e libc \u0437\u0430\u0432\u0438\u0441\u0438\u043c\u043e\u0441\u0442\u0435\u0439 \u043d\u0435\u0442):<\/p>\n<pre><code>curl -fsSL https:\/\/github.com\/MrModelOS\/ROUGE-V\/releases\/download\/v0.1\/rouge-v-0.1-linux-x86\\\\_64.tar.gz | tar xz -C ~\/.localexport PATH=\"$HOME\/.local\/bin:$PATH\"ptx2ir kernel.ptx kernel.ll<\/code><div class=\"code-explainer\"><a href=\"https:\/\/sourcecraft.dev\/\" class=\"tm-button code-explainer__link\" style=\"visibility: hidden;\"><span class=\"sc-logo-wrapper icon-only\"><\/span><\/a><\/div><\/pre>\n<p>\u0418\u0437\u00a0\u0438\u0441\u0445\u043e\u0434\u043d\u0438\u043a\u043e\u0432:<\/p>\n<pre><code>git clone https:\/\/github.com\/MrModelOS\/ROUGE-V.git &amp;&amp; cd ROUGE-Vcmake -B build -G Ninja -DCMAKE_BUILD_TYPE=Release &amp;&amp; cmake --build buildctest --test-dir build --output-on-failure<\/code><div class=\"code-explainer\"><a href=\"https:\/\/sourcecraft.dev\/\" class=\"tm-button code-explainer__link\" style=\"visibility: hidden;\"><span class=\"sc-logo-wrapper icon-only\"><\/span><\/a><\/div><\/pre>\n<p>\u0410\u0440\u0445\u0438\u0442\u0435\u043a\u0442\u0443\u0440\u0430\u00a0\u2014 <a href=\"https:\/\/github.com\/MrModelOS\/ROUGE-V\/blob\/v0.1\/docs\/06-compiler-architecture.md\" rel=\"noopener nofollow\"><code>docs\/06-compiler-architecture.md<\/code><\/a>.\u00a0\u041b\u0438\u0446\u0435\u043d\u0437\u0438\u044f\u00a0\u2014 Apache 2.0\u00a0WITH LLVM\u2011exception.<\/p>\n<h3>\u041f\u0440\u043e\u00a0\u044e\u0440\u0438\u0434\u0438\u0447\u0435\u0441\u043a\u0443\u044e \u0441\u0442\u043e\u0440\u043e\u043d\u0443\u00a0\u2014 \u043f\u0440\u044f\u043c\u043e<\/h3>\n<ul>\n<li>\n<p>\u042f <strong>\u043d\u0435\u00a0\u043a\u043e\u043f\u0438\u0440\u0443\u044e \u0438 \u043d\u0435\u00a0\u0440\u0435\u0432\u0435\u0440\u0441\u2011\u0438\u043d\u0436\u0438\u043d\u0438\u0440\u044e<\/strong> \u043f\u0440\u043e\u043f\u0440\u0438\u0435\u0442\u0430\u0440\u043d\u044b\u0435 \u0431\u0438\u043d\u0430\u0440\u043d\u0438\u043a\u0438 \u0438 \u0437\u0430\u0433\u043e\u043b\u043e\u0432\u043a\u0438 NVIDIA.<\/p>\n<\/li>\n<li>\n<p>PTX, \u043a\u043e\u0442\u043e\u0440\u044b\u0439 \u044f \u043f\u0435\u0440\u0435\u0432\u043e\u0436\u0443, <strong>\u0432\u044b\u043f\u0443\u0441\u043a\u0430\u0435\u0442 <\/strong><code><strong>nvcc<\/strong><\/code><strong> \u043f\u043e\u043b\u044c\u0437\u043e\u0432\u0430\u0442\u0435\u043b\u044f<\/strong>\u00a0\u2014 \u0438\u0437\u00a0\u0435\u0433\u043e \u0441\u043e\u0431\u0441\u0442\u0432\u0435\u043d\u043d\u043e\u0439\u00a0\u043b\u0438\u0446\u0435\u043d\u0437\u0438\u0438, \u043f\u043e\u00a0\u043e\u0442\u043a\u0440\u044b\u0442\u043e\u0439 \u0441\u043f\u0435\u0446\u0438\u0444\u0438\u043a\u0430\u0446\u0438\u0438 PTX ISA.<\/p>\n<\/li>\n<li>\n<p>ROUGE\u2011V \u043d\u0435\u00a0\u0441\u0432\u044f\u0437\u0430\u043d \u0441\u00a0NVIDIA \u0438 \u043d\u0435\u00a0\u043e\u0434\u043e\u0431\u0440\u0435\u043d \u0438\u043c\u0438. \u0422\u043e\u0440\u0433\u043e\u0432\u044b\u0435 \u043c\u0430\u0440\u043a\u0438\u00a0\u2014 \u0442\u043e\u043b\u044c\u043a\u043e \u0447\u0442\u043e\u0431\u044b \u043e\u043f\u0438\u0441\u0430\u0442\u044c \u0441\u043e\u0432\u043c\u0435\u0441\u0442\u0438\u043c\u043e\u0441\u0442\u044c.<\/p>\n<\/li>\n<\/ul>\n<hr\/>\n<h3>P.\u202fS.<\/h3>\n<p>\u0415\u0441\u043b\u0438 \u0432\u043e\u0437\u044c\u043c\u0451\u0442\u0435 PTX \u043e\u0442\u00a0\u0440\u0435\u0430\u043b\u044c\u043d\u043e\u0433\u043e <code>nvcc<\/code> \u0438 \u043c\u043e\u0439 \u0442\u0440\u0430\u043d\u0441\u043b\u044f\u0442\u043e\u0440 \u0441\u043a\u0430\u0436\u0435\u0442 \u043e\u0442\u043a\u0430\u0437\u00a0\u2014 \u0440\u0430\u0441\u0441\u043a\u0430\u0436\u0438\u0442\u0435, \u0447\u0442\u043e\u00a0\u0438\u043c\u0435\u043d\u043d\u043e: \u044d\u0442\u043e \u0441\u0430\u043c\u044b\u0439 \u043f\u043e\u043b\u0435\u0437\u043d\u044b\u0439 \u0431\u0430\u0433\u2011\u0440\u0435\u043f\u043e\u0440\u0442 \u0434\u043b\u044f\u00a0\u043f\u0440\u043e\u0435\u043a\u0442\u0430.<\/p>\n<\/div>\n<p>\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\/1089146\/\">https:\/\/habr.com\/ru\/articles\/1089146\/<\/a><\/p>\n","protected":false},"excerpt":{"rendered":"<p>TL;DR: \u041f\u0440\u043e\u0431\u043b\u0435\u043c\u0430 \u043f\u0435\u0440\u0435\u043d\u043e\u0441\u0438\u043c\u043e\u0441\u0442\u0438 CUDA \u0434\u043b\u044f\u00a0\u043c\u0435\u043d\u044f \u043e\u043a\u0430\u0437\u0430\u043b\u0430\u0441\u044c \u043d\u0435\u00a0\u0441\u0442\u043e\u043b\u044c\u043a\u043e \u0432\u00a0\u0441\u0430\u043c\u0438\u0445 \u0430\u0440\u0438\u0444\u043c\u0435\u0442\u0438\u0447\u0435\u0441\u043a\u0438\u0445 \u043e\u043f\u0435\u0440\u0430\u0446\u0438\u044f\u0445, \u0441\u043a\u043e\u043b\u044c\u043a\u043e \u0432\u00a0\u043e\u0433\u0440\u043e\u043c\u043d\u043e\u043c \u043a\u043e\u043b\u0438\u0447\u0435\u0441\u0442\u0432\u0435 \u0438\u043d\u0444\u0440\u0430\u0441\u0442\u0440\u0443\u043a\u0442\u0443\u0440\u044b \u0432\u043e\u043a\u0440\u0443\u0433 \u043d\u0438\u0445. \u041f\u043e\u044d\u0442\u043e\u043c\u0443 \u0432\u043c\u0435\u0441\u0442\u043e \u043f\u043e\u043f\u044b\u0442\u043a\u0438 \u043f\u043e\u0432\u0442\u043e\u0440\u044f\u0442\u044c CUDA runtime \u044f \u0440\u0435\u0448\u0438\u043b \u043f\u0435\u0440\u0435\u0445\u0432\u0430\u0442\u0438\u0442\u044c \u043f\u0440\u043e\u0433\u0440\u0430\u043c\u043c\u0443 \u0440\u0430\u043d\u044c\u0448\u0435\u00a0\u2014 \u043d\u0430\u00a0\u0443\u0440\u043e\u0432\u043d\u0435 PTX. \u042f \u043d\u0430\u043f\u0438\u0441\u0430\u043b ROUGE\u2011V\u00a0\u2014 \u043e\u0442\u043a\u0440\u044b\u0442\u044b\u0439 AOT\u2011\u043a\u043e\u043c\u043f\u0438\u043b\u044f\u0442\u043e\u0440, \u0440\u0430\u0437\u0440\u0430\u0431\u043e\u0442\u0430\u043d\u043d\u044b\u0439 \u043f\u043e\u00a0\u043f\u0440\u0438\u043d\u0446\u0438\u043f\u0443 clean\u2011room: \u0431\u0435\u0440\u0451\u0442 PTX, \u0441\u0433\u0435\u043d\u0435\u0440\u0438\u0440\u043e\u0432\u0430\u043d\u043d\u044b\u0439 \u043e\u0444\u0438\u0446\u0438\u0430\u043b\u044c\u043d\u044b\u043c nvcc 13.4, \u0442\u0440\u0430\u043d\u0441\u043b\u0438\u0440\u0443\u0435\u0442 \u0435\u0433\u043e \u0432\u00a0LLVM IR, \u0434\u0430\u043b\u044c\u0448\u0435 \u0441\u043e\u0431\u0438\u0440\u0430\u0435\u0442 \u043e\u0431\u044b\u0447\u043d\u044b\u0439 clang. \u0418\u0441\u043f\u043e\u043b\u043d\u044f\u0435\u0442\u0441\u044f \u043d\u0430\u00a0x86-64\u00a0\u0438 \u043d\u0430\u00a0\u0436\u0438\u0432\u043e\u0439 GeForce MX450\u00a0(\u043f\u043e\u0431\u0438\u0442\u043e\u0432\u043e \u0441\u00a0\u0445\u043e\u0441\u0442\u2011\u044d\u0442\u0430\u043b\u043e\u043d\u043e\u043c); AMD \u0438 RISC\u2011V\u00a0\u2014 \u043f\u043e\u043a\u0430 \u0442\u043e\u043b\u044c\u043a\u043e \u0441\u0431\u043e\u0440\u043a\u0430, \u0431\u0435\u0437\u00a0\u0432\u0440\u0430\u043d\u044c\u044f \u043f\u0440\u043e \u00ab\u043b\u044e\u0431\u043e\u0435\u00a0\u0436\u0435\u043b\u0435\u0437\u043e\u00bb.\u041c\u043e\u0439 GitHub: github.com\/MrModelOS\/ROUGE\u2011V (\u0430\u0432\u0442\u043e\u0440: MrModel).\u041f\u043e\u043a\u0440\u044b\u0442\u0438\u0435: ctest\u00a0\u2014 21\/21; 8\u00a0\u044f\u0434\u0435\u0440 (7\u00a0\u043a\u0430\u043d\u043e\u043d\u0438\u0447\u0435\u0441\u043a\u0438\u0445 + 1\u00a0\u043d\u0430\u0441\u0442\u043e\u044f\u0449\u0435\u0435 \u043e\u0442 nvcc) \u0438\u0441\u043f\u043e\u043b\u043d\u0435\u043d\u044b \u043d\u0430\u00a0GeForce MX450.1. \u0417\u0430\u0447\u0435\u043c \u044f \u0441\u0434\u0435\u043b\u0430\u043b \u0441\u0432\u043e\u0439 \u0442\u0440\u0430\u043d\u0441\u043b\u044f\u0442\u043e\u0440, \u043a\u043e\u0433\u0434\u0430 \u0435\u0441\u0442\u044c \u0434\u0440\u0443\u0433\u0438\u0435 \u043f\u0440\u043e\u0435\u043a\u0442\u044b?\u041e\u0434\u0438\u043d \u0438\u0437\u00a0\u043e\u0447\u0435\u0432\u0438\u0434\u043d\u044b\u0445 \u0441\u043f\u043e\u0441\u043e\u0431\u043e\u0432 \u0431\u043e\u0440\u043e\u0442\u044c\u0441\u044f \u0441\u00a0vendor lock\u2011in NVIDIA\u00a0\u2014 \u0440\u0435\u0438\u043c\u043f\u043b\u0435\u043c\u0435\u043d\u0442\u0438\u0440\u043e\u0432\u0430\u0442\u044c runtime \u0438\u043b\u0438\u00a0\u043f\u0435\u0440\u0435\u0445\u0432\u0430\u0442\u044b\u0432\u0430\u0442\u044c API \u043d\u0430\u00a0\u043b\u0435\u0442\u0443. \u042d\u0442\u043e \u0432\u0435\u0447\u043d\u0430\u044f \u0433\u043e\u043d\u043a\u0430 \u0437\u0430\u00a0\u0447\u0443\u0436\u0438\u043c \u0438\u043d\u0442\u0435\u0440\u0444\u0435\u0439\u0441\u043e\u043c: \u0432\u0435\u043d\u0434\u043e\u0440 \u043f\u043e\u043c\u0435\u043d\u044f\u043b \u043f\u043e\u0432\u0435\u0434\u0435\u043d\u0438\u0435\u00a0\u2014 \u0442\u044b \u0434\u043e\u0433\u043e\u043d\u044f\u0435\u0448\u044c.\u042f \u043f\u043e\u0448\u0451\u043b \u043f\u043e\u00a0\u043f\u0443\u0442\u0438 Clean\u2011Room \u0440\u0430\u0437\u0440\u0430\u0431\u043e\u0442\u043a\u0438.\u041a\u043e\u043c\u043f\u0438\u043b\u044f\u0442\u043e\u0440 nvcc \u0433\u0435\u043d\u0435\u0440\u0438\u0440\u0443\u0435\u0442 \u043d\u0435\u00a0\u043c\u0430\u0448\u0438\u043d\u043d\u044b\u0439 \u043a\u043e\u0434, \u0430\u00a0PTX\u00a0\u2014 \u0442\u0435\u043a\u0441\u0442\u043e\u0432\u044b\u0439 \u043f\u0440\u043e\u043c\u0435\u0436\u0443\u0442\u043e\u0447\u043d\u044b\u0439 \u0430\u0441\u0441\u0435\u043c\u0431\u043b\u0435\u0440 \u0441\u00a0\u043f\u0443\u0431\u043b\u0438\u0447\u043d\u043e \u043e\u043f\u0443\u0431\u043b\u0438\u043a\u043e\u0432\u0430\u043d\u043d\u043e\u0439 \u0441\u043f\u0435\u0446\u0438\u0444\u0438\u043a\u0430\u0446\u0438\u0435\u0439 (\u0432\u044b\u0445\u043b\u043e\u043f \u043d\u0430\u0448\u0435\u0433\u043e nvcc 13.4\u00a0\u2014 PTX ISA 9.4). \u042f \u0432\u044b\u0431\u0440\u0430\u043b \u0435\u0433\u043e \u043a\u0430\u043a\u00a0\u0442\u043e\u0447\u043a\u0443 \u043f\u0435\u0440\u0435\u0445\u0432\u0430\u0442\u0430:\u0411\u0435\u0440\u0443 \u0441\u0442\u0430\u043d\u0434\u0430\u0440\u0442\u043d\u044b\u0439 .cu\u2011\u043a\u043e\u0434 \u0438 \u043a\u043e\u043c\u043f\u0438\u043b\u0438\u0440\u0443\u044e \u0435\u0433\u043e \u0447\u0435\u0440\u0435\u0437 nvcc \u0432 .ptx.\u041c\u043e\u0439 \u0442\u0440\u0430\u043d\u0441\u043b\u044f\u0442\u043e\u0440 ptx2ir \u043f\u0430\u0440\u0441\u0438\u0442 PTX \u0438 \u0441\u0442\u0440\u043e\u0438\u0442 \u044d\u043a\u0432\u0438\u0432\u0430\u043b\u0435\u043d\u0442\u043d\u043e\u0435 \u043f\u0440\u0435\u0434\u0441\u0442\u0430\u0432\u043b\u0435\u043d\u0438\u0435 \u0432\u00a0LLVM IR.\u0421\u0442\u0430\u043d\u0434\u0430\u0440\u0442\u043d\u044b\u0439 clang \/ LLVM \u0441\u043e\u0431\u0438\u0440\u0430\u0435\u0442 \u043a\u043e\u043d\u0435\u0447\u043d\u044b\u0439 \u0431\u0438\u043d\u0430\u0440\u043d\u0438\u043a \u043f\u043e\u0434\u00a0\u0446\u0435\u043b\u0435\u0432\u0443\u044e \u0430\u0440\u0445\u0438\u0442\u0435\u043a\u0442\u0443\u0440\u0443.\u041d\u0438\u043a\u0430\u043a\u0438\u0445 \u043f\u0440\u043e\u043f\u0440\u0438\u0435\u0442\u0430\u0440\u043d\u044b\u0445 \u0431\u0438\u043d\u0430\u0440\u043d\u0438\u043a\u043e\u0432 \u0438 \u0437\u0430\u0433\u043e\u043b\u043e\u0432\u043a\u043e\u0432 NVIDIA \u044f \u043d\u0435\u00a0\u043a\u043e\u043f\u0438\u0440\u0443\u044e \u0438 \u043d\u0435\u00a0\u0440\u0435\u0432\u0435\u0440\u0441\u0438\u0440\u0443\u044e\u00a0\u2014 \u043f\u0440\u0430\u0432\u0438\u043b\u043e clean\u2011room \u0437\u0430\u043a\u0440\u0435\u043f\u043b\u0435\u043d\u043e \u0432 CONTRIBUTING.md.2. \u0410\u0440\u0445\u0438\u0442\u0435\u043a\u0442\u0443\u0440\u0430 \u0442\u0440\u0430\u043d\u0441\u043b\u044f\u0442\u043e\u0440\u0430 \u0438 4\u00a0\u0448\u0430\u0433\u0430 \u0434\u043e\u00a0\u0437\u0430\u043f\u0443\u0441\u043a\u0430\u041f\u0430\u0439\u043f\u043b\u0430\u0439\u043d \u0443\u00a0\u043c\u0435\u043d\u044f \u0432\u044b\u0433\u043b\u044f\u0434\u0438\u0442 \u0442\u0430\u043a:CUDA \/ .cu    \u2502    \u25bcnvcc &#8212;ptx    \u2502    \u25bcPTX    \u2502    \u25bcptx2ir    \u2502    \u25bcLLVM IR    \u2502    \u25bcLLVM \/ clang    \u2502    \u25bcnative codeNVPTX \u2192 GeForce MX450 \u2192 bit-exact check\u0427\u0435\u0441\u0442\u043d\u0430\u044f \u043b\u0435\u0433\u0435\u043d\u0434\u0430: AMD (amdgcn) \u0438 RISC\u2011V (rv64gcv)\u00a0\u2014 \u043a\u0440\u043e\u0441\u0441\u2011\u043a\u043e\u043c\u043f\u0438\u043b\u044f\u0446\u0438\u044f, \u0438\u0441\u043f\u043e\u043b\u043d\u0435\u043d\u0438\u044f \u043d\u0430\u00a0\u0436\u0435\u043b\u0435\u0437\u0435 \u0443\u00a0\u043c\u0435\u043d\u044f \u043d\u0435\u00a0\u0431\u044b\u043b\u043e. \u0418\u0441\u043f\u043e\u043b\u043d\u044f\u0435\u0442\u0441\u044f\u00a0\u2014 \u0445\u043e\u0441\u0442 \u0438 NVIDIA.\u041d\u0430\u00a0\u043f\u0440\u0430\u043a\u0442\u0438\u043a\u0435 \u0432\u0435\u0441\u044c \u043f\u0440\u043e\u0446\u0435\u0441\u0441 \u0437\u0430\u043d\u0438\u043c\u0430\u0435\u0442 4\u00a0\u043a\u043e\u043c\u0430\u043d\u0434\u044b (\u043f\u0440\u043e\u0432\u0435\u0440\u0435\u043d\u043e \u043d\u0430\u00a0\u0436\u0438\u0432\u043e\u043c \u0441\u0442\u0435\u043d\u0434\u0435):# 1. \u0421\u043a\u043e\u043c\u043f\u0438\u043b\u0438\u0440\u043e\u0432\u0430\u043b .cu \u0432 PTX \u0432\u0435\u043d\u0434\u043e\u0440\u043d\u044b\u043c nvcc 13.4nvcc -arch=sm_75 -ptx sq.cu -o sq.ptx# 2. \u0422\u0440\u0430\u043d\u0441\u043b\u0438\u0440\u043e\u0432\u0430\u043b PTX \u0432 LLVM IR \u0441\u0432\u043e\u0438\u043c ptx2irptx2ir &#8212;target nvptx64-nvidia-cuda sq.ptx sq.ll# 3. \u0421\u043e\u0431\u0440\u0430\u043b \u043d\u0430\u0441\u0442\u043e\u044f\u0449\u0438\u043c \u0431\u044d\u043a\u0435\u043d\u0434\u043e\u043c clang \u0432 PTX \u044f\u0434\u0440\u0430clang &#8212;target=nvptx64-nvidia-cuda -march=sm_75 -S sq.ll -o sq_gpu.ptx# 4. \u0417\u0430\u043f\u0443\u0441\u0442\u0438\u043b \u043d\u0430 \u0432\u0438\u0434\u0435\u043e\u043a\u0430\u0440\u0442\u0435 \u0447\u0435\u0440\u0435\u0437 CUDA Driver API.\/gosq# Output: OK: 1024 \u043a\u0432\u0430\u0434\u0440\u0430\u0442\u0430 \u043f\u043e\u0441\u0447\u0438\u0442\u0430\u043d\u044b \u043d\u0430 GeForce MX450\u0427\u0442\u043e\u00a0\u0432\u043d\u0443\u0442\u0440\u0438 sq_gpu.ptx\u00a0\u2014 \u043d\u0430\u0441\u0442\u043e\u044f\u0449\u0435\u0435 \u044f\u0434\u0440\u043e, \u0430\u00a0\u043d\u0435\u00a0\u0437\u0430\u0433\u043b\u0443\u0448\u043a\u0430:.visible .entry _Z2sqPKfPfj(    mul.lo.s32  %r7, %r6, %r5;    ld.global.b32   %r12, [%rd9];    mul.rn.f32  %r14, %r13, %r13;    st.global.b32   [%rd14], %r15;3. \u0413\u0440\u0430\u0431\u043b\u0438: \u043a\u0430\u043a\u00a0\u044f \u043b\u043e\u0432\u0438\u043b \u0442\u0438\u0445\u0443\u044e \u043f\u043e\u0440\u0447\u0443 \u0434\u0430\u043d\u043d\u044b\u0445\u0440\u0443\u0447\u043d\u043e\u0439 PTX \u2192 \u0432\u0441\u0451 \u0437\u0435\u043b\u0451\u043d\u043e\u0435 \u2192 \u043d\u0430\u0441\u0442\u043e\u044f\u0449\u0438\u0439 nvcc \u2192 \u043f\u0430\u0440\u0441\u0435\u0440 \u043b\u043e\u043c\u0430\u0435\u0442\u0441\u044f    \u2192 \u0447\u0438\u043d\u044e \u0433\u0440\u0430\u043c\u043c\u0430\u0442\u0438\u043a\u0443 \u2192 \u0440\u0435\u0430\u043b\u044c\u043d\u0430\u044f GPU \u2192 \u0432\u044b\u043b\u0430\u0437\u044f\u0442 address-space \u0431\u0430\u0433\u0438    \u2192 \u0438\u0441\u043f\u0440\u0430\u0432\u043b\u044f\u044e pipeline\u041d\u0430\u043f\u0438\u0441\u0430\u0442\u044c \u043f\u0430\u0440\u0441\u0435\u0440 \u0434\u043b\u044f add\/mul\u00a0\u2014 \u043f\u0440\u043e\u0441\u0442\u043e. \u0421\u043b\u043e\u0436\u043d\u043e\u0441\u0442\u0438 \u043d\u0430\u0447\u0430\u043b\u0438\u0441\u044c \u043d\u0430\u00a0\u0432\u044b\u0445\u043b\u043e\u043f\u0435 \u0441\u0432\u0435\u0436\u0435\u0433\u043e nvcc.\u041a\u0430\u043c\u0435\u043d\u044c 1: \u043e\u0434\u043d\u043e\u0441\u0442\u0440\u043e\u0447\u043d\u044b\u0435 \u0441\u0438\u0433\u043d\u0430\u0442\u0443\u0440\u044b. \u0421\u043e\u0432\u0440\u0435\u043c\u0435\u043d\u043d\u044b\u0439 nvcc \u043f\u0438\u0448\u0435\u0442 .visible .entry _Z&#8230;(\u2026) \u0432\u00a0\u043e\u0434\u043d\u0443 \u0441\u0442\u0440\u043e\u043a\u0443, \u0430\u00a0\u043c\u043e\u0439 \u043f\u0430\u0440\u0441\u0435\u0440 \u0436\u0434\u0430\u043b \u0441\u043a\u043e\u0431\u043a\u0443 \u043d\u0430\u00a0\u043e\u0442\u0434\u0435\u043b\u044c\u043d\u043e\u0439 \u0441\u0442\u0440\u043e\u043a\u0435\u00a0\u2014 \u0438 \u0442\u0435\u043b\u043e \u044f\u0434\u0440\u0430 \u043e\u0441\u0442\u0430\u0432\u0430\u043b\u043e\u0441\u044c \u043f\u0443\u0441\u0442\u044b\u043c. \u041f\u043b\u044e\u0441 37\u00a0\u043c\u0435\u043b\u043a\u0438\u0445 \u043f\u0440\u043e\u0431\u0435\u043b\u043e\u0432 \u0432\u00a0\u0433\u0440\u0430\u043c\u043c\u0430\u0442\u0438\u043a\u0435: and\/or\/xor, \u0441\u0434\u0432\u0438\u0433\u0438, selp, div\/rem, mad.lo.s32, bfi\/bfe, f16x2, \u0432\u0430\u0440\u043f\u2011shfl. \u0411\u0430\u0437\u043e\u0432\u0430\u044f \u0430\u0440\u0438\u0444\u043c\u0435\u0442\u0438\u043a\u0430, \u043a\u043e\u0442\u043e\u0440\u0443\u044e \u043a\u043e\u043c\u043f\u0438\u043b\u044f\u0442\u043e\u0440 \u0433\u0435\u043d\u0435\u0440\u0438\u0440\u0443\u0435\u0442 \u0432\u00a0\u043a\u0430\u0436\u0434\u043e\u043c \u044f\u0434\u0440\u0435.\u041a\u0430\u043c\u0435\u043d\u044c 2: \u0438\u043d\u0432\u0435\u0440\u0441\u0438\u044f \u043f\u0440\u0435\u0434\u0438\u043a\u0430\u0442\u043e\u0432. \u0417\u0430\u043f\u0438\u0441\u044c @!%p1 bra LABEL \u043e\u0437\u043d\u0430\u0447\u0430\u0435\u0442 \u00ab\u043f\u0435\u0440\u0435\u0439\u0442\u0438, \u0435\u0441\u043b\u0438 \u043f\u0440\u0435\u0434\u0438\u043a\u0430\u0442 \u043b\u043e\u0436\u0435\u043d\u00bb\u00a0\u2014 \u0444\u043b\u0430\u0433 ! \u043e\u0431\u044f\u0437\u0430\u043d \u044f\u0432\u043d\u043e \u0438\u043d\u0432\u0435\u0440\u0442\u0438\u0440\u043e\u0432\u0430\u0442\u044c\u0441\u044f \u0432\u00a0IR. \u0423\u00a0\u043c\u0435\u043d\u044f \u044d\u0442\u043e \u0437\u0430\u043a\u0440\u044b\u0442\u043e \u043d\u0430\u00a0\u0443\u0440\u043e\u0432\u043d\u0435 \u0447\u0442\u0435\u043d\u0438\u044f \u043f\u0440\u0435\u0434\u0438\u043a\u0430\u0442\u043d\u044b\u0445 \u043e\u043f\u0435\u0440\u0430\u043d\u0434\u043e\u0432: \u0432\u0435\u0434\u0443\u0449\u0438\u0439 ! \u043f\u0440\u0435\u0432\u0440\u0430\u0449\u0430\u0435\u0442\u0441\u044f \u0432 xor i1 \u2026, true, \u0438 \u0442\u043e\u00a0\u0436\u0435 \u043f\u0440\u0430\u0432\u0438\u043b\u043e \u0434\u0435\u0439\u0441\u0442\u0432\u0443\u0435\u0442 \u0432\u00a0\u0438\u043d\u0442\u0435\u0440\u043f\u0440\u0435\u0442\u0430\u0442\u043e\u0440\u0435. \u0422\u0438\u0445\u043e\u0439 \u043f\u043e\u0440\u0447\u0438 \u0437\u0434\u0435\u0441\u044c \u043d\u0435\u0442 \u0438\u043c\u0435\u043d\u043d\u043e \u043f\u043e\u0442\u043e\u043c\u0443, \u0447\u0442\u043e\u00a0\u0438\u043d\u0432\u0435\u0440\u0441\u0438\u044f\u00a0\u2014 \u0447\u0430\u0441\u0442\u044c \u043a\u043e\u043d\u0442\u0440\u0430\u043a\u0442\u0430 \u0447\u0442\u0435\u043d\u0438\u044f \u043e\u043f\u0435\u0440\u0430\u043d\u0434\u0430, \u0430\u00a0\u043d\u0435 \u00ab\u043d\u0430\u0434\u0435\u0436\u0434\u0430, \u0447\u0442\u043e\u00a0\u0431\u044d\u043a\u0435\u043d\u0434 \u0443\u0433\u0430\u0434\u0430\u0435\u0442\u00bb.\u041a\u0430\u043c\u0435\u043d\u044c 3: setp \u0438 div. setp.lt.f32 \u0443\u00a0\u043c\u0435\u043d\u044f \u0441\u0440\u0430\u0432\u043d\u0438\u0432\u0430\u043b \u0431\u0438\u0442\u044b \u0446\u0435\u043b\u043e\u0447\u0438\u0441\u043b\u0435\u043d\u043d\u044b\u043c icmp \u0432\u043c\u0435\u0441\u0442\u043e fcmp, \u0430 div.s32 \u044d\u043c\u0438\u0442\u0438\u043b mul. \u041e\u0431\u0430 \u043d\u0435\u00a0\u043f\u0430\u0434\u0430\u043b\u0438\u00a0\u2014 \u043c\u043e\u043b\u0447\u0430 \u0441\u0447\u0438\u0442\u0430\u043b\u0438 \u043d\u0435\u043f\u0440\u0430\u0432\u0438\u043b\u044c\u043d\u043e. \u041d\u0430\u0448\u043b\u0438\u0441\u044c \u0442\u043e\u043b\u044c\u043a\u043e \u0441\u0432\u0435\u0440\u043a\u043e\u0439 \u0441\u00a0\u043d\u0430\u0441\u0442\u043e\u044f\u0449\u0438\u043c \u0432\u0445\u043e\u0434\u043e\u043c \u043e\u0442 nvcc.\u041c\u043e\u0440\u0430\u043b\u044c, \u043a\u043e\u0442\u043e\u0440\u0443\u044e \u044f \u0432\u044b\u043d\u0435\u0441: \u0442\u0435\u0441\u0442 \u043d\u0430\u00a0\u0441\u043e\u0431\u0441\u0442\u0432\u0435\u043d\u043d\u044b\u0445 \u0434\u0430\u043d\u043d\u044b\u0445 \u0434\u043e\u043a\u0430\u0437\u044b\u0432\u0430\u0435\u0442, \u0447\u0442\u043e\u00a0\u043a\u043e\u0434 \u0440\u0430\u0431\u043e\u0442\u0430\u0435\u0442 \u043d\u0430\u00a0\u0442\u0432\u043e\u0438\u0445 \u0434\u0430\u043d\u043d\u044b\u0445. \u041d\u0443\u0436\u0435\u043d \u0447\u0443\u0436\u043e\u0439 \u0432\u0445\u043e\u0434\u00a0\u2014 \u043d\u0430\u0441\u0442\u043e\u044f\u0449\u0438\u0439 PTX \u0438\u0437\u00a0CI (compiler_nvcc_ptx).\u0410\u00a0\u043f\u043e\u0442\u043e\u043c \u0432\u0438\u0434\u0435\u043e\u043a\u0430\u0440\u0442\u0430 \u043d\u0430\u0448\u043b\u0430 \u0435\u0449\u0451 \u0434\u0432\u0430 \u0431\u0430\u0433\u0430, \u043d\u0435\u0432\u0438\u0434\u0438\u043c\u044b\u0445 \u043d\u0430\u00a0\u0445\u043e\u0441\u0442\u0435: \u0432\u044b\u0432\u043e\u0434 \u043c\u0430\u0442\u0440\u0438\u0446\u044b \u043f\u0438\u0441\u0430\u043b\u0441\u044f \u0432\u00a0shared \u0432\u043c\u0435\u0441\u0442\u043e global (\u043c\u0435\u0442\u043a\u0430 \u0440\u0435\u0433\u0438\u0441\u0442\u0440\u0430 \u043f\u0435\u0440\u0435\u0436\u0438\u0432\u0430\u043b\u0430 \u043f\u0435\u0440\u0435\u0437\u0430\u043f\u0438\u0441\u044c\u00a0\u2014 \u0442\u0435\u043f\u0435\u0440\u044c \u043f\u0440\u043e\u0441\u0442\u0440\u0430\u043d\u0441\u0442\u0432\u043e \u0431\u0435\u0440\u0451\u0442\u0441\u044f \u0438\u0437\u00a0\u0441\u0443\u0444\u0444\u0438\u043a\u0441\u0430 \u0438\u043d\u0441\u0442\u0440\u0443\u043a\u0446\u0438\u0438), \u0430\u00a0\u0430\u0434\u0440\u0435\u0441 device\u2011\u0433\u043b\u043e\u0431\u0430\u043b\u0430 \u0447\u0438\u0442\u0430\u043b\u0441\u044f \u043a\u0430\u043a\u00a0\u0437\u0430\u0433\u0440\u0443\u0437\u043a\u0430 \u0438\u0437\u00a0\u043d\u0435\u0433\u043e. \u041d\u0430\u00a0\u043f\u043b\u043e\u0441\u043a\u043e\u0439 \u0445\u043e\u0441\u0442\u2011\u043f\u0430\u043c\u044f\u0442\u0438 \u043e\u0431\u0430 \u043c\u043e\u043b\u0447\u0430\u043b\u0438.4. \u041c\u043e\u0438 \u0440\u0435\u0437\u0443\u043b\u044c\u0442\u0430\u0442\u044b \u0438 \u0442\u0435\u0441\u0442\u044bctest\u00a0\u2014 21\/21, \u043d\u043e\u043b\u044c \u0432\u043e\u0440\u043d\u0438\u043d\u0433\u043e\u0432. \u0418\u0437\u00a0\u043d\u0438\u0445:7\u00a0\u043a\u0430\u043d\u043e\u043d\u0438\u0447\u0435\u0441\u043a\u0438\u0445 \u044f\u0434\u0435\u0440 (vadd, block_reduce, atomic_reduce, fp16_reduce, gemm_tile, shfl_reduce, atom_cas) \u0438\u0441\u043f\u043e\u043b\u043d\u0435\u043d\u044b \u043d\u0430\u00a0GeForce MX450\u00a0\u0438 \u0441\u043e\u0448\u043b\u0438\u0441\u044c \u043f\u043e\u0431\u0438\u0442\u043e\u0432\u043e \u0441\u00a0\u0445\u043e\u0441\u0442\u2011\u044d\u0442\u0430\u043b\u043e\u043d\u043e\u043c;\u043d\u0430\u0441\u0442\u043e\u044f\u0449\u0435\u0435 \u044f\u0434\u0440\u043e \u043e\u0442 nvcc (\u043c\u0430\u043d\u0433\u043b\u0438\u0440\u043e\u0432\u0430\u043d\u043d\u043e\u0435 \u0438\u043c\u044f, device\u2011\u0433\u043b\u043e\u0431\u0430\u043b, syncthreads)\u00a0\u2014 \u043d\u0430\u0441\u043a\u0432\u043e\u0437\u044c nvcc \u2192 ptx2ir \u2192 GPU. \u0427\u0435\u0441\u0442\u043d\u0430\u044f \u043e\u0433\u043e\u0432\u043e\u0440\u043a\u0430: \u0432\u00a0\u0435\u0433\u043e \u0438\u0441\u0445\u043e\u0434\u043d\u0438\u043a\u0435 \u0433\u043e\u043d\u043a\u0430 (\u0440\u0435\u0434\u0443\u043a\u0446\u0438\u044f \u0431\u0435\u0437 syncthreads), \u0432\u0435\u043d\u0434\u043e\u0440\u043d\u0430\u044f \u0441\u0431\u043e\u0440\u043a\u0430 \u0433\u043e\u043d\u0438\u0442\u0441\u044f \u0442\u0430\u043a\u00a0\u0436\u0435\u00a0\u2014 \u044d\u0442\u0430\u043b\u043e\u043d \u043f\u0440\u0438\u043d\u0438\u043c\u0430\u0435\u0442 \u0434\u043e\u043a\u0430\u0437\u0430\u043d\u043d\u044b\u0439 \u0438\u043d\u0442\u0435\u0440\u0432\u0430\u043b, \u0430\u00a0\u043d\u0435\u00a0\u043e\u0434\u043d\u043e \u0447\u0438\u0441\u043b\u043e;shfl.sync.bfly \u043f\u0440\u043e\u0432\u0435\u0440\u0435\u043d \u043e\u0442\u0434\u0435\u043b\u044c\u043d\u044b\u043c NVPTX\u2011\u0442\u0435\u0441\u0442\u043e\u043c (shfl_reduce) \u043d\u0430\u00a0\u0440\u0435\u0430\u043b\u044c\u043d\u043e\u0439 GPU\u00a0\u2014 \u0447\u0435\u0440\u0435\u0437 \u043d\u0430\u0441\u0442\u043e\u044f\u0449\u0438\u0439 NVVM\u2011\u0438\u043d\u0442\u0440\u0438\u043d\u0441\u0438\u043a; \u043e\u0441\u0442\u0430\u043b\u044c\u043d\u044b\u0435 warp\u2011\u043a\u043e\u043b\u043b\u0435\u043a\u0442\u0438\u0432\u044b \u043f\u043e\u043a\u0440\u044b\u0442\u044b\u043c\u0438 \u0442\u0435\u0441\u0442\u0430\u043c\u0438 \u043d\u0435\u00a0\u0441\u0447\u0438\u0442\u0430\u044e: vote \/ activemask \/ \u043f\u0440\u0435\u0434\u0438\u043a\u0430\u0442\u043d\u044b\u0439 shfl\u00a0\u2014 \u044f\u0432\u043d\u044b\u0439 \u043e\u0442\u043a\u0430\u0437, \u0430\u00a0\u043d\u0435\u00a0\u0442\u0438\u0445\u0438\u0439 \u0444\u043e\u043b\u0431\u044d\u043a;\u0431\u0435\u0437\u00a0\u043a\u0430\u0440\u0442\u044b sm_75+ GPU\u2011\u0442\u0435\u0441\u0442\u044b \u0433\u043e\u0432\u043e\u0440\u044f\u0442 \u0447\u0435\u0441\u0442\u043d\u044b\u0439 SKIP, \u0430\u00a0\u043d\u0435 \u00ab\u0443\u0441\u043f\u0435\u0445\u00bb.\u041f\u0440\u043e\u0438\u0437\u0432\u043e\u0434\u0438\u0442\u0435\u043b\u044c\u043d\u043e\u0441\u0442\u0438 \u044f \u043d\u0435\u00a0\u043c\u0435\u0440\u044f\u043b\u00a0\u2014 \u0442\u043e\u043b\u044c\u043a\u043e \u0431\u0438\u0442\u2011\u0438\u0434\u0435\u043d\u0442\u0438\u0447\u043d\u043e\u0441\u0442\u044c. \u041a\u0442\u043e \u043e\u0431\u0435\u0449\u0430\u0435\u0442 \u0441\u043a\u043e\u0440\u043e\u0441\u0442\u044c \u0431\u0435\u0437\u00a0\u0437\u0430\u043c\u0435\u0440\u043e\u0432, \u0442\u043e\u0442 \u0432\u0440\u0451\u0442; \u044f \u043d\u0435\u00a0\u043e\u0431\u0435\u0449\u0430\u044e.5. \u0427\u0442\u043e\u00a0\u0434\u0430\u043b\u044c\u0448\u0435?SIMT \u2192 RVV \u0432\u0435\u043a\u0442\u043e\u0440\u0438\u0437\u0430\u0446\u0438\u044f: \u0441\u0432\u043e\u0440\u0430\u0447\u0438\u0432\u0430\u0442\u044c 32\u00a0\u043f\u043e\u0442\u043e\u043a\u0430 \u0432\u00a0\u0432\u0435\u043a\u0442\u043e\u0440\u043d\u044b\u0435\u00a0\u043b\u0438\u043d\u0438\u0438 (MLIR\u2011\u043a\u043e\u043d\u0442\u0443\u0440 \u0443\u0436\u0435 \u0430\u043d\u0430\u043b\u0438\u0437\u0438\u0440\u0443\u0435\u0442 \u043f\u0430\u0442\u0442\u0435\u0440\u043d\u044b, \u043f\u043e\u043b\u043d\u043e\u0446\u0435\u043d\u043d\u043e\u0439 \u0432\u0435\u043a\u0442\u043e\u0440\u0438\u0437\u0430\u0446\u0438\u0438 \u043f\u043e\u043a\u0430 \u043d\u0435\u0442);mma \u0438 \u0442\u0435\u043d\u0437\u043e\u0440\u043d\u044b\u0435 \u0438\u043d\u0441\u0442\u0440\u0443\u043a\u0446\u0438\u0438\u00a0\u2014 \u043f\u043e\u043a\u0430 \u043e\u0441\u043e\u0437\u043d\u0430\u043d\u043d\u044b\u0439 \u043e\u0442\u043a\u0430\u0437: \u00ab\u043f\u0440\u0438\u043c\u0435\u0440\u043d\u043e \u043f\u0440\u0430\u0432\u0438\u043b\u044c\u043d\u043e\u00bb \u043e\u043f\u0430\u0441\u043d\u0435\u0435, \u0447\u0435\u043c \u043d\u0438\u043a\u0430\u043a;\u0431\u043e\u043b\u044c\u0448\u0435 \u043d\u0430\u0441\u0442\u043e\u044f\u0449\u0438\u0445 CUDA\u2011\u0431\u0438\u043d\u0430\u0440\u043d\u0438\u043a\u043e\u0432 \u043d\u0430\u00a0\u0432\u0445\u043e\u0434\u0435 \u0432\u043c\u0435\u0441\u0442\u043e \u0441\u0438\u043d\u0442\u0435\u0442\u0438\u043a\u0438;\u043f\u0443\u0431\u043b\u0438\u0447\u043d\u044b\u0435 \u0431\u0435\u043d\u0447\u043c\u0430\u0440\u043a\u0438\u00a0\u2014 \u043d\u043e\u00a0\u0441\u043d\u0430\u0447\u0430\u043b\u0430 \u043f\u043e\u0432\u0442\u043e\u0440\u044f\u0435\u043c\u044b\u0439 \u0441\u0442\u0435\u043d\u0434.\u0413\u0434\u0435 \u043f\u043e\u0449\u0443\u043f\u0430\u0442\u044c\u0411\u0435\u0437\u00a0\u0441\u0431\u043e\u0440\u043a\u0438, \u0434\u0432\u0430 \u0431\u0438\u043d\u0430\u0440\u043d\u0438\u043a\u0430 (\u043a\u0440\u043e\u043c\u0435 \u0441\u0438\u0441\u0442\u0435\u043c\u043d\u043e\u0433\u043e libc \u0437\u0430\u0432\u0438\u0441\u0438\u043c\u043e\u0441\u0442\u0435\u0439 \u043d\u0435\u0442):curl -fsSL https:\/\/github.com\/MrModelOS\/ROUGE-V\/releases\/download\/v0.1\/rouge-v-0.1-linux-x86\\\\_64.tar.gz | tar xz -C ~\/.localexport PATH=&#187;$HOME\/.local\/bin:$PATH&#187;ptx2ir kernel.ptx kernel.ll\u0418\u0437\u00a0\u0438\u0441\u0445\u043e\u0434\u043d\u0438\u043a\u043e\u0432:git clone https:\/\/github.com\/MrModelOS\/ROUGE-V.git &amp;&amp; cd ROUGE-Vcmake -B build -G Ninja -DCMAKE_BUILD_TYPE=Release &amp;&amp; cmake &#8212;build buildctest &#8212;test-dir build &#8212;output-on-failure\u0410\u0440\u0445\u0438\u0442\u0435\u043a\u0442\u0443\u0440\u0430\u00a0\u2014 docs\/06-compiler-architecture.md.\u00a0\u041b\u0438\u0446\u0435\u043d\u0437\u0438\u044f\u00a0\u2014 Apache 2.0\u00a0WITH LLVM\u2011exception.\u041f\u0440\u043e\u00a0\u044e\u0440\u0438\u0434\u0438\u0447\u0435\u0441\u043a\u0443\u044e \u0441\u0442\u043e\u0440\u043e\u043d\u0443\u00a0\u2014 \u043f\u0440\u044f\u043c\u043e\u042f \u043d\u0435\u00a0\u043a\u043e\u043f\u0438\u0440\u0443\u044e \u0438 \u043d\u0435\u00a0\u0440\u0435\u0432\u0435\u0440\u0441\u2011\u0438\u043d\u0436\u0438\u043d\u0438\u0440\u044e \u043f\u0440\u043e\u043f\u0440\u0438\u0435\u0442\u0430\u0440\u043d\u044b\u0435 \u0431\u0438\u043d\u0430\u0440\u043d\u0438\u043a\u0438 \u0438 \u0437\u0430\u0433\u043e\u043b\u043e\u0432\u043a\u0438 NVIDIA.PTX, \u043a\u043e\u0442\u043e\u0440\u044b\u0439 \u044f \u043f\u0435\u0440\u0435\u0432\u043e\u0436\u0443, \u0432\u044b\u043f\u0443\u0441\u043a\u0430\u0435\u0442 nvcc \u043f\u043e\u043b\u044c\u0437\u043e\u0432\u0430\u0442\u0435\u043b\u044f\u00a0\u2014 \u0438\u0437\u00a0\u0435\u0433\u043e \u0441\u043e\u0431\u0441\u0442\u0432\u0435\u043d\u043d\u043e\u0439\u00a0\u043b\u0438\u0446\u0435\u043d\u0437\u0438\u0438, \u043f\u043e\u00a0\u043e\u0442\u043a\u0440\u044b\u0442\u043e\u0439 \u0441\u043f\u0435\u0446\u0438\u0444\u0438\u043a\u0430\u0446\u0438\u0438 PTX ISA.ROUGE\u2011V \u043d\u0435\u00a0\u0441\u0432\u044f\u0437\u0430\u043d \u0441\u00a0NVIDIA \u0438 \u043d\u0435\u00a0\u043e\u0434\u043e\u0431\u0440\u0435\u043d \u0438\u043c\u0438. \u0422\u043e\u0440\u0433\u043e\u0432\u044b\u0435 \u043c\u0430\u0440\u043a\u0438\u00a0\u2014 \u0442\u043e\u043b\u044c\u043a\u043e \u0447\u0442\u043e\u0431\u044b \u043e\u043f\u0438\u0441\u0430\u0442\u044c \u0441\u043e\u0432\u043c\u0435\u0441\u0442\u0438\u043c\u043e\u0441\u0442\u044c.P.\u202fS.\u0415\u0441\u043b\u0438 \u0432\u043e\u0437\u044c\u043c\u0451\u0442\u0435 PTX \u043e\u0442\u00a0\u0440\u0435\u0430\u043b\u044c\u043d\u043e\u0433\u043e nvcc \u0438 \u043c\u043e\u0439 \u0442\u0440\u0430\u043d\u0441\u043b\u044f\u0442\u043e\u0440 \u0441\u043a\u0430\u0436\u0435\u0442 \u043e\u0442\u043a\u0430\u0437\u00a0\u2014 \u0440\u0430\u0441\u0441\u043a\u0430\u0436\u0438\u0442\u0435, \u0447\u0442\u043e\u00a0\u0438\u043c\u0435\u043d\u043d\u043e: \u044d\u0442\u043e \u0441\u0430\u043c\u044b\u0439 \u043f\u043e\u043b\u0435\u0437\u043d\u044b\u0439 \u0431\u0430\u0433\u2011\u0440\u0435\u043f\u043e\u0440\u0442 \u0434\u043b\u044f\u00a0\u043f\u0440\u043e\u0435\u043a\u0442\u0430.\u0441\u0441\u044b\u043b\u043a\u0430 \u043d\u0430 \u043e\u0440\u0438\u0433\u0438\u043d\u0430\u043b \u0441\u0442\u0430\u0442\u044c\u0438 https:\/\/habr.com\/ru\/articles\/1089146\/<\/p>\n","protected":false},"author":1,"featured_media":0,"comment_status":"closed","ping_status":"closed","sticky":false,"template":"","format":"standard","meta":{"footnotes":""},"categories":[],"tags":[],"class_list":["post-496918","post","type-post","status-publish","format-standard","hentry"],"_links":{"self":[{"href":"https:\/\/savepearlharbor.com\/index.php?rest_route=\/wp\/v2\/posts\/496918","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=496918"}],"version-history":[{"count":0,"href":"https:\/\/savepearlharbor.com\/index.php?rest_route=\/wp\/v2\/posts\/496918\/revisions"}],"wp:attachment":[{"href":"https:\/\/savepearlharbor.com\/index.php?rest_route=%2Fwp%2Fv2%2Fmedia&parent=496918"}],"wp:term":[{"taxonomy":"category","embeddable":true,"href":"https:\/\/savepearlharbor.com\/index.php?rest_route=%2Fwp%2Fv2%2Fcategories&post=496918"},{"taxonomy":"post_tag","embeddable":true,"href":"https:\/\/savepearlharbor.com\/index.php?rest_route=%2Fwp%2Fv2%2Ftags&post=496918"}],"curies":[{"name":"wp","href":"https:\/\/api.w.org\/{rel}","templated":true}]}}