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/03/04 19:27:35 UTC
[GitHub] [incubator-mxnet] rondogency opened a new pull request #17762:
Custom Operator Random Number Generator Support
rondogency opened a new pull request #17762: Custom Operator Random Number Generator Support
URL: https://github.com/apache/incubator-mxnet/pull/17762
## Description ##
Add random number generator support for custom operator libraries.
Custom library now can call MXNet random number generator to get random numbers through OpResource class in forward/backward computation.
## 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 ###
- change lib_api.h OpResource class to support all random number generator API calls and returns the generated number
- change c_api.cc to init the RandomGenerator object when custom library invokes any API calls, and eventually invokes MXNet functions to generator random numbers
- currently we use GPU RNG whenever possible
## Comments ##
this is a continuation of the custom operator project #15921 and #17270
----------------------------------------------------------------
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 #17762: Custom Operator
Random Number Generator Support
Posted by GitBox <gi...@apache.org>.
wkcn commented on issue #17762: Custom Operator Random Number Generator Support
URL: https://github.com/apache/incubator-mxnet/pull/17762#issuecomment-609115979
@mxnet-bot run ci [unix-gpu]
----------------------------------------------------------------
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
#17762: Custom Operator Random Number Generator Support
Posted by GitBox <gi...@apache.org>.
samskalicky commented on a change in pull request #17762: Custom Operator Random Number Generator Support
URL: https://github.com/apache/incubator-mxnet/pull/17762#discussion_r399628432
##########
File path: src/c_api/c_api.cc
##########
@@ -809,11 +822,10 @@ int MXLoadLib(const char *path) {
regOp.set_attr<mxnet::FInferShape>("FInferShape", DefaultSubgraphOpShape, plevel);
regOp.set_attr<FInferStorageType>("FInferStorageType",
DefaultSubgraphOpStorageType, plevel);
- regOp.set_attr<FResourceRequest>("FResourceRequest",
- DefaultSubgraphOpResourceRequest, plevel);
regOp.set_attr<nnvm::FMutateInputs>("FMutateInputs",
DefaultSubgraphOpMutableInputs, plevel);
}
+ regOp.set_attr<FResourceRequest>("FResourceRequest", resc_req, plevel);
Review comment:
can we put this after line 798 instead? its kinda weird hanging out here in the middle
----------------------------------------------------------------
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
#17762: Custom Operator Random Number Generator Support
Posted by GitBox <gi...@apache.org>.
samskalicky commented on a change in pull request #17762: Custom Operator Random Number Generator Support
URL: https://github.com/apache/incubator-mxnet/pull/17762#discussion_r399431812
##########
File path: include/mxnet/lib_api.h
##########
@@ -1026,7 +1057,8 @@ typedef int (*opCallFStatefulComp_t)(int is_forward, void* state_op,
void** in_indices, void** out_indices,
void** in_indptr, void** out_indptr,
int64_t* in_indices_shapes, int64_t* out_indices_shapes,
- int64_t* in_indptr_shapes, int64_t* out_indptr_shapes);
+ int64_t* in_indptr_shapes, int64_t* out_indptr_shapes,
+ void* cpu_states, void* gpu_states);
Review comment:
can we rename it here 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] rondogency commented on issue #17762: Custom
Operator Random Number Generator Support
Posted by GitBox <gi...@apache.org>.
rondogency commented on issue #17762: Custom Operator Random Number Generator Support
URL: https://github.com/apache/incubator-mxnet/pull/17762#issuecomment-604864330
@wkcn @eric-haibin-lin this PR is ready for review. I have to reimplement the workflow to get rid of some segfault error so the code is different from last time you reviewed
----------------------------------------------------------------
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
#17762: Custom Operator Random Number Generator Support
Posted by GitBox <gi...@apache.org>.
samskalicky commented on a change in pull request #17762: Custom Operator Random Number Generator Support
URL: https://github.com/apache/incubator-mxnet/pull/17762#discussion_r399432221
##########
File path: include/mxnet/lib_api.h
##########
@@ -1419,7 +1452,8 @@ extern "C" {
int* instypes, int* outstypes, void** in_indices, void** out_indices,
void** in_indptr, void** out_indptr,
int64_t* in_indices_shapes, int64_t* out_indices_shapes,
- int64_t* in_indptr_shapes, int64_t* out_indptr_shapes) {
+ int64_t* in_indptr_shapes, int64_t* out_indptr_shapes,
+ void* cpu_states, void* gpu_states) {
Review comment:
can we rename it here 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] leezu merged pull request #17762: Custom Operator
Random Number Generator Support
Posted by GitBox <gi...@apache.org>.
leezu merged pull request #17762: Custom Operator Random Number Generator Support
URL: https://github.com/apache/incubator-mxnet/pull/17762
----------------------------------------------------------------
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
#17762: Custom Operator Random Number Generator Support
Posted by GitBox <gi...@apache.org>.
samskalicky commented on a change in pull request #17762: Custom Operator Random Number Generator Support
URL: https://github.com/apache/incubator-mxnet/pull/17762#discussion_r399046049
##########
File path: include/mxnet/lib_api.h
##########
@@ -997,7 +1015,8 @@ typedef int (*opCallFComp_t)(fcomp_t fcomp, const char* const* keys,
void** in_indices, void** out_indices,
void** in_indptr, void** out_indptr,
int64_t* in_indices_shapes, int64_t* out_indices_shapes,
- int64_t* in_indptr_shapes, int64_t* out_indptr_shapes);
+ int64_t* in_indptr_shapes, int64_t* out_indptr_shapes,
+ void* cpu_states, void* gpu_states);
Review comment:
can we call these rand_ cpu_states and rand_gpu_states so its more clear that these are for RNG?
----------------------------------------------------------------
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
#17762: Custom Operator Random Number Generator Support
Posted by GitBox <gi...@apache.org>.
rondogency commented on a change in pull request #17762: Custom Operator Random Number Generator Support
URL: https://github.com/apache/incubator-mxnet/pull/17762#discussion_r399085405
##########
File path: example/extensions/lib_custom_op/relu_lib.cu
##########
@@ -180,6 +182,75 @@ REGISTER_OP(my_state_relu)
.setCreateOpState(createOpStateCPU, "cpu")
.setCreateOpState(createOpStateGPU, "gpu");
+
+
+/* ------------------------ Below is noisy relu operator example ---------------------*/
+
+#include <curand_kernel.h>
+#include <random>
+
+#define NumRandomPerThread 64 // mxnet recommended random numbers generated per thread
+
+__global__ void noisy_relu_gpu_forward(float *out, float *in, int64_t N, void *states, int step) {
+ int tid = blockIdx.x * blockDim.x + threadIdx.x;
+ curandStatePhilox4_32_10_t* global_states = (curandStatePhilox4_32_10_t*)states;
Review comment:
Added to comments to explain the logic, and made some type tweaks
----------------------------------------------------------------
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
#17762: Custom Operator Random Number Generator Support
Posted by GitBox <gi...@apache.org>.
samskalicky commented on a change in pull request #17762: Custom Operator Random Number Generator Support
URL: https://github.com/apache/incubator-mxnet/pull/17762#discussion_r399043126
##########
File path: example/extensions/lib_custom_op/relu_lib.cu
##########
@@ -180,6 +182,75 @@ REGISTER_OP(my_state_relu)
.setCreateOpState(createOpStateCPU, "cpu")
.setCreateOpState(createOpStateGPU, "gpu");
+
+
+/* ------------------------ Below is noisy relu operator example ---------------------*/
+
+#include <curand_kernel.h>
+#include <random>
+
+#define NumRandomPerThread 64 // mxnet recommended random numbers generated per thread
+
+__global__ void noisy_relu_gpu_forward(float *out, float *in, int64_t N, void *states, int step) {
+ int tid = blockIdx.x * blockDim.x + threadIdx.x;
+ curandStatePhilox4_32_10_t* global_states = (curandStatePhilox4_32_10_t*)states;
+ curandStatePhilox4_32_10_t thread_state = global_states[tid];
+
+ int start = tid * step;
+ int end = start + step;
+ for (int i=start; i<end && i<N; ++i) {
+ float noise = curand_normal(&thread_state);
+ out[i] = in[i] + noise > 0 ? in[i] + noise : 0;
Review comment:
can you add some comments to explain how/why you're calculating this for noisy relu?
----------------------------------------------------------------
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
#17762: Custom Operator Random Number Generator Support
Posted by GitBox <gi...@apache.org>.
samskalicky commented on a change in pull request #17762: Custom Operator Random Number Generator Support
URL: https://github.com/apache/incubator-mxnet/pull/17762#discussion_r399430932
##########
File path: include/mxnet/lib_api.h
##########
@@ -426,21 +432,30 @@ typedef void (*sparse_malloc_t)(void*, int, int, int, void**, int64_t**, int64_t
#if defined(__NVCC__)
typedef cudaStream_t mx_stream_t;
+ typedef curandStatePhilox4_32_10_t* mx_gpu_rand_pt;
Review comment:
since this is a typedef, lets follow the naming style and call it mx_gpu_rand_t
should we typedef std::mt19937 to mx_cpu_rand_t too? or is that going too far?
----------------------------------------------------------------
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
#17762: Custom Operator Random Number Generator Support
Posted by GitBox <gi...@apache.org>.
samskalicky commented on a change in pull request #17762: Custom Operator Random Number Generator Support
URL: https://github.com/apache/incubator-mxnet/pull/17762#discussion_r399431922
##########
File path: include/mxnet/lib_api.h
##########
@@ -1284,7 +1316,8 @@ extern "C" {
int* instypes, int* outstypes, void** in_indices, void** out_indices,
void** in_indptr, void** out_indptr,
int64_t* in_indices_shapes, int64_t* out_indices_shapes,
- int64_t* in_indptr_shapes, int64_t* out_indptr_shapes) {
+ int64_t* in_indptr_shapes, int64_t* out_indptr_shapes,
+ void* cpu_states, void* gpu_states) {
Review comment:
can we rename it here 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
#17762: Custom Operator Random Number Generator Support
Posted by GitBox <gi...@apache.org>.
samskalicky commented on a change in pull request #17762: Custom Operator Random Number Generator Support
URL: https://github.com/apache/incubator-mxnet/pull/17762#discussion_r399440931
##########
File path: example/extensions/lib_custom_op/relu_lib.cu
##########
@@ -180,6 +182,80 @@ REGISTER_OP(my_state_relu)
.setCreateOpState(createOpStateCPU, "cpu")
.setCreateOpState(createOpStateGPU, "gpu");
+/*
+ * Below is noisy ReLU operator example
+ * noisy ReLU is made from ReLU extended to include Gaussian noise
+ * forward - add Gaussian noise generated from normal distribution to each unit
+ * backward - gradient doesn't need to change since noise is constant
+ */
+
+#define NumRandomPerThread 64 // mxnet recommended random numbers generated per thread
+
+__global__ void noisy_relu_gpu_forward(float *out, float *in, int64_t N, mx_gpu_rand_pt states, int step) {
+ // the launcher logic ensures tid less than NumGPURandomStates
+ int tid = blockIdx.x * blockDim.x + threadIdx.x;
+ // each thread generates unique sequence of random numbers
+ curandStatePhilox4_32_10_t thread_state = states[tid];
+ // each thread works on <step> number of calculation
+ int start = tid * step;
+ int end = start + step;
+ for (int i=start; i<end && i<N; ++i) {
+ float noise = curand_normal(&thread_state);
+ out[i] = in[i] + noise > 0 ? in[i] + noise : 0;
+ }
+}
+
+MXReturnValue noisyForwardCPU(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>();
+
+ std::mt19937 *states = res.get_cpu_rand_states();
+ std::normal_distribution<float> dist_normal;
+
+ for (int i=0; i<inputs[0].size(); ++i) {
+ float noise = dist_normal(*states);
+ out_data[i] = in_data[i] + noise > 0 ? in_data[i] + noise : 0;
+ }
+ return MX_SUCCESS;
+}
+
+MXReturnValue noisyForwardGPU(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();
+
+ // below is mxnet recommended workflow to parallel random number generating
+ int num_thread = (N + NumRandomPerThread - 1) / NumRandomPerThread;
+ // we should not launch more threads than mxnet supported random number GPU states
+ int num_thread_need = num_thread < NumGPURandomStates ? num_thread : NumGPURandomStates;
+ // each cuda thread processes [step * tid, step * id + step) snippet of input tensor
+ int step = (N + num_thread_need - 1) / num_thread_need;
+ // this can ensure number of parallel threads less than mxnet supported random number states
+ int num_block = (num_thread_need + NumThreadPerBlock - 1) / NumThreadPerBlock;
Review comment:
this is great! will be super helpful for users.
----------------------------------------------------------------
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
#17762: Custom Operator Random Number Generator Support
Posted by GitBox <gi...@apache.org>.
samskalicky commented on a change in pull request #17762: Custom Operator Random Number Generator Support
URL: https://github.com/apache/incubator-mxnet/pull/17762#discussion_r399431539
##########
File path: include/mxnet/lib_api.h
##########
@@ -997,7 +1015,8 @@ typedef int (*opCallFComp_t)(fcomp_t fcomp, const char* const* keys,
void** in_indices, void** out_indices,
void** in_indptr, void** out_indptr,
int64_t* in_indices_shapes, int64_t* out_indices_shapes,
- int64_t* in_indptr_shapes, int64_t* out_indptr_shapes);
+ int64_t* in_indptr_shapes, int64_t* out_indptr_shapes,
+ void* cpu_states, void* gpu_states);
Review comment:
its still called cpu_states here, can we rename it here 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
#17762: Custom Operator Random Number Generator Support
Posted by GitBox <gi...@apache.org>.
samskalicky commented on a change in pull request #17762: Custom Operator Random Number Generator Support
URL: https://github.com/apache/incubator-mxnet/pull/17762#discussion_r399493839
##########
File path: src/c_api/c_api.cc
##########
@@ -248,11 +251,24 @@ void CustomFComputeDispatcher(const std::string op_name,
// get actual cudaStream_t out of mxnet gpu stream and pass to lib_api.h
void *cuda_stream = nullptr;
#if MXNET_USE_CUDA
- if (inputs[0].ctx().dev_mask() == Context::kGPU) {
+ if (inputs.size() > 0 && inputs[0].ctx().dev_mask() == Context::kGPU) {
cuda_stream = static_cast<void*>(gpu_stream->stream_);
}
#endif
+ // get mxnet initialized and seeded RNG states and pass to lib_api.h
+ void *rng_cpu_states = nullptr, *rng_gpu_states = nullptr;
+ if (!is_subgraph_op) {
Review comment:
DefaultSubgraphOpResourceRequest calls DefaultSubgraphOpResourceRequestHelper which goes through the ops in the subgraph and sets all resources requested by any op in the subgraph. One of these could be a custom op who's defining the resource request.
You have a good point though, we should probably just have subgraph ops define all supported resources anyway.
----------------------------------------------------------------
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
#17762: Custom Operator Random Number Generator Support
Posted by GitBox <gi...@apache.org>.
rondogency commented on a change in pull request #17762: Custom Operator Random Number Generator Support
URL: https://github.com/apache/incubator-mxnet/pull/17762#discussion_r399441495
##########
File path: src/c_api/c_api.cc
##########
@@ -248,11 +251,24 @@ void CustomFComputeDispatcher(const std::string op_name,
// get actual cudaStream_t out of mxnet gpu stream and pass to lib_api.h
void *cuda_stream = nullptr;
#if MXNET_USE_CUDA
- if (inputs[0].ctx().dev_mask() == Context::kGPU) {
+ if (inputs.size() > 0 && inputs[0].ctx().dev_mask() == Context::kGPU) {
cuda_stream = static_cast<void*>(gpu_stream->stream_);
}
#endif
+ // get mxnet initialized and seeded RNG states and pass to lib_api.h
+ void *rng_cpu_states = nullptr, *rng_gpu_states = nullptr;
+ if (!is_subgraph_op) {
Review comment:
subgraph op is using DefaultSubgraphOpResourceRequest, so if no operator in the subgraph resource request pass has parallel random, resource request will error out
but I think I should check initialized resource types instead of check subgraph op, I will make the change
----------------------------------------------------------------
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 #17762: Custom
Operator Random Number Generator Support
Posted by GitBox <gi...@apache.org>.
rondogency commented on issue #17762: Custom Operator Random Number Generator Support
URL: https://github.com/apache/incubator-mxnet/pull/17762#issuecomment-602386557
in the cmake, the line "target_compile_options(subgraph_lib PUBLIC -shared)" can be removed, as we added the library as shared already
----------------------------------------------------------------
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
#17762: Custom Operator Random Number Generator Support
Posted by GitBox <gi...@apache.org>.
rondogency commented on a change in pull request #17762: Custom Operator Random Number Generator Support
URL: https://github.com/apache/incubator-mxnet/pull/17762#discussion_r387893988
##########
File path: include/mxnet/lib_api.h
##########
@@ -39,7 +39,8 @@
#include <utility>
#include <stdexcept>
-#define MX_LIBRARY_VERSION 3
+/* Make sure to update the version number everytime you make changes */
+#define MX_LIBRARY_VERSION 4
Review comment:
I know that's why I put that note and will resolve conflict when merging
----------------------------------------------------------------
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] mxnet-bot commented on issue #17762: Custom
Operator Random Number Generator Support
Posted by GitBox <gi...@apache.org>.
mxnet-bot commented on issue #17762: Custom Operator Random Number Generator Support
URL: https://github.com/apache/incubator-mxnet/pull/17762#issuecomment-609475765
Jenkins CI successfully triggered : [unix-gpu]
----------------------------------------------------------------
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
#17762: Custom Operator Random Number Generator Support
Posted by GitBox <gi...@apache.org>.
samskalicky commented on a change in pull request #17762: Custom Operator Random Number Generator Support
URL: https://github.com/apache/incubator-mxnet/pull/17762#discussion_r399043504
##########
File path: example/extensions/lib_custom_op/relu_lib.cu
##########
@@ -180,6 +182,75 @@ REGISTER_OP(my_state_relu)
.setCreateOpState(createOpStateCPU, "cpu")
.setCreateOpState(createOpStateGPU, "gpu");
+
+
+/* ------------------------ Below is noisy relu operator example ---------------------*/
+
+#include <curand_kernel.h>
+#include <random>
+
+#define NumRandomPerThread 64 // mxnet recommended random numbers generated per thread
+
+__global__ void noisy_relu_gpu_forward(float *out, float *in, int64_t N, void *states, int step) {
+ int tid = blockIdx.x * blockDim.x + threadIdx.x;
+ curandStatePhilox4_32_10_t* global_states = (curandStatePhilox4_32_10_t*)states;
Review comment:
can you add a comment to explain why you're doing this? where do the states get allocated? how can you guarantee the states pointer is pointing to a big enough array to index with tid?
----------------------------------------------------------------
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 #17762: Custom
Operator Random Number Generator Support
Posted by GitBox <gi...@apache.org>.
rondogency commented on issue #17762: Custom Operator Random Number Generator Support
URL: https://github.com/apache/incubator-mxnet/pull/17762#issuecomment-610092971
@mxnet-bot run ci [unix-gpu]
----------------------------------------------------------------
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 #17762: Custom Operator
Random Number Generator Support
Posted by GitBox <gi...@apache.org>.
wkcn commented on issue #17762: Custom Operator Random Number Generator Support
URL: https://github.com/apache/incubator-mxnet/pull/17762#issuecomment-609008854
@mxnet-bot run ci [unix-gpu, centos-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
#17762: Custom Operator Random Number Generator Support
Posted by GitBox <gi...@apache.org>.
samskalicky commented on a change in pull request #17762: Custom Operator Random Number Generator Support
URL: https://github.com/apache/incubator-mxnet/pull/17762#discussion_r387889850
##########
File path: include/mxnet/lib_api.h
##########
@@ -39,7 +39,8 @@
#include <utility>
#include <stdexcept>
-#define MX_LIBRARY_VERSION 3
+/* Make sure to update the version number everytime you make changes */
+#define MX_LIBRARY_VERSION 4
Review comment:
Just heads up, #17569 and #17623 are all racing to see who can change from 3 to 4 first. The other two will need to change from 4 to 5 afterwards...
----------------------------------------------------------------
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 #17762: Custom
Operator Random Number Generator Support
Posted by GitBox <gi...@apache.org>.
rondogency commented on issue #17762: Custom Operator Random Number Generator Support
URL: https://github.com/apache/incubator-mxnet/pull/17762#issuecomment-610516730
@mxnet-bot run ci [unix-gpu, windows-gpu]
----------------------------------------------------------------
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
#17762: Custom Operator Random Number Generator Support
Posted by GitBox <gi...@apache.org>.
samskalicky commented on a change in pull request #17762: Custom Operator Random Number Generator Support
URL: https://github.com/apache/incubator-mxnet/pull/17762#discussion_r399420491
##########
File path: include/mxnet/lib_api.h
##########
@@ -426,21 +432,30 @@ typedef void (*sparse_malloc_t)(void*, int, int, int, void**, int64_t**, int64_t
#if defined(__NVCC__)
typedef cudaStream_t mx_stream_t;
+ typedef curandStatePhilox4_32_10_t* mx_gpu_rand_pt;
#else
typedef void* mx_stream_t;
+ typedef void* mx_gpu_rand_pt;
#endif
+/*! \brief MXNet initialized random states for each device, used for parallelism */
+/* Each thread should generate random number unique sequence out of different states */
+#define NumCPURandomStates 1024
+#define NumGPURandomStates 32768
Review comment:
Can we prefix these with MX_ and use all caps?
MX_CPU_RANDOM_STATES
MX_GPU_RANDOM_STATES
----------------------------------------------------------------
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] mxnet-bot commented on issue #17762: Custom
Operator Random Number Generator Support
Posted by GitBox <gi...@apache.org>.
mxnet-bot commented on issue #17762: Custom Operator Random Number Generator Support
URL: https://github.com/apache/incubator-mxnet/pull/17762#issuecomment-609008876
Jenkins CI successfully triggered : [centos-cpu, unix-gpu]
----------------------------------------------------------------
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] mxnet-bot commented on issue #17762: Custom
Operator Random Number Generator Support
Posted by GitBox <gi...@apache.org>.
mxnet-bot commented on issue #17762: Custom Operator Random Number Generator Support
URL: https://github.com/apache/incubator-mxnet/pull/17762#issuecomment-609115991
Jenkins CI successfully triggered : [unix-gpu]
----------------------------------------------------------------
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
#17762: Custom Operator Random Number Generator Support
Posted by GitBox <gi...@apache.org>.
samskalicky commented on a change in pull request #17762: Custom Operator Random Number Generator Support
URL: https://github.com/apache/incubator-mxnet/pull/17762#discussion_r387889850
##########
File path: include/mxnet/lib_api.h
##########
@@ -39,7 +39,8 @@
#include <utility>
#include <stdexcept>
-#define MX_LIBRARY_VERSION 3
+/* Make sure to update the version number everytime you make changes */
+#define MX_LIBRARY_VERSION 4
Review comment:
Just heads up, #17569 and #17623 are also racing to see who can change from 3 to 4 first. The other two will need to change from 4 to 5 afterwards...
----------------------------------------------------------------
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
#17762: Custom Operator Random Number Generator Support
Posted by GitBox <gi...@apache.org>.
samskalicky commented on a change in pull request #17762: Custom Operator Random Number Generator Support
URL: https://github.com/apache/incubator-mxnet/pull/17762#discussion_r399628115
##########
File path: src/c_api/c_api.cc
##########
@@ -248,11 +250,21 @@ void CustomFComputeDispatcher(const std::string op_name,
// get actual cudaStream_t out of mxnet gpu stream and pass to lib_api.h
void *cuda_stream = nullptr;
#if MXNET_USE_CUDA
- if (inputs[0].ctx().dev_mask() == Context::kGPU) {
+ if (inputs.size() > 0 && inputs[0].ctx().dev_mask() == Context::kGPU) {
Review comment:
while you're making this change, its possible an operator might have zero inputs and some outputs (ie. generator operator, like random), should we also check for outputs if inputs is zero like:
```
if ((inputs.size() > 0 && inputs[0].ctx().dev_mask() == Context::kGPU) ||
(outputs.size() > 0 && inputs[0].ctx().dev_mask() == Context::kGPU)) {
```
----------------------------------------------------------------
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
#17762: Custom Operator Random Number Generator Support
Posted by GitBox <gi...@apache.org>.
samskalicky commented on a change in pull request #17762: Custom Operator Random Number Generator Support
URL: https://github.com/apache/incubator-mxnet/pull/17762#discussion_r399432283
##########
File path: include/mxnet/lib_api.h
##########
@@ -1476,7 +1510,7 @@ extern "C" {
}
OpResource res(cpu_malloc, cpu_alloc, gpu_malloc, gpu_alloc,
- stream, sparse_malloc, sparse_alloc);
+ stream, sparse_malloc, sparse_alloc, cpu_states, gpu_states);
Review comment:
can we rename it here 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
#17762: Custom Operator Random Number Generator Support
Posted by GitBox <gi...@apache.org>.
samskalicky commented on a change in pull request #17762: Custom Operator Random Number Generator Support
URL: https://github.com/apache/incubator-mxnet/pull/17762#discussion_r399045512
##########
File path: include/mxnet/lib_api.h
##########
@@ -437,10 +438,12 @@ class OpResource {
public:
OpResource(xpu_malloc_t cpu_malloc_fp, void* cpu_alloc_fp,
xpu_malloc_t gpu_malloc_fp, void* gpu_alloc_fp, void* stream,
- sparse_malloc_t sparse_malloc_fp, void* sparse_alloc_fp)
+ sparse_malloc_t sparse_malloc_fp, void* sparse_alloc_fp,
+ void* cpu_states, void* gpu_states)
Review comment:
can we call these cpu_rand_states and gpu_rand_states so that its clear that these are for the RNG?
----------------------------------------------------------------
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
#17762: Custom Operator Random Number Generator Support
Posted by GitBox <gi...@apache.org>.
rondogency commented on a change in pull request #17762: Custom Operator Random Number Generator Support
URL: https://github.com/apache/incubator-mxnet/pull/17762#discussion_r399085026
##########
File path: example/extensions/lib_custom_op/relu_lib.cu
##########
@@ -180,6 +182,75 @@ REGISTER_OP(my_state_relu)
.setCreateOpState(createOpStateCPU, "cpu")
.setCreateOpState(createOpStateGPU, "gpu");
+
+
+/* ------------------------ Below is noisy relu operator example ---------------------*/
+
+#include <curand_kernel.h>
+#include <random>
+
+#define NumRandomPerThread 64 // mxnet recommended random numbers generated per thread
+
+__global__ void noisy_relu_gpu_forward(float *out, float *in, int64_t N, void *states, int step) {
+ int tid = blockIdx.x * blockDim.x + threadIdx.x;
+ curandStatePhilox4_32_10_t* global_states = (curandStatePhilox4_32_10_t*)states;
+ curandStatePhilox4_32_10_t thread_state = global_states[tid];
+
+ int start = tid * step;
+ int end = start + step;
+ for (int i=start; i<end && i<N; ++i) {
+ float noise = curand_normal(&thread_state);
+ out[i] = in[i] + noise > 0 ? in[i] + noise : 0;
Review comment:
Added comments with details
----------------------------------------------------------------
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 #17762: Custom
Operator Random Number Generator Support
Posted by GitBox <gi...@apache.org>.
rondogency commented on issue #17762: Custom Operator Random Number Generator Support
URL: https://github.com/apache/incubator-mxnet/pull/17762#issuecomment-605319052
@samskalicky resolved all your comments
----------------------------------------------------------------
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
#17762: Custom Operator Random Number Generator Support
Posted by GitBox <gi...@apache.org>.
samskalicky commented on a change in pull request #17762: Custom Operator Random Number Generator Support
URL: https://github.com/apache/incubator-mxnet/pull/17762#discussion_r399045189
##########
File path: example/extensions/lib_custom_op/test_relu.py
##########
@@ -63,22 +63,33 @@
print(in_grad_base)
print("--------start testing larger ndarray---------")
-a = mx.nd.uniform(shape=(100,100,100), ctx=mx.cpu())
b = mx.nd.uniform(shape=(100,100,100), ctx=mx.gpu())
mx.nd.waitall()
t1 = time.time()
-r1 = mx.nd.my_relu(a)
+r1 = mx.nd.my_relu(b)
mx.nd.waitall()
t2 = time.time()
-r2 = mx.nd.my_relu(b)
+r2 = mx.nd.relu(b)
mx.nd.waitall()
t3 = time.time()
-r3 = mx.nd.relu(b)
-mx.nd.waitall()
-t4 = time.time()
-print("CPU running time:")
+print("Custom GPU running time:")
Review comment:
Custom GPU ==> Custom relu
----------------------------------------------------------------
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
#17762: Custom Operator Random Number Generator Support
Posted by GitBox <gi...@apache.org>.
wkcn commented on a change in pull request #17762: Custom Operator Random Number Generator Support
URL: https://github.com/apache/incubator-mxnet/pull/17762#discussion_r388646210
##########
File path: src/c_api/c_api.cc
##########
@@ -177,10 +177,39 @@ void CustomFComputeDispatcher(const std::string op_name,
return static_cast<void*>((*gpualloc)(size));
};
+ // MXNet random number generator lambda to be invoked by custom library
+ // custom library will use gpu rng whenever possible
+ auto rng_caller = [&](RandGenType rand_type, int seed) {
+#if MXNET_USE_CUDA
Review comment:
Agree to making a default choice : )
I think it should depend on the context of custom operator, rather that the version (GPU or only CPU) of MXNet.
If users build a CPU custom operator and use the GPU version of MXNet, on a machine with only CPU, it still work since the CPU random number generator is used.
----------------------------------------------------------------
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
#17762: Custom Operator Random Number Generator Support
Posted by GitBox <gi...@apache.org>.
rondogency commented on a change in pull request #17762: Custom Operator Random Number Generator Support
URL: https://github.com/apache/incubator-mxnet/pull/17762#discussion_r399935574
##########
File path: src/c_api/c_api.cc
##########
@@ -809,11 +822,10 @@ int MXLoadLib(const char *path) {
regOp.set_attr<mxnet::FInferShape>("FInferShape", DefaultSubgraphOpShape, plevel);
regOp.set_attr<FInferStorageType>("FInferStorageType",
DefaultSubgraphOpStorageType, plevel);
- regOp.set_attr<FResourceRequest>("FResourceRequest",
- DefaultSubgraphOpResourceRequest, plevel);
regOp.set_attr<nnvm::FMutateInputs>("FMutateInputs",
DefaultSubgraphOpMutableInputs, plevel);
}
+ regOp.set_attr<FResourceRequest>("FResourceRequest", resc_req, plevel);
Review comment:
put it to 806 for readability
----------------------------------------------------------------
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
#17762: Custom Operator Random Number Generator Support
Posted by GitBox <gi...@apache.org>.
samskalicky commented on a change in pull request #17762: Custom Operator Random Number Generator Support
URL: https://github.com/apache/incubator-mxnet/pull/17762#discussion_r399045224
##########
File path: example/extensions/lib_custom_op/test_relu.py
##########
@@ -63,22 +63,33 @@
print(in_grad_base)
print("--------start testing larger ndarray---------")
-a = mx.nd.uniform(shape=(100,100,100), ctx=mx.cpu())
b = mx.nd.uniform(shape=(100,100,100), ctx=mx.gpu())
mx.nd.waitall()
t1 = time.time()
-r1 = mx.nd.my_relu(a)
+r1 = mx.nd.my_relu(b)
mx.nd.waitall()
t2 = time.time()
-r2 = mx.nd.my_relu(b)
+r2 = mx.nd.relu(b)
mx.nd.waitall()
t3 = time.time()
-r3 = mx.nd.relu(b)
-mx.nd.waitall()
-t4 = time.time()
-print("CPU running time:")
+print("Custom GPU running time:")
print(t2 - t1)
-print("GPU running time:")
+print("Native GPU running time:")
Review comment:
Native GPU ==> Native relu
----------------------------------------------------------------
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
#17762: Custom Operator Random Number Generator Support
Posted by GitBox <gi...@apache.org>.
samskalicky commented on a change in pull request #17762: Custom Operator Random Number Generator Support
URL: https://github.com/apache/incubator-mxnet/pull/17762#discussion_r399045773
##########
File path: include/mxnet/lib_api.h
##########
@@ -463,6 +466,19 @@ class OpResource {
&(sparse->data), &(sparse->indices), &(sparse->indptr));
}
+ /*! \brief get pointer to cpu random number inited and seeded states */
+ /* this global states are located on cpu, of type std::mt19937 */
+ void* get_cpu_rand_states() {
+ return cpu_rand_states;
Review comment:
is there any reason not to return already casted to std::mt19937?
----------------------------------------------------------------
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
#17762: Custom Operator Random Number Generator Support
Posted by GitBox <gi...@apache.org>.
samskalicky commented on a change in pull request #17762: Custom Operator Random Number Generator Support
URL: https://github.com/apache/incubator-mxnet/pull/17762#discussion_r399493839
##########
File path: src/c_api/c_api.cc
##########
@@ -248,11 +251,24 @@ void CustomFComputeDispatcher(const std::string op_name,
// get actual cudaStream_t out of mxnet gpu stream and pass to lib_api.h
void *cuda_stream = nullptr;
#if MXNET_USE_CUDA
- if (inputs[0].ctx().dev_mask() == Context::kGPU) {
+ if (inputs.size() > 0 && inputs[0].ctx().dev_mask() == Context::kGPU) {
cuda_stream = static_cast<void*>(gpu_stream->stream_);
}
#endif
+ // get mxnet initialized and seeded RNG states and pass to lib_api.h
+ void *rng_cpu_states = nullptr, *rng_gpu_states = nullptr;
+ if (!is_subgraph_op) {
Review comment:
DefaultSubgraphOpResourceRequest calls DefaultSubgraphOpResourceRequestHelper which goes through the ops in the subgraph and sets all resources requested by any op in the subgraph. One of these could be a custom op who's defining the resource request.
https://github.com/apache/incubator-mxnet/blob/eceb5f2c1c494094c1a697286f2c4560b7ca472e/src/operator/subgraph/common.h#L242-L258
You have a good point though, we should probably just have subgraph ops define all supported resources anyway.
----------------------------------------------------------------
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 #17762: Custom
Operator Random Number Generator Support
Posted by GitBox <gi...@apache.org>.
samskalicky commented on issue #17762: Custom Operator Random Number Generator Support
URL: https://github.com/apache/incubator-mxnet/pull/17762#issuecomment-594782225
Thanks for this @rondogency! this is going to be a great addition for customOps.
A couple suggestions on the PR description:
Add random number generator support for custom operator**s** ~libraries~.
Custom ~library~ **operators** now can call MXNet random number generator to get random numbers through OpResource class in forward/backward computation.
----------------------------------------------------------------
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
#17762: Custom Operator Random Number Generator Support
Posted by GitBox <gi...@apache.org>.
samskalicky commented on a change in pull request #17762: Custom Operator Random Number Generator Support
URL: https://github.com/apache/incubator-mxnet/pull/17762#discussion_r399628115
##########
File path: src/c_api/c_api.cc
##########
@@ -248,11 +250,21 @@ void CustomFComputeDispatcher(const std::string op_name,
// get actual cudaStream_t out of mxnet gpu stream and pass to lib_api.h
void *cuda_stream = nullptr;
#if MXNET_USE_CUDA
- if (inputs[0].ctx().dev_mask() == Context::kGPU) {
+ if (inputs.size() > 0 && inputs[0].ctx().dev_mask() == Context::kGPU) {
Review comment:
while you're making this change, its possible an operator might have zero inputs and some outputs (ie. generator operator, like random), should we also check for outputs if inputs is zero like:
```
if ((inputs.size() > 0 && inputs[0].ctx().dev_mask() == Context::kGPU) ||
(outputs.size() > 0 && inputs[0].ctx().dev_mask() == Context::kGPU)) {
```
or does it really not matter, if compiled with MXNET_USE_CUDA should we just always set
```
cuda_stream = static_cast<void*>(gpu_stream->stream_);
```
----------------------------------------------------------------
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
#17762: Custom Operator Random Number Generator Support
Posted by GitBox <gi...@apache.org>.
samskalicky commented on a change in pull request #17762: Custom Operator Random Number Generator Support
URL: https://github.com/apache/incubator-mxnet/pull/17762#discussion_r399046049
##########
File path: include/mxnet/lib_api.h
##########
@@ -997,7 +1015,8 @@ typedef int (*opCallFComp_t)(fcomp_t fcomp, const char* const* keys,
void** in_indices, void** out_indices,
void** in_indptr, void** out_indptr,
int64_t* in_indices_shapes, int64_t* out_indices_shapes,
- int64_t* in_indptr_shapes, int64_t* out_indptr_shapes);
+ int64_t* in_indptr_shapes, int64_t* out_indptr_shapes,
+ void* cpu_states, void* gpu_states);
Review comment:
can we call these rand_ cpu_states and rand_gpu_states (or something with rng/rand in it) so its more clear that these are for RNG?
----------------------------------------------------------------
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
#17762: Custom Operator Random Number Generator Support
Posted by GitBox <gi...@apache.org>.
samskalicky commented on a change in pull request #17762: Custom Operator Random Number Generator Support
URL: https://github.com/apache/incubator-mxnet/pull/17762#discussion_r399045512
##########
File path: include/mxnet/lib_api.h
##########
@@ -437,10 +438,12 @@ class OpResource {
public:
OpResource(xpu_malloc_t cpu_malloc_fp, void* cpu_alloc_fp,
xpu_malloc_t gpu_malloc_fp, void* gpu_alloc_fp, void* stream,
- sparse_malloc_t sparse_malloc_fp, void* sparse_alloc_fp)
+ sparse_malloc_t sparse_malloc_fp, void* sparse_alloc_fp,
+ void* cpu_states, void* gpu_states)
Review comment:
can we call these cpu_rand_states and gpu_rand_states so that its clear that these are for the RNG?
----------------------------------------------------------------
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
#17762: Custom Operator Random Number Generator Support
Posted by GitBox <gi...@apache.org>.
samskalicky commented on a change in pull request #17762: Custom Operator Random Number Generator Support
URL: https://github.com/apache/incubator-mxnet/pull/17762#discussion_r388647625
##########
File path: src/c_api/c_api.cc
##########
@@ -177,10 +177,39 @@ void CustomFComputeDispatcher(const std::string op_name,
return static_cast<void*>((*gpualloc)(size));
};
+ // MXNet random number generator lambda to be invoked by custom library
+ // custom library will use gpu rng whenever possible
+ auto rng_caller = [&](RandGenType rand_type, int seed) {
+#if MXNET_USE_CUDA
Review comment:
I dont think we should have a default, user should be explicit which one they want to use. otherwise we introduce possible confusion. this follows the existing style:
https://github.com/apache/incubator-mxnet/blob/7c61cdadeb29cb9280697b245156eeea7501e001/include/mxnet/lib_api.h#L381-L389
Theres no default "alloc" you have to say which context to 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
#17762: Custom Operator Random Number Generator Support
Posted by GitBox <gi...@apache.org>.
samskalicky commented on a change in pull request #17762: Custom Operator Random Number Generator Support
URL: https://github.com/apache/incubator-mxnet/pull/17762#discussion_r399432140
##########
File path: include/mxnet/lib_api.h
##########
@@ -1345,7 +1378,7 @@ extern "C" {
}
OpResource res(cpu_malloc, cpu_alloc, gpu_malloc, gpu_alloc,
- cuda_stream, sparse_malloc, sparse_alloc);
+ cuda_stream, sparse_malloc, sparse_alloc, cpu_states, gpu_states);
Review comment:
can we rename it here too?
Just dont want "states" to be confused with stateful ops
----------------------------------------------------------------
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 #17762: Custom
Operator Random Number Generator Support
Posted by GitBox <gi...@apache.org>.
samskalicky commented on issue #17762: Custom Operator Random Number Generator Support
URL: https://github.com/apache/incubator-mxnet/pull/17762#issuecomment-605405522
@mxnet-label-bot add [pr-awaiting-merge]
----------------------------------------------------------------
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 #17762: Custom
Operator Random Number Generator Support
Posted by GitBox <gi...@apache.org>.
rondogency commented on issue #17762: Custom Operator Random Number Generator Support
URL: https://github.com/apache/incubator-mxnet/pull/17762#issuecomment-609475750
@mxnet-bot run ci [unix-gpu]
----------------------------------------------------------------
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
#17762: Custom Operator Random Number Generator Support
Posted by GitBox <gi...@apache.org>.
rondogency commented on a change in pull request #17762: Custom Operator Random Number Generator Support
URL: https://github.com/apache/incubator-mxnet/pull/17762#discussion_r399935679
##########
File path: src/c_api/c_api.cc
##########
@@ -248,11 +250,21 @@ void CustomFComputeDispatcher(const std::string op_name,
// get actual cudaStream_t out of mxnet gpu stream and pass to lib_api.h
void *cuda_stream = nullptr;
#if MXNET_USE_CUDA
- if (inputs[0].ctx().dev_mask() == Context::kGPU) {
+ if (inputs.size() > 0 && inputs[0].ctx().dev_mask() == Context::kGPU) {
Review comment:
I think we just check for output
----------------------------------------------------------------
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] mxnet-bot commented on issue #17762: Custom
Operator Random Number Generator Support
Posted by GitBox <gi...@apache.org>.
mxnet-bot commented on issue #17762: Custom Operator Random Number Generator Support
URL: https://github.com/apache/incubator-mxnet/pull/17762#issuecomment-610093020
Jenkins CI successfully triggered : [unix-gpu]
----------------------------------------------------------------
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
#17762: Custom Operator Random Number Generator Support
Posted by GitBox <gi...@apache.org>.
samskalicky commented on a change in pull request #17762: Custom Operator Random Number Generator Support
URL: https://github.com/apache/incubator-mxnet/pull/17762#discussion_r388647625
##########
File path: src/c_api/c_api.cc
##########
@@ -177,10 +177,39 @@ void CustomFComputeDispatcher(const std::string op_name,
return static_cast<void*>((*gpualloc)(size));
};
+ // MXNet random number generator lambda to be invoked by custom library
+ // custom library will use gpu rng whenever possible
+ auto rng_caller = [&](RandGenType rand_type, int seed) {
+#if MXNET_USE_CUDA
Review comment:
I dont think we should have a default, user should be explicit which one they want to use. otherwise we introduce possible confusion. this follows the existing style:
https://github.com/apache/incubator-mxnet/blob/7c61cdadeb29cb9280697b245156eeea7501e001/include/mxnet/lib_api.h#L381-L389
----------------------------------------------------------------
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
#17762: Custom Operator Random Number Generator Support
Posted by GitBox <gi...@apache.org>.
samskalicky commented on a change in pull request #17762: Custom Operator Random Number Generator Support
URL: https://github.com/apache/incubator-mxnet/pull/17762#discussion_r399046416
##########
File path: src/c_api/c_api.cc
##########
@@ -248,11 +249,22 @@ void CustomFComputeDispatcher(const std::string op_name,
// get actual cudaStream_t out of mxnet gpu stream and pass to lib_api.h
void *cuda_stream = nullptr;
#if MXNET_USE_CUDA
- if (inputs[0].ctx().dev_mask() == Context::kGPU) {
+ if (inputs.size() > 0 && inputs[0].ctx().dev_mask() == Context::kGPU) {
cuda_stream = static_cast<void*>(gpu_stream->stream_);
}
#endif
+ // get mxnet inited and seeded rng states and pass to lib_api.h
Review comment:
inited ==> 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] samskalicky commented on a change in pull request
#17762: Custom Operator Random Number Generator Support
Posted by GitBox <gi...@apache.org>.
samskalicky commented on a change in pull request #17762: Custom Operator Random Number Generator Support
URL: https://github.com/apache/incubator-mxnet/pull/17762#discussion_r399045838
##########
File path: include/mxnet/lib_api.h
##########
@@ -463,6 +466,19 @@ class OpResource {
&(sparse->data), &(sparse->indices), &(sparse->indptr));
}
+ /*! \brief get pointer to cpu random number inited and seeded states */
+ /* this global states are located on cpu, of type std::mt19937 */
+ void* get_cpu_rand_states() {
+ return cpu_rand_states;
+ }
+
+ /*! \brief get pointer to gpu random number inited and seeded states */
+ /* this global states are located on gpu, of type curandStatePhilox4_32_10_t */
+ /* note that if you are usign cpu build, then it will return a nullptr */
+ void* get_gpu_rand_states() {
+ return gpu_rand_states;
Review comment:
is there any reason not to return already casted to curandStatePhilox4_32_10_t?
----------------------------------------------------------------
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
#17762: Custom Operator Random Number Generator Support
Posted by GitBox <gi...@apache.org>.
samskalicky commented on a change in pull request #17762: Custom Operator Random Number Generator Support
URL: https://github.com/apache/incubator-mxnet/pull/17762#discussion_r399433340
##########
File path: src/c_api/c_api.cc
##########
@@ -248,11 +251,24 @@ void CustomFComputeDispatcher(const std::string op_name,
// get actual cudaStream_t out of mxnet gpu stream and pass to lib_api.h
void *cuda_stream = nullptr;
#if MXNET_USE_CUDA
- if (inputs[0].ctx().dev_mask() == Context::kGPU) {
+ if (inputs.size() > 0 && inputs[0].ctx().dev_mask() == Context::kGPU) {
cuda_stream = static_cast<void*>(gpu_stream->stream_);
}
#endif
+ // get mxnet initialized and seeded RNG states and pass to lib_api.h
+ void *rng_cpu_states = nullptr, *rng_gpu_states = nullptr;
+ if (!is_subgraph_op) {
Review comment:
why are we only doing this for non-subgraph ops? Subgraph ops cant get RNG?
----------------------------------------------------------------
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 #17762: Custom
Operator Random Number Generator Support
Posted by GitBox <gi...@apache.org>.
samskalicky commented on issue #17762: Custom Operator Random Number Generator Support
URL: https://github.com/apache/incubator-mxnet/pull/17762#issuecomment-595091428
> Great. how are we testing this feature?
I think @rondogency's plan is to write a dropout 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] wkcn commented on a change in pull request
#17762: Custom Operator Random Number Generator Support
Posted by GitBox <gi...@apache.org>.
wkcn commented on a change in pull request #17762: Custom Operator Random Number Generator Support
URL: https://github.com/apache/incubator-mxnet/pull/17762#discussion_r388013014
##########
File path: src/c_api/c_api.cc
##########
@@ -177,10 +177,39 @@ void CustomFComputeDispatcher(const std::string op_name,
return static_cast<void*>((*gpualloc)(size));
};
+ // MXNet random number generator lambda to be invoked by custom library
+ // custom library will use gpu rng whenever possible
+ auto rng_caller = [&](RandGenType rand_type, int seed) {
+#if MXNET_USE_CUDA
Review comment:
Some users may use MXNet built with `MXNET_USE_CUDA=1` on the machine with only CPU. In this case, the GPU random number generator is not available.
It is better to decide which device to choose in the user-side code, namely the code of custom operators.
----------------------------------------------------------------
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] mxnet-bot commented on issue #17762: Custom
Operator Random Number Generator Support
Posted by GitBox <gi...@apache.org>.
mxnet-bot commented on issue #17762: Custom Operator Random Number Generator Support
URL: https://github.com/apache/incubator-mxnet/pull/17762#issuecomment-610516774
Jenkins CI successfully triggered : [windows-gpu, unix-gpu]
----------------------------------------------------------------
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 #17762: Custom
Operator Random Number Generator Support
Posted by GitBox <gi...@apache.org>.
rondogency commented on issue #17762: Custom Operator Random Number Generator Support
URL: https://github.com/apache/incubator-mxnet/pull/17762#issuecomment-604864937
@sxjscience @yzhliu I will appreciate if you can take a look at this PR and point out if any random number related use cases that need to be covered or tested
----------------------------------------------------------------
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
#17762: Custom Operator Random Number Generator Support
Posted by GitBox <gi...@apache.org>.
samskalicky commented on a change in pull request #17762: Custom Operator Random Number Generator Support
URL: https://github.com/apache/incubator-mxnet/pull/17762#discussion_r399042740
##########
File path: example/extensions/lib_custom_op/relu_lib.cu
##########
@@ -180,6 +182,75 @@ REGISTER_OP(my_state_relu)
.setCreateOpState(createOpStateCPU, "cpu")
.setCreateOpState(createOpStateGPU, "gpu");
+
+
+/* ------------------------ Below is noisy relu operator example ---------------------*/
+
+#include <curand_kernel.h>
+#include <random>
+
+#define NumRandomPerThread 64 // mxnet recommended random numbers generated per thread
+
+__global__ void noisy_relu_gpu_forward(float *out, float *in, int64_t N, void *states, int step) {
+ int tid = blockIdx.x * blockDim.x + threadIdx.x;
+ curandStatePhilox4_32_10_t* global_states = (curandStatePhilox4_32_10_t*)states;
Review comment:
use static_cast<curandStatePhilox4_32_10_t*>(states)
----------------------------------------------------------------
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 #17762: Custom
Operator Random Number Generator Support
Posted by GitBox <gi...@apache.org>.
rondogency commented on issue #17762: Custom Operator Random Number Generator Support
URL: https://github.com/apache/incubator-mxnet/pull/17762#issuecomment-604863884
@samskalicky resolved your comments, please take another look
----------------------------------------------------------------
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
#17762: Custom Operator Random Number Generator Support
Posted by GitBox <gi...@apache.org>.
rondogency commented on a change in pull request #17762: Custom Operator Random Number Generator Support
URL: https://github.com/apache/incubator-mxnet/pull/17762#discussion_r388083204
##########
File path: src/c_api/c_api.cc
##########
@@ -177,10 +177,39 @@ void CustomFComputeDispatcher(const std::string op_name,
return static_cast<void*>((*gpualloc)(size));
};
+ // MXNet random number generator lambda to be invoked by custom library
+ // custom library will use gpu rng whenever possible
+ auto rng_caller = [&](RandGenType rand_type, int seed) {
+#if MXNET_USE_CUDA
Review comment:
thanks for the timely review. we can let user pass in a "cpu" or "gpu" string to decide it, just like the forward/backward function
Also I am just wondering if we should make a default choice of 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 a change in pull request
#17762: Custom Operator Random Number Generator Support
Posted by GitBox <gi...@apache.org>.
samskalicky commented on a change in pull request #17762: Custom Operator Random Number Generator Support
URL: https://github.com/apache/incubator-mxnet/pull/17762#discussion_r388647625
##########
File path: src/c_api/c_api.cc
##########
@@ -177,10 +177,39 @@ void CustomFComputeDispatcher(const std::string op_name,
return static_cast<void*>((*gpualloc)(size));
};
+ // MXNet random number generator lambda to be invoked by custom library
+ // custom library will use gpu rng whenever possible
+ auto rng_caller = [&](RandGenType rand_type, int seed) {
+#if MXNET_USE_CUDA
Review comment:
I dont think we should have a default, user should be explicit which one they want to use. otherwise we introduce possible confusion
----------------------------------------------------------------
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