Google Git
Sign in
apache / mxnet-site / refs/heads/test-revert / . / _modules / mxnet / rtc.html
blob: b5aafb363c025f552879ff8d2a6d9ca870df4ab7 [file] [log] [blame]
<!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"/>
<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: ''
};
</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("/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://cdn.mathjax.org/mathjax/latest/MathJax.js?config=TeX-AMS-MML_HTMLorMML"></script> -->
<!-- -->
<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></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="../../install/index.html">Install</a>
<a class="main-nav-link" href="../../tutorials/index.html">Tutorials</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="../../gluon/index.html">About</a></li>
<li><a class="main-nav-link" href="http://gluon.mxnet.io">Tutorials</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="../../api/python/index.html">Python</a></li>
<li><a class="main-nav-link" href="../../api/scala/index.html">Scala</a></li>
<li><a class="main-nav-link" href="../../api/r/index.html">R</a></li>
<li><a class="main-nav-link" href="../../api/julia/index.html">Julia</a></li>
<li><a class="main-nav-link" href="../../api/c++/index.html">C++</a></li>
<li><a class="main-nav-link" href="../../api/perl/index.html">Perl</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="../../faq/index.html">FAQ</a></li>
<li><a class="main-nav-link" href="../../architecture/index.html">Architecture</a></li>
<li><a class="main-nav-link" href="https://github.com/apache/incubator-mxnet/tree/1.0.0/example">Examples</a></li>
<li><a class="main-nav-link" href="../../model_zoo/index.html">Model Zoo</a></li>
</ul>
</span>
<a class="main-nav-link" href="https://github.com/dmlc/mxnet">Github</a>
<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="../../community/index.html">Community</a></li>
<li><a class="main-nav-link" href="../../community/contribute.html">Contribute</a></li>
<li><a class="main-nav-link" href="../../community/powered_by.html">Powered By</a></li>
</ul>
</span>
<a class="main-nav-link" href="http://discuss.mxnet.io">Discuss</a>
<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">Versions(1.0.0)<span class="caret"></span></a><ul id="package-dropdown-menu" class="dropdown-menu"><li><a class="main-nav-link" href=https://mxnet.incubator.apache.org/>1.0.0</a></li><li><a class="main-nav-link" href=https://mxnet.incubator.apache.org/versions/0.12.1/index.html>0.12.1</a></li><li><a class="main-nav-link" href=https://mxnet.incubator.apache.org/versions/0.12.0/index.html>0.12.0</a></li><li><a class="main-nav-link" href=https://mxnet.incubator.apache.org/versions/0.11.0/index.html>0.11.0</a></li><li><a class="main-nav-link" href=https://mxnet.incubator.apache.org/versions/master/index.html>master</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="../../install/index.html">Install</a></li>
<li><a class="main-nav-link" href="../../tutorials/index.html">Tutorials</a></li>
<li class="dropdown-submenu">
<a href="#" tabindex="-1">Community</a>
<ul class="dropdown-menu">
<li><a href="../../community/index.html" tabindex="-1">Community</a></li>
<li><a href="../../community/contribute.html" tabindex="-1">Contribute</a></li>
<li><a href="../../community/powered_by.html" tabindex="-1">Powered By</a></li>
</ul>
</li>
<li class="dropdown-submenu">
<a href="#" tabindex="-1">API</a>
<ul class="dropdown-menu">
<li><a href="../../api/python/index.html" tabindex="-1">Python</a>
</li>
<li><a href="../../api/scala/index.html" tabindex="-1">Scala</a>
</li>
<li><a href="../../api/r/index.html" tabindex="-1">R</a>
</li>
<li><a href="../../api/julia/index.html" tabindex="-1">Julia</a>
</li>
<li><a href="../../api/c++/index.html" tabindex="-1">C++</a>
</li>
<li><a href="../../api/perl/index.html" tabindex="-1">Perl</a>
</li>
</ul>
</li>
<li class="dropdown-submenu">
<a href="#" tabindex="-1">Docs</a>
<ul class="dropdown-menu">
<li><a href="../../tutorials/index.html" tabindex="-1">Tutorials</a></li>
<li><a href="../../faq/index.html" tabindex="-1">FAQ</a></li>
<li><a href="../../architecture/index.html" tabindex="-1">Architecture</a></li>
<li><a href="https://github.com/apache/incubator-mxnet/tree/1.0.0/example" tabindex="-1">Examples</a></li>
<li><a href="../../model_zoo/index.html" tabindex="-1">Model Zoo</a></li>
</ul>
</li>
<li><a href="../../architecture/index.html">Architecture</a></li>
<li><a class="main-nav-link" href="https://github.com/dmlc/mxnet">Github</a></li>
<li id="dropdown-menu-position-anchor-version-mobile" class="dropdown-submenu" style="position: relative"><a href="#" tabindex="-1">Versions(1.0.0)</a><ul class="dropdown-menu"><li><a tabindex="-1" href=https://mxnet.incubator.apache.org/>1.0.0</a></li><li><a tabindex="-1" href=https://mxnet.incubator.apache.org/versions/0.12.1/index.html>0.12.1</a></li><li><a tabindex="-1" href=https://mxnet.incubator.apache.org/versions/0.12.0/index.html>0.12.0</a></li><li><a tabindex="-1" href=https://mxnet.incubator.apache.org/versions/0.11.0/index.html>0.11.0</a></li><li><a tabindex="-1" href=https://mxnet.incubator.apache.org/versions/master/index.html>master</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/python/index.html">Python Documents</a></li>
<li class="toctree-l1"><a class="reference internal" href="../../api/r/index.html">R Documents</a></li>
<li class="toctree-l1"><a class="reference internal" href="../../api/julia/index.html">Julia Documents</a></li>
<li class="toctree-l1"><a class="reference internal" href="../../api/c++/index.html">C++ Documents</a></li>
<li class="toctree-l1"><a class="reference internal" href="../../api/scala/index.html">Scala Documents</a></li>
<li class="toctree-l1"><a class="reference internal" href="../../api/perl/index.html">Perl Documents</a></li>
<li class="toctree-l1"><a class="reference internal" href="../../faq/index.html">HowTo Documents</a></li>
<li class="toctree-l1"><a class="reference internal" href="../../architecture/index.html">System Documents</a></li>
<li class="toctree-l1"><a class="reference internal" href="../../tutorials/index.html">Tutorials</a></li>
<li class="toctree-l1"><a class="reference internal" href="../../community/index.html">Community</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="kn">import</span> <span class="n">absolute_import</span>
<span class="kn">from</span> <span class="nn">array</span> <span class="kn">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="kn">as</span> <span class="nn">np</span>
<span class="kn">from</span> <span class="nn">.base</span> <span class="kn">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="kn">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="kn">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="fm">__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="fm">__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="fm">__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="fm">__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, 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 type="text/javascript">
$('body').ready(function () {
$('body').css('visibility', 'visible');
});
</script>
</body>
</html>