Skip to content
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.

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

[Rocm] fix test of random_crop_op & logsumexp #32824

Merged
merged 7 commits into from
May 19, 2021
Merged
Show file tree
Hide file tree
Changes from 5 commits
Commits
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
17 changes: 17 additions & 0 deletions paddle/fluid/operators/random_crop_op.h
Original file line number Diff line number Diff line change
Expand Up @@ -59,6 +59,7 @@ HOSTDEVICE inline void StridedMemcpy(const T* x, const size_t* x_dims, T* out,
size_t offset_i = offsets[i];

if (i == rank - 1) {
#ifndef PADDLE_WITH_HIP
PADDLE_ENFORCE(x_stride == 1,
"When i:%d == rank:%d - 1, x_stride of random_crop_op "
"expected to be 1, but got %ld. Please check input "
Expand All @@ -69,6 +70,22 @@ HOSTDEVICE inline void StridedMemcpy(const T* x, const size_t* x_dims, T* out,
"expected to be 1, but got %ld. Please check input "
"value.",
i, rank, out_stride);
#else
if (x_stride != 1) {
printf(
"When i:%d == rank:%d - 1, x_stride of random_crop_op expected to be "
"1, but got %ld. Please check input value.\n",
i, rank, x_stride);
abort();
}
if (out_stride != 1) {
printf(
"When i:%d == rank:%d - 1, out_stride of random_crop_op expected to "
"be 1, but got %ld. Please check input value.\n",
i, rank, out_stride);
abort();
}
#endif
Copy link
Contributor

@qili93 qili93 May 13, 2021

Choose a reason for hiding this comment

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

看了下你的分析文档,但当前这个修复方式不能满足OP要,我们OP修复的原则是

  • 不可以因为这两段代码会挂就不检查x_stride和out_stride,而是需要找到能工作的代码替换这两段代码,或者修复PADDLE_ENFORCE的逻辑使其能够工作

我们修复OP的目的不是为了单测可以通过,而是保证功能的正确性,当前的diff只是简单跳过了x_stride和out_stride的检查,不符合我们OP修复的目的。

Copy link
Contributor Author

Choose a reason for hiding this comment

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

好的,已修改。

Copy link
Contributor

Choose a reason for hiding this comment

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

这个直接printf不符合paddle msg的标准,请参考这个文档写关于paddle的error msg
https://github.com/PaddlePaddle/Paddle/wiki/Paddle-Error-Message-Writing-Specification

Copy link
Contributor Author

Choose a reason for hiding this comment

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

已删除

x += offset_i;
for (size_t j = 0; j < out_dim_i; ++j) {
*out++ = *x++;
Expand Down
27 changes: 26 additions & 1 deletion python/paddle/fluid/tests/unittests/test_logsumexp.py
Original file line number Diff line number Diff line change
Expand Up @@ -50,15 +50,30 @@ def setUp(self):
'keepdim': self.keepdim,
'reduce_all': self.reduce_all
}
self.user_defined_grads = None
self.user_defined_grad_outputs = None
self.set_attrs_addition()

def set_attrs(self):
pass

def set_attrs_addition(self):
pass

def test_check_output(self):
self.check_output()

def test_check_grad(self):
self.check_grad(['X'], ['Out'])
self.check_grad(
['X'], ['Out'],
user_defined_grads=self.user_defined_grads,
user_defined_grad_outputs=self.user_defined_grad_outputs)

def calc_grad(self):
dy = np.ones(1, dtype=self.dtype)
x = self.inputs['X']
y = self.outputs['Out']
return dy * np.exp(x - y)
Copy link
Contributor

Choose a reason for hiding this comment

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

这里能解释一下为什么get_numeric_gradient得到的结果在ROCm上无法对齐,需要定义user_defined_grads呢?CUDA上为什么可以直接使用get_numeric_gradient?

Copy link
Contributor Author

Choose a reason for hiding this comment

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

logsumexp这个op本身运算出来的analytic_grads是没问题的,有问题的是numeric_grads(cuda下该计算没问题),所以干脆自己计算numeric_grads



class TestLogsumexp_shape(TestLogsumexp):
Expand All @@ -75,6 +90,11 @@ class TestLogsumexp_axis_all(TestLogsumexp):
def set_attrs(self):
self.axis = [0, 1, 2, 3]

def set_attrs_addition(self):
if paddle.fluid.core.is_compiled_with_rocm():
self.user_defined_grads = [self.calc_grad()]
self.user_defined_grad_outputs = [np.ones(1, dtype=self.dtype)]


class TestLogsumexp_keepdim(TestLogsumexp):
def set_attrs(self):
Expand All @@ -85,6 +105,11 @@ class TestLogsumexp_reduce_all(TestLogsumexp):
def set_attrs(self):
self.reduce_all = True

def set_attrs_addition(self):
if paddle.fluid.core.is_compiled_with_rocm():
self.user_defined_grads = [self.calc_grad()]
self.user_defined_grad_outputs = [np.ones(1, dtype=self.dtype)]


class TestLogsumexpError(unittest.TestCase):
def test_errors(self):
Expand Down
16 changes: 13 additions & 3 deletions python/paddle/fluid/tests/unittests/test_simple_rnn_op.py
Original file line number Diff line number Diff line change
Expand Up @@ -44,8 +44,11 @@ def get_weight_names(self):

def setUp(self):
self.op_type = "rnn"
self.dtype = np.float64
self.sequence_length = np.array([12, 11, 10, 9, 8], dtype=np.int32)
self.dtype = np.float32 if fluid.core.is_compiled_with_rocm(
) else np.float64
self.sequence_length = None if fluid.core.is_compiled_with_rocm(
) else np.array(
[12, 11, 10, 9, 8], dtype=np.int32)
Copy link
Contributor

Choose a reason for hiding this comment

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

MIOPEN不支持sequence_length的话,SimpleRNN的API也同样需要修改,对于sequence_length非空的输入需要抛出正确的error msg.

Copy link
Contributor Author

Choose a reason for hiding this comment

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

simpleRNN的修改以及该单测的修改从本pr中删除,将和其他rnn单测修改统一放到其他的pr中。

self.num_layers = 1
self.is_bidirec = False
self.is_test = False
Expand Down Expand Up @@ -76,12 +79,19 @@ def setUp(self):
time_major=True,
direction=direction,
dropout=self.dropout,
nonlinearity=self.mode)
nonlinearity=self.mode,
dtype=self.dtype)

flat_w = get_params_for_net(rnn1)

output, last_hidden = rnn1(input, sequence_length=self.sequence_length)

if fluid.core.is_compiled_with_rocm():
self._get_places = lambda: [fluid.core.CUDAPlace(0)]
if self.is_bidirec:
for i in range(0, len(flat_w), 4):
flat_w[i + 1], flat_w[i + 2] = flat_w[i + 2], flat_w[i + 1]

Copy link
Contributor

Choose a reason for hiding this comment

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

这里能解释一下为什么需要对flat_w的数据做交换吗?是miopenSetRNNDescriptor_V2 里面的 miopenRNNbidirection 和 cudnnSetRNNDescriptor_v6 的 CUDNN_BIDIRECTIONAL 的行为不一致?

Copy link
Contributor Author

Choose a reason for hiding this comment

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

miopen相关api的实现中,layout并没有和cuda相关pai对齐,需要手动调节。

init_h = np.zeros((self.num_layers * self.direction_num, batch_size,
hidden_size)).astype(self.dtype)

Expand Down