Skip to content

Commit 5591292

Browse files
author
Haonan
committed
modifications according to comments
1 parent 728defb commit 5591292

File tree

10 files changed

+87
-99
lines changed

10 files changed

+87
-99
lines changed

paddle/cuda/src/hl_cuda_matrix.cu

Lines changed: 4 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -346,6 +346,7 @@ void hl_matrix_multi_binary_cross_entropy(real* output,
346346
CHECK_NOTNULL(output);
347347
CHECK_NOTNULL(entropy);
348348
CHECK_NOTNULL(csr_mat);
349+
CHECK_EQ(csr_mat->format, HL_SPARSE_CSR);
349350
int n_threads = 1024;
350351
int blocks = (dimM + n_threads - 1) / n_threads;
351352
dim3 threads(n_threads);
@@ -385,6 +386,7 @@ void hl_matrix_multi_binary_cross_entropy_bp(real* output,
385386
CHECK_NOTNULL(output);
386387
CHECK_NOTNULL(grad);
387388
CHECK_NOTNULL(csr_mat);
389+
CHECK_EQ(csr_mat->format, HL_SPARSE_CSR);
388390
int n_threads = 1024;
389391
int blocks = (dimM + n_threads - 1) / n_threads;
390392
dim3 threads(n_threads);
@@ -763,7 +765,7 @@ __global__ void KeMatrixAddSharedBias(real* A,
763765
int dim = N / channel;
764766
if (index < M * N) {
765767
int i = index % N;
766-
i = i / dim;
768+
i = i / dim;
767769
A[index] += scale * B[i];
768770
}
769771
}
@@ -791,7 +793,7 @@ __global__ void KeMatrixCollectSharedBias(real *B,
791793
const int dim,
792794
const int limit,
793795
real scale) {
794-
if (dim < limit) {
796+
if (dim < limit) {
795797
int index = blockIdx.x * blockDim.x + threadIdx.x;
796798
if (index < channel) {
797799
real sum = 0.0;

paddle/gserver/layers/CostLayer.cpp

Lines changed: 2 additions & 8 deletions
Original file line numberDiff line numberDiff line change
@@ -465,10 +465,7 @@ void MultiBinaryLabelCrossEntropy::forwardImp(Matrix& output, Argument& label,
465465
MatrixPtr value = nullptr;
466466
if (label.ids) {
467467
CHECK(!label.value);
468-
value = Matrix::createSparseMatrix(
469-
label.ids->getSize(), output.getWidth(), label.ids->getSize(),
470-
NO_VALUE, SPARSE_CSR, false, useGpu_);
471-
label.idsToSparseMatrix(value);
468+
value = label.ids->toOneHotSparseMatrix(output.getWidth(), useGpu_);
472469
} else {
473470
CHECK(label.value);
474471
value = label.value;
@@ -491,10 +488,7 @@ void MultiBinaryLabelCrossEntropy::backwardImp(
491488
MatrixPtr value = nullptr;
492489
if (label.ids) {
493490
CHECK(!value);
494-
value = Matrix::createSparseMatrix(
495-
label.ids->getSize(), output.getWidth(), label.ids->getSize(),
496-
NO_VALUE, SPARSE_CSR, false, useGpu_);
497-
label.idsToSparseMatrix(value);
491+
value = label.ids->toOneHotSparseMatrix(output.getWidth(), useGpu_);
498492
} else {
499493
CHECK(label.value);
500494
value = label.value;

paddle/gserver/tests/test_LayerGrad.cpp

Lines changed: 17 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -528,7 +528,7 @@ TEST(Layer, multi_cross) {
528528
}
529529
}
530530

531-
TEST(Layer, multi_binary_label) {
531+
TEST(Layer, multi_binary_label_sparse_mat) {
532532
TestConfig config;
533533
config.layerConfig.set_type("multi_binary_label_cross_entropy");
534534
config.biasSize = 0;
@@ -544,6 +544,22 @@ TEST(Layer, multi_binary_label) {
544544
}
545545
}
546546

547+
TEST(layer, multi_binary_label_id) {
548+
TestConfig config;
549+
config.layerConfig.set_type("multi_binary_label_cross_entropy");
550+
config.biasSize = 0;
551+
552+
config.inputDefs.push_back({INPUT_DATA, "layer_0", 50, 0});
553+
config.inputDefs.push_back({INPUT_LABEL, "layer_1", 10, 0});
554+
config.layerConfig.add_inputs();
555+
config.layerConfig.add_inputs();
556+
557+
for (auto useGpu : {false, true}) {
558+
testLayerGrad(config, "multi_binary_label_cross_entropy", 100,
559+
/* trans */ false, useGpu);
560+
}
561+
}
562+
547563
TEST(Layer, multi_cross_with_selfnorm) {
548564
TestConfig config;
549565
config.layerConfig.set_type("multi_class_cross_entropy_with_selfnorm");

paddle/math/CpuSparseMatrix.cpp

Lines changed: 0 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -409,9 +409,6 @@ void CpuSparseMatrix::setRow(size_t row, size_t colNum,
409409
if (format_ == SPARSE_CSR) {
410410
CHECK_LT(row, height_);
411411
CHECK(NULL != cols);
412-
for (size_t i = row; i < height_; i++) {
413-
CHECK_EQ(rows_[i + 1], rows_[i]);
414-
}
415412
if (0 == row) {
416413
rows_[row] = 0;
417414
}

paddle/math/Matrix.cpp

Lines changed: 21 additions & 21 deletions
Original file line numberDiff line numberDiff line change
@@ -1269,37 +1269,37 @@ void GpuMatrix::bilinearBackward(const Matrix& out,
12691269
}
12701270

12711271
void GpuMatrix::multiBinaryLabelCrossEntropy(Matrix& output, Matrix& label) {
1272-
GpuMatrix* output_ptr = dynamic_cast<GpuMatrix*>(&output);
1273-
auto label_ptr = dynamic_cast<GpuSparseMatrix*>(&label);
1274-
1275-
CHECK(output_ptr && label_ptr) << "Invalid argument pointer";
1276-
CHECK(label_ptr->format_ == SPARSE_CSR) << "Matrix format not supported";
1277-
CHECK(height_ == output_ptr->height_ && width_ == 1
1278-
&& output_ptr->width_ == label_ptr->getWidth()
1279-
&& output_ptr->height_ == label_ptr->getHeight())
1272+
GpuMatrix* outputPtr = dynamic_cast<GpuMatrix*>(&output);
1273+
auto labelPtr = dynamic_cast<GpuSparseMatrix*>(&label);
1274+
1275+
CHECK(outputPtr && labelPtr) << "Invalid argument pointer";
1276+
CHECK(labelPtr->format_ == SPARSE_CSR) << "Matrix format not supported";
1277+
CHECK(height_ == outputPtr->height_ && width_ == 1
1278+
&& outputPtr->width_ == labelPtr->getWidth()
1279+
&& outputPtr->height_ == labelPtr->getHeight())
12801280
<< "Matrix dimensions are not equal";
12811281

1282-
real* output_d = output_ptr->data_;
1282+
real* output_d = outputPtr->data_;
12831283
real* entropy_d = data_;
1284-
hl_sparse_matrix_s mat_d = label_ptr->sMatrix_.get();
1284+
hl_sparse_matrix_s mat_d = labelPtr->sMatrix_.get();
12851285
hl_matrix_multi_binary_cross_entropy(
1286-
output_d, entropy_d, mat_d, height_, output_ptr->width_);
1286+
output_d, entropy_d, mat_d, height_, outputPtr->width_);
12871287
}
12881288

12891289
void GpuMatrix::multiBinaryLabelCrossEntropyBp(Matrix &output, Matrix &label) {
1290-
GpuMatrix* output_ptr = dynamic_cast<GpuMatrix*>(&output);
1291-
auto label_ptr = dynamic_cast<GpuSparseMatrix*>(&label);
1292-
1293-
CHECK(output_ptr && label_ptr) << "Invalid argument pointer";
1294-
CHECK(label_ptr->format_ == SPARSE_CSR) << "Matrix format not supported";
1295-
CHECK(height_ == output_ptr->height_ && width_ == output_ptr->width_
1296-
&& output_ptr->width_ == label_ptr->getWidth()
1297-
&& output_ptr->height_ == label_ptr->getHeight())
1290+
GpuMatrix* outputPtr = dynamic_cast<GpuMatrix*>(&output);
1291+
auto labelPtr = dynamic_cast<GpuSparseMatrix*>(&label);
1292+
1293+
CHECK(outputPtr && labelPtr) << "Invalid argument pointer";
1294+
CHECK(labelPtr->format_ == SPARSE_CSR) << "Matrix format not supported";
1295+
CHECK(height_ == outputPtr->height_ && width_ == outputPtr->width_
1296+
&& outputPtr->width_ == labelPtr->getWidth()
1297+
&& outputPtr->height_ == labelPtr->getHeight())
12981298
<< "Matrix dimensions are not equal";
12991299

1300-
real* output_d = output_ptr->data_;
1300+
real* output_d = outputPtr->data_;
13011301
real* grad_d = data_;
1302-
hl_sparse_matrix_s mat_d = label_ptr->sMatrix_.get();
1302+
hl_sparse_matrix_s mat_d = labelPtr->sMatrix_.get();
13031303
hl_matrix_multi_binary_cross_entropy_bp(
13041304
output_d, grad_d, mat_d, height_, width_);
13051305
}

paddle/math/Vector.cpp

Lines changed: 26 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -21,6 +21,7 @@ limitations under the License. */
2121
#include "paddle/utils/ThreadLocal.h"
2222
#include "paddle/utils/Thread.h"
2323
#include "paddle/utils/Flags.h"
24+
#include "Matrix.h"
2425
#include "hl_gpu.h"
2526
#include "hl_table_apply.h"
2627

@@ -73,6 +74,31 @@ std::shared_ptr<VectorT<T>> VectorT<T>::create(size_t size,
7374
}
7475
}
7576

77+
template <>
78+
MatrixPtr VectorT<real>::toOneHotSparseMatrix(size_t idRange, bool useGpu) {
79+
LOG(FATAL) << "Wrong for real vector";
80+
return nullptr;
81+
}
82+
83+
template <>
84+
MatrixPtr VectorT<int>::toOneHotSparseMatrix(size_t idRange, bool useGpu) {
85+
int height = getSize();
86+
int width = idRange;
87+
MatrixPtr mat = Matrix::createSparseMatrix(
88+
height, idRange, height, NO_VALUE, SPARSE_CSR, false, useGpu);
89+
90+
CpuIVector cpuIds(height);
91+
cpuIds.copyFrom(*this);
92+
int *idData = cpuIds.getData();
93+
94+
for (int i = 0; i < height; i ++) {
95+
const unsigned int id = idData[i];
96+
CHECK_LT(id, width);
97+
mat->setRow(i, 1, &id, nullptr);
98+
}
99+
return mat;
100+
}
101+
76102
template <class T>
77103
GpuVectorT<T>::GpuVectorT(size_t size)
78104
: VectorT<T>(size, std::make_shared<GpuMemoryHandle>(sizeof(T) * size),

paddle/math/Vector.h

Lines changed: 8 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -37,6 +37,8 @@ class BaseVector;
3737

3838
class SyncThreadPool;
3939

40+
class Matrix;
41+
4042
template<class T>
4143
class BaseVector : public BaseMatrixT<T> {
4244
public:
@@ -155,6 +157,12 @@ class VectorT : public BaseVector<T> {
155157
subVecFrom(src, interval.first, interval.second - interval.first);
156158
}
157159

160+
/**
161+
* convert the vector to a sparse one_hot matrix of width idRange
162+
* only applies to IVector
163+
*/
164+
std::shared_ptr<Matrix> toOneHotSparseMatrix(size_t idRange, bool useGpu);
165+
158166
/**
159167
* This function will crash if the size of src and dest is different.
160168
*/

paddle/math/tests/test_matrixCompare.cpp

Lines changed: 9 additions & 20 deletions
Original file line numberDiff line numberDiff line change
@@ -2232,26 +2232,15 @@ void testMultiBinaryLabelCrossEntropy(int numSamples, int dim) {
22322232
MatrixPtr cpuGrad = std::make_shared<CpuMatrix>(numSamples, dim);
22332233
MatrixPtr gpuGrad = std::make_shared<GpuMatrix>(numSamples, dim);
22342234

2235-
auto cpuRows = IVector::create(numSamples + 1, false);
2236-
auto cpuCols = IVector::create(numSamples, false);
2237-
auto gpuRows = IVector::create(numSamples + 1, true);
2238-
auto gpuCols = IVector::create(numSamples, true);
2239-
cpuRows->setElement(0, 0);
2240-
gpuRows->setElement(0, 0);
2241-
for (int i = 0; i < numSamples; i ++) {
2242-
int id = rand() % dim; // NOLINT
2243-
cpuRows->setElement(i + 1, i + 1);
2244-
gpuRows->setElement(i + 1, i + 1);
2245-
cpuCols->setElement(i, id);
2246-
gpuCols->setElement(i, id);
2247-
}
2248-
22492235
MatrixPtr cpuLabel = std::make_shared<CpuSparseMatrix>
2250-
(nullptr, cpuRows->getData(), cpuCols->getData(),
2251-
numSamples, dim, numSamples, NO_VALUE, SPARSE_CSR, false);
2236+
(numSamples, dim, numSamples, NO_VALUE, SPARSE_CSR, false);
22522237
MatrixPtr gpuLabel = std::make_shared<GpuSparseMatrix>
2253-
(nullptr, gpuRows->getData(), gpuCols->getData(),
2254-
numSamples, dim, numSamples, NO_VALUE, SPARSE_CSR, false);
2238+
(numSamples, dim, numSamples, NO_VALUE, SPARSE_CSR, false);
2239+
for (int i = 0; i < numSamples; i ++) {
2240+
const unsigned int id = rand() % dim; // NOLINT
2241+
cpuLabel->setRow(i, 1, &id, nullptr);
2242+
gpuLabel->setRow(i, 1, &id, nullptr);
2243+
}
22552244

22562245
output->randomizeUniform();
22572246
cpuOutput->zeroMem();
@@ -2278,8 +2267,8 @@ void testMultiBinaryLabelCrossEntropy(int numSamples, int dim) {
22782267
}
22792268

22802269
TEST(Matrix, multiBinaryCrossEntropy) {
2281-
for (auto numSamples : {1, 100, 500}) {
2282-
for (auto dim : {1000, 10000, 100000}) {
2270+
for (auto numSamples : {100, 1000, 10000}) {
2271+
for (auto dim : {100, 1000, 10000}) {
22832272
VLOG(3) << " numSamples=" << numSamples << " dim=" << dim;
22842273
testMultiBinaryLabelCrossEntropy(numSamples, dim);
22852274
}

paddle/parameter/Argument.cpp

Lines changed: 0 additions & 38 deletions
Original file line numberDiff line numberDiff line change
@@ -572,42 +572,4 @@ void Argument::subArgFrom(const Argument& input, size_t offset, size_t height,
572572
}
573573
}
574574

575-
void Argument::idsToSparseMatrix(MatrixPtr sparse_mat) {
576-
int height = ids->getSize();
577-
int width = sparse_mat->getWidth();
578-
579-
CpuIVector cpu_ids(height);
580-
cpu_ids.copyFrom(*ids);
581-
int *id_data = cpu_ids.getData();
582-
583-
int *rows = nullptr;
584-
int *cols = nullptr;
585-
if (sparse_mat->useGpu()) {
586-
auto gpu_sparse_mat =
587-
dynamic_cast<GpuSparseMatrix*>(sparse_mat.get());
588-
rows = gpu_sparse_mat->rows_;
589-
cols = gpu_sparse_mat->cols_;
590-
} else {
591-
rows = sparse_mat->getRows();
592-
cols = sparse_mat->getCols();
593-
}
594-
595-
rows[0] = 0;
596-
for (int i = 0; i < height; i ++) {
597-
int id = id_data[i];
598-
CHECK_LT(id, width);
599-
rows[i + 1] = i + 1;
600-
cols[i] = id;
601-
}
602-
603-
if (sparse_mat->useGpu()) {
604-
auto gpu_sparse_mat =
605-
dynamic_cast<GpuSparseMatrix*>(sparse_mat.get());
606-
hl_memcpy_csr_matrix(gpu_sparse_mat->sMatrix_.get(),
607-
nullptr, rows, cols,
608-
HPPL_STREAM_DEFAULT);
609-
hl_stream_synchronize(HPPL_STREAM_DEFAULT);
610-
}
611-
}
612-
613575
} // namespace paddle

paddle/parameter/Argument.h

Lines changed: 0 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -286,12 +286,6 @@ struct Argument {
286286
sequence has sub-sequence degrades to a sequence.
287287
*/
288288
void degradeSequence(const Argument& input, bool useGpu);
289-
290-
/*
291-
@brief convert the ids vector to value as a sparse matrix
292-
@param[out] the output sparse_mat (already allocated)
293-
*/
294-
void idsToSparseMatrix(MatrixPtr sparse_mat);
295289
};
296290

297291
} // namespace paddle

0 commit comments

Comments
 (0)