rondogency opened a new pull request #17270: [WIP] Dynamic custom operator GPU support URL: https://github.com/apache/incubator-mxnet/pull/17270 ## Description ## Add custom operator GPU support to enable users to write custom operator running on GPU. This is a continuation of custom operator project https://github.com/apache/incubator-mxnet/pull/15921 ## Design ## The main constraint is to make operator CUDA code and custom operator registration being compiled by NVCC together. Working backward from the user, user will create a single .cu file, register a single operator to contain both CPU and GPU computation logic. The registration is the same as CPU operators. `REGISTER_OP(my_relu)` User should dispatch kernel function by checking MXTensor context in custom operator forward/backward function. Here we use a simple relu example: ` if (inputs[0].ctx.dev_type == MX_GPU){ cudaStream_t gpu_stream = reinterpret_cast<cudaStream_t>(res.get_gpu_stream()); int64_t N = inputs[0].size(); int grid = (N + 255) / 256; int block = 256; relu_gpu_forward<<<grid,block,0,gpu_stream>>>(out_data, in_data, N); } else { relu_cpu_forward(out_data, in_data, inputs[0].size()); }` Then user should write CUDA code snippet in that file for GPU kernel function. All computation here will be run in GPU. `__global__ void relu_gpu_forward(float *out, float *in, int64_t N) { int tid = blockIdx.x * blockDim.x + threadIdx.x; if (tid < N){ out[tid] = in[tid] > 0 ? in[tid] : 0; } }` ## Checklist ## ### Essentials ### - [ ] 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 ### - Add Fcompute<gpu> registration, and pass NDArray context to custom library in c_api.cc - Add context info to MXTensor class in lib_api.h - Add lib_custom_op/relu.cu example file containing full registration of custom operator "my_relu", and add both CPU and GPU kernel functions in that file - Modify lib_custom_op/Makefile to compile .cu file using nvcc to custom library ## 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: [email protected] With regards, Apache Git Services
