提交 514cbeff 编写于 作者: Y Yu Yang

Merge branch 'develop' of github.com:baidu/Paddle into feature/refine_doc_drnn

...@@ -7,18 +7,14 @@ ...@@ -7,18 +7,14 @@
hooks: hooks:
- id: yapf - id: yapf
- repo: https://github.com/pre-commit/pre-commit-hooks - repo: https://github.com/pre-commit/pre-commit-hooks
sha: 4ef03c4223ad322c7adaa6c6c0efb26b57df3b71 sha: 7539d8bd1a00a3c1bfd34cdb606d3a6372e83469
hooks: hooks:
- id: check-added-large-files - id: check-added-large-files
- id: check-merge-conflict - id: check-merge-conflict
- id: check-symlinks - id: check-symlinks
- id: detect-private-key - id: detect-private-key
- id: end-of-file-fixer - id: end-of-file-fixer
# TODO(yuyang): trailing whitespace has some bugs on markdown - repo: https://github.com/PaddlePaddle/clang-format-pre-commit-hook.git
# files now, please not add it to pre-commit hook now sha: 28c0ea8a67a3e2dbbf4822ef44e85b63a0080a29
# - id: trailing-whitespace hooks:
# - id: clang-formater
# TODO(yuyang): debug-statements not fit for Paddle, because
# not all of our python code is runnable. Some are used for
# documenation
# - id: debug-statements
# PaddlePaddle # PaddlePaddle
[![Build Status](https://travis-ci.org/baidu/Paddle.svg?branch=master)](https://travis-ci.org/baidu/Paddle) [![Build Status](https://travis-ci.org/PaddlePaddle/Paddle.svg?branch=develop)](https://travis-ci.org/PaddlePaddle/Paddle)
[![Coverage Status](https://coveralls.io/repos/github/baidu/Paddle/badge.svg?branch=develop)](https://coveralls.io/github/baidu/Paddle?branch=develop) [![Documentation Status](https://img.shields.io/badge/docs-latest-brightgreen.svg?style=flat)](http://www.paddlepaddle.org/)
[![Join the chat at https://gitter.im/PaddlePaddle/Deep_Learning](https://badges.gitter.im/Join%20Chat.svg)](https://gitter.im/PaddlePaddle/Deep_Learning?utm_source=badge&utm_medium=badge&utm_campaign=pr-badge&utm_content=badge) [![Documentation Status](https://img.shields.io/badge/中文文档-最新-brightgreen.svg)](http://www.paddlepaddle.org/cn/index.html)
[![License](https://img.shields.io/badge/license-Apache%202.0-green.svg)](LICENSE) [![Coverage Status](https://coveralls.io/repos/github/PaddlePaddle/Paddle/badge.svg?branch=develop)](https://coveralls.io/github/PaddlePaddle/Paddle?branch=develop)
[![Release](https://img.shields.io/github/release/PaddlePaddle/Paddle.svg)](https://github.com/PaddlePaddle/Paddle/releases)
[![License](https://img.shields.io/badge/license-Apache%202-blue.svg)](LICENSE)
Welcome to the PaddlePaddle GitHub. Welcome to the PaddlePaddle GitHub.
...@@ -14,7 +17,7 @@ developed by Baidu scientists and engineers for the purpose of applying deep ...@@ -14,7 +17,7 @@ developed by Baidu scientists and engineers for the purpose of applying deep
learning to many products at Baidu. learning to many products at Baidu.
Our vision is to enable deep learning for everyone via PaddlePaddle. Our vision is to enable deep learning for everyone via PaddlePaddle.
Please refer to our [release announcement](https://github.com/baidu/Paddle/releases) to track the latest feature of PaddlePaddle. Please refer to our [release announcement](https://github.com/PaddlePaddle/Paddle/releases) to track the latest feature of PaddlePaddle.
## Features ## Features
...@@ -89,7 +92,7 @@ Both [English Docs](http://paddlepaddle.org/doc/) and [Chinese Docs](http://padd ...@@ -89,7 +92,7 @@ Both [English Docs](http://paddlepaddle.org/doc/) and [Chinese Docs](http://padd
## Ask Questions ## Ask Questions
You are welcome to submit questions and bug reports as [Github Issues](https://github.com/baidu/paddle/issues). You are welcome to submit questions and bug reports as [Github Issues](https://github.com/PaddlePaddle/Paddle/issues).
## Copyright and License ## Copyright and License
PaddlePaddle is provided under the [Apache-2.0 license](LICENSE). PaddlePaddle is provided under the [Apache-2.0 license](LICENSE).
...@@ -17,24 +17,15 @@ import os ...@@ -17,24 +17,15 @@ import os
from optparse import OptionParser from optparse import OptionParser
def extract_dict_features(pair_file, feature_file, src_dict_file, def extract_dict_features(pair_file, feature_file):
tgt_dict_file):
src_dict = set() with open(pair_file) as fin, open(feature_file, 'w') as feature_out:
tgt_dict = set()
with open(pair_file) as fin, open(feature_file, 'w') as feature_out, open(
src_dict_file, 'w') as src_dict_out, open(tgt_dict_file,
'w') as tgt_dict_out:
for line in fin: for line in fin:
sentence, labels = line.strip().split('\t') sentence, predicate, labels = line.strip().split('\t')
sentence_list = sentence.split() sentence_list = sentence.split()
labels_list = labels.split() labels_list = labels.split()
src_dict.update(sentence_list)
tgt_dict.update(labels_list)
verb_index = labels_list.index('B-V') verb_index = labels_list.index('B-V')
verb_feature = sentence_list[verb_index]
mark = [0] * len(labels_list) mark = [0] * len(labels_list)
if verb_index > 0: if verb_index > 0:
...@@ -42,47 +33,50 @@ def extract_dict_features(pair_file, feature_file, src_dict_file, ...@@ -42,47 +33,50 @@ def extract_dict_features(pair_file, feature_file, src_dict_file,
ctx_n1 = sentence_list[verb_index - 1] ctx_n1 = sentence_list[verb_index - 1]
else: else:
ctx_n1 = 'bos' ctx_n1 = 'bos'
ctx_n1_feature = ctx_n1
if verb_index > 1:
mark[verb_index - 2] = 1
ctx_n2 = sentence_list[verb_index - 2]
else:
ctx_n2 = 'bos'
mark[verb_index] = 1 mark[verb_index] = 1
ctx_0_feature = sentence_list[verb_index] ctx_0 = sentence_list[verb_index]
if verb_index < len(labels_list) - 2: if verb_index < len(labels_list) - 2:
mark[verb_index + 1] = 1 mark[verb_index + 1] = 1
ctx_p1 = sentence_list[verb_index + 1] ctx_p1 = sentence_list[verb_index + 1]
else: else:
ctx_p1 = 'eos' ctx_p1 = 'eos'
ctx_p1_feature = ctx_p1
if verb_index < len(labels_list) - 3:
mark[verb_index + 2] = 1
ctx_p2 = sentence_list[verb_index + 2]
else:
ctx_p2 = 'eos'
feature_str = sentence + '\t' \ feature_str = sentence + '\t' \
+ verb_feature + '\t' \ + predicate + '\t' \
+ ctx_n1_feature + '\t' \ + ctx_n2 + '\t' \
+ ctx_0_feature + '\t' \ + ctx_n1 + '\t' \
+ ctx_p1_feature + '\t' \ + ctx_0 + '\t' \
+ ctx_p1 + '\t' \
+ ctx_p2 + '\t' \
+ ' '.join([str(i) for i in mark]) + '\t' \ + ' '.join([str(i) for i in mark]) + '\t' \
+ labels + labels
feature_out.write(feature_str + '\n') feature_out.write(feature_str + '\n')
src_dict_out.write('<unk>\n')
src_dict_out.write('\n'.join(list(src_dict)))
tgt_dict_out.write('\n'.join(list(tgt_dict)))
if __name__ == '__main__': if __name__ == '__main__':
usage = '-p pair_file -f feature_file -s source dictionary -t target dictionary ' usage = '-p pair_file -f feature_file'
parser = OptionParser(usage) parser = OptionParser(usage)
parser.add_option('-p', dest='pair_file', help='the pair file') parser.add_option('-p', dest='pair_file', help='the pair file')
parser.add_option( parser.add_option('-f', dest='feature_file', help='the feature file')
'-f', dest='feature_file', help='the file to store feature')
parser.add_option(
'-s', dest='src_dict', help='the file to store source dictionary')
parser.add_option(
'-t', dest='tgt_dict', help='the file to store target dictionary')
(options, args) = parser.parse_args() (options, args) = parser.parse_args()
extract_dict_features(options.pair_file, options.feature_file, extract_dict_features(options.pair_file, options.feature_file)
options.src_dict, options.tgt_dict)
...@@ -51,7 +51,7 @@ def read_sentences(words_file): ...@@ -51,7 +51,7 @@ def read_sentences(words_file):
for line in fin: for line in fin:
line = line.strip() line = line.strip()
if line == '': if line == '':
sentences.append(s.lower()) sentences.append(s)
s = '' s = ''
else: else:
s += line + ' ' s += line + ' '
...@@ -64,6 +64,11 @@ def transform_labels(sentences, labels): ...@@ -64,6 +64,11 @@ def transform_labels(sentences, labels):
if len(labels[i]) == 1: if len(labels[i]) == 1:
continue continue
else: else:
verb_list = []
for x in labels[i][0]:
if x !='-':
verb_list.append(x)
for j in xrange(1, len(labels[i])): for j in xrange(1, len(labels[i])):
label_list = labels[i][j] label_list = labels[i][j]
current_tag = 'O' current_tag = 'O'
...@@ -88,8 +93,7 @@ def transform_labels(sentences, labels): ...@@ -88,8 +93,7 @@ def transform_labels(sentences, labels):
is_in_bracket = True is_in_bracket = True
else: else:
print 'error:', ll print 'error:', ll
sen_lab_pair.append((sentences[i], verb_list[j-1], label_seq))
sen_lab_pair.append((sentences[i], label_seq))
return sen_lab_pair return sen_lab_pair
...@@ -97,9 +101,9 @@ def write_file(sen_lab_pair, output_file): ...@@ -97,9 +101,9 @@ def write_file(sen_lab_pair, output_file):
with open(output_file, 'w') as fout: with open(output_file, 'w') as fout:
for x in sen_lab_pair: for x in sen_lab_pair:
sentence = x[0] sentence = x[0]
label_seq = ' '.join(x[1]) label_seq = ' '.join(x[2])
assert len(sentence.split()) == len(x[1]) assert len(sentence.split()) == len(x[2])
fout.write(sentence + '\t' + label_seq + '\n') fout.write(sentence + '\t' + x[1]+'\t' +label_seq + '\n')
if __name__ == '__main__': if __name__ == '__main__':
......
...@@ -14,6 +14,10 @@ ...@@ -14,6 +14,10 @@
# limitations under the License. # limitations under the License.
set -e set -e
wget http://www.cs.upc.edu/~srlconll/conll05st-tests.tar.gz wget http://www.cs.upc.edu/~srlconll/conll05st-tests.tar.gz
wget https://www.googledrive.com/host/0B7Q8d52jqeI9ejh6Q1RpMTFQT1k/semantic_role_labeling/verbDict.txt --no-check-certificate
wget https://www.googledrive.com/host/0B7Q8d52jqeI9ejh6Q1RpMTFQT1k/semantic_role_labeling/targetDict.txt --no-check-certificate
wget https://www.googledrive.com/host/0B7Q8d52jqeI9ejh6Q1RpMTFQT1k/semantic_role_labeling/wordDict.txt --no-check-certificate
wget https://www.googledrive.com/host/0B7Q8d52jqeI9ejh6Q1RpMTFQT1k/semantic_role_labeling/emb --no-check-certificate
tar -xzvf conll05st-tests.tar.gz tar -xzvf conll05st-tests.tar.gz
rm conll05st-tests.tar.gz rm conll05st-tests.tar.gz
cp ./conll05st-release/test.wsj/words/test.wsj.words.gz . cp ./conll05st-release/test.wsj/words/test.wsj.words.gz .
...@@ -22,4 +26,4 @@ gunzip test.wsj.words.gz ...@@ -22,4 +26,4 @@ gunzip test.wsj.words.gz
gunzip test.wsj.props.gz gunzip test.wsj.props.gz
python extract_pairs.py -w test.wsj.words -p test.wsj.props -o test.wsj.seq_pair python extract_pairs.py -w test.wsj.words -p test.wsj.props -o test.wsj.seq_pair
python extract_dict_feature.py -p test.wsj.seq_pair -f feature -s src.dict -t tgt.dict python extract_dict_feature.py -p test.wsj.seq_pair -f feature
...@@ -17,11 +17,15 @@ from paddle.trainer.PyDataProvider2 import * ...@@ -17,11 +17,15 @@ from paddle.trainer.PyDataProvider2 import *
UNK_IDX = 0 UNK_IDX = 0
def hook(settings, word_dict, label_dict, **kwargs): def hook(settings, word_dict, label_dict, predicate_dict, **kwargs):
settings.word_dict = word_dict settings.word_dict = word_dict
settings.label_dict = label_dict settings.label_dict = label_dict
settings.predicate_dict = predicate_dict
#all inputs are integral and sequential type #all inputs are integral and sequential type
settings.slots = [ settings.slots = [
integer_value_sequence(len(word_dict)),
integer_value_sequence(len(predicate_dict)),
integer_value_sequence(len(word_dict)), integer_value_sequence(len(word_dict)),
integer_value_sequence(len(word_dict)), integer_value_sequence(len(word_dict)),
integer_value_sequence(len(word_dict)), integer_value_sequence(len(word_dict)),
...@@ -31,27 +35,33 @@ def hook(settings, word_dict, label_dict, **kwargs): ...@@ -31,27 +35,33 @@ def hook(settings, word_dict, label_dict, **kwargs):
] ]
@provider(init_hook=hook) def get_batch_size(yeild_data):
def process(obj, file_name): return len(yeild_data[0])
@provider(init_hook=hook, should_shuffle=True, calc_batch_size=get_batch_size,
can_over_batch_size=False, cache=CacheType.CACHE_PASS_IN_MEM)
def process(settings, file_name):
with open(file_name, 'r') as fdata: with open(file_name, 'r') as fdata:
for line in fdata: for line in fdata:
sentence, predicate, ctx_n1, ctx_0, ctx_p1, mark, label = \ sentence, predicate, ctx_n2, ctx_n1, ctx_0, ctx_p1, ctx_p2, mark, label = \
line.strip().split('\t') line.strip().split('\t')
words = sentence.split() words = sentence.split()
sen_len = len(words) sen_len = len(words)
word_slot = [obj.word_dict.get(w, UNK_IDX) for w in words] word_slot = [settings.word_dict.get(w, UNK_IDX) for w in words]
predicate_slot = [obj.word_dict.get(predicate, UNK_IDX)] * sen_len predicate_slot = [settings.predicate_dict.get(predicate)] * sen_len
ctx_n1_slot = [obj.word_dict.get(ctx_n1, UNK_IDX)] * sen_len ctx_n2_slot = [settings.word_dict.get(ctx_n2, UNK_IDX)] * sen_len
ctx_0_slot = [obj.word_dict.get(ctx_0, UNK_IDX)] * sen_len ctx_n1_slot = [settings.word_dict.get(ctx_n1, UNK_IDX)] * sen_len
ctx_p1_slot = [obj.word_dict.get(ctx_p1, UNK_IDX)] * sen_len ctx_0_slot = [settings.word_dict.get(ctx_0, UNK_IDX)] * sen_len
ctx_p1_slot = [settings.word_dict.get(ctx_p1, UNK_IDX)] * sen_len
ctx_p2_slot = [settings.word_dict.get(ctx_p2, UNK_IDX)] * sen_len
marks = mark.split() marks = mark.split()
mark_slot = [int(w) for w in marks] mark_slot = [int(w) for w in marks]
label_list = label.split() label_list = label.split()
label_slot = [obj.label_dict.get(w) for w in label_list] label_slot = [settings.label_dict.get(w) for w in label_list]
yield word_slot, predicate_slot, ctx_n2_slot, ctx_n1_slot, \
yield word_slot, predicate_slot, ctx_n1_slot, \ ctx_0_slot, ctx_p1_slot, ctx_p2_slot, mark_slot, label_slot
ctx_0_slot, ctx_p1_slot, mark_slot, label_slot
...@@ -18,8 +18,9 @@ import sys ...@@ -18,8 +18,9 @@ import sys
from paddle.trainer_config_helpers import * from paddle.trainer_config_helpers import *
#file paths #file paths
word_dict_file = './data/src.dict' word_dict_file = './data/wordDict.txt'
label_dict_file = './data/tgt.dict' label_dict_file = './data/targetDict.txt'
predicate_file= './data/verbDict.txt'
train_list_file = './data/train.list' train_list_file = './data/train.list'
test_list_file = './data/test.list' test_list_file = './data/test.list'
...@@ -30,8 +31,10 @@ if not is_predict: ...@@ -30,8 +31,10 @@ if not is_predict:
#load dictionaries #load dictionaries
word_dict = dict() word_dict = dict()
label_dict = dict() label_dict = dict()
predicate_dict = dict()
with open(word_dict_file, 'r') as f_word, \ with open(word_dict_file, 'r') as f_word, \
open(label_dict_file, 'r') as f_label: open(label_dict_file, 'r') as f_label, \
open(predicate_file, 'r') as f_pre:
for i, line in enumerate(f_word): for i, line in enumerate(f_word):
w = line.strip() w = line.strip()
word_dict[w] = i word_dict[w] = i
...@@ -40,6 +43,11 @@ if not is_predict: ...@@ -40,6 +43,11 @@ if not is_predict:
w = line.strip() w = line.strip()
label_dict[w] = i label_dict[w] = i
for i, line in enumerate(f_pre):
w = line.strip()
predicate_dict[w] = i
if is_test: if is_test:
train_list_file = None train_list_file = None
...@@ -50,91 +58,157 @@ if not is_predict: ...@@ -50,91 +58,157 @@ if not is_predict:
module='dataprovider', module='dataprovider',
obj='process', obj='process',
args={'word_dict': word_dict, args={'word_dict': word_dict,
'label_dict': label_dict}) 'label_dict': label_dict,
'predicate_dict': predicate_dict })
word_dict_len = len(word_dict) word_dict_len = len(word_dict)
label_dict_len = len(label_dict) label_dict_len = len(label_dict)
pred_len = len(predicate_dict)
else: else:
word_dict_len = get_config_arg('dict_len', int) word_dict_len = get_config_arg('dict_len', int)
label_dict_len = get_config_arg('label_len', int) label_dict_len = get_config_arg('label_len', int)
pred_len = get_config_arg('pred_len', int)
############################## Hyper-parameters ##################################
mark_dict_len = 2 mark_dict_len = 2
word_dim = 32 word_dim = 32
mark_dim = 5 mark_dim = 5
hidden_dim = 128 hidden_dim = 512
depth = 8 depth = 8
emb_lr = 1e-2
fc_lr = 1e-2
lstm_lr = 2e-2
########################### Optimizer #######################################
settings( settings(
batch_size=150, batch_size=150,
learning_method=AdamOptimizer(), learning_method=MomentumOptimizer(momentum=0),
learning_rate=1e-3, learning_rate=2e-2,
regularization=L2Regularization(8e-4), regularization=L2Regularization(8e-4),
gradient_clipping_threshold=25) is_async=False,
model_average=ModelAverage(average_window=0.5,
max_average_window=10000),
)
#6 features
####################################### network ##############################
#8 features and 1 target
word = data_layer(name='word_data', size=word_dict_len) word = data_layer(name='word_data', size=word_dict_len)
predicate = data_layer(name='verb_data', size=word_dict_len) predicate = data_layer(name='verb_data', size=pred_len)
ctx_n2 = data_layer(name='ctx_n2_data', size=word_dict_len)
ctx_n1 = data_layer(name='ctx_n1_data', size=word_dict_len) ctx_n1 = data_layer(name='ctx_n1_data', size=word_dict_len)
ctx_0 = data_layer(name='ctx_0_data', size=word_dict_len) ctx_0 = data_layer(name='ctx_0_data', size=word_dict_len)
ctx_p1 = data_layer(name='ctx_p1_data', size=word_dict_len) ctx_p1 = data_layer(name='ctx_p1_data', size=word_dict_len)
ctx_p2 = data_layer(name='ctx_p2_data', size=word_dict_len)
mark = data_layer(name='mark_data', size=mark_dict_len) mark = data_layer(name='mark_data', size=mark_dict_len)
if not is_predict: if not is_predict:
target = data_layer(name='target', size=label_dict_len) target = data_layer(name='target', size=label_dict_len)
ptt = ParameterAttribute(name='src_emb', learning_rate=emb_lr)
layer_attr = ExtraLayerAttribute(drop_rate=0.5)
fc_para_attr = ParameterAttribute(learning_rate=fc_lr)
lstm_para_attr = ParameterAttribute(initial_std=0., learning_rate=lstm_lr)
para_attr = [fc_para_attr, lstm_para_attr]
word_embedding = embedding_layer(size=word_dim, input=word, param_attr=ptt) default_std=1/math.sqrt(hidden_dim)/3.0
predicate_embedding = embedding_layer(
size=word_dim, input=predicate, param_attr=ptt) emb_para = ParameterAttribute(name='emb', initial_std=0., learning_rate=0.)
ctx_n1_embedding = embedding_layer(size=word_dim, input=ctx_n1, param_attr=ptt) std_0 = ParameterAttribute(initial_std=0.)
ctx_0_embedding = embedding_layer(size=word_dim, input=ctx_0, param_attr=ptt) std_default = ParameterAttribute(initial_std=default_std)
ctx_p1_embedding = embedding_layer(size=word_dim, input=ctx_p1, param_attr=ptt)
mark_embedding = embedding_layer(size=mark_dim, input=mark) predicate_embedding = embedding_layer(size=word_dim, input=predicate, param_attr=ParameterAttribute(name='vemb',initial_std=default_std))
mark_embedding = embedding_layer(name='word_ctx-in_embedding', size=mark_dim, input=mark, param_attr=std_0)
word_input=[word, ctx_n2, ctx_n1, ctx_0, ctx_p1, ctx_p2]
emb_layers = [embedding_layer(size=word_dim, input=x, param_attr=emb_para) for x in word_input]
emb_layers.append(predicate_embedding)
emb_layers.append(mark_embedding)
hidden_0 = mixed_layer( hidden_0 = mixed_layer(
name='hidden0',
size=hidden_dim, size=hidden_dim,
input=[ bias_attr=std_default,
full_matrix_projection(input=word_embedding), input=[ full_matrix_projection(input=emb, param_attr=std_default ) for emb in emb_layers ])
full_matrix_projection(input=predicate_embedding),
full_matrix_projection(input=ctx_n1_embedding),
full_matrix_projection(input=ctx_0_embedding),
full_matrix_projection(input=ctx_p1_embedding),
full_matrix_projection(input=mark_embedding),
])
lstm_0 = lstmemory(input=hidden_0, layer_attr=layer_attr) mix_hidden_lr = 1e-3
lstm_para_attr = ParameterAttribute(initial_std=0.0, learning_rate=1.0)
hidden_para_attr = ParameterAttribute(initial_std=default_std, learning_rate=mix_hidden_lr)
lstm_0 = lstmemory(name='lstm0',
input=hidden_0,
act=ReluActivation(),
gate_act=SigmoidActivation(),
state_act=SigmoidActivation(),
bias_attr=std_0,
param_attr=lstm_para_attr)
#stack L-LSTM and R-LSTM with direct edges #stack L-LSTM and R-LSTM with direct edges
input_tmp = [hidden_0, lstm_0] input_tmp = [hidden_0, lstm_0]
for i in range(1, depth): for i in range(1, depth):
fc = fc_layer(input=input_tmp, size=hidden_dim, param_attr=para_attr) mix_hidden = mixed_layer(name='hidden'+str(i),
size=hidden_dim,
bias_attr=std_default,
input=[full_matrix_projection(input=input_tmp[0], param_attr=hidden_para_attr),
full_matrix_projection(input=input_tmp[1], param_attr=lstm_para_attr)
]
)
lstm = lstmemory(name='lstm'+str(i),
input=mix_hidden,
act=ReluActivation(),
gate_act=SigmoidActivation(),
state_act=SigmoidActivation(),
reverse=((i % 2)==1),
bias_attr=std_0,
param_attr=lstm_para_attr)
input_tmp = [mix_hidden, lstm]
feature_out = mixed_layer(name='output',
size=label_dict_len,
bias_attr=std_default,
input=[full_matrix_projection(input=input_tmp[0], param_attr=hidden_para_attr),
full_matrix_projection(input=input_tmp[1], param_attr=lstm_para_attr)
],
)
lstm = lstmemory(
input=fc,
act=ReluActivation(),
reverse=(i % 2) == 1,
layer_attr=layer_attr)
input_tmp = [fc, lstm]
prob = fc_layer(
input=input_tmp,
size=label_dict_len,
act=SoftmaxActivation(),
param_attr=para_attr)
if not is_predict: if not is_predict:
cls = classification_cost(input=prob, label=target) crf_l = crf_layer( name = 'crf',
outputs(cls) size = label_dict_len,
input = feature_out,
label = target,
param_attr=ParameterAttribute(name='crfw',initial_std=default_std, learning_rate=mix_hidden_lr)
)
crf_dec_l = crf_decoding_layer(name = 'crf_dec_l',
size = label_dict_len,
input = feature_out,
label = target,
param_attr=ParameterAttribute(name='crfw')
)
eval = sum_evaluator(input=crf_dec_l)
outputs(crf_l)
else: else:
outputs(prob) crf_dec_l = crf_decoding_layer(name = 'crf_dec_l',
size = label_dict_len,
input = feature_out,
param_attr=ParameterAttribute(name='crfw')
)
outputs(crf_dec_l)
...@@ -26,7 +26,7 @@ UNK_IDX = 0 ...@@ -26,7 +26,7 @@ UNK_IDX = 0
class Prediction(): class Prediction():
def __init__(self, train_conf, dict_file, model_dir, label_file): def __init__(self, train_conf, dict_file, model_dir, label_file, predicate_dict_file):
""" """
train_conf: trainer configure. train_conf: trainer configure.
dict_file: word dictionary file name. dict_file: word dictionary file name.
...@@ -35,26 +35,41 @@ class Prediction(): ...@@ -35,26 +35,41 @@ class Prediction():
self.dict = {} self.dict = {}
self.labels = {} self.labels = {}
self.predicate_dict={}
self.labels_reverse = {} self.labels_reverse = {}
self.load_dict_label(dict_file, label_file) self.load_dict_label(dict_file, label_file, predicate_dict_file)
len_dict = len(self.dict) len_dict = len(self.dict)
len_label = len(self.labels) len_label = len(self.labels)
len_pred = len(self.predicate_dict)
conf = parse_config(train_conf, 'dict_len=' + str(len_dict) +
',label_len=' + str(len_label) + ',is_predict=True') conf = parse_config(
train_conf,
'dict_len=' + str(len_dict) +
',label_len=' + str(len_label) +
',pred_len=' + str(len_pred) +
',is_predict=True')
self.network = swig_paddle.GradientMachine.createFromConfigProto( self.network = swig_paddle.GradientMachine.createFromConfigProto(
conf.model_config) conf.model_config)
self.network.loadParameters(model_dir) self.network.loadParameters(model_dir)
slots = [ slots = [
integer_value_sequence(len_dict),
integer_value_sequence(len_pred),
integer_value_sequence(len_dict),
integer_value_sequence(len_dict),
integer_value_sequence(len_dict),
integer_value_sequence(len_dict),
integer_value_sequence(len_dict),
integer_value_sequence(2)
]
integer_value_sequence(len_dict), integer_value_sequence(len_dict), integer_value_sequence(len_dict), integer_value_sequence(len_dict),
integer_value_sequence(len_dict), integer_value_sequence(len_dict), integer_value_sequence(len_dict), integer_value_sequence(len_dict),
integer_value_sequence(len_dict), integer_value_sequence(2) integer_value_sequence(len_dict), integer_value_sequence(2)
] ]
self.converter = DataProviderConverter(slots) self.converter = DataProviderConverter(slots)
def load_dict_label(self, dict_file, label_file): def load_dict_label(self, dict_file, label_file, predicate_dict_file):
""" """
Load dictionary from self.dict_file. Load dictionary from self.dict_file.
""" """
...@@ -65,39 +80,42 @@ class Prediction(): ...@@ -65,39 +80,42 @@ class Prediction():
self.labels[line.strip()] = line_count self.labels[line.strip()] = line_count
self.labels_reverse[line_count] = line.strip() self.labels_reverse[line_count] = line.strip()
for line_count, line in enumerate(open(predicate_dict_file, 'r')):
self.predicate_dict[line.strip()] = line_count
def get_data(self, data_file): def get_data(self, data_file):
""" """
Get input data of paddle format. Get input data of paddle format.
""" """
with open(data_file, 'r') as fdata: with open(data_file, 'r') as fdata:
for line in fdata: for line in fdata:
sentence, predicate, ctx_n1, ctx_0, ctx_p1, mark, label = line.strip( sentence, predicate, ctx_n2, ctx_n1, ctx_0, ctx_p1, ctx_p2, mark, label = line.strip(
).split('\t') ).split('\t')
words = sentence.split() words = sentence.split()
sen_len = len(words) sen_len = len(words)
word_slot = [self.dict.get(w, UNK_IDX) for w in words] word_slot = [self.dict.get(w, UNK_IDX) for w in words]
predicate_slot = [self.dict.get(predicate, UNK_IDX)] * sen_len predicate_slot = [self.predicate_dict.get(predicate, UNK_IDX)] * sen_len
ctx_n2_slot = [self.dict.get(ctx_n2, UNK_IDX)] * sen_len
ctx_n1_slot = [self.dict.get(ctx_n1, UNK_IDX)] * sen_len ctx_n1_slot = [self.dict.get(ctx_n1, UNK_IDX)] * sen_len
ctx_0_slot = [self.dict.get(ctx_0, UNK_IDX)] * sen_len ctx_0_slot = [self.dict.get(ctx_0, UNK_IDX)] * sen_len
ctx_p1_slot = [self.dict.get(ctx_p1, UNK_IDX)] * sen_len ctx_p1_slot = [self.dict.get(ctx_p1, UNK_IDX)] * sen_len
ctx_p2_slot = [self.dict.get(ctx_p2, UNK_IDX)] * sen_len
marks = mark.split() marks = mark.split()
mark_slot = [int(w) for w in marks] mark_slot = [int(w) for w in marks]
yield word_slot, predicate_slot, ctx_n2_slot, ctx_n1_slot, \
ctx_0_slot, ctx_p1_slot, ctx_p2_slot, mark_slot
yield word_slot, predicate_slot, ctx_n1_slot, \ def predict(self, data_file, output_file):
ctx_0_slot, ctx_p1_slot, mark_slot
def predict(self, data_file):
""" """
data_file: file name of input data. data_file: file name of input data.
""" """
input = self.converter(self.get_data(data_file)) input = self.converter(self.get_data(data_file))
output = self.network.forwardTest(input) output = self.network.forwardTest(input)
prob = output[0]["value"] lab = output[0]["id"].tolist()
lab = list(np.argsort(-prob)[:, 0])
with open(data_file, 'r') as fin, open('predict.res', 'w') as fout: with open(data_file, 'r') as fin, open(output_file, 'w') as fout:
index = 0 index = 0
for line in fin: for line in fin:
sen = line.split('\t')[0] sen = line.split('\t')[0]
...@@ -109,8 +127,8 @@ class Prediction(): ...@@ -109,8 +127,8 @@ class Prediction():
def option_parser(): def option_parser():
usage = ("python predict.py -c config -w model_dir " usage = ("python predict.py -c config -w model_dir "
"-d word dictionary -l label_file -i input_file") "-d word dictionary -l label_file -i input_file -p pred_dict_file")
parser = OptionParser(usage="usage: %s [options]" % usage) parser = OptionParser(usage="usage: %s [options]" % usage)
parser.add_option( parser.add_option(
"-c", "-c",
...@@ -131,6 +149,13 @@ def option_parser(): ...@@ -131,6 +149,13 @@ def option_parser():
dest="label_file", dest="label_file",
default=None, default=None,
help="label file") help="label file")
parser.add_option(
"-p",
"--predict_dict_file",
action="store",
dest="predict_dict_file",
default=None,
help="predict_dict_file")
parser.add_option( parser.add_option(
"-i", "-i",
"--data", "--data",
...@@ -144,6 +169,14 @@ def option_parser(): ...@@ -144,6 +169,14 @@ def option_parser():
dest="model_path", dest="model_path",
default=None, default=None,
help="model path") help="model path")
parser.add_option(
"-o",
"--output_file",
action="store",
dest="output_file",
default=None,
help="output file")
return parser.parse_args() return parser.parse_args()
...@@ -154,10 +187,12 @@ def main(): ...@@ -154,10 +187,12 @@ def main():
dict_file = options.dict_file dict_file = options.dict_file
model_path = options.model_path model_path = options.model_path
label_file = options.label_file label_file = options.label_file
predict_dict_file = options.predict_dict_file
output_file = options.output_file
swig_paddle.initPaddle("--use_gpu=0") swig_paddle.initPaddle("--use_gpu=0")
predict = Prediction(train_conf, dict_file, model_path, label_file) predict = Prediction(train_conf, dict_file, model_path, label_file, predict_dict_file)
predict.predict(data_file) predict.predict(data_file,output_file)
if __name__ == '__main__': if __name__ == '__main__':
......
...@@ -26,15 +26,18 @@ LOG=`get_best_pass $log` ...@@ -26,15 +26,18 @@ LOG=`get_best_pass $log`
LOG=(${LOG}) LOG=(${LOG})
best_model_path="output/pass-${LOG[1]}" best_model_path="output/pass-${LOG[1]}"
config_file=db_lstm.py config_file=db_lstm.py
dict_file=./data/src.dict dict_file=./data/wordDict.txt
label_file=./data/tgt.dict label_file=./data/targetDict.txt
predicate_dict_file=./data/verbDict.txt
input_file=./data/feature input_file=./data/feature
output_file=predict.res
python predict.py \ python predict.py \
-c $config_file \ -c $config_file \
-w $best_model_path \ -w $best_model_path \
-l $label_file \ -l $label_file \
-p $predicate_dict_file \
-d $dict_file \ -d $dict_file \
-i $input_file -i $input_file \
-o $output_file
...@@ -36,4 +36,5 @@ paddle train \ ...@@ -36,4 +36,5 @@ paddle train \
--job=test \ --job=test \
--use_gpu=false \ --use_gpu=false \
--config_args=is_test=1 \ --config_args=is_test=1 \
--test_all_data_in_one_period=1 \
2>&1 | tee 'test.log' 2>&1 | tee 'test.log'
...@@ -16,11 +16,14 @@ ...@@ -16,11 +16,14 @@
set -e set -e
paddle train \ paddle train \
--config=./db_lstm.py \ --config=./db_lstm.py \
--use_gpu=0 \
--log_period=5000 \
--trainer_count=1 \
--show_parameter_stats_period=5000 \
--save_dir=./output \ --save_dir=./output \
--trainer_count=4 \ --num_passes=10000 \
--log_period=10 \ --average_test_period=10000000 \
--num_passes=500 \ --init_model_path=./data \
--use_gpu=false \ --load_missing_parameter_strategy=rand \
--show_parameter_stats_period=10 \
--test_all_data_in_one_period=1 \ --test_all_data_in_one_period=1 \
2>&1 | tee 'train.log' 2>&1 | tee 'train.log'
# Semantic Role labeling Tutorial # # Semantic Role labeling Tutorial #
Semantic role labeling (SRL) is a form of shallow semantic parsing whose goal is to discover the predicate-argument structure of each predicate in a given input sentence. SRL is useful as an intermediate step in a wide range of natural language processing tasks, such as information extraction. automatic document categorization and question answering. An instance is as following [1]: Semantic role labeling (SRL) is a form of shallow semantic parsing whose goal is to discover the predicate-argument structure of each predicate in a given input sentence. SRL is useful as an intermediate step in a wide range of natural language processing tasks, such as information extraction. automatic document categorization and question answering. An instance is as following [1]:
[ <sub>A0</sub> He ] [ <sub>AM-MOD</sub> would ][ <sub>AM-NEG</sub> n’t ] [ <sub>V</sub> accept] [ <sub>A1</sub> anything of value ] from [<sub>A2</sub> those he was writing about ]. [ <sub>A0</sub> He ] [ <sub>AM-MOD</sub> would ][ <sub>AM-NEG</sub> n’t ] [ <sub>V</sub> accept] [ <sub>A1</sub> anything of value ] from [<sub>A2</sub> those he was writing about ].
- V: verb - V: verb
- A0: acceptor - A0: acceptor
- A1: thing accepted - A1: thing accepted
- A2: accepted-from - A2: accepted-from
- A3: Attribute - A3: Attribute
- AM-MOD: modal - AM-MOD: modal
- AM-NEG: negation - AM-NEG: negation
Given the verb "accept", the chunks in sentence would play certain semantic roles. Here, the label scheme is from Penn Proposition Bank. Given the verb "accept", the chunks in sentence would play certain semantic roles. Here, the label scheme is from Penn Proposition Bank.
To this date, most of the successful SRL systems are built on top of some form of parsing results where pre-defined feature templates over the syntactic structure are used. This tutorial will present an end-to-end system using deep bidirectional long short-term memory (DB-LSTM)[2] for solving the SRL task, which largely outperforms the previous state-of-the-art systems. The system regards SRL task as the sequence labelling problem. To this date, most of the successful SRL systems are built on top of some form of parsing results where pre-defined feature templates over the syntactic structure are used. This tutorial will present an end-to-end system using deep bidirectional long short-term memory (DB-LSTM)[2] for solving the SRL task, which largely outperforms the previous state-of-the-art systems. The system regards SRL task as the sequence labelling problem.
## Data Description ## Data Description
The relevant paper[2] takes the data set in CoNLL-2005&2012 Shared Task for training and testing. Accordingto data license, the demo adopts the test data set of CoNLL-2005, which can be reached on website. The relevant paper[2] takes the data set in CoNLL-2005&2012 Shared Task for training and testing. Accordingto data license, the demo adopts the test data set of CoNLL-2005, which can be reached on website.
To download and process the original data, user just need to execute the following command: To download and process the original data, user just need to execute the following command:
```bash ```bash
cd data cd data
./get_data.sh ./get_data.sh
``` ```
Several new files appear in the `data `directory as follows. Several new files appear in the `data `directory as follows.
```bash ```bash
conll05st-release:the test data set of CoNll-2005 shared task conll05st-release:the test data set of CoNll-2005 shared task
test.wsj.words:the Wall Street Journal data sentences test.wsj.words:the Wall Street Journal data sentences
test.wsj.props: the propositional arguments test.wsj.props: the propositional arguments
src.dict:the dictionary of words in sentences feature: the extracted features from data set
tgt.dict:the labels dictionary ```
feature: the extracted features from data set
``` ## Training
### DB-LSTM
## Training Please refer to the Sentiment Analysis demo to learn more about the long short-term memory unit.
### DB-LSTM
Please refer to the Sentiment Analysis demo to learn more about the long short-term memory unit. Unlike Bidirectional-LSTM that used in Sentiment Analysis demo, the DB-LSTM adopts another way to stack LSTM layer. First a standard LSTM processes the sequence in forward direction. The input and output of this LSTM layer are taken by the next LSTM layer as input, processed in reversed direction. These two standard LSTM layers compose a pair of LSTM. Then we stack LSTM layers pair after pair to obtain the deep LSTM model.
Unlike Bidirectional-LSTM that used in Sentiment Analysis demo, the DB-LSTM adopts another way to stack LSTM layer. First a standard LSTM processes the sequence in forward direction. The input and output of this LSTM layer are taken by the next LSTM layer as input, processed in reversed direction. These two standard LSTM layers compose a pair of LSTM. Then we stack LSTM layers pair after pair to obtain the deep LSTM model. The following figure shows a temporal expanded 2-layer DB-LSTM network.
<center>
The following figure shows a temporal expanded 2-layer DB-LSTM network. ![pic](./network_arch.png)
<center> </center>
![pic](./network_arch.png)
</center> ### Features
Two input features play an essential role in this pipeline: predicate (pred) and argument (argu). Two other features: predicate context (ctx-p) and region mark (mr) are also adopted. Because a single predicate word can not exactly describe the predicate information, especially when the same words appear more than one times in a sentence. With the predicate context, the ambiguity can be largely eliminated. Similarly, we use region mark m<sub>r</sub> = 1 to denote the argument position if it locates in the predicate context region, or m<sub>r</sub> = 0 if does not. These four simple features are all we need for our SRL system. Features of one sample with context size set to 1 is showed as following[2]:
### Features <center>
Two input features play an essential role in this pipeline: predicate (pred) and argument (argu). Two other features: predicate context (ctx-p) and region mark (mr) are also adopted. Because a single predicate word can not exactly describe the predicate information, especially when the same words appear more than one times in a sentence. With the predicate context, the ambiguity can be largely eliminated. Similarly, we use region mark m<sub>r</sub> = 1 to denote the argument position if it locates in the predicate context region, or m<sub>r</sub> = 0 if does not. These four simple features are all we need for our SRL system. Features of one sample with context size set to 1 is showed as following[2]: ![pic](./feature.jpg)
<center> </center>
![pic](./feature.jpg)
</center> In this sample, the coresponding labelled sentence is:
In this sample, the coresponding labelled sentence is: [ <sub>A1</sub> A record date ] has [ <sub>AM-NEG</sub> n't ] been [ <sub>V</sub> set ] .
[ <sub>A1</sub> A record date ] has [ <sub>AM-NEG</sub> n't ] been [ <sub>V</sub> set ] . In the demo, we adopt the feature template as above, consists of : `argument`, `predicate`, `ctx-p (p=-1,0,1)`, `mark` and use `B/I/O` scheme to label each argument. These features and labels are stored in `feature` file, and separated by `\t`.
In the demo, we adopt the feature template as above, consists of : `argument`, `predicate`, `ctx-p (p=-1,0,1)`, `mark` and use `B/I/O` scheme to label each argument. These features and labels are stored in `feature` file, and separated by `\t`. ### Data Provider
### Data Provider `dataprovider.py` is the python file to wrap data. `hook()` function is to define the data slots for network. The Six features and label are all IndexSlots.
```
`dataprovider.py` is the python file to wrap data. `hook()` function is to define the data slots for network. The Six features and label are all IndexSlots. def hook(settings, word_dict, label_dict, **kwargs):
``` settings.word_dict = word_dict
def hook(settings, word_dict, label_dict, **kwargs): settings.label_dict = label_dict
settings.word_dict = word_dict #all inputs are integral and sequential type
settings.label_dict = label_dict settings.slots = [
#all inputs are integral and sequential type integer_value_sequence(len(word_dict)),
settings.slots = [ integer_value_sequence(len(predicate_dict)),
integer_value_sequence(len(word_dict)), integer_value_sequence(len(word_dict)),
integer_value_sequence(len(word_dict)), integer_value_sequence(len(word_dict)),
integer_value_sequence(len(word_dict)), integer_value_sequence(len(word_dict)),
integer_value_sequence(len(word_dict)), integer_value_sequence(len(word_dict)),
integer_value_sequence(len(word_dict)), integer_value_sequence(len(word_dict)),
integer_value_sequence(2), integer_value_sequence(2),
integer_value_sequence(len(label_dict))] integer_value_sequence(len(label_dict))]
``` ```
The corresponding data iterator is as following: The corresponding data iterator is as following:
``` ```
@provider(use_seq=True, init_hook=hook) @provider(init_hook=hook, should_shuffle=True, calc_batch_size=get_batch_size,
def process(obj, file_name): can_over_batch_size=False, cache=CacheType.CACHE_PASS_IN_MEM)
with open(file_name, 'r') as fdata: def process(settings, file_name):
for line in fdata: with open(file_name, 'r') as fdata:
sentence, predicate, ctx_n1, ctx_0, ctx_p1, mark, label = line.strip().split('\t') for line in fdata:
words = sentence.split() sentence, predicate, ctx_n2, ctx_n1, ctx_0, ctx_p1, ctx_p2, mark, label = \
sen_len = len(words) line.strip().split('\t')
word_slot = [obj.word_dict.get(w, UNK_IDX) for w in words]
words = sentence.split()
predicate_slot = [obj.word_dict.get(predicate, UNK_IDX)] * sen_len sen_len = len(words)
ctx_n1_slot = [obj.word_dict.get(ctx_n1, UNK_IDX) ] * sen_len word_slot = [settings.word_dict.get(w, UNK_IDX) for w in words]
ctx_0_slot = [obj.word_dict.get(ctx_0, UNK_IDX) ] * sen_len
ctx_p1_slot = [obj.word_dict.get(ctx_p1, UNK_IDX) ] * sen_len predicate_slot = [settings.predicate_dict.get(predicate)] * sen_len
ctx_n2_slot = [settings.word_dict.get(ctx_n2, UNK_IDX)] * sen_len
marks = mark.split() ctx_n1_slot = [settings.word_dict.get(ctx_n1, UNK_IDX)] * sen_len
mark_slot = [int(w) for w in marks] ctx_0_slot = [settings.word_dict.get(ctx_0, UNK_IDX)] * sen_len
ctx_p1_slot = [settings.word_dict.get(ctx_p1, UNK_IDX)] * sen_len
label_list = label.split() ctx_p2_slot = [settings.word_dict.get(ctx_p2, UNK_IDX)] * sen_len
label_slot = [obj.label_dict.get(w) for w in label_list]
marks = mark.split()
yield word_slot, predicate_slot, ctx_n1_slot, ctx_0_slot, ctx_p1_slot, mark_slot, label_slot mark_slot = [int(w) for w in marks]
```
The `process`function yield 7 lists which are six features and labels. label_list = label.split()
label_slot = [settings.label_dict.get(w) for w in label_list]
### Neural Network Config yield word_slot, predicate_slot, ctx_n2_slot, ctx_n1_slot, \
`db_lstm.py` is the neural network config file to load the dictionaries and define the data provider module and network architecture during the training procedure. ctx_0_slot, ctx_p1_slot, ctx_p2_slot, mark_slot, label_slot
```
Seven `data_layer` load instances from data provider. Six features are transformed into embedddings respectively, and mixed by `mixed_layer` . Deep bidirectional LSTM layers extract features for the softmax layer. The objective function is cross entropy of labels. The `process`function yield 9 lists which are 8 features and label.
### Run Training ### Neural Network Config
The script for training is `train.sh`, user just need to execute: `db_lstm.py` is the neural network config file to load the dictionaries and define the data provider module and network architecture during the training procedure.
```bash
./train.sh Nine `data_layer` load instances from data provider. Eight features are transformed into embedddings respectively, and mixed by `mixed_layer` . Deep bidirectional LSTM layers extract features for the softmax layer. The objective function is cross entropy of labels.
```
The content in `train.sh`: ### Run Training
``` The script for training is `train.sh`, user just need to execute:
paddle train \ ```bash
--config=./db_lstm.py \ ./train.sh
--save_dir=./output \ ```
--trainer_count=4 \ The content in `train.sh`:
--log_period=10 \ ```
--num_passes=500 \ paddle train \
--use_gpu=false \ --config=./db_lstm.py \
--show_parameter_stats_period=10 \ --use_gpu=0 \
--test_all_data_in_one_period=1 \ --log_period=5000 \
2>&1 | tee 'train.log' --trainer_count=1 \
``` --show_parameter_stats_period=5000 \
--save_dir=./output \
- \--config=./db_lstm.py : network config file. --num_passes=10000 \
- \--save_di=./output: output path to save models. --average_test_period=10000000 \
- \--trainer_count=4 : set thread number (or GPU count). --init_model_path=./data \
- \--log_period=10 : print log every 20 batches. --load_missing_parameter_strategy=rand \
- \--num_passes=500: set pass number, one pass in PaddlePaddle means training all samples in dataset one time. --test_all_data_in_one_period=1 \
- \--use_gpu=false: use CPU to train, set true, if you install GPU version of PaddlePaddle and want to use GPU to train. 2>&1 | tee 'train.log'
- \--show_parameter_stats_period=10: show parameter statistic every 100 batches. ```
- \--test_all_data_in_one_period=1: test all data in every testing.
- \--config=./db_lstm.py : network config file.
- \--use_gpu=false: use CPU to train, set true, if you install GPU version of PaddlePaddle and want to use GPU to train, until now crf_layer do not support GPU
After training, the models will be saved in directory `output`. - \--log_period=500: print log every 20 batches.
- \--trainer_count=1: set thread number (or GPU count).
### Run testing - \--show_parameter_stats_period=5000: show parameter statistic every 100 batches.
The script for testing is `test.sh`, user just need to execute: - \--save_dir=./output: output path to save models.
```bash - \--num_passes=10000: set pass number, one pass in PaddlePaddle means training all samples in dataset one time.
./test.sh - \--average_test_period=10000000: do test on average parameter every average_test_period batches
``` - \--init_model_path=./data: parameter initialization path
The main part in `tesh.sh` - \--load_missing_parameter_strategy=rand: random initialization unexisted parameters
``` - \--test_all_data_in_one_period=1: test all data in one period
paddle train \
--config=./db_lstm.py \
--model_list=$model_list \ After training, the models will be saved in directory `output`. Our training curve is as following:
--job=test \ <center>
--config_args=is_test=1 \ ![pic](./curve.jpg)
``` </center>
- \--config=./db_lstm.py: network config file ### Run testing
- \--model_list=$model_list.list: model list file The script for testing is `test.sh`, user just need to execute:
- \--job=test: indicate the test job ```bash
- \--config_args=is_test=1: flag to indicate test ./test.sh
```
The main part in `tesh.sh`
### Run prediction ```
The script for prediction is `predict.sh`, user just need to execute: paddle train \
```bash --config=./db_lstm.py \
./predict.sh --model_list=$model_list \
--job=test \
``` --config_args=is_test=1 \
In `predict.sh`, user should offer the network config file, model path, label file, word dictionary file, feature file ```
```
python predict.py - \--config=./db_lstm.py: network config file
-c $config_file - \--model_list=$model_list.list: model list file
-w $model_path - \--job=test: indicate the test job
-l $label_file - \--config_args=is_test=1: flag to indicate test
-d $dict_file - \--test_all_data_in_one_period=1: test all data in 1 period
-i $input_file
```
### Run prediction
`predict.py` is the main executable python script, which includes functions: load model, load data, data prediction. The network model will output the probability distribution of labels. In the demo, we take the label with maximum probability as result. User can also implement the beam search or viterbi decoding upon the probability distribution matrix. The script for prediction is `predict.sh`, user just need to execute:
```bash
After prediction, the result is saved in `predict.res`. ./predict.sh
## Reference ```
[1] Martha Palmer, Dan Gildea, and Paul Kingsbury. The Proposition Bank: An Annotated Corpus of Semantic Roles , Computational Linguistics, 31(1), 2005. In `predict.sh`, user should offer the network config file, model path, label file, word dictionary file, feature file
```
[2] Zhou, Jie, and Wei Xu. "End-to-end learning of semantic role labeling using recurrent neural networks." Proceedings of the Annual Meeting of the Association for Computational Linguistics. 2015. python predict.py
-c $config_file \
-w $best_model_path \
-l $label_file \
-p $predicate_dict_file \
-d $dict_file \
-i $input_file \
-o $output_file
```
`predict.py` is the main executable python script, which includes functions: load model, load data, data prediction. The network model will output the probability distribution of labels. In the demo, we take the label with maximum probability as result. User can also implement the beam search or viterbi decoding upon the probability distribution matrix.
After prediction, the result is saved in `predict.res`.
## Reference
[1] Martha Palmer, Dan Gildea, and Paul Kingsbury. The Proposition Bank: An Annotated Corpus of Semantic Roles , Computational Linguistics, 31(1), 2005.
[2] Zhou, Jie, and Wei Xu. "End-to-end learning of semantic role labeling using recurrent neural networks." Proceedings of the Annual Meeting of the Association for Computational Linguistics. 2015.
...@@ -12,7 +12,6 @@ WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. ...@@ -12,7 +12,6 @@ WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
See the License for the specific language governing permissions and See the License for the specific language governing permissions and
limitations under the License. */ limitations under the License. */
#include "PaddleAPI.h" #include "PaddleAPI.h"
#include "PaddleAPIPrivate.h" #include "PaddleAPIPrivate.h"
...@@ -112,7 +111,7 @@ void Arguments::setSlotSequenceStartPositions(size_t idx, ...@@ -112,7 +111,7 @@ void Arguments::setSlotSequenceStartPositions(size_t idx,
} }
void Arguments::setSlotSubSequenceStartPositions( void Arguments::setSlotSubSequenceStartPositions(
size_t idx, IVector *vec) throw(RangeError) { size_t idx, IVector* vec) throw(RangeError) {
auto& a = m->getArg(idx); auto& a = m->getArg(idx);
auto& v = m->cast<paddle::IVector>(vec->getSharedPtr()); auto& v = m->cast<paddle::IVector>(vec->getSharedPtr());
a.subSequenceStartPositions = std::make_shared<paddle::ICpuGpuVector>(v); a.subSequenceStartPositions = std::make_shared<paddle::ICpuGpuVector>(v);
......
...@@ -12,7 +12,6 @@ WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. ...@@ -12,7 +12,6 @@ WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
See the License for the specific language governing permissions and See the License for the specific language governing permissions and
limitations under the License. */ limitations under the License. */
#include "PaddleAPI.h" #include "PaddleAPI.h"
#include "PaddleAPIPrivate.h" #include "PaddleAPIPrivate.h"
#include "paddle/trainer/Trainer.h" #include "paddle/trainer/Trainer.h"
...@@ -44,8 +43,7 @@ TrainerConfig* TrainerConfig::createFromTrainerConfigFile( ...@@ -44,8 +43,7 @@ TrainerConfig* TrainerConfig::createFromTrainerConfigFile(
return retv; return retv;
} }
TrainerConfig* TrainerConfig::createFromProtoString( TrainerConfig* TrainerConfig::createFromProtoString(const std::string& str) {
const std::string& str) {
auto retv = new TrainerConfig(); auto retv = new TrainerConfig();
paddle::TrainerConfig trainerConfigProto; paddle::TrainerConfig trainerConfigProto;
auto conf = std::make_shared<paddle::TrainerConfigHelper>(trainerConfigProto); auto conf = std::make_shared<paddle::TrainerConfigHelper>(trainerConfigProto);
......
...@@ -12,7 +12,6 @@ WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. ...@@ -12,7 +12,6 @@ WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
See the License for the specific language governing permissions and See the License for the specific language governing permissions and
limitations under the License. */ limitations under the License. */
#include "PaddleAPI.h" #include "PaddleAPI.h"
#include "PaddleAPIPrivate.h" #include "PaddleAPIPrivate.h"
...@@ -27,7 +26,8 @@ GradientMachine::GradientMachine() : m(new GradientMachinePrivate()) {} ...@@ -27,7 +26,8 @@ GradientMachine::GradientMachine() : m(new GradientMachinePrivate()) {}
GradientMachine::~GradientMachine() { delete m; } GradientMachine::~GradientMachine() { delete m; }
GradientMachine* GradientMachine::createFromPaddleModelPtr( GradientMachine* GradientMachine::createFromPaddleModelPtr(
const void* confPtr, GradientMatchineCreateMode mode, const void* confPtr,
GradientMatchineCreateMode mode,
const std::vector<int>& types) { const std::vector<int>& types) {
auto& conf = *(const paddle::ModelConfig*)(confPtr); auto& conf = *(const paddle::ModelConfig*)(confPtr);
std::vector<ParameterType> realTypes; std::vector<ParameterType> realTypes;
...@@ -44,7 +44,8 @@ GradientMachine* GradientMachine::createFromPaddleModelPtr( ...@@ -44,7 +44,8 @@ GradientMachine* GradientMachine::createFromPaddleModelPtr(
} }
GradientMachine* GradientMachine::createByConfigProtoStr( GradientMachine* GradientMachine::createByConfigProtoStr(
const std::string& protoStr, GradientMatchineCreateMode mode, const std::string& protoStr,
GradientMatchineCreateMode mode,
const std::vector<int>& types) { const std::vector<int>& types) {
paddle::ModelConfig conf; paddle::ModelConfig conf;
conf.ParseFromString(protoStr); conf.ParseFromString(protoStr);
...@@ -56,13 +57,15 @@ GradientMachine* GradientMachine::createByConfigProtoStr( ...@@ -56,13 +57,15 @@ GradientMachine* GradientMachine::createByConfigProtoStr(
} }
GradientMachine* GradientMachine::createByModelConfig( GradientMachine* GradientMachine::createByModelConfig(
ModelConfig* conf, GradientMatchineCreateMode mode, ModelConfig* conf,
GradientMatchineCreateMode mode,
const std::vector<int>& types) { const std::vector<int>& types) {
auto confPtr = &conf->m->conf->getModelConfig(); auto confPtr = &conf->m->conf->getModelConfig();
return GradientMachine::createFromPaddleModelPtr(confPtr, mode, types); return GradientMachine::createFromPaddleModelPtr(confPtr, mode, types);
} }
void GradientMachine::forward(const Arguments& inArgs, Arguments* outArgs, void GradientMachine::forward(const Arguments& inArgs,
Arguments* outArgs,
PassType passType) { PassType passType) {
auto& in = auto& in =
m->cast<std::vector<paddle::Argument>>(inArgs.getInternalArgumentsPtr()); m->cast<std::vector<paddle::Argument>>(inArgs.getInternalArgumentsPtr());
...@@ -99,7 +102,8 @@ void GradientMachine::backward(const UpdateCallback& callback) { ...@@ -99,7 +102,8 @@ void GradientMachine::backward(const UpdateCallback& callback) {
} }
void GradientMachine::forwardBackward(const Arguments& inArgs, void GradientMachine::forwardBackward(const Arguments& inArgs,
Arguments* outArgs, PassType passType, Arguments* outArgs,
PassType passType,
const UpdateCallback& callback) { const UpdateCallback& callback) {
auto& in = auto& in =
m->cast<std::vector<paddle::Argument>>(inArgs.getInternalArgumentsPtr()); m->cast<std::vector<paddle::Argument>>(inArgs.getInternalArgumentsPtr());
...@@ -129,7 +133,7 @@ Parameter* GradientMachine::getParameter(size_t i) throw(RangeError) { ...@@ -129,7 +133,7 @@ Parameter* GradientMachine::getParameter(size_t i) throw(RangeError) {
void GradientMachine::randParameters() { m->machine->randParameters(); } void GradientMachine::randParameters() { m->machine->randParameters(); }
Matrix* GradientMachine::getLayerOutput(const std::string& layerName) const Matrix* GradientMachine::getLayerOutput(const std::string& layerName) const
throw(UnsupportError) { throw(UnsupportError) {
auto nn = std::dynamic_pointer_cast<paddle::NeuralNetwork>(m->machine); auto nn = std::dynamic_pointer_cast<paddle::NeuralNetwork>(m->machine);
if (nn) { if (nn) {
auto mat = nn->getLayerOutput(layerName); auto mat = nn->getLayerOutput(layerName);
...@@ -140,8 +144,11 @@ Matrix* GradientMachine::getLayerOutput(const std::string& layerName) const ...@@ -140,8 +144,11 @@ Matrix* GradientMachine::getLayerOutput(const std::string& layerName) const
} }
SequenceGenerator* GradientMachine::asSequenceGenerator( SequenceGenerator* GradientMachine::asSequenceGenerator(
const std::vector<std::string>& dict, size_t begin_id, size_t end_id, const std::vector<std::string>& dict,
size_t max_length, size_t beam_size) { size_t begin_id,
size_t end_id,
size_t max_length,
size_t beam_size) {
SequenceGenerator* r = SequenceGenerator* r =
SequenceGenerator::createByGradientMachineSharedPtr(&m->machine); SequenceGenerator::createByGradientMachineSharedPtr(&m->machine);
r->setDict(dict); r->setDict(dict);
......
...@@ -12,7 +12,6 @@ WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. ...@@ -12,7 +12,6 @@ WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
See the License for the specific language governing permissions and See the License for the specific language governing permissions and
limitations under the License. */ limitations under the License. */
#pragma once #pragma once
#include "PaddleAPI.h" #include "PaddleAPI.h"
...@@ -23,7 +22,8 @@ limitations under the License. */ ...@@ -23,7 +22,8 @@ limitations under the License. */
template <typename T1, typename T2> template <typename T1, typename T2>
void staticCastVector(std::vector<T2>* dest, const std::vector<T1>& src) { void staticCastVector(std::vector<T2>* dest, const std::vector<T1>& src) {
dest->resize(src.size()); dest->resize(src.size());
std::transform(src.begin(), src.end(), dest->begin(), [](T1 t){ std::transform(src.begin(),
return static_cast<T2>(t); src.end(),
}); dest->begin(),
[](T1 t) { return static_cast<T2>(t); });
} }
...@@ -12,7 +12,6 @@ WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. ...@@ -12,7 +12,6 @@ WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
See the License for the specific language governing permissions and See the License for the specific language governing permissions and
limitations under the License. */ limitations under the License. */
#include "PaddleAPI.h" #include "PaddleAPI.h"
#include "paddle/math/Matrix.h" #include "paddle/math/Matrix.h"
#include "paddle/math/SparseMatrix.h" #include "paddle/math/SparseMatrix.h"
...@@ -44,17 +43,21 @@ Matrix* Matrix::createZero(size_t height, size_t width, bool useGpu) { ...@@ -44,17 +43,21 @@ Matrix* Matrix::createZero(size_t height, size_t width, bool useGpu) {
return m; return m;
} }
Matrix* Matrix::createDense(const std::vector<float>& data, size_t height, Matrix* Matrix::createDense(const std::vector<float>& data,
size_t width, bool useGpu) { size_t height,
size_t width,
bool useGpu) {
auto m = new Matrix(); auto m = new Matrix();
m->m->mat = paddle::Matrix::create(height, width, useGpu); m->m->mat = paddle::Matrix::create(height, width, useGpu);
m->m->mat->copyFrom(data.data(), data.size()); m->m->mat->copyFrom(data.data(), data.size());
return m; return m;
} }
Matrix* Matrix::createDenseFromNumpy(float* data, int dim1, int dim2, Matrix* Matrix::createDenseFromNumpy(float* data,
bool copy, bool useGpu) int dim1,
throw (UnsupportError) { int dim2,
bool copy,
bool useGpu) throw(UnsupportError) {
if (useGpu) { if (useGpu) {
/// Gpu mode only supports copy=True /// Gpu mode only supports copy=True
if (!copy) { if (!copy) {
...@@ -66,7 +69,9 @@ Matrix* Matrix::createDenseFromNumpy(float* data, int dim1, int dim2, ...@@ -66,7 +69,9 @@ Matrix* Matrix::createDenseFromNumpy(float* data, int dim1, int dim2,
} }
} }
Matrix* Matrix::createCpuDenseFromNumpy(float* data, int dim1, int dim2, Matrix* Matrix::createCpuDenseFromNumpy(float* data,
int dim1,
int dim2,
bool copy) { bool copy) {
auto m = new Matrix(); auto m = new Matrix();
if (copy) { if (copy) {
...@@ -85,12 +90,20 @@ Matrix* Matrix::createGpuDenseFromNumpy(float* data, int dim1, int dim2) { ...@@ -85,12 +90,20 @@ Matrix* Matrix::createGpuDenseFromNumpy(float* data, int dim1, int dim2) {
return m; return m;
} }
Matrix* Matrix::createSparse(size_t height, size_t width, size_t nnz, Matrix* Matrix::createSparse(size_t height,
bool isNonVal, bool isTrans, bool useGpu) { size_t width,
size_t nnz,
bool isNonVal,
bool isTrans,
bool useGpu) {
auto m = new Matrix(); auto m = new Matrix();
m->m->mat = paddle::Matrix::createSparseMatrix( m->m->mat = paddle::Matrix::createSparseMatrix(
height, width, nnz, isNonVal ? paddle::NO_VALUE : paddle::FLOAT_VALUE, height,
isTrans, useGpu); width,
nnz,
isNonVal ? paddle::NO_VALUE : paddle::FLOAT_VALUE,
isTrans,
useGpu);
return m; return m;
} }
...@@ -221,7 +234,8 @@ FloatArray Matrix::getData() const { ...@@ -221,7 +234,8 @@ FloatArray Matrix::getData() const {
} }
void Matrix::sparseCopyFrom( void Matrix::sparseCopyFrom(
const std::vector<int>& rows, const std::vector<int>& cols, const std::vector<int>& rows,
const std::vector<int>& cols,
const std::vector<float>& vals) throw(UnsupportError) { const std::vector<float>& vals) throw(UnsupportError) {
auto cpuSparseMat = auto cpuSparseMat =
std::dynamic_pointer_cast<paddle::CpuSparseMatrix>(m->mat); std::dynamic_pointer_cast<paddle::CpuSparseMatrix>(m->mat);
...@@ -240,7 +254,8 @@ void Matrix::sparseCopyFrom( ...@@ -240,7 +254,8 @@ void Matrix::sparseCopyFrom(
void* Matrix::getSharedPtr() const { return &m->mat; } void* Matrix::getSharedPtr() const { return &m->mat; }
void Matrix::toNumpyMatInplace(float** view_data, int* dim1, void Matrix::toNumpyMatInplace(float** view_data,
int* dim1,
int* dim2) throw(UnsupportError) { int* dim2) throw(UnsupportError) {
auto cpuMat = std::dynamic_pointer_cast<paddle::CpuMatrix>(m->mat); auto cpuMat = std::dynamic_pointer_cast<paddle::CpuMatrix>(m->mat);
if (cpuMat) { if (cpuMat) {
...@@ -251,7 +266,8 @@ void Matrix::toNumpyMatInplace(float** view_data, int* dim1, ...@@ -251,7 +266,8 @@ void Matrix::toNumpyMatInplace(float** view_data, int* dim1,
throw UnsupportError(); throw UnsupportError();
} }
} }
void Matrix::copyToNumpyMat(float** view_m_data, int* dim1, void Matrix::copyToNumpyMat(float** view_m_data,
int* dim1,
int* dim2) throw(UnsupportError) { int* dim2) throw(UnsupportError) {
static_assert(sizeof(paddle::real) == sizeof(float), static_assert(sizeof(paddle::real) == sizeof(float),
"Currently PaddleAPI only support for single " "Currently PaddleAPI only support for single "
...@@ -269,8 +285,8 @@ void Matrix::copyToNumpyMat(float** view_m_data, int* dim1, ...@@ -269,8 +285,8 @@ void Matrix::copyToNumpyMat(float** view_m_data, int* dim1,
} else if (auto gpuMat = dynamic_cast<paddle::GpuMatrix*>(m->mat.get())) { } else if (auto gpuMat = dynamic_cast<paddle::GpuMatrix*>(m->mat.get())) {
auto src = gpuMat->getData(); auto src = gpuMat->getData();
auto dest = *view_m_data; auto dest = *view_m_data;
hl_memcpy_device2host(dest, src, hl_memcpy_device2host(
sizeof(paddle::real) * (*dim1) * (*dim2)); dest, src, sizeof(paddle::real) * (*dim1) * (*dim2));
} else { } else {
LOG(WARNING) << "Unexpected Situation"; LOG(WARNING) << "Unexpected Situation";
throw UnsupportError(); throw UnsupportError();
...@@ -278,7 +294,8 @@ void Matrix::copyToNumpyMat(float** view_m_data, int* dim1, ...@@ -278,7 +294,8 @@ void Matrix::copyToNumpyMat(float** view_m_data, int* dim1,
} }
} }
void Matrix::copyFromNumpyMat(float* data, int dim1, void Matrix::copyFromNumpyMat(float* data,
int dim1,
int dim2) throw(UnsupportError, RangeError) { int dim2) throw(UnsupportError, RangeError) {
if (isSparse()) { if (isSparse()) {
throw UnsupportError(); throw UnsupportError();
......
...@@ -12,7 +12,6 @@ WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. ...@@ -12,7 +12,6 @@ WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
See the License for the specific language governing permissions and See the License for the specific language governing permissions and
limitations under the License. */ limitations under the License. */
#pragma once #pragma once
#include <stddef.h> #include <stddef.h>
...@@ -61,8 +60,8 @@ class RangeError {}; ...@@ -61,8 +60,8 @@ class RangeError {};
/// Not support Error, such as access GPU memory directly, etc. /// Not support Error, such as access GPU memory directly, etc.
class UnsupportError : public std::runtime_error { class UnsupportError : public std::runtime_error {
public: public:
UnsupportError() : std::runtime_error(" ") {}; UnsupportError() : std::runtime_error(" "){};
UnsupportError(const std::string& message) : std::runtime_error(message) {}; UnsupportError(const std::string& message) : std::runtime_error(message){};
}; };
/// This type will map to python's list of float. /// This type will map to python's list of float.
...@@ -112,7 +111,8 @@ public: ...@@ -112,7 +111,8 @@ public:
/** /**
* Create A Matrix with height,width, which is filled by zero. * Create A Matrix with height,width, which is filled by zero.
*/ */
static Matrix* createZero(size_t height, size_t width, static Matrix* createZero(size_t height,
size_t width,
bool useGpu = isUsingGpu()); bool useGpu = isUsingGpu());
/** /**
...@@ -124,8 +124,11 @@ public: ...@@ -124,8 +124,11 @@ public:
* *
* @note the default sparse type is SPARSE_CSR. * @note the default sparse type is SPARSE_CSR.
*/ */
static Matrix* createSparse(size_t height, size_t width, size_t nnz, static Matrix* createSparse(size_t height,
bool isNonVal = true, bool trans = false, size_t width,
size_t nnz,
bool isNonVal = true,
bool trans = false,
bool useGpu = isUsingGpu()); bool useGpu = isUsingGpu());
/** /**
...@@ -134,13 +137,17 @@ public: ...@@ -134,13 +137,17 @@ public:
* @param data list of float should be passed in python. * @param data list of float should be passed in python.
* @note the value will be copy into a new matrix. * @note the value will be copy into a new matrix.
*/ */
static Matrix* createDense(const std::vector<float>& data, size_t height, static Matrix* createDense(const std::vector<float>& data,
size_t width, bool useGpu = isUsingGpu()); size_t height,
size_t width,
static Matrix* createDenseFromNumpy(float* data, int dim1, int dim2, bool useGpu = isUsingGpu());
bool copy = true,
bool useGpu = isUsingGpu()) static Matrix* createDenseFromNumpy(
throw (UnsupportError); float* data,
int dim1,
int dim2,
bool copy = true,
bool useGpu = isUsingGpu()) throw(UnsupportError);
/** /**
* Create Cpu Dense Matrix from numpy matrix, dtype=float32 * Create Cpu Dense Matrix from numpy matrix, dtype=float32
...@@ -151,7 +158,9 @@ public: ...@@ -151,7 +158,9 @@ public:
* @param copy true if copy into a new matrix, false will create * @param copy true if copy into a new matrix, false will create
* matrix inplace. * matrix inplace.
*/ */
static Matrix* createCpuDenseFromNumpy(float* data, int dim1, int dim2, static Matrix* createCpuDenseFromNumpy(float* data,
int dim1,
int dim2,
bool copy = false); bool copy = false);
/// Create Gpu Dense Matrix from numpy matrix, dtype=float32 /// Create Gpu Dense Matrix from numpy matrix, dtype=float32
...@@ -171,11 +180,13 @@ public: ...@@ -171,11 +180,13 @@ public:
* numpy_mat = m.toNumpyMat() * numpy_mat = m.toNumpyMat()
* @endcode * @endcode
*/ */
void toNumpyMatInplace(float** view_data, int* dim1, void toNumpyMatInplace(float** view_data,
int* dim1,
int* dim2) throw(UnsupportError); int* dim2) throw(UnsupportError);
/// Copy To numpy mat. /// Copy To numpy mat.
void copyToNumpyMat(float** view_m_data, int* dim1, void copyToNumpyMat(float** view_m_data,
int* dim1,
int* dim2) throw(UnsupportError); int* dim2) throw(UnsupportError);
/// Copy From Numpy Mat /// Copy From Numpy Mat
...@@ -248,15 +259,18 @@ public: ...@@ -248,15 +259,18 @@ public:
static Vector* create(const std::vector<float>& data, static Vector* create(const std::vector<float>& data,
bool useGpu = isUsingGpu()); bool useGpu = isUsingGpu());
static Vector* createVectorFromNumpy(float* data, int dim, bool copy = true, static Vector* createVectorFromNumpy(
bool useGpu = isUsingGpu()) float* data,
throw (UnsupportError); int dim,
bool copy = true,
bool useGpu = isUsingGpu()) throw(UnsupportError);
/** /**
* Create Cpu Vector from numpy array, which dtype=float32 * Create Cpu Vector from numpy array, which dtype=float32
* *
* If copy is false, it will create vector inplace. * If copy is false, it will create vector inplace.
*/ */
static Vector* createCpuVectorFromNumpy(float* data, int dim, static Vector* createCpuVectorFromNumpy(float* data,
int dim,
bool copy = false); bool copy = false);
/// Create Gpu Vector from numpy array, which dtype=float32 /// Create Gpu Vector from numpy array, which dtype=float32
...@@ -312,16 +326,19 @@ public: ...@@ -312,16 +326,19 @@ public:
static IVector* create(const std::vector<int>& data, static IVector* create(const std::vector<int>& data,
bool useGpu = isUsingGpu()); bool useGpu = isUsingGpu());
static IVector* createVectorFromNumpy(int* data, int dim, bool copy = true, static IVector* createVectorFromNumpy(
bool useGpu = isUsingGpu()) int* data,
throw (UnsupportError); int dim,
bool copy = true,
bool useGpu = isUsingGpu()) throw(UnsupportError);
/** /**
* Create Cpu IVector from numpy array, which dtype=int32 * Create Cpu IVector from numpy array, which dtype=int32
* *
* If copy is false, it will create vector inplace * If copy is false, it will create vector inplace
*/ */
static IVector* createCpuVectorFromNumpy(int* data, int dim, static IVector* createCpuVectorFromNumpy(int* data,
int dim,
bool copy = false); bool copy = false);
/** /**
* Create Gpu IVector from numpy array, which dtype=int32 * Create Gpu IVector from numpy array, which dtype=int32
...@@ -605,7 +622,8 @@ class ParameterTraverseCallback { ...@@ -605,7 +622,8 @@ class ParameterTraverseCallback {
public: public:
~ParameterTraverseCallback(); ~ParameterTraverseCallback();
void apply(const std::vector<Vector*>& vecs, const ParameterConfig& config, void apply(const std::vector<Vector*>& vecs,
const ParameterConfig& config,
size_t sparseId); size_t sparseId);
private: private:
...@@ -638,7 +656,8 @@ public: ...@@ -638,7 +656,8 @@ public:
void finishBatch(); void finishBatch();
void update(const std::vector<Vector*>& vecs, const ParameterConfig& conf, void update(const std::vector<Vector*>& vecs,
const ParameterConfig& conf,
size_t sparseId = NO_SPARSE_ID); size_t sparseId = NO_SPARSE_ID);
std::vector<int> getParameterTypes() const; std::vector<int> getParameterTypes() const;
...@@ -678,7 +697,8 @@ public: ...@@ -678,7 +697,8 @@ public:
* model config by TrainerConfig * model config by TrainerConfig
*/ */
static GradientMachine* createByModelConfig( static GradientMachine* createByModelConfig(
ModelConfig* conf, GradientMatchineCreateMode mode = CREATE_MODE_NORMAL, ModelConfig* conf,
GradientMatchineCreateMode mode = CREATE_MODE_NORMAL,
const std::vector<int>& parameterTypes = defaultParamTypes); const std::vector<int>& parameterTypes = defaultParamTypes);
/** /**
...@@ -701,7 +721,8 @@ public: ...@@ -701,7 +721,8 @@ public:
/** /**
* Combine forward/backward * Combine forward/backward
*/ */
void forwardBackward(const Arguments& inArgs, Arguments* outArgs, void forwardBackward(const Arguments& inArgs,
Arguments* outArgs,
PassType passType, PassType passType,
const UpdateCallback& callback = UpdateCallback()); const UpdateCallback& callback = UpdateCallback());
...@@ -722,14 +743,17 @@ public: ...@@ -722,14 +743,17 @@ public:
*/ */
SequenceGenerator* asSequenceGenerator( SequenceGenerator* asSequenceGenerator(
const std::vector<std::string>& dict = std::vector<std::string>(), const std::vector<std::string>& dict = std::vector<std::string>(),
size_t begin_id = 0UL, size_t end_id = 0UL, size_t max_length = 100UL, size_t begin_id = 0UL,
size_t end_id = 0UL,
size_t max_length = 100UL,
size_t beam_size = -1UL); size_t beam_size = -1UL);
private: private:
GradientMachinePrivate* m; GradientMachinePrivate* m;
static GradientMachine* createFromPaddleModelPtr( static GradientMachine* createFromPaddleModelPtr(
const void* confPtr, GradientMatchineCreateMode mode, const void* confPtr,
GradientMatchineCreateMode mode,
const std::vector<int>& types); const std::vector<int>& types);
// Not to use c++ 11 init-list, so we use static var as function default arg. // Not to use c++ 11 init-list, so we use static var as function default arg.
...@@ -751,8 +775,8 @@ public: ...@@ -751,8 +775,8 @@ public:
/// Create A Trainer By TrainerConfig. using paddle command line. /// Create A Trainer By TrainerConfig. using paddle command line.
static Trainer* createByCommandLine() throw(IOError); static Trainer* createByCommandLine() throw(IOError);
static Trainer* create(TrainerConfig* optConfig, GradientMachine* gm) static Trainer* create(TrainerConfig* optConfig,
throw(IOError); GradientMachine* gm) throw(IOError);
/// Start training /// Start training
void startTrain(); void startTrain();
......
...@@ -12,7 +12,6 @@ WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. ...@@ -12,7 +12,6 @@ WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
See the License for the specific language governing permissions and See the License for the specific language governing permissions and
limitations under the License. */ limitations under the License. */
#include "PaddleAPI.h" #include "PaddleAPI.h"
#include "paddle/parameter/Parameter.h" #include "paddle/parameter/Parameter.h"
......
...@@ -12,7 +12,6 @@ WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. ...@@ -12,7 +12,6 @@ WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
See the License for the specific language governing permissions and See the License for the specific language governing permissions and
limitations under the License. */ limitations under the License. */
#include "PaddleAPI.h" #include "PaddleAPI.h"
#include "PaddleAPIPrivate.h" #include "PaddleAPIPrivate.h"
#include "paddle/parameter/ParameterOptimizer.h" #include "paddle/parameter/ParameterOptimizer.h"
...@@ -32,17 +31,21 @@ struct ParameterTraverseCallbackPrivate { ...@@ -32,17 +31,21 @@ struct ParameterTraverseCallbackPrivate {
const paddle::ParameterOptimizer::TraverseCallback& callback) const paddle::ParameterOptimizer::TraverseCallback& callback)
: callback(callback) {} : callback(callback) {}
void apply(const std::vector<Vector*>& vecs, const ParameterConfig& conf, void apply(const std::vector<Vector*>& vecs,
const ParameterConfig& conf,
size_t sparseId) { size_t sparseId) {
std::vector<paddle::VectorPtr> real_vecs; std::vector<paddle::VectorPtr> real_vecs;
real_vecs.resize(vecs.size()); real_vecs.resize(vecs.size());
std::transform(vecs.begin(), vecs.end(), real_vecs.begin(), [](Vector* v) { std::transform(vecs.begin(),
if (v) { vecs.end(),
return *(paddle::VectorPtr*)(v->getSharedPtr()); real_vecs.begin(),
} else { [](Vector* v) {
return paddle::VectorPtr(); if (v) {
} return *(paddle::VectorPtr*)(v->getSharedPtr());
}); } else {
return paddle::VectorPtr();
}
});
paddle::ParameterConfig& real_conf = paddle::ParameterConfig& real_conf =
*(paddle::ParameterConfig*)(const_cast<ParameterConfig&>(conf) *(paddle::ParameterConfig*)(const_cast<ParameterConfig&>(conf)
...@@ -86,10 +89,12 @@ void ParameterOptimizer::startBatch(size_t numSamplesProcessed) { ...@@ -86,10 +89,12 @@ void ParameterOptimizer::startBatch(size_t numSamplesProcessed) {
void ParameterOptimizer::finishBatch() { m->optimizer->finishBatch(); } void ParameterOptimizer::finishBatch() { m->optimizer->finishBatch(); }
void ParameterOptimizer::update(const std::vector<Vector*>& vecs, void ParameterOptimizer::update(const std::vector<Vector*>& vecs,
const ParameterConfig& conf, size_t sparseId) { const ParameterConfig& conf,
ParameterTraverseCallbackPrivate invoker([&]( size_t sparseId) {
const paddle::VectorPtr _vecs[], const paddle::ParameterConfig& config, ParameterTraverseCallbackPrivate invoker(
size_t sid = -1UL) { m->optimizer->update(_vecs, config, sid); }); [&](const paddle::VectorPtr _vecs[],
const paddle::ParameterConfig& config,
size_t sid = -1UL) { m->optimizer->update(_vecs, config, sid); });
invoker.apply(vecs, conf, sparseId); invoker.apply(vecs, conf, sparseId);
} }
...@@ -116,8 +121,9 @@ void ParameterTraverseCallback::apply(const std::vector<Vector*>& vecs, ...@@ -116,8 +121,9 @@ void ParameterTraverseCallback::apply(const std::vector<Vector*>& vecs,
ParameterTraverseCallback* ParameterOptimizer::needSpecialTraversal( ParameterTraverseCallback* ParameterOptimizer::needSpecialTraversal(
const ParameterConfig& config) const { const ParameterConfig& config) const {
auto& param_config = *(paddle::ParameterConfig*)const_cast<ParameterConfig&>( auto& param_config =
config).getRawPtr(); *(paddle::ParameterConfig*)const_cast<ParameterConfig&>(config)
.getRawPtr();
auto callback = m->optimizer->needSpecialTraversal(param_config); auto callback = m->optimizer->needSpecialTraversal(param_config);
if (callback) { if (callback) {
auto retCallback = new ParameterTraverseCallback(); auto retCallback = new ParameterTraverseCallback();
......
...@@ -12,7 +12,6 @@ WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. ...@@ -12,7 +12,6 @@ WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
See the License for the specific language governing permissions and See the License for the specific language governing permissions and
limitations under the License. */ limitations under the License. */
#include "PaddleAPI.h" #include "PaddleAPI.h"
#include "paddle/gserver/gradientmachines/GradientMachine.h" #include "paddle/gserver/gradientmachines/GradientMachine.h"
#include "paddle/parameter/Argument.h" #include "paddle/parameter/Argument.h"
...@@ -42,8 +41,10 @@ struct Path { ...@@ -42,8 +41,10 @@ struct Path {
// position // position
static void findNBest(paddle::GradientMachine* gradMachine, static void findNBest(paddle::GradientMachine* gradMachine,
std::vector<paddle::Argument>& inArgs, std::vector<paddle::Argument>& inArgs,
std::vector<Path>& finalPaths, size_t bos_id, std::vector<Path>& finalPaths,
size_t eos_id, size_t max_length) { size_t bos_id,
size_t eos_id,
size_t max_length) {
std::vector<Path> paths; std::vector<Path> paths;
Path emptyPath; Path emptyPath;
paths.push_back(emptyPath); paths.push_back(emptyPath);
...@@ -166,7 +167,8 @@ public: ...@@ -166,7 +167,8 @@ public:
if (id < getSize()) { if (id < getSize()) {
Path& p = (*path_)[id]; Path& p = (*path_)[id];
std::ostringstream sout; std::ostringstream sout;
std::transform(p.ids.begin(), p.ids.end(), std::transform(p.ids.begin(),
p.ids.end(),
std::ostream_iterator<std::string>(sout, split ? " " : ""), std::ostream_iterator<std::string>(sout, split ? " " : ""),
[&](int id) { return (*dict_)[id]; }); [&](int id) { return (*dict_)[id]; });
return sout.str(); return sout.str();
......
...@@ -64,12 +64,11 @@ Trainer* Trainer::createByCommandLine() throw(IOError) { ...@@ -64,12 +64,11 @@ Trainer* Trainer::createByCommandLine() throw(IOError) {
Trainer::Trainer(TrainerConfig* config, GradientMachine* gm) Trainer::Trainer(TrainerConfig* config, GradientMachine* gm)
: m(new TrainerPrivate()) { : m(new TrainerPrivate()) {
m->init(config->m->conf, /* testing= */false, gm ? gm->m->machine : nullptr); m->init(config->m->conf, /* testing= */ false, gm ? gm->m->machine : nullptr);
} }
Trainer* Trainer::create(TrainerConfig* config, GradientMachine* gm) Trainer* Trainer::create(TrainerConfig* config,
throw(IOError) GradientMachine* gm) throw(IOError) {
{
auto retv = new Trainer(config, gm); auto retv = new Trainer(config, gm);
if (retv->m->getConfig().IsInitialized()) { if (retv->m->getConfig().IsInitialized()) {
return retv; return retv;
...@@ -134,15 +133,17 @@ void Trainer::finishTestPeriod() { m->finishTestPeriod(); } ...@@ -134,15 +133,17 @@ void Trainer::finishTestPeriod() { m->finishTestPeriod(); }
Matrix* Trainer::getLayerOutput(const std::string& layerName) { Matrix* Trainer::getLayerOutput(const std::string& layerName) {
auto nn = std::dynamic_pointer_cast<paddle::NeuralNetwork>( auto nn = std::dynamic_pointer_cast<paddle::NeuralNetwork>(
this->m->getGradientMachine()); this->m->getGradientMachine());
CHECK(nn) << "trainerInternal_.getGradientMachine() is not NeuralNetwork"; CHECK(nn) << "trainerInternal_.getGradientMachine() is not NeuralNetwork";
auto m = nn->getLayerOutput(layerName); auto m = nn->getLayerOutput(layerName);
return Matrix::createByPaddleMatrixPtr(&m); return Matrix::createByPaddleMatrixPtr(&m);
} }
void Trainer::forwardOneBatch(size_t batchSize) { m->forwardOneBatch(batchSize); } void Trainer::forwardOneBatch(size_t batchSize) {
m->forwardOneBatch(batchSize);
}
bool TrainerPrivate::forwardOneBatch(size_t batchSize) { bool TrainerPrivate::forwardOneBatch(size_t batchSize) {
CHECK(dataProvider_) << "data_provider is not specified"; CHECK(dataProvider_) << "data_provider is not specified";
paddle::DataBatch dataBatch; paddle::DataBatch dataBatch;
int num = dataProvider_->getNextBatch(batchSize, &dataBatch); int num = dataProvider_->getNextBatch(batchSize, &dataBatch);
...@@ -156,7 +157,6 @@ bool TrainerPrivate::forwardOneBatch(size_t batchSize) { ...@@ -156,7 +157,6 @@ bool TrainerPrivate::forwardOneBatch(size_t batchSize) {
void TrainerPrivate::forwardOneDataBatch( void TrainerPrivate::forwardOneDataBatch(
const std::vector<paddle::Argument>& inArgs) { const std::vector<paddle::Argument>& inArgs) {
std::vector<paddle::Argument>& outArgs = forwardOutput_; std::vector<paddle::Argument>& outArgs = forwardOutput_;
if (config_->getOptConfig().use_sparse_remote_updater()) { if (config_->getOptConfig().use_sparse_remote_updater()) {
......
...@@ -37,13 +37,15 @@ FloatArray::FloatArray(const float* b, const size_t l) ...@@ -37,13 +37,15 @@ FloatArray::FloatArray(const float* b, const size_t l)
IntArray::IntArray(const int* b, const size_t l, bool f) IntArray::IntArray(const int* b, const size_t l, bool f)
: buf(b), length(l), needFree(f) {} : buf(b), length(l), needFree(f) {}
IntWithFloatArray::IntWithFloatArray(const float* v, const int* i, size_t l, IntWithFloatArray::IntWithFloatArray(const float* v,
const int* i,
size_t l,
bool f) bool f)
: valBuf(v), idxBuf(i), length(l), needFree(f) {} : valBuf(v), idxBuf(i), length(l), needFree(f) {}
bool isUsingGpu() {return FLAGS_use_gpu;} bool isUsingGpu() { return FLAGS_use_gpu; }
void setUseGpu(bool useGpu) {FLAGS_use_gpu = useGpu;} void setUseGpu(bool useGpu) { FLAGS_use_gpu = useGpu; }
bool isGpuVersion() { bool isGpuVersion() {
#ifdef PADDLE_ONLY_CPU #ifdef PADDLE_ONLY_CPU
......
...@@ -12,7 +12,6 @@ WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. ...@@ -12,7 +12,6 @@ WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
See the License for the specific language governing permissions and See the License for the specific language governing permissions and
limitations under the License. */ limitations under the License. */
#include "PaddleAPI.h" #include "PaddleAPI.h"
#include "paddle/math/Vector.h" #include "paddle/math/Vector.h"
...@@ -39,8 +38,10 @@ IVector* IVector::create(const std::vector<int>& data, bool useGpu) { ...@@ -39,8 +38,10 @@ IVector* IVector::create(const std::vector<int>& data, bool useGpu) {
return v; return v;
} }
IVector* IVector::createVectorFromNumpy(int* data, int dim, bool copy, IVector* IVector::createVectorFromNumpy(int* data,
bool useGpu) throw (UnsupportError){ int dim,
bool copy,
bool useGpu) throw(UnsupportError) {
if (useGpu) { if (useGpu) {
/// if use gpu only copy=true is supported /// if use gpu only copy=true is supported
if (!copy) { if (!copy) {
...@@ -137,8 +138,8 @@ void IVector::copyToNumpyArray(int** view_m_data, int* dim1) { ...@@ -137,8 +138,8 @@ void IVector::copyToNumpyArray(int** view_m_data, int* dim1) {
if (auto cpuVec = dynamic_cast<paddle::CpuIVector*>(m->vec.get())) { if (auto cpuVec = dynamic_cast<paddle::CpuIVector*>(m->vec.get())) {
std::memcpy(*view_m_data, cpuVec->getData(), sizeof(int) * (*dim1)); std::memcpy(*view_m_data, cpuVec->getData(), sizeof(int) * (*dim1));
} else if (auto gpuVec = dynamic_cast<paddle::GpuIVector*>(m->vec.get())) { } else if (auto gpuVec = dynamic_cast<paddle::GpuIVector*>(m->vec.get())) {
hl_memcpy_device2host(*view_m_data, gpuVec->getData(), hl_memcpy_device2host(
sizeof(int) * (*dim1)); *view_m_data, gpuVec->getData(), sizeof(int) * (*dim1));
} else { } else {
LOG(INFO) << "Unexpected situation"; LOG(INFO) << "Unexpected situation";
} }
...@@ -201,8 +202,10 @@ Vector* Vector::createByPaddleVectorPtr(void* ptr) { ...@@ -201,8 +202,10 @@ Vector* Vector::createByPaddleVectorPtr(void* ptr) {
} }
} }
Vector* Vector::createVectorFromNumpy(float* data, int dim, bool copy, Vector* Vector::createVectorFromNumpy(float* data,
bool useGpu) throw (UnsupportError){ int dim,
bool copy,
bool useGpu) throw(UnsupportError) {
if (useGpu) { if (useGpu) {
/// if use gpu only copy=True is supported /// if use gpu only copy=True is supported
if (!copy) { if (!copy) {
...@@ -251,8 +254,8 @@ void Vector::copyToNumpyArray(float** view_m_data, int* dim1) { ...@@ -251,8 +254,8 @@ void Vector::copyToNumpyArray(float** view_m_data, int* dim1) {
if (auto cpuVec = dynamic_cast<paddle::CpuVector*>(m->vec.get())) { if (auto cpuVec = dynamic_cast<paddle::CpuVector*>(m->vec.get())) {
std::memcpy(*view_m_data, cpuVec->getData(), sizeof(float) * (*dim1)); std::memcpy(*view_m_data, cpuVec->getData(), sizeof(float) * (*dim1));
} else if (auto gpuVec = dynamic_cast<paddle::CpuVector*>(m->vec.get())) { } else if (auto gpuVec = dynamic_cast<paddle::CpuVector*>(m->vec.get())) {
hl_memcpy_device2host(*view_m_data, gpuVec->getData(), hl_memcpy_device2host(
sizeof(float) * (*dim1)); *view_m_data, gpuVec->getData(), sizeof(float) * (*dim1));
} else { } else {
LOG(INFO) << "Unexpected situation"; LOG(INFO) << "Unexpected situation";
} }
......
...@@ -81,5 +81,8 @@ else() ...@@ -81,5 +81,8 @@ else()
add_library(paddle_cuda ${CUDA_SOURCES}) add_library(paddle_cuda ${CUDA_SOURCES})
endif() endif()
add_style_check_target(paddle_cuda ${CUDA_SOURCES}) add_style_check_target(paddle_cuda
add_style_check_target(paddle_cuda ${CUDA_HEADERS}) ${CUDA_SOURCES}
${CUDA_HEADERS}
${CUDA_DSO_SOURCES}
${CUDA_CXX_WITH_GPU_SOURCES})
...@@ -12,7 +12,6 @@ WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. ...@@ -12,7 +12,6 @@ WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
See the License for the specific language governing permissions and See the License for the specific language governing permissions and
limitations under the License. */ limitations under the License. */
#ifndef HL_ACTIVATION_FUNCTIONS_H_ #ifndef HL_ACTIVATION_FUNCTIONS_H_
#define HL_ACTIVATION_FUNCTIONS_H_ #define HL_ACTIVATION_FUNCTIONS_H_
...@@ -21,11 +20,8 @@ limitations under the License. */ ...@@ -21,11 +20,8 @@ limitations under the License. */
/** /**
* Active functions: sigmoid, relu, tanh and linear. * Active functions: sigmoid, relu, tanh and linear.
*/ */
#define HPPL_ACTIVE_FUNCTION {hppl::sigmoid, \ #define HPPL_ACTIVE_FUNCTION \
hppl::relu, \ { hppl::sigmoid, hppl::relu, hppl::tanh, hppl::linear }
hppl::tanh, \
hppl::linear \
}
namespace hppl { namespace hppl {
...@@ -42,18 +38,18 @@ public: ...@@ -42,18 +38,18 @@ public:
#ifdef __NVCC__ #ifdef __NVCC__
namespace gpu { namespace gpu {
static __device__ Active<real>::forward forward[] = HPPL_ACTIVE_FUNCTION; static __device__ Active<real>::forward forward[] = HPPL_ACTIVE_FUNCTION;
static __device__ Active<real>::backward backward[] = HPPL_ACTIVE_FUNCTION; static __device__ Active<real>::backward backward[] = HPPL_ACTIVE_FUNCTION;
} }
#else #else
namespace cpu { namespace cpu {
static Active<real>::forward forward[] = HPPL_ACTIVE_FUNCTION; static Active<real>::forward forward[] = HPPL_ACTIVE_FUNCTION;
static Active<real>::backward backward[] = HPPL_ACTIVE_FUNCTION; static Active<real>::backward backward[] = HPPL_ACTIVE_FUNCTION;
} }
#ifdef __AVX__ #ifdef __AVX__
namespace avx { namespace avx {
static Active<__m256>::forward forward[] = HPPL_ACTIVE_FUNCTION; static Active<__m256>::forward forward[] = HPPL_ACTIVE_FUNCTION;
static Active<__m256>::backward backward[] = HPPL_ACTIVE_FUNCTION; static Active<__m256>::backward backward[] = HPPL_ACTIVE_FUNCTION;
} }
#endif #endif
......
...@@ -12,7 +12,6 @@ WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. ...@@ -12,7 +12,6 @@ WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
See the License for the specific language governing permissions and See the License for the specific language governing permissions and
limitations under the License. */ limitations under the License. */
#ifndef HL_AGGREGATE_H_ #ifndef HL_AGGREGATE_H_
#define HL_AGGREGATE_H_ #define HL_AGGREGATE_H_
......
...@@ -12,22 +12,21 @@ WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. ...@@ -12,22 +12,21 @@ WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
See the License for the specific language governing permissions and See the License for the specific language governing permissions and
limitations under the License. */ limitations under the License. */
#ifndef HL_AVX_FUNCTIONS_H_ #ifndef HL_AVX_FUNCTIONS_H_
#define HL_AVX_FUNCTIONS_H_ #define HL_AVX_FUNCTIONS_H_
#include <immintrin.h> #include <immintrin.h>
namespace hppl { namespace hppl {
__m256 relu(const __m256 a); __m256 relu(const __m256 a);
__m256 sigmoid(const __m256 a); __m256 sigmoid(const __m256 a);
__m256 tanh(const __m256 a); __m256 tanh(const __m256 a);
__m256 linear(const __m256 a); __m256 linear(const __m256 a);
__m256 relu(const __m256 a, const __m256 b); __m256 relu(const __m256 a, const __m256 b);
__m256 sigmoid(const __m256 a, const __m256 b); __m256 sigmoid(const __m256 a, const __m256 b);
__m256 tanh(const __m256 a, const __m256 b); __m256 tanh(const __m256 a, const __m256 b);
__m256 linear(const __m256 a, const __m256 b); __m256 linear(const __m256 a, const __m256 b);
} // namespace hppl } // namespace hppl
#endif // HL_AVX_FUNCTIONS_H_ #endif // HL_AVX_FUNCTIONS_H_
...@@ -12,8 +12,6 @@ WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. ...@@ -12,8 +12,6 @@ WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
See the License for the specific language governing permissions and See the License for the specific language governing permissions and
limitations under the License. */ limitations under the License. */
#ifndef HL_BASE_H_ #ifndef HL_BASE_H_
#define HL_BASE_H_ #define HL_BASE_H_
...@@ -33,36 +31,36 @@ limitations under the License. */ ...@@ -33,36 +31,36 @@ limitations under the License. */
* HPPL_STREAM_DEFAULT is HPPL default stream. * HPPL_STREAM_DEFAULT is HPPL default stream.
*/ */
typedef enum { typedef enum {
HPPL_STREAM_DEFAULT = 0, /* Thread Default Stream*/ HPPL_STREAM_DEFAULT = 0, /* Thread Default Stream*/
HPPL_STREAM_1 = 1, HPPL_STREAM_1 = 1,
HPPL_STREAM_2 = 2, HPPL_STREAM_2 = 2,
HPPL_STREAM_3 = 3, HPPL_STREAM_3 = 3,
HPPL_STREAM_4 = 4, HPPL_STREAM_4 = 4,
HPPL_THREAD_STREAM_1 = 5, HPPL_THREAD_STREAM_1 = 5,
HPPL_THREAD_STREAM_2 = 6, HPPL_THREAD_STREAM_2 = 6,
HPPL_THREAD_STREAM_3 = 7, HPPL_THREAD_STREAM_3 = 7,
HPPL_THREAD_STREAM_4 = 8, HPPL_THREAD_STREAM_4 = 8,
HPPL_STREAM_END HPPL_STREAM_END
} hl_stream_t; } hl_stream_t;
/** /**
* @brief HPPL activation mode. * @brief HPPL activation mode.
*/ */
typedef enum { typedef enum {
HL_ACTIVATION_SIGMOID = 0, HL_ACTIVATION_SIGMOID = 0,
HL_ACTIVATION_RELU = 1, HL_ACTIVATION_RELU = 1,
HL_ACTIVATION_TANH = 2, HL_ACTIVATION_TANH = 2,
HL_ACTIVATION_LINEAR = 3, HL_ACTIVATION_LINEAR = 3,
HL_ACTIVATION_END HL_ACTIVATION_END
} hl_activation_mode_t; } hl_activation_mode_t;
/** /**
* @brief Transpose type. * @brief Transpose type.
*/ */
typedef enum { typedef enum {
HPPL_OP_N = 0, /* transpose */ HPPL_OP_N = 0, /* transpose */
HPPL_OP_T = 1, /* non transpose */ HPPL_OP_T = 1, /* non transpose */
HPPL_OP_END HPPL_OP_END
} hl_trans_op_t; } hl_trans_op_t;
/** /**
...@@ -148,23 +146,21 @@ typedef struct { ...@@ -148,23 +146,21 @@ typedef struct {
* @brief Sparse matrix value type. * @brief Sparse matrix value type.
*/ */
typedef enum { typedef enum {
HL_NO_VALUE = 0, /* matrix values only 0 or 1 */ HL_NO_VALUE = 0, /* matrix values only 0 or 1 */
HL_FLOAT_VALUE = 1, HL_FLOAT_VALUE = 1,
HL_VALUE_END HL_VALUE_END
} hl_matrix_value_t; } hl_matrix_value_t;
/** /**
* @brief HPPL matrix format. * @brief HPPL matrix format.
*/ */
typedef enum { typedef enum {
HL_SPARSE_CSR = 0, HL_SPARSE_CSR = 0,
HL_SPARSE_CSC = 1, HL_SPARSE_CSC = 1,
HL_SPARSE_END HL_SPARSE_END
} hl_matrix_format_t; } hl_matrix_format_t;
typedef struct _hl_matrix_s *hl_matrix_s;
typedef struct _hl_matrix_s * hl_matrix_s;
/** /**
* @brief HPPL sparse matrix. * @brief HPPL sparse matrix.
...@@ -177,12 +173,12 @@ typedef struct _hl_matrix_s * hl_matrix_s; ...@@ -177,12 +173,12 @@ typedef struct _hl_matrix_s * hl_matrix_s;
* @param nnz nonzero values of sparse matrix. * @param nnz nonzero values of sparse matrix.
*/ */
typedef struct { typedef struct {
hl_matrix_s matrix; hl_matrix_s matrix;
hl_matrix_format_t format; hl_matrix_format_t format;
hl_matrix_value_t type; hl_matrix_value_t type;
int rows; int rows;
int cols; int cols;
size_t nnz; size_t nnz;
} _hl_sparse_matrix_s, *hl_sparse_matrix_s; } _hl_sparse_matrix_s, *hl_sparse_matrix_s;
#ifndef PADDLE_TYPE_DOUBLE #ifndef PADDLE_TYPE_DOUBLE
...@@ -195,7 +191,7 @@ typedef struct { ...@@ -195,7 +191,7 @@ typedef struct {
* *
* HL_FLOAT_MIN: 1.17549435e-38F * HL_FLOAT_MIN: 1.17549435e-38F
*/ */
#define HL_FLOAT_MAX 3.40282347e+38F #define HL_FLOAT_MAX 3.40282347e+38F
/** /**
* if real == double * if real == double
* *
...@@ -203,20 +199,18 @@ typedef struct { ...@@ -203,20 +199,18 @@ typedef struct {
* *
* HL_FLOAT_MIN: 2.2250738585072014e-308 * HL_FLOAT_MIN: 2.2250738585072014e-308
*/ */
#define HL_FLOAT_MIN 1.17549435e-38F #define HL_FLOAT_MIN 1.17549435e-38F
#else #else
#define HL_FLOAT_MAX 1.7976931348623157e+308 #define HL_FLOAT_MAX 1.7976931348623157e+308
#define HL_FLOAT_MIN 2.2250738585072014e-308 #define HL_FLOAT_MIN 2.2250738585072014e-308
#endif #endif
/** /**
* The maximum input value for exp, used to avoid overflow problem. * The maximum input value for exp, used to avoid overflow problem.
* *
* Currently only used for tanh function. * Currently only used for tanh function.
*/ */
#define EXP_MAX_INPUT 40.0 #define EXP_MAX_INPUT 40.0
/** /**
* @brief DIVUP(x, y) is similar to ceil(x / y). * @brief DIVUP(x, y) is similar to ceil(x / y).
...@@ -224,7 +218,7 @@ typedef struct { ...@@ -224,7 +218,7 @@ typedef struct {
* the size of blockDim. * the size of blockDim.
*/ */
#ifndef DIVUP #ifndef DIVUP
#define DIVUP(x, y) (((x) + (y) - 1) / (y)) #define DIVUP(x, y) (((x) + (y)-1) / (y))
#endif #endif
#ifdef __NVCC__ #ifdef __NVCC__
...@@ -233,7 +227,7 @@ typedef struct { ...@@ -233,7 +227,7 @@ typedef struct {
#include "hl_cuda.h" #include "hl_cuda.h"
#include "cuda_runtime.h" #include "cuda_runtime.h"
extern __thread bool g_sync_flag; extern __thread bool g_sync_flag;
extern __thread cudaStream_t default_stream; extern __thread cudaStream_t default_stream;
#define STREAM_DEFAULT default_stream #define STREAM_DEFAULT default_stream
...@@ -241,16 +235,15 @@ extern __thread cudaStream_t default_stream; ...@@ -241,16 +235,15 @@ extern __thread cudaStream_t default_stream;
* @brief Check cuda kernel execution. * @brief Check cuda kernel execution.
* @param msg error string * @param msg error string
*/ */
#define CHECK_SYNC(msg) \ #define CHECK_SYNC(msg) \
if (true == g_sync_flag) { \ if (true == g_sync_flag) { \
hl_stream_synchronize(HPPL_STREAM_DEFAULT); \ hl_stream_synchronize(HPPL_STREAM_DEFAULT); \
cudaError_t err \ cudaError_t err = (cudaError_t)hl_get_device_last_error(); \
= (cudaError_t)hl_get_device_last_error(); \ CHECK_EQ(cudaSuccess, err) \
CHECK_EQ(cudaSuccess, err) << "[" << msg << "] " \ << "[" << msg << "] " \
<< "CUDA error: " \ << "CUDA error: " << hl_get_device_error_string((size_t)err); \
<< hl_get_device_error_string((size_t)err); \
} }
#endif /* __NVCC__ */ #endif /* __NVCC__ */
#endif /* HL_BASE_H_ */ #endif /* HL_BASE_H_ */
...@@ -12,7 +12,6 @@ WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. ...@@ -12,7 +12,6 @@ WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
See the License for the specific language governing permissions and See the License for the specific language governing permissions and
limitations under the License. */ limitations under the License. */
#ifndef HL_BATCH_TRANSPOSE_H_ #ifndef HL_BATCH_TRANSPOSE_H_
#define HL_BATCH_TRANSPOSE_H_ #define HL_BATCH_TRANSPOSE_H_
...@@ -31,10 +30,7 @@ limitations under the License. */ ...@@ -31,10 +30,7 @@ limitations under the License. */
* order. Each batch has height * width data, which are * order. Each batch has height * width data, which are
* arranged in height-first (or row-first) manner. * arranged in height-first (or row-first) manner.
*/ */
extern void batchTranspose(const real* input, extern void batchTranspose(
real* output, const real* input, real* output, int width, int height, int batchSize);
int width,
int height,
int batchSize);
#endif // HL_BATCH_TRANSPOSE_H_ #endif // HL_BATCH_TRANSPOSE_H_
...@@ -12,7 +12,6 @@ WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. ...@@ -12,7 +12,6 @@ WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
See the License for the specific language governing permissions and See the License for the specific language governing permissions and
limitations under the License. */ limitations under the License. */
#ifndef HL_CNN_H_ #ifndef HL_CNN_H_
#define HL_CNN_H_ #define HL_CNN_H_
...@@ -37,15 +36,21 @@ limitations under the License. */ ...@@ -37,15 +36,21 @@ limitations under the License. */
* @param[in] alpha * @param[in] alpha
* @param[in] beta * @param[in] beta
*/ */
extern void hl_shrink_col2feature( extern void hl_shrink_col2feature(const real* dataCol,
const real * dataCol, size_t channels, size_t channels,
size_t height, size_t width, size_t height,
size_t blockH, size_t blockW, size_t width,
size_t strideH, size_t strideW, size_t blockH,
size_t paddingH, size_t paddingW, size_t blockW,
size_t outputH, size_t outputW, size_t strideH,
real* dataIm, size_t strideW,
real alpha = 1.0f, real beta = 0.0f); size_t paddingH,
size_t paddingW,
size_t outputH,
size_t outputW,
real* dataIm,
real alpha = 1.0f,
real beta = 0.0f);
/** /**
* @brief Expand feature to column. * @brief Expand feature to column.
...@@ -65,14 +70,19 @@ extern void hl_shrink_col2feature( ...@@ -65,14 +70,19 @@ extern void hl_shrink_col2feature(
* @param[out] dataCol expand data. * @param[out] dataCol expand data.
* *
*/ */
extern void hl_expand_feature2col( extern void hl_expand_feature2col(const real* dataIm,
const real* dataIm, size_t channels, size_t channels,
size_t height, size_t width, size_t height,
size_t blockH, size_t blockW, size_t width,
size_t strideH, size_t strideW, size_t blockH,
size_t paddingH, size_t paddingW, size_t blockW,
size_t outputH, size_t outputW, size_t strideH,
real* dataCol); size_t strideW,
size_t paddingH,
size_t paddingW,
size_t outputH,
size_t outputW,
real* dataCol);
/** /**
* @brief Maximum pool forward. * @brief Maximum pool forward.
...@@ -94,15 +104,21 @@ extern void hl_expand_feature2col( ...@@ -94,15 +104,21 @@ extern void hl_expand_feature2col(
* @param[in] tgtStride stride between output data samples. * @param[in] tgtStride stride between output data samples.
* *
*/ */
extern void hl_maxpool_forward( extern void hl_maxpool_forward(const int frameCnt,
const int frameCnt, const real* inputData, const real* inputData,
const int channels, const int channels,
const int height, const int width, const int height,
const int pooledH, const int pooledW, const int width,
const int sizeX, const int sizeY, const int pooledH,
const int strideH, const int strideW, const int pooledW,
const int paddingH, const int paddingW, const int sizeX,
real* tgtData, const int tgtStride); const int sizeY,
const int strideH,
const int strideW,
const int paddingH,
const int paddingW,
real* tgtData,
const int tgtStride);
/** /**
* @brief Maximum pool backward. * @brief Maximum pool backward.
...@@ -125,20 +141,28 @@ extern void hl_maxpool_forward( ...@@ -125,20 +141,28 @@ extern void hl_maxpool_forward(
* @param[in] paddingH padding height. * @param[in] paddingH padding height.
* @param[in] paddingW padding width. * @param[in] paddingW padding width.
* @param[out] targetGrad output grad. * @param[out] targetGrad output grad.
* @param[in] outStride stride between output data samples. * @param[in] outStride stride between output data samples.
* *
*/ */
extern void hl_maxpool_backward( extern void hl_maxpool_backward(const int frameCnt,
const int frameCnt, const real* inputData, const real* inputData,
const real* outData, const real* outGrad, const real* outData,
const int channels, const int height, const real* outGrad,
const int width, const int channels,
const int pooledH, const int pooledW, const int height,
const int sizeX, const int sizeY, const int width,
const int strideH, const int strideW, const int pooledH,
const int paddingH, const int paddingW, const int pooledW,
real scaleA, real scaleB, const int sizeX,
real* targetGrad, const int outStride); const int sizeY,
const int strideH,
const int strideW,
const int paddingH,
const int paddingW,
real scaleA,
real scaleB,
real* targetGrad,
const int outStride);
/** /**
* @brief Averge pool forward. * @brief Averge pool forward.
...@@ -160,15 +184,21 @@ extern void hl_maxpool_backward( ...@@ -160,15 +184,21 @@ extern void hl_maxpool_backward(
* @param[in] tgtStride stride between output data samples. * @param[in] tgtStride stride between output data samples.
* *
*/ */
extern void hl_avgpool_forward( extern void hl_avgpool_forward(const int frameCnt,
const int frameCnt, const real* inputData, const real* inputData,
const int channels, const int channels,
const int height, const int width, const int height,
const int pooledH, const int pooledW, const int width,
const int sizeX, const int sizeY, const int pooledH,
const int strideH, const int strideW, const int pooledW,
const int paddingH, const int paddingW, const int sizeX,
real* tgtData, const int tgtStride); const int sizeY,
const int strideH,
const int strideW,
const int paddingH,
const int paddingW,
real* tgtData,
const int tgtStride);
/** /**
* @brief Maximum pool backward. * @brief Maximum pool backward.
...@@ -189,19 +219,26 @@ extern void hl_avgpool_forward( ...@@ -189,19 +219,26 @@ extern void hl_avgpool_forward(
* @param[in] scaleA scale. * @param[in] scaleA scale.
* @param[in] scaleB scale. * @param[in] scaleB scale.
* @param[out] backGrad output grad. * @param[out] backGrad output grad.
* @param[in] outStride stride between output data samples. * @param[in] outStride stride between output data samples.
* *
*/ */
extern void hl_avgpool_backward( extern void hl_avgpool_backward(const int frameCnt,
const int frameCnt, const real* outGrad, const real* outGrad,
const int channels, const int height, const int channels,
const int width, const int height,
const int pooledH, const int pooledW, const int width,
const int sizeX, const int sizeY, const int pooledH,
const int strideH, const int strideW, const int pooledW,
int paddingH, int paddingW, const int sizeX,
real scaleA, real scaleB, const int sizeY,
real* backGrad, const int outStride); const int strideH,
const int strideW,
int paddingH,
int paddingW,
real scaleA,
real scaleB,
real* backGrad,
const int outStride);
/** /**
* @brief Cross-map-respose normalize forward. * @brief Cross-map-respose normalize forward.
...@@ -218,10 +255,16 @@ extern void hl_avgpool_backward( ...@@ -218,10 +255,16 @@ extern void hl_avgpool_backward(
* @param[in] beta scale. * @param[in] beta scale.
* *
*/ */
extern void hl_CMRNorm_forward( extern void hl_CMRNorm_forward(size_t frameCnt,
size_t frameCnt, const real* in, real* scale, real* out, const real* in,
size_t channels, size_t height, size_t width, size_t sizeX, real* scale,
real alpha, real beta); real* out,
size_t channels,
size_t height,
size_t width,
size_t sizeX,
real alpha,
real beta);
/** /**
* @brief Cross-map-respose normalize backward. * @brief Cross-map-respose normalize backward.
...@@ -240,11 +283,18 @@ extern void hl_CMRNorm_forward( ...@@ -240,11 +283,18 @@ extern void hl_CMRNorm_forward(
* @param[in] beta scale. * @param[in] beta scale.
* *
*/ */
extern void hl_CMRNorm_backward( extern void hl_CMRNorm_backward(size_t frameCnt,
size_t frameCnt, const real* inV, const real* scale, const real* inV,
const real* outV, const real* outDiff, real *inDiff, const real* scale,
size_t channels, size_t height, size_t width, size_t sizeX, const real* outV,
real alpha, real beta); const real* outDiff,
real* inDiff,
size_t channels,
size_t height,
size_t width,
size_t sizeX,
real alpha,
real beta);
/** /**
* @brief Bilinear interpolation forward. * @brief Bilinear interpolation forward.
...@@ -278,24 +328,24 @@ extern void hl_bilinear_forward(const real* inData, ...@@ -278,24 +328,24 @@ extern void hl_bilinear_forward(const real* inData,
const real ratioH, const real ratioH,
const real ratioW); const real ratioW);
/** /**
* @brief Bilinear interpolation backward. * @brief Bilinear interpolation backward.
* *
* @param[out] inGrad input gradient. * @param[out] inGrad input gradient.
* @param[in] inImgH input image height. * @param[in] inImgH input image height.
* @param[in] inImgW input image width. * @param[in] inImgW input image width.
* @param[in] inputH input batchSize. * @param[in] inputH input batchSize.
* @param[in] inputW input image data dim. * @param[in] inputW input image data dim.
* @param[in] outGrad output gradient. * @param[in] outGrad output gradient.
* @param[in] outImgH output image height. * @param[in] outImgH output image height.
* @param[in] outImgW output image width. * @param[in] outImgW output image width.
* @param[in] outputH output batchSize. * @param[in] outputH output batchSize.
* @param[in] outputW output image data dim. * @param[in] outputW output image data dim.
* @param[in] numChannels number of channels. * @param[in] numChannels number of channels.
* @param[in] ratioH inImgH / outImgH. * @param[in] ratioH inImgH / outImgH.
* @param[in] ratioW inImgW / outImgW. * @param[in] ratioW inImgW / outImgW.
* *
*/ */
extern void hl_bilinear_backward(real* inGrad, extern void hl_bilinear_backward(real* inGrad,
const size_t inImgH, const size_t inImgH,
const size_t inImgW, const size_t inImgW,
...@@ -321,9 +371,13 @@ extern void hl_bilinear_backward(real* inGrad, ...@@ -321,9 +371,13 @@ extern void hl_bilinear_backward(real* inGrad,
* @param[in] featLen feature length = image height * image width. * @param[in] featLen feature length = image height * image width.
* @param[in] groups number of groups. * @param[in] groups number of groups.
*/ */
extern void hl_maxout_forward( extern void hl_maxout_forward(const real* inData,
const real* inData, real* outData, int* idData, real* outData,
size_t batchSize, size_t size, size_t featLen, size_t groups); int* idData,
size_t batchSize,
size_t size,
size_t featLen,
size_t groups);
/** /**
* @brief MaxOut backward. * @brief MaxOut backward.
...@@ -336,8 +390,12 @@ extern void hl_maxout_forward( ...@@ -336,8 +390,12 @@ extern void hl_maxout_forward(
* @param[in] featLen feature length = image height * image width. * @param[in] featLen feature length = image height * image width.
* @param[in] groups number of groups. * @param[in] groups number of groups.
*/ */
extern void hl_maxout_backward( extern void hl_maxout_backward(real* inGrad,
real* inGrad, const real* outGrad, const int* idData, const real* outGrad,
size_t batchSize, size_t size, size_t featLen, size_t groups); const int* idData,
size_t batchSize,
size_t size,
size_t featLen,
size_t groups);
#endif /* HL_CNN_H_ */ #endif /* HL_CNN_H_ */
...@@ -12,7 +12,6 @@ WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. ...@@ -12,7 +12,6 @@ WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
See the License for the specific language governing permissions and See the License for the specific language governing permissions and
limitations under the License. */ limitations under the License. */
#ifndef HL_CUDA_H_ #ifndef HL_CUDA_H_
#define HL_CUDA_H_ #define HL_CUDA_H_
...@@ -22,8 +21,7 @@ limitations under the License. */ ...@@ -22,8 +21,7 @@ limitations under the License. */
/** /**
* @brief HPPL event. * @brief HPPL event.
*/ */
typedef struct _hl_event_st * hl_event_t; typedef struct _hl_event_st *hl_event_t;
/** /**
* @brief return cuda runtime api version. * @brief return cuda runtime api version.
...@@ -42,7 +40,7 @@ extern void hl_start(); ...@@ -42,7 +40,7 @@ extern void hl_start();
* if device is NULL, will start all GPU. * if device is NULL, will start all GPU.
* @param[in] number number of devices. * @param[in] number number of devices.
*/ */
extern void hl_specify_devices_start(int* device, int number); extern void hl_specify_devices_start(int *device, int number);
/** /**
* @brief Queries if a device may directly access a peer device's memory. * @brief Queries if a device may directly access a peer device's memory.
...@@ -126,7 +124,7 @@ extern int hl_get_device(); ...@@ -126,7 +124,7 @@ extern int hl_get_device();
* *
* @return dest_d pointer to device memory. * @return dest_d pointer to device memory.
*/ */
extern void* hl_malloc_device(size_t size); extern void *hl_malloc_device(size_t size);
/** /**
* @brief Free device memory. * @brief Free device memory.
...@@ -143,7 +141,7 @@ extern void hl_free_mem_device(void *dest_d); ...@@ -143,7 +141,7 @@ extern void hl_free_mem_device(void *dest_d);
* *
* @return dest_h pointer to host memory. * @return dest_h pointer to host memory.
*/ */
extern void* hl_malloc_host(size_t size); extern void *hl_malloc_host(size_t size);
/** /**
* @brief Free host page-lock memory. * @brief Free host page-lock memory.
...@@ -228,9 +226,9 @@ extern void hl_srand(unsigned int seed); ...@@ -228,9 +226,9 @@ extern void hl_srand(unsigned int seed);
* @param[in] stream stream id. * @param[in] stream stream id.
*/ */
extern void hl_memcpy_async(void *dst, extern void hl_memcpy_async(void *dst,
void *src, void *src,
size_t size, size_t size,
hl_stream_t stream); hl_stream_t stream);
/** /**
* @brief Waits for stream tasks to complete. * @brief Waits for stream tasks to complete.
...@@ -261,8 +259,7 @@ extern void hl_destroy_event(hl_event_t event); ...@@ -261,8 +259,7 @@ extern void hl_destroy_event(hl_event_t event);
* *
* @return time Time between start and end in ms. * @return time Time between start and end in ms.
*/ */
extern float hl_event_elapsed_time(hl_event_t start, extern float hl_event_elapsed_time(hl_event_t start, hl_event_t end);
hl_event_t end);
/** /**
* @brief Records an event. * @brief Records an event.
...@@ -300,7 +297,7 @@ extern void hl_set_device_flags_block(); ...@@ -300,7 +297,7 @@ extern void hl_set_device_flags_block();
/** /**
* @brief Returns the last error string from a cuda runtime call. * @brief Returns the last error string from a cuda runtime call.
*/ */
extern const char* hl_get_device_error_string(); extern const char *hl_get_device_error_string();
/** /**
* @brief Returns the last error string from a cuda runtime call. * @brief Returns the last error string from a cuda runtime call.
...@@ -309,7 +306,7 @@ extern const char* hl_get_device_error_string(); ...@@ -309,7 +306,7 @@ extern const char* hl_get_device_error_string();
* *
* @see hl_get_device_last_error() * @see hl_get_device_last_error()
*/ */
extern const char* hl_get_device_error_string(size_t err); extern const char *hl_get_device_error_string(size_t err);
/** /**
* @brief Returns the last error number. * @brief Returns the last error number.
......
...@@ -12,7 +12,6 @@ WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. ...@@ -12,7 +12,6 @@ WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
See the License for the specific language governing permissions and See the License for the specific language governing permissions and
limitations under the License. */ limitations under the License. */
#ifndef HL_CUDA_CUBLAS_H_ #ifndef HL_CUDA_CUBLAS_H_
#define HL_CUDA_CUBLAS_H_ #define HL_CUDA_CUBLAS_H_
...@@ -29,12 +28,8 @@ limitations under the License. */ ...@@ -29,12 +28,8 @@ limitations under the License. */
* @param[in] ldc the first dimension of C_d. * @param[in] ldc the first dimension of C_d.
* *
*/ */
extern void hl_matrix_transpose(real *A_d, extern void hl_matrix_transpose(
real *C_d, real *A_d, real *C_d, int dimM, int dimN, int lda, int ldc);
int dimM,
int dimN,
int lda,
int ldc);
/* /*
* @brief Matrix transpose, while lda = dimN, ldc = dimM. * @brief Matrix transpose, while lda = dimN, ldc = dimM.
...@@ -45,10 +40,7 @@ extern void hl_matrix_transpose(real *A_d, ...@@ -45,10 +40,7 @@ extern void hl_matrix_transpose(real *A_d,
* @param[in] dimN matrix width. * @param[in] dimN matrix width.
* *
*/ */
extern void hl_matrix_transpose(real *A_d, extern void hl_matrix_transpose(real *A_d, real *C_d, int dimM, int dimN);
real *C_d,
int dimM,
int dimN);
/* /*
* @brief Matrix inverse * @brief Matrix inverse
...@@ -60,11 +52,7 @@ extern void hl_matrix_transpose(real *A_d, ...@@ -60,11 +52,7 @@ extern void hl_matrix_transpose(real *A_d,
* @param[in] ldc the first dimension of C_d * @param[in] ldc the first dimension of C_d
* *
*/ */
extern void hl_matrix_inverse(real *A_d, extern void hl_matrix_inverse(real *A_d, real *C_d, int dimN, int lda, int ldc);
real *C_d,
int dimN,
int lda,
int ldc);
/** /**
* @brief C_d = alpha*(op(A_d) * op(B_d)) + beta*C_d * @brief C_d = alpha*(op(A_d) * op(B_d)) + beta*C_d
...@@ -84,12 +72,19 @@ extern void hl_matrix_inverse(real *A_d, ...@@ -84,12 +72,19 @@ extern void hl_matrix_inverse(real *A_d,
* @param[in] ldc the first dimension of C_d. * @param[in] ldc the first dimension of C_d.
* *
*/ */
extern void hl_matrix_mul(real *A_d, hl_trans_op_t transa, extern void hl_matrix_mul(real *A_d,
real *B_d, hl_trans_op_t transb, hl_trans_op_t transa,
real *B_d,
hl_trans_op_t transb,
real *C_d, real *C_d,
int dimM, int dimN, int dimK, int dimM,
real alpha, real beta, int dimN,
int lda, int ldb, int ldc); int dimK,
real alpha,
real beta,
int lda,
int ldb,
int ldc);
/** /**
* @brief C_d = alpha*(op(A_d) * op(B_d)) + beta*C_d * @brief C_d = alpha*(op(A_d) * op(B_d)) + beta*C_d
...@@ -106,11 +101,16 @@ extern void hl_matrix_mul(real *A_d, hl_trans_op_t transa, ...@@ -106,11 +101,16 @@ extern void hl_matrix_mul(real *A_d, hl_trans_op_t transa,
* @param[in] beta scalar used for multiplication. * @param[in] beta scalar used for multiplication.
* *
*/ */
extern void hl_matrix_mul(real *A_d, hl_trans_op_t transa, extern void hl_matrix_mul(real *A_d,
real *B_d, hl_trans_op_t transb, hl_trans_op_t transa,
real *B_d,
hl_trans_op_t transb,
real *C_d, real *C_d,
int dimM, int dimN, int dimK, int dimM,
real alpha, real beta); int dimN,
int dimK,
real alpha,
real beta);
/** /**
* @brief This function performs the matrix-vector multiplication. * @brief This function performs the matrix-vector multiplication.
...@@ -132,11 +132,17 @@ extern void hl_matrix_mul(real *A_d, hl_trans_op_t transa, ...@@ -132,11 +132,17 @@ extern void hl_matrix_mul(real *A_d, hl_trans_op_t transa,
* *
*/ */
extern void hl_matrix_mul_vector(real *A_d, hl_trans_op_t trans, extern void hl_matrix_mul_vector(real *A_d,
real *B_d, real *C_d, hl_trans_op_t trans,
int dimM, int dimN, real *B_d,
real alpha, real beta, real *C_d,
int lda, int incb, int incc); int dimM,
int dimN,
real alpha,
real beta,
int lda,
int incb,
int incc);
/** /**
* @brief This function performs the matrix-vector multiplication. * @brief This function performs the matrix-vector multiplication.
...@@ -154,9 +160,13 @@ extern void hl_matrix_mul_vector(real *A_d, hl_trans_op_t trans, ...@@ -154,9 +160,13 @@ extern void hl_matrix_mul_vector(real *A_d, hl_trans_op_t trans,
* @param[in] beta scalar used for multiplication. * @param[in] beta scalar used for multiplication.
* *
*/ */
extern void hl_matrix_mul_vector(real *A_d, hl_trans_op_t trans, extern void hl_matrix_mul_vector(real *A_d,
real *B_d, real *C_d, hl_trans_op_t trans,
int dimM, int dimN, real *B_d,
real alpha, real beta); real *C_d,
int dimM,
int dimN,
real alpha,
real beta);
#endif /* HL_CUDA_CUBLAS_H_ */ #endif /* HL_CUDA_CUBLAS_H_ */
...@@ -12,7 +12,6 @@ WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. ...@@ -12,7 +12,6 @@ WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
See the License for the specific language governing permissions and See the License for the specific language governing permissions and
limitations under the License. */ limitations under the License. */
#ifndef HL_CUDA_CUDNN_H_ #ifndef HL_CUDA_CUDNN_H_
#define HL_CUDA_CUDNN_H_ #define HL_CUDA_CUDNN_H_
...@@ -22,7 +21,7 @@ limitations under the License. */ ...@@ -22,7 +21,7 @@ limitations under the License. */
* hppl pooling mode * hppl pooling mode
*/ */
typedef enum { typedef enum {
HL_POOLING_MAX = 0, HL_POOLING_MAX = 0,
// average includes padded values // average includes padded values
HL_POOLING_AVERAGE = 1, HL_POOLING_AVERAGE = 1,
// average does not include padded values // average does not include padded values
...@@ -324,17 +323,16 @@ extern void hl_convolution_forward_add_bias(hl_tensor_descriptor bias, ...@@ -324,17 +323,16 @@ extern void hl_convolution_forward_add_bias(hl_tensor_descriptor bias,
* @param[in] sizeInBytes gpu workspace size (bytes). * @param[in] sizeInBytes gpu workspace size (bytes).
* @param[in] convBwdFilterAlgo backward filter algorithm. * @param[in] convBwdFilterAlgo backward filter algorithm.
*/ */
extern void hl_convolution_backward_filter( extern void hl_convolution_backward_filter(hl_tensor_descriptor input,
hl_tensor_descriptor input, real* input_data,
real* input_data, hl_tensor_descriptor output,
hl_tensor_descriptor output, real* output_grad_data,
real* output_grad_data, hl_filter_descriptor filter,
hl_filter_descriptor filter, real* filter_grad_data,
real* filter_grad_data, hl_convolution_descriptor conv,
hl_convolution_descriptor conv, void* gpuWorkSpace,
void* gpuWorkSpace, size_t sizeInBytes,
size_t sizeInBytes, int convBwdFilterAlgo);
int convBwdFilterAlgo);
/** /**
* @brief convolution backward data(calculate input image grad data). * @brief convolution backward data(calculate input image grad data).
...@@ -350,17 +348,16 @@ extern void hl_convolution_backward_filter( ...@@ -350,17 +348,16 @@ extern void hl_convolution_backward_filter(
* @param[in] sizeInBytes gpu workspace size (bytes). * @param[in] sizeInBytes gpu workspace size (bytes).
* @param[in] convBwdDataAlgo backward data algorithm. * @param[in] convBwdDataAlgo backward data algorithm.
*/ */
extern void hl_convolution_backward_data( extern void hl_convolution_backward_data(hl_tensor_descriptor input,
hl_tensor_descriptor input, real* input_data_grad,
real* input_data_grad, hl_tensor_descriptor output,
hl_tensor_descriptor output, real* output_grad_data,
real* output_grad_data, hl_filter_descriptor filter,
hl_filter_descriptor filter, real* filter_data,
real* filter_data, hl_convolution_descriptor conv,
hl_convolution_descriptor conv, void* gpuWorkSpace,
void* gpuWorkSpace, size_t sizeInBytes,
size_t sizeInBytes, int convBwdDataAlgo);
int convBwdDataAlgo);
/** /**
* @brief convolution backward bias(calculate bias grad data). * @brief convolution backward bias(calculate bias grad data).
...@@ -383,8 +380,8 @@ extern void hl_convolution_backward_bias(hl_tensor_descriptor bias, ...@@ -383,8 +380,8 @@ extern void hl_convolution_backward_bias(hl_tensor_descriptor bias,
* @param[in] height matrix height. * @param[in] height matrix height.
* @param[in] width matrix width. * @param[in] width matrix width.
*/ */
extern void hl_softmax_forward(real *input, extern void hl_softmax_forward(real* input,
real *output, real* output,
int height, int height,
int width); int width);
...@@ -396,8 +393,8 @@ extern void hl_softmax_forward(real *input, ...@@ -396,8 +393,8 @@ extern void hl_softmax_forward(real *input,
* @param[in] height matrix height. * @param[in] height matrix height.
* @param[in] width matrix width. * @param[in] width matrix width.
*/ */
extern void hl_softmax_backward(real *output_value, extern void hl_softmax_backward(real* output_value,
real *output_grad, real* output_grad,
int height, int height,
int width); int width);
...@@ -426,18 +423,18 @@ extern void hl_softmax_backward(real *output_value, ...@@ -426,18 +423,18 @@ extern void hl_softmax_backward(real *output_value,
* *
*/ */
extern void hl_batch_norm_forward_training(hl_tensor_descriptor inputDesc, extern void hl_batch_norm_forward_training(hl_tensor_descriptor inputDesc,
real *input, real* input,
hl_tensor_descriptor outputDesc, hl_tensor_descriptor outputDesc,
real *output, real* output,
hl_tensor_descriptor bnParamDesc, hl_tensor_descriptor bnParamDesc,
real *scale, real* scale,
real *bias, real* bias,
double factor, double factor,
real *runningMean, real* runningMean,
real *runningInvVar, real* runningInvVar,
double epsilon, double epsilon,
real *savedMean, real* savedMean,
real *savedVar); real* savedVar);
/** /**
* @brief cudnn batch norm forward. * @brief cudnn batch norm forward.
...@@ -463,14 +460,14 @@ extern void hl_batch_norm_forward_training(hl_tensor_descriptor inputDesc, ...@@ -463,14 +460,14 @@ extern void hl_batch_norm_forward_training(hl_tensor_descriptor inputDesc,
* *
*/ */
extern void hl_batch_norm_forward_inference(hl_tensor_descriptor inputDesc, extern void hl_batch_norm_forward_inference(hl_tensor_descriptor inputDesc,
real *input, real* input,
hl_tensor_descriptor outputDesc, hl_tensor_descriptor outputDesc,
real *output, real* output,
hl_tensor_descriptor bnParamDesc, hl_tensor_descriptor bnParamDesc,
real *scale, real* scale,
real *bias, real* bias,
real *estimatedMean, real* estimatedMean,
real *estimatedVar, real* estimatedVar,
double epsilon); double epsilon);
/** /**
...@@ -483,7 +480,8 @@ extern void hl_batch_norm_forward_inference(hl_tensor_descriptor inputDesc, ...@@ -483,7 +480,8 @@ extern void hl_batch_norm_forward_inference(hl_tensor_descriptor inputDesc,
* @param[in] inGradDesc input tensor descriptor desc. * @param[in] inGradDesc input tensor descriptor desc.
* @param[in] inGrad input data. * @param[in] inGrad input data.
* @param[in] dBnParamDesc tensor descriptor desc. * @param[in] dBnParamDesc tensor descriptor desc.
* bnScale, bnBias, running mean/var, save_mean/var. * bnScale, bnBias, running mean/var,
* save_mean/var.
* @param[in] scale batch normalization scale parameter (in original * @param[in] scale batch normalization scale parameter (in original
* paper scale is referred to as gamma). * paper scale is referred to as gamma).
* @param[in] scaleGrad batch normalization scale parameter (in original * @param[in] scaleGrad batch normalization scale parameter (in original
...@@ -497,17 +495,17 @@ extern void hl_batch_norm_forward_inference(hl_tensor_descriptor inputDesc, ...@@ -497,17 +495,17 @@ extern void hl_batch_norm_forward_inference(hl_tensor_descriptor inputDesc,
* *
*/ */
extern void hl_batch_norm_backward(hl_tensor_descriptor inputDesc, extern void hl_batch_norm_backward(hl_tensor_descriptor inputDesc,
real *input, real* input,
hl_tensor_descriptor outGradDesc, hl_tensor_descriptor outGradDesc,
real *outGrad, real* outGrad,
hl_tensor_descriptor inGradDesc, hl_tensor_descriptor inGradDesc,
real *inGrad, real* inGrad,
hl_tensor_descriptor dBnParamDesc, hl_tensor_descriptor dBnParamDesc,
real *scale, real* scale,
real *scaleGrad, real* scaleGrad,
real *biasGrad, real* biasGrad,
double epsilon, double epsilon,
real *savedMean, real* savedMean,
real *savedInvVar); real* savedInvVar);
#endif // HL_CUDA_CUDNN_H_ #endif // HL_CUDA_CUDNN_H_
...@@ -12,7 +12,6 @@ WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. ...@@ -12,7 +12,6 @@ WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
See the License for the specific language governing permissions and See the License for the specific language governing permissions and
limitations under the License. */ limitations under the License. */
#ifndef HL_DSO_LOADER_H_ #ifndef HL_DSO_LOADER_H_
#define HL_DSO_LOADER_H_ #define HL_DSO_LOADER_H_
......
...@@ -12,7 +12,6 @@ WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. ...@@ -12,7 +12,6 @@ WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
See the License for the specific language governing permissions and See the License for the specific language governing permissions and
limitations under the License. */ limitations under the License. */
#ifndef HL_FUNCTIONS_H_ #ifndef HL_FUNCTIONS_H_
#define HL_FUNCTIONS_H_ #define HL_FUNCTIONS_H_
...@@ -21,30 +20,30 @@ limitations under the License. */ ...@@ -21,30 +20,30 @@ limitations under the License. */
/** /**
* sigmoid threshold maximum * sigmoid threshold maximum
*/ */
#define SIGMOID_THRESHOLD_MIN -40.0 #define SIGMOID_THRESHOLD_MIN -40.0
/** /**
* sigmoid threshold minimum * sigmoid threshold minimum
*/ */
#define SIGMOID_THRESHOLD_MAX 13.0 #define SIGMOID_THRESHOLD_MAX 13.0
#ifndef __NVCC__ #ifndef __NVCC__
namespace hppl { namespace hppl {
/* /*
* forward activation * forward activation
*/ */
real relu(const real a); real relu(const real a);
real sigmoid(const real a); real sigmoid(const real a);
real tanh(const real a); real tanh(const real a);
real linear(const real a); real linear(const real a);
/* /*
* backward activation * backward activation
*/ */
real relu(const real a, const real b); real relu(const real a, const real b);
real sigmoid(const real a, const real b); real sigmoid(const real a, const real b);
real tanh(const real a, const real b); real tanh(const real a, const real b);
real linear(const real a, const real b); real linear(const real a, const real b);
} // namespace hppl } // namespace hppl
#ifdef __AVX__ #ifdef __AVX__
......
...@@ -12,7 +12,6 @@ WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. ...@@ -12,7 +12,6 @@ WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
See the License for the specific language governing permissions and See the License for the specific language governing permissions and
limitations under the License. */ limitations under the License. */
#ifndef HL_GPU_H_ #ifndef HL_GPU_H_
#define HL_GPU_H_ #define HL_GPU_H_
......
...@@ -12,7 +12,6 @@ WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. ...@@ -12,7 +12,6 @@ WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
See the License for the specific language governing permissions and See the License for the specific language governing permissions and
limitations under the License. */ limitations under the License. */
#ifndef HL_LSTM_H_ #ifndef HL_LSTM_H_
#define HL_LSTM_H_ #define HL_LSTM_H_
......
...@@ -12,7 +12,6 @@ WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. ...@@ -12,7 +12,6 @@ WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
See the License for the specific language governing permissions and See the License for the specific language governing permissions and
limitations under the License. */ limitations under the License. */
#ifndef HL_MATRIX_H_ #ifndef HL_MATRIX_H_
#define HL_MATRIX_H_ #define HL_MATRIX_H_
...@@ -30,13 +29,8 @@ limitations under the License. */ ...@@ -30,13 +29,8 @@ limitations under the License. */
* @param[in] beta scalar used for addition. * @param[in] beta scalar used for addition.
* *
*/ */
extern void hl_matrix_add(real* A_d, extern void hl_matrix_add(
real* B_d, real* A_d, real* B_d, real* C_d, int dimM, int dimN, real alpha, real beta);
real* C_d,
int dimM,
int dimN,
real alpha,
real beta);
/** /**
* @brief Matrix Softmax. * @brief Matrix Softmax.
* *
...@@ -46,7 +40,7 @@ extern void hl_matrix_add(real* A_d, ...@@ -46,7 +40,7 @@ extern void hl_matrix_add(real* A_d,
* @param[in] dimN matrix width. * @param[in] dimN matrix width.
* *
*/ */
extern void hl_matrix_softmax(real *A_d, real *C_d, int dimM, int dimN); extern void hl_matrix_softmax(real* A_d, real* C_d, int dimM, int dimN);
/** /**
* @brief Matrix softmax derivative. * @brief Matrix softmax derivative.
...@@ -58,11 +52,8 @@ extern void hl_matrix_softmax(real *A_d, real *C_d, int dimM, int dimN); ...@@ -58,11 +52,8 @@ extern void hl_matrix_softmax(real *A_d, real *C_d, int dimM, int dimN);
* @param[in] dimN matrix width. * @param[in] dimN matrix width.
* *
*/ */
extern void hl_matrix_softmax_derivative(real* grad_d, extern void hl_matrix_softmax_derivative(
real* output_d, real* grad_d, real* output_d, real* sftmaxSum_d, int dimM, int dimN);
real* sftmaxSum_d,
int dimM,
int dimN);
/** /**
* @brief Sequence softmax. * @brief Sequence softmax.
...@@ -73,8 +64,8 @@ extern void hl_matrix_softmax_derivative(real* grad_d, ...@@ -73,8 +64,8 @@ extern void hl_matrix_softmax_derivative(real* grad_d,
* @param[in] numSequence sequence number. * @param[in] numSequence sequence number.
* *
*/ */
extern void hl_sequence_softmax_forward(real *A_d, extern void hl_sequence_softmax_forward(real* A_d,
real *C_d, real* C_d,
const int* index, const int* index,
int numSequence); int numSequence);
...@@ -88,11 +79,8 @@ extern void hl_sequence_softmax_forward(real *A_d, ...@@ -88,11 +79,8 @@ extern void hl_sequence_softmax_forward(real *A_d,
* @param[in] dimN matrix width. * @param[in] dimN matrix width.
* *
*/ */
extern void hl_matrix_classification_error(real* A_d, extern void hl_matrix_classification_error(
int* B_d, real* A_d, int* B_d, real* C_d, int dimM, int dimN);
real* C_d,
int dimM,
int dimN);
/** /**
* @brief Matrix cross entropy. * @brief Matrix cross entropy.
...@@ -104,11 +92,8 @@ extern void hl_matrix_classification_error(real* A_d, ...@@ -104,11 +92,8 @@ extern void hl_matrix_classification_error(real* A_d,
* @param[in] dimN matrix width. * @param[in] dimN matrix width.
* *
*/ */
extern void hl_matrix_cross_entropy(real* A_d, extern void hl_matrix_cross_entropy(
real* C_d, real* A_d, real* C_d, int* label_d, int dimM, int dimN);
int* label_d,
int dimM,
int dimN);
/** /**
* @brief Matrix cross entropy back propagation. * @brief Matrix cross entropy back propagation.
...@@ -120,11 +105,8 @@ extern void hl_matrix_cross_entropy(real* A_d, ...@@ -120,11 +105,8 @@ extern void hl_matrix_cross_entropy(real* A_d,
* @param[in] dimN matrix width. * @param[in] dimN matrix width.
* *
*/ */
extern void hl_matrix_cross_entropy_bp(real* grad_d, extern void hl_matrix_cross_entropy_bp(
real* output_d, real* grad_d, real* output_d, int* label_d, int dimM, int dimN);
int* label_d,
int dimM,
int dimN);
/** /**
* @brief Matrix multi-binary label cross entropy * @brief Matrix multi-binary label cross entropy
...@@ -135,11 +117,8 @@ extern void hl_matrix_cross_entropy_bp(real* grad_d, ...@@ -135,11 +117,8 @@ extern void hl_matrix_cross_entropy_bp(real* grad_d,
* @param[in] dimM matrix height. * @param[in] dimM matrix height.
* @param[in] dimN matrix width. * @param[in] dimN matrix width.
*/ */
extern void hl_matrix_multi_binary_cross_entropy(real* output, extern void hl_matrix_multi_binary_cross_entropy(
real* entropy, real* output, real* entropy, hl_sparse_matrix_s mat, int dimM, int dimN);
hl_sparse_matrix_s mat,
int dimM,
int dimN);
/** /**
* @brief Matrix multi-binary label cross entropy backprop * @brief Matrix multi-binary label cross entropy backprop
...@@ -150,11 +129,8 @@ extern void hl_matrix_multi_binary_cross_entropy(real* output, ...@@ -150,11 +129,8 @@ extern void hl_matrix_multi_binary_cross_entropy(real* output,
* @param[in] dimM matrix height. * @param[in] dimM matrix height.
* @param[in] dimN matrix width. * @param[in] dimN matrix width.
*/ */
extern void hl_matrix_multi_binary_cross_entropy_bp(real* output, extern void hl_matrix_multi_binary_cross_entropy_bp(
real* grad, real* output, real* grad, hl_sparse_matrix_s mat, int dimM, int dimN);
hl_sparse_matrix_s mat,
int dimM,
int dimN);
/** /**
* @brief Matrix zero memory. * @brief Matrix zero memory.
...@@ -176,12 +152,8 @@ extern void hl_matrix_zero_mem(real* data, int num); ...@@ -176,12 +152,8 @@ extern void hl_matrix_zero_mem(real* data, int num);
* @param[in] partial_sum * @param[in] partial_sum
*/ */
extern void hl_param_relu_forward(real* output, extern void hl_param_relu_forward(
real* input, real* output, real* input, real* w, int width, int height, int partial_sum);
real* w,
int width,
int height,
int partial_sum);
/** /**
* @brief parameter relu backward w * @brief parameter relu backward w
* *
......
...@@ -12,7 +12,6 @@ WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. ...@@ -12,7 +12,6 @@ WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
See the License for the specific language governing permissions and See the License for the specific language governing permissions and
limitations under the License. */ limitations under the License. */
#ifndef HL_SEQUENCE_H_ #ifndef HL_SEQUENCE_H_
#define HL_SEQUENCE_H_ #define HL_SEQUENCE_H_
...@@ -32,7 +31,7 @@ limitations under the License. */ ...@@ -32,7 +31,7 @@ limitations under the License. */
extern void hl_max_sequence_forward(real* input, extern void hl_max_sequence_forward(real* input,
const int* sequence, const int* sequence,
real* output, real* output,
int *index, int* index,
int numSequences, int numSequences,
int dim); int dim);
...@@ -46,11 +45,8 @@ extern void hl_max_sequence_forward(real* input, ...@@ -46,11 +45,8 @@ extern void hl_max_sequence_forward(real* input,
* @param[in] dim input dimension. * @param[in] dim input dimension.
* *
*/ */
extern void hl_max_sequence_backward(real* outputGrad, extern void hl_max_sequence_backward(
int *index, real* outputGrad, int* index, real* inputGrad, int numSequences, int dim);
real* inputGrad,
int numSequences,
int dim);
/** /**
* @brief Context projection forward. * @brief Context projection forward.
...@@ -63,7 +59,8 @@ extern void hl_max_sequence_backward(real* outputGrad, ...@@ -63,7 +59,8 @@ extern void hl_max_sequence_backward(real* outputGrad,
* @param[in] inputDim input sequence dimension. * @param[in] inputDim input sequence dimension.
* @param[in] contextLength context length. * @param[in] contextLength context length.
* @param[in] contextStart context start. * @param[in] contextStart context start.
* @param[in] beginPad number of extra timesteps added at the beginning. * @param[in] beginPad number of extra timesteps added at the
* beginning.
* @param[in] isPadding trainable padding. * @param[in] isPadding trainable padding.
* *
*/ */
...@@ -109,7 +106,8 @@ extern void hl_context_projection_backward_data(real* outputGrad, ...@@ -109,7 +106,8 @@ extern void hl_context_projection_backward_data(real* outputGrad,
* @param[in] totalPad number of extra timesteps. * @param[in] totalPad number of extra timesteps.
* @param[in] contextLength context length. * @param[in] contextLength context length.
* @param[in] contextStart context start. * @param[in] contextStart context start.
* @param[in] beginPad number of extra timesteps added at the beginning. * @param[in] beginPad number of extra timesteps added at the
* beginning.
* *
*/ */
extern void hl_context_projection_backward_weight(real* outputGrad, extern void hl_context_projection_backward_weight(real* outputGrad,
...@@ -141,9 +139,9 @@ extern void hl_context_projection_backward_weight(real* outputGrad, ...@@ -141,9 +139,9 @@ extern void hl_context_projection_backward_weight(real* outputGrad,
* @param[in] seq2batch copy direction. * @param[in] seq2batch copy direction.
* *
*/ */
extern void hl_sequence2batch_copy(real *batch, extern void hl_sequence2batch_copy(real* batch,
real *sequence, real* sequence,
const int *batchIndex, const int* batchIndex,
int seqWidth, int seqWidth,
int batchCount, int batchCount,
bool seq2batch); bool seq2batch);
...@@ -167,9 +165,9 @@ extern void hl_sequence2batch_copy(real *batch, ...@@ -167,9 +165,9 @@ extern void hl_sequence2batch_copy(real *batch,
* @param[in] seq2batch copy direction. * @param[in] seq2batch copy direction.
* *
*/ */
extern void hl_sequence2batch_add(real *batch, extern void hl_sequence2batch_add(real* batch,
real *sequence, real* sequence,
int *batchIndex, int* batchIndex,
int seqWidth, int seqWidth,
int batchCount, int batchCount,
bool seq2batch); bool seq2batch);
......
...@@ -12,7 +12,6 @@ WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. ...@@ -12,7 +12,6 @@ WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
See the License for the specific language governing permissions and See the License for the specific language governing permissions and
limitations under the License. */ limitations under the License. */
#ifndef HL_SPARSE_H_ #ifndef HL_SPARSE_H_
#define HL_SPARSE_H_ #define HL_SPARSE_H_
...@@ -31,7 +30,7 @@ limitations under the License. */ ...@@ -31,7 +30,7 @@ limitations under the License. */
*/ */
extern void hl_malloc_sparse_matrix(hl_sparse_matrix_s *A_d, extern void hl_malloc_sparse_matrix(hl_sparse_matrix_s *A_d,
hl_matrix_format_t format, hl_matrix_format_t format,
hl_matrix_value_t value_type, hl_matrix_value_t value_type,
int dimM, int dimM,
int dimN, int dimN,
int nnz); int nnz);
...@@ -60,10 +59,10 @@ extern void hl_free_sparse_matrix(hl_sparse_matrix_s A_d); ...@@ -60,10 +59,10 @@ extern void hl_free_sparse_matrix(hl_sparse_matrix_s A_d);
* *
*/ */
extern void hl_construct_sparse_matrix(hl_sparse_matrix_s *A_d, extern void hl_construct_sparse_matrix(hl_sparse_matrix_s *A_d,
void * dest_d, void *dest_d,
size_t size, size_t size,
hl_matrix_format_t format, hl_matrix_format_t format,
hl_matrix_value_t value_type, hl_matrix_value_t value_type,
int dimM, int dimM,
int dimN, int dimN,
int nnz); int nnz);
...@@ -94,11 +93,11 @@ extern void hl_construct_sparse_matrix(hl_sparse_matrix_s *A_d, ...@@ -94,11 +93,11 @@ extern void hl_construct_sparse_matrix(hl_sparse_matrix_s *A_d,
* *
*/ */
extern void hl_construct_sparse_matrix(hl_sparse_matrix_s *A_d, extern void hl_construct_sparse_matrix(hl_sparse_matrix_s *A_d,
real* value_d, real *value_d,
int* rows_d, int *rows_d,
int* cols_d, int *cols_d,
hl_matrix_format_t format, hl_matrix_format_t format,
hl_matrix_value_t value_type, hl_matrix_value_t value_type,
int dimM, int dimM,
int dimN, int dimN,
int nnz); int nnz);
...@@ -259,10 +258,14 @@ extern void hl_matrix_csr_mul_dense(hl_sparse_matrix_s A_d, ...@@ -259,10 +258,14 @@ extern void hl_matrix_csr_mul_dense(hl_sparse_matrix_s A_d,
*/ */
extern void hl_matrix_csc_mul_dense(hl_sparse_matrix_s A_d, extern void hl_matrix_csc_mul_dense(hl_sparse_matrix_s A_d,
hl_trans_op_t transa, hl_trans_op_t transa,
real *B_d, hl_trans_op_t transb, real *B_d,
hl_trans_op_t transb,
real *C_d, real *C_d,
int dimM, int dimN, int dimK, int dimM,
real alpha, real beta); int dimN,
int dimK,
real alpha,
real beta);
/** /**
* @brief C_d = alpha*(op(A_d) * op(B_d)) + beta*C_d. * @brief C_d = alpha*(op(A_d) * op(B_d)) + beta*C_d.
...@@ -311,11 +314,16 @@ extern void hl_matrix_dense_mul_csc(real *A_d, ...@@ -311,11 +314,16 @@ extern void hl_matrix_dense_mul_csc(real *A_d,
* @note transb is not support HPPL_OP_T. * @note transb is not support HPPL_OP_T.
* *
*/ */
extern void hl_sparse_matrix_mul(real* A_d, hl_trans_op_t transa, extern void hl_sparse_matrix_mul(real *A_d,
real *B_d, hl_trans_op_t transb, hl_trans_op_t transa,
real *B_d,
hl_trans_op_t transb,
hl_sparse_matrix_s C_d, hl_sparse_matrix_s C_d,
int dimM, int dimN, int dimK, int dimM,
real alpha, real beta); int dimN,
int dimK,
real alpha,
real beta);
/** /**
* @brief C_d = alpha*(op(A_d) * op(B_d)) + beta*C_d * @brief C_d = alpha*(op(A_d) * op(B_d)) + beta*C_d
...@@ -336,12 +344,16 @@ extern void hl_sparse_matrix_mul(real* A_d, hl_trans_op_t transa, ...@@ -336,12 +344,16 @@ extern void hl_sparse_matrix_mul(real* A_d, hl_trans_op_t transa,
* @note transa is not support HPPL_OP_T. * @note transa is not support HPPL_OP_T.
* *
*/ */
extern void hl_matrix_dense_mul_csr(real *A_d, hl_trans_op_t transa, extern void hl_matrix_dense_mul_csr(real *A_d,
hl_trans_op_t transa,
hl_sparse_matrix_s B_d, hl_sparse_matrix_s B_d,
hl_trans_op_t transb, hl_trans_op_t transb,
real *C_d, real *C_d,
int dimM, int dimN, int dimK, int dimM,
real alpha, real beta); int dimN,
int dimK,
real alpha,
real beta);
/** /**
* @brief Memcpy csc_matrix to host. * @brief Memcpy csc_matrix to host.
...@@ -412,7 +424,6 @@ extern void hl_memcpy_from_csr_matrix(real *csr_val, ...@@ -412,7 +424,6 @@ extern void hl_memcpy_from_csr_matrix(real *csr_val,
hl_sparse_matrix_s csr_matrix, hl_sparse_matrix_s csr_matrix,
hl_stream_t stream); hl_stream_t stream);
/** /**
* @brief A_d[j] += B_d[i,j] for i in range(height) * @brief A_d[j] += B_d[i,j] for i in range(height)
* *
...@@ -423,19 +434,13 @@ extern void hl_memcpy_from_csr_matrix(real *csr_val, ...@@ -423,19 +434,13 @@ extern void hl_memcpy_from_csr_matrix(real *csr_val,
* @param[in] scale scale of B_d * @param[in] scale scale of B_d
* *
*/ */
extern void hl_sparse_matrix_column_sum(real* A_d, extern void hl_sparse_matrix_column_sum(
hl_sparse_matrix_s B_d, real *A_d, hl_sparse_matrix_s B_d, int dimM, int dimN, real scale);
int dimM,
int dimN,
real scale);
/** /**
* @brief implementation of csr sparse matrix in hl_sparse_matirx_column_sum * @brief implementation of csr sparse matrix in hl_sparse_matirx_column_sum
*/ */
extern void hl_matrix_csr_column_sum(real* A_d, extern void hl_matrix_csr_column_sum(
hl_sparse_matrix_s B_d, real *A_d, hl_sparse_matrix_s B_d, int dimM, int dimN, real scale);
int dimM,
int dimN,
real scale);
/** /**
* @brief A_d[i,j] += B_d[j] * @brief A_d[i,j] += B_d[j]
...@@ -446,13 +451,13 @@ extern void hl_matrix_csr_column_sum(real* A_d, ...@@ -446,13 +451,13 @@ extern void hl_matrix_csr_column_sum(real* A_d,
* *
*/ */
extern void hl_sparse_matrix_add_bias(hl_sparse_matrix_s A_d, extern void hl_sparse_matrix_add_bias(hl_sparse_matrix_s A_d,
real* B_d, real *B_d,
real scale); real scale);
/** /**
* @brief implementation of csr sparse matrix in hl_sparse_matrix_add_bias * @brief implementation of csr sparse matrix in hl_sparse_matrix_add_bias
*/ */
extern void hl_matrix_csr_add_bias(hl_sparse_matrix_s A_d, extern void hl_matrix_csr_add_bias(hl_sparse_matrix_s A_d,
real* B_d, real *B_d,
real scale); real scale);
/** /**
...@@ -470,7 +475,7 @@ extern void hl_matrix_csr_add_bias(hl_sparse_matrix_s A_d, ...@@ -470,7 +475,7 @@ extern void hl_matrix_csr_add_bias(hl_sparse_matrix_s A_d,
* *
*/ */
extern void hl_sparse_matrix_add_dense(hl_sparse_matrix_s A_d, extern void hl_sparse_matrix_add_dense(hl_sparse_matrix_s A_d,
real* B_d, real *B_d,
int dimM, int dimM,
int dimN, int dimN,
real alpha, real alpha,
...@@ -479,7 +484,7 @@ extern void hl_sparse_matrix_add_dense(hl_sparse_matrix_s A_d, ...@@ -479,7 +484,7 @@ extern void hl_sparse_matrix_add_dense(hl_sparse_matrix_s A_d,
* @brief implementation of csr sparse matrix in hl_sparse_matrix_add_dense * @brief implementation of csr sparse matrix in hl_sparse_matrix_add_dense
*/ */
extern void hl_matrix_csr_add_dense(hl_sparse_matrix_s A_d, extern void hl_matrix_csr_add_dense(hl_sparse_matrix_s A_d,
real* B_d, real *B_d,
int dimM, int dimM,
int dimN, int dimN,
real alpha, real alpha,
...@@ -493,7 +498,7 @@ extern void hl_matrix_csr_add_dense(hl_sparse_matrix_s A_d, ...@@ -493,7 +498,7 @@ extern void hl_matrix_csr_add_dense(hl_sparse_matrix_s A_d,
* @return return rows pointer, which is gpu address * @return return rows pointer, which is gpu address
* *
*/ */
extern int* hl_sparse_matrix_get_rows(hl_sparse_matrix_s sMat); extern int *hl_sparse_matrix_get_rows(hl_sparse_matrix_s sMat);
/** /**
* @brief get cols pionter of GpuSparseMatrix * @brief get cols pionter of GpuSparseMatrix
...@@ -503,7 +508,7 @@ extern int* hl_sparse_matrix_get_rows(hl_sparse_matrix_s sMat); ...@@ -503,7 +508,7 @@ extern int* hl_sparse_matrix_get_rows(hl_sparse_matrix_s sMat);
* @return return cols pointer, which is gpu address * @return return cols pointer, which is gpu address
* *
*/ */
extern int* hl_sparse_matrix_get_cols(hl_sparse_matrix_s sMat); extern int *hl_sparse_matrix_get_cols(hl_sparse_matrix_s sMat);
/** /**
* @brief get value pionter of GpuSparseMatrix * @brief get value pionter of GpuSparseMatrix
...@@ -513,7 +518,6 @@ extern int* hl_sparse_matrix_get_cols(hl_sparse_matrix_s sMat); ...@@ -513,7 +518,6 @@ extern int* hl_sparse_matrix_get_cols(hl_sparse_matrix_s sMat);
* @return return value pointer, which is gpu address * @return return value pointer, which is gpu address
* *
*/ */
extern real* hl_sparse_matrix_get_value(hl_sparse_matrix_s sMat); extern real *hl_sparse_matrix_get_value(hl_sparse_matrix_s sMat);
#endif /* HL_SPARSE_H_ */ #endif /* HL_SPARSE_H_ */
...@@ -12,7 +12,6 @@ WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. ...@@ -12,7 +12,6 @@ WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
See the License for the specific language governing permissions and See the License for the specific language governing permissions and
limitations under the License. */ limitations under the License. */
#ifndef HL_TABLE_APPLY_H_ #ifndef HL_TABLE_APPLY_H_
#define HL_TABLE_APPLY_H_ #define HL_TABLE_APPLY_H_
...@@ -31,8 +30,10 @@ limitations under the License. */ ...@@ -31,8 +30,10 @@ limitations under the License. */
* @param[in] dim width of table. * @param[in] dim width of table.
* *
*/ */
extern void hl_matrix_select_rows(real* output, int ldo, extern void hl_matrix_select_rows(real* output,
real* table, int ldt, int ldo,
real* table,
int ldt,
int* ids, int* ids,
int numSamples, int numSamples,
int tableSize, int tableSize,
...@@ -53,8 +54,10 @@ extern void hl_matrix_select_rows(real* output, int ldo, ...@@ -53,8 +54,10 @@ extern void hl_matrix_select_rows(real* output, int ldo,
* @param[in] dim width of table. * @param[in] dim width of table.
* *
*/ */
extern void hl_matrix_add_to_rows(real* table, int ldt, extern void hl_matrix_add_to_rows(real* table,
real* input, int ldi, int ldt,
real* input,
int ldi,
int* ids, int* ids,
int numSamples, int numSamples,
int tableSize, int tableSize,
...@@ -72,8 +75,7 @@ extern void hl_matrix_add_to_rows(real* table, int ldt, ...@@ -72,8 +75,7 @@ extern void hl_matrix_add_to_rows(real* table, int ldt,
* *
*/ */
template <class T> template <class T>
extern void hl_vector_select_from(T* dst, int sized, extern void hl_vector_select_from(
const T* src, int sizes, T* dst, int sized, const T* src, int sizes, const int* ids, int sizei);
const int* ids, int sizei);
#endif /* HL_TABLE_APPLY_H_ */ #endif /* HL_TABLE_APPLY_H_ */
...@@ -12,7 +12,6 @@ WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. ...@@ -12,7 +12,6 @@ WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
See the License for the specific language governing permissions and See the License for the specific language governing permissions and
limitations under the License. */ limitations under the License. */
#ifndef HL_TIME_H_ #ifndef HL_TIME_H_
#define HL_TIME_H_ #define HL_TIME_H_
......
...@@ -12,7 +12,6 @@ WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. ...@@ -12,7 +12,6 @@ WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
See the License for the specific language governing permissions and See the License for the specific language governing permissions and
limitations under the License. */ limitations under the License. */
#ifndef HL_TOP_K_H_ #ifndef HL_TOP_K_H_
#define HL_TOP_K_H_ #define HL_TOP_K_H_
...@@ -31,9 +30,11 @@ limitations under the License. */ ...@@ -31,9 +30,11 @@ limitations under the License. */
* @param[in] numSamples height of input value. * @param[in] numSamples height of input value.
* *
*/ */
extern void hl_matrix_top_k(real* topVal, int ldv, extern void hl_matrix_top_k(real* topVal,
int * topIds, int ldv,
real* src, int lds, int* topIds,
real* src,
int lds,
int dim, int dim,
int beamSize, int beamSize,
int numSamples); int numSamples);
...@@ -50,8 +51,9 @@ extern void hl_matrix_top_k(real* topVal, int ldv, ...@@ -50,8 +51,9 @@ extern void hl_matrix_top_k(real* topVal, int ldv,
* *
* @note Only support HL_SPARSE_CSR format. * @note Only support HL_SPARSE_CSR format.
*/ */
extern void hl_sparse_matrix_top_k(real* topVal, int ldv, extern void hl_sparse_matrix_top_k(real* topVal,
int * topIds, int ldv,
int* topIds,
hl_sparse_matrix_s src, hl_sparse_matrix_s src,
int beamSize, int beamSize,
int numSamples); int numSamples);
......
...@@ -12,29 +12,22 @@ WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. ...@@ -12,29 +12,22 @@ WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
See the License for the specific language governing permissions and See the License for the specific language governing permissions and
limitations under the License. */ limitations under the License. */
#ifndef HL_AGGREGATE_STUB_H_ #ifndef HL_AGGREGATE_STUB_H_
#define HL_AGGREGATE_STUB_H_ #define HL_AGGREGATE_STUB_H_
#include "hl_aggregate.h" #include "hl_aggregate.h"
inline void hl_matrix_row_sum(real *A_d, real *C_d, inline void hl_matrix_row_sum(real *A_d, real *C_d, int dimM, int dimN) {}
int dimM, int dimN) {}
inline void hl_matrix_row_max(real *A_d, real *C_d, inline void hl_matrix_row_max(real *A_d, real *C_d, int dimM, int dimN) {}
int dimM, int dimN) {}
inline void hl_matrix_row_min(real *A_d, real *C_d, inline void hl_matrix_row_min(real *A_d, real *C_d, int dimM, int dimN) {}
int dimM, int dimN) {}
inline void hl_matrix_column_sum(real *A_d, real *C_d, inline void hl_matrix_column_sum(real *A_d, real *C_d, int dimM, int dimN) {}
int dimM, int dimN) {}
inline void hl_matrix_column_max(real *A_d, real *C_d, inline void hl_matrix_column_max(real *A_d, real *C_d, int dimM, int dimN) {}
int dimM, int dimN) {}
inline void hl_matrix_column_min(real *A_d, real *C_d, inline void hl_matrix_column_min(real *A_d, real *C_d, int dimM, int dimN) {}
int dimM, int dimN) {}
inline void hl_vector_sum(real *A_d, real *C_h, int dimM) {} inline void hl_vector_sum(real *A_d, real *C_h, int dimM) {}
......
...@@ -12,84 +12,134 @@ WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. ...@@ -12,84 +12,134 @@ WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
See the License for the specific language governing permissions and See the License for the specific language governing permissions and
limitations under the License. */ limitations under the License. */
#ifndef HL_CNN_STUB_H_ #ifndef HL_CNN_STUB_H_
#define HL_CNN_STUB_H_ #define HL_CNN_STUB_H_
#include "hl_cnn.h" #include "hl_cnn.h"
inline void hl_shrink_col2feature( inline void hl_shrink_col2feature(const real* dataCol,
const real * dataCol, size_t channels, size_t channels,
size_t height, size_t width, size_t height,
size_t blockH, size_t blockW, size_t width,
size_t strideH, size_t strideW, size_t blockH,
size_t paddingH, size_t paddingW, size_t blockW,
size_t outputH, size_t outputW, size_t strideH,
real* dataIm, size_t strideW,
real alpha, real beta) {} size_t paddingH,
size_t paddingW,
inline void hl_expand_feature2col( size_t outputH,
const real* dataIm, size_t channels, size_t outputW,
size_t height, size_t width, real* dataIm,
size_t blockH, size_t blockW, real alpha,
size_t strideH, size_t strideW, real beta) {}
size_t paddingH, size_t paddingW,
size_t outputH, size_t outputW, inline void hl_expand_feature2col(const real* dataIm,
real* dataCol) {} size_t channels,
size_t height,
inline void hl_maxpool_forward( size_t width,
const int frameCnt, const real* inputData, size_t blockH,
const int channels, size_t blockW,
const int height, const int width, size_t strideH,
const int pooledH, const int pooledW, size_t strideW,
const int sizeX, const int sizeY, size_t paddingH,
const int strideH, const int strideW, size_t paddingW,
const int paddingH, const int paddingW, size_t outputH,
real* tgtData, const int tgtStride) {} size_t outputW,
real* dataCol) {}
inline void hl_maxpool_backward(
const int frameCnt, const real* inputData, inline void hl_maxpool_forward(const int frameCnt,
const real* outData, const real* outGrad, const real* inputData,
const int channels, const int height, const int channels,
const int width, const int height,
const int pooledH, const int pooledW, const int width,
const int sizeX, const int sizeY, const int pooledH,
const int strideH, const int strideW, const int pooledW,
const int paddingH, const int paddingW, const int sizeX,
real scaleA, real scaleB, const int sizeY,
real* targetGrad, const int outStride) {} const int strideH,
const int strideW,
inline void hl_avgpool_forward( const int paddingH,
const int frameCnt, const real* inputData, const int paddingW,
const int channels, real* tgtData,
const int height, const int width, const int tgtStride) {}
const int pooledH, const int pooledW,
const int sizeX, const int sizeY, inline void hl_maxpool_backward(const int frameCnt,
const int strideH, const int strideW, const real* inputData,
const int paddingH, const int paddingW, const real* outData,
real* tgtData, const int tgtStride) {} const real* outGrad,
const int channels,
inline void hl_avgpool_backward( const int height,
const int frameCnt, const real* outGrad, const int width,
const int channels, const int height, const int pooledH,
const int width, const int pooledW,
const int pooledH, const int pooledW, const int sizeX,
const int sizeX, const int sizeY, const int sizeY,
const int strideH, const int strideW, const int strideH,
int paddingH, int paddingW, const int strideW,
real scaleA, real scaleB, const int paddingH,
real* backGrad, const int outStride) {} const int paddingW,
real scaleA,
inline void hl_CMRNorm_forward( real scaleB,
size_t frameCnt, const real* in, real* scale, real* out, real* targetGrad,
size_t channels, size_t height, size_t width, size_t sizeX, const int outStride) {}
real alpha, real beta) {}
inline void hl_avgpool_forward(const int frameCnt,
inline void hl_CMRNorm_backward( const real* inputData,
size_t frameCnt, const real* inV, const real* scale, const int channels,
const real* outV, const real* outDiff, real *inDiff, const int height,
size_t channels, size_t height, size_t width, size_t sizeX, const int width,
real alpha, real beta) {} const int pooledH,
const int pooledW,
const int sizeX,
const int sizeY,
const int strideH,
const int strideW,
const int paddingH,
const int paddingW,
real* tgtData,
const int tgtStride) {}
inline void hl_avgpool_backward(const int frameCnt,
const real* outGrad,
const int channels,
const int height,
const int width,
const int pooledH,
const int pooledW,
const int sizeX,
const int sizeY,
const int strideH,
const int strideW,
int paddingH,
int paddingW,
real scaleA,
real scaleB,
real* backGrad,
const int outStride) {}
inline void hl_CMRNorm_forward(size_t frameCnt,
const real* in,
real* scale,
real* out,
size_t channels,
size_t height,
size_t width,
size_t sizeX,
real alpha,
real beta) {}
inline void hl_CMRNorm_backward(size_t frameCnt,
const real* inV,
const real* scale,
const real* outV,
const real* outDiff,
real* inDiff,
size_t channels,
size_t height,
size_t width,
size_t sizeX,
real alpha,
real beta) {}
inline void hl_bilinear_forward(const real* inData, inline void hl_bilinear_forward(const real* inData,
const size_t inImgH, const size_t inImgH,
...@@ -106,25 +156,33 @@ inline void hl_bilinear_forward(const real* inData, ...@@ -106,25 +156,33 @@ inline void hl_bilinear_forward(const real* inData,
const real ratioW) {} const real ratioW) {}
inline void hl_bilinear_backward(real* inGrad, inline void hl_bilinear_backward(real* inGrad,
const size_t inImgH, const size_t inImgH,
const size_t inImgW, const size_t inImgW,
const size_t inputH, const size_t inputH,
const size_t inputW, const size_t inputW,
const real* outGrad, const real* outGrad,
const size_t outImgH, const size_t outImgH,
const size_t outImgW, const size_t outImgW,
const size_t outputH, const size_t outputH,
const size_t outputW, const size_t outputW,
const size_t numChannels, const size_t numChannels,
const real ratioH, const real ratioH,
const real ratioW) {} const real ratioW) {}
inline void hl_maxout_forward( inline void hl_maxout_forward(const real* inData,
const real* inData, real* outData, int* idData, real* outData,
size_t batchSize, size_t size, size_t featLen, size_t group) {} int* idData,
size_t batchSize,
inline void hl_maxout_backward( size_t size,
real* inGrad, const real* outGrad, const int* idData, size_t featLen,
size_t batchSize, size_t size, size_t featLen, size_t group) {} size_t group) {}
inline void hl_maxout_backward(real* inGrad,
const real* outGrad,
const int* idData,
size_t batchSize,
size_t size,
size_t featLen,
size_t group) {}
#endif // HL_CNN_STUB_H_ #endif // HL_CNN_STUB_H_
...@@ -12,41 +12,42 @@ WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. ...@@ -12,41 +12,42 @@ WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
See the License for the specific language governing permissions and See the License for the specific language governing permissions and
limitations under the License. */ limitations under the License. */
#ifndef HL_CUDA_CUBLAS_STUB_H_ #ifndef HL_CUDA_CUBLAS_STUB_H_
#define HL_CUDA_CUBLAS_STUB_H_ #define HL_CUDA_CUBLAS_STUB_H_
#include "hl_cuda_cublas.h" #include "hl_cuda_cublas.h"
inline void hl_matrix_transpose(real *A_d, inline void hl_matrix_transpose(
real *C_d, real *A_d, real *C_d, int dimM, int dimN, int lda, int ldc) {}
int dimM,
int dimN, inline void hl_matrix_transpose(real *A_d, real *C_d, int dimM, int dimN) {}
int lda,
int ldc) {}
inline void hl_matrix_transpose(real *A_d,
real *C_d,
int dimM,
int dimN) {}
inline void hl_matrix_inverse(real *A_d,
real *C_d,
int dimN,
int lda,
int ldc) {}
inline void hl_matrix_mul(real *A_d, hl_trans_op_t transa,
real *B_d, hl_trans_op_t transb,
real *C_d,
int dimM, int dimN, int dimK,
real alpha, real beta,
int lda, int ldb, int ldc) {}
inline void hl_matrix_mul(real *A_d, hl_trans_op_t transa, inline void hl_matrix_inverse(
real *B_d, hl_trans_op_t transb, real *A_d, real *C_d, int dimN, int lda, int ldc) {}
inline void hl_matrix_mul(real *A_d,
hl_trans_op_t transa,
real *B_d,
hl_trans_op_t transb,
real *C_d,
int dimM,
int dimN,
int dimK,
real alpha,
real beta,
int lda,
int ldb,
int ldc) {}
inline void hl_matrix_mul(real *A_d,
hl_trans_op_t transa,
real *B_d,
hl_trans_op_t transb,
real *C_d, real *C_d,
int dimM, int dimN, int dimK, int dimM,
real alpha, real beta) {} int dimN,
int dimK,
real alpha,
real beta) {}
#endif // HL_CUDA_CUBLAS_STUB_H_ #endif // HL_CUDA_CUBLAS_STUB_H_
...@@ -12,15 +12,12 @@ WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. ...@@ -12,15 +12,12 @@ WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
See the License for the specific language governing permissions and See the License for the specific language governing permissions and
limitations under the License. */ limitations under the License. */
#ifndef HL_CUDA_CUDNN_STUB_H_ #ifndef HL_CUDA_CUDNN_STUB_H_
#define HL_CUDA_CUDNN_STUB_H_ #define HL_CUDA_CUDNN_STUB_H_
#include "hl_cuda_cudnn.h" #include "hl_cuda_cudnn.h"
inline int hl_get_cudnn_lib_version() { inline int hl_get_cudnn_lib_version() { return 0; }
return 0;
}
inline void hl_create_tensor_descriptor(hl_tensor_descriptor* image_desc) {} inline void hl_create_tensor_descriptor(hl_tensor_descriptor* image_desc) {}
...@@ -68,41 +65,41 @@ inline void hl_pooling_backward(hl_tensor_descriptor input, ...@@ -68,41 +65,41 @@ inline void hl_pooling_backward(hl_tensor_descriptor input,
hl_pooling_descriptor pooling) {} hl_pooling_descriptor pooling) {}
inline void hl_create_filter_descriptor(hl_filter_descriptor* filter, inline void hl_create_filter_descriptor(hl_filter_descriptor* filter,
int input_feature_maps, int input_feature_maps,
int output_feature_maps, int output_feature_maps,
int height, int height,
int width) {} int width) {}
inline void hl_destroy_filter_descriptor(hl_filter_descriptor filter) {} inline void hl_destroy_filter_descriptor(hl_filter_descriptor filter) {}
inline void hl_create_convolution_descriptor(hl_convolution_descriptor* conv, inline void hl_create_convolution_descriptor(hl_convolution_descriptor* conv,
hl_tensor_descriptor image, hl_tensor_descriptor image,
hl_filter_descriptor filter, hl_filter_descriptor filter,
int padding_height, int padding_height,
int padding_width, int padding_width,
int stride_height, int stride_height,
int stride_width) {} int stride_width) {}
inline void hl_reset_convolution_descriptor(hl_convolution_descriptor conv, inline void hl_reset_convolution_descriptor(hl_convolution_descriptor conv,
hl_tensor_descriptor image, hl_tensor_descriptor image,
hl_filter_descriptor filter, hl_filter_descriptor filter,
int padding_height, int padding_height,
int padding_width, int padding_width,
int stride_height, int stride_height,
int stride_width) {} int stride_width) {}
inline void hl_destroy_convolution_descriptor(hl_convolution_descriptor conv) {} inline void hl_destroy_convolution_descriptor(hl_convolution_descriptor conv) {}
inline void hl_conv_workspace(hl_tensor_descriptor input, inline void hl_conv_workspace(hl_tensor_descriptor input,
hl_tensor_descriptor output, hl_tensor_descriptor output,
hl_filter_descriptor filter, hl_filter_descriptor filter,
hl_convolution_descriptor conv, hl_convolution_descriptor conv,
int* convFwdAlgo, int* convFwdAlgo,
size_t* fwdLimitBytes, size_t* fwdLimitBytes,
int* convBwdDataAlgo, int* convBwdDataAlgo,
size_t* bwdDataLimitBytes, size_t* bwdDataLimitBytes,
int* convBwdFilterAlgo, int* convBwdFilterAlgo,
size_t* bwdFilterLimitBytes) {} size_t* bwdFilterLimitBytes) {}
inline void hl_convolution_forward(hl_tensor_descriptor input, inline void hl_convolution_forward(hl_tensor_descriptor input,
real* input_data, real* input_data,
...@@ -116,86 +113,84 @@ inline void hl_convolution_forward(hl_tensor_descriptor input, ...@@ -116,86 +113,84 @@ inline void hl_convolution_forward(hl_tensor_descriptor input,
int convFwdAlgo) {} int convFwdAlgo) {}
inline void hl_convolution_forward_add_bias(hl_tensor_descriptor bias, inline void hl_convolution_forward_add_bias(hl_tensor_descriptor bias,
real* bias_data, real* bias_data,
hl_tensor_descriptor output, hl_tensor_descriptor output,
real* output_data) {} real* output_data) {}
inline void hl_convolution_backward_filter( inline void hl_convolution_backward_filter(hl_tensor_descriptor input,
hl_tensor_descriptor input, real* input_data,
real* input_data, hl_tensor_descriptor output,
hl_tensor_descriptor output, real* output_grad_data,
real* output_grad_data, hl_filter_descriptor filter,
hl_filter_descriptor filter, real* filter_grad_data,
real* filter_grad_data, hl_convolution_descriptor conv,
hl_convolution_descriptor conv, void* gpuWorkSpace,
void* gpuWorkSpace, size_t sizeInBytes,
size_t sizeInBytes, int convBwdFilterAlgo) {}
int convBwdFilterAlgo) {}
inline void hl_convolution_backward_data(hl_tensor_descriptor input,
inline void hl_convolution_backward_data( real* input_data_grad,
hl_tensor_descriptor input, hl_tensor_descriptor output,
real* input_data_grad, real* output_grad_data,
hl_tensor_descriptor output, hl_filter_descriptor filter,
real* output_grad_data, real* filter_data,
hl_filter_descriptor filter, hl_convolution_descriptor conv,
real* filter_data, void* gpuWorkSpace,
hl_convolution_descriptor conv, size_t sizeInBytes,
void* gpuWorkSpace, int convBwdDataAlgo) {}
size_t sizeInBytes,
int convBwdDataAlgo) {}
inline void hl_convolution_backward_bias(hl_tensor_descriptor bias, inline void hl_convolution_backward_bias(hl_tensor_descriptor bias,
real* bias_grad_data, real* bias_grad_data,
hl_tensor_descriptor output, hl_tensor_descriptor output,
real* output_grad_data) {} real* output_grad_data) {}
inline void hl_softmax_forward(real *input, inline void hl_softmax_forward(real* input,
real *output, real* output,
int height,
int width) {}
inline void hl_softmax_backward(real *output_value,
real *output_grad,
int height, int height,
int width) {} int width) {}
inline void hl_softmax_backward(real* output_value,
real* output_grad,
int height,
int width) {}
inline void hl_batch_norm_forward_training(hl_tensor_descriptor inputDesc, inline void hl_batch_norm_forward_training(hl_tensor_descriptor inputDesc,
real *input, real* input,
hl_tensor_descriptor outputDesc, hl_tensor_descriptor outputDesc,
real *output, real* output,
hl_tensor_descriptor bnParamDesc, hl_tensor_descriptor bnParamDesc,
real *scale, real* scale,
real *bias, real* bias,
double factor, double factor,
real *runningMean, real* runningMean,
real *runningInvVar, real* runningInvVar,
double epsilon, double epsilon,
real *savedMean, real* savedMean,
real *savedVar) {} real* savedVar) {}
inline void hl_batch_norm_forward_inference(hl_tensor_descriptor inputDesc, inline void hl_batch_norm_forward_inference(hl_tensor_descriptor inputDesc,
real *input, real* input,
hl_tensor_descriptor outputDesc, hl_tensor_descriptor outputDesc,
real *output, real* output,
hl_tensor_descriptor bnParamDesc, hl_tensor_descriptor bnParamDesc,
real *scale, real* scale,
real *bias, real* bias,
real *estimatedMean, real* estimatedMean,
real *estimatedVar, real* estimatedVar,
double epsilon) {} double epsilon) {}
inline void hl_batch_norm_backward(hl_tensor_descriptor inputDesc, inline void hl_batch_norm_backward(hl_tensor_descriptor inputDesc,
real *input, real* input,
hl_tensor_descriptor outGradDesc, hl_tensor_descriptor outGradDesc,
real *outGrad, real* outGrad,
hl_tensor_descriptor inGradDesc, hl_tensor_descriptor inGradDesc,
real *inGrad, real* inGrad,
hl_tensor_descriptor dBnParamDesc, hl_tensor_descriptor dBnParamDesc,
real *scale, real* scale,
real *scaleGrad, real* scaleGrad,
real *biasGrad, real* biasGrad,
double epsilon, double epsilon,
real *savedMean, real* savedMean,
real *savedInvVar) {} real* savedInvVar) {}
#endif // HL_CUDA_CUDNN_STUB_H_ #endif // HL_CUDA_CUDNN_STUB_H_
...@@ -12,7 +12,6 @@ WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. ...@@ -12,7 +12,6 @@ WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
See the License for the specific language governing permissions and See the License for the specific language governing permissions and
limitations under the License. */ limitations under the License. */
#ifndef HL_CUDA_STUB_H_ #ifndef HL_CUDA_STUB_H_
#define HL_CUDA_STUB_H_ #define HL_CUDA_STUB_H_
...@@ -24,29 +23,25 @@ inline void hl_specify_devices_start(int *device, int number) {} ...@@ -24,29 +23,25 @@ inline void hl_specify_devices_start(int *device, int number) {}
inline void hl_init(int device) {} inline void hl_init(int device) {}
inline int hl_get_cuda_lib_version(int device) { inline int hl_get_cuda_lib_version(int device) { return 0; }
return 0;
}
inline void hl_fini() {} inline void hl_fini() {}
inline void hl_set_sync_flag(bool flag) {} inline void hl_set_sync_flag(bool flag) {}
inline bool hl_get_sync_flag() { inline bool hl_get_sync_flag() { return false; }
return false;
}
inline int hl_get_device_count() { return 0; } inline int hl_get_device_count() { return 0; }
inline void hl_set_device(int device) {} inline void hl_set_device(int device) {}
inline int hl_get_device() { return 0; } inline int hl_get_device() { return 0; }
inline void* hl_malloc_device(size_t size) { return NULL; } inline void *hl_malloc_device(size_t size) { return NULL; }
inline void hl_free_mem_device(void *dest_d) {} inline void hl_free_mem_device(void *dest_d) {}
inline void* hl_malloc_host(size_t size) { return NULL; } inline void *hl_malloc_host(size_t size) { return NULL; }
inline void hl_free_mem_host(void *dest_h) {} inline void hl_free_mem_host(void *dest_h) {}
...@@ -64,7 +59,9 @@ inline void hl_rand(real *dest_d, size_t num) {} ...@@ -64,7 +59,9 @@ inline void hl_rand(real *dest_d, size_t num) {}
inline void hl_srand(unsigned int seed) {} inline void hl_srand(unsigned int seed) {}
inline void hl_memcpy_async(void *dst, void *src, size_t size, inline void hl_memcpy_async(void *dst,
void *src,
size_t size,
hl_stream_t stream) {} hl_stream_t stream) {}
inline void hl_stream_synchronize(hl_stream_t stream) {} inline void hl_stream_synchronize(hl_stream_t stream) {}
...@@ -83,11 +80,11 @@ inline void hl_stream_wait_event(hl_stream_t stream, hl_event_t event) {} ...@@ -83,11 +80,11 @@ inline void hl_stream_wait_event(hl_stream_t stream, hl_event_t event) {}
inline void hl_event_synchronize(hl_event_t event) {} inline void hl_event_synchronize(hl_event_t event) {}
inline int hl_get_device_last_error() { return 0; } inline int hl_get_device_last_error() { return 0; }
inline const char* hl_get_device_error_string() { return NULL; } inline const char *hl_get_device_error_string() { return NULL; }
inline const char* hl_get_device_error_string(size_t err) { return NULL; } inline const char *hl_get_device_error_string(size_t err) { return NULL; }
inline bool hl_cuda_event_is_ready(hl_event_t event) { return true; } inline bool hl_cuda_event_is_ready(hl_event_t event) { return true; }
......
...@@ -12,7 +12,6 @@ WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. ...@@ -12,7 +12,6 @@ WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
See the License for the specific language governing permissions and See the License for the specific language governing permissions and
limitations under the License. */ limitations under the License. */
#ifndef HL_LSTM_STUB_H_ #ifndef HL_LSTM_STUB_H_
#define HL_LSTM_STUB_H_ #define HL_LSTM_STUB_H_
......
...@@ -12,7 +12,6 @@ WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. ...@@ -12,7 +12,6 @@ WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
See the License for the specific language governing permissions and See the License for the specific language governing permissions and
limitations under the License. */ limitations under the License. */
#ifndef HL_MATRIX_STUB_H_ #ifndef HL_MATRIX_STUB_H_
#define HL_MATRIX_STUB_H_ #define HL_MATRIX_STUB_H_
...@@ -26,48 +25,30 @@ inline void hl_matrix_add(real* A_d, ...@@ -26,48 +25,30 @@ inline void hl_matrix_add(real* A_d,
real alpha, real alpha,
real beta) {} real beta) {}
inline void hl_matrix_softmax(real *A_d, real *C_d, int dimM, int dimN) {} inline void hl_matrix_softmax(real* A_d, real* C_d, int dimM, int dimN) {}
inline void hl_sequence_softmax_forward(real *A_d, inline void hl_sequence_softmax_forward(real* A_d,
real *C_d, real* C_d,
const int* index, const int* index,
int numSequence) {} int numSequence) {}
inline void hl_matrix_softmax_derivative(real* grad_d, inline void hl_matrix_softmax_derivative(
real* output_d, real* grad_d, real* output_d, real* sftmaxSum_d, int dimM, int dimN) {}
real* sftmaxSum_d,
int dimM, inline void hl_matrix_classification_error(
int dimN) {} real* A_d, int* B_d, real* C_d, int dimM, int dimN) {}
inline void hl_matrix_classification_error(real* A_d, inline void hl_matrix_cross_entropy(
int* B_d, real* A_d, real* C_d, int* label_d, int dimM, int dimN) {}
real* C_d,
int dimM, inline void hl_matrix_cross_entropy_bp(
int dimN) {} real* grad_d, real* output_d, int* label_d, int dimM, int dimN) {}
inline void hl_matrix_cross_entropy(real* A_d, inline void hl_matrix_multi_binary_cross_entropy(
real* C_d, real* output, real* entropy, hl_sparse_matrix_s mat, int dimM, int dimN) {}
int* label_d,
int dimM, inline void hl_matrix_multi_binary_cross_entropy_bp(
int dimN) {} real* output, real* grad, hl_sparse_matrix_s mat, int dimM, int dimN) {}
inline void hl_matrix_cross_entropy_bp(real* grad_d,
real* output_d,
int* label_d,
int dimM,
int dimN) {}
inline void hl_matrix_multi_binary_cross_entropy(real* output,
real* entropy,
hl_sparse_matrix_s mat,
int dimM,
int dimN) {}
inline void hl_matrix_multi_binary_cross_entropy_bp(real* output,
real* grad,
hl_sparse_matrix_s mat,
int dimM,
int dimN) {}
inline void hl_matrix_zero_mem(real* data, int num) {} inline void hl_matrix_zero_mem(real* data, int num) {}
...@@ -101,7 +82,6 @@ inline void hl_cossim(real* output, ...@@ -101,7 +82,6 @@ inline void hl_cossim(real* output,
int input2_height, int input2_height,
real scale) {} real scale) {}
inline void hl_cossim_derivative(real* grad, inline void hl_cossim_derivative(real* grad,
real* output, real* output,
real* prevOutX, real* prevOutX,
......
...@@ -12,7 +12,6 @@ WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. ...@@ -12,7 +12,6 @@ WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
See the License for the specific language governing permissions and See the License for the specific language governing permissions and
limitations under the License. */ limitations under the License. */
#ifndef HL_SEQUENCE_STUB_H_ #ifndef HL_SEQUENCE_STUB_H_
#define HL_SEQUENCE_STUB_H_ #define HL_SEQUENCE_STUB_H_
...@@ -21,15 +20,12 @@ limitations under the License. */ ...@@ -21,15 +20,12 @@ limitations under the License. */
inline void hl_max_sequence_forward(real* input, inline void hl_max_sequence_forward(real* input,
const int* sequence, const int* sequence,
real* output, real* output,
int *index, int* index,
int numSequences, int numSequences,
int dim) {} int dim) {}
inline void hl_max_sequence_backward(real* outputGrad, inline void hl_max_sequence_backward(
int *index, real* outputGrad, int* index, real* inputGrad, int numSequences, int dim) {}
real* inputGrad,
int numSequences,
int dim) {}
inline void hl_context_projection_forward(real* input, inline void hl_context_projection_forward(real* input,
const int* sequence, const int* sequence,
...@@ -60,16 +56,16 @@ inline void hl_context_projection_backward_weight(real* outputGrad, ...@@ -60,16 +56,16 @@ inline void hl_context_projection_backward_weight(real* outputGrad,
int contextStart, int contextStart,
int beginPad) {} int beginPad) {}
inline void hl_sequence2batch_copy(real *batch, inline void hl_sequence2batch_copy(real* batch,
real *sequence, real* sequence,
const int *batchIndex, const int* batchIndex,
int seqWidth, int seqWidth,
int batchCount, int batchCount,
bool seq2batch) {} bool seq2batch) {}
inline void hl_sequence2batch_add(real *batch, inline void hl_sequence2batch_add(real* batch,
real *sequence, real* sequence,
int *batchIndex, int* batchIndex,
int seqWidth, int seqWidth,
int batchCount, int batchCount,
bool seq2batch) {} bool seq2batch) {}
......
...@@ -12,7 +12,6 @@ WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. ...@@ -12,7 +12,6 @@ WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
See the License for the specific language governing permissions and See the License for the specific language governing permissions and
limitations under the License. */ limitations under the License. */
#ifndef HL_SPARSE_STUB_H_ #ifndef HL_SPARSE_STUB_H_
#define HL_SPARSE_STUB_H_ #define HL_SPARSE_STUB_H_
...@@ -20,7 +19,7 @@ limitations under the License. */ ...@@ -20,7 +19,7 @@ limitations under the License. */
inline void hl_malloc_sparse_matrix(hl_sparse_matrix_s *A_d, inline void hl_malloc_sparse_matrix(hl_sparse_matrix_s *A_d,
hl_matrix_format_t format, hl_matrix_format_t format,
hl_matrix_value_t value_type, hl_matrix_value_t value_type,
int dimM, int dimM,
int dimN, int dimN,
int nnz) {} int nnz) {}
...@@ -28,20 +27,20 @@ inline void hl_malloc_sparse_matrix(hl_sparse_matrix_s *A_d, ...@@ -28,20 +27,20 @@ inline void hl_malloc_sparse_matrix(hl_sparse_matrix_s *A_d,
inline void hl_free_sparse_matrix(hl_sparse_matrix_s A_d) {} inline void hl_free_sparse_matrix(hl_sparse_matrix_s A_d) {}
inline void hl_construct_sparse_matrix(hl_sparse_matrix_s *A_d, inline void hl_construct_sparse_matrix(hl_sparse_matrix_s *A_d,
void * dest_d, void *dest_d,
size_t size, size_t size,
hl_matrix_format_t format, hl_matrix_format_t format,
hl_matrix_value_t value_type, hl_matrix_value_t value_type,
int dimM, int dimM,
int dimN, int dimN,
int nnz) {} int nnz) {}
inline void hl_construct_sparse_matrix(hl_sparse_matrix_s *A_d, inline void hl_construct_sparse_matrix(hl_sparse_matrix_s *A_d,
real* value_d, real *value_d,
int* rows_d, int *rows_d,
int* cols_d, int *cols_d,
hl_matrix_format_t format, hl_matrix_format_t format,
hl_matrix_value_t value_type, hl_matrix_value_t value_type,
int dimM, int dimM,
int dimN, int dimN,
int nnz) {} int nnz) {}
...@@ -87,10 +86,14 @@ inline void hl_matrix_csr_mul_dense(hl_sparse_matrix_s A_d, ...@@ -87,10 +86,14 @@ inline void hl_matrix_csr_mul_dense(hl_sparse_matrix_s A_d,
inline void hl_matrix_csc_mul_dense(hl_sparse_matrix_s A_d, inline void hl_matrix_csc_mul_dense(hl_sparse_matrix_s A_d,
hl_trans_op_t transa, hl_trans_op_t transa,
real *B_d, hl_trans_op_t transb, real *B_d,
hl_trans_op_t transb,
real *C_d, real *C_d,
int dimM, int dimN, int dimK, int dimM,
real alpha, real beta) {} int dimN,
int dimK,
real alpha,
real beta) {}
inline void hl_matrix_dense_mul_csc(real *A_d, inline void hl_matrix_dense_mul_csc(real *A_d,
hl_trans_op_t transa, hl_trans_op_t transa,
...@@ -103,18 +106,27 @@ inline void hl_matrix_dense_mul_csc(real *A_d, ...@@ -103,18 +106,27 @@ inline void hl_matrix_dense_mul_csc(real *A_d,
real alpha, real alpha,
real beta) {} real beta) {}
inline void hl_sparse_matrix_mul(real* A_d, hl_trans_op_t transa, inline void hl_sparse_matrix_mul(real *A_d,
real *B_d, hl_trans_op_t transb, hl_trans_op_t transa,
real *B_d,
hl_trans_op_t transb,
hl_sparse_matrix_s C_d, hl_sparse_matrix_s C_d,
int dimM, int dimN, int dimK, int dimM,
real alpha, real beta) {} int dimN,
int dimK,
real alpha,
real beta) {}
inline void hl_matrix_dense_mul_csr(real *A_d, hl_trans_op_t transa, inline void hl_matrix_dense_mul_csr(real *A_d,
hl_trans_op_t transa,
hl_sparse_matrix_s B_d, hl_sparse_matrix_s B_d,
hl_trans_op_t transb, hl_trans_op_t transb,
real *C_d, real *C_d,
int dimM, int dimN, int dimK, int dimM,
real alpha, real beta) {} int dimN,
int dimK,
real alpha,
real beta) {}
inline void hl_memcpy_from_csc_matrix(real *csc_val, inline void hl_memcpy_from_csc_matrix(real *csc_val,
size_t val_size, size_t val_size,
...@@ -134,49 +146,39 @@ inline void hl_memcpy_from_csr_matrix(real *csr_val, ...@@ -134,49 +146,39 @@ inline void hl_memcpy_from_csr_matrix(real *csr_val,
hl_sparse_matrix_s csr_matrix, hl_sparse_matrix_s csr_matrix,
hl_stream_t stream) {} hl_stream_t stream) {}
inline void hl_sparse_matrix_column_sum(real* A_d, inline void hl_sparse_matrix_column_sum(
hl_sparse_matrix_s B_d, real *A_d, hl_sparse_matrix_s B_d, int dimM, int dimN, real scale) {}
int dimM,
int dimN,
real scale) {}
inline void hl_matrix_csr_column_sum(real* A_d, inline void hl_matrix_csr_column_sum(
hl_sparse_matrix_s B_d, real *A_d, hl_sparse_matrix_s B_d, int dimM, int dimN, real scale) {}
int dimM,
int dimN,
real scale) {}
inline void hl_sparse_matrix_add_bias(hl_sparse_matrix_s A_d, inline void hl_sparse_matrix_add_bias(hl_sparse_matrix_s A_d,
real* B_d, real *B_d,
real scale) {} real scale) {}
inline void hl_matrix_csr_add_bias(hl_sparse_matrix_s A_d, inline void hl_matrix_csr_add_bias(hl_sparse_matrix_s A_d,
real* B_d, real *B_d,
real scale) {} real scale) {}
inline void hl_sparse_matrix_add_dense(hl_sparse_matrix_s A_d, inline void hl_sparse_matrix_add_dense(hl_sparse_matrix_s A_d,
real* B_d, real *B_d,
int dimM, int dimM,
int dimN, int dimN,
real alpha, real alpha,
real beta) {} real beta) {}
inline void hl_matrix_csr_add_dense(hl_sparse_matrix_s A_d, inline void hl_matrix_csr_add_dense(hl_sparse_matrix_s A_d,
real* B_d, real *B_d,
int dimM, int dimM,
int dimN, int dimN,
real alpha, real alpha,
real beta) {} real beta) {}
inline int* hl_sparse_matrix_get_rows(hl_sparse_matrix_s sMat) { inline int *hl_sparse_matrix_get_rows(hl_sparse_matrix_s sMat) { return NULL; }
return NULL;
}
inline int* hl_sparse_matrix_get_cols(hl_sparse_matrix_s sMat) { inline int *hl_sparse_matrix_get_cols(hl_sparse_matrix_s sMat) { return NULL; }
return NULL;
}
inline real* hl_sparse_matrix_get_value(hl_sparse_matrix_s sMat) { inline real *hl_sparse_matrix_get_value(hl_sparse_matrix_s sMat) {
return NULL; return NULL;
} }
......
此差异已折叠。
...@@ -12,62 +12,58 @@ WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. ...@@ -12,62 +12,58 @@ WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
See the License for the specific language governing permissions and See the License for the specific language governing permissions and
limitations under the License. */ limitations under the License. */
#include <immintrin.h> #include <immintrin.h>
#include "hl_functions.h" #include "hl_functions.h"
namespace hppl { namespace hppl {
extern __m256 exp(__m256 a); extern __m256 exp(__m256 a);
__m256 relu(const __m256 a) { __m256 relu(const __m256 a) {
__m256 tmp = _mm256_set1_ps(0.0f); __m256 tmp = _mm256_set1_ps(0.0f);
return _mm256_max_ps(a, tmp); return _mm256_max_ps(a, tmp);
} }
__m256 sigmoid(const __m256 a) { __m256 sigmoid(const __m256 a) {
__m256 max = _mm256_set1_ps(SIGMOID_THRESHOLD_MAX); __m256 max = _mm256_set1_ps(SIGMOID_THRESHOLD_MAX);
__m256 min = _mm256_set1_ps(SIGMOID_THRESHOLD_MIN); __m256 min = _mm256_set1_ps(SIGMOID_THRESHOLD_MIN);
__m256 tmp = _mm256_max_ps(a, min); __m256 tmp = _mm256_max_ps(a, min);
tmp = _mm256_min_ps(tmp, max); tmp = _mm256_min_ps(tmp, max);
tmp = _mm256_sub_ps(_mm256_set1_ps(0.0f), tmp); tmp = _mm256_sub_ps(_mm256_set1_ps(0.0f), tmp);
tmp = exp(tmp); tmp = exp(tmp);
tmp = _mm256_add_ps(_mm256_set1_ps(1.0f), tmp); tmp = _mm256_add_ps(_mm256_set1_ps(1.0f), tmp);
tmp = _mm256_div_ps(_mm256_set1_ps(1.0f), tmp); tmp = _mm256_div_ps(_mm256_set1_ps(1.0f), tmp);
return tmp; return tmp;
} }
__m256 tanh(const __m256 a) { __m256 tanh(const __m256 a) {
__m256 max = _mm256_set1_ps(EXP_MAX_INPUT); __m256 max = _mm256_set1_ps(EXP_MAX_INPUT);
__m256 tmp = _mm256_mul_ps(_mm256_set1_ps(-2.0f), a); __m256 tmp = _mm256_mul_ps(_mm256_set1_ps(-2.0f), a);
tmp = _mm256_min_ps(tmp, max); tmp = _mm256_min_ps(tmp, max);
tmp = exp(tmp); tmp = exp(tmp);
return _mm256_sub_ps( return _mm256_sub_ps(_mm256_div_ps(_mm256_set1_ps(2.0f),
_mm256_div_ps(_mm256_set1_ps(2.0f), _mm256_add_ps(_mm256_set1_ps(1.0f), tmp)),
_mm256_add_ps(_mm256_set1_ps(1.0f), tmp)), _mm256_set1_ps(1.0f)); _mm256_set1_ps(1.0f));
} }
__m256 linear(const __m256 a) { __m256 linear(const __m256 a) { return a; }
return a;
}
__m256 relu(const __m256 a, const __m256 b) { __m256 relu(const __m256 a, const __m256 b) {
return _mm256_mul_ps(a, return _mm256_mul_ps(
a,
_mm256_and_ps(_mm256_cmp_ps(b, _mm256_set1_ps(0.0f), _CMP_GT_OS), _mm256_and_ps(_mm256_cmp_ps(b, _mm256_set1_ps(0.0f), _CMP_GT_OS),
_mm256_set1_ps(1.0f))); _mm256_set1_ps(1.0f)));
} }
__m256 sigmoid(const __m256 a, const __m256 b) { __m256 sigmoid(const __m256 a, const __m256 b) {
return _mm256_mul_ps(_mm256_mul_ps(a, b), return _mm256_mul_ps(_mm256_mul_ps(a, b),
_mm256_sub_ps(_mm256_set1_ps(1.0f), b)); _mm256_sub_ps(_mm256_set1_ps(1.0f), b));
} }
__m256 tanh(const __m256 a, const __m256 b) { __m256 tanh(const __m256 a, const __m256 b) {
return _mm256_mul_ps(a, return _mm256_mul_ps(
_mm256_sub_ps(_mm256_set1_ps(1.0f), _mm256_mul_ps(b, b))); a, _mm256_sub_ps(_mm256_set1_ps(1.0f), _mm256_mul_ps(b, b)));
} }
__m256 linear(const __m256 a, const __m256 b) { __m256 linear(const __m256 a, const __m256 b) { return a; }
return a;
}
} // namespace hppl } // namespace hppl
...@@ -12,46 +12,33 @@ WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. ...@@ -12,46 +12,33 @@ WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
See the License for the specific language governing permissions and See the License for the specific language governing permissions and
limitations under the License. */ limitations under the License. */
#include <math.h> #include <math.h>
#include "hl_functions.h" #include "hl_functions.h"
namespace hppl { namespace hppl {
real relu(const real a) { real relu(const real a) { return a > 0.0f ? a : 0.0f; }
return a > 0.0f ? a : 0.0f;
} real sigmoid(const real a) {
const real min = SIGMOID_THRESHOLD_MIN;
real sigmoid(const real a) { const real max = SIGMOID_THRESHOLD_MAX;
const real min = SIGMOID_THRESHOLD_MIN; real tmp = (a < min) ? min : ((a > max) ? max : a);
const real max = SIGMOID_THRESHOLD_MAX; return 1.0 / (1.0 + exp(-tmp));
real tmp = (a < min) ? min : ((a > max) ? max : a); }
return 1.0 / (1.0 + exp(-tmp));
} real tanh(const real a) {
real tmp = -2.0 * a;
real tanh(const real a) { tmp = (tmp > EXP_MAX_INPUT) ? EXP_MAX_INPUT : tmp;
real tmp = -2.0 * a; return (2.0 / (1.0 + exp(tmp))) - 1.0;
tmp = (tmp > EXP_MAX_INPUT) ? EXP_MAX_INPUT : tmp; }
return (2.0 / (1.0 + exp(tmp))) - 1.0;
} real linear(const real a) { return a; }
real linear(const real a) { real relu(const real a, const real b) { return a * (b > 0.0f ? 1.0f : 0.0f); }
return a;
} real sigmoid(const real a, const real b) { return a * b * (1 - b); }
real relu(const real a, const real b) { real tanh(const real a, const real b) { return a * (1.0f - b * b); }
return a * (b > 0.0f ? 1.0f : 0.0f);
} real linear(const real a, const real b) { return a; }
real sigmoid(const real a, const real b) {
return a * b * (1 - b);
}
real tanh(const real a, const real b) {
return a * (1.0f - b * b);
}
real linear(const real a, const real b) {
return a;
}
} // namespace hppl } // namespace hppl
...@@ -12,7 +12,6 @@ WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. ...@@ -12,7 +12,6 @@ WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
See the License for the specific language governing permissions and See the License for the specific language governing permissions and
limitations under the License. */ limitations under the License. */
#include <sys/time.h> #include <sys/time.h>
#include <mutex> #include <mutex>
#include "hl_cuda.h" #include "hl_cuda.h"
...@@ -24,7 +23,7 @@ limitations under the License. */ ...@@ -24,7 +23,7 @@ limitations under the License. */
namespace dynload { namespace dynload {
std::once_flag cublas_dso_flag; std::once_flag cublas_dso_flag;
void* cublas_dso_handle = nullptr; void *cublas_dso_handle = nullptr;
/** /**
* The following macro definition can generate structs * The following macro definition can generate structs
...@@ -34,31 +33,30 @@ void* cublas_dso_handle = nullptr; ...@@ -34,31 +33,30 @@ void* cublas_dso_handle = nullptr;
* note: default dynamic linked libs * note: default dynamic linked libs
*/ */
#ifdef PADDLE_USE_DSO #ifdef PADDLE_USE_DSO
#define DYNAMIC_LOAD_CUBLAS_WRAP(__name) \ #define DYNAMIC_LOAD_CUBLAS_WRAP(__name) \
struct DynLoad__##__name { \ struct DynLoad__##__name { \
template <typename... Args> \ template <typename... Args> \
cublasStatus_t operator()(Args... args) { \ cublasStatus_t operator()(Args... args) { \
typedef cublasStatus_t (*cublasFunc)(Args...); \ typedef cublasStatus_t (*cublasFunc)(Args...); \
std::call_once(cublas_dso_flag, GetCublasDsoHandle, \ std::call_once(cublas_dso_flag, GetCublasDsoHandle, &cublas_dso_handle); \
&cublas_dso_handle); \ void *p_##__name = dlsym(cublas_dso_handle, #__name); \
void* p_##__name = dlsym(cublas_dso_handle, #__name); \ return reinterpret_cast<cublasFunc>(p_##__name)(args...); \
return reinterpret_cast<cublasFunc>(p_##__name)(args...); \ } \
} \
} __name; // struct DynLoad__##__name } __name; // struct DynLoad__##__name
#else #else
#define DYNAMIC_LOAD_CUBLAS_WRAP(__name) \ #define DYNAMIC_LOAD_CUBLAS_WRAP(__name) \
struct DynLoad__##__name { \ struct DynLoad__##__name { \
template <typename... Args> \ template <typename... Args> \
cublasStatus_t operator()(Args... args) { \ cublasStatus_t operator()(Args... args) { \
return __name(args...); \ return __name(args...); \
} \ } \
} __name; // struct DynLoad__##__name } __name; // struct DynLoad__##__name
#endif #endif
#define DYNAMIC_LOAD_CUBLAS_V2_WRAP(__name) \ #define DYNAMIC_LOAD_CUBLAS_V2_WRAP(__name) DYNAMIC_LOAD_CUBLAS_WRAP(__name)
DYNAMIC_LOAD_CUBLAS_WRAP(__name)
// include all needed cublas functions in HPPL // include all needed cublas functions in HPPL
// clang-format off
#define CUBLAS_BLAS_ROUTINE_EACH(__macro) \ #define CUBLAS_BLAS_ROUTINE_EACH(__macro) \
__macro(cublasSgemv) \ __macro(cublasSgemv) \
__macro(cublasDgemv) \ __macro(cublasDgemv) \
...@@ -88,41 +86,41 @@ CUBLAS_BLAS_ROUTINE_EACH(DYNAMIC_LOAD_CUBLAS_V2_WRAP) ...@@ -88,41 +86,41 @@ CUBLAS_BLAS_ROUTINE_EACH(DYNAMIC_LOAD_CUBLAS_V2_WRAP)
} /* namespace dynload */ } /* namespace dynload */
// clang-format on
#ifndef PADDLE_TYPE_DOUBLE #ifndef PADDLE_TYPE_DOUBLE
#define CUBLAS_GEAM dynload::cublasSgeam #define CUBLAS_GEAM dynload::cublasSgeam
#define CUBLAS_GEMV dynload::cublasSgemv #define CUBLAS_GEMV dynload::cublasSgemv
#define CUBLAS_GEMM dynload::cublasSgemm #define CUBLAS_GEMM dynload::cublasSgemm
#define CUBLAS_GETRF dynload::cublasSgetrfBatched #define CUBLAS_GETRF dynload::cublasSgetrfBatched
#define CUBLAS_GETRI dynload::cublasSgetriBatched #define CUBLAS_GETRI dynload::cublasSgetriBatched
#else #else
#define CUBLAS_GEAM dynload::cublasDgeam #define CUBLAS_GEAM dynload::cublasDgeam
#define CUBLAS_GEMV dynload::cublasDgemv #define CUBLAS_GEMV dynload::cublasDgemv
#define CUBLAS_GEMM dynload::cublasDgemm #define CUBLAS_GEMM dynload::cublasDgemm
#define CUBLAS_GETRF dynload::cublasDgetrfBatched #define CUBLAS_GETRF dynload::cublasDgetrfBatched
#define CUBLAS_GETRI dynload::cublasDgetriBatched #define CUBLAS_GETRI dynload::cublasDgetriBatched
#endif #endif
const char* hl_cublas_get_error_string(cublasStatus_t status) { const char *hl_cublas_get_error_string(cublasStatus_t status) {
switch(status) { switch (status) {
case CUBLAS_STATUS_NOT_INITIALIZED: case CUBLAS_STATUS_NOT_INITIALIZED:
return "[cublas status]: not initialized"; return "[cublas status]: not initialized";
case CUBLAS_STATUS_ALLOC_FAILED: case CUBLAS_STATUS_ALLOC_FAILED:
return "[cublas status]: allocate failed"; return "[cublas status]: allocate failed";
case CUBLAS_STATUS_INVALID_VALUE: case CUBLAS_STATUS_INVALID_VALUE:
return "[cublas status]: invalid value"; return "[cublas status]: invalid value";
case CUBLAS_STATUS_ARCH_MISMATCH: case CUBLAS_STATUS_ARCH_MISMATCH:
return "[cublas status]: arch mismatch"; return "[cublas status]: arch mismatch";
case CUBLAS_STATUS_MAPPING_ERROR: case CUBLAS_STATUS_MAPPING_ERROR:
return "[cublas status]: mapping error"; return "[cublas status]: mapping error";
case CUBLAS_STATUS_EXECUTION_FAILED: case CUBLAS_STATUS_EXECUTION_FAILED:
return "[cublas status]: execution failed"; return "[cublas status]: execution failed";
case CUBLAS_STATUS_INTERNAL_ERROR: case CUBLAS_STATUS_INTERNAL_ERROR:
return "[cublas status]: internal error"; return "[cublas status]: internal error";
case CUBLAS_STATUS_SUCCESS: case CUBLAS_STATUS_SUCCESS:
return "[cublas status]: success"; return "[cublas status]: success";
default: default:
return "[cublas status]: unknown error"; return "[cublas status]: unknown error";
} }
} }
...@@ -131,27 +129,21 @@ const char* hl_cublas_get_error_string(cublasStatus_t status) { ...@@ -131,27 +129,21 @@ const char* hl_cublas_get_error_string(cublasStatus_t status) {
* support << operator for more details error info. * support << operator for more details error info.
*/ */
cublasStatus_t g_cublasStat; cublasStatus_t g_cublasStat;
#define CHECK_CUBLAS(cublas_func) \ #define CHECK_CUBLAS(cublas_func) \
g_cublasStat = cublas_func; \ g_cublasStat = cublas_func; \
CHECK_EQ(CUBLAS_STATUS_SUCCESS, g_cublasStat) \ CHECK_EQ(CUBLAS_STATUS_SUCCESS, g_cublasStat) \
<< "Cublas Error: " \ << "Cublas Error: " << hl_cublas_get_error_string(g_cublasStat) << " "
<< hl_cublas_get_error_string(g_cublasStat) \
<< " "
void hl_cublas_init(cublasHandle_t *cublas_handle, cudaStream_t stream) { void hl_cublas_init(cublasHandle_t *cublas_handle, cudaStream_t stream) {
CHECK_CUBLAS(dynload::cublasCreate(cublas_handle)) CHECK_CUBLAS(dynload::cublasCreate(cublas_handle))
<< "[cublas init] Cublas create handle faild!"; << "[cublas init] Cublas create handle faild!";
CHECK_CUBLAS(dynload::cublasSetStream(*cublas_handle, stream)) CHECK_CUBLAS(dynload::cublasSetStream(*cublas_handle, stream))
<< "[cublas init] Cublas set stream faild!"; << "[cublas init] Cublas set stream faild!";
} }
void hl_matrix_transpose(real *A_d, void hl_matrix_transpose(
real *C_d, real *A_d, real *C_d, int dimM, int dimN, int lda, int ldc) {
int dimM,
int dimN,
int lda,
int ldc) {
real alpha = 1.0; real alpha = 1.0;
real beta = 0.0; real beta = 0.0;
...@@ -159,11 +151,18 @@ void hl_matrix_transpose(real *A_d, ...@@ -159,11 +151,18 @@ void hl_matrix_transpose(real *A_d,
CHECK_NOTNULL(C_d); CHECK_NOTNULL(C_d);
CHECK_CUBLAS(CUBLAS_GEAM(t_resource.handle, CHECK_CUBLAS(CUBLAS_GEAM(t_resource.handle,
CUBLAS_OP_T, CUBLAS_OP_N, CUBLAS_OP_T,
dimM, dimN, CUBLAS_OP_N,
&alpha, A_d, lda, dimM,
&beta, nullptr, dimM, dimN,
C_d, ldc)); &alpha,
A_d,
lda,
&beta,
nullptr,
dimM,
C_d,
ldc));
CHECK_SYNC("hl_matrix_transpose failed"); CHECK_SYNC("hl_matrix_transpose failed");
} }
...@@ -181,21 +180,20 @@ void hl_matrix_inverse(real *A_d, real *C_d, int dimN, int lda, int ldc) { ...@@ -181,21 +180,20 @@ void hl_matrix_inverse(real *A_d, real *C_d, int dimN, int lda, int ldc) {
real **inout_d = (real **)hl_malloc_device(sizeof(real *)); real **inout_d = (real **)hl_malloc_device(sizeof(real *));
hl_memcpy(inout_d, inout_h, sizeof(real *)); hl_memcpy(inout_d, inout_h, sizeof(real *));
int *pivot_d = (int *)hl_malloc_device(dimN*sizeof(int)); int *pivot_d = (int *)hl_malloc_device(dimN * sizeof(int));
int *info_d = (int *)t_resource.gpu_mem; int *info_d = (int *)t_resource.gpu_mem;
/* Note: cublasSgetrfBatched is used to calculate a number of /* Note: cublasSgetrfBatched is used to calculate a number of
small-sized matrices. There may be a better way to reconstruct small-sized matrices. There may be a better way to reconstruct
the API for better performance. the API for better performance.
*/ */
CHECK_CUBLAS(CUBLAS_GETRF(t_resource.handle, CHECK_CUBLAS(
dimN, inout_d, lda, pivot_d, CUBLAS_GETRF(t_resource.handle, dimN, inout_d, lda, pivot_d, info_d, 1));
info_d, 1));
int info_h; int info_h;
hl_memcpy(&info_h, info_d, sizeof(int)); hl_memcpy(&info_h, info_d, sizeof(int));
if (info_h != 0) { if (info_h != 0) {
LOG(FATAL) << "Factorization of matrix failed: matrix may be singular.\n"; LOG(FATAL) << "Factorization of matrix failed: matrix may be singular.\n";
} }
/* Step 2: Compute the inverse of the matrix given its LU decomposition */ /* Step 2: Compute the inverse of the matrix given its LU decomposition */
...@@ -204,27 +202,40 @@ void hl_matrix_inverse(real *A_d, real *C_d, int dimN, int lda, int ldc) { ...@@ -204,27 +202,40 @@ void hl_matrix_inverse(real *A_d, real *C_d, int dimN, int lda, int ldc) {
hl_memcpy(out_d, out_h, sizeof(real *)); hl_memcpy(out_d, out_h, sizeof(real *));
CHECK_CUBLAS(CUBLAS_GETRI(t_resource.handle, CHECK_CUBLAS(CUBLAS_GETRI(t_resource.handle,
dimN, (const real **)inout_d, lda, pivot_d, dimN,
out_d, ldc, info_d, 1)); (const real **)inout_d,
lda,
pivot_d,
out_d,
ldc,
info_d,
1));
hl_memcpy(&info_h, info_d, sizeof(int)); hl_memcpy(&info_h, info_d, sizeof(int));
if (info_h != 0) { if (info_h != 0) {
LOG(FATAL) << "Inversion of matrix failed: matrix may be singular.\n"; LOG(FATAL) << "Inversion of matrix failed: matrix may be singular.\n";
} }
hl_free_mem_device(inout_d); hl_free_mem_device(inout_d);
hl_free_mem_device(pivot_d); hl_free_mem_device(pivot_d);
hl_free_mem_device(out_d); hl_free_mem_device(out_d);
CHECK_SYNC("hl_matrix_inverse failed"); CHECK_SYNC("hl_matrix_inverse failed");
} }
void hl_matrix_mul(real *A_d, hl_trans_op_t transa, void hl_matrix_mul(real *A_d,
real *B_d, hl_trans_op_t transb, hl_trans_op_t transa,
real *B_d,
hl_trans_op_t transb,
real *C_d, real *C_d,
int dimM, int dimN, int dimK, int dimM,
real alpha, real beta, int dimN,
int lda, int ldb, int ldc) { int dimK,
real alpha,
real beta,
int lda,
int ldb,
int ldc) {
CHECK_NOTNULL(A_d); CHECK_NOTNULL(A_d);
CHECK_NOTNULL(B_d); CHECK_NOTNULL(B_d);
CHECK_NOTNULL(C_d); CHECK_NOTNULL(C_d);
...@@ -232,8 +243,8 @@ void hl_matrix_mul(real *A_d, hl_trans_op_t transa, ...@@ -232,8 +243,8 @@ void hl_matrix_mul(real *A_d, hl_trans_op_t transa,
if (dimN == 1 && dimM != 1 && dimK != 1 && transb == HPPL_OP_N) { if (dimN == 1 && dimM != 1 && dimK != 1 && transb == HPPL_OP_N) {
int m = (transa == HPPL_OP_N) ? dimM : dimK; int m = (transa == HPPL_OP_N) ? dimM : dimK;
int n = (transa == HPPL_OP_N) ? dimK : dimM; int n = (transa == HPPL_OP_N) ? dimK : dimM;
hl_matrix_mul_vector(A_d, transa, B_d, C_d, m, n, hl_matrix_mul_vector(
alpha, beta, lda, ldb, ldc); A_d, transa, B_d, C_d, m, n, alpha, beta, lda, ldb, ldc);
return; return;
} }
...@@ -241,8 +252,7 @@ void hl_matrix_mul(real *A_d, hl_trans_op_t transa, ...@@ -241,8 +252,7 @@ void hl_matrix_mul(real *A_d, hl_trans_op_t transa,
int m = (transb == HPPL_OP_N) ? dimK : dimN; int m = (transb == HPPL_OP_N) ? dimK : dimN;
int n = (transb == HPPL_OP_N) ? dimN : dimK; int n = (transb == HPPL_OP_N) ? dimN : dimK;
hl_trans_op_t trans = (transb == HPPL_OP_N) ? HPPL_OP_T : HPPL_OP_N; hl_trans_op_t trans = (transb == HPPL_OP_N) ? HPPL_OP_T : HPPL_OP_N;
hl_matrix_mul_vector(B_d, trans, A_d, C_d, m, n, hl_matrix_mul_vector(B_d, trans, A_d, C_d, m, n, alpha, beta, ldb, 1, 1);
alpha, beta, ldb, 1, 1);
return; return;
} }
...@@ -251,26 +261,47 @@ void hl_matrix_mul(real *A_d, hl_trans_op_t transa, ...@@ -251,26 +261,47 @@ void hl_matrix_mul(real *A_d, hl_trans_op_t transa,
stat = CUBLAS_GEMM(t_resource.handle, stat = CUBLAS_GEMM(t_resource.handle,
CUBLAS_OP_N, CUBLAS_OP_N,
CUBLAS_OP_N, CUBLAS_OP_N,
dimN, dimM, dimK, dimN,
&alpha, B_d, ldb, dimM,
A_d, lda, dimK,
&beta, C_d, ldc); &alpha,
B_d,
ldb,
A_d,
lda,
&beta,
C_d,
ldc);
} else if ((HPPL_OP_T == transa) && (HPPL_OP_N == transb)) { } else if ((HPPL_OP_T == transa) && (HPPL_OP_N == transb)) {
stat = CUBLAS_GEMM(t_resource.handle, stat = CUBLAS_GEMM(t_resource.handle,
CUBLAS_OP_N, CUBLAS_OP_N,
CUBLAS_OP_T, CUBLAS_OP_T,
dimN, dimM, dimK, dimN,
&alpha, B_d, ldb, dimM,
A_d, lda, dimK,
&beta, C_d, ldc); &alpha,
B_d,
ldb,
A_d,
lda,
&beta,
C_d,
ldc);
} else if ((HPPL_OP_N == transa) && (HPPL_OP_T == transb)) { } else if ((HPPL_OP_N == transa) && (HPPL_OP_T == transb)) {
stat = CUBLAS_GEMM(t_resource.handle, stat = CUBLAS_GEMM(t_resource.handle,
CUBLAS_OP_T, CUBLAS_OP_T,
CUBLAS_OP_N, CUBLAS_OP_N,
dimN, dimM, dimK, dimN,
&alpha, B_d, ldb, dimM,
A_d, lda, dimK,
&beta, C_d, ldc); &alpha,
B_d,
ldb,
A_d,
lda,
&beta,
C_d,
ldc);
} else { } else {
LOG(FATAL) << "parameter transa error!"; LOG(FATAL) << "parameter transa error!";
} }
...@@ -278,24 +309,46 @@ void hl_matrix_mul(real *A_d, hl_trans_op_t transa, ...@@ -278,24 +309,46 @@ void hl_matrix_mul(real *A_d, hl_trans_op_t transa,
CHECK_SYNC("hl_matrix_mul failed"); CHECK_SYNC("hl_matrix_mul failed");
} }
void hl_matrix_mul(real *A_d, hl_trans_op_t transa, void hl_matrix_mul(real *A_d,
real *B_d, hl_trans_op_t transb, hl_trans_op_t transa,
real *B_d,
hl_trans_op_t transb,
real *C_d, real *C_d,
int dimM, int dimN, int dimK, int dimM,
real alpha, real beta) { int dimN,
int dimK,
real alpha,
real beta) {
int lda = (HPPL_OP_N == transa) ? dimK : dimM; int lda = (HPPL_OP_N == transa) ? dimK : dimM;
int ldb = (HPPL_OP_N == transb) ? dimN : dimK; int ldb = (HPPL_OP_N == transb) ? dimN : dimK;
int ldc = dimN; int ldc = dimN;
hl_matrix_mul(A_d, transa, B_d, transb, C_d, dimM, dimN, hl_matrix_mul(A_d,
dimK, alpha, beta, lda, ldb, ldc); transa,
B_d,
transb,
C_d,
dimM,
dimN,
dimK,
alpha,
beta,
lda,
ldb,
ldc);
} }
void hl_matrix_mul_vector(real *A_d, hl_trans_op_t trans, void hl_matrix_mul_vector(real *A_d,
real *B_d, real *C_d, hl_trans_op_t trans,
int dimM, int dimN, real *B_d,
real alpha, real beta, real *C_d,
int lda, int incb, int incc) { int dimM,
int dimN,
real alpha,
real beta,
int lda,
int incb,
int incc) {
CHECK_NOTNULL(A_d); CHECK_NOTNULL(A_d);
CHECK_NOTNULL(B_d); CHECK_NOTNULL(B_d);
CHECK_NOTNULL(C_d); CHECK_NOTNULL(C_d);
...@@ -304,21 +357,29 @@ void hl_matrix_mul_vector(real *A_d, hl_trans_op_t trans, ...@@ -304,21 +357,29 @@ void hl_matrix_mul_vector(real *A_d, hl_trans_op_t trans,
if (HPPL_OP_N == trans) { if (HPPL_OP_N == trans) {
stat = CUBLAS_GEMV(t_resource.handle, stat = CUBLAS_GEMV(t_resource.handle,
CUBLAS_OP_T, CUBLAS_OP_T,
dimN, dimM, dimN,
dimM,
&alpha, &alpha,
A_d, lda, A_d,
B_d, incb, lda,
B_d,
incb,
&beta, &beta,
C_d, incc); C_d,
incc);
} else if (HPPL_OP_T == trans) { } else if (HPPL_OP_T == trans) {
stat = CUBLAS_GEMV(t_resource.handle, stat = CUBLAS_GEMV(t_resource.handle,
CUBLAS_OP_N, CUBLAS_OP_N,
dimN, dimM, dimN,
dimM,
&alpha, &alpha,
A_d, lda, A_d,
B_d, incb, lda,
B_d,
incb,
&beta, &beta,
C_d, incc); C_d,
incc);
} else { } else {
LOG(FATAL) << "parameter transa error!"; LOG(FATAL) << "parameter transa error!";
} }
...@@ -327,10 +388,14 @@ void hl_matrix_mul_vector(real *A_d, hl_trans_op_t trans, ...@@ -327,10 +388,14 @@ void hl_matrix_mul_vector(real *A_d, hl_trans_op_t trans,
CHECK_SYNC("hl_matrix_mul_vector"); CHECK_SYNC("hl_matrix_mul_vector");
} }
void hl_matrix_mul_vector(real *A_d, hl_trans_op_t trans, void hl_matrix_mul_vector(real *A_d,
real *B_d, real *C_d, hl_trans_op_t trans,
int dimM, int dimN, real *B_d,
real alpha, real beta) { real *C_d,
hl_matrix_mul_vector(A_d, trans, B_d, C_d, dimM, dimN, int dimM,
alpha, beta, dimN, 1, 1); int dimN,
real alpha,
real beta) {
hl_matrix_mul_vector(
A_d, trans, B_d, C_d, dimM, dimN, alpha, beta, dimN, 1, 1);
} }
此差异已折叠。
此差异已折叠。
此差异已折叠。
此差异已折叠。
...@@ -12,24 +12,15 @@ WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. ...@@ -12,24 +12,15 @@ WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
See the License for the specific language governing permissions and See the License for the specific language governing permissions and
limitations under the License. */ limitations under the License. */
#include "avx_mathfun.h" #include "avx_mathfun.h"
namespace hppl { namespace hppl {
__m256 exp(__m256 a) { __m256 exp(__m256 a) { return exp256_ps(a); }
return exp256_ps(a);
}
__m256 log(__m256 a) { __m256 log(__m256 a) { return log256_ps(a); }
return log256_ps(a);
}
__m256 sin(__m256 a) { __m256 sin(__m256 a) { return sin256_ps(a); }
return sin256_ps(a);
}
__m256 cos(__m256 a) { __m256 cos(__m256 a) { return cos256_ps(a); }
return cos256_ps(a);
}
} // namespace hppl } // namespace hppl
...@@ -12,7 +12,6 @@ WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. ...@@ -12,7 +12,6 @@ WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
See the License for the specific language governing permissions and See the License for the specific language governing permissions and
limitations under the License. */ limitations under the License. */
#include <chrono> #include <chrono>
#include <stdlib.h> #include <stdlib.h>
#include <iostream> #include <iostream>
...@@ -21,8 +20,7 @@ limitations under the License. */ ...@@ -21,8 +20,7 @@ limitations under the License. */
using std::chrono::high_resolution_clock; using std::chrono::high_resolution_clock;
int64_t getCurrentTimeStick() { int64_t getCurrentTimeStick() {
high_resolution_clock::time_point tp = high_resolution_clock::now(); high_resolution_clock::time_point tp = high_resolution_clock::now();
high_resolution_clock::duration dtn = tp.time_since_epoch(); high_resolution_clock::duration dtn = tp.time_since_epoch();
return dtn.count(); return dtn.count();
} }
...@@ -12,7 +12,6 @@ WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. ...@@ -12,7 +12,6 @@ WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
See the License for the specific language governing permissions and See the License for the specific language governing permissions and
limitations under the License. */ limitations under the License. */
#pragma once #pragma once
#include <string> #include <string>
#include <vector> #include <vector>
......
...@@ -12,7 +12,6 @@ WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. ...@@ -12,7 +12,6 @@ WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
See the License for the specific language governing permissions and See the License for the specific language governing permissions and
limitations under the License. */ limitations under the License. */
#include "DataProvider.h" #include "DataProvider.h"
#include "paddle/utils/Util.h" #include "paddle/utils/Util.h"
...@@ -57,7 +56,7 @@ void BufferBatch::clone(DataBatch* srcBatch, bool useGpu) { ...@@ -57,7 +56,7 @@ void BufferBatch::clone(DataBatch* srcBatch, bool useGpu) {
} }
} }
DoubleBuffer::DoubleBuffer(DataProvider *dataPool, DoubleBuffer::DoubleBuffer(DataProvider* dataPool,
bool useGpu, bool useGpu,
int64_t batchSize) { int64_t batchSize) {
batchSize_ = batchSize; batchSize_ = batchSize;
...@@ -155,7 +154,7 @@ void DoubleBuffer::startAsyncLoad() { ...@@ -155,7 +154,7 @@ void DoubleBuffer::startAsyncLoad() {
} }
ClassRegistrar<DataProvider, DataConfig, ModelConfig, bool> ClassRegistrar<DataProvider, DataConfig, ModelConfig, bool>
DataProvider::registrar_; DataProvider::registrar_;
DataProvider* DataProvider::create(const DataConfig& config, DataProvider* DataProvider::create(const DataConfig& config,
const ModelConfig& modelConfig, const ModelConfig& modelConfig,
...@@ -182,7 +181,8 @@ int64_t DataProvider::getNextBatch(int64_t size, DataBatch* batch) { ...@@ -182,7 +181,8 @@ int64_t DataProvider::getNextBatch(int64_t size, DataBatch* batch) {
for (int i = 0; i < config_.constant_slots_size(); ++i) { for (int i = 0; i < config_.constant_slots_size(); ++i) {
MemoryHandlePtr handle = MemoryHandlePtr handle =
constantSlots[i] ? constantSlots[i]->getMemoryHandle() : nullptr; constantSlots[i] ? constantSlots[i]->getMemoryHandle() : nullptr;
Matrix::resizeOrCreate(constantSlots[i], batchSize, Matrix::resizeOrCreate(constantSlots[i],
batchSize,
1, // = width 1, // = width
false, // = trans false, // = trans
useGpu_); // = useGpu useGpu_); // = useGpu
...@@ -216,7 +216,8 @@ void DataProvider::initAsyncLoader() { ...@@ -216,7 +216,8 @@ void DataProvider::initAsyncLoader() {
} }
SimpleDataProviderBase::SimpleDataProviderBase(const DataConfig& config, SimpleDataProviderBase::SimpleDataProviderBase(const DataConfig& config,
bool useGpu, bool withInfo) bool useGpu,
bool withInfo)
: DataProvider(config, useGpu) { : DataProvider(config, useGpu) {
/* initialize the size of a sample, and the buffer */ /* initialize the size of a sample, and the buffer */
sampleDim_ = config_.feat_dim() * (2 * config_.context_len() + 1); sampleDim_ = config_.feat_dim() * (2 * config_.context_len() + 1);
...@@ -337,7 +338,8 @@ int64_t SimpleDataProviderBase::fillBuffer() { ...@@ -337,7 +338,8 @@ int64_t SimpleDataProviderBase::fillBuffer() {
sampleNumInBuf_ = sampleNumInBuf_ =
n + fillBufferImp(hInputDataBuf_->getData() + n * sampleDim_, n + fillBufferImp(hInputDataBuf_->getData() + n * sampleDim_,
hInputLabelBuf_->getData() + n, hInputLabelBuf_->getData() + n,
hInputInfoBuf_->getData() + n, bufferCapacity_ - n); hInputInfoBuf_->getData() + n,
bufferCapacity_ - n);
/* for stachastic gradient training */ /* for stachastic gradient training */
if (!skipShuffle_) { if (!skipShuffle_) {
...@@ -357,11 +359,14 @@ SimpleDataProvider::SimpleDataProvider(const DataConfig& config, bool useGpu) ...@@ -357,11 +359,14 @@ SimpleDataProvider::SimpleDataProvider(const DataConfig& config, bool useGpu)
SimpleDataProvider::~SimpleDataProvider() {} SimpleDataProvider::~SimpleDataProvider() {}
int64_t SimpleDataProvider::fillBufferImp(real* data, int* label, int* info, int64_t SimpleDataProvider::fillBufferImp(real* data,
int* label,
int* info,
int64_t size) { int64_t size) {
(void)info; (void)info;
int64_t n = std::min<int64_t>(labels_.size() - currentSampleIndex_, size); int64_t n = std::min<int64_t>(labels_.size() - currentSampleIndex_, size);
memcpy(data, &data_[currentSampleIndex_ * sampleDim_], memcpy(data,
&data_[currentSampleIndex_ * sampleDim_],
n * sampleDim_ * sizeof(real)); n * sampleDim_ * sizeof(real));
memcpy(label, &labels_[currentSampleIndex_], sizeof(int) * n); memcpy(label, &labels_[currentSampleIndex_], sizeof(int) * n);
currentSampleIndex_ += n; currentSampleIndex_ += n;
......
...@@ -12,7 +12,6 @@ WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. ...@@ -12,7 +12,6 @@ WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
See the License for the specific language governing permissions and See the License for the specific language governing permissions and
limitations under the License. */ limitations under the License. */
#pragma once #pragma once
#include "DataProvider.h" #include "DataProvider.h"
...@@ -65,8 +64,8 @@ void DataProviderGroup<T>::reset() { ...@@ -65,8 +64,8 @@ void DataProviderGroup<T>::reset() {
provider_ = nullptr; provider_ = nullptr;
// shuffle file list // shuffle file list
std::shuffle(fileList_.begin(), fileList_.end(), std::shuffle(
ThreadLocalRandomEngine::get()); fileList_.begin(), fileList_.end(), ThreadLocalRandomEngine::get());
startLoader(); startLoader();
DataProvider::reset(); DataProvider::reset();
...@@ -113,8 +112,9 @@ void DataProviderGroup<T>::startLoader() { ...@@ -113,8 +112,9 @@ void DataProviderGroup<T>::startLoader() {
size_t endPos = std::min(fileList_.size(), startPos + loadFileCount); size_t endPos = std::min(fileList_.size(), startPos + loadFileCount);
std::vector<std::string> fileVec(fileList_.begin() + startPos, std::vector<std::string> fileVec(fileList_.begin() + startPos,
fileList_.begin() + endPos); fileList_.begin() + endPos);
loader_->addJob([this, fileVec]() loader_->addJob([this, fileVec]() -> ProviderPtrType {
-> ProviderPtrType { return this->loadFile(fileVec); }); return this->loadFile(fileVec);
});
} }
loader_->stopAddJob(); loader_->stopAddJob();
} }
......
...@@ -12,7 +12,6 @@ WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. ...@@ -12,7 +12,6 @@ WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
See the License for the specific language governing permissions and See the License for the specific language governing permissions and
limitations under the License. */ limitations under the License. */
#pragma once #pragma once
#include "DataProvider.h" #include "DataProvider.h"
......
...@@ -12,7 +12,6 @@ WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. ...@@ -12,7 +12,6 @@ WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
See the License for the specific language governing permissions and See the License for the specific language governing permissions and
limitations under the License. */ limitations under the License. */
#include "AddtoLayer.h" #include "AddtoLayer.h"
#include "paddle/utils/Logging.h" #include "paddle/utils/Logging.h"
......
此差异已折叠。
此差异已折叠。
此差异已折叠。
此差异已折叠。
此差异已折叠。
此差异已折叠。
此差异已折叠。
此差异已折叠。
此差异已折叠。
此差异已折叠。
此差异已折叠。
此差异已折叠。
此差异已折叠。
此差异已折叠。
此差异已折叠。
此差异已折叠。
此差异已折叠。
此差异已折叠。
此差异已折叠。
此差异已折叠。
此差异已折叠。
此差异已折叠。
此差异已折叠。
此差异已折叠。
此差异已折叠。
此差异已折叠。
此差异已折叠。
此差异已折叠。
此差异已折叠。
此差异已折叠。
此差异已折叠。
此差异已折叠。
此差异已折叠。
此差异已折叠。
此差异已折叠。
此差异已折叠。
此差异已折叠。
此差异已折叠。
此差异已折叠。
此差异已折叠。
此差异已折叠。
此差异已折叠。
此差异已折叠。
此差异已折叠。
此差异已折叠。
此差异已折叠。
此差异已折叠。
此差异已折叠。
此差异已折叠。
此差异已折叠。
此差异已折叠。
此差异已折叠。
此差异已折叠。
此差异已折叠。
此差异已折叠。
此差异已折叠。
此差异已折叠。
此差异已折叠。
此差异已折叠。
此差异已折叠。
此差异已折叠。
此差异已折叠。
此差异已折叠。
此差异已折叠。
此差异已折叠。
此差异已折叠。
此差异已折叠。
此差异已折叠。
此差异已折叠。
此差异已折叠。
此差异已折叠。
此差异已折叠。
此差异已折叠。
此差异已折叠。
此差异已折叠。
此差异已折叠。
此差异已折叠。
此差异已折叠。
此差异已折叠。
此差异已折叠。
此差异已折叠。
此差异已折叠。
此差异已折叠。
此差异已折叠。
此差异已折叠。
此差异已折叠。
此差异已折叠。
此差异已折叠。
此差异已折叠。
此差异已折叠。
此差异已折叠。
此差异已折叠。
此差异已折叠。
此差异已折叠。
此差异已折叠。
此差异已折叠。
此差异已折叠。
此差异已折叠。
此差异已折叠。
此差异已折叠。
此差异已折叠。
此差异已折叠。
此差异已折叠。
此差异已折叠。
此差异已折叠。
此差异已折叠。
此差异已折叠。
此差异已折叠。
此差异已折叠。
此差异已折叠。
此差异已折叠。
此差异已折叠。
此差异已折叠。
此差异已折叠。
此差异已折叠。
此差异已折叠。
此差异已折叠。
此差异已折叠。
此差异已折叠。
此差异已折叠。
此差异已折叠。
此差异已折叠。
此差异已折叠。
此差异已折叠。
此差异已折叠。
此差异已折叠。
此差异已折叠。
此差异已折叠。
此差异已折叠。
Markdown is supported
0% .
You are about to add 0 people to the discussion. Proceed with caution.
先完成此消息的编辑!
想要评论请 注册