piiswrong commented on a change in pull request #10018: Bug Fix and performance 
optimized for rtc
URL: https://github.com/apache/incubator-mxnet/pull/10018#discussion_r172750653
 
 

 ##########
 File path: example/numpy-ops/custom_softmax_rtc.py
 ##########
 @@ -23,78 +23,91 @@
 
 class Softmax(mx.operator.CustomOp):
     def __init__(self):
-        self.fwd_kernel_mod = None
-        self.bwd_kernel_mod = None
-        super().__init__()
+        super(Softmax,self).__init__()
+        # Each thread processes a row (a sample in the batch).
+        fwd_src = r"""
+            template<class DType>
+            __global__ void fwd(const DType* x, DType* y, const int row_size, 
const int req) {
+                const int offset = row_size * threadIdx.x;
+                DType max = x[offset];
+                for(int i = 1; i < row_size; ++i) {
+                    if(max < x[offset + i]) {
+                        max = x[offset + i];
+                    }
+                }
+                DType sum = 0;
+                for(int i = 0; i < row_size; ++i) {
+                    sum += exp(x[offset + i] - max);
+                }
+                switch(req) {
+                    case 1:
+                        for(int i = 0; i < row_size; ++i) {
+                            y[offset + i] = exp(x[offset + i] - max) / sum;
+                        }
+                        break;
+                    case 2:
+                        for(int i = 0; i < row_size; ++i) {
+                            y[offset + i] += exp(x[offset + i] - max) / sum;
+                        }
+                        break;
+                }
+            }
+        """
+
+        # Each block processes a row and each thread in a block calculate an 
element of `dx`.
+        bwd_src = r"""
+            template<class DType>
+            __global__ void bwd(const DType* l, const DType* y, DType* dx, 
const int req) {
+                const int z = static_cast<int>(l[blockIdx.x]);
+                const int i = threadIdx.x + blockDim.x * blockIdx.x;
+                if(req == 1) {
+                    dx[i]  = threadIdx.x == z ? y[i] - 1 : y[i];
+                } else {
+                    dx[i] += threadIdx.x == z ? y[i] - 1 : y[i];
+                }
+            }
+        """
+        fwd_kernel_mod = mx.rtc.CudaModule(fwd_src, exports=["fwd<float>", 
"fwd<double>"])
+        bwd_kernel_mod = mx.rtc.CudaModule(bwd_src, exports=["bwd<float>", 
"bwd<double>"])
+
+        fwd_kernel_float_signature = "const {0}*, const {0}*, const int, const 
int".format("float")
 
 Review comment:
   why use format here? seems unnecessary

----------------------------------------------------------------
This is an automated message from the Apache Git Service.
To respond to the message, please log on GitHub and use the
URL above to go to the specific comment.
 
For queries about this service, please contact Infrastructure at:
us...@infra.apache.org


With regards,
Apache Git Services

Reply via email to