Skip to content

optimize cross entropy kernel by using reduce. #4237

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
lcy-seso merged 6 commits into from
Sep 26, 2017
Merged

optimize cross entropy kernel by using reduce. #4237

Show file tree
Hide file tree
Changes from 1 commit
Commits
Show all changes
6 commits
Select commit Hold shift + click to select a range
a3a8a09
optimize cross entropy kernel by using reduce.
lcy-seso Sep 20, 2017
141b8db
update the backward kernel.
lcy-seso Sep 21, 2017
30bfaab
Merge branch 'develop' into optimize_cross_entropy_kernel
lcy-seso Sep 22, 2017
6735585
fix cpu kernel with soft labels.
lcy-seso Sep 22, 2017
201c2bc
delete redundant codes.
lcy-seso Sep 23, 2017
000d751
fix backward op.
lcy-seso Sep 25, 2017
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
36 changes: 27 additions & 9 deletions paddle/operators/cross_entropy_op.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -32,16 +32,33 @@ __global__ void CrossEntropyKernel(T* Y, const T* X, const int* label,
}
}

template <typename T>
template <typename T, int blockSize>
__global__ void SoftCrossEntropyKernel(T* Y, const T* X, const T* label,
const int N, const int D) {
for (int i = blockIdx.x * blockDim.x + threadIdx.x; i < N;
i += blockDim.x * gridDim.x) {
T sum = static_cast<T>(0);
for (int j = 0; j < D; j++) {
sum += label[i * D + j] * tolerable_value(log(X[i * D + j]));
int tid = threadIdx.x;
__shared__ T d_sum[blockSize];
int next_idx = blockIdx.x * D + tid;

d_sum[tid] = 0;
int cur_idx = tid;
while (cur_idx < D) {
d_sum[tid] += tolerable_value(std::log(X[next_idx])) * label[next_idx];
next_idx += blockSize;
cur_idx += blockSize;
}
__syncthreads();

for (int stride = blockSize >> 1; stride > 0; stride >>= 1) {
__syncthreads();
if (tid < stride) {
next_idx = tid + stride;
d_sum[tid] += d_sum[next_idx];
}
Y[i] = -sum;
}
__syncthreads();

if (tid == 0) {
Y[blockIdx.x] = -d_sum[0];
}
}

Expand Down Expand Up @@ -104,8 +121,9 @@ class CrossEntropyOpCUDAKernel : public framework::OpKernel {
// base on ExecutionContext.
if (ctx.Attr<int>("soft_label") == 1) {
auto* label_data = ctx.Input<Tensor>("Label")->data<T>();
SoftCrossEntropyKernel<T><<<grid, block>>>(y_data, x_data, label_data, n,
d);
grid = d;
SoftCrossEntropyKernel<T, 512><<<grid, block>>>(y_data, x_data,
Copy link
Contributor

@qingqing01 qingqing01 Sep 22, 2017

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

去掉line 123， 写成：

SoftCrossEntropyKernel<T, 512><<<d, block>>>

Copy link
Contributor Author

@lcy-seso lcy-seso Sep 22, 2017

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

done.

label_data, n, d);
} else {
auto* label_data = ctx.Input<Tensor>("Label")->data<int>();
CrossEntropyKernel<T><<<grid, block>>>(y_data, x_data, label_data, n, d);
Expand Down
27 changes: 14 additions & 13 deletions python/paddle/v2/framework/tests/test_cross_entropy_op.py
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -19,7 +19,7 @@ def setUp(self):
dtype="float32")
self.inputs = {"X": X, "Label": label}
self.outputs = {"Y": cross_entropy}
self.attrs = {'soft_label': 0}
self.attrs = {"soft_label": 0}

def test_check_output(self):
self.check_output()
Expand All @@ -34,24 +34,25 @@ class TestCrossEntropyOp2(OpTest):

def setUp(self):
self.op_type = "cross_entropy"
batch_size = 10
class_num = 5
batch_size = 13
class_num = 37
X = np.random.uniform(0.1, 1.0,
[batch_size, class_num]).astype("float32")
label = np.random.uniform(0.1, 1.0,
[batch_size, class_num]).astype("float32")
label /= label.sum(axis=1, keepdims=True)
cross_entropy = (-label * np.log(X)).sum(
axis=1, keepdims=True).astype("float32")
self.inputs = {'X': X, 'Label': label}
self.outputs = {'Y': cross_entropy}
self.attrs = {'soft_label': 1}

self.inputs = {"X": X, "Label": label}
self.outputs = {"Y": cross_entropy}
self.attrs = {"soft_label": 1}

def test_check_output(self):
self.check_output()

def test_check_grad(self):
self.check_grad(['X'], 'Y')
self.check_grad(["X"], "Y", max_relative_error=0.05)


class TestCrossEntropyOp3(OpTest):
Expand All @@ -61,8 +62,8 @@ class TestCrossEntropyOp3(OpTest):

def setUp(self):
self.op_type = "cross_entropy"
batch_size = 30
class_num = 10
batch_size = 13
class_num = 37
X = np.random.uniform(0.1, 1.0,
[batch_size, class_num]).astype("float32")
label_index = np.random.randint(
Expand All @@ -74,15 +75,15 @@ def setUp(self):
dtype="float32")
cross_entropy2 = (-label * np.log(X)).sum(
axis=1, keepdims=True).astype("float32")
self.inputs = {'X': X, 'Label': label}
self.outputs = {'Y': cross_entropy}
self.attrs = {'soft_label': 1}
self.inputs = {"X": X, "Label": label}
Copy link
Collaborator

@reyoung reyoung Sep 20, 2017

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

https://www.python.org/dev/peps/pep-0008/#string-quotes

We following PEP8 as our style guide. In python, both '' and "" are fine.

Copy link
Contributor Author

@lcy-seso lcy-seso Sep 20, 2017 •
edited
Loading

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Yes, I know, only personally, I prefer to keep consistent in one file.

self.outputs = {"Y": cross_entropy}
self.attrs = {"soft_label": 1}

def test_check_output(self):
self.check_output()

def test_check_grad(self):
self.check_grad(['X'], 'Y')
self.check_grad(["X"], "Y", max_relative_error=0.05)


if __name__ == "__main__":
Expand Down