diff --git a/paddle/cuda/include/hl_sequence.h b/paddle/cuda/include/hl_sequence.h index 9bcd25b0623e569052e08c0befc8e09f937fa4bd..9f9d8f972e3a4c62e5caedcf85054be5681b96c1 100644 --- a/paddle/cuda/include/hl_sequence.h +++ b/paddle/cuda/include/hl_sequence.h @@ -48,78 +48,6 @@ extern void hl_max_sequence_forward(real* input, extern void hl_max_sequence_backward( real* outputGrad, int* index, real* inputGrad, int numSequences, int dim); -/** - * @brief Context projection forward. - * - * @param[in] input input sequence. - * @param[in] sequence sequence index. - * @param[in] weightData padding data. - * @param[out] output output sequence. - * @param[in] numSequences number of sequences. - * @param[in] inputDim input sequence dimension. - * @param[in] contextLength context length. - * @param[in] contextStart context start. - * @param[in] beginPad number of extra timesteps added at the - * beginning. - * @param[in] isPadding trainable padding. - * - */ -extern void hl_context_projection_forward(real* input, - const int* sequence, - real* weightData, - real* output, - int numSequences, - int inputDim, - int contextLength, - int contextStart, - int beginPad, - bool isPadding); - -/** - * @brief Context projection backward data. - * - * @param[in] outputGrad output gradient. - * @param[in] sequence sequence index. - * @param[out] inputGrad input gradient. - * @param[in] numSequences number of sequences. - * @param[in] inputDim input sequence dimension. - * @param[in] contextLength context length. - * @param[in] contextStart context start. - * - */ -extern void hl_context_projection_backward_data(real* outputGrad, - const int* sequence, - real* inputGrad, - int numSequences, - int inputDim, - int contextLength, - int contextStart); - -/** - * @brief Context projection backward weight. - * - * @param[in] outputGrad output gradient. - * @param[in] sequence sequence index. - * @param[out] weightGrad weight gradient. - * @param[in] numSequences number of sequences. - * @param[in] weightDim input sequence dimension. - * @param[in] totalPad number of extra timesteps. - * @param[in] contextLength context length. - * @param[in] contextStart context start. - * @param[in] beginPad number of extra timesteps added at the - * beginning. - * - */ -extern void hl_context_projection_backward_weight(real* outputGrad, - const int* sequence, - real* weightGrad, - int numSequences, - int weightDim, - int totalPad, - int contextLength, - int contextStart, - int beginPad); - /** * @brief Memory copy from sequence to batch. * diff --git a/paddle/cuda/include/stub/hl_sequence_stub.h b/paddle/cuda/include/stub/hl_sequence_stub.h index d6b07556f8958a62bd47f0b47b75bbebafeb58d3..05e51bce9e1df6fc6ef1cad891b44a9172da185d 100644 --- a/paddle/cuda/include/stub/hl_sequence_stub.h +++ b/paddle/cuda/include/stub/hl_sequence_stub.h @@ -27,35 +27,6 @@ inline void hl_max_sequence_forward(real* input, inline void hl_max_sequence_backward( real* outputGrad, int* index, real* inputGrad, int numSequences, int dim) {} -inline void hl_context_projection_forward(real* input, - const int* sequence, - real* weightData, - real* output, - int numSequences, - int inputDim, - int contextLength, - int contextStart, - int beginPad, - bool isPadding) {} - -inline void hl_context_projection_backward_data(real* outputGrad, - const int* sequence, - real* inputGrad, - int numSequences, - int inputDim, - int contextLength, - int contextStart) {} - -inline void hl_context_projection_backward_weight(real* outputGrad, - const int* sequence, - real* weightGrad, - int numSequences, - int weightDim, - int totalPad, - int contextLength, - int contextStart, - int beginPad) {} - inline void hl_sequence2batch_copy(real* batch, real* sequence, const int* batchIndex, diff --git a/paddle/cuda/src/hl_cuda_sequence.cu b/paddle/cuda/src/hl_cuda_sequence.cu index 4e33ac443c1f78b7fa50a15784875cbadfcf7497..ba823de2720336851bf9c49d8162360af93e8601 100644 --- a/paddle/cuda/src/hl_cuda_sequence.cu +++ b/paddle/cuda/src/hl_cuda_sequence.cu @@ -90,258 +90,6 @@ void hl_max_sequence_backward(real* outputGrad, CHECK_SYNC("hl_max_sequence_backward failed"); } -template -__global__ void KeContextProjectionForward(real* input, - const int* sequence, - real* weightData, - real* output, - int inputDim, - int contextLength, - int contextStart, - int beginPad) { - int idx = threadIdx.x; - int blockSize = blockDim.x; - int sequenceId = blockIdx.x; - int seqStart = sequence[sequenceId]; - int seqEnd = sequence[sequenceId+1]; - real value = 0; - - int instances = seqEnd - seqStart + contextLength - 1; - output += seqStart * inputDim * contextLength; - input += seqStart * inputDim; - for (int k = 0; k <= inputDim / blockSize; k++) { - if (idx < inputDim) { - for (int i = 0; i < instances; i++) { - // i + contextStart; - if ((i + contextStart) < 0) { - if (padding) { - value = weightData[i * inputDim + idx]; - } else { - continue; - } - } else if ((i + contextStart) >= (seqEnd - seqStart)) { - if (padding) { - value = - weightData[(beginPad + i + contextStart - (seqEnd - seqStart)) * - inputDim + idx]; - } else { - continue; - } - } else { - value = input[(i + contextStart) * inputDim + idx]; - } - - int outx = (i - contextLength) < 0 ? i : (contextLength - 1); - int outy = (i - contextLength) < 0 ? 0 : (i - (contextLength - 1)); - real* output_r = - output + outy * inputDim * contextLength + outx * inputDim; - for (int j = outy; j < seqEnd - seqStart; j++) { - output_r[idx] += value; - if (j - outy == outx) break; - output_r += (contextLength - 1) * inputDim; - } - } - } - idx += blockSize; - } -} - -void hl_context_projection_forward(real* input, - const int* sequence, - real* weightData, - real* output, - int numSequences, - int inputDim, - int contextLength, - int contextStart, - int beginPad, - bool isPadding) { - CHECK_NOTNULL(input); - CHECK_NOTNULL(sequence); - CHECK_NOTNULL(output); - CHECK(!isPadding || weightData); - - int blockSize = 128; - int blocksX = numSequences; - int blocksY = 1; - dim3 threads(blockSize, 1); - dim3 grid(blocksX, blocksY); - - if (isPadding) { - KeContextProjectionForward<<< grid, threads, 0, STREAM_DEFAULT >>> - (input, sequence, weightData, output, inputDim, - contextLength, contextStart, beginPad); - } else { - KeContextProjectionForward<<< grid, threads, 0, STREAM_DEFAULT >>> - (input, sequence, weightData, output, inputDim, - contextLength, contextStart, beginPad); - } - CHECK_SYNC("hl_context_projection_forward failed"); -} - -__global__ void KeContextProjectionBackwardData(real* outputGrad, - const int* sequence, - real* inputGrad, - int inputDim, - int contextLength, - int contextStart) { - int idx = threadIdx.x; - int blockSize = blockDim.x; - int sequenceId = blockIdx.x; - int seqStart = sequence[sequenceId]; - int seqEnd = sequence[sequenceId+1]; - real value = 0; - - int instances = seqEnd - seqStart + contextLength - 1; - outputGrad += seqStart * inputDim * contextLength; - inputGrad += seqStart * inputDim; - for (int k = 0; k <= inputDim / blockSize; k++) { - if (idx < inputDim) { - for (int i = 0; i < instances; i++) { - if ((i + contextStart) < 0) { - continue; - } else if ((i + contextStart) >= (seqEnd - seqStart)) { - continue; - } else { - // value = 0; - value = inputGrad[(i + contextStart) * inputDim + idx]; - } - - int outx = (i - contextLength) < 0 ? i : (contextLength - 1); - int outy = (i - contextLength) < 0 ? 0 : (i - (contextLength - 1)); - real* output_r = - outputGrad + outy * inputDim * contextLength + outx * inputDim; - for (int j = outy; j < seqEnd - seqStart; j++) { - value += output_r[idx]; - if (j - outy == outx) break; - output_r += (contextLength - 1) * inputDim; - } - inputGrad[(i + contextStart) * inputDim + idx] = value; - } - } - idx += blockSize; - } -} - -void hl_context_projection_backward_data(real* outputGrad, - const int* sequence, - real* inputGrad, - int numSequences, - int inputDim, - int contextLength, - int contextStart) { - CHECK_NOTNULL(outputGrad); - CHECK_NOTNULL(sequence); - CHECK_NOTNULL(inputGrad); - - int blockSize = 128; - int blocksX = numSequences; - int blocksY = 1; - dim3 threads(blockSize, 1); - dim3 grid(blocksX, blocksY); - KeContextProjectionBackwardData<<< grid, threads, 0, STREAM_DEFAULT >>> - (outputGrad, sequence, inputGrad, inputDim, contextLength, contextStart); - CHECK_SYNC("hl_context_projection_backward_data failed"); -} - -template -__global__ void KeContextProjectionBackwardWeight(real* outputGrad, - const int* sequence, - real* weightGrad, - int numSequences, - int weightDim, - int contextLength, - int contextStart, - int beginPad) { - __shared__ real sum_s[THREADS_Y][THREADS_X]; - int padOfBlock = (weightDim + THREADS_X - 1) / THREADS_X; - const int idx = threadIdx.x; - const int idy = threadIdx.y; - int padId = blockIdx.x / padOfBlock; - int weightIdx = idx + THREADS_X * (blockIdx.x % padOfBlock); - int instanceId; - real value = 0; - real* output_r; - - sum_s[idy][idx] = 0.0f; - if (weightIdx < weightDim) { - for (int seqId = idy; seqId < numSequences; seqId += THREADS_Y) { - int seqStart = sequence[seqId]; - int seqEnd = sequence[seqId+1]; - output_r = outputGrad + seqStart * weightDim * contextLength; - - if (contextStart < 0) { - if (padId + contextStart < 0) { - instanceId = padId; - } else { - // beginPad > 0; - instanceId = (padId - beginPad) + (seqEnd - seqStart) - contextStart; - } - } else { - if (padId + (seqEnd - seqStart) < contextStart) { - continue; - } else { - // beginPad == 0; - instanceId = padId + (seqEnd - seqStart) - contextStart; - } - } - - int outx = (instanceId - contextLength) < 0 ? - instanceId : (contextLength - 1); - int outy = (instanceId - contextLength) < 0 ? - 0 : (instanceId - (contextLength - 1)); - output_r += outy * weightDim * contextLength + outx * weightDim; - for (int j = outy; j < seqEnd - seqStart; j++) { - value += output_r[weightIdx]; - if (j - outy == outx) break; - output_r += (contextLength - 1) * weightDim; - } - } - sum_s[idy][idx] = value; - } - __syncthreads(); - - for (int stride = THREADS_Y/2; stride > 0; stride = stride/2) { - if (idy < stride) { - sum_s[idy][idx] += sum_s[idy + stride][idx]; - } - __syncthreads(); - } - __syncthreads(); - - if (weightIdx < weightDim) { - if (idy == 0) { - weightGrad[padId * weightDim + weightIdx] += sum_s[0][idx]; - } - } -} - -void hl_context_projection_backward_weight(real* outputGrad, - const int* sequence, - real* weightGrad, - int numSequences, - int weightDim, - int totalPad, - int contextLength, - int contextStart, - int beginPad) { - CHECK_NOTNULL(outputGrad); - CHECK_NOTNULL(sequence); - CHECK_NOTNULL(weightGrad); - - int threadsX = 32; - int threadsY = 32; - int blocksX = totalPad * ((weightDim + threadsX - 1) / threadsX); - dim3 threads(threadsX, threadsY); - dim3 grid(blocksX, 1); - - KeContextProjectionBackwardWeight<32, 32> - <<< grid, threads, 0, STREAM_DEFAULT >>> - (outputGrad, sequence, weightGrad, numSequences, weightDim, - contextLength, contextStart, beginPad); - CHECK_SYNC("hl_context_projection_backward_weight failed"); -} - template __global__ void KeMatrixAddRows(real* output, real* table, diff --git a/paddle/function/ContextProjectionOpGpu.cu b/paddle/function/ContextProjectionOpGpu.cu index 7c4ebacdbf1ff11a8c66ae83b2c8c44adf8ae4bd..1ec7058f96c8200728e5add051d5fa6a77a97e36 100644 --- a/paddle/function/ContextProjectionOpGpu.cu +++ b/paddle/function/ContextProjectionOpGpu.cu @@ -73,15 +73,30 @@ __global__ void KeContextProjectionForward(const real* input, } } +/** + * @brief Context projection forward. + * + * @param[in] input input sequence. + * @param[in] sequence sequence index. + * @param[in] weight padding data. + * @param[out] output output sequence. + * @param[in] num_sequences number of sequences. + * @param[in] input_dim input sequence dimension. + * @param[in] context_length context length. + * @param[in] context_start context start. + * @param[in] begin_pad number of extra timesteps added at the + * beginning. + * + */ void hl_context_projection_forward(const real* input, const int* sequence, const real* weight, real* output, - int num_sequences, - int input_dim, - int context_length, + size_t num_sequences, + size_t input_dim, + size_t context_length, int context_start, - int begin_pad) { + size_t begin_pad) { CHECK_NOTNULL(input); CHECK_NOTNULL(sequence); CHECK_NOTNULL(output); @@ -168,12 +183,24 @@ __global__ void KeContextProjectionBackwardData(real* out_grad, } } +/** + * @brief Context projection backward data. + * + * @param[in] out_grad output gradient. + * @param[in] sequence sequence index. + * @param[out] input_grad input gradient. + * @param[in] num_sequences number of sequences. + * @param[in] input_dim input sequence dimension. + * @param[in] context_length context length. + * @param[in] context_start context start. + * + */ void hl_context_projection_backward_data(real* out_grad, const int* sequence, real* input_grad, - int num_sequences, - int input_dim, - int context_length, + size_t num_sequences, + size_t input_dim, + size_t context_length, int context_start) { CHECK_NOTNULL(out_grad); CHECK_NOTNULL(sequence); @@ -278,15 +305,30 @@ __global__ void KeContextProjectionBackwardWeight(real* out_grad, } } +/** + * @brief Context projection backward weight. + * + * @param[in] out_grad output gradient. + * @param[in] sequence sequence index. + * @param[out] w_grad weight gradient. + * @param[in] num_sequences number of sequences. + * @param[in] w_dim input sequence dimension. + * @param[in] total_pad number of extra timesteps. + * @param[in] context_length context length. + * @param[in] context_start context start. + * @param[in] begin_pad number of extra timesteps added at the + * beginning. + * + */ void hl_context_projection_backward_weight(real* out_grad, const int* sequence, real* w_grad, - int num_sequences, - int w_dim, + size_t num_sequences, + size_t w_dim, size_t total_pad, - int context_length, + size_t context_length, int context_start, - int begin_pad) { + size_t begin_pad) { CHECK_NOTNULL(out_grad); CHECK_NOTNULL(sequence); CHECK_NOTNULL(w_grad); diff --git a/paddle/math/Matrix.cpp b/paddle/math/Matrix.cpp index 50d2e3eb671028c8169321fcd85fe25735c11a14..90813a89969c2525f7029f1c2609bed116c910c4 100644 --- a/paddle/math/Matrix.cpp +++ b/paddle/math/Matrix.cpp @@ -1304,68 +1304,6 @@ void GpuMatrix::maxSequenceBackward(Matrix& outputGrad, hl_max_sequence_backward(outGrad, maxIndex, inputGrad, numSequences, dim); } -void GpuMatrix::contextProjectionForward(Matrix& input, - Matrix* weight, - const IVector& sequence, - int contextLength, - int contextStart, - size_t beginPad, - bool isPadding) { - CHECK(dynamic_cast(&input)); - CHECK(dynamic_cast(&sequence)); - if (weight) CHECK(dynamic_cast(weight)); - CHECK_EQ(getWidth(), input.getWidth() * contextLength); - - hl_context_projection_forward(input.getData(), - sequence.getData(), - isPadding ? weight->getData() : NULL, - getData(), - sequence.getSize() - 1, - input.getWidth(), - contextLength, - contextStart, - beginPad, - isPadding); -} - -void GpuMatrix::contextProjectionBackwardData(Matrix& inputGrad, - const IVector& sequence, - int contextLength, - int contextStart) { - CHECK(dynamic_cast(&inputGrad)); - CHECK(dynamic_cast(&sequence)); - CHECK_EQ(getWidth(), inputGrad.getWidth() * contextLength); - - hl_context_projection_backward_data(getData(), - sequence.getData(), - inputGrad.getData(), - sequence.getSize() - 1, - inputGrad.getWidth(), - contextLength, - contextStart); -} - -void GpuMatrix::contextProjectionBackwardWeight(Matrix& weightGrad, - const IVector& sequence, - int contextLength, - int contextStart, - int totalPad, - size_t beginPad) { - CHECK(dynamic_cast(&weightGrad)); - CHECK(dynamic_cast(&sequence)); - CHECK_EQ(getWidth(), weightGrad.getWidth() * contextLength); - - hl_context_projection_backward_weight(getData(), - sequence.getData(), - weightGrad.getData(), - sequence.getSize() - 1, - weightGrad.getWidth(), - totalPad, - contextLength, - contextStart, - beginPad); -} - void GpuMatrix::paramReluForward(Matrix& data, Matrix& W) { CHECK(data.useGpu_ == true && W.useGpu_ == true) << "Matrix type are not equal"; @@ -2203,113 +2141,6 @@ void CpuMatrix::maxSequenceBackward(Matrix& outputGrad, } } -void CpuMatrix::contextProjectionForward(Matrix& input, - Matrix* weight, - const IVector& sequence, - int contextLength, - int contextStart, - size_t beginPad, - bool isPadding) { - auto input_ptr = dynamic_cast(&input); - auto seq_ptr = dynamic_cast(&sequence); - CHECK(input_ptr && seq_ptr); - if (weight) CHECK(dynamic_cast(weight)); - CHECK_EQ(getWidth(), input_ptr->getWidth() * contextLength); - - const int* starts = seq_ptr->getData(); - size_t numSequences = seq_ptr->getSize() - 1; - for (size_t i = 0; i < numSequences; ++i) { - for (int j = 0; j < contextLength; ++j) { - int begin = starts[i] + contextStart + j; - int end = starts[i + 1] + contextStart + j; - int dstBegin = starts[i]; - int dstEnd = starts[i + 1]; - if (begin < starts[i]) { - int64_t padSize = - std::min(starts[i] - begin, starts[i + 1] - starts[i]); - MatrixPtr mat = this->subMatrix(starts[i], padSize); - if (isPadding) { - MatrixPtr sub = weight->subMatrix(j, padSize); - mat->addAtOffset(*sub, j * input_ptr->getWidth()); - } - dstBegin = starts[i] + padSize; - begin = starts[i]; - } - if (end > starts[i + 1]) { - int64_t padSize = - std::min(end - starts[i + 1], starts[i + 1] - starts[i]); - MatrixPtr mat = this->subMatrix(starts[i + 1] - padSize, padSize); - if (isPadding) { - MatrixPtr sub = - weight->subMatrix(beginPad + contextStart + j - padSize, padSize); - mat->addAtOffset(*sub, j * input_ptr->getWidth()); - } - dstEnd = starts[i + 1] - padSize; - end = starts[i + 1]; - } - if (end <= begin) continue; - MatrixPtr src = input_ptr->subMatrix(begin, end - begin); - MatrixPtr dst = this->subMatrix(dstBegin, dstEnd - dstBegin); - dst->addAtOffset(*src, j * input_ptr->getWidth()); - } - } -} - -void CpuMatrix::contextProjectionBackward(Matrix* inputGrad, - Matrix* weightGrad, - const IVector& sequence, - int contextLength, - int contextStart, - size_t beginPad, - bool isPadding) { - if (inputGrad) CHECK(dynamic_cast(inputGrad)); - if (weightGrad) CHECK(dynamic_cast(weightGrad)); - CHECK(dynamic_cast(&sequence)); - - int64_t inputDim = inputGrad ? inputGrad->getWidth() - : weightGrad ? weightGrad->getWidth() : 0; - CHECK_EQ(getWidth(), inputDim * contextLength); - - const int* starts = sequence.getData(); - size_t numSequences = sequence.getSize() - 1; - for (size_t i = 0; i < numSequences; ++i) { - for (int j = 0; j < contextLength; ++j) { - int begin = starts[i] + contextStart + j; - int end = starts[i + 1] + contextStart + j; - int dstBegin = starts[i]; - int dstEnd = starts[i + 1]; - if (begin < starts[i]) { - int64_t padSize = - std::min(starts[i] - begin, starts[i + 1] - starts[i]); - if (isPadding && weightGrad) { - MatrixPtr mat = this->subMatrix(starts[i], padSize); - MatrixPtr sub = weightGrad->subMatrix(j, padSize); - sub->addAtOffset(*mat, j * inputDim); - } - dstBegin = starts[i] + padSize; - begin = starts[i]; - } - if (end > starts[i + 1]) { - int64_t padSize = - std::min(end - starts[i + 1], starts[i + 1] - starts[i]); - if (isPadding && weightGrad) { - MatrixPtr mat = this->subMatrix(starts[i + 1] - padSize, padSize); - MatrixPtr sub = weightGrad->subMatrix( - beginPad + contextStart + j - padSize, padSize); - sub->addAtOffset(*mat, j * inputDim); - } - dstEnd = starts[i + 1] - padSize; - end = starts[i + 1]; - } - if (end <= begin) continue; - if (!inputGrad) continue; - MatrixPtr src = inputGrad->subMatrix(begin, end - begin); - MatrixPtr dst = this->subMatrix(dstBegin, dstEnd - dstBegin); - src->addAtOffset(*dst, j * inputDim); - } - } -} - inline void vecAddTo(real* a, const real* b, size_t len) { for (unsigned int i = 0; i < len; ++i) { a[i] += b[i]; diff --git a/paddle/math/Matrix.h b/paddle/math/Matrix.h index bda863de38675fe481544a7e82b69f445df361bd..4865a081a5aaa010d5b3ce0127ffc6f8330d4a68 100644 --- a/paddle/math/Matrix.h +++ b/paddle/math/Matrix.h @@ -972,42 +972,6 @@ public: LOG(FATAL) << "Not implemeted"; } - virtual void contextProjectionForward(Matrix& input, - Matrix* weight, - const IVector& sequence, - int contextLength, - int contextStart, - size_t beginPad, - bool isPadding) { - LOG(FATAL) << "Not implemeted"; - } - - virtual void contextProjectionBackward(Matrix* inputGrad, - Matrix* weightGrad, - const IVector& sequence, - int contextLength, - int contextStart, - size_t beginPad, - bool isPadding) { - LOG(FATAL) << "Not implemeted"; - } - - virtual void contextProjectionBackwardData(Matrix& inputGrad, - const IVector& sequence, - int contextLength, - int contextStart) { - LOG(FATAL) << "Not implemeted"; - } - - virtual void contextProjectionBackwardWeight(Matrix& weightGrad, - const IVector& sequence, - int contextLength, - int contextStart, - int totalPad, - size_t beginPad) { - LOG(FATAL) << "Not implemeted"; - } - /** * @code * this.row[i] += table.row[ids[i]] @@ -1442,26 +1406,6 @@ public: const IVector& sequence, IVector& index); - void contextProjectionForward(Matrix& input, - Matrix* weight, - const IVector& sequence, - int contextLength, - int contextStart, - size_t beginPad, - bool isPadding); - - void contextProjectionBackwardData(Matrix& inputGrad, - const IVector& sequence, - int contextLength, - int contextStart); - - void contextProjectionBackwardWeight(Matrix& weightGrad, - const IVector& sequence, - int contextLength, - int contextStart, - int totalPad, - size_t beginPad); - void bilinearForward(const Matrix& in, const size_t inImgH, const size_t inImgW, @@ -1648,22 +1592,6 @@ public: const IVector& sequence, IVector& index); - void contextProjectionForward(Matrix& input, - Matrix* weight, - const IVector& sequence, - int contextLength, - int contextStart, - size_t beginPad, - bool isPadding); - - void contextProjectionBackward(Matrix* inputGrad, - Matrix* weightGrad, - const IVector& sequence, - int contextLength, - int contextStart, - size_t beginPad, - bool isPadding); - real* getRow(size_t row) { return BaseMatrix::rowBuf(row); } virtual real* getRowBuf(size_t row) { return getRow(row); } diff --git a/paddle/math/tests/test_matrixCompare.cpp b/paddle/math/tests/test_matrixCompare.cpp index efda4ff27be92916228b4bcc9170107ff707d6cf..98d63438a57b48340bc3b05ac7ac3d6c5cd90fb0 100644 --- a/paddle/math/tests/test_matrixCompare.cpp +++ b/paddle/math/tests/test_matrixCompare.cpp @@ -29,148 +29,6 @@ using namespace std; // NOLINT using autotest::TensorCheckEqual; using autotest::TensorCheckErr; -void testMatrixProjectionForward(int contextStart, - int contextLength, - bool padding, - int batchSize, - int inputDim) { - MatrixPtr cpuInput = std::make_shared(batchSize, inputDim); - MatrixPtr gpuInput = std::make_shared(batchSize, inputDim); - cpuInput->randomizeUniform(); - gpuInput->copyFrom(*cpuInput); - - int pad = std::max(0, -contextStart) + - std::max(0, contextStart + contextLength - 1); - if (pad == 0) padding = false; - MatrixPtr cpuWeight = nullptr; - MatrixPtr gpuWeight = nullptr; - if (padding) { - cpuWeight = std::make_shared(pad, inputDim); - gpuWeight = std::make_shared(pad, inputDim); - cpuWeight->randomizeUniform(); - gpuWeight->copyFrom(*cpuWeight); - } - - IVectorPtr cpuSequence; - generateSequenceStartPositions(batchSize, cpuSequence); - IVectorPtr gpuSequence = IVector::create(cpuSequence->getSize(), true); - gpuSequence->copyFrom(*cpuSequence); - - MatrixPtr cpuOutput = - std::make_shared(batchSize, inputDim * contextLength); - MatrixPtr gpuOutput = - std::make_shared(batchSize, inputDim * contextLength); - cpuOutput->randomizeUniform(); - gpuOutput->copyFrom(*cpuOutput); - - // calculate - int beginPad = std::max(0, -contextStart); - cpuOutput->contextProjectionForward(*cpuInput, - cpuWeight.get(), - *cpuSequence, - contextLength, - contextStart, - beginPad, - padding); - - gpuOutput->contextProjectionForward(*gpuInput, - gpuWeight.get(), - *gpuSequence, - contextLength, - contextStart, - beginPad, - padding); - - TensorCheckEqual(*cpuOutput, *gpuOutput); -} - -void testMatrixProjectionBackward(int contextStart, - int contextLength, - bool padding, - int batchSize, - int inputDim) { - MatrixPtr cpuOutputGrad = - std::make_shared(batchSize, inputDim * contextLength); - MatrixPtr gpuOutputGrad = - std::make_shared(batchSize, inputDim * contextLength); - cpuOutputGrad->randomizeUniform(); - gpuOutputGrad->copyFrom(*cpuOutputGrad); - - IVectorPtr cpuSequence; - generateSequenceStartPositions(batchSize, cpuSequence); - IVectorPtr gpuSequence = IVector::create(cpuSequence->getSize(), true); - gpuSequence->copyFrom(*cpuSequence); - - MatrixPtr cpuInputGrad = std::make_shared(batchSize, inputDim); - MatrixPtr gpuInputGrad = std::make_shared(batchSize, inputDim); - cpuInputGrad->randomizeUniform(); - gpuInputGrad->copyFrom(*cpuInputGrad); - - int pad = std::max(0, -contextStart) + - std::max(0, contextStart + contextLength - 1); - if (pad == 0) padding = false; - MatrixPtr cpuWeightGrad = nullptr; - MatrixPtr gpuWeightGrad = nullptr; - if (padding) { - cpuWeightGrad = std::make_shared(pad, inputDim); - gpuWeightGrad = std::make_shared(pad, inputDim); - cpuWeightGrad->randomizeUniform(); - gpuWeightGrad->copyFrom(*cpuWeightGrad); - } - - // calculate - int beginPad = std::max(0, -contextStart); - cpuOutputGrad->contextProjectionBackward(cpuInputGrad.get(), - cpuWeightGrad.get(), - *cpuSequence, - contextLength, - contextStart, - beginPad, - padding); - gpuOutputGrad->contextProjectionBackwardData( - *gpuInputGrad, *gpuSequence, contextLength, contextStart); - if (padding) { - gpuOutputGrad->contextProjectionBackwardWeight(*gpuWeightGrad, - *gpuSequence, - contextLength, - contextStart, - pad, - beginPad); - } - - TensorCheckErr(*cpuInputGrad, *gpuInputGrad); - if (padding) { - TensorCheckErr(*cpuWeightGrad, *gpuWeightGrad); - } -} - -TEST(Matrix, projection) { - for (auto contextStart : {-5, -3, -1, 0, 3}) { - for (auto contextLength : {1, 2, 5, 7}) { - for (auto trainablePadding : {false, true}) { - for (auto batchSize : {1, 2, 5, 20, 100}) { - for (auto inputDim : {15, 32, 63, 128, 200}) { - VLOG(3) << " contextStart=" << contextStart - << " contextLength=" << contextLength - << " trainablePadding=" << trainablePadding - << " batchSize=" << batchSize << " inputDim=" << inputDim; - testMatrixProjectionForward(contextStart, - contextLength, - trainablePadding, - batchSize, - inputDim); - testMatrixProjectionBackward(contextStart, - contextLength, - trainablePadding, - batchSize, - inputDim); - } - } - } - } - } -} - void testMatrixMaxSequence(int batchSize, int inputDim) { // forward MatrixPtr cpuInput = std::make_shared(batchSize, inputDim);