| <!DOCTYPE html> |
| |
| <html lang="en"> |
| <head> |
| <meta charset="utf-8"/> |
| <meta content="IE=edge" http-equiv="X-UA-Compatible"/> |
| <meta content="width=device-width, initial-scale=1" name="viewport"/> |
| <meta content="mxnet.rtc" property="og:title"> |
| <meta content="https://raw.githubusercontent.com/dmlc/web-data/master/mxnet/image/og-logo.png" property="og:image"> |
| <meta content="https://raw.githubusercontent.com/dmlc/web-data/master/mxnet/image/og-logo.png" property="og:image:secure_url"> |
| <meta content="mxnet.rtc" property="og:description"/> |
| <title>mxnet.rtc — mxnet documentation</title> |
| <link crossorigin="anonymous" href="https://maxcdn.bootstrapcdn.com/bootstrap/3.3.6/css/bootstrap.min.css" integrity="sha384-1q8mTJOASx8j1Au+a5WDVnPi2lkFfwwEAa8hDDdjZlpLegxhjVME1fgjWPGmkzs7" rel="stylesheet"/> |
| <link href="https://maxcdn.bootstrapcdn.com/font-awesome/4.5.0/css/font-awesome.min.css" rel="stylesheet"/> |
| <link href="../../_static/basic.css" rel="stylesheet" type="text/css"> |
| <link href="../../_static/pygments.css" rel="stylesheet" type="text/css"> |
| <link href="../../_static/mxnet.css" rel="stylesheet" type="text/css"/> |
| <script type="text/javascript"> |
| var DOCUMENTATION_OPTIONS = { |
| URL_ROOT: '../../', |
| VERSION: '', |
| COLLAPSE_INDEX: false, |
| FILE_SUFFIX: '.html', |
| HAS_SOURCE: true, |
| SOURCELINK_SUFFIX: '.txt' |
| }; |
| </script> |
| <script src="https://code.jquery.com/jquery-1.11.1.min.js" type="text/javascript"></script> |
| <script src="../../_static/underscore.js" type="text/javascript"></script> |
| <script src="../../_static/searchtools_custom.js" type="text/javascript"></script> |
| <script src="../../_static/doctools.js" type="text/javascript"></script> |
| <script src="../../_static/selectlang.js" type="text/javascript"></script> |
| <script src="https://cdnjs.cloudflare.com/ajax/libs/mathjax/2.7.1/MathJax.js?config=TeX-AMS-MML_HTMLorMML" type="text/javascript"></script> |
| <script type="text/javascript"> jQuery(function() { Search.loadIndex("/versions/1.5.0/searchindex.js"); Search.init();}); </script> |
| <script> |
| (function(i,s,o,g,r,a,m){i['GoogleAnalyticsObject']=r;i[r]=i[r]||function(){ |
| (i[r].q=i[r].q||[]).push(arguments)},i[r].l=1*new |
| Date();a=s.createElement(o), |
| m=s.getElementsByTagName(o)[0];a.async=1;a.src=g;m.parentNode.insertBefore(a,m) |
| })(window,document,'script','https://www.google-analytics.com/analytics.js','ga'); |
| |
| ga('create', 'UA-96378503-1', 'auto'); |
| ga('send', 'pageview'); |
| |
| </script> |
| <!-- --> |
| <!-- <script type="text/javascript" src="../../_static/jquery.js"></script> --> |
| <!-- --> |
| <!-- <script type="text/javascript" src="../../_static/underscore.js"></script> --> |
| <!-- --> |
| <!-- <script type="text/javascript" src="../../_static/doctools.js"></script> --> |
| <!-- --> |
| <!-- <script type="text/javascript" src="https://cdnjs.cloudflare.com/ajax/libs/mathjax/2.7.0/MathJax.js?config=TeX-AMS-MML_HTMLorMML"></script> --> |
| <!-- --> |
| <link href="../../genindex.html" rel="index" title="Index"> |
| <link href="../../search.html" rel="search" title="Search"/> |
| <link href="../index.html" rel="up" title="Module code"/> |
| <link href="https://raw.githubusercontent.com/dmlc/web-data/master/mxnet/image/mxnet-icon.png" rel="icon" type="image/png"/> |
| </link></link></link></meta></meta></meta></head> |
| <body background="https://raw.githubusercontent.com/dmlc/web-data/master/mxnet/image/mxnet-background-compressed.jpeg" role="document"> |
| <div class="content-block"><div class="navbar navbar-fixed-top"> |
| <div class="container" id="navContainer"> |
| <div class="innder" id="header-inner"> |
| <h1 id="logo-wrap"> |
| <a href="../../" id="logo"><img src="https://raw.githubusercontent.com/dmlc/web-data/master/mxnet/image/mxnet_logo.png"/></a> |
| </h1> |
| <nav class="nav-bar" id="main-nav"> |
| <a class="main-nav-link" href="/versions/1.5.0/install/index.html">Install</a> |
| <span id="dropdown-menu-position-anchor"> |
| <a aria-expanded="true" aria-haspopup="true" class="main-nav-link dropdown-toggle" data-toggle="dropdown" href="#" role="button">Gluon <span class="caret"></span></a> |
| <ul class="dropdown-menu navbar-menu" id="package-dropdown-menu"> |
| <li><a class="main-nav-link" href="/versions/1.5.0/tutorials/gluon/gluon.html">About</a></li> |
| <li><a class="main-nav-link" href="https://www.d2l.ai/">Dive into Deep Learning</a></li> |
| <li><a class="main-nav-link" href="https://gluon-cv.mxnet.io">GluonCV Toolkit</a></li> |
| <li><a class="main-nav-link" href="https://gluon-nlp.mxnet.io/">GluonNLP Toolkit</a></li> |
| </ul> |
| </span> |
| <span id="dropdown-menu-position-anchor"> |
| <a aria-expanded="true" aria-haspopup="true" class="main-nav-link dropdown-toggle" data-toggle="dropdown" href="#" role="button">API <span class="caret"></span></a> |
| <ul class="dropdown-menu navbar-menu" id="package-dropdown-menu"> |
| <li><a class="main-nav-link" href="/versions/1.5.0/api/python/index.html">Python</a></li> |
| <li><a class="main-nav-link" href="/versions/1.5.0/api/c++/index.html">C++</a></li> |
| <li><a class="main-nav-link" href="/versions/1.5.0/api/clojure/index.html">Clojure</a></li> |
| <li><a class="main-nav-link" href="/versions/1.5.0/api/java/index.html">Java</a></li> |
| <li><a class="main-nav-link" href="/versions/1.5.0/api/julia/index.html">Julia</a></li> |
| <li><a class="main-nav-link" href="/versions/1.5.0/api/perl/index.html">Perl</a></li> |
| <li><a class="main-nav-link" href="/versions/1.5.0/api/r/index.html">R</a></li> |
| <li><a class="main-nav-link" href="/versions/1.5.0/api/scala/index.html">Scala</a></li> |
| </ul> |
| </span> |
| <span id="dropdown-menu-position-anchor-docs"> |
| <a aria-expanded="true" aria-haspopup="true" class="main-nav-link dropdown-toggle" data-toggle="dropdown" href="#" role="button">Docs <span class="caret"></span></a> |
| <ul class="dropdown-menu navbar-menu" id="package-dropdown-menu-docs"> |
| <li><a class="main-nav-link" href="/versions/1.5.0/faq/index.html">FAQ</a></li> |
| <li><a class="main-nav-link" href="/versions/1.5.0/tutorials/index.html">Tutorials</a> |
| <li><a class="main-nav-link" href="https://github.com/apache/incubator-mxnet/tree/1.5.0/example">Examples</a></li> |
| <li><a class="main-nav-link" href="/versions/1.5.0/architecture/index.html">Architecture</a></li> |
| <li><a class="main-nav-link" href="https://cwiki.apache.org/confluence/display/MXNET/Apache+MXNet+Home">Developer Wiki</a></li> |
| <li><a class="main-nav-link" href="/versions/1.5.0/model_zoo/index.html">Model Zoo</a></li> |
| <li><a class="main-nav-link" href="https://github.com/onnx/onnx-mxnet">ONNX</a></li> |
| </li></ul> |
| </span> |
| <span id="dropdown-menu-position-anchor-community"> |
| <a aria-expanded="true" aria-haspopup="true" class="main-nav-link dropdown-toggle" data-toggle="dropdown" href="#" role="button">Community <span class="caret"></span></a> |
| <ul class="dropdown-menu navbar-menu" id="package-dropdown-menu-community"> |
| <li><a class="main-nav-link" href="http://discuss.mxnet.io">Forum</a></li> |
| <li><a class="main-nav-link" href="https://github.com/apache/incubator-mxnet/tree/1.5.0">Github</a></li> |
| <li><a class="main-nav-link" href="/versions/1.5.0/community/contribute.html">Contribute</a></li> |
| <li><a class="main-nav-link" href="/versions/1.5.0/community/ecosystem.html">Ecosystem</a></li> |
| <li><a class="main-nav-link" href="/versions/1.5.0/community/powered_by.html">Powered By</a></li> |
| </ul> |
| </span> |
| <span id="dropdown-menu-position-anchor-version" style="position: relative"><a href="#" class="main-nav-link dropdown-toggle" data-toggle="dropdown" role="button" aria-haspopup="true" aria-expanded="true">1.5.0<span class="caret"></span></a><ul id="package-dropdown-menu" class="dropdown-menu"><li><a href="/">master</a></li><li><a href=/versions/1.6/>1.6</a></li><li><a href=/versions/1.5.0/>1.5.0</a></li><li><a href=/versions/1.4.1/>1.4.1</a></li><li><a href=/versions/1.3.1/>1.3.1</a></li><li><a href=/versions/1.2.1/>1.2.1</a></li><li><a href=/versions/1.1.0/>1.1.0</a></li><li><a href=/versions/1.0.0/>1.0.0</a></li><li><a href=/versions/0.12.1/>0.12.1</a></li><li><a href=/versions/0.11.0/>0.11.0</a></li></ul></span></nav> |
| <script> function getRootPath(){ return "../../" } </script> |
| <div class="burgerIcon dropdown"> |
| <a class="dropdown-toggle" data-toggle="dropdown" href="#" role="button">☰</a> |
| <ul class="dropdown-menu" id="burgerMenu"> |
| <li><a href="/versions/1.5.0/install/index.html">Install</a></li> |
| <li><a class="main-nav-link" href="/versions/1.5.0/tutorials/index.html">Tutorials</a></li> |
| <li class="dropdown-submenu dropdown"> |
| <a aria-expanded="true" aria-haspopup="true" class="dropdown-toggle burger-link" data-toggle="dropdown" href="#" tabindex="-1">Gluon</a> |
| <ul class="dropdown-menu navbar-menu" id="package-dropdown-menu"> |
| <li><a class="main-nav-link" href="/versions/1.5.0/tutorials/gluon/gluon.html">About</a></li> |
| <li><a class="main-nav-link" href="http://gluon.mxnet.io">The Straight Dope (Tutorials)</a></li> |
| <li><a class="main-nav-link" href="https://gluon-cv.mxnet.io">GluonCV Toolkit</a></li> |
| <li><a class="main-nav-link" href="https://gluon-nlp.mxnet.io/">GluonNLP Toolkit</a></li> |
| </ul> |
| </li> |
| <li class="dropdown-submenu"> |
| <a aria-expanded="true" aria-haspopup="true" class="dropdown-toggle burger-link" data-toggle="dropdown" href="#" tabindex="-1">API</a> |
| <ul class="dropdown-menu"> |
| <li><a class="main-nav-link" href="/versions/1.5.0/api/python/index.html">Python</a></li> |
| <li><a class="main-nav-link" href="/versions/1.5.0/api/c++/index.html">C++</a></li> |
| <li><a class="main-nav-link" href="/versions/1.5.0/api/clojure/index.html">Clojure</a></li> |
| <li><a class="main-nav-link" href="/versions/1.5.0/api/java/index.html">Java</a></li> |
| <li><a class="main-nav-link" href="/versions/1.5.0/api/julia/index.html">Julia</a></li> |
| <li><a class="main-nav-link" href="/versions/1.5.0/api/perl/index.html">Perl</a></li> |
| <li><a class="main-nav-link" href="/versions/1.5.0/api/r/index.html">R</a></li> |
| <li><a class="main-nav-link" href="/versions/1.5.0/api/scala/index.html">Scala</a></li> |
| </ul> |
| </li> |
| <li class="dropdown-submenu"> |
| <a aria-expanded="true" aria-haspopup="true" class="dropdown-toggle burger-link" data-toggle="dropdown" href="#" tabindex="-1">Docs</a> |
| <ul class="dropdown-menu"> |
| <li><a href="/versions/1.5.0/faq/index.html" tabindex="-1">FAQ</a></li> |
| <li><a href="/versions/1.5.0/tutorials/index.html" tabindex="-1">Tutorials</a></li> |
| <li><a href="https://github.com/apache/incubator-mxnet/tree/1.5.0/example" tabindex="-1">Examples</a></li> |
| <li><a href="/versions/1.5.0/architecture/index.html" tabindex="-1">Architecture</a></li> |
| <li><a href="https://cwiki.apache.org/confluence/display/MXNET/Apache+MXNet+Home" tabindex="-1">Developer Wiki</a></li> |
| <li><a href="/versions/1.5.0/model_zoo/index.html" tabindex="-1">Gluon Model Zoo</a></li> |
| <li><a href="https://github.com/onnx/onnx-mxnet" tabindex="-1">ONNX</a></li> |
| </ul> |
| </li> |
| <li class="dropdown-submenu dropdown"> |
| <a aria-haspopup="true" class="dropdown-toggle burger-link" data-toggle="dropdown" href="#" role="button" tabindex="-1">Community</a> |
| <ul class="dropdown-menu"> |
| <li><a href="http://discuss.mxnet.io" tabindex="-1">Forum</a></li> |
| <li><a href="https://github.com/apache/incubator-mxnet/tree/1.5.0" tabindex="-1">Github</a></li> |
| <li><a href="/versions/1.5.0/community/contribute.html" tabindex="-1">Contribute</a></li> |
| <li><a href="/versions/1.5.0/community/ecosystem.html" tabindex="-1">Ecosystem</a></li> |
| <li><a href="/versions/1.5.0/community/powered_by.html" tabindex="-1">Powered By</a></li> |
| </ul> |
| </li> |
| <li id="dropdown-menu-position-anchor-version-mobile" class="dropdown-submenu" style="position: relative"><a href="#" tabindex="-1">1.5.0</a><ul class="dropdown-menu"><li><a tabindex="-1" href=/>master</a></li><li><a tabindex="-1" href=/versions/1.6/>1.6</a></li><li><a tabindex="-1" href=/versions/1.5.0/>1.5.0</a></li><li><a tabindex="-1" href=/versions/1.4.1/>1.4.1</a></li><li><a tabindex="-1" href=/versions/1.3.1/>1.3.1</a></li><li><a tabindex="-1" href=/versions/1.2.1/>1.2.1</a></li><li><a tabindex="-1" href=/versions/1.1.0/>1.1.0</a></li><li><a tabindex="-1" href=/versions/1.0.0/>1.0.0</a></li><li><a tabindex="-1" href=/versions/0.12.1/>0.12.1</a></li><li><a tabindex="-1" href=/versions/0.11.0/>0.11.0</a></li></ul></li></ul> |
| </div> |
| <div class="plusIcon dropdown"> |
| <a class="dropdown-toggle" data-toggle="dropdown" href="#" role="button"><span aria-hidden="true" class="glyphicon glyphicon-plus"></span></a> |
| <ul class="dropdown-menu dropdown-menu-right" id="plusMenu"></ul> |
| </div> |
| <div id="search-input-wrap"> |
| <form action="../../search.html" autocomplete="off" class="" method="get" role="search"> |
| <div class="form-group inner-addon left-addon"> |
| <i class="glyphicon glyphicon-search"></i> |
| <input class="form-control" name="q" placeholder="Search" type="text"/> |
| </div> |
| <input name="check_keywords" type="hidden" value="yes"> |
| <input name="area" type="hidden" value="default"/> |
| </input></form> |
| <div id="search-preview"></div> |
| </div> |
| <div id="searchIcon"> |
| <span aria-hidden="true" class="glyphicon glyphicon-search"></span> |
| </div> |
| <!-- <div id="lang-select-wrap"> --> |
| <!-- <label id="lang-select-label"> --> |
| <!-- <\!-- <i class="fa fa-globe"></i> -\-> --> |
| <!-- <span></span> --> |
| <!-- </label> --> |
| <!-- <select id="lang-select"> --> |
| <!-- <option value="en">Eng</option> --> |
| <!-- <option value="zh">中文</option> --> |
| <!-- </select> --> |
| <!-- </div> --> |
| <!-- <a id="mobile-nav-toggle"> |
| <span class="mobile-nav-toggle-bar"></span> |
| <span class="mobile-nav-toggle-bar"></span> |
| <span class="mobile-nav-toggle-bar"></span> |
| </a> --> |
| </div> |
| </div> |
| </div> |
| <script type="text/javascript"> |
| $('body').css('background', 'white'); |
| </script> |
| <div class="container"> |
| <div class="row"> |
| <div aria-label="main navigation" class="sphinxsidebar leftsidebar" role="navigation"> |
| <div class="sphinxsidebarwrapper"> |
| <ul> |
| <li class="toctree-l1"><a class="reference internal" href="../../api/index.html">MXNet APIs</a></li> |
| <li class="toctree-l1"><a class="reference internal" href="../../architecture/index.html">MXNet Architecture</a></li> |
| <li class="toctree-l1"><a class="reference internal" href="../../community/index.html">MXNet Community</a></li> |
| <li class="toctree-l1"><a class="reference internal" href="../../faq/index.html">MXNet FAQ</a></li> |
| <li class="toctree-l1"><a class="reference internal" href="../../gluon/index.html">About Gluon</a></li> |
| <li class="toctree-l1"><a class="reference internal" href="../../install/index.html">Installing MXNet</a></li> |
| <li class="toctree-l1"><a class="reference internal" href="../../install/index.html#nvidia-jetson-tx-family">Nvidia Jetson TX family</a></li> |
| <li class="toctree-l1"><a class="reference internal" href="../../install/index.html#source-download">Source Download</a></li> |
| <li class="toctree-l1"><a class="reference internal" href="../../model_zoo/index.html">MXNet Model Zoo</a></li> |
| <li class="toctree-l1"><a class="reference internal" href="../../tutorials/index.html">Tutorials</a></li> |
| </ul> |
| </div> |
| </div> |
| <div class="content"> |
| <div class="page-tracker"></div> |
| <h1>Source code for mxnet.rtc</h1><div class="highlight"><pre> |
| <span></span><span class="c1"># Licensed to the Apache Software Foundation (ASF) under one</span> |
| <span class="c1"># or more contributor license agreements. See the NOTICE file</span> |
| <span class="c1"># distributed with this work for additional information</span> |
| <span class="c1"># regarding copyright ownership. The ASF licenses this file</span> |
| <span class="c1"># to you under the Apache License, Version 2.0 (the</span> |
| <span class="c1"># "License"); you may not use this file except in compliance</span> |
| <span class="c1"># with the License. You may obtain a copy of the License at</span> |
| <span class="c1">#</span> |
| <span class="c1"># http://www.apache.org/licenses/LICENSE-2.0</span> |
| <span class="c1">#</span> |
| <span class="c1"># Unless required by applicable law or agreed to in writing,</span> |
| <span class="c1"># software distributed under the License is distributed on an</span> |
| <span class="c1"># "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY</span> |
| <span class="c1"># KIND, either express or implied. See the License for the</span> |
| <span class="c1"># specific language governing permissions and limitations</span> |
| <span class="c1"># under the License.</span> |
| |
| <span class="sd">"""Interface to runtime cuda kernel compile module."""</span> |
| <span class="kn">from</span> <span class="nn">__future__</span> <span class="k">import</span> <span class="n">absolute_import</span> |
| |
| <span class="kn">from</span> <span class="nn">array</span> <span class="k">import</span> <span class="n">array</span> |
| <span class="kn">import</span> <span class="nn">re</span> |
| <span class="kn">import</span> <span class="nn">ctypes</span> |
| <span class="kn">import</span> <span class="nn">numpy</span> <span class="k">as</span> <span class="nn">np</span> |
| |
| <span class="kn">from</span> <span class="nn">.base</span> <span class="k">import</span> <span class="n">_LIB</span><span class="p">,</span> <span class="n">mx_uint</span><span class="p">,</span> <span class="n">c_array</span><span class="p">,</span> <span class="n">c_array_buf</span><span class="p">,</span> <span class="n">c_str_array</span><span class="p">,</span> <span class="n">check_call</span> |
| <span class="kn">from</span> <span class="nn">.base</span> <span class="k">import</span> <span class="n">c_str</span><span class="p">,</span> <span class="n">CudaModuleHandle</span><span class="p">,</span> <span class="n">CudaKernelHandle</span><span class="p">,</span> <span class="n">numeric_types</span><span class="p">,</span> <span class="n">string_types</span> |
| <span class="kn">from</span> <span class="nn">.ndarray</span> <span class="k">import</span> <span class="n">_DTYPE_NP_TO_MX</span><span class="p">,</span> <span class="n">_DTYPE_MX_TO_NP</span><span class="p">,</span> <span class="n">NDArray</span> |
| |
| <span class="n">_DTYPE_CPP_TO_NP</span> <span class="o">=</span> <span class="p">{</span> |
| <span class="s1">'float'</span><span class="p">:</span> <span class="n">np</span><span class="o">.</span><span class="n">float32</span><span class="p">,</span> |
| <span class="s1">'double'</span><span class="p">:</span> <span class="n">np</span><span class="o">.</span><span class="n">float64</span><span class="p">,</span> |
| <span class="s1">'__half'</span><span class="p">:</span> <span class="n">np</span><span class="o">.</span><span class="n">float16</span><span class="p">,</span> |
| <span class="s1">'uint8_t'</span><span class="p">:</span> <span class="n">np</span><span class="o">.</span><span class="n">uint8</span><span class="p">,</span> |
| <span class="s1">'int'</span><span class="p">:</span> <span class="n">np</span><span class="o">.</span><span class="n">int32</span><span class="p">,</span> |
| <span class="s1">'int32_t'</span><span class="p">:</span> <span class="n">np</span><span class="o">.</span><span class="n">int32</span><span class="p">,</span> |
| <span class="s1">'int8_t'</span><span class="p">:</span> <span class="n">np</span><span class="o">.</span><span class="n">int8</span><span class="p">,</span> |
| <span class="s1">'char'</span><span class="p">:</span> <span class="n">np</span><span class="o">.</span><span class="n">int8</span><span class="p">,</span> |
| <span class="s1">'int64_t'</span><span class="p">:</span> <span class="n">np</span><span class="o">.</span><span class="n">int64</span><span class="p">,</span> |
| <span class="p">}</span> |
| |
| <div class="viewcode-block" id="CudaModule"><a class="viewcode-back" href="../../api/python/rtc/rtc.html#mxnet.rtc.CudaModule">[docs]</a><span class="k">class</span> <span class="nc">CudaModule</span><span class="p">(</span><span class="nb">object</span><span class="p">):</span> |
| <span class="sa">r</span><span class="sd">"""Compile and run CUDA code from Python.</span> |
| |
| <span class="sd"> In CUDA 7.5, you need to prepend your kernel definitions</span> |
| <span class="sd"> with 'extern "C"' to avoid name mangling::</span> |
| |
| <span class="sd"> source = r'''</span> |
| <span class="sd"> extern "C" __global__ void axpy(const float *x, float *y, float alpha) {</span> |
| <span class="sd"> int i = threadIdx.x + blockIdx.x * blockDim.x;</span> |
| <span class="sd"> y[i] += alpha * x[i];</span> |
| <span class="sd"> }</span> |
| <span class="sd"> '''</span> |
| <span class="sd"> module = mx.rtc.CudaModule(source)</span> |
| <span class="sd"> func = module.get_kernel("axpy", "const float *x, float *y, float alpha")</span> |
| <span class="sd"> x = mx.nd.ones((10,), ctx=mx.gpu(0))</span> |
| <span class="sd"> y = mx.nd.zeros((10,), ctx=mx.gpu(0))</span> |
| <span class="sd"> func.launch([x, y, 3.0], mx.gpu(0), (1, 1, 1), (10, 1, 1))</span> |
| <span class="sd"> print(y)</span> |
| |
| <span class="sd"> Starting from CUDA 8.0, you can instead export functions by name.</span> |
| <span class="sd"> This also allows you to use templates::</span> |
| |
| <span class="sd"> source = r'''</span> |
| <span class="sd"> template<typename DType></span> |
| <span class="sd"> __global__ void axpy(const DType *x, DType *y, DType alpha) {</span> |
| <span class="sd"> int i = threadIdx.x + blockIdx.x * blockDim.x;</span> |
| <span class="sd"> y[i] += alpha * x[i];</span> |
| <span class="sd"> }</span> |
| <span class="sd"> '''</span> |
| <span class="sd"> module = mx.rtc.CudaModule(source, exports=['axpy<float>', 'axpy<double>'])</span> |
| <span class="sd"> func32 = module.get_kernel("axpy<float>", "const float *x, float *y, float alpha")</span> |
| <span class="sd"> x = mx.nd.ones((10,), dtype='float32', ctx=mx.gpu(0))</span> |
| <span class="sd"> y = mx.nd.zeros((10,), dtype='float32', ctx=mx.gpu(0))</span> |
| <span class="sd"> func32.launch([x, y, 3.0], mx.gpu(0), (1, 1, 1), (10, 1, 1))</span> |
| <span class="sd"> print(y)</span> |
| |
| <span class="sd"> func64 = module.get_kernel("axpy<double>", "const double *x, double *y, double alpha")</span> |
| <span class="sd"> x = mx.nd.ones((10,), dtype='float64', ctx=mx.gpu(0))</span> |
| <span class="sd"> y = mx.nd.zeros((10,), dtype='float64', ctx=mx.gpu(0))</span> |
| <span class="sd"> func32.launch([x, y, 3.0], mx.gpu(0), (1, 1, 1), (10, 1, 1))</span> |
| <span class="sd"> print(y)</span> |
| |
| |
| <span class="sd"> Parameters</span> |
| <span class="sd"> ----------</span> |
| <span class="sd"> source : str</span> |
| <span class="sd"> Complete source code.</span> |
| <span class="sd"> options : tuple of str</span> |
| <span class="sd"> Compiler flags. For example, use "-I/usr/local/cuda/include" to</span> |
| <span class="sd"> add cuda headers to include path.</span> |
| <span class="sd"> exports : tuple of str</span> |
| <span class="sd"> Export kernel names.</span> |
| <span class="sd"> """</span> |
| <span class="k">def</span> <span class="nf">__init__</span><span class="p">(</span><span class="bp">self</span><span class="p">,</span> <span class="n">source</span><span class="p">,</span> <span class="n">options</span><span class="o">=</span><span class="p">(),</span> <span class="n">exports</span><span class="o">=</span><span class="p">()):</span> |
| <span class="k">if</span> <span class="nb">isinstance</span><span class="p">(</span><span class="n">options</span><span class="p">,</span> <span class="n">string_types</span><span class="p">):</span> |
| <span class="n">options</span> <span class="o">=</span> <span class="p">(</span><span class="n">options</span><span class="p">,)</span> |
| <span class="k">if</span> <span class="nb">isinstance</span><span class="p">(</span><span class="n">exports</span><span class="p">,</span> <span class="n">string_types</span><span class="p">):</span> |
| <span class="n">exports</span> <span class="o">=</span> <span class="p">(</span><span class="n">exports</span><span class="p">,)</span> |
| <span class="bp">self</span><span class="o">.</span><span class="n">handle</span> <span class="o">=</span> <span class="n">CudaModuleHandle</span><span class="p">()</span> |
| <span class="n">check_call</span><span class="p">(</span><span class="n">_LIB</span><span class="o">.</span><span class="n">MXRtcCudaModuleCreate</span><span class="p">(</span> |
| <span class="n">c_str</span><span class="p">(</span><span class="n">source</span><span class="p">),</span> |
| <span class="nb">len</span><span class="p">(</span><span class="n">options</span><span class="p">),</span> |
| <span class="n">c_str_array</span><span class="p">(</span><span class="n">options</span><span class="p">),</span> |
| <span class="nb">len</span><span class="p">(</span><span class="n">exports</span><span class="p">),</span> |
| <span class="n">c_str_array</span><span class="p">(</span><span class="n">exports</span><span class="p">),</span> |
| <span class="n">ctypes</span><span class="o">.</span><span class="n">byref</span><span class="p">(</span><span class="bp">self</span><span class="o">.</span><span class="n">handle</span><span class="p">)))</span> |
| |
| <span class="k">def</span> <span class="nf">__del__</span><span class="p">(</span><span class="bp">self</span><span class="p">):</span> |
| <span class="n">check_call</span><span class="p">(</span><span class="n">_LIB</span><span class="o">.</span><span class="n">MXRtcCudaModuleFree</span><span class="p">(</span><span class="bp">self</span><span class="o">.</span><span class="n">handle</span><span class="p">))</span> |
| |
| <div class="viewcode-block" id="CudaModule.get_kernel"><a class="viewcode-back" href="../../api/python/rtc/rtc.html#mxnet.rtc.CudaModule.get_kernel">[docs]</a> <span class="k">def</span> <span class="nf">get_kernel</span><span class="p">(</span><span class="bp">self</span><span class="p">,</span> <span class="n">name</span><span class="p">,</span> <span class="n">signature</span><span class="p">):</span> |
| <span class="sa">r</span><span class="sd">"""Get CUDA kernel from compiled module.</span> |
| |
| <span class="sd"> Parameters</span> |
| <span class="sd"> ----------</span> |
| <span class="sd"> name : str</span> |
| <span class="sd"> String name of the kernel.</span> |
| <span class="sd"> signature : str</span> |
| <span class="sd"> Function signature for the kernel. For example, if a kernel is</span> |
| <span class="sd"> declared as::</span> |
| |
| <span class="sd"> extern "C" __global__ void axpy(const float *x, double *y, int alpha)</span> |
| |
| <span class="sd"> Then its signature should be::</span> |
| |
| <span class="sd"> const float *x, double *y, int alpha</span> |
| |
| <span class="sd"> or::</span> |
| |
| <span class="sd"> const float *, double *, int</span> |
| |
| <span class="sd"> Note that `*` in signature marks an argument as array and</span> |
| <span class="sd"> `const` marks an argument as constant (input) array.</span> |
| |
| <span class="sd"> Returns</span> |
| <span class="sd"> -------</span> |
| <span class="sd"> CudaKernel</span> |
| <span class="sd"> CUDA kernels that can be launched on GPUs.</span> |
| <span class="sd"> """</span> |
| <span class="n">hdl</span> <span class="o">=</span> <span class="n">CudaKernelHandle</span><span class="p">()</span> |
| <span class="n">is_ndarray</span> <span class="o">=</span> <span class="p">[]</span> |
| <span class="n">is_const</span> <span class="o">=</span> <span class="p">[]</span> |
| <span class="n">dtypes</span> <span class="o">=</span> <span class="p">[]</span> |
| <span class="n">pattern</span> <span class="o">=</span> <span class="n">re</span><span class="o">.</span><span class="n">compile</span><span class="p">(</span><span class="sa">r</span><span class="s2">"""^\s*(const)?\s*([\w_]+)\s*(\*)?\s*([\w_]+)?\s*$"""</span><span class="p">)</span> |
| <span class="n">args</span> <span class="o">=</span> <span class="n">re</span><span class="o">.</span><span class="n">sub</span><span class="p">(</span><span class="sa">r</span><span class="s2">"\s+"</span><span class="p">,</span> <span class="s2">" "</span><span class="p">,</span> <span class="n">signature</span><span class="p">)</span><span class="o">.</span><span class="n">split</span><span class="p">(</span><span class="s2">","</span><span class="p">)</span> |
| <span class="k">for</span> <span class="n">arg</span> <span class="ow">in</span> <span class="n">args</span><span class="p">:</span> |
| <span class="n">match</span> <span class="o">=</span> <span class="n">pattern</span><span class="o">.</span><span class="n">match</span><span class="p">(</span><span class="n">arg</span><span class="p">)</span> |
| <span class="k">if</span> <span class="ow">not</span> <span class="n">match</span> <span class="ow">or</span> <span class="n">match</span><span class="o">.</span><span class="n">groups</span><span class="p">()[</span><span class="mi">1</span><span class="p">]</span> <span class="o">==</span> <span class="s1">'const'</span><span class="p">:</span> |
| <span class="k">raise</span> <span class="ne">ValueError</span><span class="p">(</span> |
| <span class="s1">'Invalid function prototype "</span><span class="si">%s</span><span class="s1">". Must be in the '</span> |
| <span class="s1">'form of "(const) type (*) (name)"'</span><span class="o">%</span><span class="n">arg</span><span class="p">)</span> |
| <span class="n">is_const</span><span class="o">.</span><span class="n">append</span><span class="p">(</span><span class="nb">bool</span><span class="p">(</span><span class="n">match</span><span class="o">.</span><span class="n">groups</span><span class="p">()[</span><span class="mi">0</span><span class="p">]))</span> |
| <span class="n">dtype</span> <span class="o">=</span> <span class="n">match</span><span class="o">.</span><span class="n">groups</span><span class="p">()[</span><span class="mi">1</span><span class="p">]</span> |
| <span class="n">is_ndarray</span><span class="o">.</span><span class="n">append</span><span class="p">(</span><span class="nb">bool</span><span class="p">(</span><span class="n">match</span><span class="o">.</span><span class="n">groups</span><span class="p">()[</span><span class="mi">2</span><span class="p">]))</span> |
| <span class="k">if</span> <span class="n">dtype</span> <span class="ow">not</span> <span class="ow">in</span> <span class="n">_DTYPE_CPP_TO_NP</span><span class="p">:</span> |
| <span class="k">raise</span> <span class="ne">TypeError</span><span class="p">(</span> |
| <span class="s2">"Unsupported kernel argument type </span><span class="si">%s</span><span class="s2">. Supported types are: </span><span class="si">%s</span><span class="s2">."</span><span class="o">%</span><span class="p">(</span> |
| <span class="n">arg</span><span class="p">,</span> <span class="s1">','</span><span class="o">.</span><span class="n">join</span><span class="p">(</span><span class="n">_DTYPE_CPP_TO_NP</span><span class="o">.</span><span class="n">keys</span><span class="p">())))</span> |
| <span class="n">dtypes</span><span class="o">.</span><span class="n">append</span><span class="p">(</span><span class="n">_DTYPE_NP_TO_MX</span><span class="p">[</span><span class="n">_DTYPE_CPP_TO_NP</span><span class="p">[</span><span class="n">dtype</span><span class="p">]])</span> |
| |
| <span class="n">check_call</span><span class="p">(</span><span class="n">_LIB</span><span class="o">.</span><span class="n">MXRtcCudaKernelCreate</span><span class="p">(</span> |
| <span class="bp">self</span><span class="o">.</span><span class="n">handle</span><span class="p">,</span> |
| <span class="n">c_str</span><span class="p">(</span><span class="n">name</span><span class="p">),</span> |
| <span class="nb">len</span><span class="p">(</span><span class="n">dtypes</span><span class="p">),</span> |
| <span class="n">c_array_buf</span><span class="p">(</span><span class="n">ctypes</span><span class="o">.</span><span class="n">c_int</span><span class="p">,</span> <span class="n">array</span><span class="p">(</span><span class="s1">'i'</span><span class="p">,</span> <span class="n">is_ndarray</span><span class="p">)),</span> |
| <span class="n">c_array_buf</span><span class="p">(</span><span class="n">ctypes</span><span class="o">.</span><span class="n">c_int</span><span class="p">,</span> <span class="n">array</span><span class="p">(</span><span class="s1">'i'</span><span class="p">,</span> <span class="n">is_const</span><span class="p">)),</span> |
| <span class="n">c_array_buf</span><span class="p">(</span><span class="n">ctypes</span><span class="o">.</span><span class="n">c_int</span><span class="p">,</span> <span class="n">array</span><span class="p">(</span><span class="s1">'i'</span><span class="p">,</span> <span class="n">dtypes</span><span class="p">)),</span> |
| <span class="n">ctypes</span><span class="o">.</span><span class="n">byref</span><span class="p">(</span><span class="n">hdl</span><span class="p">)))</span> |
| |
| <span class="k">return</span> <span class="n">CudaKernel</span><span class="p">(</span><span class="n">hdl</span><span class="p">,</span> <span class="n">name</span><span class="p">,</span> <span class="n">is_ndarray</span><span class="p">,</span> <span class="n">dtypes</span><span class="p">)</span></div></div> |
| |
| <div class="viewcode-block" id="CudaKernel"><a class="viewcode-back" href="../../api/python/rtc/rtc.html#mxnet.rtc.CudaKernel">[docs]</a><span class="k">class</span> <span class="nc">CudaKernel</span><span class="p">(</span><span class="nb">object</span><span class="p">):</span> |
| <span class="sd">"""Constructs CUDA kernel. Should be created by `CudaModule.get_kernel`,</span> |
| <span class="sd"> not intended to be used by users."""</span> |
| <span class="k">def</span> <span class="nf">__init__</span><span class="p">(</span><span class="bp">self</span><span class="p">,</span> <span class="n">handle</span><span class="p">,</span> <span class="n">name</span><span class="p">,</span> <span class="n">is_ndarray</span><span class="p">,</span> <span class="n">dtypes</span><span class="p">):</span> |
| <span class="bp">self</span><span class="o">.</span><span class="n">handle</span> <span class="o">=</span> <span class="n">handle</span> |
| <span class="bp">self</span><span class="o">.</span><span class="n">_name</span> <span class="o">=</span> <span class="n">name</span> |
| <span class="bp">self</span><span class="o">.</span><span class="n">_is_ndarray</span> <span class="o">=</span> <span class="n">is_ndarray</span> |
| <span class="bp">self</span><span class="o">.</span><span class="n">_dtypes</span> <span class="o">=</span> <span class="p">[</span><span class="n">_DTYPE_MX_TO_NP</span><span class="p">[</span><span class="n">i</span><span class="p">]</span> <span class="k">for</span> <span class="n">i</span> <span class="ow">in</span> <span class="n">dtypes</span><span class="p">]</span> |
| |
| <span class="k">def</span> <span class="nf">__del__</span><span class="p">(</span><span class="bp">self</span><span class="p">):</span> |
| <span class="n">check_call</span><span class="p">(</span><span class="n">_LIB</span><span class="o">.</span><span class="n">MXRtcCudaKernelFree</span><span class="p">(</span><span class="bp">self</span><span class="o">.</span><span class="n">handle</span><span class="p">))</span> |
| |
| <div class="viewcode-block" id="CudaKernel.launch"><a class="viewcode-back" href="../../api/python/rtc/rtc.html#mxnet.rtc.CudaKernel.launch">[docs]</a> <span class="k">def</span> <span class="nf">launch</span><span class="p">(</span><span class="bp">self</span><span class="p">,</span> <span class="n">args</span><span class="p">,</span> <span class="n">ctx</span><span class="p">,</span> <span class="n">grid_dims</span><span class="p">,</span> <span class="n">block_dims</span><span class="p">,</span> <span class="n">shared_mem</span><span class="o">=</span><span class="mi">0</span><span class="p">):</span> |
| <span class="sd">"""Launch cuda kernel.</span> |
| |
| <span class="sd"> Parameters</span> |
| <span class="sd"> ----------</span> |
| <span class="sd"> args : tuple of NDArray or numbers</span> |
| <span class="sd"> List of arguments for kernel. NDArrays are expected for pointer</span> |
| <span class="sd"> types (e.g. `float*`, `double*`) while numbers are expected for</span> |
| <span class="sd"> non-pointer types (e.g. `int`, `float`).</span> |
| <span class="sd"> ctx : Context</span> |
| <span class="sd"> The context to launch kernel on. Must be GPU context.</span> |
| <span class="sd"> grid_dims : tuple of 3 integers</span> |
| <span class="sd"> Grid dimensions for CUDA kernel.</span> |
| <span class="sd"> block_dims : tuple of 3 integers</span> |
| <span class="sd"> Block dimensions for CUDA kernel.</span> |
| <span class="sd"> shared_mem : integer, optional</span> |
| <span class="sd"> Size of dynamically allocated shared memory. Defaults to 0.</span> |
| <span class="sd"> """</span> |
| <span class="k">assert</span> <span class="n">ctx</span><span class="o">.</span><span class="n">device_type</span> <span class="o">==</span> <span class="s1">'gpu'</span><span class="p">,</span> <span class="s2">"Cuda kernel can only be launched on GPU"</span> |
| <span class="k">assert</span> <span class="nb">len</span><span class="p">(</span><span class="n">grid_dims</span><span class="p">)</span> <span class="o">==</span> <span class="mi">3</span><span class="p">,</span> <span class="s2">"grid_dims must be a tuple of 3 integers"</span> |
| <span class="k">assert</span> <span class="nb">len</span><span class="p">(</span><span class="n">block_dims</span><span class="p">)</span> <span class="o">==</span> <span class="mi">3</span><span class="p">,</span> <span class="s2">"grid_dims must be a tuple of 3 integers"</span> |
| <span class="k">assert</span> <span class="nb">len</span><span class="p">(</span><span class="n">args</span><span class="p">)</span> <span class="o">==</span> <span class="nb">len</span><span class="p">(</span><span class="bp">self</span><span class="o">.</span><span class="n">_dtypes</span><span class="p">),</span> \ |
| <span class="s2">"CudaKernel(</span><span class="si">%s</span><span class="s2">) expects </span><span class="si">%d</span><span class="s2"> arguments but got </span><span class="si">%d</span><span class="s2">"</span><span class="o">%</span><span class="p">(</span> |
| <span class="bp">self</span><span class="o">.</span><span class="n">_name</span><span class="p">,</span> <span class="nb">len</span><span class="p">(</span><span class="bp">self</span><span class="o">.</span><span class="n">_dtypes</span><span class="p">),</span> <span class="nb">len</span><span class="p">(</span><span class="n">args</span><span class="p">))</span> |
| <span class="n">void_args</span> <span class="o">=</span> <span class="p">[]</span> |
| <span class="n">ref_holder</span> <span class="o">=</span> <span class="p">[]</span> |
| <span class="k">for</span> <span class="n">i</span><span class="p">,</span> <span class="p">(</span><span class="n">arg</span><span class="p">,</span> <span class="n">is_nd</span><span class="p">,</span> <span class="n">dtype</span><span class="p">)</span> <span class="ow">in</span> <span class="nb">enumerate</span><span class="p">(</span><span class="nb">zip</span><span class="p">(</span><span class="n">args</span><span class="p">,</span> <span class="bp">self</span><span class="o">.</span><span class="n">_is_ndarray</span><span class="p">,</span> <span class="bp">self</span><span class="o">.</span><span class="n">_dtypes</span><span class="p">)):</span> |
| <span class="k">if</span> <span class="n">is_nd</span><span class="p">:</span> |
| <span class="k">assert</span> <span class="nb">isinstance</span><span class="p">(</span><span class="n">arg</span><span class="p">,</span> <span class="n">NDArray</span><span class="p">),</span> \ |
| <span class="s2">"The </span><span class="si">%d</span><span class="s2">-th argument is expected to be a NDArray but got </span><span class="si">%s</span><span class="s2">"</span><span class="o">%</span><span class="p">(</span> |
| <span class="n">i</span><span class="p">,</span> <span class="nb">type</span><span class="p">(</span><span class="n">arg</span><span class="p">))</span> |
| <span class="n">void_args</span><span class="o">.</span><span class="n">append</span><span class="p">(</span><span class="n">arg</span><span class="o">.</span><span class="n">handle</span><span class="p">)</span> |
| <span class="k">else</span><span class="p">:</span> |
| <span class="k">assert</span> <span class="nb">isinstance</span><span class="p">(</span><span class="n">arg</span><span class="p">,</span> <span class="n">numeric_types</span><span class="p">),</span> \ |
| <span class="s2">"The </span><span class="si">%d</span><span class="s2">-th argument is expected to be a number, but got </span><span class="si">%s</span><span class="s2">"</span><span class="o">%</span><span class="p">(</span> |
| <span class="n">i</span><span class="p">,</span> <span class="nb">type</span><span class="p">(</span><span class="n">arg</span><span class="p">))</span> |
| <span class="n">ref_holder</span><span class="o">.</span><span class="n">append</span><span class="p">(</span><span class="n">np</span><span class="o">.</span><span class="n">array</span><span class="p">(</span><span class="n">arg</span><span class="p">,</span> <span class="n">dtype</span><span class="o">=</span><span class="n">dtype</span><span class="p">))</span> |
| <span class="n">void_args</span><span class="o">.</span><span class="n">append</span><span class="p">(</span><span class="n">ref_holder</span><span class="p">[</span><span class="o">-</span><span class="mi">1</span><span class="p">]</span><span class="o">.</span><span class="n">ctypes</span><span class="o">.</span><span class="n">data_as</span><span class="p">(</span><span class="n">ctypes</span><span class="o">.</span><span class="n">c_void_p</span><span class="p">))</span> |
| |
| <span class="n">check_call</span><span class="p">(</span><span class="n">_LIB</span><span class="o">.</span><span class="n">MXRtcCudaKernelCall</span><span class="p">(</span> |
| <span class="bp">self</span><span class="o">.</span><span class="n">handle</span><span class="p">,</span> |
| <span class="n">ctx</span><span class="o">.</span><span class="n">device_id</span><span class="p">,</span> |
| <span class="n">c_array</span><span class="p">(</span><span class="n">ctypes</span><span class="o">.</span><span class="n">c_void_p</span><span class="p">,</span> <span class="n">void_args</span><span class="p">),</span> |
| <span class="n">mx_uint</span><span class="p">(</span><span class="n">grid_dims</span><span class="p">[</span><span class="mi">0</span><span class="p">]),</span> <span class="n">mx_uint</span><span class="p">(</span><span class="n">grid_dims</span><span class="p">[</span><span class="mi">1</span><span class="p">]),</span> <span class="n">mx_uint</span><span class="p">(</span><span class="n">grid_dims</span><span class="p">[</span><span class="mi">2</span><span class="p">]),</span> |
| <span class="n">mx_uint</span><span class="p">(</span><span class="n">block_dims</span><span class="p">[</span><span class="mi">0</span><span class="p">]),</span> <span class="n">mx_uint</span><span class="p">(</span><span class="n">block_dims</span><span class="p">[</span><span class="mi">1</span><span class="p">]),</span> <span class="n">mx_uint</span><span class="p">(</span><span class="n">block_dims</span><span class="p">[</span><span class="mi">2</span><span class="p">]),</span> |
| <span class="n">mx_uint</span><span class="p">(</span><span class="n">shared_mem</span><span class="p">)))</span></div></div> |
| </pre></div> |
| </div> |
| </div> |
| <div aria-label="main navigation" class="sphinxsidebar rightsidebar" role="navigation"> |
| <div class="sphinxsidebarwrapper"> |
| </div> |
| </div> |
| </div><div class="footer"> |
| <div class="section-disclaimer"> |
| <div class="container"> |
| <div> |
| <img height="60" src="https://raw.githubusercontent.com/dmlc/web-data/master/mxnet/image/apache_incubator_logo.png"/> |
| <p> |
| Apache MXNet is an effort undergoing incubation at The Apache Software Foundation (ASF), <strong>sponsored by the <i>Apache Incubator</i></strong>. Incubation is required of all newly accepted projects until a further review indicates that the infrastructure, communications, and decision making process have stabilized in a manner consistent with other successful ASF projects. While incubation status is not necessarily a reflection of the completeness or stability of the code, it does indicate that the project has yet to be fully endorsed by the ASF. |
| </p> |
| <p> |
| "Copyright © 2017-2018, The Apache Software Foundation |
| Apache MXNet, MXNet, Apache, the Apache feather, and the Apache MXNet project logo are either registered trademarks or trademarks of the Apache Software Foundation." |
| </p> |
| </div> |
| </div> |
| </div> |
| </div> <!-- pagename != index --> |
| </div> |
| <script crossorigin="anonymous" integrity="sha384-0mSbJDEHialfmuBBQP6A4Qrprq5OVfW37PRR3j5ELqxss1yVqOtnepnHVP9aJ7xS" src="https://maxcdn.bootstrapcdn.com/bootstrap/3.3.6/js/bootstrap.min.js"></script> |
| <script src="../../_static/js/sidebar.js" type="text/javascript"></script> |
| <script src="../../_static/js/search.js" type="text/javascript"></script> |
| <script src="../../_static/js/navbar.js" type="text/javascript"></script> |
| <script src="../../_static/js/clipboard.min.js" type="text/javascript"></script> |
| <script src="../../_static/js/copycode.js" type="text/javascript"></script> |
| <script src="../../_static/js/page.js" type="text/javascript"></script> |
| <script src="../../_static/js/docversion.js" type="text/javascript"></script> |
| <script type="text/javascript"> |
| $('body').ready(function () { |
| $('body').css('visibility', 'visible'); |
| }); |
| </script> |
| </body> |
| </html> |
