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));
}
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];
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>
int32_t channels_;
int32_t count_;
bool use_slm_;
+ bool log_softmax_;
UMat scale_data_;
};
#endif // HAVE_OPENCL
{
softmax_axis_ = config.axis;
channels_ = config.channels;
+ log_softmax_ = config.logsoftmax;
inner_num_ = 1;
outer_num_ = 1;
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
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;
}
}
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;
}
}