Merge pull request opencv#9829 from pengli:ocl4dnn

alalek · alalek · commit df5b2224d729 · 2017-10-12T11:26:20.000Z
diff --git a/modules/dnn/src/layers/softmax_layer.cpp b/modules/dnn/src/layers/softmax_layer.cpp
@@ -100,6 +100,7 @@ class SoftMaxLayerImpl : public SoftmaxLayer
             config.in_shape = shape(*inputs[0]);
             config.axis = axisRaw;
             config.channels = inputs[0]->size[axisRaw];
+            config.logsoftmax = logSoftMax;
 
             softmaxOp = Ptr<OCL4DNNSoftmax<float> >(new OCL4DNNSoftmax<float>(config));
         }
@@ -108,7 +109,7 @@ class SoftMaxLayerImpl : public SoftmaxLayer
         srcMat = inputs[0]->getUMat(ACCESS_READ);
         dstMat = outputs[0].getUMat(ACCESS_WRITE);
 
-        if (!logSoftMax && softmaxOp->Forward(srcMat, dstMat))
+        if (softmaxOp->Forward(srcMat, dstMat))
             return true;
 
         const Mat &src = *inputs[0];
diff --git a/modules/dnn/src/ocl4dnn/include/ocl4dnn.hpp b/modules/dnn/src/ocl4dnn/include/ocl4dnn.hpp
@@ -445,11 +445,12 @@ class OCL4DNNLRN
 
 struct OCL4DNNSoftmaxConfig
 {
-    OCL4DNNSoftmaxConfig() : axis(0), channels(0)
+    OCL4DNNSoftmaxConfig() : axis(0), channels(0), logsoftmax(false)
     {}
     MatShape in_shape;
     int axis;
     int channels;
+    bool logsoftmax;
 };
 
 template<typename Dtype>
@@ -467,6 +468,7 @@ class OCL4DNNSoftmax
         int32_t channels_;
         int32_t count_;
         bool use_slm_;
+        bool log_softmax_;
         UMat scale_data_;
 };
 #endif // HAVE_OPENCL
diff --git a/modules/dnn/src/ocl4dnn/src/ocl4dnn_softmax.cpp b/modules/dnn/src/ocl4dnn/src/ocl4dnn_softmax.cpp
@@ -52,6 +52,7 @@ OCL4DNNSoftmax<Dtype>::OCL4DNNSoftmax(OCL4DNNSoftmaxConfig config)
 {
     softmax_axis_ = config.axis;
     channels_ = config.channels;
+    log_softmax_ = config.logsoftmax;
 
     inner_num_ = 1;
     outer_num_ = 1;
@@ -90,6 +91,7 @@ bool OCL4DNNSoftmax<Dtype>::Forward(const UMat& bottom, UMat& top)
         String kname;
         ocl::Kernel oclk_softmax_forward_kernel;
 
+        if (log_softmax_) opts += " -DLOG_SOFTMAX ";
         if (use_slm_)
             kname = CL_KERNEL_SELECT("softmax_forward_slm");
         else
diff --git a/modules/dnn/src/opencl/softmax_loss.cl b/modules/dnn/src/opencl/softmax_loss.cl
@@ -112,7 +112,11 @@ __kernel void TEMPLATE(softmax_forward_slm,Dtype)(const int num, const int chann
   for (int index = get_global_id(0); index < channels * spatial_dim;
       index += get_global_size(0)) {
     int s = index % spatial_dim;
-    out[n * channels * spatial_dim + index] = out_tmp[index] / scale_tmp[s];
+    Dtype v = out_tmp[index] / scale_tmp[s];
+#ifdef LOG_SOFTMAX
+    v = log(v);
+#endif
+    out[n * channels * spatial_dim + index] = v;
   }
 }
 
@@ -177,6 +181,10 @@ __kernel void TEMPLATE(softmax_forward,Dtype)(const int num, const int channels,
   for (int index = get_global_id(0); index < channels * spatial_dim;
       index += get_global_size(0)) {
     int s = index % spatial_dim;
-    out[n * channels * spatial_dim + index] /= scale[n * spatial_dim + s];
+    Dtype v = out[n * channels * spatial_dim + index] / scale[n * spatial_dim + s];
+#ifdef LOG_SOFTMAX
+    v = log(v);
+#endif
+    out[n * channels * spatial_dim + index] = v;
   }
 }

Original file line number	Diff line number	Diff line change
`@@ -112,7 +112,11 @@ __kernel void TEMPLATE(softmax_forward_slm,Dtype)(const int num, const int chann`
`112`	`112`	`for (int index = get_global_id(0); index < channels * spatial_dim;`
`113`	`113`	`index += get_global_size(0)) {`
`114`	`114`	`int s = index % spatial_dim;`
`115`		`- out[n * channels * spatial_dim + index] = out_tmp[index] / scale_tmp[s];`
	`115`	`+ Dtype v = out_tmp[index] / scale_tmp[s];`
	`116`	`+#ifdef LOG_SOFTMAX`
	`117`	`+ v = log(v);`
	`118`	`+#endif`
	`119`	`+ out[n * channels * spatial_dim + index] = v;`
`116`	`120`	`}`
`117`	`121`	`}`
`118`	`122`
`@@ -177,6 +181,10 @@ __kernel void TEMPLATE(softmax_forward,Dtype)(const int num, const int channels,`
`177`	`181`	`for (int index = get_global_id(0); index < channels * spatial_dim;`
`178`	`182`	`index += get_global_size(0)) {`
`179`	`183`	`int s = index % spatial_dim;`
`180`		`- out[n * channels * spatial_dim + index] /= scale[n * spatial_dim + s];`
	`184`	`+ Dtype v = out[n * channels * spatial_dim + index] / scale[n * spatial_dim + s];`
	`185`	`+#ifdef LOG_SOFTMAX`
	`186`	`+ v = log(v);`
	`187`	`+#endif`
	`188`	`+ out[n * channels * spatial_dim + index] = v;`
`181`	`189`	`}`
`182`	`190`	`}`