You are viewing a plain text version of this content. The canonical link for it is here.
Posted to commits@mxnet.apache.org by GitBox <gi...@apache.org> on 2020/01/30 19:51:36 UTC
[GitHub] [incubator-mxnet] rondogency opened a new pull request #17486:
Update CustomOp doc with changes for GPU support
rondogency opened a new pull request #17486: Update CustomOp doc with changes for GPU support
URL: https://github.com/apache/incubator-mxnet/pull/17486
## Description
Improve the custom operator example doc by adding how to write GPU custom operator.
This is a continuation of the CustomOp tutorial doc #17241
## This PR is not
- Updating docs on the MXNet website on extending MXNet
- Consolidating operator related docs on the website
Those will be addressed in the coming PRs
## Checklist ##
### Essentials ###
Please feel free to remove inapplicable items for your PR.
- [ ] The PR title starts with [MXNET-$JIRA_ID], where $JIRA_ID refers to the relevant [JIRA issue](https://issues.apache.org/jira/projects/MXNET/issues) created (except PRs with tiny changes)
- [ ] Changes are complete (i.e. I finished coding on this PR)
- [ ] All changes have test coverage:
- Unit tests are added for small changes to verify correctness (e.g. adding a new operator)
- Nightly tests are added for complicated/long-running ones (e.g. changing distributed kvstore)
- Build tests will be added for build configuration changes (e.g. adding a new build option with NCCL)
- [ ] Code is well-documented:
- For user-facing API changes, API doc string has been updated.
- For new C++ functions in header files, their functionalities and arguments are documented.
- For new examples, README.md is added to explain the what the example does, the source of the dataset, expected performance on test set and reference to the original paper if applicable
- Check the API doc at https://mxnet-ci-doc.s3-accelerate.dualstack.amazonaws.com/PR-$PR_ID/$BUILD_ID/index.html
- [ ] To the best of my knowledge, examples are either not affected by this change, or have been fixed to be compatible with this change
### Changes ###
- [ ] Feature1, tests, (and when applicable, API doc)
- [ ] Feature2, tests, (and when applicable, API doc)
## Comments ##
- If this change is a backward incompatible change, why must this change be made.
- Interesting edge cases to note here
----------------------------------------------------------------
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
With regards,
Apache Git Services
[GitHub] [incubator-mxnet] aaronmarkham commented on a change in pull
request #17486: Update CustomOp doc with changes for GPU support
Posted by GitBox <gi...@apache.org>.
aaronmarkham commented on a change in pull request #17486: Update CustomOp doc with changes for GPU support
URL: https://github.com/apache/incubator-mxnet/pull/17486#discussion_r373216524
##########
File path: example/extensions/lib_custom_op/README.md
##########
@@ -67,76 +80,85 @@ mx.library.load(‘libmyop_lib.so’)
mx.nd.my_op(...)
```
-### Writing Regular Custom Operator:
+### Writing Regular Custom Operator
Review comment:
```suggestion
### Writing a Regular Custom Operator
```
----------------------------------------------------------------
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
With regards,
Apache Git Services
[GitHub] [incubator-mxnet] rondogency edited a comment on issue #17486:
Update CustomOp doc with changes for GPU support
Posted by GitBox <gi...@apache.org>.
rondogency edited a comment on issue #17486: Update CustomOp doc with changes for GPU support
URL: https://github.com/apache/incubator-mxnet/pull/17486#issuecomment-586023537
@a550461053 Thanks for the questions. Currently you cannot call built-in operators from MXNet. We are making a plan for adding this support, but it cannot be easily implemented and may be a separate project. Also I added a snippet of how to use mutable input, hope it would help.
----------------------------------------------------------------
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
With regards,
Apache Git Services
[GitHub] [incubator-mxnet] aaronmarkham commented on a change in pull
request #17486: Update CustomOp doc with changes for GPU support
Posted by GitBox <gi...@apache.org>.
aaronmarkham commented on a change in pull request #17486: Update CustomOp doc with changes for GPU support
URL: https://github.com/apache/incubator-mxnet/pull/17486#discussion_r373216296
##########
File path: example/extensions/lib_custom_op/README.md
##########
@@ -28,26 +28,39 @@ Custom operators (CustomOp) enable users to write new operators without compilin
### Have MXNet Ready
-First you should install MXNet either from compiling from source code or download from nightly build. It doesn’t matter if the build comes with CUDA or MKLDNN. The custom operator doesn’t interact with the execution of other native MXNet operators.
+Custom Operator support was recently merged (#15921, #17270) and is not available in a released version of MXNet yet. It will be part of the forthcoming 1.7 release, until then please install MXNet by compiling from source or downloading one of the nightly builds. For running the example below, it doesn’t matter if it is a CUDA, MKLDNN or Vanila build. The custom operator doesn’t interact with the execution of other native MXNet operators. Note that if you want to run GPU examples and write your custom operators running on GPU, you still need MXNet CUDA build.
-### Run An Example:
+### Run An Example
-You can start getting familiar with custom operators by running some examples provided in the **example/extensions/lib_custom_op** directory. Start with a common linear algebra operator like `gemm` (Generalized Matrix Multiplication). Go to `lib_custom_op` directory and follow these steps:
+You can start getting familiar with custom operators by running some examples provided in the `example/extensions/lib_custom_op` directory. Start with a common linear algebra operator like `gemm` (Generalized Matrix Multiplication). Go to `lib_custom_op` directory and follow these steps:
1. Run `make gemm_lib`. The Makefile will generate a dynamic library **libgemm_lib.so** compiled from `gemm_lib.cc`. This is the library you are going to load that contains everything for the custom gemm operator.
-2. Run `python test_gemm.py`. It’ll first load the above .so library, find the operators, register them in the MXNet backend, print "Found x operators", then invoke the operator like a regular MXNet operator and output the result.
+2. Run `python test_gemm.py`. It’ll first load the library compiled from step 1, find the operators, register them in the MXNet backend, then invoke the operator like a regular MXNet operator and output the result.
+```
+[19:22:02] ../src/c_api/c_api.cc:286: Found 2 operators in library
+[19:22:02] ../src/c_api/c_api.cc:350: Op[0] my_gemm
+[19:22:02] ../src/c_api/c_api.cc:350: Op[1] state_gemm
+[19:22:02] ../src/c_api/c_api.cc:785: Found 0 partitioners in library
+--------start ndarray compute---------
+[[ 50.]
+ [122.]]
+<NDArray 2x1 @cpu(0)>
+...
+```
+
+Note that you can safely ignore the `Found 0 partitioners` info as it is not related to the custom operator.
-### Basic Files For Gemm Library:
+### Basic Files For Gemm Library
* **lib_custom_op/gemm_lib.cc**: This file has a source code implementation of all required components of a custom operator, as well as the registration of the custom operator.
* **lib_custom_op/Makefile**: Compile source code to a dynamic shared library, with a header file `include/mxnet/lib_api.h` from MXNet source code. Currently the custom operator is compatible with C++11 onwards.
* **lib_custom_op/test_gemm.py**: This file calls `mx.library.load(‘libgemm_lib.so’)` to load the library containing the custom operator, invokes the operator using both NDArray and Symbol APIs, and prints outputs of the forward and backward passes. The outputs should be the same as the regular MXNet `gemm` operator.
-## Writing Custom Operator Library:
+## Writing Custom CPU Operator Library
Review comment:
```suggestion
## Writing a Custom CPU Operator Library
```
----------------------------------------------------------------
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
With regards,
Apache Git Services
[GitHub] [incubator-mxnet] samskalicky commented on a change in pull request
#17486: Update CustomOp doc with changes for GPU support
Posted by GitBox <gi...@apache.org>.
samskalicky commented on a change in pull request #17486: Update CustomOp doc with changes for GPU support
URL: https://github.com/apache/incubator-mxnet/pull/17486#discussion_r375023931
##########
File path: example/extensions/lib_custom_op/README.md
##########
@@ -146,21 +186,160 @@ Let’s take a closer look at those registry functions:
* **inferShape**: This function is similar to the `inferType` function, except it is used for populating the output data shapes. You need to figure out the shapes of each output tensors for this computation. For example, if the inputs are images with shape (224,224,3) and you write a padding operator to make 10px borders for the images, then your output shape will be (234,234,3).
-* **forward**: This function executes the main forward computation. It takes four arguments. The 1st argument is the attributes. The 2nd argument is the input `MXTensors` which stores all data and info of input ndarrays. The 3rd argument is the output `MXTensors`. The 4th argument is the `OpResource` object for memory allocation and other utilities. Additionally, you can use a `dltensor` tensor structure stored in the `MXTensor` as a more standardized data structure for computing.
+* **forward**: This function executes the main forward computation. It takes four arguments. The 1st argument is the attributes. The 2nd argument is the input `MXTensors` which stores all data and info of input ndarrays. The 3rd argument is the output `MXTensors`. The 4th argument is the `OpResource` object for memory allocation and other utilities. The details of `OpResource` are covered in the below section.
+Additionally, you can use a `dltensor` tensor structure stored in the `MXTensor` as a more standardized data structure for computing.
* **backward**: This function is doing the backward gradient computation. It will be similar to the forward function. And you need to figure out the formula of the backward gradient computation.
* **mutateInputs**: This function is for marking mutable inputs. It takes two arguments. The 1st argument is the attributes. The 2nd argument is a list of input indices that are mutable among all input tensors. It is useful when some inputs are auxiliary model parameters and might be altered during forward/backward computation. Remember, the index number of `input_indices` should not exceed the number of inputs.
-### Writing Stateful Custom Operator:
+### Writing Stateful Custom Operator
A stateful custom operator is useful when a forward/backward call needs some data or ‘state’ from previous forward/backward calls. Normally we create a class, and make instance variables store the states used for computing or caching.
Most of the building blocks for making a stateful custom operator is the same as regular custom operator, except it’ll register `createOpState` instead of a `forward` function for the computation.
* [createOpState](./gemm_lib.cc#L204) - Create stateful operator instance:
* This function takes two arguments. The 1st argument is attributes. The 2nd argument is a placeholder for `CustomStatefulOp` object. You must [define a class that inherits CustomStatefulOp](./gemm_lib.cc#L178) and override the forward function (optionally the backward function). Then you need to create an instance of your class and assign it to the placeholder. In this way, all of the forward/backward calls will use the same methods in that instance, and the instance is able to keep the state of the operator.
+```c++
+ MXReturnValue createOpState(
+ std::map<std::string, std::string> attrs,
+ CustomStatefulOp** op_inst)
+```
+
+* The operator registering function will look like this:
+```c++
+ REGISTER_OP(my_state_op)
+ ...
+ .setCreateOpState(createOpState, "cpu");
+```
+
+## Writing A Custom GPU Operator Library
+
+Most of the building blocks for registering GPU custom operators are the exactly same as CPU ones, except you need to specify the `"gpu"` context name when registering `forward`, `backward` or `createOpState` function.
+
+### Run A GPU Example
+
+For illustration purposes, we provided a `ReLU` (Rectified Linear Unit) activation operator that can run on GPU. Make sure you have installed a CUDA compatible MXNet build. Go to `lib_custom_op` directory and follow these steps:
+
+1. Run `make relu_lib`. The Makefile will invoke `NVCC` compiler to compile the CUDA kernel along with regular custom operator functions from `relu_lib.cu` to generate `librelu_lib.so` library.
+2. Run `python test_relu.py`. It’ll register the GPU `ReLU` operator in the MXNet backend, then invoke the operator by feeding an `NDArray` input with GPU context, and output the result tensor with GPU context.
Review comment:
feeding --> passing
NDArray --> MXTensor
----------------------------------------------------------------
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
With regards,
Apache Git Services
[GitHub] [incubator-mxnet] samskalicky commented on issue #17486: Update
CustomOp doc with changes for GPU support
Posted by GitBox <gi...@apache.org>.
samskalicky commented on issue #17486: Update CustomOp doc with changes for GPU support
URL: https://github.com/apache/incubator-mxnet/pull/17486#issuecomment-586640766
> But now I can only build mxnet to create c++ customOp and I have a problem on compiling into python wheel at #17577
@a550461053 just to clarify, you do not need to recompile MXNet from source to use a custom operator, you can dynamically load your library into MXNet at runtime using one of the nightly builds from here: https://repo.mxnet.io/dist/index.html
----------------------------------------------------------------
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
With regards,
Apache Git Services
[GitHub] [incubator-mxnet] aaronmarkham commented on a change in pull
request #17486: Update CustomOp doc with changes for GPU support
Posted by GitBox <gi...@apache.org>.
aaronmarkham commented on a change in pull request #17486: Update CustomOp doc with changes for GPU support
URL: https://github.com/apache/incubator-mxnet/pull/17486#discussion_r373216396
##########
File path: example/extensions/lib_custom_op/README.md
##########
@@ -28,26 +28,39 @@ Custom operators (CustomOp) enable users to write new operators without compilin
### Have MXNet Ready
-First you should install MXNet either from compiling from source code or download from nightly build. It doesn’t matter if the build comes with CUDA or MKLDNN. The custom operator doesn’t interact with the execution of other native MXNet operators.
+Custom Operator support was recently merged (#15921, #17270) and is not available in a released version of MXNet yet. It will be part of the forthcoming 1.7 release, until then please install MXNet by compiling from source or downloading one of the nightly builds. For running the example below, it doesn’t matter if it is a CUDA, MKLDNN or Vanila build. The custom operator doesn’t interact with the execution of other native MXNet operators. Note that if you want to run GPU examples and write your custom operators running on GPU, you still need MXNet CUDA build.
-### Run An Example:
+### Run An Example
-You can start getting familiar with custom operators by running some examples provided in the **example/extensions/lib_custom_op** directory. Start with a common linear algebra operator like `gemm` (Generalized Matrix Multiplication). Go to `lib_custom_op` directory and follow these steps:
+You can start getting familiar with custom operators by running some examples provided in the `example/extensions/lib_custom_op` directory. Start with a common linear algebra operator like `gemm` (Generalized Matrix Multiplication). Go to `lib_custom_op` directory and follow these steps:
1. Run `make gemm_lib`. The Makefile will generate a dynamic library **libgemm_lib.so** compiled from `gemm_lib.cc`. This is the library you are going to load that contains everything for the custom gemm operator.
-2. Run `python test_gemm.py`. It’ll first load the above .so library, find the operators, register them in the MXNet backend, print "Found x operators", then invoke the operator like a regular MXNet operator and output the result.
+2. Run `python test_gemm.py`. It’ll first load the library compiled from step 1, find the operators, register them in the MXNet backend, then invoke the operator like a regular MXNet operator and output the result.
+```
+[19:22:02] ../src/c_api/c_api.cc:286: Found 2 operators in library
+[19:22:02] ../src/c_api/c_api.cc:350: Op[0] my_gemm
+[19:22:02] ../src/c_api/c_api.cc:350: Op[1] state_gemm
+[19:22:02] ../src/c_api/c_api.cc:785: Found 0 partitioners in library
+--------start ndarray compute---------
+[[ 50.]
+ [122.]]
+<NDArray 2x1 @cpu(0)>
+...
+```
+
+Note that you can safely ignore the `Found 0 partitioners` info as it is not related to the custom operator.
-### Basic Files For Gemm Library:
+### Basic Files For Gemm Library
Review comment:
```suggestion
### Basic Files For a Gemm Library
```
----------------------------------------------------------------
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
With regards,
Apache Git Services
[GitHub] [incubator-mxnet] aaronmarkham commented on a change in pull
request #17486: Update CustomOp doc with changes for GPU support
Posted by GitBox <gi...@apache.org>.
aaronmarkham commented on a change in pull request #17486: Update CustomOp doc with changes for GPU support
URL: https://github.com/apache/incubator-mxnet/pull/17486#discussion_r373217041
##########
File path: example/extensions/lib_custom_op/README.md
##########
@@ -146,21 +168,154 @@ Let’s take a closer look at those registry functions:
* **inferShape**: This function is similar to the `inferType` function, except it is used for populating the output data shapes. You need to figure out the shapes of each output tensors for this computation. For example, if the inputs are images with shape (224,224,3) and you write a padding operator to make 10px borders for the images, then your output shape will be (234,234,3).
-* **forward**: This function executes the main forward computation. It takes four arguments. The 1st argument is the attributes. The 2nd argument is the input `MXTensors` which stores all data and info of input ndarrays. The 3rd argument is the output `MXTensors`. The 4th argument is the `OpResource` object for memory allocation and other utilities. Additionally, you can use a `dltensor` tensor structure stored in the `MXTensor` as a more standardized data structure for computing.
+* **forward**: This function executes the main forward computation. It takes four arguments. The 1st argument is the attributes. The 2nd argument is the input `MXTensors` which stores all data and info of input ndarrays. The 3rd argument is the output `MXTensors`. The 4th argument is the `OpResource` object for memory allocation and other utilities. The details of `OpResource` are covered in the below section.
+Additionally, you can use a `dltensor` tensor structure stored in the `MXTensor` as a more standardized data structure for computing.
* **backward**: This function is doing the backward gradient computation. It will be similar to the forward function. And you need to figure out the formula of the backward gradient computation.
* **mutateInputs**: This function is for marking mutable inputs. It takes two arguments. The 1st argument is the attributes. The 2nd argument is a list of input indices that are mutable among all input tensors. It is useful when some inputs are auxiliary model parameters and might be altered during forward/backward computation. Remember, the index number of `input_indices` should not exceed the number of inputs.
-### Writing Stateful Custom Operator:
+### Writing Stateful Custom Operator
A stateful custom operator is useful when a forward/backward call needs some data or ‘state’ from previous forward/backward calls. Normally we create a class, and make instance variables store the states used for computing or caching.
Most of the building blocks for making a stateful custom operator is the same as regular custom operator, except it’ll register `createOpState` instead of a `forward` function for the computation.
* [createOpState](./gemm_lib.cc#L204) - Create stateful operator instance:
* This function takes two arguments. The 1st argument is attributes. The 2nd argument is a placeholder for `CustomStatefulOp` object. You must [define a class that inherits CustomStatefulOp](./gemm_lib.cc#L178) and override the forward function (optionally the backward function). Then you need to create an instance of your class and assign it to the placeholder. In this way, all of the forward/backward calls will use the same methods in that instance, and the instance is able to keep the state of the operator.
+```c++
Review comment:
```suggestion
```c++
```
give it some space
----------------------------------------------------------------
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
With regards,
Apache Git Services
[GitHub] [incubator-mxnet] samskalicky commented on a change in pull request
#17486: Update CustomOp doc with changes for GPU support
Posted by GitBox <gi...@apache.org>.
samskalicky commented on a change in pull request #17486: Update CustomOp doc with changes for GPU support
URL: https://github.com/apache/incubator-mxnet/pull/17486#discussion_r375011251
##########
File path: example/extensions/lib_custom_op/README.md
##########
@@ -28,26 +28,42 @@ Custom operators (CustomOp) enable users to write new operators without compilin
### Have MXNet Ready
-First you should install MXNet either from compiling from source code or download from nightly build. It doesn’t matter if the build comes with CUDA or MKLDNN. The custom operator doesn’t interact with the execution of other native MXNet operators.
+Custom Operator support was merged (#15921, #17270) and is not available in versions of MXNet prior to v1.7.0. To access the feature now, please install MXNet by compiling from source using master or using the previously mentioned commits, downloading one of the nightly builds, or from a release of MXNet 1.7.0+. For running the following example, it doesn’t matter if it is a CUDA, MKLDNN or plain MXNet build; the custom operator doesn’t interact with the execution of other native MXNet operators. Note that if you want to run GPU examples and write your custom operators running on GPU, you still need an MXNet CUDA build.
-### Run An Example:
+### Run An Example
-You can start getting familiar with custom operators by running some examples provided in the **example/extensions/lib_custom_op** directory. Start with a common linear algebra operator like `gemm` (Generalized Matrix Multiplication). Go to `lib_custom_op` directory and follow these steps:
+You can start getting familiar with custom operators by running some examples provided in the `example/extensions/lib_custom_op` directory. Start with a common linear algebra operator like `gemm` (Generalized Matrix Multiplication). Go to `lib_custom_op` directory and follow these steps:
1. Run `make gemm_lib`. The Makefile will generate a dynamic library **libgemm_lib.so** compiled from `gemm_lib.cc`. This is the library you are going to load that contains everything for the custom gemm operator.
-2. Run `python test_gemm.py`. It’ll first load the above .so library, find the operators, register them in the MXNet backend, print "Found x operators", then invoke the operator like a regular MXNet operator and output the result.
+2. Run `python test_gemm.py`. It’ll first load the library compiled from step 1, find the operators, register them in the MXNet backend, then invoke the operator like a regular MXNet operator and output the result.
Review comment:
This is good. Lets add another sentence introducing the log that is coming next. I suggest:
"Below is the output when running the `python test_gemm.py` command. Notice that it loaded 2 operators: `my_gemm` and `state_gemm`.
----------------------------------------------------------------
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
With regards,
Apache Git Services
[GitHub] [incubator-mxnet] rondogency commented on a change in pull request
#17486: Update CustomOp doc with changes for GPU support
Posted by GitBox <gi...@apache.org>.
rondogency commented on a change in pull request #17486: Update CustomOp doc with changes for GPU support
URL: https://github.com/apache/incubator-mxnet/pull/17486#discussion_r375602215
##########
File path: example/extensions/lib_custom_op/README.md
##########
@@ -146,21 +186,160 @@ Let’s take a closer look at those registry functions:
* **inferShape**: This function is similar to the `inferType` function, except it is used for populating the output data shapes. You need to figure out the shapes of each output tensors for this computation. For example, if the inputs are images with shape (224,224,3) and you write a padding operator to make 10px borders for the images, then your output shape will be (234,234,3).
-* **forward**: This function executes the main forward computation. It takes four arguments. The 1st argument is the attributes. The 2nd argument is the input `MXTensors` which stores all data and info of input ndarrays. The 3rd argument is the output `MXTensors`. The 4th argument is the `OpResource` object for memory allocation and other utilities. Additionally, you can use a `dltensor` tensor structure stored in the `MXTensor` as a more standardized data structure for computing.
+* **forward**: This function executes the main forward computation. It takes four arguments. The 1st argument is the attributes. The 2nd argument is the input `MXTensors` which stores all data and info of input ndarrays. The 3rd argument is the output `MXTensors`. The 4th argument is the `OpResource` object for memory allocation and other utilities. The details of `OpResource` are covered in the below section.
+Additionally, you can use a `dltensor` tensor structure stored in the `MXTensor` as a more standardized data structure for computing.
* **backward**: This function is doing the backward gradient computation. It will be similar to the forward function. And you need to figure out the formula of the backward gradient computation.
* **mutateInputs**: This function is for marking mutable inputs. It takes two arguments. The 1st argument is the attributes. The 2nd argument is a list of input indices that are mutable among all input tensors. It is useful when some inputs are auxiliary model parameters and might be altered during forward/backward computation. Remember, the index number of `input_indices` should not exceed the number of inputs.
-### Writing Stateful Custom Operator:
+### Writing Stateful Custom Operator
A stateful custom operator is useful when a forward/backward call needs some data or ‘state’ from previous forward/backward calls. Normally we create a class, and make instance variables store the states used for computing or caching.
Most of the building blocks for making a stateful custom operator is the same as regular custom operator, except it’ll register `createOpState` instead of a `forward` function for the computation.
* [createOpState](./gemm_lib.cc#L204) - Create stateful operator instance:
* This function takes two arguments. The 1st argument is attributes. The 2nd argument is a placeholder for `CustomStatefulOp` object. You must [define a class that inherits CustomStatefulOp](./gemm_lib.cc#L178) and override the forward function (optionally the backward function). Then you need to create an instance of your class and assign it to the placeholder. In this way, all of the forward/backward calls will use the same methods in that instance, and the instance is able to keep the state of the operator.
+```c++
+ MXReturnValue createOpState(
+ std::map<std::string, std::string> attrs,
+ CustomStatefulOp** op_inst)
+```
+
+* The operator registering function will look like this:
+```c++
+ REGISTER_OP(my_state_op)
+ ...
+ .setCreateOpState(createOpState, "cpu");
+```
+
+## Writing A Custom GPU Operator Library
+
+Most of the building blocks for registering GPU custom operators are the exactly same as CPU ones, except you need to specify the `"gpu"` context name when registering `forward`, `backward` or `createOpState` function.
+
+### Run A GPU Example
+
+For illustration purposes, we provided a `ReLU` (Rectified Linear Unit) activation operator that can run on GPU. Make sure you have installed a CUDA compatible MXNet build. Go to `lib_custom_op` directory and follow these steps:
+
+1. Run `make relu_lib`. The Makefile will invoke `NVCC` compiler to compile the CUDA kernel along with regular custom operator functions from `relu_lib.cu` to generate `librelu_lib.so` library.
+2. Run `python test_relu.py`. It’ll register the GPU `ReLU` operator in the MXNet backend, then invoke the operator by feeding an `NDArray` input with GPU context, and output the result tensor with GPU context.
+
+### Writing A Regular GPU Custom Operator
+
+Since most of the building blocks for registering GPU custom operators are the exactly same as CPU ones, the registering function for an operator supporting both GPU and CPU will look like this:
+
+```c++
+ REGISTER_OP(my_op_gpu)
+ ...
+ .setForward(forwardCPU, "cpu")
+ .setForward(forwardGPU, "gpu")
+ .setBackward(backwardCPU, "cpu")
+ .setBackward(backwardGPU, "gpu");
+```
+
+Note that you can write only the GPU funcitons, as operators don’t have to support both CPU and GPU functions (can be GPU only).
+
+After you register forward or backward functions with context name `“gpu”`, MXNet will dispatch forward or backward GPU functions you just registered when the operator is invoked with GPU context.
+
+In the registered `forwardGPU` function, specify the grid and block size and launch the CUDA kernel. In the GPU operators input and output tensors are pre-allocated on the GPU memory, just like in the CPU operator tensors are pre-allocated on the CPU memory. As a result, you don’t need to call `cudaMemcpy` to move the tensor data to GPU device.
+
+```c++
+ MXReturnValue forwardGPU(std::map<std::string, std::string> attrs,
+ std::vector<MXTensor> inputs,
+ std::vector<MXTensor> outputs,
+ OpResource res) {
+ float* in_data = inputs[0].data<float>();
+ float* out_data = outputs[0].data<float>();
+ mx_stream_t cuda_stream = res.get_cuda_stream();
+ int64_t N = inputs[0].size();
+ int block = 256;
+ int grid = (N + (block - 1)) / block;
+ my_op_gpu_forward<<<grid,block,0,cuda_stream>>>(out_data, in_data, N);
+ return MX_SUCCESS;
+ }
+```
+
+Note that the `cuda_stream` object used for launching kernels is passed from MXNet backend via `OpResource` object. See below for details of `Operator Resource`.
+
+At this point all the attribute functions for each operator (`parseAttrs`, `inferShape`, etc.) run on the CPU, including the `forwardGPU` function. The only part that will actually run on the GPU is the CUDA kernel function that is launched from `forwardGPU` function (ie. my_op_gpu_forward).
+
+```c++
+ __global__ void my_op_gpu_forward(float *out, float *in, int64_t N) {
+ int tid = blockIdx.x * blockDim.x + threadIdx.x;
Review comment:
I just made another pass of the doc, and I am taking sam's suggestion to make it very generic. I will take your suggestion for doing the improvement on the code example in the next PR, thanks for it!
----------------------------------------------------------------
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
With regards,
Apache Git Services
[GitHub] [incubator-mxnet] samskalicky commented on a change in pull request
#17486: Update CustomOp doc with changes for GPU support
Posted by GitBox <gi...@apache.org>.
samskalicky commented on a change in pull request #17486: Update CustomOp doc with changes for GPU support
URL: https://github.com/apache/incubator-mxnet/pull/17486#discussion_r375022979
##########
File path: example/extensions/lib_custom_op/README.md
##########
@@ -56,87 +72,111 @@ For building a library containing your own custom operator, compose a C++ source
- `forward` - Forward Computation (can be replace with `createOpState`, see below for details)
Then compile it to `libmyop_lib.so` dynamic library using the following command:
+
```bash
g++ -shared -fPIC -std=c++11 myop_lib.cc -o libmyop_lib.so -I ../../../include/mxnet
```
+If you don't want to download MXNet source and choose only using `lib_api.h` header, you can copy the header over to the same folder of `myop_lib.cc` and run:
+
+```bash
+g++ -shared -fPIC -std=c++11 myop_lib.cc -o libmyop_lib.so
+```
+
Finally, you can write a Python script to load the library and run your custom operator:
+
```python
import mxnet as mx
mx.library.load(‘libmyop_lib.so’)
mx.nd.my_op(...)
```
-### Writing Regular Custom Operator:
+### Writing A Regular Custom Operator
-There are several essential building blocks for making a (stateless) custom operator:
+There are several essential building blocks for making a custom operator:
* [initialize](./gemm_lib.cc#L227):
* This function is the library initialization function necessary for any dynamic libraries. It checks if you are using a compatible version of MXNet. Note that this `version` parameter is passed from MXNet when library is loaded.
- MXReturnValue initialize(int version)
+```c++
+ MXReturnValue initialize(int version)
+```
* [parseAttrs](./gemm_lib.cc#L118):
* This function specifies number of input and output tensors for the custom operator; also this is where a custom operator can validate the attributes (ie. options) specified by the user.
- MXReturnValue parseAttrs(
- std::map<std::string,
- std::string> attrs,
- int* num_in,
- int* num_out)
+```c++
+ MXReturnValue parseAttrs(
+ std::map<std::string,
+ std::string> attrs,
+ int* num_in,
+ int* num_out)
+```
* [inferType](./gemm_lib.cc#L124):
* This function specifies how the custom operator infers output data types using input data types.
- MXReturnValue inferType(
- std::map<std::string, std::string> attrs,
- std::vector<int> &intypes,
- std::vector<int> &outtypes)
+```c++
+ MXReturnValue inferType(
+ std::map<std::string, std::string> attrs,
+ std::vector<int> &intypes,
+ std::vector<int> &outtypes)
+```
* [inferShape](./gemm_lib.cc#L143):
* This function specifies how the custom operator infers output tensor shape using input shape.
- MXReturnValue inferShape(
- std::map<std::string, std::string> attrs,
- std::vector<std::vector<unsigned int>> &inshapes,
- std::vector<std::vector<unsigned int>> &outshapes)
+```c++
+ MXReturnValue inferShape(
+ std::map<std::string, std::string> attrs,
+ std::vector<std::vector<unsigned int>> &inshapes,
+ std::vector<std::vector<unsigned int>> &outshapes)
+```
* [forward](./gemm_lib.cc#L56):
* This function specifies the computation of the forward pass of the operator.
- MXReturnValue forward(
- std::map<std::string, std::string> attrs,
- std::vector<MXTensor> inputs,
- std::vector<MXTensor> outputs,
- OpResource res)
+```c++
+ MXReturnValue forward(
+ std::map<std::string, std::string> attrs,
+ std::vector<MXTensor> inputs,
+ std::vector<MXTensor> outputs,
+ OpResource res)
+```
* [REGISTER_OP(my_op_name)](./gemm_lib.cc#L169):
- * This macro registers the custom operator and its properties to MXNet NDArray and Symbol APIs by its name.
-
- REGISTER_OP(my_op_name)
- .setForward(forward)
- .setParseAttrs(parseAttrs)
- .setInferType(inferType)
- .setInferShape(inferShape);
+ * This macro registers the custom operator and its properties to MXNet NDArray and Symbol APIs by its name. Note that for operator running on CPU, you need pass the name of the context `"cpu"` when registering forward or backward function.
+```c++
+ REGISTER_OP(my_op_name)
+ .setForward(forward, "cpu")
+ .setBackward(backward, "cpu")
Review comment:
remove backward
----------------------------------------------------------------
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
With regards,
Apache Git Services
[GitHub] [incubator-mxnet] samskalicky commented on a change in pull request
#17486: Update CustomOp doc with changes for GPU support
Posted by GitBox <gi...@apache.org>.
samskalicky commented on a change in pull request #17486: Update CustomOp doc with changes for GPU support
URL: https://github.com/apache/incubator-mxnet/pull/17486#discussion_r375019792
##########
File path: example/extensions/lib_custom_op/README.md
##########
@@ -56,87 +72,111 @@ For building a library containing your own custom operator, compose a C++ source
- `forward` - Forward Computation (can be replace with `createOpState`, see below for details)
Review comment:
"can be replace" --> "can be replaced"
----------------------------------------------------------------
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
With regards,
Apache Git Services
[GitHub] [incubator-mxnet] wkcn edited a comment on issue #17486: Update
CustomOp doc with changes for GPU support
Posted by GitBox <gi...@apache.org>.
wkcn edited a comment on issue #17486: Update CustomOp doc with changes for GPU support
URL: https://github.com/apache/incubator-mxnet/pull/17486#issuecomment-586661767
The PR has been merged. Thanks for your contribution! : )
----------------------------------------------------------------
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
With regards,
Apache Git Services
[GitHub] [incubator-mxnet] a550461053 commented on issue #17486: Update
CustomOp doc with changes for GPU support
Posted by GitBox <gi...@apache.org>.
a550461053 commented on issue #17486: Update CustomOp doc with changes for GPU support
URL: https://github.com/apache/incubator-mxnet/pull/17486#issuecomment-586058857
> > Ok, thank you. I want to create a custom operator calling another operator which input NDArray. Both operator is async pushed to engine, I think this way is possible. Also, I am not clear on using mutateInputs API, how and when to use it? If you can provide a example of the mutateInputs API, I will be grateful to you~
>
> Hi @a550461053 currently the custom operator design focused on separating the custom operator from the MXNet backend source code complexity. This means that your custom operator can (must) be entirely separated from MXNet. So you cannot call a regular built-in MXNet operator from your custom operator. We have an item here #17006 for adding support in the future to be able to do this, but it is not implemented yet.
>
> As for how to use mutateInputs, it works exactly as the doc describes:
>
> > This function allows you to mark some inputs to be mutable inputs. It is useful when using aux parameters for BatchNorm-like operators.
>
> So lets say you have an operator with 5 inputs, you can mark the indices of the inputs that you want to be mutable like this (for example mark the last two inputs as mutable):
>
> ```
> MXReturnValue batchNorm_mutateInputs(std::map<std::string, std::string> attrs,
> std::vector<int> &input_indices) {
> // mark mutable inputs
> input_indices.push_back(3);
> input_indices.push_back(4);
> return MX_SUCCESS;
> }
> ```
Thank you and @rondogency . I am clear on use of mutateInputs API.
But now I can only build mxnet to create c++ customOp and I have a problem on compiling into python wheel at #17577
----------------------------------------------------------------
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
With regards,
Apache Git Services
[GitHub] [incubator-mxnet] samskalicky commented on a change in pull request
#17486: Update CustomOp doc with changes for GPU support
Posted by GitBox <gi...@apache.org>.
samskalicky commented on a change in pull request #17486: Update CustomOp doc with changes for GPU support
URL: https://github.com/apache/incubator-mxnet/pull/17486#discussion_r375019642
##########
File path: example/extensions/lib_custom_op/README.md
##########
@@ -28,26 +28,42 @@ Custom operators (CustomOp) enable users to write new operators without compilin
### Have MXNet Ready
-First you should install MXNet either from compiling from source code or download from nightly build. It doesn’t matter if the build comes with CUDA or MKLDNN. The custom operator doesn’t interact with the execution of other native MXNet operators.
+Custom Operator support was merged (#15921, #17270) and is not available in versions of MXNet prior to v1.7.0. To access the feature now, please install MXNet by compiling from source using master or using the previously mentioned commits, downloading one of the nightly builds, or from a release of MXNet 1.7.0+. For running the following example, it doesn’t matter if it is a CUDA, MKLDNN or plain MXNet build; the custom operator doesn’t interact with the execution of other native MXNet operators. Note that if you want to run GPU examples and write your custom operators running on GPU, you still need an MXNet CUDA build.
-### Run An Example:
+### Run An Example
-You can start getting familiar with custom operators by running some examples provided in the **example/extensions/lib_custom_op** directory. Start with a common linear algebra operator like `gemm` (Generalized Matrix Multiplication). Go to `lib_custom_op` directory and follow these steps:
+You can start getting familiar with custom operators by running some examples provided in the `example/extensions/lib_custom_op` directory. Start with a common linear algebra operator like `gemm` (Generalized Matrix Multiplication). Go to `lib_custom_op` directory and follow these steps:
1. Run `make gemm_lib`. The Makefile will generate a dynamic library **libgemm_lib.so** compiled from `gemm_lib.cc`. This is the library you are going to load that contains everything for the custom gemm operator.
-2. Run `python test_gemm.py`. It’ll first load the above .so library, find the operators, register them in the MXNet backend, print "Found x operators", then invoke the operator like a regular MXNet operator and output the result.
+2. Run `python test_gemm.py`. It’ll first load the library compiled from step 1, find the operators, register them in the MXNet backend, then invoke the operator like a regular MXNet operator and output the result.
-### Basic Files For Gemm Library:
+```
+[19:22:02] ../src/c_api/c_api.cc:286: Found 2 operators in library
+[19:22:02] ../src/c_api/c_api.cc:350: Op[0] my_gemm
+[19:22:02] ../src/c_api/c_api.cc:350: Op[1] state_gemm
+[19:22:02] ../src/c_api/c_api.cc:785: Found 0 partitioners in library
+--------start ndarray compute---------
+[[ 50.]
+ [122.]]
+<NDArray 2x1 @cpu(0)>
+...
+```
+
+Note that you can safely ignore the `Found 0 partitioners` info as it is not related to the custom operator.
+
+### Basic Files For A GeMM Library
* **lib_custom_op/gemm_lib.cc**: This file has a source code implementation of all required components of a custom operator, as well as the registration of the custom operator.
-* **lib_custom_op/Makefile**: Compile source code to a dynamic shared library, with a header file `include/mxnet/lib_api.h` from MXNet source code. Currently the custom operator is compatible with C++11 onwards.
+* **lib_custom_op/Makefile**: Compile `gemm_lib.cc` to a dynamic shared library named `libgemm_lib.so`, with a include path of the header file `include/mxnet/lib_api.h` from MXNet source code. Currently the custom operator is compatible with C++11 onwards.
* **lib_custom_op/test_gemm.py**: This file calls `mx.library.load(‘libgemm_lib.so’)` to load the library containing the custom operator, invokes the operator using both NDArray and Symbol APIs, and prints outputs of the forward and backward passes. The outputs should be the same as the regular MXNet `gemm` operator.
-## Writing Custom Operator Library:
+* **include/mxnet/lib_api.h**: This file from MXNet source code is the one-stop header to include all necessary data types and function prototypes for writing a custom operator library. You can either specify the include path in `Makefile`, or copy the header file over to `example/extensions/lib_custom_op` folder for `gemm_lib.cc` to include. Note that apart from this header, the custom operator library is independent of MXNet source.
-For building a library containing your own custom operator, compose a C++ source file like `myop_lib.cc`, include `lib_api.h` header file, and write your custom operator implementation with these essential functions:
+## Writing A Custom CPU Operator Library
+
+For building your own library containing custom CPU operator, compose a C++ source file like `myop_lib.cc`, include `lib_api.h` header file, and write your custom operator implementation with these essential functions:
Review comment:
"For building your own" --> "To build your own"
"implementation with these essential functions" --> "implementation with these required functions:"
----------------------------------------------------------------
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
With regards,
Apache Git Services
[GitHub] [incubator-mxnet] a550461053 commented on issue #17486: Update
CustomOp doc with changes for GPU support
Posted by GitBox <gi...@apache.org>.
a550461053 commented on issue #17486: Update CustomOp doc with changes for GPU support
URL: https://github.com/apache/incubator-mxnet/pull/17486#issuecomment-585356310
> > Hi, @rondogency I am trying to use the dynamic operator and have some questions.
> > How can we use `mxnet::NDArray` or `mshadow_op` used in xxx-inl.h of mxnet-backend operator? And I want to know how to use the `mutateInputs` function like in BatchNorm-like operators.
> > Thank you very much.
>
> Hi @a550461053, MXNet custom operators have been designed not to integrate with any MXNet backend classes (mxnet::NDArray) to simplify building. We've designed MXTensor to be equivalent. Is there some other functionality you were wanting from NDArray thats not in MXTensor?
>
> The mutateInputs API allows you to do the same as the backend APIs, so you can specify the index of the input that you want to be mutable to make a BatchNorm-like operator.
Ok, thank you. I want to create a custom operator calling another operator which input NDArray. Both operator is async pushed to engine, I think this way is possible. Also, I am not clear on using mutateInputs API, how and when to use it? If you can provide a example of the mutateInputs API, I will be grateful to you~
----------------------------------------------------------------
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
With regards,
Apache Git Services
[GitHub] [incubator-mxnet] wkcn merged pull request #17486: Update CustomOp
doc with changes for GPU support
Posted by GitBox <gi...@apache.org>.
wkcn merged pull request #17486: Update CustomOp doc with changes for GPU support
URL: https://github.com/apache/incubator-mxnet/pull/17486
----------------------------------------------------------------
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
With regards,
Apache Git Services
[GitHub] [incubator-mxnet] samskalicky commented on a change in pull request
#17486: Update CustomOp doc with changes for GPU support
Posted by GitBox <gi...@apache.org>.
samskalicky commented on a change in pull request #17486: Update CustomOp doc with changes for GPU support
URL: https://github.com/apache/incubator-mxnet/pull/17486#discussion_r375023931
##########
File path: example/extensions/lib_custom_op/README.md
##########
@@ -146,21 +186,160 @@ Let’s take a closer look at those registry functions:
* **inferShape**: This function is similar to the `inferType` function, except it is used for populating the output data shapes. You need to figure out the shapes of each output tensors for this computation. For example, if the inputs are images with shape (224,224,3) and you write a padding operator to make 10px borders for the images, then your output shape will be (234,234,3).
-* **forward**: This function executes the main forward computation. It takes four arguments. The 1st argument is the attributes. The 2nd argument is the input `MXTensors` which stores all data and info of input ndarrays. The 3rd argument is the output `MXTensors`. The 4th argument is the `OpResource` object for memory allocation and other utilities. Additionally, you can use a `dltensor` tensor structure stored in the `MXTensor` as a more standardized data structure for computing.
+* **forward**: This function executes the main forward computation. It takes four arguments. The 1st argument is the attributes. The 2nd argument is the input `MXTensors` which stores all data and info of input ndarrays. The 3rd argument is the output `MXTensors`. The 4th argument is the `OpResource` object for memory allocation and other utilities. The details of `OpResource` are covered in the below section.
+Additionally, you can use a `dltensor` tensor structure stored in the `MXTensor` as a more standardized data structure for computing.
* **backward**: This function is doing the backward gradient computation. It will be similar to the forward function. And you need to figure out the formula of the backward gradient computation.
* **mutateInputs**: This function is for marking mutable inputs. It takes two arguments. The 1st argument is the attributes. The 2nd argument is a list of input indices that are mutable among all input tensors. It is useful when some inputs are auxiliary model parameters and might be altered during forward/backward computation. Remember, the index number of `input_indices` should not exceed the number of inputs.
-### Writing Stateful Custom Operator:
+### Writing Stateful Custom Operator
A stateful custom operator is useful when a forward/backward call needs some data or ‘state’ from previous forward/backward calls. Normally we create a class, and make instance variables store the states used for computing or caching.
Most of the building blocks for making a stateful custom operator is the same as regular custom operator, except it’ll register `createOpState` instead of a `forward` function for the computation.
* [createOpState](./gemm_lib.cc#L204) - Create stateful operator instance:
* This function takes two arguments. The 1st argument is attributes. The 2nd argument is a placeholder for `CustomStatefulOp` object. You must [define a class that inherits CustomStatefulOp](./gemm_lib.cc#L178) and override the forward function (optionally the backward function). Then you need to create an instance of your class and assign it to the placeholder. In this way, all of the forward/backward calls will use the same methods in that instance, and the instance is able to keep the state of the operator.
+```c++
+ MXReturnValue createOpState(
+ std::map<std::string, std::string> attrs,
+ CustomStatefulOp** op_inst)
+```
+
+* The operator registering function will look like this:
+```c++
+ REGISTER_OP(my_state_op)
+ ...
+ .setCreateOpState(createOpState, "cpu");
+```
+
+## Writing A Custom GPU Operator Library
+
+Most of the building blocks for registering GPU custom operators are the exactly same as CPU ones, except you need to specify the `"gpu"` context name when registering `forward`, `backward` or `createOpState` function.
+
+### Run A GPU Example
+
+For illustration purposes, we provided a `ReLU` (Rectified Linear Unit) activation operator that can run on GPU. Make sure you have installed a CUDA compatible MXNet build. Go to `lib_custom_op` directory and follow these steps:
+
+1. Run `make relu_lib`. The Makefile will invoke `NVCC` compiler to compile the CUDA kernel along with regular custom operator functions from `relu_lib.cu` to generate `librelu_lib.so` library.
+2. Run `python test_relu.py`. It’ll register the GPU `ReLU` operator in the MXNet backend, then invoke the operator by feeding an `NDArray` input with GPU context, and output the result tensor with GPU context.
Review comment:
feeding --> passing
----------------------------------------------------------------
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
With regards,
Apache Git Services
[GitHub] [incubator-mxnet] samskalicky commented on issue #17486: Update
CustomOp doc with changes for GPU support
Posted by GitBox <gi...@apache.org>.
samskalicky commented on issue #17486: Update CustomOp doc with changes for GPU support
URL: https://github.com/apache/incubator-mxnet/pull/17486#issuecomment-586650041
@aaronmarkham or @wkcn can we merge this please?
----------------------------------------------------------------
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
With regards,
Apache Git Services
[GitHub] [incubator-mxnet] samskalicky commented on a change in pull request
#17486: Update CustomOp doc with changes for GPU support
Posted by GitBox <gi...@apache.org>.
samskalicky commented on a change in pull request #17486: Update CustomOp doc with changes for GPU support
URL: https://github.com/apache/incubator-mxnet/pull/17486#discussion_r375025352
##########
File path: example/extensions/lib_custom_op/README.md
##########
@@ -146,21 +186,160 @@ Let’s take a closer look at those registry functions:
* **inferShape**: This function is similar to the `inferType` function, except it is used for populating the output data shapes. You need to figure out the shapes of each output tensors for this computation. For example, if the inputs are images with shape (224,224,3) and you write a padding operator to make 10px borders for the images, then your output shape will be (234,234,3).
-* **forward**: This function executes the main forward computation. It takes four arguments. The 1st argument is the attributes. The 2nd argument is the input `MXTensors` which stores all data and info of input ndarrays. The 3rd argument is the output `MXTensors`. The 4th argument is the `OpResource` object for memory allocation and other utilities. Additionally, you can use a `dltensor` tensor structure stored in the `MXTensor` as a more standardized data structure for computing.
+* **forward**: This function executes the main forward computation. It takes four arguments. The 1st argument is the attributes. The 2nd argument is the input `MXTensors` which stores all data and info of input ndarrays. The 3rd argument is the output `MXTensors`. The 4th argument is the `OpResource` object for memory allocation and other utilities. The details of `OpResource` are covered in the below section.
+Additionally, you can use a `dltensor` tensor structure stored in the `MXTensor` as a more standardized data structure for computing.
* **backward**: This function is doing the backward gradient computation. It will be similar to the forward function. And you need to figure out the formula of the backward gradient computation.
* **mutateInputs**: This function is for marking mutable inputs. It takes two arguments. The 1st argument is the attributes. The 2nd argument is a list of input indices that are mutable among all input tensors. It is useful when some inputs are auxiliary model parameters and might be altered during forward/backward computation. Remember, the index number of `input_indices` should not exceed the number of inputs.
-### Writing Stateful Custom Operator:
+### Writing Stateful Custom Operator
A stateful custom operator is useful when a forward/backward call needs some data or ‘state’ from previous forward/backward calls. Normally we create a class, and make instance variables store the states used for computing or caching.
Most of the building blocks for making a stateful custom operator is the same as regular custom operator, except it’ll register `createOpState` instead of a `forward` function for the computation.
* [createOpState](./gemm_lib.cc#L204) - Create stateful operator instance:
* This function takes two arguments. The 1st argument is attributes. The 2nd argument is a placeholder for `CustomStatefulOp` object. You must [define a class that inherits CustomStatefulOp](./gemm_lib.cc#L178) and override the forward function (optionally the backward function). Then you need to create an instance of your class and assign it to the placeholder. In this way, all of the forward/backward calls will use the same methods in that instance, and the instance is able to keep the state of the operator.
+```c++
+ MXReturnValue createOpState(
+ std::map<std::string, std::string> attrs,
+ CustomStatefulOp** op_inst)
+```
+
+* The operator registering function will look like this:
+```c++
+ REGISTER_OP(my_state_op)
+ ...
+ .setCreateOpState(createOpState, "cpu");
+```
+
+## Writing A Custom GPU Operator Library
+
+Most of the building blocks for registering GPU custom operators are the exactly same as CPU ones, except you need to specify the `"gpu"` context name when registering `forward`, `backward` or `createOpState` function.
+
+### Run A GPU Example
+
+For illustration purposes, we provided a `ReLU` (Rectified Linear Unit) activation operator that can run on GPU. Make sure you have installed a CUDA compatible MXNet build. Go to `lib_custom_op` directory and follow these steps:
+
+1. Run `make relu_lib`. The Makefile will invoke `NVCC` compiler to compile the CUDA kernel along with regular custom operator functions from `relu_lib.cu` to generate `librelu_lib.so` library.
+2. Run `python test_relu.py`. It’ll register the GPU `ReLU` operator in the MXNet backend, then invoke the operator by feeding an `NDArray` input with GPU context, and output the result tensor with GPU context.
+
+### Writing A Regular GPU Custom Operator
+
+Since most of the building blocks for registering GPU custom operators are the exactly same as CPU ones, the registering function for an operator supporting both GPU and CPU will look like this:
+
+```c++
+ REGISTER_OP(my_op_gpu)
+ ...
+ .setForward(forwardCPU, "cpu")
+ .setForward(forwardGPU, "gpu")
+ .setBackward(backwardCPU, "cpu")
+ .setBackward(backwardGPU, "gpu");
+```
+
+Note that you can write only the GPU funcitons, as operators don’t have to support both CPU and GPU functions (can be GPU only).
+
+After you register forward or backward functions with context name `“gpu”`, MXNet will dispatch forward or backward GPU functions you just registered when the operator is invoked with GPU context.
+
+In the registered `forwardGPU` function, specify the grid and block size and launch the CUDA kernel. In the GPU operators input and output tensors are pre-allocated on the GPU memory, just like in the CPU operator tensors are pre-allocated on the CPU memory. As a result, you don’t need to call `cudaMemcpy` to move the tensor data to GPU device.
Review comment:
In the `forwardGPU` function, you'll specify the grid and block size and launch your CUDA kernel. MXNet pre-allocates the memory for input and output tensors on the GPU, just like for CPU operators tensors are pre-allocated on the CPU. As a result, you don’t need to call `cudaMemcpy` to move the tensor data to the GPU device.
----------------------------------------------------------------
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
With regards,
Apache Git Services
[GitHub] [incubator-mxnet] samskalicky commented on issue #17486: Update
CustomOp doc with changes for GPU support
Posted by GitBox <gi...@apache.org>.
samskalicky commented on issue #17486: Update CustomOp doc with changes for GPU support
URL: https://github.com/apache/incubator-mxnet/pull/17486#issuecomment-580872254
@rondogency please add the fix for the gpu example test program here from https://github.com/apache/incubator-mxnet/pull/17270#discussion_r372603452
----------------------------------------------------------------
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
With regards,
Apache Git Services
[GitHub] [incubator-mxnet] wkcn commented on a change in pull request
#17486: Update CustomOp doc with changes for GPU support
Posted by GitBox <gi...@apache.org>.
wkcn commented on a change in pull request #17486: Update CustomOp doc with changes for GPU support
URL: https://github.com/apache/incubator-mxnet/pull/17486#discussion_r375044803
##########
File path: example/extensions/lib_custom_op/README.md
##########
@@ -146,21 +186,160 @@ Let’s take a closer look at those registry functions:
* **inferShape**: This function is similar to the `inferType` function, except it is used for populating the output data shapes. You need to figure out the shapes of each output tensors for this computation. For example, if the inputs are images with shape (224,224,3) and you write a padding operator to make 10px borders for the images, then your output shape will be (234,234,3).
-* **forward**: This function executes the main forward computation. It takes four arguments. The 1st argument is the attributes. The 2nd argument is the input `MXTensors` which stores all data and info of input ndarrays. The 3rd argument is the output `MXTensors`. The 4th argument is the `OpResource` object for memory allocation and other utilities. Additionally, you can use a `dltensor` tensor structure stored in the `MXTensor` as a more standardized data structure for computing.
+* **forward**: This function executes the main forward computation. It takes four arguments. The 1st argument is the attributes. The 2nd argument is the input `MXTensors` which stores all data and info of input ndarrays. The 3rd argument is the output `MXTensors`. The 4th argument is the `OpResource` object for memory allocation and other utilities. The details of `OpResource` are covered in the below section.
+Additionally, you can use a `dltensor` tensor structure stored in the `MXTensor` as a more standardized data structure for computing.
* **backward**: This function is doing the backward gradient computation. It will be similar to the forward function. And you need to figure out the formula of the backward gradient computation.
* **mutateInputs**: This function is for marking mutable inputs. It takes two arguments. The 1st argument is the attributes. The 2nd argument is a list of input indices that are mutable among all input tensors. It is useful when some inputs are auxiliary model parameters and might be altered during forward/backward computation. Remember, the index number of `input_indices` should not exceed the number of inputs.
-### Writing Stateful Custom Operator:
+### Writing Stateful Custom Operator
A stateful custom operator is useful when a forward/backward call needs some data or ‘state’ from previous forward/backward calls. Normally we create a class, and make instance variables store the states used for computing or caching.
Most of the building blocks for making a stateful custom operator is the same as regular custom operator, except it’ll register `createOpState` instead of a `forward` function for the computation.
* [createOpState](./gemm_lib.cc#L204) - Create stateful operator instance:
* This function takes two arguments. The 1st argument is attributes. The 2nd argument is a placeholder for `CustomStatefulOp` object. You must [define a class that inherits CustomStatefulOp](./gemm_lib.cc#L178) and override the forward function (optionally the backward function). Then you need to create an instance of your class and assign it to the placeholder. In this way, all of the forward/backward calls will use the same methods in that instance, and the instance is able to keep the state of the operator.
+```c++
+ MXReturnValue createOpState(
+ std::map<std::string, std::string> attrs,
+ CustomStatefulOp** op_inst)
+```
+
+* The operator registering function will look like this:
+```c++
+ REGISTER_OP(my_state_op)
+ ...
+ .setCreateOpState(createOpState, "cpu");
+```
+
+## Writing A Custom GPU Operator Library
+
+Most of the building blocks for registering GPU custom operators are the exactly same as CPU ones, except you need to specify the `"gpu"` context name when registering `forward`, `backward` or `createOpState` function.
+
+### Run A GPU Example
+
+For illustration purposes, we provided a `ReLU` (Rectified Linear Unit) activation operator that can run on GPU. Make sure you have installed a CUDA compatible MXNet build. Go to `lib_custom_op` directory and follow these steps:
+
+1. Run `make relu_lib`. The Makefile will invoke `NVCC` compiler to compile the CUDA kernel along with regular custom operator functions from `relu_lib.cu` to generate `librelu_lib.so` library.
+2. Run `python test_relu.py`. It’ll register the GPU `ReLU` operator in the MXNet backend, then invoke the operator by feeding an `NDArray` input with GPU context, and output the result tensor with GPU context.
+
+### Writing A Regular GPU Custom Operator
+
+Since most of the building blocks for registering GPU custom operators are the exactly same as CPU ones, the registering function for an operator supporting both GPU and CPU will look like this:
+
+```c++
+ REGISTER_OP(my_op_gpu)
+ ...
+ .setForward(forwardCPU, "cpu")
+ .setForward(forwardGPU, "gpu")
+ .setBackward(backwardCPU, "cpu")
+ .setBackward(backwardGPU, "gpu");
+```
+
+Note that you can write only the GPU funcitons, as operators don’t have to support both CPU and GPU functions (can be GPU only).
+
+After you register forward or backward functions with context name `“gpu”`, MXNet will dispatch forward or backward GPU functions you just registered when the operator is invoked with GPU context.
+
+In the registered `forwardGPU` function, specify the grid and block size and launch the CUDA kernel. In the GPU operators input and output tensors are pre-allocated on the GPU memory, just like in the CPU operator tensors are pre-allocated on the CPU memory. As a result, you don’t need to call `cudaMemcpy` to move the tensor data to GPU device.
+
+```c++
+ MXReturnValue forwardGPU(std::map<std::string, std::string> attrs,
+ std::vector<MXTensor> inputs,
+ std::vector<MXTensor> outputs,
+ OpResource res) {
+ float* in_data = inputs[0].data<float>();
+ float* out_data = outputs[0].data<float>();
+ mx_stream_t cuda_stream = res.get_cuda_stream();
+ int64_t N = inputs[0].size();
+ int block = 256;
+ int grid = (N + (block - 1)) / block;
+ my_op_gpu_forward<<<grid,block,0,cuda_stream>>>(out_data, in_data, N);
+ return MX_SUCCESS;
+ }
+```
+
+Note that the `cuda_stream` object used for launching kernels is passed from MXNet backend via `OpResource` object. See below for details of `Operator Resource`.
+
+At this point all the attribute functions for each operator (`parseAttrs`, `inferShape`, etc.) run on the CPU, including the `forwardGPU` function. The only part that will actually run on the GPU is the CUDA kernel function that is launched from `forwardGPU` function (ie. my_op_gpu_forward).
+
+```c++
+ __global__ void my_op_gpu_forward(float *out, float *in, int64_t N) {
+ int tid = blockIdx.x * blockDim.x + threadIdx.x;
Review comment:
I have a question about the balance between code readability and correctness.
In CUDA code, the number of block and thread is limited. This code works for small N, but it doesn't work when `N` is larger than the maximum number of block and thread. However, it is good to read for users.
The correct code shows as below, but it is unfriendly to read for newbies.
```c++
for (int64_t i = blockIdx.x * blockDim.x + threadIdx.x; i < N; i += blockDim.x * gridDim.x) {
// code your computational logic here
}
```
I'm not sure which one is better.
----------------------------------------------------------------
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
With regards,
Apache Git Services
[GitHub] [incubator-mxnet] samskalicky commented on issue #17486: Update
CustomOp doc with changes for GPU support
Posted by GitBox <gi...@apache.org>.
samskalicky commented on issue #17486: Update CustomOp doc with changes for GPU support
URL: https://github.com/apache/incubator-mxnet/pull/17486#issuecomment-584816097
> Hi, @rondogency I am trying to use the dynamic operator and have some questions.
> How can we use `mxnet::NDArray` or `mshadow_op` used in xxx-inl.h of mxnet-backend operator? And I want to know how to use the `mutateInputs` function like in BatchNorm-like operators.
> Thank you very much.
Hi @a550461053, MXNet custom operators have been designed not to integrate with any MXNet backend classes (mxnet::NDArray) to simplify building. We've designed MXTensor to be equivalent. Is there some other functionality you were wanting from NDArray thats not in MXTensor?
The mutateInputs API allows you to do the same as the backend APIs, so you can specify the index of the input that you want to be mutable to make a BatchNorm-like operator.
----------------------------------------------------------------
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
With regards,
Apache Git Services
[GitHub] [incubator-mxnet] samskalicky commented on a change in pull request
#17486: Update CustomOp doc with changes for GPU support
Posted by GitBox <gi...@apache.org>.
samskalicky commented on a change in pull request #17486: Update CustomOp doc with changes for GPU support
URL: https://github.com/apache/incubator-mxnet/pull/17486#discussion_r375020045
##########
File path: example/extensions/lib_custom_op/README.md
##########
@@ -56,87 +72,111 @@ For building a library containing your own custom operator, compose a C++ source
- `forward` - Forward Computation (can be replace with `createOpState`, see below for details)
Then compile it to `libmyop_lib.so` dynamic library using the following command:
+
```bash
g++ -shared -fPIC -std=c++11 myop_lib.cc -o libmyop_lib.so -I ../../../include/mxnet
```
+If you don't want to download MXNet source and choose only using `lib_api.h` header, you can copy the header over to the same folder of `myop_lib.cc` and run:
Review comment:
"choose only using" --> "choose to only use"
----------------------------------------------------------------
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
With regards,
Apache Git Services
[GitHub] [incubator-mxnet] samskalicky commented on a change in pull request
#17486: Update CustomOp doc with changes for GPU support
Posted by GitBox <gi...@apache.org>.
samskalicky commented on a change in pull request #17486: Update CustomOp doc with changes for GPU support
URL: https://github.com/apache/incubator-mxnet/pull/17486#discussion_r375018587
##########
File path: example/extensions/lib_custom_op/README.md
##########
@@ -28,26 +28,42 @@ Custom operators (CustomOp) enable users to write new operators without compilin
### Have MXNet Ready
-First you should install MXNet either from compiling from source code or download from nightly build. It doesn’t matter if the build comes with CUDA or MKLDNN. The custom operator doesn’t interact with the execution of other native MXNet operators.
+Custom Operator support was merged (#15921, #17270) and is not available in versions of MXNet prior to v1.7.0. To access the feature now, please install MXNet by compiling from source using master or using the previously mentioned commits, downloading one of the nightly builds, or from a release of MXNet 1.7.0+. For running the following example, it doesn’t matter if it is a CUDA, MKLDNN or plain MXNet build; the custom operator doesn’t interact with the execution of other native MXNet operators. Note that if you want to run GPU examples and write your custom operators running on GPU, you still need an MXNet CUDA build.
-### Run An Example:
+### Run An Example
-You can start getting familiar with custom operators by running some examples provided in the **example/extensions/lib_custom_op** directory. Start with a common linear algebra operator like `gemm` (Generalized Matrix Multiplication). Go to `lib_custom_op` directory and follow these steps:
+You can start getting familiar with custom operators by running some examples provided in the `example/extensions/lib_custom_op` directory. Start with a common linear algebra operator like `gemm` (Generalized Matrix Multiplication). Go to `lib_custom_op` directory and follow these steps:
1. Run `make gemm_lib`. The Makefile will generate a dynamic library **libgemm_lib.so** compiled from `gemm_lib.cc`. This is the library you are going to load that contains everything for the custom gemm operator.
-2. Run `python test_gemm.py`. It’ll first load the above .so library, find the operators, register them in the MXNet backend, print "Found x operators", then invoke the operator like a regular MXNet operator and output the result.
+2. Run `python test_gemm.py`. It’ll first load the library compiled from step 1, find the operators, register them in the MXNet backend, then invoke the operator like a regular MXNet operator and output the result.
-### Basic Files For Gemm Library:
+```
+[19:22:02] ../src/c_api/c_api.cc:286: Found 2 operators in library
+[19:22:02] ../src/c_api/c_api.cc:350: Op[0] my_gemm
+[19:22:02] ../src/c_api/c_api.cc:350: Op[1] state_gemm
+[19:22:02] ../src/c_api/c_api.cc:785: Found 0 partitioners in library
+--------start ndarray compute---------
+[[ 50.]
+ [122.]]
+<NDArray 2x1 @cpu(0)>
+...
+```
+
+Note that you can safely ignore the `Found 0 partitioners` info as it is not related to the custom operator.
+
+### Basic Files For A GeMM Library
* **lib_custom_op/gemm_lib.cc**: This file has a source code implementation of all required components of a custom operator, as well as the registration of the custom operator.
-* **lib_custom_op/Makefile**: Compile source code to a dynamic shared library, with a header file `include/mxnet/lib_api.h` from MXNet source code. Currently the custom operator is compatible with C++11 onwards.
+* **lib_custom_op/Makefile**: Compile `gemm_lib.cc` to a dynamic shared library named `libgemm_lib.so`, with a include path of the header file `include/mxnet/lib_api.h` from MXNet source code. Currently the custom operator is compatible with C++11 onwards.
* **lib_custom_op/test_gemm.py**: This file calls `mx.library.load(‘libgemm_lib.so’)` to load the library containing the custom operator, invokes the operator using both NDArray and Symbol APIs, and prints outputs of the forward and backward passes. The outputs should be the same as the regular MXNet `gemm` operator.
-## Writing Custom Operator Library:
+* **include/mxnet/lib_api.h**: This file from MXNet source code is the one-stop header to include all necessary data types and function prototypes for writing a custom operator library. You can either specify the include path in `Makefile`, or copy the header file over to `example/extensions/lib_custom_op` folder for `gemm_lib.cc` to include. Note that apart from this header, the custom operator library is independent of MXNet source.
Review comment:
"is the one-stop header to include all necessary" --> "is the single header file needed to include all necessary"
"include path in `Makefile`" --> "include path in the `Makefile`"
"`example/extensions/lib_custom_op` folder for `gemm_lib.cc` to include. " --> "`example/extensions/lib_custom_op` folder. "
----------------------------------------------------------------
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
With regards,
Apache Git Services
[GitHub] [incubator-mxnet] samskalicky commented on a change in pull request
#17486: Update CustomOp doc with changes for GPU support
Posted by GitBox <gi...@apache.org>.
samskalicky commented on a change in pull request #17486: Update CustomOp doc with changes for GPU support
URL: https://github.com/apache/incubator-mxnet/pull/17486#discussion_r375024885
##########
File path: example/extensions/lib_custom_op/README.md
##########
@@ -146,21 +186,160 @@ Let’s take a closer look at those registry functions:
* **inferShape**: This function is similar to the `inferType` function, except it is used for populating the output data shapes. You need to figure out the shapes of each output tensors for this computation. For example, if the inputs are images with shape (224,224,3) and you write a padding operator to make 10px borders for the images, then your output shape will be (234,234,3).
-* **forward**: This function executes the main forward computation. It takes four arguments. The 1st argument is the attributes. The 2nd argument is the input `MXTensors` which stores all data and info of input ndarrays. The 3rd argument is the output `MXTensors`. The 4th argument is the `OpResource` object for memory allocation and other utilities. Additionally, you can use a `dltensor` tensor structure stored in the `MXTensor` as a more standardized data structure for computing.
+* **forward**: This function executes the main forward computation. It takes four arguments. The 1st argument is the attributes. The 2nd argument is the input `MXTensors` which stores all data and info of input ndarrays. The 3rd argument is the output `MXTensors`. The 4th argument is the `OpResource` object for memory allocation and other utilities. The details of `OpResource` are covered in the below section.
+Additionally, you can use a `dltensor` tensor structure stored in the `MXTensor` as a more standardized data structure for computing.
* **backward**: This function is doing the backward gradient computation. It will be similar to the forward function. And you need to figure out the formula of the backward gradient computation.
* **mutateInputs**: This function is for marking mutable inputs. It takes two arguments. The 1st argument is the attributes. The 2nd argument is a list of input indices that are mutable among all input tensors. It is useful when some inputs are auxiliary model parameters and might be altered during forward/backward computation. Remember, the index number of `input_indices` should not exceed the number of inputs.
-### Writing Stateful Custom Operator:
+### Writing Stateful Custom Operator
A stateful custom operator is useful when a forward/backward call needs some data or ‘state’ from previous forward/backward calls. Normally we create a class, and make instance variables store the states used for computing or caching.
Most of the building blocks for making a stateful custom operator is the same as regular custom operator, except it’ll register `createOpState` instead of a `forward` function for the computation.
* [createOpState](./gemm_lib.cc#L204) - Create stateful operator instance:
* This function takes two arguments. The 1st argument is attributes. The 2nd argument is a placeholder for `CustomStatefulOp` object. You must [define a class that inherits CustomStatefulOp](./gemm_lib.cc#L178) and override the forward function (optionally the backward function). Then you need to create an instance of your class and assign it to the placeholder. In this way, all of the forward/backward calls will use the same methods in that instance, and the instance is able to keep the state of the operator.
+```c++
+ MXReturnValue createOpState(
+ std::map<std::string, std::string> attrs,
+ CustomStatefulOp** op_inst)
+```
+
+* The operator registering function will look like this:
+```c++
+ REGISTER_OP(my_state_op)
+ ...
+ .setCreateOpState(createOpState, "cpu");
+```
+
+## Writing A Custom GPU Operator Library
+
+Most of the building blocks for registering GPU custom operators are the exactly same as CPU ones, except you need to specify the `"gpu"` context name when registering `forward`, `backward` or `createOpState` function.
+
+### Run A GPU Example
+
+For illustration purposes, we provided a `ReLU` (Rectified Linear Unit) activation operator that can run on GPU. Make sure you have installed a CUDA compatible MXNet build. Go to `lib_custom_op` directory and follow these steps:
+
+1. Run `make relu_lib`. The Makefile will invoke `NVCC` compiler to compile the CUDA kernel along with regular custom operator functions from `relu_lib.cu` to generate `librelu_lib.so` library.
+2. Run `python test_relu.py`. It’ll register the GPU `ReLU` operator in the MXNet backend, then invoke the operator by feeding an `NDArray` input with GPU context, and output the result tensor with GPU context.
+
+### Writing A Regular GPU Custom Operator
+
+Since most of the building blocks for registering GPU custom operators are the exactly same as CPU ones, the registering function for an operator supporting both GPU and CPU will look like this:
+
+```c++
+ REGISTER_OP(my_op_gpu)
+ ...
+ .setForward(forwardCPU, "cpu")
+ .setForward(forwardGPU, "gpu")
+ .setBackward(backwardCPU, "cpu")
+ .setBackward(backwardGPU, "gpu");
+```
+
+Note that you can write only the GPU funcitons, as operators don’t have to support both CPU and GPU functions (can be GPU only).
+
+After you register forward or backward functions with context name `“gpu”`, MXNet will dispatch forward or backward GPU functions you just registered when the operator is invoked with GPU context.
Review comment:
register forward --> register the forward
with context name `“gpu”` --> with `“gpu”` context
MXNet will dispatch forward --> MXNet will call the appropriate forward
backward GPU functions --> backward functions
----------------------------------------------------------------
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
With regards,
Apache Git Services
[GitHub] [incubator-mxnet] wkcn commented on a change in pull request
#17486: Update CustomOp doc with changes for GPU support
Posted by GitBox <gi...@apache.org>.
wkcn commented on a change in pull request #17486: Update CustomOp doc with changes for GPU support
URL: https://github.com/apache/incubator-mxnet/pull/17486#discussion_r375044803
##########
File path: example/extensions/lib_custom_op/README.md
##########
@@ -146,21 +186,160 @@ Let’s take a closer look at those registry functions:
* **inferShape**: This function is similar to the `inferType` function, except it is used for populating the output data shapes. You need to figure out the shapes of each output tensors for this computation. For example, if the inputs are images with shape (224,224,3) and you write a padding operator to make 10px borders for the images, then your output shape will be (234,234,3).
-* **forward**: This function executes the main forward computation. It takes four arguments. The 1st argument is the attributes. The 2nd argument is the input `MXTensors` which stores all data and info of input ndarrays. The 3rd argument is the output `MXTensors`. The 4th argument is the `OpResource` object for memory allocation and other utilities. Additionally, you can use a `dltensor` tensor structure stored in the `MXTensor` as a more standardized data structure for computing.
+* **forward**: This function executes the main forward computation. It takes four arguments. The 1st argument is the attributes. The 2nd argument is the input `MXTensors` which stores all data and info of input ndarrays. The 3rd argument is the output `MXTensors`. The 4th argument is the `OpResource` object for memory allocation and other utilities. The details of `OpResource` are covered in the below section.
+Additionally, you can use a `dltensor` tensor structure stored in the `MXTensor` as a more standardized data structure for computing.
* **backward**: This function is doing the backward gradient computation. It will be similar to the forward function. And you need to figure out the formula of the backward gradient computation.
* **mutateInputs**: This function is for marking mutable inputs. It takes two arguments. The 1st argument is the attributes. The 2nd argument is a list of input indices that are mutable among all input tensors. It is useful when some inputs are auxiliary model parameters and might be altered during forward/backward computation. Remember, the index number of `input_indices` should not exceed the number of inputs.
-### Writing Stateful Custom Operator:
+### Writing Stateful Custom Operator
A stateful custom operator is useful when a forward/backward call needs some data or ‘state’ from previous forward/backward calls. Normally we create a class, and make instance variables store the states used for computing or caching.
Most of the building blocks for making a stateful custom operator is the same as regular custom operator, except it’ll register `createOpState` instead of a `forward` function for the computation.
* [createOpState](./gemm_lib.cc#L204) - Create stateful operator instance:
* This function takes two arguments. The 1st argument is attributes. The 2nd argument is a placeholder for `CustomStatefulOp` object. You must [define a class that inherits CustomStatefulOp](./gemm_lib.cc#L178) and override the forward function (optionally the backward function). Then you need to create an instance of your class and assign it to the placeholder. In this way, all of the forward/backward calls will use the same methods in that instance, and the instance is able to keep the state of the operator.
+```c++
+ MXReturnValue createOpState(
+ std::map<std::string, std::string> attrs,
+ CustomStatefulOp** op_inst)
+```
+
+* The operator registering function will look like this:
+```c++
+ REGISTER_OP(my_state_op)
+ ...
+ .setCreateOpState(createOpState, "cpu");
+```
+
+## Writing A Custom GPU Operator Library
+
+Most of the building blocks for registering GPU custom operators are the exactly same as CPU ones, except you need to specify the `"gpu"` context name when registering `forward`, `backward` or `createOpState` function.
+
+### Run A GPU Example
+
+For illustration purposes, we provided a `ReLU` (Rectified Linear Unit) activation operator that can run on GPU. Make sure you have installed a CUDA compatible MXNet build. Go to `lib_custom_op` directory and follow these steps:
+
+1. Run `make relu_lib`. The Makefile will invoke `NVCC` compiler to compile the CUDA kernel along with regular custom operator functions from `relu_lib.cu` to generate `librelu_lib.so` library.
+2. Run `python test_relu.py`. It’ll register the GPU `ReLU` operator in the MXNet backend, then invoke the operator by feeding an `NDArray` input with GPU context, and output the result tensor with GPU context.
+
+### Writing A Regular GPU Custom Operator
+
+Since most of the building blocks for registering GPU custom operators are the exactly same as CPU ones, the registering function for an operator supporting both GPU and CPU will look like this:
+
+```c++
+ REGISTER_OP(my_op_gpu)
+ ...
+ .setForward(forwardCPU, "cpu")
+ .setForward(forwardGPU, "gpu")
+ .setBackward(backwardCPU, "cpu")
+ .setBackward(backwardGPU, "gpu");
+```
+
+Note that you can write only the GPU funcitons, as operators don’t have to support both CPU and GPU functions (can be GPU only).
+
+After you register forward or backward functions with context name `“gpu”`, MXNet will dispatch forward or backward GPU functions you just registered when the operator is invoked with GPU context.
+
+In the registered `forwardGPU` function, specify the grid and block size and launch the CUDA kernel. In the GPU operators input and output tensors are pre-allocated on the GPU memory, just like in the CPU operator tensors are pre-allocated on the CPU memory. As a result, you don’t need to call `cudaMemcpy` to move the tensor data to GPU device.
+
+```c++
+ MXReturnValue forwardGPU(std::map<std::string, std::string> attrs,
+ std::vector<MXTensor> inputs,
+ std::vector<MXTensor> outputs,
+ OpResource res) {
+ float* in_data = inputs[0].data<float>();
+ float* out_data = outputs[0].data<float>();
+ mx_stream_t cuda_stream = res.get_cuda_stream();
+ int64_t N = inputs[0].size();
+ int block = 256;
+ int grid = (N + (block - 1)) / block;
+ my_op_gpu_forward<<<grid,block,0,cuda_stream>>>(out_data, in_data, N);
+ return MX_SUCCESS;
+ }
+```
+
+Note that the `cuda_stream` object used for launching kernels is passed from MXNet backend via `OpResource` object. See below for details of `Operator Resource`.
+
+At this point all the attribute functions for each operator (`parseAttrs`, `inferShape`, etc.) run on the CPU, including the `forwardGPU` function. The only part that will actually run on the GPU is the CUDA kernel function that is launched from `forwardGPU` function (ie. my_op_gpu_forward).
+
+```c++
+ __global__ void my_op_gpu_forward(float *out, float *in, int64_t N) {
+ int tid = blockIdx.x * blockDim.x + threadIdx.x;
Review comment:
I have a question about the balance between code readability and correctness.
In CUDA code, the number of block and thread is limited. This code works for small N, but it doesn't work when `N` is larger than the maximum number of block and thread. However, it is good to read for users.
The correct code shows as below, but it is unfriendly to read for newbies.
```c++
for (int64_t i = blockIdx.x * blockDim.x + threadIdx.x; i < N; i += blockDim.x * gridDim.x) {
// code your computational logic here
}
```
I'm not sure which one is better.
A solution is to provide the `CUDA_KERNEL_LOOP` macro.
```c++
#define CUDA_KERNEL_LOOP(i, n) \
for (int i = blockIdx.x * blockDim.x + threadIdx.x; \
i < (n); \
i += blockDim.x * gridDim.x)
```
----------------------------------------------------------------
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
With regards,
Apache Git Services
[GitHub] [incubator-mxnet] samskalicky commented on a change in pull request
#17486: Update CustomOp doc with changes for GPU support
Posted by GitBox <gi...@apache.org>.
samskalicky commented on a change in pull request #17486: Update CustomOp doc with changes for GPU support
URL: https://github.com/apache/incubator-mxnet/pull/17486#discussion_r378678068
##########
File path: example/extensions/lib_custom_op/README.md
##########
@@ -28,139 +28,340 @@ Custom operators (CustomOp) enable users to write new operators without compilin
### Have MXNet Ready
-First you should install MXNet either from compiling from source code or download from nightly build. It doesn’t matter if the build comes with CUDA or MKLDNN. The custom operator doesn’t interact with the execution of other native MXNet operators.
+Custom Operator support was merged (#15921, #17270) and is not available in versions of MXNet prior to v1.7.0.
+To access the feature now, please install MXNet by compiling from source using master or using the previously mentioned commits, downloading one of the nightly builds, or from a release of MXNet 1.7.0+.
+For running the following example, it doesn’t matter if it is a CUDA, MKLDNN or plain MXNet build; the custom operator doesn’t interact with the execution of other native MXNet operators.
+Note that if you want to run GPU examples and write your custom operators running on GPU, you still need an MXNet CUDA build.
-### Run An Example:
+### Run An Example
-You can start getting familiar with custom operators by running some examples provided in the **example/extensions/lib_custom_op** directory. Start with a common linear algebra operator like `gemm` (Generalized Matrix Multiplication). Go to `lib_custom_op` directory and follow these steps:
+You can start getting familiar with custom operators by running some examples provided in the `example/extensions/lib_custom_op` directory. Start with a common linear algebra operator like `gemm` (Generalized Matrix Multiplication). Go to `lib_custom_op` directory and follow these steps:
1. Run `make gemm_lib`. The Makefile will generate a dynamic library **libgemm_lib.so** compiled from `gemm_lib.cc`. This is the library you are going to load that contains everything for the custom gemm operator.
-2. Run `python test_gemm.py`. It’ll first load the above .so library, find the operators, register them in the MXNet backend, print "Found x operators", then invoke the operator like a regular MXNet operator and output the result.
+2. Run `python test_gemm.py`. It’ll first load the library compiled from step 1, find the operators, register them in the MXNet backend, then invoke the operator like a regular MXNet operator and output the result.
+Below is the output when running the python `test_gemm.py` command. Notice that it loads 2 operators: `my_gemm` and `state_gemm`.
-### Basic Files For Gemm Library:
+```
+[19:22:02] ../src/c_api/c_api.cc:286: Found 2 operators in library
+[19:22:02] ../src/c_api/c_api.cc:350: Op[0] my_gemm
+[19:22:02] ../src/c_api/c_api.cc:350: Op[1] state_gemm
+[19:22:02] ../src/c_api/c_api.cc:785: Found 0 partitioners in library
+--------start ndarray compute---------
+[[ 50.]
+ [122.]]
+<NDArray 2x1 @cpu(0)>
+...
+```
+
+Note that you can safely ignore the `Found 0 partitioners` info as it is not related to the custom operator.
+
+### Basic Files For A GeMM Library
* **lib_custom_op/gemm_lib.cc**: This file has a source code implementation of all required components of a custom operator, as well as the registration of the custom operator.
-* **lib_custom_op/Makefile**: Compile source code to a dynamic shared library, with a header file `include/mxnet/lib_api.h` from MXNet source code. Currently the custom operator is compatible with C++11 onwards.
+* **lib_custom_op/Makefile**: This file compiles `gemm_lib.cc` to a dynamic shared library named `libgemm_lib.so`. It includes the header file `include/mxnet/lib_api.h` from MXNet source code. Currently the custom operator APIs require C++11 onwards.
* **lib_custom_op/test_gemm.py**: This file calls `mx.library.load(‘libgemm_lib.so’)` to load the library containing the custom operator, invokes the operator using both NDArray and Symbol APIs, and prints outputs of the forward and backward passes. The outputs should be the same as the regular MXNet `gemm` operator.
-## Writing Custom Operator Library:
+* **include/mxnet/lib_api.h**: This file from MXNet source code is the single header file needed to include all necessary data types and function prototypes for writing a custom operator library.
+You can either specify the include path in the `Makefile`, or copy the header file over to `example/extensions/lib_custom_op` folder.
+Note that apart from this header, the custom operator library is independent of MXNet source.
-For building a library containing your own custom operator, compose a C++ source file like `myop_lib.cc`, include `lib_api.h` header file, and write your custom operator implementation with these essential functions:
+## Writing A Custom CPU Operator Library
+
+To build your own library containing custom CPU operator, compose a C++ source file like `myop_lib.cc`, include `lib_api.h` header file, and write your custom operator implementation with these required functions:
- `initialize` - Library Initialization Function
- `REGISTER_OP` - Operator Registration Marco
Review comment:
@rondogency I know you agree with me in appreciating @marcoabreu, but lets call this "Macro" here instead
----------------------------------------------------------------
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
With regards,
Apache Git Services
[GitHub] [incubator-mxnet] samskalicky commented on a change in pull request
#17486: Update CustomOp doc with changes for GPU support
Posted by GitBox <gi...@apache.org>.
samskalicky commented on a change in pull request #17486: Update CustomOp doc with changes for GPU support
URL: https://github.com/apache/incubator-mxnet/pull/17486#discussion_r375020328
##########
File path: example/extensions/lib_custom_op/README.md
##########
@@ -56,87 +72,111 @@ For building a library containing your own custom operator, compose a C++ source
- `forward` - Forward Computation (can be replace with `createOpState`, see below for details)
Then compile it to `libmyop_lib.so` dynamic library using the following command:
+
```bash
g++ -shared -fPIC -std=c++11 myop_lib.cc -o libmyop_lib.so -I ../../../include/mxnet
```
+If you don't want to download MXNet source and choose only using `lib_api.h` header, you can copy the header over to the same folder of `myop_lib.cc` and run:
+
+```bash
+g++ -shared -fPIC -std=c++11 myop_lib.cc -o libmyop_lib.so
+```
+
Finally, you can write a Python script to load the library and run your custom operator:
Review comment:
Finally, you can write some code to load the library and run your custom operator in any language binding. Heres an example in Python (but C Predict API and C++ APIs work too):
----------------------------------------------------------------
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
With regards,
Apache Git Services
[GitHub] [incubator-mxnet] samskalicky commented on a change in pull request
#17486: Update CustomOp doc with changes for GPU support
Posted by GitBox <gi...@apache.org>.
samskalicky commented on a change in pull request #17486: Update CustomOp doc with changes for GPU support
URL: https://github.com/apache/incubator-mxnet/pull/17486#discussion_r375022531
##########
File path: example/extensions/lib_custom_op/README.md
##########
@@ -56,87 +72,111 @@ For building a library containing your own custom operator, compose a C++ source
- `forward` - Forward Computation (can be replace with `createOpState`, see below for details)
Then compile it to `libmyop_lib.so` dynamic library using the following command:
+
```bash
g++ -shared -fPIC -std=c++11 myop_lib.cc -o libmyop_lib.so -I ../../../include/mxnet
```
+If you don't want to download MXNet source and choose only using `lib_api.h` header, you can copy the header over to the same folder of `myop_lib.cc` and run:
+
+```bash
+g++ -shared -fPIC -std=c++11 myop_lib.cc -o libmyop_lib.so
+```
+
Finally, you can write a Python script to load the library and run your custom operator:
+
```python
import mxnet as mx
mx.library.load(‘libmyop_lib.so’)
mx.nd.my_op(...)
```
-### Writing Regular Custom Operator:
+### Writing A Regular Custom Operator
-There are several essential building blocks for making a (stateless) custom operator:
+There are several essential building blocks for making a custom operator:
* [initialize](./gemm_lib.cc#L227):
* This function is the library initialization function necessary for any dynamic libraries. It checks if you are using a compatible version of MXNet. Note that this `version` parameter is passed from MXNet when library is loaded.
- MXReturnValue initialize(int version)
+```c++
+ MXReturnValue initialize(int version)
+```
* [parseAttrs](./gemm_lib.cc#L118):
* This function specifies number of input and output tensors for the custom operator; also this is where a custom operator can validate the attributes (ie. options) specified by the user.
- MXReturnValue parseAttrs(
- std::map<std::string,
- std::string> attrs,
- int* num_in,
- int* num_out)
+```c++
+ MXReturnValue parseAttrs(
+ std::map<std::string,
+ std::string> attrs,
+ int* num_in,
+ int* num_out)
+```
* [inferType](./gemm_lib.cc#L124):
* This function specifies how the custom operator infers output data types using input data types.
- MXReturnValue inferType(
- std::map<std::string, std::string> attrs,
- std::vector<int> &intypes,
- std::vector<int> &outtypes)
+```c++
+ MXReturnValue inferType(
+ std::map<std::string, std::string> attrs,
+ std::vector<int> &intypes,
+ std::vector<int> &outtypes)
+```
* [inferShape](./gemm_lib.cc#L143):
* This function specifies how the custom operator infers output tensor shape using input shape.
- MXReturnValue inferShape(
- std::map<std::string, std::string> attrs,
- std::vector<std::vector<unsigned int>> &inshapes,
- std::vector<std::vector<unsigned int>> &outshapes)
+```c++
+ MXReturnValue inferShape(
+ std::map<std::string, std::string> attrs,
+ std::vector<std::vector<unsigned int>> &inshapes,
+ std::vector<std::vector<unsigned int>> &outshapes)
+```
* [forward](./gemm_lib.cc#L56):
* This function specifies the computation of the forward pass of the operator.
- MXReturnValue forward(
- std::map<std::string, std::string> attrs,
- std::vector<MXTensor> inputs,
- std::vector<MXTensor> outputs,
- OpResource res)
+```c++
+ MXReturnValue forward(
+ std::map<std::string, std::string> attrs,
+ std::vector<MXTensor> inputs,
+ std::vector<MXTensor> outputs,
+ OpResource res)
+```
* [REGISTER_OP(my_op_name)](./gemm_lib.cc#L169):
- * This macro registers the custom operator and its properties to MXNet NDArray and Symbol APIs by its name.
-
- REGISTER_OP(my_op_name)
- .setForward(forward)
- .setParseAttrs(parseAttrs)
- .setInferType(inferType)
- .setInferShape(inferShape);
+ * This macro registers the custom operator and its properties to MXNet NDArray and Symbol APIs by its name. Note that for operator running on CPU, you need pass the name of the context `"cpu"` when registering forward or backward function.
+```c++
+ REGISTER_OP(my_op_name)
+ .setForward(forward, "cpu")
+ .setBackward(backward, "cpu")
+ .setParseAttrs(parseAttrs)
+ .setInferType(inferType)
+ .setInferShape(inferShape);
+```
Also there are some optional functions you can specify:
* [backward](./gemm_lib.cc#L90) - Backward gradient function:
* This function specifies the computation of the backward pass of the operator.
- MXReturnValue backward(
- std::map<std::string, std::string> attrs,
- std::vector<MXTensor> inputs,
- std::vector<MXTensor> outputs,
- OpResource res)
+```c++
+ MXReturnValue backward(
+ std::map<std::string, std::string> attrs,
+ std::vector<MXTensor> inputs,
+ std::vector<MXTensor> outputs,
+ OpResource res)
+```
* [mutateInputs](./gemm_lib.cc#L214) - Specify mutable input:
* This function allows you to mark some inputs to be mutable inputs. It is useful when using aux parameters for BatchNorm-like operators.
- MXReturnValue mutateInputs(
- std::map<std::string, std::string> attrs,
- std::vector<int> &input_indices)
+```c++
+ MXReturnValue mutateInputs(
+ std::map<std::string, std::string> attrs,
+ std::vector<int> &input_indices)
+```
Review comment:
put register_op here
----------------------------------------------------------------
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
With regards,
Apache Git Services
[GitHub] [incubator-mxnet] aaronmarkham commented on a change in pull
request #17486: Update CustomOp doc with changes for GPU support
Posted by GitBox <gi...@apache.org>.
aaronmarkham commented on a change in pull request #17486: Update CustomOp doc with changes for GPU support
URL: https://github.com/apache/incubator-mxnet/pull/17486#discussion_r373215293
##########
File path: example/extensions/lib_custom_op/README.md
##########
@@ -28,26 +28,39 @@ Custom operators (CustomOp) enable users to write new operators without compilin
### Have MXNet Ready
-First you should install MXNet either from compiling from source code or download from nightly build. It doesn’t matter if the build comes with CUDA or MKLDNN. The custom operator doesn’t interact with the execution of other native MXNet operators.
+Custom Operator support was recently merged (#15921, #17270) and is not available in a released version of MXNet yet. It will be part of the forthcoming 1.7 release, until then please install MXNet by compiling from source or downloading one of the nightly builds. For running the example below, it doesn’t matter if it is a CUDA, MKLDNN or Vanila build. The custom operator doesn’t interact with the execution of other native MXNet operators. Note that if you want to run GPU examples and write your custom operators running on GPU, you still need MXNet CUDA build.
Review comment:
```suggestion
Custom Operator support was merged (#15921, #17270) and is not available in versions of MXNet prior to v1.7.0. To access the feature now, please install MXNet by compiling from source using master or using the previously mentioned commits, downloading one of the nightly builds, or from a release of MXNet 1.7.0+. For running the following example, it doesn’t matter if it is a CUDA, MKLDNN or plain MXNet build; the custom operator doesn’t interact with the execution of other native MXNet operators. Note that if you want to run GPU examples and write your custom operators running on GPU, you still need an MXNet CUDA build.
```
Adjusted so you don't have to update it later.
----------------------------------------------------------------
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
With regards,
Apache Git Services
[GitHub] [incubator-mxnet] aaronmarkham commented on a change in pull
request #17486: Update CustomOp doc with changes for GPU support
Posted by GitBox <gi...@apache.org>.
aaronmarkham commented on a change in pull request #17486: Update CustomOp doc with changes for GPU support
URL: https://github.com/apache/incubator-mxnet/pull/17486#discussion_r373216823
##########
File path: example/extensions/lib_custom_op/README.md
##########
@@ -146,21 +168,154 @@ Let’s take a closer look at those registry functions:
* **inferShape**: This function is similar to the `inferType` function, except it is used for populating the output data shapes. You need to figure out the shapes of each output tensors for this computation. For example, if the inputs are images with shape (224,224,3) and you write a padding operator to make 10px borders for the images, then your output shape will be (234,234,3).
-* **forward**: This function executes the main forward computation. It takes four arguments. The 1st argument is the attributes. The 2nd argument is the input `MXTensors` which stores all data and info of input ndarrays. The 3rd argument is the output `MXTensors`. The 4th argument is the `OpResource` object for memory allocation and other utilities. Additionally, you can use a `dltensor` tensor structure stored in the `MXTensor` as a more standardized data structure for computing.
+* **forward**: This function executes the main forward computation. It takes four arguments. The 1st argument is the attributes. The 2nd argument is the input `MXTensors` which stores all data and info of input ndarrays. The 3rd argument is the output `MXTensors`. The 4th argument is the `OpResource` object for memory allocation and other utilities. The details of `OpResource` are covered in the below section.
+Additionally, you can use a `dltensor` tensor structure stored in the `MXTensor` as a more standardized data structure for computing.
* **backward**: This function is doing the backward gradient computation. It will be similar to the forward function. And you need to figure out the formula of the backward gradient computation.
* **mutateInputs**: This function is for marking mutable inputs. It takes two arguments. The 1st argument is the attributes. The 2nd argument is a list of input indices that are mutable among all input tensors. It is useful when some inputs are auxiliary model parameters and might be altered during forward/backward computation. Remember, the index number of `input_indices` should not exceed the number of inputs.
-### Writing Stateful Custom Operator:
+### Writing Stateful Custom Operator
A stateful custom operator is useful when a forward/backward call needs some data or ‘state’ from previous forward/backward calls. Normally we create a class, and make instance variables store the states used for computing or caching.
Most of the building blocks for making a stateful custom operator is the same as regular custom operator, except it’ll register `createOpState` instead of a `forward` function for the computation.
* [createOpState](./gemm_lib.cc#L204) - Create stateful operator instance:
* This function takes two arguments. The 1st argument is attributes. The 2nd argument is a placeholder for `CustomStatefulOp` object. You must [define a class that inherits CustomStatefulOp](./gemm_lib.cc#L178) and override the forward function (optionally the backward function). Then you need to create an instance of your class and assign it to the placeholder. In this way, all of the forward/backward calls will use the same methods in that instance, and the instance is able to keep the state of the operator.
+```c++
+ MXReturnValue createOpState(
+ std::map<std::string, std::string> attrs,
+ CustomStatefulOp** op_inst)
+```
+
+* The operator registering function will look like this:
+```c++
+ REGISTER_OP(my_state_op)
+ ...
+ .setCreateOpState(createOpState, "cpu");
+```
+
+## Writing Custom GPU Operator Library
+
+Most of the building blocks for registering GPU custom operators are the exactly same as CPU ones, except you need to specify the `"gpu"` context name when registering `forward`, `backward` or `createOpState` function.
+
+### Run A GPU Example
+
+For illustration purposes, we provided a `ReLU` (Rectified Linear Unit) activation operator that can run on GPU. Make sure you have installed a CUDA compatible MXNet build. Go to `lib_custom_op` directory and follow these steps:
+
+1. Run `make relu_lib`. The Makefile will invoke `NVCC` compiler to compile the CUDA kernel along with regular custom operator functions from `relu_lib.cu` to generate `librelu_lib.so` library.
+2. Run `python test_relu.py`. It’ll register the GPU `ReLU` operator in the MXNet backend, then invoke the operator by feeding an `NDArray` input with GPU context, and output the result tensor with GPU context.
+
+### Regular GPU Custom Operator
+
+Since most of the building blocks for registering GPU custom operators are the exactly same as CPU ones, the registering function for an operator supporting both GPU and CPU will look like this:
+```c++
+ REGISTER_OP(my_op_gpu)
+ ...
+ .setForward(forwardCPU, "cpu")
+ .setForward(forwardGPU, "gpu")
+ .setBackward(backwardCPU, "cpu")
+ .setBackward(backwardGPU, "gpu");
+```
+
+Note that you can write only the GPU funcitons, as operators don’t have to support both CPU and GPU functions (can be GPU only).
+
+After you register forward or backward functions with context name `“gpu”`, MXNet will dispatch forward or backward GPU functions you just registered when the operator is invoked with GPU context.
+
+In the registered `forwardGPU` function, specify the grid and block size and launch the CUDA kernel. In the GPU operators input and output tensors are pre-allocated on the GPU memory, just like in the CPU operator tensors are pre-allocated on the CPU memory. As a result, you don’t need to call `cudaMemcpy` to move the tensor data to GPU device.
+```c++
Review comment:
```suggestion
```c++
```
give it some space
----------------------------------------------------------------
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
With regards,
Apache Git Services
[GitHub] [incubator-mxnet] samskalicky commented on a change in pull request
#17486: Update CustomOp doc with changes for GPU support
Posted by GitBox <gi...@apache.org>.
samskalicky commented on a change in pull request #17486: Update CustomOp doc with changes for GPU support
URL: https://github.com/apache/incubator-mxnet/pull/17486#discussion_r375023372
##########
File path: example/extensions/lib_custom_op/README.md
##########
@@ -146,21 +186,160 @@ Let’s take a closer look at those registry functions:
* **inferShape**: This function is similar to the `inferType` function, except it is used for populating the output data shapes. You need to figure out the shapes of each output tensors for this computation. For example, if the inputs are images with shape (224,224,3) and you write a padding operator to make 10px borders for the images, then your output shape will be (234,234,3).
-* **forward**: This function executes the main forward computation. It takes four arguments. The 1st argument is the attributes. The 2nd argument is the input `MXTensors` which stores all data and info of input ndarrays. The 3rd argument is the output `MXTensors`. The 4th argument is the `OpResource` object for memory allocation and other utilities. Additionally, you can use a `dltensor` tensor structure stored in the `MXTensor` as a more standardized data structure for computing.
+* **forward**: This function executes the main forward computation. It takes four arguments. The 1st argument is the attributes. The 2nd argument is the input `MXTensors` which stores all data and info of input ndarrays. The 3rd argument is the output `MXTensors`. The 4th argument is the `OpResource` object for memory allocation and other utilities. The details of `OpResource` are covered in the below section.
+Additionally, you can use a `dltensor` tensor structure stored in the `MXTensor` as a more standardized data structure for computing.
* **backward**: This function is doing the backward gradient computation. It will be similar to the forward function. And you need to figure out the formula of the backward gradient computation.
* **mutateInputs**: This function is for marking mutable inputs. It takes two arguments. The 1st argument is the attributes. The 2nd argument is a list of input indices that are mutable among all input tensors. It is useful when some inputs are auxiliary model parameters and might be altered during forward/backward computation. Remember, the index number of `input_indices` should not exceed the number of inputs.
-### Writing Stateful Custom Operator:
+### Writing Stateful Custom Operator
Review comment:
Writing A Stateful Custom Operator
----------------------------------------------------------------
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
With regards,
Apache Git Services
[GitHub] [incubator-mxnet] samskalicky commented on issue #17486: Update
CustomOp doc with changes for GPU support
Posted by GitBox <gi...@apache.org>.
samskalicky commented on issue #17486: Update CustomOp doc with changes for GPU support
URL: https://github.com/apache/incubator-mxnet/pull/17486#issuecomment-583613040
@aaronmarkham it looks fine to me here:
https://github.com/apache/incubator-mxnet/blob/9b93000c4a8f8da05e51547312bd605ec3798b3f/example/extensions/lib_custom_op/README.md
Looks like its just weird GitHub source-diff showing odd italics.
----------------------------------------------------------------
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
With regards,
Apache Git Services
[GitHub] [incubator-mxnet] rondogency edited a comment on issue #17486:
Update CustomOp doc with changes for GPU support
Posted by GitBox <gi...@apache.org>.
rondogency edited a comment on issue #17486: Update CustomOp doc with changes for GPU support
URL: https://github.com/apache/incubator-mxnet/pull/17486#issuecomment-586023537
@a550461053 Thanks for the questions. Currently you cannot call MXNet built-in operators from a custom operator. We are making a plan for adding this support, but it cannot be easily implemented and may be a separate project. Also I added a snippet of how to use mutable input, hope it would help.
----------------------------------------------------------------
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
With regards,
Apache Git Services
[GitHub] [incubator-mxnet] samskalicky commented on a change in pull request
#17486: Update CustomOp doc with changes for GPU support
Posted by GitBox <gi...@apache.org>.
samskalicky commented on a change in pull request #17486: Update CustomOp doc with changes for GPU support
URL: https://github.com/apache/incubator-mxnet/pull/17486#discussion_r375025679
##########
File path: example/extensions/lib_custom_op/README.md
##########
@@ -146,21 +186,160 @@ Let’s take a closer look at those registry functions:
* **inferShape**: This function is similar to the `inferType` function, except it is used for populating the output data shapes. You need to figure out the shapes of each output tensors for this computation. For example, if the inputs are images with shape (224,224,3) and you write a padding operator to make 10px borders for the images, then your output shape will be (234,234,3).
-* **forward**: This function executes the main forward computation. It takes four arguments. The 1st argument is the attributes. The 2nd argument is the input `MXTensors` which stores all data and info of input ndarrays. The 3rd argument is the output `MXTensors`. The 4th argument is the `OpResource` object for memory allocation and other utilities. Additionally, you can use a `dltensor` tensor structure stored in the `MXTensor` as a more standardized data structure for computing.
+* **forward**: This function executes the main forward computation. It takes four arguments. The 1st argument is the attributes. The 2nd argument is the input `MXTensors` which stores all data and info of input ndarrays. The 3rd argument is the output `MXTensors`. The 4th argument is the `OpResource` object for memory allocation and other utilities. The details of `OpResource` are covered in the below section.
+Additionally, you can use a `dltensor` tensor structure stored in the `MXTensor` as a more standardized data structure for computing.
* **backward**: This function is doing the backward gradient computation. It will be similar to the forward function. And you need to figure out the formula of the backward gradient computation.
* **mutateInputs**: This function is for marking mutable inputs. It takes two arguments. The 1st argument is the attributes. The 2nd argument is a list of input indices that are mutable among all input tensors. It is useful when some inputs are auxiliary model parameters and might be altered during forward/backward computation. Remember, the index number of `input_indices` should not exceed the number of inputs.
-### Writing Stateful Custom Operator:
+### Writing Stateful Custom Operator
A stateful custom operator is useful when a forward/backward call needs some data or ‘state’ from previous forward/backward calls. Normally we create a class, and make instance variables store the states used for computing or caching.
Most of the building blocks for making a stateful custom operator is the same as regular custom operator, except it’ll register `createOpState` instead of a `forward` function for the computation.
* [createOpState](./gemm_lib.cc#L204) - Create stateful operator instance:
* This function takes two arguments. The 1st argument is attributes. The 2nd argument is a placeholder for `CustomStatefulOp` object. You must [define a class that inherits CustomStatefulOp](./gemm_lib.cc#L178) and override the forward function (optionally the backward function). Then you need to create an instance of your class and assign it to the placeholder. In this way, all of the forward/backward calls will use the same methods in that instance, and the instance is able to keep the state of the operator.
+```c++
+ MXReturnValue createOpState(
+ std::map<std::string, std::string> attrs,
+ CustomStatefulOp** op_inst)
+```
+
+* The operator registering function will look like this:
+```c++
+ REGISTER_OP(my_state_op)
+ ...
+ .setCreateOpState(createOpState, "cpu");
+```
+
+## Writing A Custom GPU Operator Library
+
+Most of the building blocks for registering GPU custom operators are the exactly same as CPU ones, except you need to specify the `"gpu"` context name when registering `forward`, `backward` or `createOpState` function.
+
+### Run A GPU Example
+
+For illustration purposes, we provided a `ReLU` (Rectified Linear Unit) activation operator that can run on GPU. Make sure you have installed a CUDA compatible MXNet build. Go to `lib_custom_op` directory and follow these steps:
+
+1. Run `make relu_lib`. The Makefile will invoke `NVCC` compiler to compile the CUDA kernel along with regular custom operator functions from `relu_lib.cu` to generate `librelu_lib.so` library.
+2. Run `python test_relu.py`. It’ll register the GPU `ReLU` operator in the MXNet backend, then invoke the operator by feeding an `NDArray` input with GPU context, and output the result tensor with GPU context.
+
+### Writing A Regular GPU Custom Operator
+
+Since most of the building blocks for registering GPU custom operators are the exactly same as CPU ones, the registering function for an operator supporting both GPU and CPU will look like this:
+
+```c++
+ REGISTER_OP(my_op_gpu)
+ ...
+ .setForward(forwardCPU, "cpu")
+ .setForward(forwardGPU, "gpu")
+ .setBackward(backwardCPU, "cpu")
+ .setBackward(backwardGPU, "gpu");
+```
+
+Note that you can write only the GPU funcitons, as operators don’t have to support both CPU and GPU functions (can be GPU only).
+
+After you register forward or backward functions with context name `“gpu”`, MXNet will dispatch forward or backward GPU functions you just registered when the operator is invoked with GPU context.
+
+In the registered `forwardGPU` function, specify the grid and block size and launch the CUDA kernel. In the GPU operators input and output tensors are pre-allocated on the GPU memory, just like in the CPU operator tensors are pre-allocated on the CPU memory. As a result, you don’t need to call `cudaMemcpy` to move the tensor data to GPU device.
+
+```c++
+ MXReturnValue forwardGPU(std::map<std::string, std::string> attrs,
+ std::vector<MXTensor> inputs,
+ std::vector<MXTensor> outputs,
+ OpResource res) {
+ float* in_data = inputs[0].data<float>();
+ float* out_data = outputs[0].data<float>();
+ mx_stream_t cuda_stream = res.get_cuda_stream();
+ int64_t N = inputs[0].size();
+ int block = 256;
+ int grid = (N + (block - 1)) / block;
+ my_op_gpu_forward<<<grid,block,0,cuda_stream>>>(out_data, in_data, N);
+ return MX_SUCCESS;
+ }
+```
+
Review comment:
Notice that `in_data` and `out_data` are pointers to memory allocated on the GPU. So you can pass them directly to your CUDA kernel.
----------------------------------------------------------------
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
With regards,
Apache Git Services
[GitHub] [incubator-mxnet] samskalicky commented on a change in pull request
#17486: Update CustomOp doc with changes for GPU support
Posted by GitBox <gi...@apache.org>.
samskalicky commented on a change in pull request #17486: Update CustomOp doc with changes for GPU support
URL: https://github.com/apache/incubator-mxnet/pull/17486#discussion_r375017581
##########
File path: example/extensions/lib_custom_op/README.md
##########
@@ -28,26 +28,42 @@ Custom operators (CustomOp) enable users to write new operators without compilin
### Have MXNet Ready
-First you should install MXNet either from compiling from source code or download from nightly build. It doesn’t matter if the build comes with CUDA or MKLDNN. The custom operator doesn’t interact with the execution of other native MXNet operators.
+Custom Operator support was merged (#15921, #17270) and is not available in versions of MXNet prior to v1.7.0. To access the feature now, please install MXNet by compiling from source using master or using the previously mentioned commits, downloading one of the nightly builds, or from a release of MXNet 1.7.0+. For running the following example, it doesn’t matter if it is a CUDA, MKLDNN or plain MXNet build; the custom operator doesn’t interact with the execution of other native MXNet operators. Note that if you want to run GPU examples and write your custom operators running on GPU, you still need an MXNet CUDA build.
-### Run An Example:
+### Run An Example
-You can start getting familiar with custom operators by running some examples provided in the **example/extensions/lib_custom_op** directory. Start with a common linear algebra operator like `gemm` (Generalized Matrix Multiplication). Go to `lib_custom_op` directory and follow these steps:
+You can start getting familiar with custom operators by running some examples provided in the `example/extensions/lib_custom_op` directory. Start with a common linear algebra operator like `gemm` (Generalized Matrix Multiplication). Go to `lib_custom_op` directory and follow these steps:
1. Run `make gemm_lib`. The Makefile will generate a dynamic library **libgemm_lib.so** compiled from `gemm_lib.cc`. This is the library you are going to load that contains everything for the custom gemm operator.
-2. Run `python test_gemm.py`. It’ll first load the above .so library, find the operators, register them in the MXNet backend, print "Found x operators", then invoke the operator like a regular MXNet operator and output the result.
+2. Run `python test_gemm.py`. It’ll first load the library compiled from step 1, find the operators, register them in the MXNet backend, then invoke the operator like a regular MXNet operator and output the result.
-### Basic Files For Gemm Library:
+```
+[19:22:02] ../src/c_api/c_api.cc:286: Found 2 operators in library
+[19:22:02] ../src/c_api/c_api.cc:350: Op[0] my_gemm
+[19:22:02] ../src/c_api/c_api.cc:350: Op[1] state_gemm
+[19:22:02] ../src/c_api/c_api.cc:785: Found 0 partitioners in library
+--------start ndarray compute---------
+[[ 50.]
+ [122.]]
+<NDArray 2x1 @cpu(0)>
+...
+```
+
+Note that you can safely ignore the `Found 0 partitioners` info as it is not related to the custom operator.
+
+### Basic Files For A GeMM Library
* **lib_custom_op/gemm_lib.cc**: This file has a source code implementation of all required components of a custom operator, as well as the registration of the custom operator.
-* **lib_custom_op/Makefile**: Compile source code to a dynamic shared library, with a header file `include/mxnet/lib_api.h` from MXNet source code. Currently the custom operator is compatible with C++11 onwards.
+* **lib_custom_op/Makefile**: Compile `gemm_lib.cc` to a dynamic shared library named `libgemm_lib.so`, with a include path of the header file `include/mxnet/lib_api.h` from MXNet source code. Currently the custom operator is compatible with C++11 onwards.
Review comment:
This file compiles `gemm_lib.cc` to a dynamic shared library named `libgemm_lib.so`. It includes the header file `include/mxnet/lib_api.h` from MXNet source code. Currently the custom operator APIs require C++11 onwards.
----------------------------------------------------------------
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
With regards,
Apache Git Services
[GitHub] [incubator-mxnet] wkcn commented on a change in pull request
#17486: Update CustomOp doc with changes for GPU support
Posted by GitBox <gi...@apache.org>.
wkcn commented on a change in pull request #17486: Update CustomOp doc with changes for GPU support
URL: https://github.com/apache/incubator-mxnet/pull/17486#discussion_r375044803
##########
File path: example/extensions/lib_custom_op/README.md
##########
@@ -146,21 +186,160 @@ Let’s take a closer look at those registry functions:
* **inferShape**: This function is similar to the `inferType` function, except it is used for populating the output data shapes. You need to figure out the shapes of each output tensors for this computation. For example, if the inputs are images with shape (224,224,3) and you write a padding operator to make 10px borders for the images, then your output shape will be (234,234,3).
-* **forward**: This function executes the main forward computation. It takes four arguments. The 1st argument is the attributes. The 2nd argument is the input `MXTensors` which stores all data and info of input ndarrays. The 3rd argument is the output `MXTensors`. The 4th argument is the `OpResource` object for memory allocation and other utilities. Additionally, you can use a `dltensor` tensor structure stored in the `MXTensor` as a more standardized data structure for computing.
+* **forward**: This function executes the main forward computation. It takes four arguments. The 1st argument is the attributes. The 2nd argument is the input `MXTensors` which stores all data and info of input ndarrays. The 3rd argument is the output `MXTensors`. The 4th argument is the `OpResource` object for memory allocation and other utilities. The details of `OpResource` are covered in the below section.
+Additionally, you can use a `dltensor` tensor structure stored in the `MXTensor` as a more standardized data structure for computing.
* **backward**: This function is doing the backward gradient computation. It will be similar to the forward function. And you need to figure out the formula of the backward gradient computation.
* **mutateInputs**: This function is for marking mutable inputs. It takes two arguments. The 1st argument is the attributes. The 2nd argument is a list of input indices that are mutable among all input tensors. It is useful when some inputs are auxiliary model parameters and might be altered during forward/backward computation. Remember, the index number of `input_indices` should not exceed the number of inputs.
-### Writing Stateful Custom Operator:
+### Writing Stateful Custom Operator
A stateful custom operator is useful when a forward/backward call needs some data or ‘state’ from previous forward/backward calls. Normally we create a class, and make instance variables store the states used for computing or caching.
Most of the building blocks for making a stateful custom operator is the same as regular custom operator, except it’ll register `createOpState` instead of a `forward` function for the computation.
* [createOpState](./gemm_lib.cc#L204) - Create stateful operator instance:
* This function takes two arguments. The 1st argument is attributes. The 2nd argument is a placeholder for `CustomStatefulOp` object. You must [define a class that inherits CustomStatefulOp](./gemm_lib.cc#L178) and override the forward function (optionally the backward function). Then you need to create an instance of your class and assign it to the placeholder. In this way, all of the forward/backward calls will use the same methods in that instance, and the instance is able to keep the state of the operator.
+```c++
+ MXReturnValue createOpState(
+ std::map<std::string, std::string> attrs,
+ CustomStatefulOp** op_inst)
+```
+
+* The operator registering function will look like this:
+```c++
+ REGISTER_OP(my_state_op)
+ ...
+ .setCreateOpState(createOpState, "cpu");
+```
+
+## Writing A Custom GPU Operator Library
+
+Most of the building blocks for registering GPU custom operators are the exactly same as CPU ones, except you need to specify the `"gpu"` context name when registering `forward`, `backward` or `createOpState` function.
+
+### Run A GPU Example
+
+For illustration purposes, we provided a `ReLU` (Rectified Linear Unit) activation operator that can run on GPU. Make sure you have installed a CUDA compatible MXNet build. Go to `lib_custom_op` directory and follow these steps:
+
+1. Run `make relu_lib`. The Makefile will invoke `NVCC` compiler to compile the CUDA kernel along with regular custom operator functions from `relu_lib.cu` to generate `librelu_lib.so` library.
+2. Run `python test_relu.py`. It’ll register the GPU `ReLU` operator in the MXNet backend, then invoke the operator by feeding an `NDArray` input with GPU context, and output the result tensor with GPU context.
+
+### Writing A Regular GPU Custom Operator
+
+Since most of the building blocks for registering GPU custom operators are the exactly same as CPU ones, the registering function for an operator supporting both GPU and CPU will look like this:
+
+```c++
+ REGISTER_OP(my_op_gpu)
+ ...
+ .setForward(forwardCPU, "cpu")
+ .setForward(forwardGPU, "gpu")
+ .setBackward(backwardCPU, "cpu")
+ .setBackward(backwardGPU, "gpu");
+```
+
+Note that you can write only the GPU funcitons, as operators don’t have to support both CPU and GPU functions (can be GPU only).
+
+After you register forward or backward functions with context name `“gpu”`, MXNet will dispatch forward or backward GPU functions you just registered when the operator is invoked with GPU context.
+
+In the registered `forwardGPU` function, specify the grid and block size and launch the CUDA kernel. In the GPU operators input and output tensors are pre-allocated on the GPU memory, just like in the CPU operator tensors are pre-allocated on the CPU memory. As a result, you don’t need to call `cudaMemcpy` to move the tensor data to GPU device.
+
+```c++
+ MXReturnValue forwardGPU(std::map<std::string, std::string> attrs,
+ std::vector<MXTensor> inputs,
+ std::vector<MXTensor> outputs,
+ OpResource res) {
+ float* in_data = inputs[0].data<float>();
+ float* out_data = outputs[0].data<float>();
+ mx_stream_t cuda_stream = res.get_cuda_stream();
+ int64_t N = inputs[0].size();
+ int block = 256;
+ int grid = (N + (block - 1)) / block;
+ my_op_gpu_forward<<<grid,block,0,cuda_stream>>>(out_data, in_data, N);
+ return MX_SUCCESS;
+ }
+```
+
+Note that the `cuda_stream` object used for launching kernels is passed from MXNet backend via `OpResource` object. See below for details of `Operator Resource`.
+
+At this point all the attribute functions for each operator (`parseAttrs`, `inferShape`, etc.) run on the CPU, including the `forwardGPU` function. The only part that will actually run on the GPU is the CUDA kernel function that is launched from `forwardGPU` function (ie. my_op_gpu_forward).
+
+```c++
+ __global__ void my_op_gpu_forward(float *out, float *in, int64_t N) {
+ int tid = blockIdx.x * blockDim.x + threadIdx.x;
Review comment:
I have a question about the balance between code readability and correctness.
In CUDA code, the number of block and thread is limited. This code works for small N, but it doesn't work when `N` is larger than the maximum number of block and thread. However, it is good to read for users.
The correct code shows as below, but it is unfriendly to read for newbies.
```c++
for (int64_t i = blockIdx.x * blockDim.x + threadIdx.x; i < N; i += blockDim.x * gridDim.x) {
// code your computational logic here
}
```
I'm not sure which one is better.
A solution is to provide the `CUDA_KERNEL_LOOP` macro.
```c++
#define CUDA_KERNEL_LOOP(i, n) \
for (int64_t i = blockIdx.x * blockDim.x + threadIdx.x; \
i < (n); \
i += blockDim.x * gridDim.x)
```
----------------------------------------------------------------
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
With regards,
Apache Git Services
[GitHub] [incubator-mxnet] samskalicky commented on a change in pull request
#17486: Update CustomOp doc with changes for GPU support
Posted by GitBox <gi...@apache.org>.
samskalicky commented on a change in pull request #17486: Update CustomOp doc with changes for GPU support
URL: https://github.com/apache/incubator-mxnet/pull/17486#discussion_r375023718
##########
File path: example/extensions/lib_custom_op/README.md
##########
@@ -146,21 +186,160 @@ Let’s take a closer look at those registry functions:
* **inferShape**: This function is similar to the `inferType` function, except it is used for populating the output data shapes. You need to figure out the shapes of each output tensors for this computation. For example, if the inputs are images with shape (224,224,3) and you write a padding operator to make 10px borders for the images, then your output shape will be (234,234,3).
-* **forward**: This function executes the main forward computation. It takes four arguments. The 1st argument is the attributes. The 2nd argument is the input `MXTensors` which stores all data and info of input ndarrays. The 3rd argument is the output `MXTensors`. The 4th argument is the `OpResource` object for memory allocation and other utilities. Additionally, you can use a `dltensor` tensor structure stored in the `MXTensor` as a more standardized data structure for computing.
+* **forward**: This function executes the main forward computation. It takes four arguments. The 1st argument is the attributes. The 2nd argument is the input `MXTensors` which stores all data and info of input ndarrays. The 3rd argument is the output `MXTensors`. The 4th argument is the `OpResource` object for memory allocation and other utilities. The details of `OpResource` are covered in the below section.
+Additionally, you can use a `dltensor` tensor structure stored in the `MXTensor` as a more standardized data structure for computing.
* **backward**: This function is doing the backward gradient computation. It will be similar to the forward function. And you need to figure out the formula of the backward gradient computation.
* **mutateInputs**: This function is for marking mutable inputs. It takes two arguments. The 1st argument is the attributes. The 2nd argument is a list of input indices that are mutable among all input tensors. It is useful when some inputs are auxiliary model parameters and might be altered during forward/backward computation. Remember, the index number of `input_indices` should not exceed the number of inputs.
-### Writing Stateful Custom Operator:
+### Writing Stateful Custom Operator
A stateful custom operator is useful when a forward/backward call needs some data or ‘state’ from previous forward/backward calls. Normally we create a class, and make instance variables store the states used for computing or caching.
Most of the building blocks for making a stateful custom operator is the same as regular custom operator, except it’ll register `createOpState` instead of a `forward` function for the computation.
* [createOpState](./gemm_lib.cc#L204) - Create stateful operator instance:
* This function takes two arguments. The 1st argument is attributes. The 2nd argument is a placeholder for `CustomStatefulOp` object. You must [define a class that inherits CustomStatefulOp](./gemm_lib.cc#L178) and override the forward function (optionally the backward function). Then you need to create an instance of your class and assign it to the placeholder. In this way, all of the forward/backward calls will use the same methods in that instance, and the instance is able to keep the state of the operator.
+```c++
+ MXReturnValue createOpState(
+ std::map<std::string, std::string> attrs,
+ CustomStatefulOp** op_inst)
+```
+
+* The operator registering function will look like this:
+```c++
+ REGISTER_OP(my_state_op)
+ ...
+ .setCreateOpState(createOpState, "cpu");
+```
+
Review comment:
Notice that you will need to register each `createOpState` function specific for each context you operator will support.
----------------------------------------------------------------
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
With regards,
Apache Git Services
[GitHub] [incubator-mxnet] samskalicky commented on a change in pull request
#17486: Update CustomOp doc with changes for GPU support
Posted by GitBox <gi...@apache.org>.
samskalicky commented on a change in pull request #17486: Update CustomOp doc with changes for GPU support
URL: https://github.com/apache/incubator-mxnet/pull/17486#discussion_r375027292
##########
File path: example/extensions/lib_custom_op/README.md
##########
@@ -146,21 +186,160 @@ Let’s take a closer look at those registry functions:
* **inferShape**: This function is similar to the `inferType` function, except it is used for populating the output data shapes. You need to figure out the shapes of each output tensors for this computation. For example, if the inputs are images with shape (224,224,3) and you write a padding operator to make 10px borders for the images, then your output shape will be (234,234,3).
-* **forward**: This function executes the main forward computation. It takes four arguments. The 1st argument is the attributes. The 2nd argument is the input `MXTensors` which stores all data and info of input ndarrays. The 3rd argument is the output `MXTensors`. The 4th argument is the `OpResource` object for memory allocation and other utilities. Additionally, you can use a `dltensor` tensor structure stored in the `MXTensor` as a more standardized data structure for computing.
+* **forward**: This function executes the main forward computation. It takes four arguments. The 1st argument is the attributes. The 2nd argument is the input `MXTensors` which stores all data and info of input ndarrays. The 3rd argument is the output `MXTensors`. The 4th argument is the `OpResource` object for memory allocation and other utilities. The details of `OpResource` are covered in the below section.
+Additionally, you can use a `dltensor` tensor structure stored in the `MXTensor` as a more standardized data structure for computing.
* **backward**: This function is doing the backward gradient computation. It will be similar to the forward function. And you need to figure out the formula of the backward gradient computation.
* **mutateInputs**: This function is for marking mutable inputs. It takes two arguments. The 1st argument is the attributes. The 2nd argument is a list of input indices that are mutable among all input tensors. It is useful when some inputs are auxiliary model parameters and might be altered during forward/backward computation. Remember, the index number of `input_indices` should not exceed the number of inputs.
-### Writing Stateful Custom Operator:
+### Writing Stateful Custom Operator
A stateful custom operator is useful when a forward/backward call needs some data or ‘state’ from previous forward/backward calls. Normally we create a class, and make instance variables store the states used for computing or caching.
Most of the building blocks for making a stateful custom operator is the same as regular custom operator, except it’ll register `createOpState` instead of a `forward` function for the computation.
* [createOpState](./gemm_lib.cc#L204) - Create stateful operator instance:
* This function takes two arguments. The 1st argument is attributes. The 2nd argument is a placeholder for `CustomStatefulOp` object. You must [define a class that inherits CustomStatefulOp](./gemm_lib.cc#L178) and override the forward function (optionally the backward function). Then you need to create an instance of your class and assign it to the placeholder. In this way, all of the forward/backward calls will use the same methods in that instance, and the instance is able to keep the state of the operator.
+```c++
+ MXReturnValue createOpState(
+ std::map<std::string, std::string> attrs,
+ CustomStatefulOp** op_inst)
+```
+
+* The operator registering function will look like this:
+```c++
+ REGISTER_OP(my_state_op)
+ ...
+ .setCreateOpState(createOpState, "cpu");
+```
+
+## Writing A Custom GPU Operator Library
+
+Most of the building blocks for registering GPU custom operators are the exactly same as CPU ones, except you need to specify the `"gpu"` context name when registering `forward`, `backward` or `createOpState` function.
+
+### Run A GPU Example
+
+For illustration purposes, we provided a `ReLU` (Rectified Linear Unit) activation operator that can run on GPU. Make sure you have installed a CUDA compatible MXNet build. Go to `lib_custom_op` directory and follow these steps:
+
+1. Run `make relu_lib`. The Makefile will invoke `NVCC` compiler to compile the CUDA kernel along with regular custom operator functions from `relu_lib.cu` to generate `librelu_lib.so` library.
+2. Run `python test_relu.py`. It’ll register the GPU `ReLU` operator in the MXNet backend, then invoke the operator by feeding an `NDArray` input with GPU context, and output the result tensor with GPU context.
+
+### Writing A Regular GPU Custom Operator
+
+Since most of the building blocks for registering GPU custom operators are the exactly same as CPU ones, the registering function for an operator supporting both GPU and CPU will look like this:
+
+```c++
+ REGISTER_OP(my_op_gpu)
+ ...
+ .setForward(forwardCPU, "cpu")
+ .setForward(forwardGPU, "gpu")
+ .setBackward(backwardCPU, "cpu")
+ .setBackward(backwardGPU, "gpu");
+```
+
+Note that you can write only the GPU funcitons, as operators don’t have to support both CPU and GPU functions (can be GPU only).
+
+After you register forward or backward functions with context name `“gpu”`, MXNet will dispatch forward or backward GPU functions you just registered when the operator is invoked with GPU context.
+
+In the registered `forwardGPU` function, specify the grid and block size and launch the CUDA kernel. In the GPU operators input and output tensors are pre-allocated on the GPU memory, just like in the CPU operator tensors are pre-allocated on the CPU memory. As a result, you don’t need to call `cudaMemcpy` to move the tensor data to GPU device.
+
+```c++
+ MXReturnValue forwardGPU(std::map<std::string, std::string> attrs,
+ std::vector<MXTensor> inputs,
+ std::vector<MXTensor> outputs,
+ OpResource res) {
+ float* in_data = inputs[0].data<float>();
+ float* out_data = outputs[0].data<float>();
+ mx_stream_t cuda_stream = res.get_cuda_stream();
+ int64_t N = inputs[0].size();
+ int block = 256;
+ int grid = (N + (block - 1)) / block;
+ my_op_gpu_forward<<<grid,block,0,cuda_stream>>>(out_data, in_data, N);
+ return MX_SUCCESS;
+ }
+```
+
+Note that the `cuda_stream` object used for launching kernels is passed from MXNet backend via `OpResource` object. See below for details of `Operator Resource`.
+
+At this point all the attribute functions for each operator (`parseAttrs`, `inferShape`, etc.) run on the CPU, including the `forwardGPU` function. The only part that will actually run on the GPU is the CUDA kernel function that is launched from `forwardGPU` function (ie. my_op_gpu_forward).
Review comment:
remove: from `forwardGPU` function (ie. my_op_gpu_forward).
----------------------------------------------------------------
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
With regards,
Apache Git Services
[GitHub] [incubator-mxnet] aaronmarkham commented on a change in pull
request #17486: Update CustomOp doc with changes for GPU support
Posted by GitBox <gi...@apache.org>.
aaronmarkham commented on a change in pull request #17486: Update CustomOp doc with changes for GPU support
URL: https://github.com/apache/incubator-mxnet/pull/17486#discussion_r373216936
##########
File path: example/extensions/lib_custom_op/README.md
##########
@@ -146,21 +168,154 @@ Let’s take a closer look at those registry functions:
* **inferShape**: This function is similar to the `inferType` function, except it is used for populating the output data shapes. You need to figure out the shapes of each output tensors for this computation. For example, if the inputs are images with shape (224,224,3) and you write a padding operator to make 10px borders for the images, then your output shape will be (234,234,3).
-* **forward**: This function executes the main forward computation. It takes four arguments. The 1st argument is the attributes. The 2nd argument is the input `MXTensors` which stores all data and info of input ndarrays. The 3rd argument is the output `MXTensors`. The 4th argument is the `OpResource` object for memory allocation and other utilities. Additionally, you can use a `dltensor` tensor structure stored in the `MXTensor` as a more standardized data structure for computing.
+* **forward**: This function executes the main forward computation. It takes four arguments. The 1st argument is the attributes. The 2nd argument is the input `MXTensors` which stores all data and info of input ndarrays. The 3rd argument is the output `MXTensors`. The 4th argument is the `OpResource` object for memory allocation and other utilities. The details of `OpResource` are covered in the below section.
+Additionally, you can use a `dltensor` tensor structure stored in the `MXTensor` as a more standardized data structure for computing.
* **backward**: This function is doing the backward gradient computation. It will be similar to the forward function. And you need to figure out the formula of the backward gradient computation.
* **mutateInputs**: This function is for marking mutable inputs. It takes two arguments. The 1st argument is the attributes. The 2nd argument is a list of input indices that are mutable among all input tensors. It is useful when some inputs are auxiliary model parameters and might be altered during forward/backward computation. Remember, the index number of `input_indices` should not exceed the number of inputs.
-### Writing Stateful Custom Operator:
+### Writing Stateful Custom Operator
A stateful custom operator is useful when a forward/backward call needs some data or ‘state’ from previous forward/backward calls. Normally we create a class, and make instance variables store the states used for computing or caching.
Most of the building blocks for making a stateful custom operator is the same as regular custom operator, except it’ll register `createOpState` instead of a `forward` function for the computation.
* [createOpState](./gemm_lib.cc#L204) - Create stateful operator instance:
* This function takes two arguments. The 1st argument is attributes. The 2nd argument is a placeholder for `CustomStatefulOp` object. You must [define a class that inherits CustomStatefulOp](./gemm_lib.cc#L178) and override the forward function (optionally the backward function). Then you need to create an instance of your class and assign it to the placeholder. In this way, all of the forward/backward calls will use the same methods in that instance, and the instance is able to keep the state of the operator.
+```c++
+ MXReturnValue createOpState(
+ std::map<std::string, std::string> attrs,
+ CustomStatefulOp** op_inst)
+```
+
+* The operator registering function will look like this:
+```c++
+ REGISTER_OP(my_state_op)
+ ...
+ .setCreateOpState(createOpState, "cpu");
+```
+
+## Writing Custom GPU Operator Library
+
+Most of the building blocks for registering GPU custom operators are the exactly same as CPU ones, except you need to specify the `"gpu"` context name when registering `forward`, `backward` or `createOpState` function.
+
+### Run A GPU Example
+
+For illustration purposes, we provided a `ReLU` (Rectified Linear Unit) activation operator that can run on GPU. Make sure you have installed a CUDA compatible MXNet build. Go to `lib_custom_op` directory and follow these steps:
+
+1. Run `make relu_lib`. The Makefile will invoke `NVCC` compiler to compile the CUDA kernel along with regular custom operator functions from `relu_lib.cu` to generate `librelu_lib.so` library.
+2. Run `python test_relu.py`. It’ll register the GPU `ReLU` operator in the MXNet backend, then invoke the operator by feeding an `NDArray` input with GPU context, and output the result tensor with GPU context.
+
+### Regular GPU Custom Operator
+
+Since most of the building blocks for registering GPU custom operators are the exactly same as CPU ones, the registering function for an operator supporting both GPU and CPU will look like this:
+```c++
Review comment:
```suggestion
```c++
```
give it some space
----------------------------------------------------------------
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
With regards,
Apache Git Services
[GitHub] [incubator-mxnet] rondogency commented on a change in pull request
#17486: Update CustomOp doc with changes for GPU support
Posted by GitBox <gi...@apache.org>.
rondogency commented on a change in pull request #17486: Update CustomOp doc with changes for GPU support
URL: https://github.com/apache/incubator-mxnet/pull/17486#discussion_r375582060
##########
File path: example/extensions/lib_custom_op/README.md
##########
@@ -146,21 +186,160 @@ Let’s take a closer look at those registry functions:
* **inferShape**: This function is similar to the `inferType` function, except it is used for populating the output data shapes. You need to figure out the shapes of each output tensors for this computation. For example, if the inputs are images with shape (224,224,3) and you write a padding operator to make 10px borders for the images, then your output shape will be (234,234,3).
-* **forward**: This function executes the main forward computation. It takes four arguments. The 1st argument is the attributes. The 2nd argument is the input `MXTensors` which stores all data and info of input ndarrays. The 3rd argument is the output `MXTensors`. The 4th argument is the `OpResource` object for memory allocation and other utilities. Additionally, you can use a `dltensor` tensor structure stored in the `MXTensor` as a more standardized data structure for computing.
+* **forward**: This function executes the main forward computation. It takes four arguments. The 1st argument is the attributes. The 2nd argument is the input `MXTensors` which stores all data and info of input ndarrays. The 3rd argument is the output `MXTensors`. The 4th argument is the `OpResource` object for memory allocation and other utilities. The details of `OpResource` are covered in the below section.
+Additionally, you can use a `dltensor` tensor structure stored in the `MXTensor` as a more standardized data structure for computing.
* **backward**: This function is doing the backward gradient computation. It will be similar to the forward function. And you need to figure out the formula of the backward gradient computation.
* **mutateInputs**: This function is for marking mutable inputs. It takes two arguments. The 1st argument is the attributes. The 2nd argument is a list of input indices that are mutable among all input tensors. It is useful when some inputs are auxiliary model parameters and might be altered during forward/backward computation. Remember, the index number of `input_indices` should not exceed the number of inputs.
-### Writing Stateful Custom Operator:
+### Writing Stateful Custom Operator
A stateful custom operator is useful when a forward/backward call needs some data or ‘state’ from previous forward/backward calls. Normally we create a class, and make instance variables store the states used for computing or caching.
Most of the building blocks for making a stateful custom operator is the same as regular custom operator, except it’ll register `createOpState` instead of a `forward` function for the computation.
* [createOpState](./gemm_lib.cc#L204) - Create stateful operator instance:
* This function takes two arguments. The 1st argument is attributes. The 2nd argument is a placeholder for `CustomStatefulOp` object. You must [define a class that inherits CustomStatefulOp](./gemm_lib.cc#L178) and override the forward function (optionally the backward function). Then you need to create an instance of your class and assign it to the placeholder. In this way, all of the forward/backward calls will use the same methods in that instance, and the instance is able to keep the state of the operator.
+```c++
+ MXReturnValue createOpState(
+ std::map<std::string, std::string> attrs,
+ CustomStatefulOp** op_inst)
+```
+
+* The operator registering function will look like this:
+```c++
+ REGISTER_OP(my_state_op)
+ ...
+ .setCreateOpState(createOpState, "cpu");
+```
+
+## Writing A Custom GPU Operator Library
+
+Most of the building blocks for registering GPU custom operators are the exactly same as CPU ones, except you need to specify the `"gpu"` context name when registering `forward`, `backward` or `createOpState` function.
+
+### Run A GPU Example
+
+For illustration purposes, we provided a `ReLU` (Rectified Linear Unit) activation operator that can run on GPU. Make sure you have installed a CUDA compatible MXNet build. Go to `lib_custom_op` directory and follow these steps:
+
+1. Run `make relu_lib`. The Makefile will invoke `NVCC` compiler to compile the CUDA kernel along with regular custom operator functions from `relu_lib.cu` to generate `librelu_lib.so` library.
+2. Run `python test_relu.py`. It’ll register the GPU `ReLU` operator in the MXNet backend, then invoke the operator by feeding an `NDArray` input with GPU context, and output the result tensor with GPU context.
+
+### Writing A Regular GPU Custom Operator
+
+Since most of the building blocks for registering GPU custom operators are the exactly same as CPU ones, the registering function for an operator supporting both GPU and CPU will look like this:
+
+```c++
+ REGISTER_OP(my_op_gpu)
+ ...
+ .setForward(forwardCPU, "cpu")
+ .setForward(forwardGPU, "gpu")
+ .setBackward(backwardCPU, "cpu")
+ .setBackward(backwardGPU, "gpu");
+```
+
+Note that you can write only the GPU funcitons, as operators don’t have to support both CPU and GPU functions (can be GPU only).
+
+After you register forward or backward functions with context name `“gpu”`, MXNet will dispatch forward or backward GPU functions you just registered when the operator is invoked with GPU context.
+
+In the registered `forwardGPU` function, specify the grid and block size and launch the CUDA kernel. In the GPU operators input and output tensors are pre-allocated on the GPU memory, just like in the CPU operator tensors are pre-allocated on the CPU memory. As a result, you don’t need to call `cudaMemcpy` to move the tensor data to GPU device.
+
+```c++
+ MXReturnValue forwardGPU(std::map<std::string, std::string> attrs,
+ std::vector<MXTensor> inputs,
+ std::vector<MXTensor> outputs,
+ OpResource res) {
+ float* in_data = inputs[0].data<float>();
+ float* out_data = outputs[0].data<float>();
+ mx_stream_t cuda_stream = res.get_cuda_stream();
+ int64_t N = inputs[0].size();
+ int block = 256;
+ int grid = (N + (block - 1)) / block;
+ my_op_gpu_forward<<<grid,block,0,cuda_stream>>>(out_data, in_data, N);
+ return MX_SUCCESS;
+ }
+```
+
+Note that the `cuda_stream` object used for launching kernels is passed from MXNet backend via `OpResource` object. See below for details of `Operator Resource`.
+
+At this point all the attribute functions for each operator (`parseAttrs`, `inferShape`, etc.) run on the CPU, including the `forwardGPU` function. The only part that will actually run on the GPU is the CUDA kernel function that is launched from `forwardGPU` function (ie. my_op_gpu_forward).
+
+```c++
+ __global__ void my_op_gpu_forward(float *out, float *in, int64_t N) {
+ int tid = blockIdx.x * blockDim.x + threadIdx.x;
Review comment:
agree, let's be generic enough for this
----------------------------------------------------------------
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
With regards,
Apache Git Services
[GitHub] [incubator-mxnet] samskalicky commented on a change in pull request
#17486: Update CustomOp doc with changes for GPU support
Posted by GitBox <gi...@apache.org>.
samskalicky commented on a change in pull request #17486: Update CustomOp doc with changes for GPU support
URL: https://github.com/apache/incubator-mxnet/pull/17486#discussion_r375022066
##########
File path: example/extensions/lib_custom_op/README.md
##########
@@ -56,87 +72,111 @@ For building a library containing your own custom operator, compose a C++ source
- `forward` - Forward Computation (can be replace with `createOpState`, see below for details)
Then compile it to `libmyop_lib.so` dynamic library using the following command:
+
```bash
g++ -shared -fPIC -std=c++11 myop_lib.cc -o libmyop_lib.so -I ../../../include/mxnet
```
+If you don't want to download MXNet source and choose only using `lib_api.h` header, you can copy the header over to the same folder of `myop_lib.cc` and run:
+
+```bash
+g++ -shared -fPIC -std=c++11 myop_lib.cc -o libmyop_lib.so
+```
+
Finally, you can write a Python script to load the library and run your custom operator:
+
```python
import mxnet as mx
mx.library.load(‘libmyop_lib.so’)
mx.nd.my_op(...)
```
-### Writing Regular Custom Operator:
+### Writing A Regular Custom Operator
-There are several essential building blocks for making a (stateless) custom operator:
+There are several essential building blocks for making a custom operator:
* [initialize](./gemm_lib.cc#L227):
* This function is the library initialization function necessary for any dynamic libraries. It checks if you are using a compatible version of MXNet. Note that this `version` parameter is passed from MXNet when library is loaded.
- MXReturnValue initialize(int version)
+```c++
+ MXReturnValue initialize(int version)
+```
* [parseAttrs](./gemm_lib.cc#L118):
* This function specifies number of input and output tensors for the custom operator; also this is where a custom operator can validate the attributes (ie. options) specified by the user.
- MXReturnValue parseAttrs(
- std::map<std::string,
- std::string> attrs,
- int* num_in,
- int* num_out)
+```c++
+ MXReturnValue parseAttrs(
+ std::map<std::string,
+ std::string> attrs,
+ int* num_in,
+ int* num_out)
+```
* [inferType](./gemm_lib.cc#L124):
* This function specifies how the custom operator infers output data types using input data types.
- MXReturnValue inferType(
- std::map<std::string, std::string> attrs,
- std::vector<int> &intypes,
- std::vector<int> &outtypes)
+```c++
+ MXReturnValue inferType(
+ std::map<std::string, std::string> attrs,
+ std::vector<int> &intypes,
+ std::vector<int> &outtypes)
+```
* [inferShape](./gemm_lib.cc#L143):
* This function specifies how the custom operator infers output tensor shape using input shape.
- MXReturnValue inferShape(
- std::map<std::string, std::string> attrs,
- std::vector<std::vector<unsigned int>> &inshapes,
- std::vector<std::vector<unsigned int>> &outshapes)
+```c++
+ MXReturnValue inferShape(
+ std::map<std::string, std::string> attrs,
+ std::vector<std::vector<unsigned int>> &inshapes,
+ std::vector<std::vector<unsigned int>> &outshapes)
+```
* [forward](./gemm_lib.cc#L56):
* This function specifies the computation of the forward pass of the operator.
- MXReturnValue forward(
- std::map<std::string, std::string> attrs,
- std::vector<MXTensor> inputs,
- std::vector<MXTensor> outputs,
- OpResource res)
+```c++
+ MXReturnValue forward(
+ std::map<std::string, std::string> attrs,
+ std::vector<MXTensor> inputs,
+ std::vector<MXTensor> outputs,
+ OpResource res)
+```
* [REGISTER_OP(my_op_name)](./gemm_lib.cc#L169):
- * This macro registers the custom operator and its properties to MXNet NDArray and Symbol APIs by its name.
-
- REGISTER_OP(my_op_name)
- .setForward(forward)
- .setParseAttrs(parseAttrs)
- .setInferType(inferType)
- .setInferShape(inferShape);
+ * This macro registers the custom operator and its properties to MXNet NDArray and Symbol APIs by its name. Note that for operator running on CPU, you need pass the name of the context `"cpu"` when registering forward or backward function.
Review comment:
This macro registers the custom operator with MXNet. You can register additional functions with the operator too. Note that you register functions for each context for Forward and Backward, here we show an example for CPU.
----------------------------------------------------------------
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
With regards,
Apache Git Services
[GitHub] [incubator-mxnet] samskalicky commented on a change in pull request
#17486: Update CustomOp doc with changes for GPU support
Posted by GitBox <gi...@apache.org>.
samskalicky commented on a change in pull request #17486: Update CustomOp doc with changes for GPU support
URL: https://github.com/apache/incubator-mxnet/pull/17486#discussion_r375020986
##########
File path: example/extensions/lib_custom_op/README.md
##########
@@ -56,87 +72,111 @@ For building a library containing your own custom operator, compose a C++ source
- `forward` - Forward Computation (can be replace with `createOpState`, see below for details)
Then compile it to `libmyop_lib.so` dynamic library using the following command:
+
```bash
g++ -shared -fPIC -std=c++11 myop_lib.cc -o libmyop_lib.so -I ../../../include/mxnet
```
+If you don't want to download MXNet source and choose only using `lib_api.h` header, you can copy the header over to the same folder of `myop_lib.cc` and run:
+
+```bash
+g++ -shared -fPIC -std=c++11 myop_lib.cc -o libmyop_lib.so
+```
+
Finally, you can write a Python script to load the library and run your custom operator:
+
```python
import mxnet as mx
mx.library.load(‘libmyop_lib.so’)
mx.nd.my_op(...)
```
-### Writing Regular Custom Operator:
+### Writing A Regular Custom Operator
-There are several essential building blocks for making a (stateless) custom operator:
+There are several essential building blocks for making a custom operator:
* [initialize](./gemm_lib.cc#L227):
* This function is the library initialization function necessary for any dynamic libraries. It checks if you are using a compatible version of MXNet. Note that this `version` parameter is passed from MXNet when library is loaded.
Review comment:
The initialize function is called when MXNet first loads the library. MXNet passes its version to this function when called. This gives the library the ability to check which version of MXNet is being used. It also provides a place where library state can be initialized.
----------------------------------------------------------------
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
With regards,
Apache Git Services
[GitHub] [incubator-mxnet] aaronmarkham commented on a change in pull
request #17486: Update CustomOp doc with changes for GPU support
Posted by GitBox <gi...@apache.org>.
aaronmarkham commented on a change in pull request #17486: Update CustomOp doc with changes for GPU support
URL: https://github.com/apache/incubator-mxnet/pull/17486#discussion_r373217146
##########
File path: example/extensions/lib_custom_op/README.md
##########
@@ -146,21 +168,154 @@ Let’s take a closer look at those registry functions:
* **inferShape**: This function is similar to the `inferType` function, except it is used for populating the output data shapes. You need to figure out the shapes of each output tensors for this computation. For example, if the inputs are images with shape (224,224,3) and you write a padding operator to make 10px borders for the images, then your output shape will be (234,234,3).
-* **forward**: This function executes the main forward computation. It takes four arguments. The 1st argument is the attributes. The 2nd argument is the input `MXTensors` which stores all data and info of input ndarrays. The 3rd argument is the output `MXTensors`. The 4th argument is the `OpResource` object for memory allocation and other utilities. Additionally, you can use a `dltensor` tensor structure stored in the `MXTensor` as a more standardized data structure for computing.
+* **forward**: This function executes the main forward computation. It takes four arguments. The 1st argument is the attributes. The 2nd argument is the input `MXTensors` which stores all data and info of input ndarrays. The 3rd argument is the output `MXTensors`. The 4th argument is the `OpResource` object for memory allocation and other utilities. The details of `OpResource` are covered in the below section.
+Additionally, you can use a `dltensor` tensor structure stored in the `MXTensor` as a more standardized data structure for computing.
* **backward**: This function is doing the backward gradient computation. It will be similar to the forward function. And you need to figure out the formula of the backward gradient computation.
* **mutateInputs**: This function is for marking mutable inputs. It takes two arguments. The 1st argument is the attributes. The 2nd argument is a list of input indices that are mutable among all input tensors. It is useful when some inputs are auxiliary model parameters and might be altered during forward/backward computation. Remember, the index number of `input_indices` should not exceed the number of inputs.
-### Writing Stateful Custom Operator:
+### Writing Stateful Custom Operator
A stateful custom operator is useful when a forward/backward call needs some data or ‘state’ from previous forward/backward calls. Normally we create a class, and make instance variables store the states used for computing or caching.
Most of the building blocks for making a stateful custom operator is the same as regular custom operator, except it’ll register `createOpState` instead of a `forward` function for the computation.
* [createOpState](./gemm_lib.cc#L204) - Create stateful operator instance:
* This function takes two arguments. The 1st argument is attributes. The 2nd argument is a placeholder for `CustomStatefulOp` object. You must [define a class that inherits CustomStatefulOp](./gemm_lib.cc#L178) and override the forward function (optionally the backward function). Then you need to create an instance of your class and assign it to the placeholder. In this way, all of the forward/backward calls will use the same methods in that instance, and the instance is able to keep the state of the operator.
+```c++
+ MXReturnValue createOpState(
+ std::map<std::string, std::string> attrs,
+ CustomStatefulOp** op_inst)
+```
+
+* The operator registering function will look like this:
+```c++
+ REGISTER_OP(my_state_op)
+ ...
+ .setCreateOpState(createOpState, "cpu");
+```
+
+## Writing Custom GPU Operator Library
+
+Most of the building blocks for registering GPU custom operators are the exactly same as CPU ones, except you need to specify the `"gpu"` context name when registering `forward`, `backward` or `createOpState` function.
+
+### Run A GPU Example
+
+For illustration purposes, we provided a `ReLU` (Rectified Linear Unit) activation operator that can run on GPU. Make sure you have installed a CUDA compatible MXNet build. Go to `lib_custom_op` directory and follow these steps:
+
+1. Run `make relu_lib`. The Makefile will invoke `NVCC` compiler to compile the CUDA kernel along with regular custom operator functions from `relu_lib.cu` to generate `librelu_lib.so` library.
+2. Run `python test_relu.py`. It’ll register the GPU `ReLU` operator in the MXNet backend, then invoke the operator by feeding an `NDArray` input with GPU context, and output the result tensor with GPU context.
+
+### Regular GPU Custom Operator
+
+Since most of the building blocks for registering GPU custom operators are the exactly same as CPU ones, the registering function for an operator supporting both GPU and CPU will look like this:
+```c++
+ REGISTER_OP(my_op_gpu)
+ ...
+ .setForward(forwardCPU, "cpu")
+ .setForward(forwardGPU, "gpu")
+ .setBackward(backwardCPU, "cpu")
+ .setBackward(backwardGPU, "gpu");
+```
+
+Note that you can write only the GPU funcitons, as operators don’t have to support both CPU and GPU functions (can be GPU only).
+
+After you register forward or backward functions with context name `“gpu”`, MXNet will dispatch forward or backward GPU functions you just registered when the operator is invoked with GPU context.
+
+In the registered `forwardGPU` function, specify the grid and block size and launch the CUDA kernel. In the GPU operators input and output tensors are pre-allocated on the GPU memory, just like in the CPU operator tensors are pre-allocated on the CPU memory. As a result, you don’t need to call `cudaMemcpy` to move the tensor data to GPU device.
+```c++
+ MXReturnValue forwardGPU(std::map<std::string, std::string> attrs,
+ std::vector<MXTensor> inputs,
+ std::vector<MXTensor> outputs,
+ OpResource res) {
+ float* in_data = inputs[0].data<float>();
+ float* out_data = outputs[0].data<float>();
+ mx_stream_t cuda_stream = res.get_cuda_stream();
+ int64_t N = inputs[0].size();
+ int block = 256;
+ int grid = (N + (block - 1)) / block;
+ my_op_gpu_forward<<<grid,block,0,cuda_stream>>>(out_data, in_data, N);
+ return MX_SUCCESS;
+ }
+```
+
+Note that the `cuda_stream` object used for launching kernels is passed from MXNet backend via `OpResource` object. See below for details of `Operator Resource`.
+
+At this point all the attribute functions for each operator (`parseAttrs`, `inferShape`, etc.) run on the CPU, including the `forwardGPU` function. The only part that will actually run on the GPU is the CUDA kernel function that is launched from `forwardGPU` function (ie. my_op_gpu_forward).
+```c++
+ __global__ void my_op_gpu_forward(float *out, float *in, int64_t N) {
+ int tid = blockIdx.x * blockDim.x + threadIdx.x;
+ if (tid < N) {
+ // code your computational logic here
+ }
+ }
+```
+
+### Stateful GPU Custom Operator
+
+Recall that for stateful custom operators, you need to define a class that inherits `CustomStatefulOp` and overrides the `forward` and `backward` functions. Stateful operators are created context-aware, so you can create different classes for GPU and CPU stateful operators separately if desired. To do so, you register a createOpState function for each context separately like this
+```c++
Review comment:
```suggestion
```c++
```
----------------------------------------------------------------
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
With regards,
Apache Git Services
[GitHub] [incubator-mxnet] a550461053 commented on issue #17486: Update
CustomOp doc with changes for GPU support
Posted by GitBox <gi...@apache.org>.
a550461053 commented on issue #17486: Update CustomOp doc with changes for GPU support
URL: https://github.com/apache/incubator-mxnet/pull/17486#issuecomment-584780995
Hi, @rondogency I am trying to use the dynamic operator and have some questions.
How can we use `mxnet::NDArray` or `mshadow_op` used in xxx-inl.h of mxnet-backend operator? And I want to know how to use the `mutateInputs` function like in BatchNorm-like operators.
Thank you very much.
----------------------------------------------------------------
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
With regards,
Apache Git Services
[GitHub] [incubator-mxnet] samskalicky commented on a change in pull request
#17486: Update CustomOp doc with changes for GPU support
Posted by GitBox <gi...@apache.org>.
samskalicky commented on a change in pull request #17486: Update CustomOp doc with changes for GPU support
URL: https://github.com/apache/incubator-mxnet/pull/17486#discussion_r375024525
##########
File path: example/extensions/lib_custom_op/README.md
##########
@@ -146,21 +186,160 @@ Let’s take a closer look at those registry functions:
* **inferShape**: This function is similar to the `inferType` function, except it is used for populating the output data shapes. You need to figure out the shapes of each output tensors for this computation. For example, if the inputs are images with shape (224,224,3) and you write a padding operator to make 10px borders for the images, then your output shape will be (234,234,3).
-* **forward**: This function executes the main forward computation. It takes four arguments. The 1st argument is the attributes. The 2nd argument is the input `MXTensors` which stores all data and info of input ndarrays. The 3rd argument is the output `MXTensors`. The 4th argument is the `OpResource` object for memory allocation and other utilities. Additionally, you can use a `dltensor` tensor structure stored in the `MXTensor` as a more standardized data structure for computing.
+* **forward**: This function executes the main forward computation. It takes four arguments. The 1st argument is the attributes. The 2nd argument is the input `MXTensors` which stores all data and info of input ndarrays. The 3rd argument is the output `MXTensors`. The 4th argument is the `OpResource` object for memory allocation and other utilities. The details of `OpResource` are covered in the below section.
+Additionally, you can use a `dltensor` tensor structure stored in the `MXTensor` as a more standardized data structure for computing.
* **backward**: This function is doing the backward gradient computation. It will be similar to the forward function. And you need to figure out the formula of the backward gradient computation.
* **mutateInputs**: This function is for marking mutable inputs. It takes two arguments. The 1st argument is the attributes. The 2nd argument is a list of input indices that are mutable among all input tensors. It is useful when some inputs are auxiliary model parameters and might be altered during forward/backward computation. Remember, the index number of `input_indices` should not exceed the number of inputs.
-### Writing Stateful Custom Operator:
+### Writing Stateful Custom Operator
A stateful custom operator is useful when a forward/backward call needs some data or ‘state’ from previous forward/backward calls. Normally we create a class, and make instance variables store the states used for computing or caching.
Most of the building blocks for making a stateful custom operator is the same as regular custom operator, except it’ll register `createOpState` instead of a `forward` function for the computation.
* [createOpState](./gemm_lib.cc#L204) - Create stateful operator instance:
* This function takes two arguments. The 1st argument is attributes. The 2nd argument is a placeholder for `CustomStatefulOp` object. You must [define a class that inherits CustomStatefulOp](./gemm_lib.cc#L178) and override the forward function (optionally the backward function). Then you need to create an instance of your class and assign it to the placeholder. In this way, all of the forward/backward calls will use the same methods in that instance, and the instance is able to keep the state of the operator.
+```c++
+ MXReturnValue createOpState(
+ std::map<std::string, std::string> attrs,
+ CustomStatefulOp** op_inst)
+```
+
+* The operator registering function will look like this:
+```c++
+ REGISTER_OP(my_state_op)
+ ...
+ .setCreateOpState(createOpState, "cpu");
+```
+
+## Writing A Custom GPU Operator Library
+
+Most of the building blocks for registering GPU custom operators are the exactly same as CPU ones, except you need to specify the `"gpu"` context name when registering `forward`, `backward` or `createOpState` function.
+
+### Run A GPU Example
+
+For illustration purposes, we provided a `ReLU` (Rectified Linear Unit) activation operator that can run on GPU. Make sure you have installed a CUDA compatible MXNet build. Go to `lib_custom_op` directory and follow these steps:
+
+1. Run `make relu_lib`. The Makefile will invoke `NVCC` compiler to compile the CUDA kernel along with regular custom operator functions from `relu_lib.cu` to generate `librelu_lib.so` library.
+2. Run `python test_relu.py`. It’ll register the GPU `ReLU` operator in the MXNet backend, then invoke the operator by feeding an `NDArray` input with GPU context, and output the result tensor with GPU context.
+
+### Writing A Regular GPU Custom Operator
+
+Since most of the building blocks for registering GPU custom operators are the exactly same as CPU ones, the registering function for an operator supporting both GPU and CPU will look like this:
+
+```c++
+ REGISTER_OP(my_op_gpu)
+ ...
+ .setForward(forwardCPU, "cpu")
+ .setForward(forwardGPU, "gpu")
+ .setBackward(backwardCPU, "cpu")
+ .setBackward(backwardGPU, "gpu");
+```
+
+Note that you can write only the GPU funcitons, as operators don’t have to support both CPU and GPU functions (can be GPU only).
Review comment:
Note that operators don’t have to support both CPU and GPU functions (can be GPU only).
----------------------------------------------------------------
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
With regards,
Apache Git Services
[GitHub] [incubator-mxnet] rondogency commented on issue #17486: Update
CustomOp doc with changes for GPU support
Posted by GitBox <gi...@apache.org>.
rondogency commented on issue #17486: Update CustomOp doc with changes for GPU support
URL: https://github.com/apache/incubator-mxnet/pull/17486#issuecomment-582185417
@samskalicky the fix is in #17516 and you can review it now
----------------------------------------------------------------
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
With regards,
Apache Git Services
[GitHub] [incubator-mxnet] samskalicky commented on a change in pull request
#17486: Update CustomOp doc with changes for GPU support
Posted by GitBox <gi...@apache.org>.
samskalicky commented on a change in pull request #17486: Update CustomOp doc with changes for GPU support
URL: https://github.com/apache/incubator-mxnet/pull/17486#discussion_r375023276
##########
File path: example/extensions/lib_custom_op/README.md
##########
@@ -146,21 +186,160 @@ Let’s take a closer look at those registry functions:
* **inferShape**: This function is similar to the `inferType` function, except it is used for populating the output data shapes. You need to figure out the shapes of each output tensors for this computation. For example, if the inputs are images with shape (224,224,3) and you write a padding operator to make 10px borders for the images, then your output shape will be (234,234,3).
-* **forward**: This function executes the main forward computation. It takes four arguments. The 1st argument is the attributes. The 2nd argument is the input `MXTensors` which stores all data and info of input ndarrays. The 3rd argument is the output `MXTensors`. The 4th argument is the `OpResource` object for memory allocation and other utilities. Additionally, you can use a `dltensor` tensor structure stored in the `MXTensor` as a more standardized data structure for computing.
+* **forward**: This function executes the main forward computation. It takes four arguments. The 1st argument is the attributes. The 2nd argument is the input `MXTensors` which stores all data and info of input ndarrays. The 3rd argument is the output `MXTensors`. The 4th argument is the `OpResource` object for memory allocation and other utilities. The details of `OpResource` are covered in the below section.
Review comment:
"in the below section" --> "in the section below"
----------------------------------------------------------------
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
With regards,
Apache Git Services
[GitHub] [incubator-mxnet] samskalicky commented on a change in pull request
#17486: Update CustomOp doc with changes for GPU support
Posted by GitBox <gi...@apache.org>.
samskalicky commented on a change in pull request #17486: Update CustomOp doc with changes for GPU support
URL: https://github.com/apache/incubator-mxnet/pull/17486#discussion_r375020986
##########
File path: example/extensions/lib_custom_op/README.md
##########
@@ -56,87 +72,111 @@ For building a library containing your own custom operator, compose a C++ source
- `forward` - Forward Computation (can be replace with `createOpState`, see below for details)
Then compile it to `libmyop_lib.so` dynamic library using the following command:
+
```bash
g++ -shared -fPIC -std=c++11 myop_lib.cc -o libmyop_lib.so -I ../../../include/mxnet
```
+If you don't want to download MXNet source and choose only using `lib_api.h` header, you can copy the header over to the same folder of `myop_lib.cc` and run:
+
+```bash
+g++ -shared -fPIC -std=c++11 myop_lib.cc -o libmyop_lib.so
+```
+
Finally, you can write a Python script to load the library and run your custom operator:
+
```python
import mxnet as mx
mx.library.load(‘libmyop_lib.so’)
mx.nd.my_op(...)
```
-### Writing Regular Custom Operator:
+### Writing A Regular Custom Operator
-There are several essential building blocks for making a (stateless) custom operator:
+There are several essential building blocks for making a custom operator:
* [initialize](./gemm_lib.cc#L227):
* This function is the library initialization function necessary for any dynamic libraries. It checks if you are using a compatible version of MXNet. Note that this `version` parameter is passed from MXNet when library is loaded.
Review comment:
This function is called when MXNet first loads the library. MXNet passes its version to this function when called. This gives the library the ability to check which version of MXNet is being used. It also provides a place where library state can be initialized.
----------------------------------------------------------------
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
With regards,
Apache Git Services
[GitHub] [incubator-mxnet] wkcn commented on issue #17486: Update CustomOp
doc with changes for GPU support
Posted by GitBox <gi...@apache.org>.
wkcn commented on issue #17486: Update CustomOp doc with changes for GPU support
URL: https://github.com/apache/incubator-mxnet/pull/17486#issuecomment-586661767
Thanks for your contribution!
----------------------------------------------------------------
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
With regards,
Apache Git Services
[GitHub] [incubator-mxnet] samskalicky commented on a change in pull request
#17486: Update CustomOp doc with changes for GPU support
Posted by GitBox <gi...@apache.org>.
samskalicky commented on a change in pull request #17486: Update CustomOp doc with changes for GPU support
URL: https://github.com/apache/incubator-mxnet/pull/17486#discussion_r375024388
##########
File path: example/extensions/lib_custom_op/README.md
##########
@@ -146,21 +186,160 @@ Let’s take a closer look at those registry functions:
* **inferShape**: This function is similar to the `inferType` function, except it is used for populating the output data shapes. You need to figure out the shapes of each output tensors for this computation. For example, if the inputs are images with shape (224,224,3) and you write a padding operator to make 10px borders for the images, then your output shape will be (234,234,3).
-* **forward**: This function executes the main forward computation. It takes four arguments. The 1st argument is the attributes. The 2nd argument is the input `MXTensors` which stores all data and info of input ndarrays. The 3rd argument is the output `MXTensors`. The 4th argument is the `OpResource` object for memory allocation and other utilities. Additionally, you can use a `dltensor` tensor structure stored in the `MXTensor` as a more standardized data structure for computing.
+* **forward**: This function executes the main forward computation. It takes four arguments. The 1st argument is the attributes. The 2nd argument is the input `MXTensors` which stores all data and info of input ndarrays. The 3rd argument is the output `MXTensors`. The 4th argument is the `OpResource` object for memory allocation and other utilities. The details of `OpResource` are covered in the below section.
+Additionally, you can use a `dltensor` tensor structure stored in the `MXTensor` as a more standardized data structure for computing.
* **backward**: This function is doing the backward gradient computation. It will be similar to the forward function. And you need to figure out the formula of the backward gradient computation.
* **mutateInputs**: This function is for marking mutable inputs. It takes two arguments. The 1st argument is the attributes. The 2nd argument is a list of input indices that are mutable among all input tensors. It is useful when some inputs are auxiliary model parameters and might be altered during forward/backward computation. Remember, the index number of `input_indices` should not exceed the number of inputs.
-### Writing Stateful Custom Operator:
+### Writing Stateful Custom Operator
A stateful custom operator is useful when a forward/backward call needs some data or ‘state’ from previous forward/backward calls. Normally we create a class, and make instance variables store the states used for computing or caching.
Most of the building blocks for making a stateful custom operator is the same as regular custom operator, except it’ll register `createOpState` instead of a `forward` function for the computation.
* [createOpState](./gemm_lib.cc#L204) - Create stateful operator instance:
* This function takes two arguments. The 1st argument is attributes. The 2nd argument is a placeholder for `CustomStatefulOp` object. You must [define a class that inherits CustomStatefulOp](./gemm_lib.cc#L178) and override the forward function (optionally the backward function). Then you need to create an instance of your class and assign it to the placeholder. In this way, all of the forward/backward calls will use the same methods in that instance, and the instance is able to keep the state of the operator.
+```c++
+ MXReturnValue createOpState(
+ std::map<std::string, std::string> attrs,
+ CustomStatefulOp** op_inst)
+```
+
+* The operator registering function will look like this:
+```c++
+ REGISTER_OP(my_state_op)
+ ...
+ .setCreateOpState(createOpState, "cpu");
+```
+
+## Writing A Custom GPU Operator Library
+
+Most of the building blocks for registering GPU custom operators are the exactly same as CPU ones, except you need to specify the `"gpu"` context name when registering `forward`, `backward` or `createOpState` function.
+
+### Run A GPU Example
+
+For illustration purposes, we provided a `ReLU` (Rectified Linear Unit) activation operator that can run on GPU. Make sure you have installed a CUDA compatible MXNet build. Go to `lib_custom_op` directory and follow these steps:
+
+1. Run `make relu_lib`. The Makefile will invoke `NVCC` compiler to compile the CUDA kernel along with regular custom operator functions from `relu_lib.cu` to generate `librelu_lib.so` library.
+2. Run `python test_relu.py`. It’ll register the GPU `ReLU` operator in the MXNet backend, then invoke the operator by feeding an `NDArray` input with GPU context, and output the result tensor with GPU context.
+
+### Writing A Regular GPU Custom Operator
+
+Since most of the building blocks for registering GPU custom operators are the exactly same as CPU ones, the registering function for an operator supporting both GPU and CPU will look like this:
+
+```c++
+ REGISTER_OP(my_op_gpu)
Review comment:
it doesnt make sense to call the op "my_op_gpu" and then register it for both CPU and GPU. lets just call it "my_op"
----------------------------------------------------------------
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
With regards,
Apache Git Services
[GitHub] [incubator-mxnet] samskalicky commented on issue #17486: Update
CustomOp doc with changes for GPU support
Posted by GitBox <gi...@apache.org>.
samskalicky commented on issue #17486: Update CustomOp doc with changes for GPU support
URL: https://github.com/apache/incubator-mxnet/pull/17486#issuecomment-585571802
> Ok, thank you. I want to create a custom operator calling another operator which input NDArray. Both operator is async pushed to engine, I think this way is possible. Also, I am not clear on using mutateInputs API, how and when to use it? If you can provide a example of the mutateInputs API, I will be grateful to you~
Hi @a550461053 currently the custom operator design focused on separating the custom operator from the MXNet backend source code complexity. This means that your custom operator can (must) be entirely separated from MXNet. So you cannot call a regular built-in MXNet operator from your custom operator. We have an item here #17006 for adding support in the future to be able to do this, but it is not implemented yet.
As for how to use mutateInputs, it works exactly as the doc describes:
> This function allows you to mark some inputs to be mutable inputs. It is useful when using aux parameters for BatchNorm-like operators.
So lets say you have an operator with 5 inputs, you can mark the indices of the inputs that you want to be mutable like this (for example mark the last two inputs as mutable):
```
MXReturnValue batchNorm_mutateInputs(std::map<std::string, std::string> attrs,
std::vector<int> &input_indices) {
// mark mutable inputs
input_indices.push_back(3);
input_indices.push_back(4);
return MX_SUCCESS;
}
```
----------------------------------------------------------------
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
With regards,
Apache Git Services
[GitHub] [incubator-mxnet] samskalicky commented on a change in pull request
#17486: Update CustomOp doc with changes for GPU support
Posted by GitBox <gi...@apache.org>.
samskalicky commented on a change in pull request #17486: Update CustomOp doc with changes for GPU support
URL: https://github.com/apache/incubator-mxnet/pull/17486#discussion_r375027065
##########
File path: example/extensions/lib_custom_op/README.md
##########
@@ -146,21 +186,160 @@ Let’s take a closer look at those registry functions:
* **inferShape**: This function is similar to the `inferType` function, except it is used for populating the output data shapes. You need to figure out the shapes of each output tensors for this computation. For example, if the inputs are images with shape (224,224,3) and you write a padding operator to make 10px borders for the images, then your output shape will be (234,234,3).
-* **forward**: This function executes the main forward computation. It takes four arguments. The 1st argument is the attributes. The 2nd argument is the input `MXTensors` which stores all data and info of input ndarrays. The 3rd argument is the output `MXTensors`. The 4th argument is the `OpResource` object for memory allocation and other utilities. Additionally, you can use a `dltensor` tensor structure stored in the `MXTensor` as a more standardized data structure for computing.
+* **forward**: This function executes the main forward computation. It takes four arguments. The 1st argument is the attributes. The 2nd argument is the input `MXTensors` which stores all data and info of input ndarrays. The 3rd argument is the output `MXTensors`. The 4th argument is the `OpResource` object for memory allocation and other utilities. The details of `OpResource` are covered in the below section.
+Additionally, you can use a `dltensor` tensor structure stored in the `MXTensor` as a more standardized data structure for computing.
* **backward**: This function is doing the backward gradient computation. It will be similar to the forward function. And you need to figure out the formula of the backward gradient computation.
* **mutateInputs**: This function is for marking mutable inputs. It takes two arguments. The 1st argument is the attributes. The 2nd argument is a list of input indices that are mutable among all input tensors. It is useful when some inputs are auxiliary model parameters and might be altered during forward/backward computation. Remember, the index number of `input_indices` should not exceed the number of inputs.
-### Writing Stateful Custom Operator:
+### Writing Stateful Custom Operator
A stateful custom operator is useful when a forward/backward call needs some data or ‘state’ from previous forward/backward calls. Normally we create a class, and make instance variables store the states used for computing or caching.
Most of the building blocks for making a stateful custom operator is the same as regular custom operator, except it’ll register `createOpState` instead of a `forward` function for the computation.
* [createOpState](./gemm_lib.cc#L204) - Create stateful operator instance:
* This function takes two arguments. The 1st argument is attributes. The 2nd argument is a placeholder for `CustomStatefulOp` object. You must [define a class that inherits CustomStatefulOp](./gemm_lib.cc#L178) and override the forward function (optionally the backward function). Then you need to create an instance of your class and assign it to the placeholder. In this way, all of the forward/backward calls will use the same methods in that instance, and the instance is able to keep the state of the operator.
+```c++
+ MXReturnValue createOpState(
+ std::map<std::string, std::string> attrs,
+ CustomStatefulOp** op_inst)
+```
+
+* The operator registering function will look like this:
+```c++
+ REGISTER_OP(my_state_op)
+ ...
+ .setCreateOpState(createOpState, "cpu");
+```
+
+## Writing A Custom GPU Operator Library
+
+Most of the building blocks for registering GPU custom operators are the exactly same as CPU ones, except you need to specify the `"gpu"` context name when registering `forward`, `backward` or `createOpState` function.
+
+### Run A GPU Example
+
+For illustration purposes, we provided a `ReLU` (Rectified Linear Unit) activation operator that can run on GPU. Make sure you have installed a CUDA compatible MXNet build. Go to `lib_custom_op` directory and follow these steps:
+
+1. Run `make relu_lib`. The Makefile will invoke `NVCC` compiler to compile the CUDA kernel along with regular custom operator functions from `relu_lib.cu` to generate `librelu_lib.so` library.
+2. Run `python test_relu.py`. It’ll register the GPU `ReLU` operator in the MXNet backend, then invoke the operator by feeding an `NDArray` input with GPU context, and output the result tensor with GPU context.
+
+### Writing A Regular GPU Custom Operator
+
+Since most of the building blocks for registering GPU custom operators are the exactly same as CPU ones, the registering function for an operator supporting both GPU and CPU will look like this:
+
+```c++
+ REGISTER_OP(my_op_gpu)
+ ...
+ .setForward(forwardCPU, "cpu")
+ .setForward(forwardGPU, "gpu")
+ .setBackward(backwardCPU, "cpu")
+ .setBackward(backwardGPU, "gpu");
+```
+
+Note that you can write only the GPU funcitons, as operators don’t have to support both CPU and GPU functions (can be GPU only).
+
+After you register forward or backward functions with context name `“gpu”`, MXNet will dispatch forward or backward GPU functions you just registered when the operator is invoked with GPU context.
+
+In the registered `forwardGPU` function, specify the grid and block size and launch the CUDA kernel. In the GPU operators input and output tensors are pre-allocated on the GPU memory, just like in the CPU operator tensors are pre-allocated on the CPU memory. As a result, you don’t need to call `cudaMemcpy` to move the tensor data to GPU device.
+
+```c++
+ MXReturnValue forwardGPU(std::map<std::string, std::string> attrs,
+ std::vector<MXTensor> inputs,
+ std::vector<MXTensor> outputs,
+ OpResource res) {
+ float* in_data = inputs[0].data<float>();
+ float* out_data = outputs[0].data<float>();
+ mx_stream_t cuda_stream = res.get_cuda_stream();
+ int64_t N = inputs[0].size();
+ int block = 256;
+ int grid = (N + (block - 1)) / block;
+ my_op_gpu_forward<<<grid,block,0,cuda_stream>>>(out_data, in_data, N);
Review comment:
shouldnt we check for CUDA success? and then return MX_FAIL or MX_SUCCESS appropriately?
```
cudaError_t code = cudaGetLastError();
if (code != cudaSuccess) {
const char * errorMessage = cudaGetErrorString(code);
fprintf(stderr, "CUDA error returned from \"%s\" at %s:%d, Error code: %d (%s)\n", func, file, line, code, errorMessage);
return MX_FAIL;
} else {
return MX_SUCCESS;
}
----------------------------------------------------------------
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
With regards,
Apache Git Services
[GitHub] [incubator-mxnet] samskalicky edited a comment on issue #17486:
Update CustomOp doc with changes for GPU support
Posted by GitBox <gi...@apache.org>.
samskalicky edited a comment on issue #17486: Update CustomOp doc with changes for GPU support
URL: https://github.com/apache/incubator-mxnet/pull/17486#issuecomment-586640766
> But now I can only build mxnet to create c++ customOp and I have a problem on compiling into python wheel at #17577
@a550461053 just to clarify, you do not need to recompile MXNet from source to use a custom operator, you can dynamically load your library into MXNet at runtime using one of the nightly builds from here: https://repo.mxnet.io/dist/index.html
----------------------------------------------------------------
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
With regards,
Apache Git Services
[GitHub] [incubator-mxnet] wkcn commented on a change in pull request
#17486: Update CustomOp doc with changes for GPU support
Posted by GitBox <gi...@apache.org>.
wkcn commented on a change in pull request #17486: Update CustomOp doc with changes for GPU support
URL: https://github.com/apache/incubator-mxnet/pull/17486#discussion_r375598920
##########
File path: example/extensions/lib_custom_op/README.md
##########
@@ -146,21 +186,160 @@ Let’s take a closer look at those registry functions:
* **inferShape**: This function is similar to the `inferType` function, except it is used for populating the output data shapes. You need to figure out the shapes of each output tensors for this computation. For example, if the inputs are images with shape (224,224,3) and you write a padding operator to make 10px borders for the images, then your output shape will be (234,234,3).
-* **forward**: This function executes the main forward computation. It takes four arguments. The 1st argument is the attributes. The 2nd argument is the input `MXTensors` which stores all data and info of input ndarrays. The 3rd argument is the output `MXTensors`. The 4th argument is the `OpResource` object for memory allocation and other utilities. Additionally, you can use a `dltensor` tensor structure stored in the `MXTensor` as a more standardized data structure for computing.
+* **forward**: This function executes the main forward computation. It takes four arguments. The 1st argument is the attributes. The 2nd argument is the input `MXTensors` which stores all data and info of input ndarrays. The 3rd argument is the output `MXTensors`. The 4th argument is the `OpResource` object for memory allocation and other utilities. The details of `OpResource` are covered in the below section.
+Additionally, you can use a `dltensor` tensor structure stored in the `MXTensor` as a more standardized data structure for computing.
* **backward**: This function is doing the backward gradient computation. It will be similar to the forward function. And you need to figure out the formula of the backward gradient computation.
* **mutateInputs**: This function is for marking mutable inputs. It takes two arguments. The 1st argument is the attributes. The 2nd argument is a list of input indices that are mutable among all input tensors. It is useful when some inputs are auxiliary model parameters and might be altered during forward/backward computation. Remember, the index number of `input_indices` should not exceed the number of inputs.
-### Writing Stateful Custom Operator:
+### Writing Stateful Custom Operator
A stateful custom operator is useful when a forward/backward call needs some data or ‘state’ from previous forward/backward calls. Normally we create a class, and make instance variables store the states used for computing or caching.
Most of the building blocks for making a stateful custom operator is the same as regular custom operator, except it’ll register `createOpState` instead of a `forward` function for the computation.
* [createOpState](./gemm_lib.cc#L204) - Create stateful operator instance:
* This function takes two arguments. The 1st argument is attributes. The 2nd argument is a placeholder for `CustomStatefulOp` object. You must [define a class that inherits CustomStatefulOp](./gemm_lib.cc#L178) and override the forward function (optionally the backward function). Then you need to create an instance of your class and assign it to the placeholder. In this way, all of the forward/backward calls will use the same methods in that instance, and the instance is able to keep the state of the operator.
+```c++
+ MXReturnValue createOpState(
+ std::map<std::string, std::string> attrs,
+ CustomStatefulOp** op_inst)
+```
+
+* The operator registering function will look like this:
+```c++
+ REGISTER_OP(my_state_op)
+ ...
+ .setCreateOpState(createOpState, "cpu");
+```
+
+## Writing A Custom GPU Operator Library
+
+Most of the building blocks for registering GPU custom operators are the exactly same as CPU ones, except you need to specify the `"gpu"` context name when registering `forward`, `backward` or `createOpState` function.
+
+### Run A GPU Example
+
+For illustration purposes, we provided a `ReLU` (Rectified Linear Unit) activation operator that can run on GPU. Make sure you have installed a CUDA compatible MXNet build. Go to `lib_custom_op` directory and follow these steps:
+
+1. Run `make relu_lib`. The Makefile will invoke `NVCC` compiler to compile the CUDA kernel along with regular custom operator functions from `relu_lib.cu` to generate `librelu_lib.so` library.
+2. Run `python test_relu.py`. It’ll register the GPU `ReLU` operator in the MXNet backend, then invoke the operator by feeding an `NDArray` input with GPU context, and output the result tensor with GPU context.
+
+### Writing A Regular GPU Custom Operator
+
+Since most of the building blocks for registering GPU custom operators are the exactly same as CPU ones, the registering function for an operator supporting both GPU and CPU will look like this:
+
+```c++
+ REGISTER_OP(my_op_gpu)
+ ...
+ .setForward(forwardCPU, "cpu")
+ .setForward(forwardGPU, "gpu")
+ .setBackward(backwardCPU, "cpu")
+ .setBackward(backwardGPU, "gpu");
+```
+
+Note that you can write only the GPU funcitons, as operators don’t have to support both CPU and GPU functions (can be GPU only).
+
+After you register forward or backward functions with context name `“gpu”`, MXNet will dispatch forward or backward GPU functions you just registered when the operator is invoked with GPU context.
+
+In the registered `forwardGPU` function, specify the grid and block size and launch the CUDA kernel. In the GPU operators input and output tensors are pre-allocated on the GPU memory, just like in the CPU operator tensors are pre-allocated on the CPU memory. As a result, you don’t need to call `cudaMemcpy` to move the tensor data to GPU device.
+
+```c++
+ MXReturnValue forwardGPU(std::map<std::string, std::string> attrs,
+ std::vector<MXTensor> inputs,
+ std::vector<MXTensor> outputs,
+ OpResource res) {
+ float* in_data = inputs[0].data<float>();
+ float* out_data = outputs[0].data<float>();
+ mx_stream_t cuda_stream = res.get_cuda_stream();
+ int64_t N = inputs[0].size();
+ int block = 256;
+ int grid = (N + (block - 1)) / block;
+ my_op_gpu_forward<<<grid,block,0,cuda_stream>>>(out_data, in_data, N);
+ return MX_SUCCESS;
+ }
+```
+
+Note that the `cuda_stream` object used for launching kernels is passed from MXNet backend via `OpResource` object. See below for details of `Operator Resource`.
+
+At this point all the attribute functions for each operator (`parseAttrs`, `inferShape`, etc.) run on the CPU, including the `forwardGPU` function. The only part that will actually run on the GPU is the CUDA kernel function that is launched from `forwardGPU` function (ie. my_op_gpu_forward).
+
+```c++
+ __global__ void my_op_gpu_forward(float *out, float *in, int64_t N) {
+ int tid = blockIdx.x * blockDim.x + threadIdx.x;
Review comment:
It is correct and readable to replace `int tid = blockIdx.x * blockDim.x + threadIdx.x;` with the `CUDA_KERNEL_LOOP` macro.
----------------------------------------------------------------
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
With regards,
Apache Git Services
[GitHub] [incubator-mxnet] samskalicky commented on a change in pull request
#17486: Update CustomOp doc with changes for GPU support
Posted by GitBox <gi...@apache.org>.
samskalicky commented on a change in pull request #17486: Update CustomOp doc with changes for GPU support
URL: https://github.com/apache/incubator-mxnet/pull/17486#discussion_r375022066
##########
File path: example/extensions/lib_custom_op/README.md
##########
@@ -56,87 +72,111 @@ For building a library containing your own custom operator, compose a C++ source
- `forward` - Forward Computation (can be replace with `createOpState`, see below for details)
Then compile it to `libmyop_lib.so` dynamic library using the following command:
+
```bash
g++ -shared -fPIC -std=c++11 myop_lib.cc -o libmyop_lib.so -I ../../../include/mxnet
```
+If you don't want to download MXNet source and choose only using `lib_api.h` header, you can copy the header over to the same folder of `myop_lib.cc` and run:
+
+```bash
+g++ -shared -fPIC -std=c++11 myop_lib.cc -o libmyop_lib.so
+```
+
Finally, you can write a Python script to load the library and run your custom operator:
+
```python
import mxnet as mx
mx.library.load(‘libmyop_lib.so’)
mx.nd.my_op(...)
```
-### Writing Regular Custom Operator:
+### Writing A Regular Custom Operator
-There are several essential building blocks for making a (stateless) custom operator:
+There are several essential building blocks for making a custom operator:
* [initialize](./gemm_lib.cc#L227):
* This function is the library initialization function necessary for any dynamic libraries. It checks if you are using a compatible version of MXNet. Note that this `version` parameter is passed from MXNet when library is loaded.
- MXReturnValue initialize(int version)
+```c++
+ MXReturnValue initialize(int version)
+```
* [parseAttrs](./gemm_lib.cc#L118):
* This function specifies number of input and output tensors for the custom operator; also this is where a custom operator can validate the attributes (ie. options) specified by the user.
- MXReturnValue parseAttrs(
- std::map<std::string,
- std::string> attrs,
- int* num_in,
- int* num_out)
+```c++
+ MXReturnValue parseAttrs(
+ std::map<std::string,
+ std::string> attrs,
+ int* num_in,
+ int* num_out)
+```
* [inferType](./gemm_lib.cc#L124):
* This function specifies how the custom operator infers output data types using input data types.
- MXReturnValue inferType(
- std::map<std::string, std::string> attrs,
- std::vector<int> &intypes,
- std::vector<int> &outtypes)
+```c++
+ MXReturnValue inferType(
+ std::map<std::string, std::string> attrs,
+ std::vector<int> &intypes,
+ std::vector<int> &outtypes)
+```
* [inferShape](./gemm_lib.cc#L143):
* This function specifies how the custom operator infers output tensor shape using input shape.
- MXReturnValue inferShape(
- std::map<std::string, std::string> attrs,
- std::vector<std::vector<unsigned int>> &inshapes,
- std::vector<std::vector<unsigned int>> &outshapes)
+```c++
+ MXReturnValue inferShape(
+ std::map<std::string, std::string> attrs,
+ std::vector<std::vector<unsigned int>> &inshapes,
+ std::vector<std::vector<unsigned int>> &outshapes)
+```
* [forward](./gemm_lib.cc#L56):
* This function specifies the computation of the forward pass of the operator.
- MXReturnValue forward(
- std::map<std::string, std::string> attrs,
- std::vector<MXTensor> inputs,
- std::vector<MXTensor> outputs,
- OpResource res)
+```c++
+ MXReturnValue forward(
+ std::map<std::string, std::string> attrs,
+ std::vector<MXTensor> inputs,
+ std::vector<MXTensor> outputs,
+ OpResource res)
+```
* [REGISTER_OP(my_op_name)](./gemm_lib.cc#L169):
- * This macro registers the custom operator and its properties to MXNet NDArray and Symbol APIs by its name.
-
- REGISTER_OP(my_op_name)
- .setForward(forward)
- .setParseAttrs(parseAttrs)
- .setInferType(inferType)
- .setInferShape(inferShape);
+ * This macro registers the custom operator and its properties to MXNet NDArray and Symbol APIs by its name. Note that for operator running on CPU, you need pass the name of the context `"cpu"` when registering forward or backward function.
Review comment:
This macro registers the custom operator with MXNet. You can register additional functions with the operator too. Note that you register functions for each context for Forward and Backward, here we show an example for CPU. These are the minimum required functions, but you can specify more as needed.
----------------------------------------------------------------
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
With regards,
Apache Git Services
[GitHub] [incubator-mxnet] aaronmarkham commented on a change in pull
request #17486: Update CustomOp doc with changes for GPU support
Posted by GitBox <gi...@apache.org>.
aaronmarkham commented on a change in pull request #17486: Update CustomOp doc with changes for GPU support
URL: https://github.com/apache/incubator-mxnet/pull/17486#discussion_r373217213
##########
File path: example/extensions/lib_custom_op/README.md
##########
@@ -146,21 +168,154 @@ Let’s take a closer look at those registry functions:
* **inferShape**: This function is similar to the `inferType` function, except it is used for populating the output data shapes. You need to figure out the shapes of each output tensors for this computation. For example, if the inputs are images with shape (224,224,3) and you write a padding operator to make 10px borders for the images, then your output shape will be (234,234,3).
-* **forward**: This function executes the main forward computation. It takes four arguments. The 1st argument is the attributes. The 2nd argument is the input `MXTensors` which stores all data and info of input ndarrays. The 3rd argument is the output `MXTensors`. The 4th argument is the `OpResource` object for memory allocation and other utilities. Additionally, you can use a `dltensor` tensor structure stored in the `MXTensor` as a more standardized data structure for computing.
+* **forward**: This function executes the main forward computation. It takes four arguments. The 1st argument is the attributes. The 2nd argument is the input `MXTensors` which stores all data and info of input ndarrays. The 3rd argument is the output `MXTensors`. The 4th argument is the `OpResource` object for memory allocation and other utilities. The details of `OpResource` are covered in the below section.
+Additionally, you can use a `dltensor` tensor structure stored in the `MXTensor` as a more standardized data structure for computing.
* **backward**: This function is doing the backward gradient computation. It will be similar to the forward function. And you need to figure out the formula of the backward gradient computation.
* **mutateInputs**: This function is for marking mutable inputs. It takes two arguments. The 1st argument is the attributes. The 2nd argument is a list of input indices that are mutable among all input tensors. It is useful when some inputs are auxiliary model parameters and might be altered during forward/backward computation. Remember, the index number of `input_indices` should not exceed the number of inputs.
-### Writing Stateful Custom Operator:
+### Writing Stateful Custom Operator
A stateful custom operator is useful when a forward/backward call needs some data or ‘state’ from previous forward/backward calls. Normally we create a class, and make instance variables store the states used for computing or caching.
Most of the building blocks for making a stateful custom operator is the same as regular custom operator, except it’ll register `createOpState` instead of a `forward` function for the computation.
* [createOpState](./gemm_lib.cc#L204) - Create stateful operator instance:
* This function takes two arguments. The 1st argument is attributes. The 2nd argument is a placeholder for `CustomStatefulOp` object. You must [define a class that inherits CustomStatefulOp](./gemm_lib.cc#L178) and override the forward function (optionally the backward function). Then you need to create an instance of your class and assign it to the placeholder. In this way, all of the forward/backward calls will use the same methods in that instance, and the instance is able to keep the state of the operator.
+```c++
+ MXReturnValue createOpState(
+ std::map<std::string, std::string> attrs,
+ CustomStatefulOp** op_inst)
+```
+
+* The operator registering function will look like this:
+```c++
+ REGISTER_OP(my_state_op)
+ ...
+ .setCreateOpState(createOpState, "cpu");
+```
+
+## Writing Custom GPU Operator Library
+
+Most of the building blocks for registering GPU custom operators are the exactly same as CPU ones, except you need to specify the `"gpu"` context name when registering `forward`, `backward` or `createOpState` function.
+
+### Run A GPU Example
+
+For illustration purposes, we provided a `ReLU` (Rectified Linear Unit) activation operator that can run on GPU. Make sure you have installed a CUDA compatible MXNet build. Go to `lib_custom_op` directory and follow these steps:
+
+1. Run `make relu_lib`. The Makefile will invoke `NVCC` compiler to compile the CUDA kernel along with regular custom operator functions from `relu_lib.cu` to generate `librelu_lib.so` library.
+2. Run `python test_relu.py`. It’ll register the GPU `ReLU` operator in the MXNet backend, then invoke the operator by feeding an `NDArray` input with GPU context, and output the result tensor with GPU context.
+
+### Regular GPU Custom Operator
+
+Since most of the building blocks for registering GPU custom operators are the exactly same as CPU ones, the registering function for an operator supporting both GPU and CPU will look like this:
+```c++
+ REGISTER_OP(my_op_gpu)
+ ...
+ .setForward(forwardCPU, "cpu")
+ .setForward(forwardGPU, "gpu")
+ .setBackward(backwardCPU, "cpu")
+ .setBackward(backwardGPU, "gpu");
+```
+
+Note that you can write only the GPU funcitons, as operators don’t have to support both CPU and GPU functions (can be GPU only).
+
+After you register forward or backward functions with context name `“gpu”`, MXNet will dispatch forward or backward GPU functions you just registered when the operator is invoked with GPU context.
+
+In the registered `forwardGPU` function, specify the grid and block size and launch the CUDA kernel. In the GPU operators input and output tensors are pre-allocated on the GPU memory, just like in the CPU operator tensors are pre-allocated on the CPU memory. As a result, you don’t need to call `cudaMemcpy` to move the tensor data to GPU device.
+```c++
+ MXReturnValue forwardGPU(std::map<std::string, std::string> attrs,
+ std::vector<MXTensor> inputs,
+ std::vector<MXTensor> outputs,
+ OpResource res) {
+ float* in_data = inputs[0].data<float>();
+ float* out_data = outputs[0].data<float>();
+ mx_stream_t cuda_stream = res.get_cuda_stream();
+ int64_t N = inputs[0].size();
+ int block = 256;
+ int grid = (N + (block - 1)) / block;
+ my_op_gpu_forward<<<grid,block,0,cuda_stream>>>(out_data, in_data, N);
+ return MX_SUCCESS;
+ }
+```
+
+Note that the `cuda_stream` object used for launching kernels is passed from MXNet backend via `OpResource` object. See below for details of `Operator Resource`.
+
+At this point all the attribute functions for each operator (`parseAttrs`, `inferShape`, etc.) run on the CPU, including the `forwardGPU` function. The only part that will actually run on the GPU is the CUDA kernel function that is launched from `forwardGPU` function (ie. my_op_gpu_forward).
+```c++
+ __global__ void my_op_gpu_forward(float *out, float *in, int64_t N) {
+ int tid = blockIdx.x * blockDim.x + threadIdx.x;
+ if (tid < N) {
+ // code your computational logic here
+ }
+ }
+```
+
+### Stateful GPU Custom Operator
+
+Recall that for stateful custom operators, you need to define a class that inherits `CustomStatefulOp` and overrides the `forward` and `backward` functions. Stateful operators are created context-aware, so you can create different classes for GPU and CPU stateful operators separately if desired. To do so, you register a createOpState function for each context separately like this
+```c++
+ REGISTER_OP(my_state_op_gpu)
+ ...
+ .setCreateOpState(createOpStateCPU, "cpu")
+ .setCreateOpState(createOpStateGPU, "gpu");
+```
+
+Then you can create different classes for CPU and GPU stateful operators. MXNet will create the stateful operator instance based on the running context when the operator is invoked, and call stateful `forward` or `backward` function from the instantiated stateful operator class.
+```c++
Review comment:
```suggestion
```c++
```
----------------------------------------------------------------
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
With regards,
Apache Git Services
[GitHub] [incubator-mxnet] samskalicky commented on a change in pull request
#17486: Update CustomOp doc with changes for GPU support
Posted by GitBox <gi...@apache.org>.
samskalicky commented on a change in pull request #17486: Update CustomOp doc with changes for GPU support
URL: https://github.com/apache/incubator-mxnet/pull/17486#discussion_r375022474
##########
File path: example/extensions/lib_custom_op/README.md
##########
@@ -56,87 +72,111 @@ For building a library containing your own custom operator, compose a C++ source
- `forward` - Forward Computation (can be replace with `createOpState`, see below for details)
Then compile it to `libmyop_lib.so` dynamic library using the following command:
+
```bash
g++ -shared -fPIC -std=c++11 myop_lib.cc -o libmyop_lib.so -I ../../../include/mxnet
```
+If you don't want to download MXNet source and choose only using `lib_api.h` header, you can copy the header over to the same folder of `myop_lib.cc` and run:
+
+```bash
+g++ -shared -fPIC -std=c++11 myop_lib.cc -o libmyop_lib.so
+```
+
Finally, you can write a Python script to load the library and run your custom operator:
+
```python
import mxnet as mx
mx.library.load(‘libmyop_lib.so’)
mx.nd.my_op(...)
```
-### Writing Regular Custom Operator:
+### Writing A Regular Custom Operator
-There are several essential building blocks for making a (stateless) custom operator:
+There are several essential building blocks for making a custom operator:
* [initialize](./gemm_lib.cc#L227):
* This function is the library initialization function necessary for any dynamic libraries. It checks if you are using a compatible version of MXNet. Note that this `version` parameter is passed from MXNet when library is loaded.
- MXReturnValue initialize(int version)
+```c++
+ MXReturnValue initialize(int version)
+```
* [parseAttrs](./gemm_lib.cc#L118):
* This function specifies number of input and output tensors for the custom operator; also this is where a custom operator can validate the attributes (ie. options) specified by the user.
- MXReturnValue parseAttrs(
- std::map<std::string,
- std::string> attrs,
- int* num_in,
- int* num_out)
+```c++
+ MXReturnValue parseAttrs(
+ std::map<std::string,
+ std::string> attrs,
+ int* num_in,
+ int* num_out)
+```
* [inferType](./gemm_lib.cc#L124):
* This function specifies how the custom operator infers output data types using input data types.
- MXReturnValue inferType(
- std::map<std::string, std::string> attrs,
- std::vector<int> &intypes,
- std::vector<int> &outtypes)
+```c++
+ MXReturnValue inferType(
+ std::map<std::string, std::string> attrs,
+ std::vector<int> &intypes,
+ std::vector<int> &outtypes)
+```
* [inferShape](./gemm_lib.cc#L143):
* This function specifies how the custom operator infers output tensor shape using input shape.
- MXReturnValue inferShape(
- std::map<std::string, std::string> attrs,
- std::vector<std::vector<unsigned int>> &inshapes,
- std::vector<std::vector<unsigned int>> &outshapes)
+```c++
+ MXReturnValue inferShape(
+ std::map<std::string, std::string> attrs,
+ std::vector<std::vector<unsigned int>> &inshapes,
+ std::vector<std::vector<unsigned int>> &outshapes)
+```
* [forward](./gemm_lib.cc#L56):
* This function specifies the computation of the forward pass of the operator.
- MXReturnValue forward(
- std::map<std::string, std::string> attrs,
- std::vector<MXTensor> inputs,
- std::vector<MXTensor> outputs,
- OpResource res)
+```c++
+ MXReturnValue forward(
+ std::map<std::string, std::string> attrs,
+ std::vector<MXTensor> inputs,
+ std::vector<MXTensor> outputs,
+ OpResource res)
+```
* [REGISTER_OP(my_op_name)](./gemm_lib.cc#L169):
- * This macro registers the custom operator and its properties to MXNet NDArray and Symbol APIs by its name.
-
- REGISTER_OP(my_op_name)
- .setForward(forward)
- .setParseAttrs(parseAttrs)
- .setInferType(inferType)
- .setInferShape(inferShape);
+ * This macro registers the custom operator and its properties to MXNet NDArray and Symbol APIs by its name. Note that for operator running on CPU, you need pass the name of the context `"cpu"` when registering forward or backward function.
Review comment:
lets move this to the end
----------------------------------------------------------------
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
With regards,
Apache Git Services
[GitHub] [incubator-mxnet] samskalicky commented on a change in pull request
#17486: Update CustomOp doc with changes for GPU support
Posted by GitBox <gi...@apache.org>.
samskalicky commented on a change in pull request #17486: Update CustomOp doc with changes for GPU support
URL: https://github.com/apache/incubator-mxnet/pull/17486#discussion_r375060298
##########
File path: example/extensions/lib_custom_op/README.md
##########
@@ -146,21 +186,160 @@ Let’s take a closer look at those registry functions:
* **inferShape**: This function is similar to the `inferType` function, except it is used for populating the output data shapes. You need to figure out the shapes of each output tensors for this computation. For example, if the inputs are images with shape (224,224,3) and you write a padding operator to make 10px borders for the images, then your output shape will be (234,234,3).
-* **forward**: This function executes the main forward computation. It takes four arguments. The 1st argument is the attributes. The 2nd argument is the input `MXTensors` which stores all data and info of input ndarrays. The 3rd argument is the output `MXTensors`. The 4th argument is the `OpResource` object for memory allocation and other utilities. Additionally, you can use a `dltensor` tensor structure stored in the `MXTensor` as a more standardized data structure for computing.
+* **forward**: This function executes the main forward computation. It takes four arguments. The 1st argument is the attributes. The 2nd argument is the input `MXTensors` which stores all data and info of input ndarrays. The 3rd argument is the output `MXTensors`. The 4th argument is the `OpResource` object for memory allocation and other utilities. The details of `OpResource` are covered in the below section.
+Additionally, you can use a `dltensor` tensor structure stored in the `MXTensor` as a more standardized data structure for computing.
* **backward**: This function is doing the backward gradient computation. It will be similar to the forward function. And you need to figure out the formula of the backward gradient computation.
* **mutateInputs**: This function is for marking mutable inputs. It takes two arguments. The 1st argument is the attributes. The 2nd argument is a list of input indices that are mutable among all input tensors. It is useful when some inputs are auxiliary model parameters and might be altered during forward/backward computation. Remember, the index number of `input_indices` should not exceed the number of inputs.
-### Writing Stateful Custom Operator:
+### Writing Stateful Custom Operator
A stateful custom operator is useful when a forward/backward call needs some data or ‘state’ from previous forward/backward calls. Normally we create a class, and make instance variables store the states used for computing or caching.
Most of the building blocks for making a stateful custom operator is the same as regular custom operator, except it’ll register `createOpState` instead of a `forward` function for the computation.
* [createOpState](./gemm_lib.cc#L204) - Create stateful operator instance:
* This function takes two arguments. The 1st argument is attributes. The 2nd argument is a placeholder for `CustomStatefulOp` object. You must [define a class that inherits CustomStatefulOp](./gemm_lib.cc#L178) and override the forward function (optionally the backward function). Then you need to create an instance of your class and assign it to the placeholder. In this way, all of the forward/backward calls will use the same methods in that instance, and the instance is able to keep the state of the operator.
+```c++
+ MXReturnValue createOpState(
+ std::map<std::string, std::string> attrs,
+ CustomStatefulOp** op_inst)
+```
+
+* The operator registering function will look like this:
+```c++
+ REGISTER_OP(my_state_op)
+ ...
+ .setCreateOpState(createOpState, "cpu");
+```
+
+## Writing A Custom GPU Operator Library
+
+Most of the building blocks for registering GPU custom operators are the exactly same as CPU ones, except you need to specify the `"gpu"` context name when registering `forward`, `backward` or `createOpState` function.
+
+### Run A GPU Example
+
+For illustration purposes, we provided a `ReLU` (Rectified Linear Unit) activation operator that can run on GPU. Make sure you have installed a CUDA compatible MXNet build. Go to `lib_custom_op` directory and follow these steps:
+
+1. Run `make relu_lib`. The Makefile will invoke `NVCC` compiler to compile the CUDA kernel along with regular custom operator functions from `relu_lib.cu` to generate `librelu_lib.so` library.
+2. Run `python test_relu.py`. It’ll register the GPU `ReLU` operator in the MXNet backend, then invoke the operator by feeding an `NDArray` input with GPU context, and output the result tensor with GPU context.
+
+### Writing A Regular GPU Custom Operator
+
+Since most of the building blocks for registering GPU custom operators are the exactly same as CPU ones, the registering function for an operator supporting both GPU and CPU will look like this:
+
+```c++
+ REGISTER_OP(my_op_gpu)
+ ...
+ .setForward(forwardCPU, "cpu")
+ .setForward(forwardGPU, "gpu")
+ .setBackward(backwardCPU, "cpu")
+ .setBackward(backwardGPU, "gpu");
+```
+
+Note that you can write only the GPU funcitons, as operators don’t have to support both CPU and GPU functions (can be GPU only).
+
+After you register forward or backward functions with context name `“gpu”`, MXNet will dispatch forward or backward GPU functions you just registered when the operator is invoked with GPU context.
+
+In the registered `forwardGPU` function, specify the grid and block size and launch the CUDA kernel. In the GPU operators input and output tensors are pre-allocated on the GPU memory, just like in the CPU operator tensors are pre-allocated on the CPU memory. As a result, you don’t need to call `cudaMemcpy` to move the tensor data to GPU device.
+
+```c++
+ MXReturnValue forwardGPU(std::map<std::string, std::string> attrs,
+ std::vector<MXTensor> inputs,
+ std::vector<MXTensor> outputs,
+ OpResource res) {
+ float* in_data = inputs[0].data<float>();
+ float* out_data = outputs[0].data<float>();
+ mx_stream_t cuda_stream = res.get_cuda_stream();
+ int64_t N = inputs[0].size();
+ int block = 256;
+ int grid = (N + (block - 1)) / block;
+ my_op_gpu_forward<<<grid,block,0,cuda_stream>>>(out_data, in_data, N);
+ return MX_SUCCESS;
+ }
+```
+
+Note that the `cuda_stream` object used for launching kernels is passed from MXNet backend via `OpResource` object. See below for details of `Operator Resource`.
+
+At this point all the attribute functions for each operator (`parseAttrs`, `inferShape`, etc.) run on the CPU, including the `forwardGPU` function. The only part that will actually run on the GPU is the CUDA kernel function that is launched from `forwardGPU` function (ie. my_op_gpu_forward).
+
+```c++
+ __global__ void my_op_gpu_forward(float *out, float *in, int64_t N) {
+ int tid = blockIdx.x * blockDim.x + threadIdx.x;
Review comment:
I agree with @wkcn, lets just replace the function body with:
```
// code your computational logic here
```
Then it can be generic and not confusing to advanced users. The function body here doesnt really give any additional meaning
----------------------------------------------------------------
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
With regards,
Apache Git Services
[GitHub] [incubator-mxnet] rondogency commented on issue #17486: Update
CustomOp doc with changes for GPU support
Posted by GitBox <gi...@apache.org>.
rondogency commented on issue #17486: Update CustomOp doc with changes for GPU support
URL: https://github.com/apache/incubator-mxnet/pull/17486#issuecomment-586023537
@a550461053 Thanks for the questions. Currently you cannot call MXNet built-in operator from MXNet. We are making a plan for adding this support, but it cannot be easily implemented and may be a separate project. Also I added a snippet of how to use mutable input, hope it would help.
----------------------------------------------------------------
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
With regards,
Apache Git Services