{"id":470994,"date":"2025-08-16T21:00:24","date_gmt":"2025-08-16T21:00:24","guid":{"rendered":"http:\/\/savepearlharbor.com\/?p=470994"},"modified":"-0001-11-30T00:00:00","modified_gmt":"-0001-11-29T21:00:00","slug":"","status":"publish","type":"post","link":"https:\/\/savepearlharbor.com\/?p=470994","title":{"rendered":"<span>\u041a\u0430\u043a \u043f\u0440\u0430\u0432\u0438\u043b\u044c\u043d\u043e \u0432\u044b\u0437\u044b\u0432\u0430\u0442\u044c CUDA<\/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<p>\u0412\u0435\u0440\u043e\u044f\u0442\u043d\u043e, \u0432\u0430\u043c \u0443\u0436\u0435 \u043f\u043e\u043f\u0430\u0434\u0430\u043b\u0438\u0441\u044c \u043f\u043e\u0434\u043e\u0431\u043d\u044b\u0435 \u0440\u0443\u043a\u043e\u0432\u043e\u0434\u0441\u0442\u0432\u0430 \u043f\u043e CUDA: \u0445\u0440\u0435\u0441\u0442\u043e\u043c\u0430\u0442\u0438\u0439\u043d\u044b\u0439 \u043f\u0440\u0438\u043c\u0435\u0440 \u00abHello World\u00bb, \u0432 \u043a\u043e\u0442\u043e\u0440\u043e\u043c \u043f\u0435\u0440\u0435\u043c\u0435\u0448\u0430\u043d \u043a\u043e\u0434 \u0434\u043b\u044f \u0426\u041f \u0438 \u0433\u0440\u0430\u0444\u0438\u0447\u0435\u0441\u043a\u043e\u0433\u043e \u043f\u0440\u043e\u0446\u0435\u0441\u0441\u043e\u0440\u0430. \u0412\u0441\u0451 \u044d\u0442\u043e \u0441\u043b\u043e\u0436\u0435\u043d\u043e \u0432 \u043e\u0434\u0438\u043d \u0433\u0435\u0442\u0435\u0440\u043e\u0433\u0435\u043d\u043d\u044b\u0439 \u0444\u0430\u0439\u043b \u0441 \u0438\u0441\u0445\u043e\u0434\u043d\u0438\u043a\u0430\u043c\u0438 \u043d\u0430 CUDA C++, \u0430 \u0434\u043b\u044f \u0437\u0430\u043f\u0443\u0441\u043a\u0430 \u044f\u0434\u0440\u0430 \u043f\u0440\u0438\u043c\u0435\u043d\u044f\u0435\u0442\u0441\u044f \u0441\u0438\u043d\u0442\u0430\u043a\u0441\u0438\u0441 NVCC \u0441 \u0442\u0440\u043e\u0439\u043d\u044b\u043c\u0438 \u0443\u0433\u043b\u043e\u0432\u044b\u043c\u0438 \u0441\u043a\u043e\u0431\u043a\u0430\u043c\u0438\u00a0<code>&lt;&lt;&lt;&gt;&gt;&gt;<\/code>, \u043a\u043e\u0442\u043e\u0440\u044b\u0439 \u0443\u0436\u0435 \u0441\u0442\u0430\u043b \u043a\u0443\u043b\u044c\u0442\u043e\u0432\u044b\u043c:<\/p>\n<pre><code class=\"cpp\">#include &lt;cuda_runtime.h&gt; #include &lt;stdio.h&gt;  __global__ void kernel() {     printf(\"Hello World from block %d, thread %d\\n\", blockIdx.x, threadIdx.x); }  int main() {     kernel&lt;&lt;&lt;1, 1&gt;&gt;&gt;(); \/\/ \u0412\u043e\u0437\u0432\u0440\u0430\u0449\u0430\u0435\u0442 `void`?!          return cudaDeviceSynchronize() == cudaSuccess ? 0 : -1; }<\/code><\/pre>\n<p>\u0412\u0440\u0435\u043c\u044f \u0438\u0434\u0451\u0442, \u0430 \u0442\u0430\u043a\u043e\u0439 \u043f\u0430\u0442\u0442\u0435\u0440\u043d \u043f\u043e-\u043f\u0440\u0435\u0436\u043d\u0435\u043c\u0443 \u043f\u043e\u043f\u0430\u0434\u0430\u0435\u0442\u0441\u044f \u043c\u043d\u0435 \u0432 \u043f\u0440\u043e\u0434\u0430\u043a\u0448\u0435\u043d-\u043a\u043e\u0434\u0435. \u041f\u0440\u0438\u0437\u043d\u0430\u044e\u0441\u044c, \u043a\u043e\u0435-\u0433\u0434\u0435 \u043e\u043d \u0432\u0441\u043f\u043b\u044b\u0432\u0430\u0435\u0442 \u0438 \u0432 \u043c\u043e\u0438\u0445 \u043b\u044e\u0431\u0438\u0442\u0435\u043b\u044c\u0441\u043a\u0438\u0445 \u043f\u0440\u043e\u0435\u043a\u0442\u0430\u0445 \u2014 <a href=\"https:\/\/github.com\/ashvardanian\/ParallelReductionsBenchmark\" rel=\"noopener noreferrer nofollow\">\u0440\u0430\u0437<\/a>,\u00a0<a href=\"https:\/\/github.com\/ashvardanian\/cuda-python-starter-kit\" rel=\"noopener noreferrer nofollow\">\u0434\u0432\u0430<\/a>, <a href=\"https:\/\/github.com\/ashvardanian\/scaling-democracy\" rel=\"noopener noreferrer nofollow\">\u0442\u0440\u0438<\/a>. \u041d\u043e \u044d\u0442\u043e \u043d\u0435 \u043b\u0443\u0447\u0448\u0430\u044f \u0438\u0434\u0435\u044f, \u043f\u043e\u043b\u0430\u0433\u0430\u0442\u044c\u0441\u044f \u0432 \u0441\u0435\u0440\u044c\u0451\u0437\u043d\u043e\u043c \u043a\u043e\u0434\u0435 \u043d\u0430 \u0437\u0430\u043f\u0443\u0441\u043a \u044f\u0434\u0440\u0430 \u0447\u0435\u0440\u0435\u0437 \u0442\u0440\u043e\u0439\u043d\u044b\u0435 \u0443\u0433\u043b\u043e\u0432\u044b\u0435 \u0441\u043a\u043e\u0431\u043a\u0438. \u0412 \u0442\u0430\u043a\u043e\u043c \u0441\u043b\u0443\u0447\u0430\u0435 \u043f\u0440\u043e\u0433\u0440\u0430\u043c\u043c\u0430 \u043d\u0435 \u0432\u043e\u0437\u0432\u0440\u0430\u0449\u0430\u0435\u0442 \u043a\u043e\u0434\u044b \u043e\u0448\u0438\u0431\u043e\u043a, \u043f\u043e\u044d\u0442\u043e\u043c\u0443 \u043c\u043e\u0436\u0435\u0442 \u043f\u043e\u043a\u0430\u0437\u0430\u0442\u044c\u0441\u044f \u043e\u0431\u043c\u0430\u043d\u0447\u0438\u0432\u043e \u043f\u0440\u043e\u0441\u0442\u043e\u0439. \u041d\u0438\u0436\u0435 \u0432\u0430\u0441 \u0436\u0434\u0443\u0442 \u043f\u0440\u0438\u043c\u0435\u0440\u043d\u043e 25 \u043a\u0438\u043b\u043e\u0431\u0430\u0439\u0442 \u0442\u0435\u043a\u0441\u0442\u0430, \u0432 \u043a\u043e\u0442\u043e\u0440\u044b\u0445 \u043c\u044b \u043e\u0431\u0441\u0443\u0434\u0438\u043c <em>\u043d\u0435 \u0441\u0430\u043c\u044b\u0435 \u043a\u043e\u0440\u044f\u0432\u044b\u0435<\/em> \u0441\u043f\u043e\u0441\u043e\u0431\u044b \u0437\u0430\u043f\u0443\u0441\u043a\u0430 \u044f\u0434\u0435\u0440.<\/p>\n<h2>\u041e\u0441\u043d\u043e\u0432\u044b \u0438 \u043a\u043e\u0440\u0440\u0435\u043a\u0442\u043d\u043e\u0441\u0442\u044c<\/h2>\n<p>\u0412\u044b\u0448\u0435\u043f\u0440\u0438\u0432\u0435\u0434\u0451\u043d\u043d\u044b\u0439 \u043a\u043e\u0434 \u0441\u043a\u043e\u043c\u043f\u0438\u043b\u0438\u0440\u0443\u0435\u0442\u0441\u044f, \u043f\u043e\u0441\u043b\u0435 \u0447\u0435\u0433\u043e \u0432\u044b\u043f\u043e\u043b\u043d\u0438\u0442 \u043e\u0436\u0438\u0434\u0430\u0435\u043c\u044b\u0439 \u0432\u044b\u0432\u043e\u0434: <\/p>\n<pre><code>$ nvcc -o hello_world hello_world.cu &amp;&amp; .\/hello_world &gt; Hello World from block 0, thread 0<\/code><\/pre>\n<p>\u0412 \u043a\u0430\u043a\u043e\u043c-\u0442\u043e \u0441\u043c\u044b\u0441\u043b\u0435 \u0435\u0433\u043e \u0443\u0436\u0435 \u043c\u043e\u0436\u043d\u043e \u0441\u0447\u0438\u0442\u0430\u0442\u044c \u00ab\u043a\u043e\u0440\u0440\u0435\u043a\u0442\u043d\u044b\u043c\u00bb.<\/p>\n<p>\u041d\u043e \u0432 \u043d\u0430\u0448\u0435 \u0432\u0440\u0435\u043c\u044f \u043e\u0431\u044b\u0447\u043d\u044b \u0441\u0438\u0441\u0442\u0435\u043c\u044b, \u0432 \u043a\u043e\u0442\u043e\u0440\u044b\u0445 \u043f\u0440\u0435\u0434\u0443\u0441\u043c\u0430\u0442\u0440\u0438\u0432\u0430\u0435\u0442\u0441\u044f \u043d\u0435 \u043c\u0435\u043d\u0435\u0435 8 GPU \u043d\u0430 \u043f\u043b\u0430\u0442\u0443 HGX, \u043f\u043e\u044d\u0442\u043e\u043c\u0443, \u0435\u0441\u0442\u0435\u0441\u0442\u0432\u0435\u043d\u043d\u043e, \u0445\u043e\u0447\u0435\u0442\u0441\u044f \u0440\u0435\u0430\u043b\u0438\u0437\u043e\u0432\u0430\u0442\u044c \u043d\u0435\u043a\u043e\u0442\u043e\u0440\u044b\u0439 \u043f\u0430\u0440\u0430\u043b\u043b\u0435\u043b\u0438\u0437\u043c \u2014 \u0447\u0442\u043e\u0431\u044b \u0433\u0430\u0440\u0430\u043d\u0442\u0438\u0440\u043e\u0432\u0430\u0442\u044c, \u0447\u0442\u043e \u0431\u0443\u0434\u0443\u0442 \u0437\u0430\u043f\u0443\u0441\u043a\u0430\u0442\u044c\u0441\u044f \u044f\u0434\u0440\u0430 \u043d\u0430 \u043a\u0430\u0436\u0434\u043e\u043c \u0438\u0437 \u043d\u0438\u0445.<\/p>\n<p> \u041e\u043f\u0443\u0441\u043a\u0430\u044f \u0431\u0430\u0437\u043e\u0432\u044b\u0435 \u0430\u0441\u043f\u0435\u043a\u0442\u044b \u044d\u043a\u0441\u043f\u043b\u0443\u0430\u0442\u0430\u0446\u0438\u0438, \u043e\u0442\u043c\u0435\u0447\u0443, \u0447\u0442\u043e \u043a\u0430\u0436\u0434\u044b\u0439 \u0443\u0437\u0435\u043b DGX H100 \u043e\u0441\u043d\u0430\u0449\u0430\u0435\u0442\u0441\u044f \u0434\u0432\u0443\u043c\u044f \u043c\u043e\u0433\u0443\u0447\u0438\u043c\u0438 \u044f\u0434\u0440\u0430\u043c\u0438 \u0426\u041f, \u043a\u043e\u0442\u043e\u0440\u044b\u0435 \u0431\u0435\u0437 \u043f\u0435\u0440\u0435\u0440\u044b\u0432\u0430 \u0436\u043e\u043d\u0433\u043b\u0438\u0440\u0443\u044e\u0442 \u0441\u043b\u043e\u0436\u043d\u044b\u043c\u0438 \u0433\u0440\u0430\u0444\u0430\u043c\u0438 \u0432\u044b\u043f\u043e\u043b\u043d\u0435\u043d\u0438\u044f, \u043e\u0434\u043d\u043e\u0432\u0440\u0435\u043c\u0435\u043d\u043d\u043e \u043f\u0435\u0440\u0435\u0434\u0430\u0432\u0430\u044f \u043d\u0430 \u0432\u0445\u043e\u0434 \u0438 \u043d\u0430 \u0432\u044b\u0445\u043e\u0434 \u0441\u043e\u0442\u043d\u0438 \u0433\u0438\u0433\u0430\u0431\u0430\u0439\u0442. \u0412\u0441\u0451 \u044d\u0442\u043e \u2014 \u0430\u0441\u0438\u043d\u0445\u0440\u043e\u043d\u043d\u043e.<\/p>\n<p>\u0412 \u0442\u0430\u043a\u043e\u0439 \u0441\u0438\u0441\u0442\u0435\u043c\u0435 \u043e\u0447\u0435\u043d\u044c \u043c\u043d\u043e\u0433\u043e\u0435 \u043c\u043e\u0436\u0435\u0442 \u0441\u0431\u0438\u0442\u044c\u0441\u044f, \u0442\u0430\u043a \u0447\u0442\u043e \u0434\u0430\u0432\u0430\u0439\u0442\u0435 \u0441\u0444\u043e\u0440\u043c\u0443\u043b\u0438\u0440\u0443\u0435\u043c \u043d\u0435\u0441\u043a\u043e\u043b\u044c\u043a\u043e \u043e\u0441\u043d\u043e\u0432\u043e\u043f\u043e\u043b\u0430\u0433\u0430\u044e\u0449\u0438\u0445 \u043f\u0440\u0430\u0432\u0438\u043b \u043e\u0440\u043a\u0435\u0441\u0442\u0440\u0430\u0446\u0438\u0438 \u044f\u0434\u0435\u0440 GPU:<\/p>\n<ul>\n<li>\n<p>\u0417\u0430\u043f\u0443\u0441\u043a \u043a\u0430\u0436\u0434\u043e\u0433\u043e \u044f\u0434\u0440\u0430 \u0441\u043e\u043f\u0440\u044f\u0436\u0451\u043d \u0441 \u0441\u0435\u0440\u044c\u0451\u0437\u043d\u043e\u0439 \u0437\u0430\u0434\u0435\u0440\u0436\u043a\u043e\u0439, \u043f\u043e\u044d\u0442\u043e\u043c\u0443 \u0442\u0430\u043a\u0438\u0435 \u043e\u043f\u0435\u0440\u0430\u0446\u0438\u0438 \u0434\u043e\u043b\u0436\u043d\u044b \u0432\u044b\u043f\u043e\u043b\u043d\u044f\u0442\u044c\u0441\u044f \u0430\u0441\u0438\u043d\u0445\u0440\u043e\u043d\u043d\u043e.<\/p>\n<\/li>\n<li>\n<p>\u0420\u0430\u0431\u043e\u0442\u0443 \u0432 \u043f\u0440\u0435\u0434\u0435\u043b\u0430\u0445 \u043f\u043e\u0442\u043e\u043a\u043e\u0432 \u0441\u043b\u0435\u0434\u0443\u0435\u0442 \u0440\u0430\u0437\u0434\u0430\u0432\u0430\u0442\u044c \u044f\u0432\u043d\u043e.<\/p>\n<\/li>\n<li>\n<p>\u0412\u044b\u0437\u043e\u0432\u044b API CUDA \u0438 \u043f\u0443\u0441\u043a\u0438 \u044f\u0434\u0440\u0430 \u0434\u043e\u043b\u0436\u043d\u044b \u0441\u043e\u043f\u0440\u043e\u0432\u043e\u0436\u0434\u0430\u0442\u044c\u0441\u044f \u043d\u0430\u0434\u0451\u0436\u043d\u043e\u0439 \u043f\u0440\u043e\u0432\u0435\u0440\u043a\u043e\u0439 \u043d\u0430 \u0441\u043b\u0443\u0447\u0430\u0439 \u0432\u043e\u0437\u043c\u043e\u0436\u043d\u044b\u0445 \u043e\u0448\u0438\u0431\u043e\u043a.<\/p>\n<\/li>\n<\/ul>\n<p>\u0412\u043e\u0442 \u043a\u0430\u043a \u0438\u043d\u0442\u0435\u0433\u0440\u0438\u0440\u043e\u0432\u0430\u0442\u044c \u043f\u043e\u0442\u043e\u043a\u0438 CUDA \u2014 \u043d\u0438\u0447\u0435\u0433\u043e \u0441\u0432\u0435\u0440\u0445\u044a\u0435\u0441\u0442\u0435\u0441\u0442\u0432\u0435\u043d\u043d\u043e\u0433\u043e:<\/p>\n<pre><code class=\"cpp\">#include &lt;cuda_runtime.h&gt; #include &lt;stdio.h&gt;  __global__ void kernel() {     extern __shared__ char shared_buffer[];     printf(\"Hello World from block %d, thread %d\\n\", blockIdx.x, threadIdx.x); }  int main() {     cudaStream_t stream;     cudaStreamCreate(&amp;stream);     uint shared_memory_size = 0;     kernel&lt;&lt;&lt;1, 1, shared_memory_size, stream&gt;&gt;&gt;(); \/\/ 4 \u0430\u0440\u0433\u0443\u043c\u0435\u043d\u0442\u0430, \u0430 \u043d\u0435 2     cudaStreamSynchronize(stream);     cudaStreamDestroy(stream);     return 0; }<\/code><\/pre>\n<p>\u041e\u0431\u0440\u0430\u0442\u0438\u0442\u0435 \u0432\u043d\u0438\u043c\u0430\u043d\u0438\u0435: \u044f\u0434\u0440\u043e \u0437\u0430\u043f\u0443\u0441\u043a\u0430\u0435\u0442\u0441\u044f \u0441 \u0447\u0435\u0442\u044b\u0440\u044c\u043c\u044f \u0430\u0440\u0433\u0443\u043c\u0435\u043d\u0442\u0430\u043c\u0438.<\/p>\n<p>\u0412\u043e\u0442 \u0431\u043e\u043b\u0435\u0435 \u0430\u043a\u043a\u0443\u0440\u0430\u0442\u043d\u0430\u044f \u0432\u0435\u0440\u0441\u0438\u044f \u0441 \u044f\u0432\u043d\u043e\u0439 \u043e\u0431\u0440\u0430\u0431\u043e\u0442\u043a\u043e\u0439 \u043e\u0448\u0438\u0431\u043e\u043a:<\/p>\n<pre><code class=\"cpp\">#include &lt;cuda_runtime.h&gt; #include &lt;stdio.h&gt;  __global__ void kernel() {     extern __shared__ char shared_buffer[];     printf(\"Hello World from block %d, thread %d\\n\", blockIdx.x, threadIdx.x); }  int main() {     cudaStream_t stream;     cudaError_t err = cudaStreamCreate(&amp;stream);     if (err != cudaSuccess) {         fprintf(stderr, \"Failed to create stream: %s\\n\", cudaGetErrorString(err));         return -1;     }     uint shared_memory_size = 1 &lt;&lt; 30; \/\/ 1 \u0413\u0411 \u2013 \u044d\u0442\u043e \u043c\u043d\u043e\u0433\u043e, \u043d\u043e \u0432 \u0434\u0435\u043c\u043e\u043d\u0441\u0442\u0440\u0430\u0446\u0438\u043e\u043d\u043d\u044b\u0445 \u0446\u0435\u043b\u044f\u0445 \u0442\u0430\u043a\u0430\u044f \u0432\u0435\u043b\u0438\u0447\u0438\u043d\u0430 \u043f\u043e\u0434\u043e\u0431\u0440\u0430\u043d\u0430 \u0441\u043f\u0435\u0446\u0438\u0430\u043b\u044c\u043d\u043e     kernel&lt;&lt;&lt;1, 1, shared_memory_size, stream&gt;&gt;&gt;();     err = cudaStreamSynchronize(stream);     if (err != cudaSuccess) {         fprintf(stderr, \"Failed to synchronize stream: %s\\n\", cudaGetErrorString(err));         cudaStreamDestroy(stream);         return -1;     }     cudaStreamDestroy(stream);     return 0; }<\/code><\/pre>\n<p>\u0418\u043c\u0435\u043d\u043d\u043e \u043d\u0430 \u044d\u0442\u043e\u043c, \u043a\u0430\u043a \u043f\u0440\u0430\u0432\u0438\u043b\u043e, \u0437\u0430\u043a\u0430\u043d\u0447\u0438\u0432\u0430\u044e\u0442\u0441\u044f \u0442\u0443\u0442\u043e\u0440\u0438\u0430\u043b\u044b \u2014 \u0438 \u043d\u0435\u0437\u0430\u043c\u0435\u0442\u043d\u043e \u043d\u0430\u0447\u0438\u043d\u0430\u044e\u0442\u0441\u044f \u043f\u0440\u043e\u0431\u043b\u0435\u043c\u044b. \u041f\u043e\u043f\u0440\u043e\u0431\u0443\u0439\u0442\u0435 \u0437\u0430\u043f\u0443\u0441\u0442\u0438\u0442\u044c \u0432\u043e\u0442 \u044d\u0442\u043e:<\/p>\n<pre><code>$ nvcc -o hello_world hello_world.cu &amp;&amp; .\/hello_world<\/code><\/pre>\n<p>\u041d\u0438\u043a\u0430\u043a\u043e\u0433\u043e \u0432\u044b\u0432\u043e\u0434\u0430, \u043d\u0438\u043a\u0430\u043a\u043e\u0433\u043e \u0441\u043e\u043e\u0431\u0449\u0435\u043d\u0438\u044f \u043e\u0431 \u043e\u0448\u0438\u0431\u043a\u0435, \u0432\u043e\u043e\u0431\u0449\u0435 \u043d\u0438\u0447\u0435\u0433\u043e. \u041c\u044b \u043d\u0435 \u043c\u043e\u0436\u0435\u043c \u0432\u044b\u0442\u044f\u043d\u0443\u0442\u044c \u043e\u0448\u0438\u0431\u043a\u0443 \u0438\u0437 \u043f\u043e\u0442\u043e\u043a\u0430, \u043f\u043e\u0441\u043a\u043e\u043b\u044c\u043a\u0443 \u043e\u0442\u043a\u0430\u0437 \u043f\u0440\u043e\u0438\u0441\u0445\u043e\u0434\u0438\u0442 \u043f\u0440\u0438 \u043e\u0442\u043f\u0440\u0430\u0432\u043a\u0435 \u0434\u0430\u043d\u043d\u044b\u0445, \u0430 \u043d\u0435 \u043f\u0440\u0438 \u0432\u044b\u043f\u043e\u043b\u043d\u0435\u043d\u0438\u0438.<\/p>\n<h2>API \u0441\u0440\u0435\u0434\u044b \u0432\u044b\u043f\u043e\u043b\u043d\u0435\u043d\u0438\u044f CUDA<\/h2>\n<p>\u0422\u0440\u043e\u0439\u043d\u044b\u0435 \u0443\u0433\u043b\u043e\u0432\u044b\u0435 \u0441\u043a\u043e\u0431\u043a\u0438 NVCC \u2014 \u044d\u0442\u043e \u0441\u0438\u043d\u0442\u0430\u043a\u0441\u0438\u0447\u0435\u0441\u043a\u0438\u0439 \u0441\u0430\u0445\u0430\u0440, \u043a\u043e\u0442\u043e\u0440\u044b\u043c \u043f\u0440\u0438\u0441\u044b\u043f\u0430\u043d API \u0441\u0440\u0435\u0434\u044b \u0432\u044b\u043f\u043e\u043b\u043d\u0435\u043d\u0438\u044f CUDA, \u0430 \u044d\u0442\u043e\u0442 \u0438\u043d\u0442\u0435\u0440\u0444\u0435\u0439\u0441, \u0432 \u0441\u0432\u043e\u044e \u043e\u0447\u0435\u0440\u0435\u0434\u044c, \u043e\u0431\u0451\u0440\u0442\u044b\u0432\u0430\u0435\u0442 \u0431\u043e\u043b\u0435\u0435 \u043d\u0438\u0437\u043a\u043e\u0443\u0440\u043e\u0432\u043d\u0435\u0432\u044b\u0439 API \u0434\u0440\u0430\u0439\u0432\u0435\u0440\u043e\u0432 CUDA. \u0414\u043b\u044f \u044d\u0444\u0444\u0435\u043a\u0442\u0438\u0432\u043d\u043e\u0433\u043e \u043e\u0442\u043b\u043e\u0432\u0430 \u043e\u0448\u0438\u0431\u043e\u043a \u043d\u0435\u043e\u0431\u0445\u043e\u0434\u0438\u043c\u043e \u0437\u0430\u0434\u0435\u0439\u0441\u0442\u0432\u043e\u0432\u0430\u0442\u044c <a href=\"https:\/\/docs.nvidia.com\/cuda\/cuda-runtime-api\/group__CUDART__EXECUTION.html\" rel=\"noopener noreferrer nofollow\">CUDA Driver Execution Control API<\/a>:<\/p>\n<pre><code class=\"cpp\">#include &lt;cuda_runtime.h&gt; #include &lt;stdio.h&gt;  __global__ void kernel() {     extern __shared__ char shared_buffer[];     printf(\"Hello World from block %d, thread %d\\n\", blockIdx.x, threadIdx.x); }  int main() {     cudaStream_t stream;     cudaError_t err;     err = cudaStreamCreate(&amp;stream);     if (err != cudaSuccess) {         fprintf(stderr, \"Failed to create stream: %s\\n\", cudaGetErrorString(err));         return -1;     }      dim3 grid(1);     dim3 block(1);     size_t shared_memory_size = 1 &lt;&lt; 30; \/\/ 1 \u0413\u0411     void *kernel_args[] = {};     err = cudaLaunchKernel((void *)kernel, grid, block, kernel_args, shared_memory_size, stream);     if (err != cudaSuccess) {         fprintf(stderr, \"Failed to launch kernel: %s\\n\", cudaGetErrorString(err));         cudaStreamDestroy(stream);         return -1;     }      err = cudaStreamSynchronize(stream);     if (err != cudaSuccess) {         fprintf(stderr, \"Kernel execution failed: %s\\n\", cudaGetErrorString(err));         cudaStreamDestroy(stream);         return -1;     }     err = cudaStreamDestroy(stream);     if (err != cudaSuccess) {         fprintf(stderr, \"Failed to destroy stream: %s\\n\", cudaGetErrorString(err));         return -1;     }     return 0; }<\/code><\/pre>\n<p>\u0421\u043a\u043e\u043c\u043f\u0438\u043b\u0438\u0440\u0443\u0435\u043c \u0438 \u0432\u044b\u043f\u043e\u043b\u043d\u0438\u043c:<\/p>\n<pre><code>$ nvcc -o hello_world hello_world.cu &amp;&amp; .\/hello_world &gt; Failed to launch kernel: invalid argument<\/code><\/pre>\n<p>\u042d\u0442\u0430 \u043e\u0448\u0438\u0431\u043a\u0430 \u043e\u0436\u0438\u0434\u0430\u0435\u043c\u0430, \u043f\u043e\u0441\u043a\u043e\u043b\u044c\u043a\u0443 \u044f\u0434\u0440\u043e \u043d\u0435\u0432\u043e\u0437\u043c\u043e\u0436\u043d\u043e \u0434\u0430\u0436\u0435 \u043e\u0442\u043f\u0440\u0430\u0432\u0438\u0442\u044c \u2014 \u0438\u0437-\u0437\u0430 \u0430\u0431\u0441\u0443\u0440\u0434\u043d\u043e\u0433\u043e \u0437\u0430\u043f\u0440\u043e\u0441\u0430 \u043a \u043f\u0430\u043c\u044f\u0442\u0438. \u041d\u043e \u043f\u0440\u043e\u0431\u043b\u0435\u043c\u0430 \u0441 \u044d\u0442\u0438\u043c API \u0432 \u0442\u043e\u043c, \u0447\u0442\u043e \u043d\u0430\u043c \u0442\u0440\u0435\u0431\u0443\u0435\u0442\u0441\u044f \u043d\u0430\u0439\u0442\u0438 \u0438\u043d\u043e\u0439 \u0441\u043f\u043e\u0441\u043e\u0431, \u0447\u0442\u043e\u0431\u044b \u043f\u0435\u0440\u0435\u0434\u0430\u0432\u0430\u0442\u044c \u0430\u0440\u0433\u0443\u043c\u0435\u043d\u0442\u044b \u044f\u0434\u0440\u0443. \u0412\u043e\u0442 \u043a\u0430\u043a \u043c\u044b \u043f\u0435\u0440\u0435\u0434\u0430\u0432\u0430\u043b\u0438 \u0431\u044b \u044f\u0434\u0440\u0443 \u043c\u0430\u0441\u0441\u0438\u0432\u044b \u0438 \u0441\u043a\u0430\u043b\u044f\u0440\u044b:<\/p>\n<pre><code class=\"cpp\">#include &lt;cuda_runtime.h&gt; #include &lt;stdio.h&gt;  __global__ void kernel(float *amount, size_t count, int power) {     size_t idx = blockIdx.x * blockDim.x + threadIdx.x;     if (idx &gt; count) return;     amount[idx] = amount[idx] * scalbln(1.0, power); \/\/ \u041f\u0440\u0438\u043c\u0435\u0440 \u0432\u0441\u0442\u0440\u043e\u0435\u043d\u043d\u043e\u0439 \u0444\u0443\u043d\u043a\u0446\u0438\u0438 CUDA ;) }  int main() {     cudaError_t err;     size_t num_elements = 1024;     int integral_power = -2;     double *data;      \/\/ \u0412\u044b\u0434\u0435\u043b\u044f\u0435\u043c \u043e\u0431\u044a\u0435\u0434\u0438\u043d\u0451\u043d\u043d\u0443\u044e \u043f\u0430\u043c\u044f\u0442\u044c     err = cudaMallocManaged(&amp;data, num_elements * sizeof(double));     if (err != cudaSuccess) {         fprintf(stderr, \"cudaMallocManaged failed: %s\\n\", cudaGetErrorString(err));         return -1;     }      \/\/ \u0418\u043d\u0438\u0446\u0438\u0430\u043b\u0438\u0437\u0438\u0440\u0443\u0435\u043c \u0434\u0430\u043d\u043d\u044b\u0435     for (size_t i = 0; i &lt; num_elements; ++i) data[i] = (double)i;      \/\/ \u0421\u043e\u0437\u0434\u0430\u0451\u043c \u043f\u043e\u0442\u043e\u043a CUDA      cudaStream_t stream;     err = cudaStreamCreate(&amp;stream);     if (err != cudaSuccess) {         fprintf(stderr, \"Failed to create stream: %s\\n\", cudaGetErrorString(err));         cudaFree(data);         return -1;     }      \/\/ \u041e\u043f\u0440\u0435\u0434\u0435\u043b\u044f\u0435\u043c \u043f\u0430\u0440\u0430\u043c\u0435\u0442\u0440\u044b \u043f\u0443\u0441\u043a\u0430 \u044f\u0434\u0440\u0430     dim3 grid((num_elements + 255) \/ 256);     dim3 block(256);     void *kernel_args[] = {         (void *)&amp;data,         (void *)&amp;num_elements,         (void *)&amp;integral_power,     };      \/\/ \u0417\u0430\u043f\u0443\u0441\u043a\u0430\u0435\u043c \u044f\u0434\u0440\u043e     err = cudaLaunchKernel((void *)kernel, grid, block, kernel_args, 0, stream);     if (err != cudaSuccess) {         fprintf(stderr, \"Failed to launch kernel: %s\\n\", cudaGetErrorString(err));         cudaStreamDestroy(stream);         cudaFree(data);         return -1;     }      \/\/ \u0421\u0438\u043d\u0445\u0440\u043e\u043d\u0438\u0437\u0438\u0440\u0443\u0435\u043c \u043f\u043e\u0442\u043e\u043a     err = cudaStreamSynchronize(stream);     if (err != cudaSuccess) {         fprintf(stderr, \"Kernel execution failed: %s\\n\", cudaGetErrorString(err));         cudaStreamDestroy(stream);         cudaFree(data);         return -1;     }      \/\/ \u0412\u044b\u0432\u043e\u0434\u0438\u043c \u0440\u0435\u0437\u0443\u043b\u044c\u0442\u0430\u0442\u044b \u043d\u0430 \u044d\u043a\u0440\u0430\u043d     for (size_t i = 0; i &lt; 5; ++i) printf(\"data[%zu] = %f\\n\", i, data[i]);     cudaStreamDestroy(stream);     cudaFree(data);     return 0; }<\/code><\/pre>\n<p>\u042f \u0432\u043e\u0441\u043f\u043e\u043b\u044c\u0437\u043e\u0432\u0430\u043b\u0441\u044f \u043e\u0431\u044a\u0435\u0434\u0438\u043d\u0451\u043d\u043d\u043e\u0439 \u043f\u0430\u043c\u044f\u0442\u044c\u044e, \u0447\u0442\u043e\u0431\u044b \u0443\u043f\u0440\u043e\u0441\u0442\u0438\u0442\u044c \u043f\u0440\u0438\u043c\u0435\u0440. \u041c\u044b \u043d\u0435 \u043e\u0431\u044f\u0437\u0430\u043d\u044b \u044f\u0432\u043d\u043e \u0432\u044b\u0434\u0435\u043b\u044f\u0442\u044c 2 \u0431\u0443\u0444\u0435\u0440\u0430 \u043e\u0434\u043d\u043e\u0432\u0440\u0435\u043c\u0435\u043d\u043d\u043e \u0432 \u0426\u041f \u0438 \u0433\u0440\u0430\u0444\u0438\u0447\u0435\u0441\u043a\u043e\u043c \u043f\u0440\u043e\u0446\u0435\u0441\u0441\u043e\u0440\u0435 \u0438 \u043a\u043e\u043f\u0438\u0440\u043e\u0432\u0430\u0442\u044c \u0434\u0430\u043d\u043d\u044b\u0435 \u043c\u0435\u0436\u0434\u0443 \u043d\u0438\u043c\u0438. \u0414\u0440\u0430\u0439\u0432\u0435\u0440 \u0434\u0435\u0440\u0436\u0438\u0442 \u043a\u043e\u043f\u0438\u0438 \u0434\u0430\u043d\u043d\u044b\u0445 \u043e\u0434\u043d\u043e\u0432\u0440\u0435\u043c\u0435\u043d\u043d\u043e \u043d\u0430 \u0445\u043e\u0441\u0442\u0435 \u0438 \u0432 \u043f\u0430\u043c\u044f\u0442\u0438 \u0443\u0441\u0442\u0440\u043e\u0439\u0441\u0442\u0432\u0430, \u0438 \u043f\u043e \u043c\u0435\u0440\u0435 \u043d\u0435\u043e\u0431\u0445\u043e\u0434\u0438\u043c\u043e\u0441\u0442\u0438 \u0430\u0432\u0442\u043e\u043c\u0430\u0442\u0438\u0447\u0435\u0441\u043a\u0438 \u043f\u0435\u0440\u0435\u0434\u0430\u0451\u0442 \u043e\u0431\u043d\u043e\u0432\u043b\u0435\u043d\u0438\u044f \u043c\u0435\u0436\u0434\u0443 \u043d\u0438\u043c\u0438.<\/p>\n<h2>\u041a\u043e\u043e\u043f\u0435\u0440\u0430\u0442\u0438\u0432\u043d\u044b\u0435 \u0433\u0440\u0443\u043f\u043f\u044b<\/h2>\n<p>\u0410\u0445, \u043a\u0430\u043a \u0431\u044b \u0445\u043e\u0442\u0435\u043b\u043e\u0441\u044c, \u0447\u0442\u043e\u0431\u044b \u043c\u043e\u0436\u043d\u043e \u0431\u044b\u043b\u043e \u043d\u0430\u043f\u0438\u0441\u0430\u0442\u044c \u043f\u0430\u0440\u0430\u043b\u043b\u0435\u043b\u044c\u043d\u044b\u0435 \u0430\u043b\u0433\u043e\u0440\u0438\u0442\u043c\u044b \u0434\u043b\u044f \u0433\u0440\u0430\u0444\u0438\u0447\u0435\u0441\u043a\u043e\u0433\u043e \u043f\u0440\u043e\u0446\u0435\u0441\u0441\u043e\u0440\u0430 \u0435\u0434\u0438\u043d\u043e\u0436\u0434\u044b \u2014 \u0441\u043e\u0431\u0440\u0430\u0442\u044c \u0441\u0442\u043e\u043f\u043a\u0443 \u0430\u0431\u0441\u0442\u0440\u0430\u043a\u0446\u0438\u0439, \u043e\u0431\u0435\u0440\u043d\u0443\u0442\u044c \u0438\u0445 \u0432 \u0448\u0430\u0431\u043b\u043e\u043d\u044b, \u0430 \u0434\u0430\u043b\u044c\u0448\u0435 \u0441\u0440\u0435\u0434\u0430 \u0432\u044b\u043f\u043e\u043b\u043d\u0435\u043d\u0438\u044f \u043f\u0443\u0441\u0442\u044c \u0441\u0430\u043c\u0430 \u0440\u0430\u0437\u0431\u0438\u0440\u0430\u0435\u0442\u0441\u044f. \u041d\u043e \u043d\u0430 \u043f\u0440\u0430\u043a\u0442\u0438\u043a\u0435 \u0442\u0430\u043a\u043e\u0435 \u0443\u0434\u0430\u0451\u0442\u0441\u044f \u0440\u0435\u0434\u043a\u043e. \u041a \u0441\u043e\u0436\u0430\u043b\u0435\u043d\u0438\u044e, \u0438 API CUDA \u0434\u043b\u044f \u043e\u0431\u0440\u0430\u0449\u0435\u043d\u0438\u044f \u0441 \u043a\u043e\u043e\u043f\u0435\u0440\u0430\u0442\u0438\u0432\u043d\u044b\u043c\u0438 \u0433\u0440\u0443\u043f\u043f\u0430\u043c\u0438 \u2014 \u043d\u0435 \u0438\u0441\u043a\u043b\u044e\u0447\u0435\u043d\u0438\u0435.<\/p>\n<p>\u041e\u043d \u043f\u0440\u043e\u0435\u043a\u0442\u0438\u0440\u043e\u0432\u0430\u043b\u0441\u044f \u043a\u0430\u043a \u0435\u0434\u0438\u043d\u043e\u043e\u0431\u0440\u0430\u0437\u043d\u0430\u044f \u0430\u0431\u0441\u0442\u0440\u0430\u043a\u0446\u0438\u044f \u0434\u043b\u044f \u043a\u043e\u043e\u0440\u0434\u0438\u043d\u0430\u0446\u0438\u0438 \u043f\u043e\u0442\u043e\u043a\u043e\u0432, \u043d\u0435 \u043e\u0433\u0440\u0430\u043d\u0438\u0447\u0435\u043d\u043d\u0430\u044f \u043f\u0440\u0435\u0434\u0435\u043b\u0430\u043c\u0438 \u043e\u0434\u043d\u043e\u0433\u043e \u0431\u043b\u043e\u043a\u0430. \u041f\u0440\u0438 \u044d\u0442\u043e\u043c \u0432\u0441\u0442\u0440\u043e\u0435\u043d\u043d\u044b\u0435 \u0444\u0443\u043d\u043a\u0446\u0438\u0438 C++ \u043f\u0440\u0438\u043c\u0435\u043d\u044f\u044e\u0442\u0441\u044f \u0434\u043b\u044f \u043d\u0430\u0437\u043d\u0430\u0447\u0435\u043d\u0438\u044f \u0441\u043b\u043e\u0436\u043d\u044b\u0445 \u0430\u043b\u0433\u043e\u0440\u0438\u0442\u043c\u043e\u0432 GPU, \u0438 \u0437\u0430\u0434\u0435\u0439\u0441\u0442\u0432\u0443\u0435\u043c\u0430\u044f \u0441 \u043d\u0438\u043c\u0438 \u0441\u0435\u043c\u0430\u043d\u0442\u0438\u043a\u0430 \u0441\u0438\u043d\u0445\u0440\u043e\u043d\u0438\u0437\u0430\u0446\u0438\u0438 \u0434\u043e\u0441\u0442\u0430\u0442\u043e\u0447\u043d\u043e \u0433\u0438\u0431\u043a\u0430\u044f. \u0422\u0435\u043e\u0440\u0435\u0442\u0438\u0447\u0435\u0441\u043a\u0438, \u0442\u0430\u043a \u0434\u043e\u043b\u0436\u043d\u0430 \u0440\u0435\u0448\u0430\u0442\u044c\u0441\u044f \u0441\u0435\u0440\u044c\u0451\u0437\u043d\u0430\u044f \u043f\u0440\u043e\u0431\u043b\u0435\u043c\u0430: \u043f\u0443\u0441\u0442\u044c \u0432\u0441\u0435 \u043f\u043e\u0442\u043e\u043a\u0438 \u043d\u0430 \u0443\u0441\u0442\u0440\u043e\u0439\u0441\u0442\u0432\u0435 \u0432\u0440\u0435\u043c\u044f \u043e\u0442 \u0432\u0440\u0435\u043c\u0435\u043d\u0438 \u0441\u0438\u043d\u0445\u0440\u043e\u043d\u0438\u0437\u0438\u0440\u0443\u044e\u0442\u0441\u044f \u043f\u0435\u0440\u0435\u0434 \u043f\u0440\u043e\u0434\u043e\u043b\u0436\u0435\u043d\u0438\u0435\u043c \u0440\u0430\u0431\u043e\u0442\u044b \u2014 \u044d\u0442\u043e \u0447\u0440\u0435\u0437\u0432\u044b\u0447\u0430\u0439\u043d\u043e \u0432\u0430\u0436\u043d\u043e \u0434\u043b\u044f \u0438\u0442\u0435\u0440\u0430\u0442\u0438\u0432\u043d\u044b\u0445 \u0430\u043b\u0433\u043e\u0440\u0438\u0442\u043c\u043e\u0432, \u043d\u0430 \u043a\u043e\u0442\u043e\u0440\u044b\u0445 \u043e\u0441\u043d\u043e\u0432\u0430\u043d\u044b, \u043d\u0430\u043f\u0440\u0438\u043c\u0435\u0440, \u0441\u0438\u043c\u0443\u043b\u044f\u0446\u0438\u0438 \u0444\u0438\u0437\u0438\u0447\u0435\u0441\u043a\u0438\u0445 \u043f\u0440\u043e\u0446\u0435\u0441\u0441\u043e\u0432 \u0438\u043b\u0438 \u0440\u0435\u0448\u0430\u043b\u043a\u0438.<\/p>\n<p>\u041f\u0440\u043e\u0441\u0442\u043e \u043d\u0430\u043f\u043e\u043c\u043d\u044e:<\/p>\n<ul>\n<li>\n<p> <code>__syncwarp()<\/code>\u00a0\u043e\u0431\u0451\u0440\u0442\u044b\u0432\u0430\u0435\u0442 32 \u043f\u043e\u0442\u043e\u043a\u0430.<\/p>\n<\/li>\n<li>\n<p> <code>__syncthreads()\u00a0<\/code>\u043e\u0431\u0451\u0440\u0442\u044b\u0432\u0430\u0435\u0442 \u043b\u043e\u0433\u0438\u0447\u0435\u0441\u043a\u0438\u0439 \u0431\u043b\u043e\u043a, \u0432\u043a\u043b\u044e\u0447\u0430\u044e\u0449\u0438\u0439 1-1024 \u043f\u043e\u0442\u043e\u043a\u043e\u0432.<\/p>\n<\/li>\n<li>\n<p>\u041d\u0430 \u0432\u0441\u0435 \u043f\u0440\u043e\u0447\u0438\u0435 \u0441\u043b\u0443\u0447\u0430\u0438 \u0435\u0441\u0442\u044c\u00a0<s>Mastercard<\/s>\u043a\u043e\u043e\u043f\u0435\u0440\u0430\u0442\u0438\u0432\u043d\u044b\u0435 \u0433\u0440\u0443\u043f\u043f\u044b.<\/p>\n<\/li>\n<\/ul>\n<p>\u0421\u0438\u043b\u044c\u043d\u0435\u0435 \u0432\u0441\u0435\u0433\u043e \u043d\u0430\u043f\u0440\u0430\u0448\u0438\u0432\u0430\u0435\u0442\u0441\u044f \u0442\u0430\u043a\u043e\u0439 \u043f\u0440\u0438\u043c\u0435\u0440: \u0441\u0438\u043d\u0445\u0440\u043e\u043d\u0438\u0437\u0438\u0440\u0443\u0435\u043c \u0432\u0435\u0441\u044c \u0433\u0440\u0438\u0434 \u0432 \u0432\u0438\u0434\u0435 \u043c\u043d\u043e\u0433\u043e\u044d\u0442\u0430\u043f\u043d\u044b\u0445 \u0438\u0442\u0435\u0440\u0430\u0442\u0438\u0432\u043d\u044b\u0445 \u0430\u043b\u0433\u043e\u0440\u0438\u0442\u043c\u043e\u0432, \u0441\u043a\u0430\u0436\u0435\u043c, \u043f\u0440\u0438 \u0441\u0438\u043c\u0443\u043b\u044f\u0446\u0438\u0438 \u0444\u0438\u0437\u0438\u0447\u0435\u0441\u043a\u043e\u0433\u043e \u043f\u0440\u043e\u0446\u0435\u0441\u0441\u0430. \u0414\u043b\u044f \u044d\u0442\u043e\u0433\u043e Nvidia \u0440\u0435\u043a\u043e\u043c\u0435\u043d\u0434\u0443\u0435\u0442 \u0437\u0430\u0434\u0435\u0439\u0441\u0442\u0432\u043e\u0432\u0430\u0442\u044c \u043d\u043e\u0432\u0443\u044e \u0444\u0443\u043d\u043a\u0446\u0438\u044e <code>cooperative_groups::sync()<\/code>:<\/p>\n<pre><code class=\"cpp\">#include &lt;cuda_runtime.h&gt; #include &lt;cooperative_groups.h&gt; #include &lt;stdio.h&gt; #include &lt;math.h&gt;  namespace cg = cooperative_groups;  __device__ float3 compute_force(float3 position_first, float3 position_second) {     float3 r;     r.x = position_second.x - position_first.x;     r.y = position_second.y - position_first.y;     r.z = position_second.z - position_first.z;      float squared_distance = r.x * r.x + r.y * r.y + r.z * r.z + 1e-6f; \/\/ avoid div by zero     float reciprocal_distance = rsqrtf(squared_distance);     float reciprocal_cube = reciprocal_distance * reciprocal_distance * reciprocal_distance;      constexpr float gravitational_constant = 1.0f;     float scale = gravitational_constant * reciprocal_cube;     r.x *= scale;     r.y *= scale;     r.z *= scale;     return r; }  __global__ void cooperative_kernel(     float3 *positions_old, float3 *positions_new,     float3 *velocities_old, float3 *velocities_new,     size_t count, size_t iterations, float dt) {      cg::grid_group grid = cg::this_grid();     size_t idx = blockIdx.x * blockDim.x + threadIdx.x;     if (idx &gt;= count) return;      for (size_t iter = 0; iter &lt; iterations; ++iter) {         float3 force = {0.0f, 0.0f, 0.0f};          \/\/ \u0410\u043a\u043a\u0443\u043c\u0443\u043b\u0438\u0440\u0443\u0435\u043c \u0441\u0438\u043b\u044b \u043e\u0442 \u0432\u0441\u0435\u0445 \u043f\u0440\u043e\u0447\u0438\u0445 \u0447\u0430\u0441\u0442\u0438\u0446         for (size_t j = 0; j &lt; count; ++j) {             if (j == idx) continue;             float3 f = compute_force(positions_old[idx], positions_old[j]);             force.x += f.x;             force.y += f.y;             force.z += f.z;         }          \/\/ \u041e\u0431\u043d\u043e\u0432\u043b\u044f\u0435\u043c \u0437\u043d\u0430\u0447\u0435\u043d\u0438\u044f \u0441\u043a\u043e\u0440\u043e\u0441\u0442\u0438 \u0438 \u043f\u043e\u043b\u043e\u0436\u0435\u043d\u0438\u044f         velocities_new[idx].x = velocities_old[idx].x + force.x * dt;         velocities_new[idx].y = velocities_old[idx].y + force.y * dt;         velocities_new[idx].z = velocities_old[idx].z + force.z * dt;         positions_new[idx].x = positions_old[idx].x + velocities_new[idx].x * dt;         positions_new[idx].y = positions_old[idx].y + velocities_new[idx].y * dt;         positions_new[idx].z = positions_old[idx].z + velocities_new[idx].z * dt;          \/\/ \u041c\u0435\u043d\u044f\u0435\u043c \u0431\u0443\u0444\u0435\u0440\u044b, \u0433\u043e\u0442\u043e\u0432\u044f\u0441\u044c \u043a \u0441\u043b\u0435\u0434\u0443\u044e\u0449\u0435\u0439 \u0438\u0442\u0435\u0440\u0430\u0446\u0438\u0438         grid.sync();         float3 *temp_pos = positions_old, *temp_vel = velocities_old;         positions_old = positions_new, positions_new = temp_pos;         velocities_old = velocities_new, velocities_new = temp_vel;         grid.sync();     } }  int main() {     cudaError_t err;     size_t num_particles = 256;     size_t iterations = 10;     float dt = 0.01f;     dim3 block;     dim3 grid;     void *kernel_args[7];     float3 *positions_old = nullptr, *positions_new = nullptr;     float3 *velocities_old = nullptr, *velocities_new = nullptr;      \/\/ \u0412\u044b\u0434\u0435\u043b\u044f\u0435\u043c \u043f\u0430\u043c\u044f\u0442\u044c     err = cudaMallocManaged(&amp;positions_old, num_particles * sizeof(float3));     if (err != cudaSuccess) goto cleanup;     err = cudaMallocManaged(&amp;positions_new, num_particles * sizeof(float3));     if (err != cudaSuccess) goto cleanup;     err = cudaMallocManaged(&amp;velocities_old, num_particles * sizeof(float3));     if (err != cudaSuccess) goto cleanup;     err = cudaMallocManaged(&amp;velocities_new, num_particles * sizeof(float3));     if (err != cudaSuccess) goto cleanup;      \/\/ \u0418\u043d\u0438\u0446\u0438\u0430\u043b\u0438\u0437\u0438\u0440\u0443\u0435\u043c \u0437\u043d\u0430\u0447\u0435\u043d\u0438\u044f \u043f\u043e\u043b\u043e\u0436\u0435\u043d\u0438\u044f \u0438 \u0441\u043a\u043e\u0440\u043e\u0441\u0442\u0438      for (size_t i = 0; i &lt; num_particles; ++i) {         float theta = (float)i * 0.01f;         float phi = (float)i * 0.005f;         float radius = 10.0f + (i % 32) * 0.1f;         positions_old[i] = {radius * cosf(theta) * sinf(phi), radius * sinf(theta) * sinf(phi), radius * cosf(phi)};         velocities_old[i] = {0.01f * sinf(phi), 0.01f * cosf(theta), 0.01f * sinf(theta + phi)};     }      \/\/ \u0423\u0431\u0435\u0434\u0438\u043c\u0441\u044f, \u0447\u0442\u043e \u043d\u0430 \u0434\u0430\u043d\u043d\u043e\u043c \u0443\u0441\u0442\u0440\u043e\u0439\u0441\u0442\u0432\u0435 \u043f\u043e\u0434\u0434\u0435\u0440\u0436\u0438\u0432\u0430\u0435\u0442\u0441\u044f \u043a\u043e\u043e\u043f\u0435\u0440\u0430\u0442\u0438\u0432\u043d\u044b\u0439 \u0437\u0430\u043f\u0443\u0441\u043a \u043f\u043e\u0442\u043e\u043a\u043e\u0432     cudaDeviceProp props;     cudaGetDeviceProperties(&amp;props, 0);     if (!props.cooperativeLaunch) {         fprintf(stderr, \"Cooperative launch not supported on this device.\\n\");         err = cudaErrorNotSupported;         goto cleanup;     }     cudaStream_t stream;     err = cudaStreamCreate(&amp;stream);     if (err != cudaSuccess) {         fprintf(stderr, \"Failed to create stream: %s\\n\", cudaGetErrorString(err));         goto cleanup;     }      block = dim3(256);     grid = dim3((num_particles + block.x - 1) \/ block.x);     kernel_args[0] = &amp;positions_old;     kernel_args[1] = &amp;positions_new;     kernel_args[2] = &amp;velocities_old;     kernel_args[3] = &amp;velocities_new;     kernel_args[4] = &amp;num_particles;     kernel_args[5] = &amp;iterations;     kernel_args[6] = &amp;dt;      \/\/ \u0417\u0430\u043f\u0443\u0441\u043a\u0430\u0435\u043c \u044f\u0434\u0440\u043e     err = cudaLaunchCooperativeKernel((void *)cooperative_kernel, grid, block, kernel_args, 0, stream);     if (err != cudaSuccess) {         fprintf(stderr, \"Failed to launch cooperative kernel: %s\\n\", cudaGetErrorString(err));         goto cleanup;     }     err = cudaStreamSynchronize(stream);     if (err != cudaSuccess) {         fprintf(stderr, \"Kernel execution failed: %s\\n\", cudaGetErrorString(err));         goto cleanup;     }      \/\/ \u0412\u044b\u0432\u043e\u0434\u0438\u043c \u043d\u0430 \u044d\u043a\u0440\u0430\u043d \u043e\u043a\u043e\u043d\u0447\u0430\u0442\u0435\u043b\u044c\u043d\u044b\u0435 \u043f\u043e\u043b\u043e\u0436\u0435\u043d\u0438\u044f     for (size_t i = 0; i &lt; num_particles; ++i)         printf(\"Final position[%zu] = (%f, %f, %f)\\n\", i, positions_old[i].x, positions_old[i].y, positions_old[i].z);  cleanup:     if (positions_old) cudaFree(positions_old);     if (positions_new) cudaFree(positions_new);     if (velocities_old) cudaFree(velocities_old);     if (velocities_new) cudaFree(velocities_new);     return (err == cudaSuccess) ? 0 : -1; }<\/code><\/pre>\n<p>\u041e\u0431\u0440\u0430\u0442\u0438\u0442\u0435 \u0432\u043d\u0438\u043c\u0430\u043d\u0438\u0435, \u043a\u0430\u043a \u044f \u0437\u0430\u043c\u0435\u043d\u0438\u043b\u00a0<code>cudaLaunchKernel<\/code>\u00a0\u043d\u0430\u00a0<code>cudaLaunchCooperativeKernel<\/code> \u0438 \u0434\u043e\u0431\u0430\u0432\u0438\u043b \u043a \u044f\u0434\u0440\u0443 \u043e\u0431\u044a\u0435\u043a\u0442\u00a0<code>cg::grid_group<\/code>. \u0415\u0441\u043b\u0438 \u0431\u044b \u043c\u044b \u0432\u043e\u0441\u043f\u043e\u043b\u044c\u0437\u043e\u0432\u0430\u043b\u0438\u0441\u044c \u0441\u0442\u0430\u0440\u044b\u043c \u043d\u0435\u043a\u043e\u043e\u043f\u0435\u0440\u0430\u0442\u0438\u0432\u043d\u044b\u043c \u043f\u0443\u0441\u043a\u043e\u0432\u044b\u043c API\u00a0<code>cudaLaunchKernel<\/code>, \u0442\u043e \u0443 \u043d\u0430\u0441 \u0431\u044b \u043f\u043e\u043b\u0443\u0447\u0438\u043b\u043e\u0441\u044c: <\/p>\n<pre><code>$ nvcc -o hello_world hello_world.cu &amp;&amp; .\/hello_world &gt; Kernel execution failed: unspecified launch failure<\/code><\/pre>\n<p>\u0422\u0430\u043a\u0438\u043c \u043e\u0431\u0440\u0430\u0437\u043e\u043c, \u0441\u043b\u0435\u0434\u0443\u0435\u0442 \u0440\u0430\u0431\u043e\u0442\u0430\u0442\u044c \u0441 \u043d\u043e\u0432\u044b\u043c \u00ab\u043a\u043e\u043e\u043f\u0435\u0440\u0430\u0442\u0438\u0432\u043d\u044b\u043c\u00bb API, \u043a\u043e\u0442\u043e\u0440\u044b\u0439 \u0441\u0440\u0430\u0437\u0443 \u0432\u044b\u0433\u043b\u044f\u0434\u0435\u043b \u043c\u043d\u043e\u0433\u043e\u043e\u0431\u0435\u0449\u0430\u044e\u0449\u0435. \u041f\u0440\u0438 \u0443\u0441\u043a\u043e\u0440\u0435\u043d\u043d\u044b\u0445 \u0441\u0432\u044f\u0437\u044f\u0445 GPU-GPU \u0438 \u0432\u043d\u0443\u0442\u0440\u0438\u0443\u0437\u043b\u043e\u0432\u044b\u0445 \u043f\u0435\u0440\u0435\u043a\u043b\u044e\u0447\u0430\u0442\u0435\u043b\u044f\u0445 NVLink \u043c\u043e\u0436\u043d\u043e \u043d\u0430\u0434\u0435\u044f\u0442\u044c\u0441\u044f, \u0447\u0442\u043e \u0443 \u043d\u0430\u0441 \u0431\u0443\u0434\u0443\u0442 \u0431\u043e\u043b\u0435\u0435 \u043d\u0430\u0434\u0451\u0436\u043d\u044b\u0435 \u043f\u0440\u0438\u043c\u0438\u0442\u0438\u0432\u044b \u0441\u0438\u043d\u0445\u0440\u043e\u043d\u0438\u0437\u0430\u0446\u0438\u0438 \u0434\u043b\u044f \u0441\u0438\u0441\u0442\u0435\u043c \u0441\u043e \u043c\u043d\u043e\u0436\u0435\u0441\u0442\u0432\u043e\u043c \u00a0GPU. \u0412 \u043a\u0430\u043a\u043e\u0439-\u0442\u043e \u043c\u043e\u043c\u0435\u043d\u0442 \u043f\u043e\u043a\u0430\u0437\u0430\u043b\u043e\u0441\u044c, \u0447\u0442\u043e \u044d\u0442\u043e \u0431\u0443\u0434\u0443\u0449\u0435\u0435 \u043f\u043e\u0447\u0442\u0438 \u043d\u0430\u0441\u0442\u0443\u043f\u0438\u043b\u043e: \u0432 \u0441\u0440\u0435\u0434\u0435 \u0432\u044b\u043f\u043e\u043b\u043d\u0435\u043d\u0438\u044f CUDA \u0432\u0432\u0435\u043b\u0438 <code>cudaLaunchCooperativeKernelMultiDevice<\/code>\u00a0\u0438 \u0430\u0431\u0441\u0442\u0440\u0430\u043a\u0446\u0438\u044e <code>cg::multi_grid_group<\/code> \u2014 \u0442\u0435 \u0441\u0430\u043c\u044b\u0435 \u043d\u0435\u0434\u043e\u0441\u0442\u0430\u044e\u0449\u0438\u0435 \u0437\u0432\u0435\u043d\u044c\u044f, \u0431\u0435\u0437 \u043a\u043e\u0442\u043e\u0440\u044b\u0445 \u0431\u044b\u043b\u043e \u0441\u043b\u043e\u0436\u043d\u043e \u043a\u043e\u043e\u0440\u0434\u0438\u043d\u0438\u0440\u043e\u0432\u0430\u0442\u044c \u0440\u0430\u0431\u043e\u0442\u0443 \u044f\u0434\u0435\u0440 \u0432 \u043c\u0430\u0441\u0448\u0442\u0430\u0431\u0435 \u043c\u043d\u043e\u0436\u0435\u0441\u0442\u0432\u0430 GPU. \u041d\u043e \u0432 CUDA 11.3 \u043e\u0431\u0430 \u044d\u0442\u0438 \u043d\u043e\u0432\u043e\u0432\u0432\u0435\u0434\u0435\u043d\u0438\u044f \u0431\u044b\u043b\u0438 \u043f\u0440\u0438\u0437\u043d\u0430\u043d\u044b \u0443\u0441\u0442\u0430\u0440\u0435\u0432\u0448\u0438\u043c\u0438 \u0438 \u0442\u0430\u043a\u0438\u043c \u043e\u0431\u0440\u0430\u0437\u043e\u043c \u0441\u0442\u0430\u043b\u0438 \u043e\u0434\u043d\u0438\u043c\u0438 \u0438\u0437 \u0441\u0430\u043c\u044b\u0445 \u043c\u0438\u043c\u043e\u043b\u0451\u0442\u043d\u044b\u0445 API \u0432 \u0438\u0441\u0442\u043e\u0440\u0438\u0438 CUDA.<\/p>\n<p>\u041f\u0440\u0438 \u0432\u0441\u0435\u0439 \u043a\u043e\u043d\u0446\u0435\u043f\u0442\u0443\u0430\u043b\u044c\u043d\u043e\u0439 \u043f\u0440\u0438\u0432\u043b\u0435\u043a\u0430\u0442\u0435\u043b\u044c\u043d\u043e\u0441\u0442\u0438 \u043a\u043e\u043e\u043f\u0435\u0440\u0430\u0442\u0438\u0432\u043d\u044b\u0445 \u0433\u0440\u0443\u043f\u043f, \u043d\u0435 \u0434\u0443\u043c\u0430\u044e, \u0447\u0442\u043e \u043e\u043d\u0438 \u0441\u043c\u043e\u0433\u0443\u0442 \u0432\u043d\u044f\u0442\u043d\u043e \u043c\u0430\u0441\u0448\u0442\u0430\u0431\u0438\u0440\u043e\u0432\u0430\u0442\u044c\u0441\u044f. \u0420\u0430\u0431\u043e\u0442\u0430, \u043a\u043e\u0442\u043e\u0440\u043e\u0439 \u044f \u0437\u0430\u043d\u0438\u043c\u0430\u044e\u0441\u044c, \u0432 \u043e\u0441\u043d\u043e\u0432\u043d\u043e\u043c \u043f\u0440\u043e\u0438\u0441\u0445\u043e\u0434\u0438\u0442 \u043d\u0430 \u0443\u0440\u043e\u0432\u043d\u0435 \u043e\u0431\u0451\u0440\u0442\u043a\u0438 \u0441 \u0438\u0441\u043f\u043e\u043b\u044c\u0437\u043e\u0432\u0430\u043d\u0438\u0435\u043c <code>__syncwarp()<\/code>, \u0430 \u0435\u0441\u043b\u0438 \u043c\u043d\u0435 \u0442\u0440\u0435\u0431\u0443\u0435\u0442\u0441\u044f \u043f\u043e\u0434\u043d\u044f\u0442\u044c\u0441\u044f \u043d\u0430 \u0443\u0440\u043e\u0432\u0435\u043d\u044c \u0432\u044b\u0448\u0435, \u044f \u043f\u0440\u0435\u0434\u043f\u043e\u0447\u0438\u0442\u0430\u044e \u043f\u0440\u0438\u0431\u0435\u0433\u0430\u0442\u044c \u043a \u0432\u0441\u0442\u0440\u0430\u0438\u0432\u0430\u0435\u043c\u043e\u043c\u0443 <a href=\"https:\/\/en.wikipedia.org\/wiki\/Parallel_Thread_Execution\" rel=\"noopener noreferrer nofollow\">PTX-\u0430\u0441\u0441\u0435\u043c\u0431\u043b\u0435\u0440\u0443<\/a>. \u041d\u0435 \u0441\u043e\u0441\u0442\u0430\u0432\u043b\u044f\u0435\u0442 \u0442\u0440\u0443\u0434\u0430 \u043f\u0440\u043e\u0432\u0435\u0440\u0438\u0442\u044c, \u0432\u043e \u0447\u0442\u043e \u043a\u043e\u043c\u043f\u0438\u043b\u0438\u0440\u0443\u0435\u0442\u0441\u044f \u0444\u0443\u043d\u043a\u0446\u0438\u044f <code>cooperative_groups::sync()<\/code>: \u043f\u0440\u043e\u0441\u0442\u043e \u0438\u0441\u043f\u043e\u043b\u044c\u0437\u0443\u0435\u043c NVCC \u0441 \u0444\u043b\u0430\u0433\u043e\u043c\u00a0<code>-ptx<\/code>:<\/p>\n<pre><code>$ nvcc -arch=sm_80 -ptx -o hello_world.ptx hello_world.cu $ grep -A 1 \"barrier.sync\" hello_world.ptx<\/code><\/pre>\n<p>\u041d\u0430\u0448\u0438 \u043f\u043e\u0434\u043e\u0437\u0440\u0435\u043d\u0438\u044f \u043e\u043f\u0440\u0430\u0432\u0434\u044b\u0432\u0430\u044e\u0442\u0441\u044f: \u043f\u043e\u0434 \u043a\u0430\u043f\u043e\u0442\u043e\u043c \u044d\u0442\u043e \u043e\u0431\u044b\u0447\u043d\u0430\u044f \u0438\u043d\u0441\u0442\u0440\u0443\u043a\u0446\u0438\u044f\u00a0<code>barrier.sync<\/code>. \u0422\u0430\u043a\u0438\u043c \u043e\u0431\u0440\u0430\u0437\u043e\u043c, \u0435\u0441\u043b\u0438 \u0432\u044b \u0445\u043e\u0440\u043e\u0448\u043e \u0432\u043b\u0430\u0434\u0435\u0435\u0442\u0435 \u0432\u0441\u0442\u0440\u0430\u0438\u0432\u0430\u0435\u043c\u044b\u043c PTX, \u0442\u043e \u043c\u043e\u0436\u0435\u0442\u0435 \u0432\u043e\u0441\u043f\u0440\u043e\u0438\u0437\u0432\u0435\u0441\u0442\u0438 \u0432\u044b\u0448\u0435\u043e\u043f\u0438\u0441\u0430\u043d\u043d\u043e\u0435 \u043f\u043e\u0432\u0435\u0434\u0435\u043d\u0438\u0435 \u0447\u0435\u0440\u0435\u0437 \u0437\u0430\u0433\u043e\u043b\u043e\u0432\u043e\u043a <code>&lt;cooperative_groups.h&gt;<\/code>. \u0410 \u0435\u0441\u043b\u0438 \u0432\u044b \u043f\u0438\u0448\u0435\u0442\u0435 \u043a\u043e\u0434, \u0432 \u043a\u043e\u0442\u043e\u0440\u043e\u043c \u0442\u0440\u0435\u0431\u0443\u0435\u0442\u0441\u044f \u043e\u0431\u0435\u0441\u043f\u0435\u0447\u0438\u0442\u044c \u043c\u0430\u043a\u0441\u0438\u043c\u0430\u043b\u044c\u043d\u0443\u044e \u043f\u0440\u043e\u0438\u0437\u0432\u043e\u0434\u0438\u0442\u0435\u043b\u044c\u043d\u043e\u0441\u0442\u044c, \u0442\u043e \u043e\u043d \u043f\u043e\u043b\u0443\u0447\u0438\u0442\u0441\u044f \u043d\u0435 \u0442\u043e\u043b\u044c\u043a\u043e \u0447\u0438\u0449\u0435, \u043d\u043e \u0438 \u043f\u0440\u043e\u0437\u0440\u0430\u0447\u043d\u0435\u0435, \u0430 \u0442\u0430\u043a\u0436\u0435 \u0431\u0443\u0434\u0435\u0442 \u043b\u0443\u0447\u0448\u0435 \u043f\u043e\u0434\u0434\u0430\u0432\u0430\u0442\u044c\u0441\u044f \u043e\u0442\u043b\u0430\u0434\u043a\u0435.<\/p>\n<pre><code class=\"cpp\">#include &lt;cuda_runtime.h&gt; #include &lt;stdio.h&gt; #include &lt;math.h&gt;  __device__ inline void grid_sync_ptx() { asm volatile(\"barrier.sync 0;\" ::); }  __device__ float3 compute_force(float3 position_first, float3 position_second) {     float3 r;     r.x = position_second.x - position_first.x;     r.y = position_second.y - position_first.y;     r.z = position_second.z - position_first.z;      float squared_distance = r.x * r.x + r.y * r.y + r.z * r.z + 1e-6f; \/\/ \u0438\u0437\u0431\u0435\u0433\u0430\u0435\u043c \u0434\u0435\u043b\u0435\u043d\u0438\u044f \u043d\u0430 \u043d\u043e\u043b\u044c      float reciprocal_distance = rsqrtf(squared_distance);     float reciprocal_cube = reciprocal_distance * reciprocal_distance * reciprocal_distance;      constexpr float gravitational_constant = 1.0f;     float scale = gravitational_constant * reciprocal_cube;     r.x *= scale;     r.y *= scale;     r.z *= scale;     return r; }  __global__ void cooperative_kernel(float3 *positions_old, float3 *positions_new, float3 *velocities_old, float3 *velocities_new, size_t count,                                    size_t iterations, float dt) {     size_t idx = blockIdx.x * blockDim.x + threadIdx.x;     if (idx &gt;= count) return;      for (size_t iter = 0; iter &lt; iterations; ++iter) {         float3 force = {0.0f, 0.0f, 0.0f};          \/\/ \u0410\u043a\u043a\u0443\u043c\u0443\u043b\u0438\u0440\u0443\u0435\u043c \u0441\u0438\u043b\u044b \u043e\u0442 \u0432\u0441\u0435\u0445 \u043f\u0440\u043e\u0447\u0438\u0445 \u0447\u0430\u0441\u0442\u0438\u0446         for (size_t j = 0; j &lt; count; ++j) {             if (j == idx) continue;             float3 f = compute_force(positions_old[idx], positions_old[j]);             force.x += f.x;             force.y += f.y;             force.z += f.z;         }          \/\/ \u041e\u0431\u043d\u043e\u0432\u043b\u044f\u0435\u043c \u0437\u043d\u0430\u0447\u0435\u043d\u0438\u044f \u0441\u043a\u043e\u0440\u043e\u0441\u0442\u0438 \u0438 \u043f\u043e\u043b\u043e\u0436\u0435\u043d\u0438\u044f         velocities_new[idx].x = velocities_old[idx].x + force.x * dt;         velocities_new[idx].y = velocities_old[idx].y + force.y * dt;         velocities_new[idx].z = velocities_old[idx].z + force.z * dt;          positions_new[idx].x = positions_old[idx].x + velocities_new[idx].x * dt;         positions_new[idx].y = positions_old[idx].y + velocities_new[idx].y * dt;         positions_new[idx].z = positions_old[idx].z + velocities_new[idx].z * dt;          grid_sync_ptx();          \/\/ \u041c\u0435\u043d\u044f\u0435\u043c \u0431\u0443\u0444\u0435\u0440\u044b, \u0433\u043e\u0442\u043e\u0432\u044f\u0441\u044c \u043a \u0441\u043b\u0435\u0434\u0443\u044e\u0449\u0435\u0439 \u0438\u0442\u0435\u0440\u0430\u0446\u0438\u0438         float3 *temp_pos = positions_old, *temp_vel = velocities_old;         positions_old = positions_new, positions_new = temp_pos;         velocities_old = velocities_new, velocities_new = temp_vel;          grid_sync_ptx();     } }  int main() {     cudaError_t err;     size_t num_particles = 256;     size_t iterations = 10;     float dt = 0.01f;      float3 *positions_old = nullptr, *positions_new = nullptr;     float3 *velocities_old = nullptr, *velocities_new = nullptr;      dim3 block;     dim3 grid;     void *kernel_args[7];      \/\/ \u0412\u044b\u0434\u0435\u043b\u044f\u0435\u043c \u043f\u0430\u043c\u044f\u0442\u044c     err = cudaMallocManaged(&amp;positions_old, num_particles * sizeof(float3));     if (err != cudaSuccess) goto cleanup;     err = cudaMallocManaged(&amp;positions_new, num_particles * sizeof(float3));     if (err != cudaSuccess) goto cleanup;     err = cudaMallocManaged(&amp;velocities_old, num_particles * sizeof(float3));     if (err != cudaSuccess) goto cleanup;     err = cudaMallocManaged(&amp;velocities_new, num_particles * sizeof(float3));     if (err != cudaSuccess) goto cleanup;      for (size_t i = 0; i &lt; num_particles; ++i) {         float theta = (float)i * 0.01f;         float phi = (float)i * 0.005f;         float radius = 10.0f + (i % 32) * 0.1f;         positions_old[i] = {radius * cosf(theta) * sinf(phi), radius * sinf(theta) * sinf(phi), radius * cosf(phi)};         velocities_old[i] = {0.01f * sinf(phi), 0.01f * cosf(theta), 0.01f * sinf(theta + phi)};     }      cudaDeviceProp props;     cudaGetDeviceProperties(&amp;props, 0);     if (!props.cooperativeLaunch) {         fprintf(stderr, \"Cooperative launch not supported on this device.\\n\");         err = cudaErrorNotSupported;         goto cleanup;     }      cudaStream_t stream;     err = cudaStreamCreate(&amp;stream);     if (err != cudaSuccess) {         fprintf(stderr, \"Failed to create stream: %s\\n\", cudaGetErrorString(err));         goto cleanup;     }      block = dim3(256);     grid = dim3((num_particles + block.x - 1) \/ block.x);     kernel_args[0] = &amp;positions_old;     kernel_args[1] = &amp;positions_new;     kernel_args[2] = &amp;velocities_old;     kernel_args[3] = &amp;velocities_new;     kernel_args[4] = &amp;num_particles;     kernel_args[5] = &amp;iterations;     kernel_args[6] = &amp;dt;      err = cudaLaunchKernel((void *)cooperative_kernel, grid, block, kernel_args, 0, stream);     if (err != cudaSuccess) {         fprintf(stderr, \"Failed to launch cooperative kernel: %s\\n\", cudaGetErrorString(err));         goto cleanup;     }      err = cudaStreamSynchronize(stream);     if (err != cudaSuccess) {         fprintf(stderr, \"Kernel execution failed: %s\\n\", cudaGetErrorString(err));         goto cleanup;     }      for (size_t i = 0; i &lt; num_particles; ++i)         printf(\"Final position[%zu] = (%f, %f, %f)\\n\", i, positions_old[i].x, positions_old[i].y, positions_old[i].z);  cleanup:     if (positions_old) cudaFree(positions_old);     if (positions_new) cudaFree(positions_new);     if (velocities_old) cudaFree(velocities_old);     if (velocities_new) cudaFree(velocities_new);     return (err == cudaSuccess) ? 0 : -1; }<\/code><\/pre>\n<p>\u041a\u0441\u0442\u0430\u0442\u0438, \u043f\u043e\u043a\u0430 \u043d\u0438\u043a\u0442\u043e \u043d\u0435 \u0432\u0438\u0434\u0435\u043b, \u043c\u044b \u0437\u0430\u043f\u0443\u0441\u0442\u0438\u043b\u0438\u0441\u044c \u043f\u043e \u0441\u0442\u0430\u0440\u0438\u043d\u043a\u0435 \u043f\u0440\u0438 \u043f\u043e\u043c\u043e\u0449\u0438 \u00a0<code>cudaLaunchKernel<\/code>, \u0430 \u043d\u0435 \u0441 <code>cudaLaunchCooperativeKernel<\/code>. \u0421\u0440\u0435\u0434\u0430 \u0432\u044b\u043f\u043e\u043b\u043d\u0435\u043d\u0438\u044f \u043d\u0430 \u044d\u0442\u043e \u0436\u0430\u043b\u043e\u0432\u0430\u0442\u044c\u0441\u044f \u043d\u0435 \u0431\u0443\u0434\u0435\u0442, \u043f\u0440\u0438 \u0443\u0441\u043b\u043e\u0432\u0438\u0438, \u0447\u0442\u043e \u0443 \u0432\u0430\u0441 \u043d\u0430 \u0443\u0441\u0442\u0440\u043e\u0439\u0441\u0442\u0432\u0435 \u043f\u043e\u0434\u0434\u0435\u0440\u0436\u0438\u0432\u0430\u0435\u0442\u0441\u044f <code>barrier.sync<\/code>. \u0415\u0441\u043b\u0438 \u0432\u044b \u0430\u0437\u0430\u0440\u0442\u043d\u044b, \u0442\u043e PTX \u043d\u0430\u0439\u0434\u0451\u0442\u0441\u044f \u0434\u043b\u044f \u0432\u0430\u0441 \u0438 \u0446\u0435\u043b\u044b\u0439 \u0441\u043f\u0435\u043a\u0442\u0440 \u0434\u0440\u0443\u0433\u0438\u0445 \u0431\u0430\u0440\u044c\u0435\u0440\u043e\u0432.<\/p>\n<h2>API \u0434\u0440\u0430\u0439\u0432\u0435\u0440\u043e\u0432 CUDA<\/h2>\n<p>\u041d\u0430\u043a\u043e\u043d\u0435\u0446, \u043d\u0435\u0437\u0430\u0441\u043b\u0443\u0436\u0435\u043d\u043d\u043e \u043e\u0431\u0445\u043e\u0434\u044f\u0442 \u0432\u043d\u0438\u043c\u0430\u043d\u0438\u0435\u043c \u0435\u0449\u0451 \u0431\u043e\u043b\u0435\u0435 \u043d\u0438\u0437\u043a\u043e\u0443\u0440\u043e\u0432\u043d\u0435\u0432\u044b\u0439 API \u0434\u0440\u0430\u0439\u0432\u0435\u0440\u043e\u0432 CUDA, \u043a\u043e\u0442\u043e\u0440\u044b\u0439 \u043c\u043e\u0436\u0435\u0442 \u0431\u044b\u0442\u044c \u0438\u0441\u043a\u043b\u044e\u0447\u0438\u0442\u0435\u043b\u044c\u043d\u043e \u043f\u043e\u043b\u0435\u0437\u0435\u043d \u043d\u0430 \u043f\u0440\u0430\u043a\u0442\u0438\u043a\u0435. \u0414\u0430, \u043a\u043e\u0434\u0430 \u0432 \u043d\u0451\u043c \u043c\u043d\u043e\u0433\u043e\u0432\u0430\u0442\u043e, \u0437\u0430\u0442\u043e \u043e\u043d \u043f\u0440\u0435\u0434\u043e\u0441\u0442\u0430\u0432\u043b\u044f\u0435\u0442 \u0432\u0430\u043c \u043f\u043e\u043b\u043d\u044b\u0439 \u043a\u043e\u043d\u0442\u0440\u043e\u043b\u044c \u043d\u0430\u0434 \u0437\u0430\u0433\u0440\u0443\u0437\u043a\u043e\u0439 \u0438 \u0437\u0430\u043f\u0443\u0441\u043a\u043e\u043c \u044f\u0434\u0440\u0430, \u0432 \u0442\u043e\u043c \u0447\u0438\u0441\u043b\u0435, \u043f\u043e\u0434\u0434\u0435\u0440\u0436\u0438\u0432\u0430\u0435\u0442 \u0434\u0438\u043d\u0430\u043c\u0438\u0447\u0435\u0441\u043a\u0443\u044e \u0437\u0430\u0433\u0440\u0443\u0437\u043a\u0443 PTX, CUBIN \u0438\u043b\u0438 SASS \u0432\u043e \u0432\u0440\u0435\u043c\u044f \u0432\u044b\u043f\u043e\u043b\u043d\u0435\u043d\u0438\u044f. \u041d\u0435\u043c\u043d\u043e\u0433\u043e \u0437\u0430\u0431\u0435\u0433\u0430\u044f \u0432\u043f\u0435\u0440\u0451\u0434, \u0440\u0435\u043a\u043e\u043c\u0435\u043d\u0434\u0443\u044e \u043f\u043e\u043b\u043d\u043e\u0441\u0442\u044c\u044e \u043e\u0442\u0433\u0440\u0430\u043d\u0438\u0447\u0438\u0442\u044c \u043a\u043e\u0434 \u044f\u0434\u0440\u0430 \u043e\u0442 \u043a\u043e\u0434\u0430 \u0445\u043e\u0441\u0442\u0430, \u043f\u0440\u0438\u0447\u0451\u043c, \u0438\u0441\u043f\u043e\u043b\u044c\u0437\u043e\u0432\u0430\u0442\u044c \u0434\u043b\u044f \u043d\u0438\u0445 \u0434\u0432\u0430 \u043e\u0442\u0434\u0435\u043b\u044c\u043d\u044b\u0445 \u043a\u043e\u043c\u043f\u0438\u043b\u044f\u0442\u043e\u0440\u0430 \u0438 \u043e\u0431\u0443\u0441\u0442\u0440\u043e\u0438\u0442\u044c \u043c\u0435\u0436\u0434\u0443 \u043d\u0438\u043c\u0438 \u0441\u0442\u0430\u0431\u0438\u043b\u044c\u043d\u044b\u0439 ABI. \u0412\u043e\u0442 \u043a\u0430\u043a \u0443\u043f\u0440\u043e\u0449\u0451\u043d\u043d\u043e\u043c \u0432\u0438\u0434\u0435 \u043c\u043e\u0433 \u0431\u044b \u0432\u044b\u0433\u043b\u044f\u0434\u0435\u0442\u044c \u043d\u0430\u0448 \u043a\u043e\u0434 \u0445\u043e\u0441\u0442\u0430 \u043d\u0430 C99:<\/p>\n<pre><code class=\"cpp\">#include &lt;cuda.h&gt; #include &lt;stdio.h&gt; #include &lt;math.h&gt;  #define CUDA_CHECK(err)                             \\     if (err != CUDA_SUCCESS) {                      \\         const char *msg;                            \\         cuGetErrorString(err, &amp;msg);                \\         fprintf(stderr, \"CUDA error: %s\\n\", msg);   \\         goto cleanup;                               \\     }  int main() {     CUresult err;     size_t num_particles = 256;     size_t iterations = 10;     float dt = 0.01f;      CUdevice device;     CUcontext context = NULL;     CUmodule module = NULL;     CUfunction kernel;     CUstream stream = NULL;     float *positions_old = NULL, *positions_new = NULL;     float *velocities_old = NULL, *velocities_new = NULL;     void *kernel_args[7];      \/\/ \u0418\u043d\u0438\u0446\u0438\u0430\u043b\u0438\u0437\u0430\u0446\u0438\u044f CUDA     err = cuInit(0);     CUDA_CHECK(err);     err = cuDeviceGet(&amp;device, 0);     CUDA_CHECK(err);     err = cuCtxCreate(&amp;context, 0, device);     CUDA_CHECK(err);     err = cuStreamCreate(&amp;stream, CU_STREAM_DEFAULT);     CUDA_CHECK(err);      \/\/ \u0417\u0430\u0433\u0440\u0443\u0437\u043a\u0430 PTX \"\u0431\u0430\u0439\u0442-\u043a\u043e\u0434\u0430\", \u043a\u043e\u0442\u043e\u0440\u044b\u0439 \u0432\u043f\u043e\u0441\u043b\u0435\u0434\u0441\u0442\u0432\u0438\u0438 \u0431\u0443\u0434\u0435\u0442 \u0434\u0438\u043d\u0430\u043c\u0438\u0447\u0435\u0441\u043a\u0438 \u043a\u043e\u043c\u043f\u0438\u043b\u0438\u0440\u043e\u0432\u0430\u0442\u044c\u0441\u044f \u0432 SASS     err = cuModuleLoad(&amp;module, \"hello_world.ptx\");     CUDA_CHECK(err);     err = cuModuleGetFunction(&amp;kernel, module, \"cooperative_kernel\");     CUDA_CHECK(err);      \/\/ \u0412\u044b\u0434\u0435\u043b\u0435\u043d\u0438\u0435 \u0443\u043f\u0440\u0430\u0432\u043b\u044f\u0435\u043c\u043e\u0439 \u043f\u0430\u043c\u044f\u0442\u0438 \u0434\u043b\u044f \u043f\u043e\u043b\u043e\u0436\u0435\u043d\u0438\u0439 \u0438 \u0441\u043a\u043e\u0440\u043e\u0441\u0442\u0435\u0439      size_t buffer_size = num_particles * sizeof(float) * 3;     err = cuMemAllocManaged((CUdeviceptr *)&amp;positions_old, buffer_size, CU_MEM_ATTACH_GLOBAL);     CUDA_CHECK(err);     err = cuMemAllocManaged((CUdeviceptr *)&amp;positions_new, buffer_size, CU_MEM_ATTACH_GLOBAL);     CUDA_CHECK(err);     err = cuMemAllocManaged((CUdeviceptr *)&amp;velocities_old, buffer_size, CU_MEM_ATTACH_GLOBAL);     CUDA_CHECK(err);     err = cuMemAllocManaged((CUdeviceptr *)&amp;velocities_new, buffer_size, CU_MEM_ATTACH_GLOBAL);     CUDA_CHECK(err);      \/\/ \u0418\u043d\u0438\u0446\u0438\u0430\u043b\u0438\u0437\u0430\u0446\u0438\u044f \u043f\u043e\u043b\u043e\u0436\u0435\u043d\u0438\u0439 \u0438 \u0441\u043a\u043e\u0440\u043e\u0441\u0442\u0435\u0439      for (size_t i = 0; i &lt; num_particles; ++i) {         float theta = (float)i * 0.01f;         float phi = (float)i * 0.005f;         float radius = 10.0f + (i % 32) * 0.1f;         positions_old[3 * i + 0] = radius * cosf(theta) * sinf(phi);         positions_old[3 * i + 1] = radius * sinf(theta) * sinf(phi);         positions_old[3 * i + 2] = radius * cosf(phi);         velocities_old[3 * i + 0] = 0.01f * sinf(phi);         velocities_old[3 * i + 1] = 0.01f * cosf(theta);         velocities_old[3 * i + 2] = 0.01f * sinf(theta + phi);     }      kernel_args[0] = &amp;positions_old;     kernel_args[1] = &amp;positions_new;     kernel_args[2] = &amp;velocities_old;     kernel_args[3] = &amp;velocities_new;     kernel_args[4] = &amp;num_particles;     kernel_args[5] = &amp;iterations;     kernel_args[6] = &amp;dt;      \/\/ \u0417\u0430\u043f\u0443\u0441\u043a \u044f\u0434\u0440\u0430     int threads_per_block = 256;     int blocks_per_grid = (num_particles + threads_per_block - 1) \/ threads_per_block;     err = cuLaunchKernel(kernel,                          blocks_per_grid, 1, 1,                          threads_per_block, 1, 1,                          0, stream,                          kernel_args, NULL);     CUDA_CHECK(err);     err = cuStreamSynchronize(stream);     CUDA_CHECK(err);      \/\/ \u041b\u043e\u0433\u0438\u0440\u043e\u0432\u0430\u043d\u0438\u0435 \u043e\u043a\u043e\u043d\u0447\u0430\u0442\u0435\u043b\u044c\u043d\u044b\u0445 \u043f\u043e\u043b\u043e\u0436\u0435\u043d\u0438\u0439     for (size_t i = 0; i &lt; num_particles; ++i)         printf(\"Final position[%zu] = (%f, %f, %f)\\n\", i,                positions_old[3 * i + 0],                positions_old[3 * i + 1],                positions_old[3 * i + 2]);  cleanup:     if (stream) cuStreamDestroy(stream);     if (positions_old) cuMemFree((CUdeviceptr)positions_old);     if (positions_new) cuMemFree((CUdeviceptr)positions_new);     if (velocities_old) cuMemFree((CUdeviceptr)velocities_old);     if (velocities_new) cuMemFree((CUdeviceptr)velocities_new);     if (module) cuModuleUnload(module);     if (context) cuCtxDestroy(context);     return (err == CUDA_SUCCESS) ? 0 : -1; }<\/code><\/pre>\n<p>\u0412 \u0434\u0430\u043d\u043d\u043e\u043c \u0441\u043b\u0443\u0447\u0430\u0435 \u043e\u0441\u043e\u0431\u043e\u0433\u043e \u0432\u043d\u0438\u043c\u0430\u043d\u0438\u044f \u0442\u0440\u0435\u0431\u0443\u0435\u0442 \u043b\u0438\u0448\u044c \u0434\u0435\u043a\u043e\u0440\u0438\u0440\u043e\u0432\u0430\u043d\u0438\u0435 \u0438\u043c\u0451\u043d. \u0427\u0442\u043e\u0431\u044b \u044d\u0442\u043e\u0442 \u043c\u0435\u0445\u0430\u043d\u0438\u0437\u043c \u0440\u0430\u0431\u043e\u0442\u0430\u043b, \u0443\u0431\u0435\u0434\u0438\u0442\u0435\u0441\u044c, \u0447\u0442\u043e \u043e\u0431\u044a\u044f\u0432\u043b\u0435\u043d\u0438\u0435 \u044f\u0434\u0440\u0430 \u0443 \u0432\u0430\u0441 \u0432 \u0444\u0430\u0439\u043b\u0435 <code>.cu<\/code>\u00a0\u043e\u0431\u0451\u0440\u043d\u0443\u0442\u043e \u0432\u00a0<code>extern \"C\"<\/code>:<\/p>\n<pre><code>extern \"C\" __global__ void cooperative_kernel(...);<\/code><\/pre>\n<p>\u0423\u043b\u0430\u0434\u0438\u0432 \u044d\u0442\u043e, \u043c\u043e\u0436\u043d\u043e \u043a\u043e\u043c\u043f\u0438\u043b\u0438\u0440\u043e\u0432\u0430\u0442\u044c \u043a\u043e\u0434 GPU \u0432 PTX \u043f\u0440\u0438 \u043f\u043e\u043c\u043e\u0449\u0438 NVCC, \u0430 \u043a\u043e\u0434 \u0445\u043e\u0441\u0442\u0430 \u043f\u0440\u0438 \u043f\u043e\u043c\u043e\u0449\u0438 GCC \u2014 \u0441\u043e\u0432\u0435\u0440\u0448\u0435\u043d\u043d\u043e \u043d\u0435\u0437\u0430\u0432\u0438\u0441\u0438\u043c\u043e:<\/p>\n<h2>\u0417\u0430\u043a\u043b\u044e\u0447\u0435\u043d\u0438\u0435<\/h2>\n<p>\u041a\u043e\u043d\u0435\u0447\u043d\u043e, \u0443 \u043d\u0430\u0441 \u043f\u043e\u043b\u0443\u0447\u0438\u043b\u043e\u0441\u044c \u0433\u043e\u0440\u0430\u0437\u0434\u043e \u0431\u043e\u043b\u044c\u0448\u0435 \u043a\u043e\u0434\u0430, \u0447\u0435\u043c \u0432 \u0438\u0441\u0445\u043e\u0434\u043d\u043e\u043c \u043f\u0440\u0438\u043c\u0435\u0440\u0435 \u0441 <code>&lt;&lt;&lt;1, 1&gt;&gt;&gt;<\/code> \u2014 \u043d\u043e \u043e\u0431\u044b\u0447\u043d\u043e \u0442\u0430\u043a \u0438 \u0431\u044b\u0432\u0430\u0435\u0442, \u0435\u0441\u043b\u0438 \u0441\u0442\u0430\u0440\u0430\u0435\u0448\u044c\u0441\u044f \u0432\u0441\u0451 \u0434\u0435\u043b\u0430\u0442\u044c \u043f\u0440\u0430\u0432\u0438\u043b\u044c\u043d\u043e. \u041f\u0440\u0438 \u044d\u0442\u043e\u043c \u043e\u0442\u043c\u0435\u0447\u0443, \u0447\u0442\u043e \u0432 \u043f\u043e\u0441\u043b\u0435\u0434\u043d\u0438\u0435 \u0433\u043e\u0434\u044b \u043f\u043e\u044f\u0432\u0438\u043b\u043e\u0441\u044c \u043c\u043d\u043e\u0436\u0435\u0441\u0442\u0432\u043e \u0438\u043d\u0441\u0442\u0440\u0443\u043c\u0435\u043d\u0442\u043e\u0432, \u0443\u043f\u0440\u043e\u0449\u0430\u044e\u0449\u0438\u0445 \u043f\u0440\u043e\u0442\u043e\u0442\u0438\u043f\u0438\u0440\u043e\u0432\u0430\u043d\u0438\u0435 \u043a\u043e\u0434\u0430 CUDA, \u0432 \u0442\u043e\u043c \u0447\u0438\u0441\u043b\u0435, \u0440\u0430\u0437\u043d\u043e\u043e\u0431\u0440\u0430\u0437\u043d\u044b\u0435 DSL \u0438 \u043a\u043e\u043c\u043f\u0438\u043b\u044f\u0442\u043e\u0440\u044b, \u043a\u043e\u0442\u043e\u0440\u044b\u043c\u0438 NVIDIA \u043b\u044e\u0431\u0438\u0442 \u043a\u043e\u0437\u044b\u0440\u043d\u0443\u0442\u044c. \u041d\u043e, \u0435\u0441\u043b\u0438 \u0433\u043e\u0442\u043e\u0432\u0438\u0442\u044c \u0438 \u0432\u044b\u043f\u0443\u0441\u043a\u0430\u0442\u044c \u044f\u0434\u0440\u0430 \u0434\u043b\u044f \u043f\u0440\u043e\u0434\u0430\u043a\u0448\u0435\u043d\u0430 \u0442\u0430\u043a, \u043a\u0430\u043a \u043e\u043f\u0438\u0441\u0430\u043d\u043e \u0437\u0434\u0435\u0441\u044c, \u0442\u043e \u043e\u043d\u0438 \u043f\u043e\u043b\u0443\u0447\u0430\u044e\u0442\u0441\u044f \u0443\u0434\u0438\u0432\u0438\u0442\u0435\u043b\u044c\u043d\u043e \u0441\u0442\u0430\u0431\u0438\u043b\u044c\u043d\u044b\u043c\u0438 \u0438 \u0440\u0430\u0431\u043e\u0442\u0430\u044e\u0442 \u043f\u043e 10 \u0438 \u0431\u043e\u043b\u0435\u0435 \u043b\u0435\u0442. <\/p>\n<p>\u0418\u0437\u043c\u0435\u043d\u0438\u043b\u0430\u0441\u044c \u043b\u0438\u0448\u044c \u0441\u043b\u043e\u0436\u043d\u043e\u0441\u0442\u044c \u0441\u0430\u043c\u0438\u0445 \u044f\u0434\u0435\u0440 \u2014 \u0442\u0435\u043f\u0435\u0440\u044c \u0432 \u043d\u0438\u0445 \u043c\u0435\u043d\u044c\u0448\u0435 \u0440\u0430\u0441\u043f\u0430\u0440\u0430\u043b\u043b\u0435\u043b\u0438\u0432\u0430\u044e\u0442\u0441\u044f \u0434\u0430\u043d\u043d\u044b\u0435, \u0438 \u044f\u0434\u0440\u0430 \u0441\u0431\u043b\u0438\u0436\u0430\u044e\u0442\u0441\u044f \u043f\u043e \u043f\u0440\u0438\u043d\u0446\u0438\u043f\u0443 \u0440\u0430\u0431\u043e\u0442\u044b \u0441 \u043a\u043e\u043d\u043a\u0443\u0440\u0435\u043d\u0442\u043d\u044b\u043c\u0438 \u0430\u043b\u0433\u043e\u0440\u0438\u0442\u043c\u0430\u043c\u0438 \u0434\u043b\u044f \u0426\u041f. \u0417\u0434\u0435\u0441\u044c \u043f\u0440\u0438\u043c\u0435\u043d\u044f\u044e\u0442\u0441\u044f \u0438 \u0430\u0442\u043e\u043c\u0430\u0440\u043d\u044b\u0435 \u043e\u043f\u0435\u0440\u0430\u0446\u0438\u0438, \u0438 \u0440\u0435\u0434\u0443\u043a\u0446\u0438\u044f \u043d\u0430 \u0443\u0440\u043e\u0432\u043d\u0435 \u0432\u0430\u0440\u043f\u043e\u0432, \u0438 \u0432\u043e \u0432\u0441\u0451 \u044d\u0442\u043e \u0437\u0430\u043f\u0435\u0447\u0435\u043d\u0430 \u043b\u043e\u0433\u0438\u043a\u0430 \u0442\u0435\u043d\u0437\u043e\u0440\u043d\u044b\u0445 \u044f\u0434\u0435\u0440. \u00a0<a href=\"https:\/\/github.com\/ashvardanian\/less_slow.cpp\/releases\/tag\/v0.9.0\" rel=\"noopener noreferrer nofollow\">\u042d\u0442\u0438 \u0440\u0435\u0430\u043b\u0438\u0437\u0430\u0446\u0438\u0438 \u0432\u044b\u0433\u043b\u044f\u0434\u044f\u0442 \u0438 \u043e\u0449\u0443\u0449\u0430\u044e\u0442\u0441\u044f \u043f\u043e-\u0440\u0430\u0437\u043d\u043e\u043c\u0443 \u0432 \u043a\u0430\u0436\u0434\u043e\u043c \u043d\u043e\u0432\u043e\u043c \u043f\u043e\u043a\u043e\u043b\u0435\u043d\u0438\u0438 GPU<\/a>, \u0432 \u0447\u0451\u043c \u043c\u043d\u0435 \u0434\u043e\u0432\u0435\u043b\u043e\u0441\u044c \u0443\u0431\u0435\u0434\u0438\u0442\u044c\u0441\u044f \u043d\u0430 \u0441\u043e\u0431\u0441\u0442\u0432\u0435\u043d\u043d\u043e\u0439 \u0448\u043a\u0443\u0440\u0435 \u043f\u0440\u0438\u00a0<a href=\"http:\/\/github.com\/ashvardanian\" rel=\"noopener noreferrer nofollow\">\u043f\u043e\u0440\u0442\u0438\u0440\u043e\u0432\u0430\u043d\u0438\u0438 \u0434\u0440\u0443\u0433\u0438\u0445 \u0431\u0438\u0431\u043b\u0438\u043e\u0442\u0435\u043a<\/a>.\u00a0\u00a0<\/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\/937826\/\"> https:\/\/habr.com\/ru\/articles\/937826\/<\/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<p>\u0412\u0435\u0440\u043e\u044f\u0442\u043d\u043e, \u0432\u0430\u043c \u0443\u0436\u0435 \u043f\u043e\u043f\u0430\u0434\u0430\u043b\u0438\u0441\u044c \u043f\u043e\u0434\u043e\u0431\u043d\u044b\u0435 \u0440\u0443\u043a\u043e\u0432\u043e\u0434\u0441\u0442\u0432\u0430 \u043f\u043e CUDA: \u0445\u0440\u0435\u0441\u0442\u043e\u043c\u0430\u0442\u0438\u0439\u043d\u044b\u0439 \u043f\u0440\u0438\u043c\u0435\u0440 \u00abHello World\u00bb, \u0432 \u043a\u043e\u0442\u043e\u0440\u043e\u043c \u043f\u0435\u0440\u0435\u043c\u0435\u0448\u0430\u043d \u043a\u043e\u0434 \u0434\u043b\u044f \u0426\u041f \u0438 \u0433\u0440\u0430\u0444\u0438\u0447\u0435\u0441\u043a\u043e\u0433\u043e \u043f\u0440\u043e\u0446\u0435\u0441\u0441\u043e\u0440\u0430. \u0412\u0441\u0451 \u044d\u0442\u043e \u0441\u043b\u043e\u0436\u0435\u043d\u043e \u0432 \u043e\u0434\u0438\u043d \u0433\u0435\u0442\u0435\u0440\u043e\u0433\u0435\u043d\u043d\u044b\u0439 \u0444\u0430\u0439\u043b \u0441 \u0438\u0441\u0445\u043e\u0434\u043d\u0438\u043a\u0430\u043c\u0438 \u043d\u0430 CUDA C++, \u0430 \u0434\u043b\u044f \u0437\u0430\u043f\u0443\u0441\u043a\u0430 \u044f\u0434\u0440\u0430 \u043f\u0440\u0438\u043c\u0435\u043d\u044f\u0435\u0442\u0441\u044f \u0441\u0438\u043d\u0442\u0430\u043a\u0441\u0438\u0441 NVCC \u0441 \u0442\u0440\u043e\u0439\u043d\u044b\u043c\u0438 \u0443\u0433\u043b\u043e\u0432\u044b\u043c\u0438 \u0441\u043a\u043e\u0431\u043a\u0430\u043c\u0438\u00a0<code>&lt;&lt;&lt;&gt;&gt;&gt;<\/code>, \u043a\u043e\u0442\u043e\u0440\u044b\u0439 \u0443\u0436\u0435 \u0441\u0442\u0430\u043b \u043a\u0443\u043b\u044c\u0442\u043e\u0432\u044b\u043c:<\/p>\n<pre><code class=\"cpp\">#include &lt;cuda_runtime.h&gt; #include &lt;stdio.h&gt;  __global__ void kernel() {     printf(\"Hello World from block %d, thread %d\\n\", blockIdx.x, threadIdx.x); }  int main() {     kernel&lt;&lt;&lt;1, 1&gt;&gt;&gt;(); \/\/ \u0412\u043e\u0437\u0432\u0440\u0430\u0449\u0430\u0435\u0442 `void`?!          return cudaDeviceSynchronize() == cudaSuccess ? 0 : -1; }<\/code><\/pre>\n<p>\u0412\u0440\u0435\u043c\u044f \u0438\u0434\u0451\u0442, \u0430 \u0442\u0430\u043a\u043e\u0439 \u043f\u0430\u0442\u0442\u0435\u0440\u043d \u043f\u043e-\u043f\u0440\u0435\u0436\u043d\u0435\u043c\u0443 \u043f\u043e\u043f\u0430\u0434\u0430\u0435\u0442\u0441\u044f \u043c\u043d\u0435 \u0432 \u043f\u0440\u043e\u0434\u0430\u043a\u0448\u0435\u043d-\u043a\u043e\u0434\u0435. \u041f\u0440\u0438\u0437\u043d\u0430\u044e\u0441\u044c, \u043a\u043e\u0435-\u0433\u0434\u0435 \u043e\u043d \u0432\u0441\u043f\u043b\u044b\u0432\u0430\u0435\u0442 \u0438 \u0432 \u043c\u043e\u0438\u0445 \u043b\u044e\u0431\u0438\u0442\u0435\u043b\u044c\u0441\u043a\u0438\u0445 \u043f\u0440\u043e\u0435\u043a\u0442\u0430\u0445 \u2014 <a href=\"https:\/\/github.com\/ashvardanian\/ParallelReductionsBenchmark\" rel=\"noopener noreferrer nofollow\">\u0440\u0430\u0437<\/a>,\u00a0<a href=\"https:\/\/github.com\/ashvardanian\/cuda-python-starter-kit\" rel=\"noopener noreferrer nofollow\">\u0434\u0432\u0430<\/a>, <a href=\"https:\/\/github.com\/ashvardanian\/scaling-democracy\" rel=\"noopener noreferrer nofollow\">\u0442\u0440\u0438<\/a>. \u041d\u043e \u044d\u0442\u043e \u043d\u0435 \u043b\u0443\u0447\u0448\u0430\u044f \u0438\u0434\u0435\u044f, \u043f\u043e\u043b\u0430\u0433\u0430\u0442\u044c\u0441\u044f \u0432 \u0441\u0435\u0440\u044c\u0451\u0437\u043d\u043e\u043c \u043a\u043e\u0434\u0435 \u043d\u0430 \u0437\u0430\u043f\u0443\u0441\u043a \u044f\u0434\u0440\u0430 \u0447\u0435\u0440\u0435\u0437 \u0442\u0440\u043e\u0439\u043d\u044b\u0435 \u0443\u0433\u043b\u043e\u0432\u044b\u0435 \u0441\u043a\u043e\u0431\u043a\u0438. \u0412 \u0442\u0430\u043a\u043e\u043c \u0441\u043b\u0443\u0447\u0430\u0435 \u043f\u0440\u043e\u0433\u0440\u0430\u043c\u043c\u0430 \u043d\u0435 \u0432\u043e\u0437\u0432\u0440\u0430\u0449\u0430\u0435\u0442 \u043a\u043e\u0434\u044b \u043e\u0448\u0438\u0431\u043e\u043a, \u043f\u043e\u044d\u0442\u043e\u043c\u0443 \u043c\u043e\u0436\u0435\u0442 \u043f\u043e\u043a\u0430\u0437\u0430\u0442\u044c\u0441\u044f \u043e\u0431\u043c\u0430\u043d\u0447\u0438\u0432\u043e \u043f\u0440\u043e\u0441\u0442\u043e\u0439. \u041d\u0438\u0436\u0435 \u0432\u0430\u0441 \u0436\u0434\u0443\u0442 \u043f\u0440\u0438\u043c\u0435\u0440\u043d\u043e 25 \u043a\u0438\u043b\u043e\u0431\u0430\u0439\u0442 \u0442\u0435\u043a\u0441\u0442\u0430, \u0432 \u043a\u043e\u0442\u043e\u0440\u044b\u0445 \u043c\u044b \u043e\u0431\u0441\u0443\u0434\u0438\u043c <em>\u043d\u0435 \u0441\u0430\u043c\u044b\u0435 \u043a\u043e\u0440\u044f\u0432\u044b\u0435<\/em> \u0441\u043f\u043e\u0441\u043e\u0431\u044b \u0437\u0430\u043f\u0443\u0441\u043a\u0430 \u044f\u0434\u0435\u0440.<\/p>\n<h2>\u041e\u0441\u043d\u043e\u0432\u044b \u0438 \u043a\u043e\u0440\u0440\u0435\u043a\u0442\u043d\u043e\u0441\u0442\u044c<\/h2>\n<p>\u0412\u044b\u0448\u0435\u043f\u0440\u0438\u0432\u0435\u0434\u0451\u043d\u043d\u044b\u0439 \u043a\u043e\u0434 \u0441\u043a\u043e\u043c\u043f\u0438\u043b\u0438\u0440\u0443\u0435\u0442\u0441\u044f, \u043f\u043e\u0441\u043b\u0435 \u0447\u0435\u0433\u043e \u0432\u044b\u043f\u043e\u043b\u043d\u0438\u0442 \u043e\u0436\u0438\u0434\u0430\u0435\u043c\u044b\u0439 \u0432\u044b\u0432\u043e\u0434: <\/p>\n<pre><code>$ nvcc -o hello_world hello_world.cu &amp;&amp; .\/hello_world &gt; Hello World from block 0, thread 0<\/code><\/pre>\n<p>\u0412 \u043a\u0430\u043a\u043e\u043c-\u0442\u043e \u0441\u043c\u044b\u0441\u043b\u0435 \u0435\u0433\u043e \u0443\u0436\u0435 \u043c\u043e\u0436\u043d\u043e \u0441\u0447\u0438\u0442\u0430\u0442\u044c \u00ab\u043a\u043e\u0440\u0440\u0435\u043a\u0442\u043d\u044b\u043c\u00bb.<\/p>\n<p>\u041d\u043e \u0432 \u043d\u0430\u0448\u0435 \u0432\u0440\u0435\u043c\u044f \u043e\u0431\u044b\u0447\u043d\u044b \u0441\u0438\u0441\u0442\u0435\u043c\u044b, \u0432 \u043a\u043e\u0442\u043e\u0440\u044b\u0445 \u043f\u0440\u0435\u0434\u0443\u0441\u043c\u0430\u0442\u0440\u0438\u0432\u0430\u0435\u0442\u0441\u044f \u043d\u0435 \u043c\u0435\u043d\u0435\u0435 8 GPU \u043d\u0430 \u043f\u043b\u0430\u0442\u0443 HGX, \u043f\u043e\u044d\u0442\u043e\u043c\u0443, \u0435\u0441\u0442\u0435\u0441\u0442\u0432\u0435\u043d\u043d\u043e, \u0445\u043e\u0447\u0435\u0442\u0441\u044f \u0440\u0435\u0430\u043b\u0438\u0437\u043e\u0432\u0430\u0442\u044c \u043d\u0435\u043a\u043e\u0442\u043e\u0440\u044b\u0439 \u043f\u0430\u0440\u0430\u043b\u043b\u0435\u043b\u0438\u0437\u043c \u2014 \u0447\u0442\u043e\u0431\u044b \u0433\u0430\u0440\u0430\u043d\u0442\u0438\u0440\u043e\u0432\u0430\u0442\u044c, \u0447\u0442\u043e \u0431\u0443\u0434\u0443\u0442 \u0437\u0430\u043f\u0443\u0441\u043a\u0430\u0442\u044c\u0441\u044f \u044f\u0434\u0440\u0430 \u043d\u0430 \u043a\u0430\u0436\u0434\u043e\u043c \u0438\u0437 \u043d\u0438\u0445.<\/p>\n<p> \u041e\u043f\u0443\u0441\u043a\u0430\u044f \u0431\u0430\u0437\u043e\u0432\u044b\u0435 \u0430\u0441\u043f\u0435\u043a\u0442\u044b \u044d\u043a\u0441\u043f\u043b\u0443\u0430\u0442\u0430\u0446\u0438\u0438, \u043e\u0442\u043c\u0435\u0447\u0443, \u0447\u0442\u043e \u043a\u0430\u0436\u0434\u044b\u0439 \u0443\u0437\u0435\u043b DGX H100 \u043e\u0441\u043d\u0430\u0449\u0430\u0435\u0442\u0441\u044f \u0434\u0432\u0443\u043c\u044f \u043c\u043e\u0433\u0443\u0447\u0438\u043c\u0438 \u044f\u0434\u0440\u0430\u043c\u0438 \u0426\u041f, \u043a\u043e\u0442\u043e\u0440\u044b\u0435 \u0431\u0435\u0437 \u043f\u0435\u0440\u0435\u0440\u044b\u0432\u0430 \u0436\u043e\u043d\u0433\u043b\u0438\u0440\u0443\u044e\u0442 \u0441\u043b\u043e\u0436\u043d\u044b\u043c\u0438 \u0433\u0440\u0430\u0444\u0430\u043c\u0438 \u0432\u044b\u043f\u043e\u043b\u043d\u0435\u043d\u0438\u044f, \u043e\u0434\u043d\u043e\u0432\u0440\u0435\u043c\u0435\u043d\u043d\u043e \u043f\u0435\u0440\u0435\u0434\u0430\u0432\u0430\u044f \u043d\u0430 \u0432\u0445\u043e\u0434 \u0438 \u043d\u0430 \u0432\u044b\u0445\u043e\u0434 \u0441\u043e\u0442\u043d\u0438 \u0433\u0438\u0433\u0430\u0431\u0430\u0439\u0442. \u0412\u0441\u0451 \u044d\u0442\u043e \u2014 \u0430\u0441\u0438\u043d\u0445\u0440\u043e\u043d\u043d\u043e.<\/p>\n<p>\u0412 \u0442\u0430\u043a\u043e\u0439 \u0441\u0438\u0441\u0442\u0435\u043c\u0435 \u043e\u0447\u0435\u043d\u044c \u043c\u043d\u043e\u0433\u043e\u0435 \u043c\u043e\u0436\u0435\u0442 \u0441\u0431\u0438\u0442\u044c\u0441\u044f, \u0442\u0430\u043a \u0447\u0442\u043e \u0434\u0430\u0432\u0430\u0439\u0442\u0435 \u0441\u0444\u043e\u0440\u043c\u0443\u043b\u0438\u0440\u0443\u0435\u043c \u043d\u0435\u0441\u043a\u043e\u043b\u044c\u043a\u043e \u043e\u0441\u043d\u043e\u0432\u043e\u043f\u043e\u043b\u0430\u0433\u0430\u044e\u0449\u0438\u0445 \u043f\u0440\u0430\u0432\u0438\u043b \u043e\u0440\u043a\u0435\u0441\u0442\u0440\u0430\u0446\u0438\u0438 \u044f\u0434\u0435\u0440 GPU:<\/p>\n<ul>\n<li>\n<p>\u0417\u0430\u043f\u0443\u0441\u043a \u043a\u0430\u0436\u0434\u043e\u0433\u043e \u044f\u0434\u0440\u0430 \u0441\u043e\u043f\u0440\u044f\u0436\u0451\u043d \u0441 \u0441\u0435\u0440\u044c\u0451\u0437\u043d\u043e\u0439 \u0437\u0430\u0434\u0435\u0440\u0436\u043a\u043e\u0439, \u043f\u043e\u044d\u0442\u043e\u043c\u0443 \u0442\u0430\u043a\u0438\u0435 \u043e\u043f\u0435\u0440\u0430\u0446\u0438\u0438 \u0434\u043e\u043b\u0436\u043d\u044b \u0432\u044b\u043f\u043e\u043b\u043d\u044f\u0442\u044c\u0441\u044f \u0430\u0441\u0438\u043d\u0445\u0440\u043e\u043d\u043d\u043e.<\/p>\n<\/li>\n<li>\n<p>\u0420\u0430\u0431\u043e\u0442\u0443 \u0432 \u043f\u0440\u0435\u0434\u0435\u043b\u0430\u0445 \u043f\u043e\u0442\u043e\u043a\u043e\u0432 \u0441\u043b\u0435\u0434\u0443\u0435\u0442 \u0440\u0430\u0437\u0434\u0430\u0432\u0430\u0442\u044c \u044f\u0432\u043d\u043e.<\/p>\n<\/li>\n<li>\n<p>\u0412\u044b\u0437\u043e\u0432\u044b API CUDA \u0438 \u043f\u0443\u0441\u043a\u0438 \u044f\u0434\u0440\u0430 \u0434\u043e\u043b\u0436\u043d\u044b \u0441\u043e\u043f\u0440\u043e\u0432\u043e\u0436\u0434\u0430\u0442\u044c\u0441\u044f \u043d\u0430\u0434\u0451\u0436\u043d\u043e\u0439 \u043f\u0440\u043e\u0432\u0435\u0440\u043a\u043e\u0439 \u043d\u0430 \u0441\u043b\u0443\u0447\u0430\u0439 \u0432\u043e\u0437\u043c\u043e\u0436\u043d\u044b\u0445 \u043e\u0448\u0438\u0431\u043e\u043a.<\/p>\n<\/li>\n<\/ul>\n<p>\u0412\u043e\u0442 \u043a\u0430\u043a \u0438\u043d\u0442\u0435\u0433\u0440\u0438\u0440\u043e\u0432\u0430\u0442\u044c \u043f\u043e\u0442\u043e\u043a\u0438 CUDA \u2014 \u043d\u0438\u0447\u0435\u0433\u043e \u0441\u0432\u0435\u0440\u0445\u044a\u0435\u0441\u0442\u0435\u0441\u0442\u0432\u0435\u043d\u043d\u043e\u0433\u043e:<\/p>\n<pre><code class=\"cpp\">#include &lt;cuda_runtime.h&gt; #include &lt;stdio.h&gt;  __global__ void kernel() {     extern __shared__ char shared_buffer[];     printf(\"Hello World from block %d, thread %d\\n\", blockIdx.x, threadIdx.x); }  int main() {     cudaStream_t stream;     cudaStreamCreate(&amp;stream);     uint shared_memory_size = 0;     kernel&lt;&lt;&lt;1, 1, shared_memory_size, stream&gt;&gt;&gt;(); \/\/ 4 \u0430\u0440\u0433\u0443\u043c\u0435\u043d\u0442\u0430, \u0430 \u043d\u0435 2     cudaStreamSynchronize(stream);     cudaStreamDestroy(stream);     return 0; }<\/code><\/pre>\n<p>\u041e\u0431\u0440\u0430\u0442\u0438\u0442\u0435 \u0432\u043d\u0438\u043c\u0430\u043d\u0438\u0435: \u044f\u0434\u0440\u043e \u0437\u0430\u043f\u0443\u0441\u043a\u0430\u0435\u0442\u0441\u044f \u0441 \u0447\u0435\u0442\u044b\u0440\u044c\u043c\u044f \u0430\u0440\u0433\u0443\u043c\u0435\u043d\u0442\u0430\u043c\u0438.<\/p>\n<p>\u0412\u043e\u0442 \u0431\u043e\u043b\u0435\u0435 \u0430\u043a\u043a\u0443\u0440\u0430\u0442\u043d\u0430\u044f \u0432\u0435\u0440\u0441\u0438\u044f \u0441 \u044f\u0432\u043d\u043e\u0439 \u043e\u0431\u0440\u0430\u0431\u043e\u0442\u043a\u043e\u0439 \u043e\u0448\u0438\u0431\u043e\u043a:<\/p>\n<pre><code class=\"cpp\">#include &lt;cuda_runtime.h&gt; #include &lt;stdio.h&gt;  __global__ void kernel() {     extern __shared__ char shared_buffer[];     printf(\"Hello World from block %d, thread %d\\n\", blockIdx.x, threadIdx.x); }  int main() {     cudaStream_t stream;     cudaError_t err = cudaStreamCreate(&amp;stream);     if (err != cudaSuccess) {         fprintf(stderr, \"Failed to create stream: %s\\n\", cudaGetErrorString(err));         return -1;     }     uint shared_memory_size = 1 &lt;&lt; 30; \/\/ 1 \u0413\u0411 \u2013 \u044d\u0442\u043e \u043c\u043d\u043e\u0433\u043e, \u043d\u043e \u0432 \u0434\u0435\u043c\u043e\u043d\u0441\u0442\u0440\u0430\u0446\u0438\u043e\u043d\u043d\u044b\u0445 \u0446\u0435\u043b\u044f\u0445 \u0442\u0430\u043a\u0430\u044f \u0432\u0435\u043b\u0438\u0447\u0438\u043d\u0430 \u043f\u043e\u0434\u043e\u0431\u0440\u0430\u043d\u0430 \u0441\u043f\u0435\u0446\u0438\u0430\u043b\u044c\u043d\u043e     kernel&lt;&lt;&lt;1, 1, shared_memory_size, stream&gt;&gt;&gt;();     err = cudaStreamSynchronize(stream);     if (err != cudaSuccess) {         fprintf(stderr, \"Failed to synchronize stream: %s\\n\", cudaGetErrorString(err));         cudaStreamDestroy(stream);         return -1;     }     cudaStreamDestroy(stream);     return 0; }<\/code><\/pre>\n<p>\u0418\u043c\u0435\u043d\u043d\u043e \u043d\u0430 \u044d\u0442\u043e\u043c, \u043a\u0430\u043a \u043f\u0440\u0430\u0432\u0438\u043b\u043e, \u0437\u0430\u043a\u0430\u043d\u0447\u0438\u0432\u0430\u044e\u0442\u0441\u044f \u0442\u0443\u0442\u043e\u0440\u0438\u0430\u043b\u044b \u2014 \u0438 \u043d\u0435\u0437\u0430\u043c\u0435\u0442\u043d\u043e \u043d\u0430\u0447\u0438\u043d\u0430\u044e\u0442\u0441\u044f \u043f\u0440\u043e\u0431\u043b\u0435\u043c\u044b. \u041f\u043e\u043f\u0440\u043e\u0431\u0443\u0439\u0442\u0435 \u0437\u0430\u043f\u0443\u0441\u0442\u0438\u0442\u044c \u0432\u043e\u0442 \u044d\u0442\u043e:<\/p>\n<pre><code>$ nvcc -o hello_world hello_world.cu &amp;&amp; .\/hello_world<\/code><\/pre>\n<p>\u041d\u0438\u043a\u0430\u043a\u043e\u0433\u043e \u0432\u044b\u0432\u043e\u0434\u0430, \u043d\u0438\u043a\u0430\u043a\u043e\u0433\u043e \u0441\u043e\u043e\u0431\u0449\u0435\u043d\u0438\u044f \u043e\u0431 \u043e\u0448\u0438\u0431\u043a\u0435, \u0432\u043e\u043e\u0431\u0449\u0435 \u043d\u0438\u0447\u0435\u0433\u043e. \u041c\u044b \u043d\u0435 \u043c\u043e\u0436\u0435\u043c \u0432\u044b\u0442\u044f\u043d\u0443\u0442\u044c \u043e\u0448\u0438\u0431\u043a\u0443 \u0438\u0437 \u043f\u043e\u0442\u043e\u043a\u0430, \u043f\u043e\u0441\u043a\u043e\u043b\u044c\u043a\u0443 \u043e\u0442\u043a\u0430\u0437 \u043f\u0440\u043e\u0438\u0441\u0445\u043e\u0434\u0438\u0442 \u043f\u0440\u0438 \u043e\u0442\u043f\u0440\u0430\u0432\u043a\u0435 \u0434\u0430\u043d\u043d\u044b\u0445, \u0430 \u043d\u0435 \u043f\u0440\u0438 \u0432\u044b\u043f\u043e\u043b\u043d\u0435\u043d\u0438\u0438.<\/p>\n<h2>API \u0441\u0440\u0435\u0434\u044b \u0432\u044b\u043f\u043e\u043b\u043d\u0435\u043d\u0438\u044f CUDA<\/h2>\n<p>\u0422\u0440\u043e\u0439\u043d\u044b\u0435 \u0443\u0433\u043b\u043e\u0432\u044b\u0435 \u0441\u043a\u043e\u0431\u043a\u0438 NVCC \u2014 \u044d\u0442\u043e \u0441\u0438\u043d\u0442\u0430\u043a\u0441\u0438\u0447\u0435\u0441\u043a\u0438\u0439 \u0441\u0430\u0445\u0430\u0440, \u043a\u043e\u0442\u043e\u0440\u044b\u043c \u043f\u0440\u0438\u0441\u044b\u043f\u0430\u043d API \u0441\u0440\u0435\u0434\u044b \u0432\u044b\u043f\u043e\u043b\u043d\u0435\u043d\u0438\u044f CUDA, \u0430 \u044d\u0442\u043e\u0442 \u0438\u043d\u0442\u0435\u0440\u0444\u0435\u0439\u0441, \u0432 \u0441\u0432\u043e\u044e \u043e\u0447\u0435\u0440\u0435\u0434\u044c, \u043e\u0431\u0451\u0440\u0442\u044b\u0432\u0430\u0435\u0442 \u0431\u043e\u043b\u0435\u0435 \u043d\u0438\u0437\u043a\u043e\u0443\u0440\u043e\u0432\u043d\u0435\u0432\u044b\u0439 API \u0434\u0440\u0430\u0439\u0432\u0435\u0440\u043e\u0432 CUDA. \u0414\u043b\u044f \u044d\u0444\u0444\u0435\u043a\u0442\u0438\u0432\u043d\u043e\u0433\u043e \u043e\u0442\u043b\u043e\u0432\u0430 \u043e\u0448\u0438\u0431\u043e\u043a \u043d\u0435\u043e\u0431\u0445\u043e\u0434\u0438\u043c\u043e \u0437\u0430\u0434\u0435\u0439\u0441\u0442\u0432\u043e\u0432\u0430\u0442\u044c <a href=\"https:\/\/docs.nvidia.com\/cuda\/cuda-runtime-api\/group__CUDART__EXECUTION.html\" rel=\"noopener noreferrer nofollow\">CUDA Driver Execution Control API<\/a>:<\/p>\n<pre><code class=\"cpp\">#include &lt;cuda_runtime.h&gt; #include &lt;stdio.h&gt;  __global__ void kernel() {     extern __shared__ char shared_buffer[];     printf(\"Hello World from block %d, thread %d\\n\", blockIdx.x, threadIdx.x); }  int main() {     cudaStream_t stream;     cudaError_t err;     err = cudaStreamCreate(&amp;stream);     if (err != cudaSuccess) {         fprintf(stderr, \"Failed to create stream: %s\\n\", cudaGetErrorString(err));         return -1;     }      dim3 grid(1);     dim3 block(1);     size_t shared_memory_size = 1 &lt;&lt; 30; \/\/ 1 \u0413\u0411     void *kernel_args[] = {};     err = cudaLaunchKernel((void *)kernel, grid, block, kernel_args, shared_memory_size, stream);     if (err != cudaSuccess) {         fprintf(stderr, \"Failed to launch kernel: %s\\n\", cudaGetErrorString(err));         cudaStreamDestroy(stream);         return -1;     }      err = cudaStreamSynchronize(stream);     if (err != cudaSuccess) {         fprintf(stderr, \"Kernel execution failed: %s\\n\", cudaGetErrorString(err));         cudaStreamDestroy(stream);         return -1;     }     err = cudaStreamDestroy(stream);     if (err != cudaSuccess) {         fprintf(stderr, \"Failed to destroy stream: %s\\n\", cudaGetErrorString(err));         return -1;     }     return 0; }<\/code><\/pre>\n<p>\u0421\u043a\u043e\u043c\u043f\u0438\u043b\u0438\u0440\u0443\u0435\u043c \u0438 \u0432\u044b\u043f\u043e\u043b\u043d\u0438\u043c:<\/p>\n<pre><code>$ nvcc -o hello_world hello_world.cu &amp;&amp; .\/hello_world &gt; Failed to launch kernel: invalid argument<\/code><\/pre>\n<p>\u042d\u0442\u0430 \u043e\u0448\u0438\u0431\u043a\u0430 \u043e\u0436\u0438\u0434\u0430\u0435\u043c\u0430, \u043f\u043e\u0441\u043a\u043e\u043b\u044c\u043a\u0443 \u044f\u0434\u0440\u043e \u043d\u0435\u0432\u043e\u0437\u043c\u043e\u0436\u043d\u043e \u0434\u0430\u0436\u0435 \u043e\u0442\u043f\u0440\u0430\u0432\u0438\u0442\u044c \u2014 \u0438\u0437-\u0437\u0430 \u0430\u0431\u0441\u0443\u0440\u0434\u043d\u043e\u0433\u043e \u0437\u0430\u043f\u0440\u043e\u0441\u0430 \u043a \u043f\u0430\u043c\u044f\u0442\u0438. \u041d\u043e \u043f\u0440\u043e\u0431\u043b\u0435\u043c\u0430 \u0441 \u044d\u0442\u0438\u043c API \u0432 \u0442\u043e\u043c, \u0447\u0442\u043e \u043d\u0430\u043c \u0442\u0440\u0435\u0431\u0443\u0435\u0442\u0441\u044f \u043d\u0430\u0439\u0442\u0438 \u0438\u043d\u043e\u0439 \u0441\u043f\u043e\u0441\u043e\u0431, \u0447\u0442\u043e\u0431\u044b \u043f\u0435\u0440\u0435\u0434\u0430\u0432\u0430\u0442\u044c \u0430\u0440\u0433\u0443\u043c\u0435\u043d\u0442\u044b \u044f\u0434\u0440\u0443. \u0412\u043e\u0442 \u043a\u0430\u043a \u043c\u044b \u043f\u0435\u0440\u0435\u0434\u0430\u0432\u0430\u043b\u0438 \u0431\u044b \u044f\u0434\u0440\u0443 \u043c\u0430\u0441\u0441\u0438\u0432\u044b \u0438 \u0441\u043a\u0430\u043b\u044f\u0440\u044b:<\/p>\n<pre><code class=\"cpp\">#include &lt;cuda_runtime.h&gt; #include &lt;stdio.h&gt;  __global__ void kernel(float *amount, size_t count, int power) {     size_t idx = blockIdx.x * blockDim.x + threadIdx.x;     if (idx &gt; count) return;     amount[idx] = amount[idx] * scalbln(1.0, power); \/\/ \u041f\u0440\u0438\u043c\u0435\u0440 \u0432\u0441\u0442\u0440\u043e\u0435\u043d\u043d\u043e\u0439 \u0444\u0443\u043d\u043a\u0446\u0438\u0438 CUDA ;) }  int main() {     cudaError_t err;     size_t num_elements = 1024;     int integral_power = -2;     double *data;      \/\/ \u0412\u044b\u0434\u0435\u043b\u044f\u0435\u043c \u043e\u0431\u044a\u0435\u0434\u0438\u043d\u0451\u043d\u043d\u0443\u044e \u043f\u0430\u043c\u044f\u0442\u044c     err = cudaMallocManaged(&amp;data, num_elements * sizeof(double));     if (err != cudaSuccess) {         fprintf(stderr, \"cudaMallocManaged failed: %s\\n\", cudaGetErrorString(err));         return -1;     }      \/\/ \u0418\u043d\u0438\u0446\u0438\u0430\u043b\u0438\u0437\u0438\u0440\u0443\u0435\u043c \u0434\u0430\u043d\u043d\u044b\u0435     for (size_t i = 0; i &lt; num_elements; ++i) data[i] = (double)i;      \/\/ \u0421\u043e\u0437\u0434\u0430\u0451\u043c \u043f\u043e\u0442\u043e\u043a CUDA      cudaStream_t stream;     err = cudaStreamCreate(&amp;stream);     if (err != cudaSuccess) {         fprintf(stderr, \"Failed to create stream: %s\\n\", cudaGetErrorString(err));         cudaFree(data);         return -1;     }      \/\/ \u041e\u043f\u0440\u0435\u0434\u0435\u043b\u044f\u0435\u043c \u043f\u0430\u0440\u0430\u043c\u0435\u0442\u0440\u044b \u043f\u0443\u0441\u043a\u0430 \u044f\u0434\u0440\u0430     dim3 grid((num_elements + 255) \/ 256);     dim3 block(256);     void *kernel_args[] = {         (void *)&amp;data,         (void *)&amp;num_elements,         (void *)&amp;integral_power,     };      \/\/ \u0417\u0430\u043f\u0443\u0441\u043a\u0430\u0435\u043c \u044f\u0434\u0440\u043e     err = cudaLaunchKernel((void *)kernel, grid, block, kernel_args, 0, stream);     if (err != cudaSuccess) {         fprintf(stderr, \"Failed to launch kernel: %s\\n\", cudaGetErrorString(err));         cudaStreamDestroy(stream);         cudaFree(data);         return -1;     }      \/\/ \u0421\u0438\u043d\u0445\u0440\u043e\u043d\u0438\u0437\u0438\u0440\u0443\u0435\u043c \u043f\u043e\u0442\u043e\u043a     err = cudaStreamSynchronize(stream);     if (err != cudaSuccess) {         fprintf(stderr, \"Kernel execution failed: %s\\n\", cudaGetErrorString(err));         cudaStreamDestroy(stream);         cudaFree(data);         return -1;     }      \/\/ \u0412\u044b\u0432\u043e\u0434\u0438\u043c \u0440\u0435\u0437\u0443\u043b\u044c\u0442\u0430\u0442\u044b \u043d\u0430 \u044d\u043a\u0440\u0430\u043d     for (size_t i = 0; i &lt; 5; ++i) printf(\"data[%zu] = %f\\n\", i, data[i]);     cudaStreamDestroy(stream);     cudaFree(data);     return 0; }<\/code><\/pre>\n<p>\u042f \u0432\u043e\u0441\u043f\u043e\u043b\u044c\u0437\u043e\u0432\u0430\u043b\u0441\u044f \u043e\u0431\u044a\u0435\u0434\u0438\u043d\u0451\u043d\u043d\u043e\u0439 \u043f\u0430\u043c\u044f\u0442\u044c\u044e, \u0447\u0442\u043e\u0431\u044b \u0443\u043f\u0440\u043e\u0441\u0442\u0438\u0442\u044c \u043f\u0440\u0438\u043c\u0435\u0440. \u041c\u044b \u043d\u0435 \u043e\u0431\u044f\u0437\u0430\u043d\u044b \u044f\u0432\u043d\u043e \u0432\u044b\u0434\u0435\u043b\u044f\u0442\u044c 2 \u0431\u0443\u0444\u0435\u0440\u0430 \u043e\u0434\u043d\u043e\u0432\u0440\u0435\u043c\u0435\u043d\u043d\u043e \u0432 \u0426\u041f \u0438 \u0433\u0440\u0430\u0444\u0438\u0447\u0435\u0441\u043a\u043e\u043c \u043f\u0440\u043e\u0446\u0435\u0441\u0441\u043e\u0440\u0435 \u0438 \u043a\u043e\u043f\u0438\u0440\u043e\u0432\u0430\u0442\u044c \u0434\u0430\u043d\u043d\u044b\u0435 \u043c\u0435\u0436\u0434\u0443 \u043d\u0438\u043c\u0438. \u0414\u0440\u0430\u0439\u0432\u0435\u0440 \u0434\u0435\u0440\u0436\u0438\u0442 \u043a\u043e\u043f\u0438\u0438 \u0434\u0430\u043d\u043d\u044b\u0445 \u043e\u0434\u043d\u043e\u0432\u0440\u0435\u043c\u0435\u043d\u043d\u043e \u043d\u0430 \u0445\u043e\u0441\u0442\u0435 \u0438 \u0432 \u043f\u0430\u043c\u044f\u0442\u0438 \u0443\u0441\u0442\u0440\u043e\u0439\u0441\u0442\u0432\u0430, \u0438 \u043f\u043e \u043c\u0435\u0440\u0435 \u043d\u0435\u043e\u0431\u0445\u043e\u0434\u0438\u043c\u043e\u0441\u0442\u0438 \u0430\u0432\u0442\u043e\u043c\u0430\u0442\u0438\u0447\u0435\u0441\u043a\u0438 \u043f\u0435\u0440\u0435\u0434\u0430\u0451\u0442 \u043e\u0431\u043d\u043e\u0432\u043b\u0435\u043d\u0438\u044f \u043c\u0435\u0436\u0434\u0443 \u043d\u0438\u043c\u0438.<\/p>\n<h2>\u041a\u043e\u043e\u043f\u0435\u0440\u0430\u0442\u0438\u0432\u043d\u044b\u0435 \u0433\u0440\u0443\u043f\u043f\u044b<\/h2>\n<p>\u0410\u0445, \u043a\u0430\u043a \u0431\u044b \u0445\u043e\u0442\u0435\u043b\u043e\u0441\u044c, \u0447\u0442\u043e\u0431\u044b \u043c\u043e\u0436\u043d\u043e \u0431\u044b\u043b\u043e \u043d\u0430\u043f\u0438\u0441\u0430\u0442\u044c \u043f\u0430\u0440\u0430\u043b\u043b\u0435\u043b\u044c\u043d\u044b\u0435 \u0430\u043b\u0433\u043e\u0440\u0438\u0442\u043c\u044b \u0434\u043b\u044f \u0433\u0440\u0430\u0444\u0438\u0447\u0435\u0441\u043a\u043e\u0433\u043e \u043f\u0440\u043e\u0446\u0435\u0441\u0441\u043e\u0440\u0430 \u0435\u0434\u0438\u043d\u043e\u0436\u0434\u044b \u2014 \u0441\u043e\u0431\u0440\u0430\u0442\u044c \u0441\u0442\u043e\u043f\u043a\u0443 \u0430\u0431\u0441\u0442\u0440\u0430\u043a\u0446\u0438\u0439, \u043e\u0431\u0435\u0440\u043d\u0443\u0442\u044c \u0438\u0445 \u0432 \u0448\u0430\u0431\u043b\u043e\u043d\u044b, \u0430 \u0434\u0430\u043b\u044c\u0448\u0435 \u0441\u0440\u0435\u0434\u0430 \u0432\u044b\u043f\u043e\u043b\u043d\u0435\u043d\u0438\u044f \u043f\u0443\u0441\u0442\u044c \u0441\u0430\u043c\u0430 \u0440\u0430\u0437\u0431\u0438\u0440\u0430\u0435\u0442\u0441\u044f. \u041d\u043e \u043d\u0430 \u043f\u0440\u0430\u043a\u0442\u0438\u043a\u0435 \u0442\u0430\u043a\u043e\u0435 \u0443\u0434\u0430\u0451\u0442\u0441\u044f \u0440\u0435\u0434\u043a\u043e. \u041a \u0441\u043e\u0436\u0430\u043b\u0435\u043d\u0438\u044e, \u0438 API CUDA \u0434\u043b\u044f \u043e\u0431\u0440\u0430\u0449\u0435\u043d\u0438\u044f \u0441 \u043a\u043e\u043e\u043f\u0435\u0440\u0430\u0442\u0438\u0432\u043d\u044b\u043c\u0438 \u0433\u0440\u0443\u043f\u043f\u0430\u043c\u0438 \u2014 \u043d\u0435 \u0438\u0441\u043a\u043b\u044e\u0447\u0435\u043d\u0438\u0435.<\/p>\n<p>\u041e\u043d \u043f\u0440\u043e\u0435\u043a\u0442\u0438\u0440\u043e\u0432\u0430\u043b\u0441\u044f \u043a\u0430\u043a \u0435\u0434\u0438\u043d\u043e\u043e\u0431\u0440\u0430\u0437\u043d\u0430\u044f \u0430\u0431\u0441\u0442\u0440\u0430\u043a\u0446\u0438\u044f \u0434\u043b\u044f \u043a\u043e\u043e\u0440\u0434\u0438\u043d\u0430\u0446\u0438\u0438 \u043f\u043e\u0442\u043e\u043a\u043e\u0432, \u043d\u0435 \u043e\u0433\u0440\u0430\u043d\u0438\u0447\u0435\u043d\u043d\u0430\u044f \u043f\u0440\u0435\u0434\u0435\u043b\u0430\u043c\u0438 \u043e\u0434\u043d\u043e\u0433\u043e \u0431\u043b\u043e\u043a\u0430. \u041f\u0440\u0438 \u044d\u0442\u043e\u043c \u0432\u0441\u0442\u0440\u043e\u0435\u043d\u043d\u044b\u0435 \u0444\u0443\u043d\u043a\u0446\u0438\u0438 C++ \u043f\u0440\u0438\u043c\u0435\u043d\u044f\u044e\u0442\u0441\u044f \u0434\u043b\u044f \u043d\u0430\u0437\u043d\u0430\u0447\u0435\u043d\u0438\u044f \u0441\u043b\u043e\u0436\u043d\u044b\u0445 \u0430\u043b\u0433\u043e\u0440\u0438\u0442\u043c\u043e\u0432 GPU, \u0438 \u0437\u0430\u0434\u0435\u0439\u0441\u0442\u0432\u0443\u0435\u043c\u0430\u044f \u0441 \u043d\u0438\u043c\u0438 \u0441\u0435\u043c\u0430\u043d\u0442\u0438\u043a\u0430 \u0441\u0438\u043d\u0445\u0440\u043e\u043d\u0438\u0437\u0430\u0446\u0438\u0438 \u0434\u043e\u0441\u0442\u0430\u0442\u043e\u0447\u043d\u043e \u0433\u0438\u0431\u043a\u0430\u044f. \u0422\u0435\u043e\u0440\u0435\u0442\u0438\u0447\u0435\u0441\u043a\u0438, \u0442\u0430\u043a \u0434\u043e\u043b\u0436\u043d\u0430 \u0440\u0435\u0448\u0430\u0442\u044c\u0441\u044f \u0441\u0435\u0440\u044c\u0451\u0437\u043d\u0430\u044f \u043f\u0440\u043e\u0431\u043b\u0435\u043c\u0430: \u043f\u0443\u0441\u0442\u044c \u0432\u0441\u0435 \u043f\u043e\u0442\u043e\u043a\u0438 \u043d\u0430 \u0443\u0441\u0442\u0440\u043e\u0439\u0441\u0442\u0432\u0435 \u0432\u0440\u0435\u043c\u044f \u043e\u0442 \u0432\u0440\u0435\u043c\u0435\u043d\u0438 \u0441\u0438\u043d\u0445\u0440\u043e\u043d\u0438\u0437\u0438\u0440\u0443\u044e\u0442\u0441\u044f \u043f\u0435\u0440\u0435\u0434 \u043f\u0440\u043e\u0434\u043e\u043b\u0436\u0435\u043d\u0438\u0435\u043c \u0440\u0430\u0431\u043e\u0442\u044b \u2014 \u044d\u0442\u043e \u0447\u0440\u0435\u0437\u0432\u044b\u0447\u0430\u0439\u043d\u043e \u0432\u0430\u0436\u043d\u043e \u0434\u043b\u044f \u0438\u0442\u0435\u0440\u0430\u0442\u0438\u0432\u043d\u044b\u0445 \u0430\u043b\u0433\u043e\u0440\u0438\u0442\u043c\u043e\u0432, \u043d\u0430 \u043a\u043e\u0442\u043e\u0440\u044b\u0445 \u043e\u0441\u043d\u043e\u0432\u0430\u043d\u044b, \u043d\u0430\u043f\u0440\u0438\u043c\u0435\u0440, \u0441\u0438\u043c\u0443\u043b\u044f\u0446\u0438\u0438 \u0444\u0438\u0437\u0438\u0447\u0435\u0441\u043a\u0438\u0445 \u043f\u0440\u043e\u0446\u0435\u0441\u0441\u043e\u0432 \u0438\u043b\u0438 \u0440\u0435\u0448\u0430\u043b\u043a\u0438.<\/p>\n<p>\u041f\u0440\u043e\u0441\u0442\u043e \u043d\u0430\u043f\u043e\u043c\u043d\u044e:<\/p>\n<ul>\n<li>\n<p> <code>__syncwarp()<\/code>\u00a0\u043e\u0431\u0451\u0440\u0442\u044b\u0432\u0430\u0435\u0442 32 \u043f\u043e\u0442\u043e\u043a\u0430.<\/p>\n<\/li>\n<li>\n<p> <code>__syncthreads()\u00a0<\/code>\u043e\u0431\u0451\u0440\u0442\u044b\u0432\u0430\u0435\u0442 \u043b\u043e\u0433\u0438\u0447\u0435\u0441\u043a\u0438\u0439 \u0431\u043b\u043e\u043a, \u0432\u043a\u043b\u044e\u0447\u0430\u044e\u0449\u0438\u0439 1-1024 \u043f\u043e\u0442\u043e\u043a\u043e\u0432.<\/p>\n<\/li>\n<li>\n<p>\u041d\u0430 \u0432\u0441\u0435 \u043f\u0440\u043e\u0447\u0438\u0435 \u0441\u043b\u0443\u0447\u0430\u0438 \u0435\u0441\u0442\u044c\u00a0<s>Mastercard<\/s>\u043a\u043e\u043e\u043f\u0435\u0440\u0430\u0442\u0438\u0432\u043d\u044b\u0435 \u0433\u0440\u0443\u043f\u043f\u044b.<\/p>\n<\/li>\n<\/ul>\n<p>\u0421\u0438\u043b\u044c\u043d\u0435\u0435 \u0432\u0441\u0435\u0433\u043e \u043d\u0430\u043f\u0440\u0430\u0448\u0438\u0432\u0430\u0435\u0442\u0441\u044f \u0442\u0430\u043a\u043e\u0439 \u043f\u0440\u0438\u043c\u0435\u0440: \u0441\u0438\u043d\u0445\u0440\u043e\u043d\u0438\u0437\u0438\u0440\u0443\u0435\u043c \u0432\u0435\u0441\u044c \u0433\u0440\u0438\u0434 \u0432 \u0432\u0438\u0434\u0435 \u043c\u043d\u043e\u0433\u043e\u044d\u0442\u0430\u043f\u043d\u044b\u0445 \u0438\u0442\u0435\u0440\u0430\u0442\u0438\u0432\u043d\u044b\u0445 \u0430\u043b\u0433\u043e\u0440\u0438\u0442\u043c\u043e\u0432, \u0441\u043a\u0430\u0436\u0435\u043c, \u043f\u0440\u0438 \u0441\u0438\u043c\u0443\u043b\u044f\u0446\u0438\u0438 \u0444\u0438\u0437\u0438\u0447\u0435\u0441\u043a\u043e\u0433\u043e \u043f\u0440\u043e\u0446\u0435\u0441\u0441\u0430. \u0414\u043b\u044f \u044d\u0442\u043e\u0433\u043e Nvidia \u0440\u0435\u043a\u043e\u043c\u0435\u043d\u0434\u0443\u0435\u0442 \u0437\u0430\u0434\u0435\u0439\u0441\u0442\u0432\u043e\u0432\u0430\u0442\u044c \u043d\u043e\u0432\u0443\u044e \u0444\u0443\u043d\u043a\u0446\u0438\u044e <code>cooperative_groups::sync()<\/code>:<\/p>\n<pre><code class=\"cpp\">#include &lt;cuda_runtime.h&gt; #include &lt;cooperative_groups.h&gt; #include &lt;stdio.h&gt; #include &lt;math.h&gt;  namespace cg = cooperative_groups;  __device__ float3 compute_force(float3 position_first, float3 position_second) {     float3 r;     r.x = position_second.x - position_first.x;     r.y = position_second.y - position_first.y;     r.z = position_second.z - position_first.z;      float squared_distance = r.x * r.x + r.y * r.y + r.z * r.z + 1e-6f; \/\/ avoid div by zero     float reciprocal_distance = rsqrtf(squared_distance);<\/code><\/pre>\n<\/div>\n<\/div>\n<\/div>\n<\/div>\n","protected":false},"author":1,"featured_media":0,"comment_status":"open","ping_status":"open","sticky":false,"template":"","format":"standard","meta":{"footnotes":""},"categories":[],"tags":[],"class_list":["post-470994","post","type-post","status-publish","format-standard","hentry"],"_links":{"self":[{"href":"https:\/\/savepearlharbor.com\/index.php?rest_route=\/wp\/v2\/posts\/470994","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=470994"}],"version-history":[{"count":0,"href":"https:\/\/savepearlharbor.com\/index.php?rest_route=\/wp\/v2\/posts\/470994\/revisions"}],"wp:attachment":[{"href":"https:\/\/savepearlharbor.com\/index.php?rest_route=%2Fwp%2Fv2%2Fmedia&parent=470994"}],"wp:term":[{"taxonomy":"category","embeddable":true,"href":"https:\/\/savepearlharbor.com\/index.php?rest_route=%2Fwp%2Fv2%2Fcategories&post=470994"},{"taxonomy":"post_tag","embeddable":true,"href":"https:\/\/savepearlharbor.com\/index.php?rest_route=%2Fwp%2Fv2%2Ftags&post=470994"}],"curies":[{"name":"wp","href":"https:\/\/api.w.org\/{rel}","templated":true}]}}