未验证 提交 de3c067a 编写于 作者: T tangwei12 提交者: GitHub

fix gpu outofrange (#29238) (#29348)

* fix gpu emb out of range

Change-Id: I5794ac73bd634d5ea069a6fbbd914274b6d6b7bf

* fix doc

Change-Id: I5a3350b2930a9ab2f52116c192b087307faf8fdf
上级 07a7cd4b
...@@ -31,16 +31,6 @@ __global__ void LookupTableV2(T *output, const T *table, const int64_t *ids, ...@@ -31,16 +31,6 @@ __global__ void LookupTableV2(T *output, const T *table, const int64_t *ids,
while (idy < K) { while (idy < K) {
int64_t id = ids[idy]; int64_t id = ids[idy];
PADDLE_ENFORCE(
id >= 0,
"Variable value (input) of OP(fluid.layers.embedding) "
"expected >= 0 and < %ld, but got %ld. Please check input value.",
N, id);
PADDLE_ENFORCE(
id < N,
"Variable value (input) of OP(fluid.layers.embedding) "
"expected >= 0 and < %ld, but got %ld. Please check input value.",
N, id);
T *out = output + idy * D; T *out = output + idy * D;
const T *tab = table + id * D; const T *tab = table + id * D;
for (int i = idx; i < D; i += BlockDimX) { for (int i = idx; i < D; i += BlockDimX) {
...@@ -66,16 +56,6 @@ __global__ void LookupTableV2Grad(T *table, const T *output, const int64_t *ids, ...@@ -66,16 +56,6 @@ __global__ void LookupTableV2Grad(T *table, const T *output, const int64_t *ids,
while (idy < K) { while (idy < K) {
int64_t id = ids[idy]; int64_t id = ids[idy];
PADDLE_ENFORCE(
id >= 0,
"Variable value (input) of OP(fluid.layers.embedding) "
"expected >= 0 and < %ld, but got %ld. Please check input value.",
N, id);
PADDLE_ENFORCE(
id < N,
"Variable value (input) of OP(fluid.layers.embedding) "
"expected >= 0 and < %ld, but got %ld. Please check input value.",
N, id);
const T *out = output + idy * D; const T *out = output + idy * D;
T *tab = table + id * D; T *tab = table + id * D;
for (int i = idx; i < D; i += BlockDimX) { for (int i = idx; i < D; i += BlockDimX) {
...@@ -127,6 +107,21 @@ class LookupTableV2CUDAKernel : public framework::OpKernel<T> { ...@@ -127,6 +107,21 @@ class LookupTableV2CUDAKernel : public framework::OpKernel<T> {
ids_p = ids_t->data<int64_t>(); ids_p = ids_t->data<int64_t>();
} }
for (int64_t i = 0; i < K; ++i) {
PADDLE_ENFORCE_GE(
ids[i], 0,
platform::errors::InvalidArgument(
"Variable value (input) of OP(paddle.nn.embedding) "
"expected >= 0 and < %ld, but got %ld. Please check input value.",
N, ids[i]));
PADDLE_ENFORCE_LT(
ids[i], N,
platform::errors::InvalidArgument(
"Variable value (input) of OP(paddle.nn.embedding) "
"expected >= 0 and < %ld, but got %ld. Please check input value.",
N, ids[i]));
}
auto *table = table_t->data<T>(); auto *table = table_t->data<T>();
auto *output = output_t->mutable_data<T>(context.GetPlace()); auto *output = output_t->mutable_data<T>(context.GetPlace());
......
...@@ -197,10 +197,7 @@ def embedding(input, ...@@ -197,10 +197,7 @@ def embedding(input,
indicates the size of the dictionary of embeddings and the size of each embedding vector respectively. indicates the size of the dictionary of embeddings and the size of each embedding vector respectively.
is_sparse(bool): The flag indicating whether to use sparse update. This parameter only is_sparse(bool): The flag indicating whether to use sparse update. This parameter only
affects the performance of the backwards gradient update. It is recommended to set affects the performance of the backwards gradient update. It is recommended to set
True because sparse update is faster. But some optimizer does not support sparse update, True because sparse update is faster. But some optimizer does not support sparse update
such as :ref:`api_fluid_optimizer_AdadeltaOptimizer` , :ref:`api_fluid_optimizer_AdamaxOptimizer` ,
:ref:`api_fluid_optimizer_DecayedAdagradOptimizer` , :ref:`api_fluid_optimizer_FtrlOptimizer` ,
:ref:`api_fluid_optimizer_LambOptimizer` and :ref:`api_fluid_optimizer_LarsMomentumOptimizer` .
In these case, is_sparse must be False. Default: False. In these case, is_sparse must be False. Default: False.
is_distributed(bool): Whether to store the embedding matrix in a distributed manner. Only used is_distributed(bool): Whether to store the embedding matrix in a distributed manner. Only used
in multi-machine distributed CPU training. Default: False. in multi-machine distributed CPU training. Default: False.
...@@ -210,11 +207,10 @@ def embedding(input, ...@@ -210,11 +207,10 @@ def embedding(input,
encounters :math:`padding\_idx` in id. And the padding data will not be updated while training. encounters :math:`padding\_idx` in id. And the padding data will not be updated while training.
If set None, it makes no effect to output. Default: None. If set None, it makes no effect to output. Default: None.
param_attr(ParamAttr): To specify the weight parameter property. Default: None, which means the param_attr(ParamAttr): To specify the weight parameter property. Default: None, which means the
default weight parameter property is used. See usage for details in :ref:`api_fluid_ParamAttr` . In addition, default weight parameter property is used. In addition,
user-defined or pre-trained word vectors can be loaded with the :attr:`param_attr` parameter. user-defined or pre-trained word vectors can be loaded with the :attr:`param_attr` parameter.
The local word vector needs to be transformed into numpy format, and the shape of local word The local word vector needs to be transformed into numpy format, and the shape of local word
vector should be consistent with :attr:`size` . Then :ref:`api_fluid_initializer_NumpyArrayInitializer` vector should be consistent with :attr:`size` .
is used to load custom or pre-trained word vectors. See code example 2 for details.
dtype(str|core.VarDesc.VarType): It refers to the data type of output Tensor. dtype(str|core.VarDesc.VarType): It refers to the data type of output Tensor.
It must be float32 or float64. Default: float32. It must be float32 or float64. Default: float32.
...@@ -267,9 +263,7 @@ def embedding(input, ...@@ -267,9 +263,7 @@ def embedding(input,
import paddle import paddle
import numpy as np import numpy as np
paddle.disable_static()
x_data = np.arange(3, 6).reshape((3, 1)).astype(np.int64) x_data = np.arange(3, 6).reshape((3, 1)).astype(np.int64)
# x is a Tensor. # x is a Tensor.
......
...@@ -168,28 +168,25 @@ def embedding(x, weight, padding_idx=None, sparse=False, name=None): ...@@ -168,28 +168,25 @@ def embedding(x, weight, padding_idx=None, sparse=False, name=None):
.. code-block:: python .. code-block:: python
import numpy as np
import paddle import paddle
import paddle.nn as nn import paddle.nn as nn
weight = prog.global_block().create_parameter( x0 = np.arange(3, 6).reshape((3, 1)).astype(np.int64)
attr=self._param_attr, w0 = np.full(shape=(10, 3), fill_value=2).astype(np.float32)
shape=param_shape,
dtype=self._dtype,
default_initializer=Constant(1.0))
prog = paddle.static.Program() # x.data = [[3], [4], [5]]
# x.shape = [3, 1]
x = paddle.to_tensor(x0, stop_gradient=False)
weight = prog.global_block().create_parameter( # w.data = [[2. 2. 2.] ... [2. 2. 2.]]
(128, 100), dtype="float32", default_initializer=Constant(1.0)) # w.shape = [10, 3]
w = paddle.to_tensor(w0, stop_gradient=False)
label = paddle.static.data( # emb.data = [[[2., 2., 2.]], [[2., 2., 2.]], [[2., 2., 2.]]]
name="label", # emb.shape = [3, 1, 3]
shape=[4], emb = nn.functional.embedding(
append_batch_size=False, x=x, weight=w, sparse=True, name="embedding")
dtype="int64")
emb = nn.embedding(
x=label, weight=weight, sparse=True, name="embedding")
""" """
padding_idx = -1 if padding_idx is None else padding_idx if padding_idx >= 0 else ( padding_idx = -1 if padding_idx is None else padding_idx if padding_idx >= 0 else (
......
...@@ -1216,7 +1216,7 @@ class Embedding(layers.Layer): ...@@ -1216,7 +1216,7 @@ class Embedding(layers.Layer):
x_data = np.arange(3, 6).reshape((3, 1)).astype(np.int64) x_data = np.arange(3, 6).reshape((3, 1)).astype(np.int64)
y_data = np.arange(6, 12).reshape((3, 2)).astype(np.float32) y_data = np.arange(6, 12).reshape((3, 2)).astype(np.float32)
paddle.disable_static(paddle.CPUPlace())
x = paddle.to_tensor(x_data, stop_gradient=False) x = paddle.to_tensor(x_data, stop_gradient=False)
y = paddle.to_tensor(y_data, stop_gradient=False) y = paddle.to_tensor(y_data, stop_gradient=False)
......
Markdown is supported
0% .
You are about to add 0 people to the discussion. Proceed with caution.
先完成此消息的编辑!
想要评论请 注册