Skip to content

Commit 3a782f5

Browse files
byshiuebyshiue
committed
Add the v2.0 codes from DeepLearningExamples/FasterTransformer/v2
1 parent 2d67b97 commit 3a782f5

119 files changed

Lines changed: 46320 additions & 9121 deletions

File tree

Some content is hidden

Large Commits have some content hidden by default. Use the searchbox below for content that may be hidden.

‎.gitignore‎

Lines changed: 4 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,4 @@
1+
*~
2+
*.o
3+
build/
4+
*.pyc

‎CMakeLists.txt‎

Lines changed: 23 additions & 11 deletions
Original file line numberDiff line numberDiff line change
@@ -1,4 +1,4 @@
1-
# Copyright (c) 2019, NVIDIA CORPORATION. All rights reserved.
1+
# Copyright (c) 2020, NVIDIA CORPORATION. All rights reserved.
22
#
33
# Licensed under the Apache License, Version 2.0 (the "License");
44
# you may not use this file except in compliance with the License.
@@ -39,6 +39,10 @@ list(APPEND CMAKE_MODULE_PATH ${CUDA_PATH}/lib64)
3939
find_package(CUDA REQUIRED)
4040

4141
# setting compiler flags
42+
set(CMAKE_C_FLAGS "${CMAKE_C_FLAGS}")
43+
set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS}")
44+
set(CMAKE_CUDA_FLAGS "${CMAKE_CUDA_FLAGS} -Xcompiler -Wall")
45+
4246
if (SM STREQUAL 70 OR
4347
SM STREQUAL 75 OR
4448
SM STREQUAL 61 OR
@@ -49,15 +53,21 @@ set(CMAKE_CUDA_FLAGS "${CMAKE_CUDA_FLAGS} -gencode=arch=compute_${SM},code=\\\"s
4953
set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -DWMMA")
5054
set(CMAKE_CUDA_FLAGS "${CMAKE_CUDA_FLAGS} -DWMMA")
5155
endif()
52-
53-
set(CMAKE_C_FLAGS "${CMAKE_C_FLAGS}")
54-
set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS}")
55-
set(CMAKE_CUDA_FLAGS "${CMAKE_CUDA_FLAGS} -Xcompiler -Wall")
5656
message("-- Assign GPU architecture (sm=${SM})")
57+
5758
else()
58-
set(CMAKE_CUDA_FLAGS "${CMAKE_CUDA_FLAGS} -gencode=arch=compute_60,code=\\\"sm_60,compute_60\\\" -rdc=true")
59-
message("-- Unknown or unsupported GPU architecture (set sm=60)")
59+
set(CMAKE_CUDA_FLAGS "${CMAKE_CUDA_FLAGS} \
60+
-gencode=arch=compute_60,code=\\\"sm_60,compute_60\\\" \
61+
-gencode=arch=compute_61,code=\\\"sm_61,compute_61\\\" \
62+
-gencode=arch=compute_70,code=\\\"sm_70,compute_70\\\" \
63+
-gencode=arch=compute_75,code=\\\"sm_75,compute_75\\\" \
64+
-rdc=true")
65+
set(CMAKE_C_FLAGS "${CMAKE_C_FLAGS} -DWMMA")
66+
set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -DWMMA")
67+
set(CMAKE_CUDA_FLAGS "${CMAKE_CUDA_FLAGS} -DWMMA")
68+
message("-- Assign GPU architecture (sm=60,61,70,75)")
6069
endif()
70+
6171
set(CMAKE_C_FLAGS_DEBUG "${CMAKE_C_FLAGS_DEBUG} -Wall -O0")
6272
set(CMAKE_CXX_FLAGS_DEBUG "${CMAKE_CXX_FLAGS_DEBUG} -Wall -O0")
6373
set(CMAKE_CUDA_FLAGS_DEBUG "${CMAKE_CUDA_FLAGS_DEBUG} -O0 -G -Xcompiler -Wall")
@@ -72,8 +82,8 @@ if(CMAKE_CXX_STANDARD STREQUAL "11")
7282
set(CMAKE_CUDA_FLAGS "${CMAKE_CUDA_FLAGS} --std=c++11")
7383
endif()
7484

75-
set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -O3")
76-
set(CMAKE_CUDA_FLAGS "${CMAKE_CUDA_FLAGS} -Xcompiler -O3")
85+
set(CMAKE_CXX_FLAGS_RELEASE "${CMAKE_CXX_FLAGS_RELEASE} -O3")
86+
set(CMAKE_CUDA_FLAGS_RELEASE "${CMAKE_CUDA_FLAGS_RELEASE} -Xcompiler -O3")
7787

7888
set(CMAKE_ARCHIVE_OUTPUT_DIRECTORY ${CMAKE_BINARY_DIR}/lib)
7989
set(CMAKE_LIBRARY_OUTPUT_DIRECTORY ${CMAKE_BINARY_DIR}/lib)
@@ -110,12 +120,14 @@ add_subdirectory(tools)
110120
add_subdirectory(fastertransformer)
111121
add_subdirectory(sample)
112122

113-
114123
if(BUILD_TF)
115124
add_custom_target(copy ALL COMMENT "Copying tensorflow test scripts")
116125
add_custom_command(TARGET copy
117126
POST_BUILD
118127
COMMAND cp ${PROJECT_SOURCE_DIR}/sample/tensorflow/*.py ${PROJECT_SOURCE_DIR}/build/
119-
)
128+
COMMAND cp ${PROJECT_SOURCE_DIR}/sample/tensorflow/utils ${PROJECT_SOURCE_DIR}/build/ -r
129+
COMMAND cp ${PROJECT_SOURCE_DIR}/sample/tensorflow/scripts ${PROJECT_SOURCE_DIR}/build/ -r
130+
COMMAND cp ${PROJECT_SOURCE_DIR}/sample/tensorflow_bert ${PROJECT_SOURCE_DIR}/build/ -r
131+
)
120132
endif()
121133

‎README.md‎

Lines changed: 867 additions & 76 deletions
Large diffs are not rendered by default.

‎fastertransformer/CMakeLists.txt‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1,4 +1,4 @@
1-
# Copyright (c) 2019, NVIDIA CORPORATION. All rights reserved.
1+
# Copyright (c) 2020, NVIDIA CORPORATION. All rights reserved.
22
#
33
# Licensed under the Apache License, Version 2.0 (the "License");
44
# you may not use this file except in compliance with the License.

‎fastertransformer/allocator.h‎

Lines changed: 49 additions & 27 deletions
Original file line numberDiff line numberDiff line change
@@ -1,5 +1,5 @@
11
/*
2-
* Copyright (c) 2019, NVIDIA CORPORATION. All rights reserved.
2+
* Copyright (c) 2020, NVIDIA CORPORATION. All rights reserved.
33
*
44
* Licensed under the Apache License, Version 2.0 (the "License");
55
* you may not use this file except in compliance with the License.
@@ -35,34 +35,39 @@
3535
#include "third_party/eigen3/unsupported/Eigen/CXX11/Tensor"
3636
#endif
3737

38-
namespace fastertransformer{
38+
namespace fastertransformer
39+
{
3940

40-
41-
class IAllocator{
42-
public:
43-
virtual void* malloc(size_t size) const = 0;
44-
virtual void free(void* ptr) const = 0;
41+
class IAllocator
42+
{
43+
public:
44+
virtual void *malloc(size_t size) const = 0;
45+
virtual void free(void *ptr) const = 0;
4546
};
4647

47-
template<AllocatorType AllocType_>
48+
template <AllocatorType AllocType_>
4849
class Allocator;
4950

50-
template<>
51-
class Allocator<AllocatorType::CUDA> : public IAllocator{
51+
template <>
52+
class Allocator<AllocatorType::CUDA> : public IAllocator
53+
{
5254
const int device_id_;
53-
public:
54-
Allocator(int device_id): device_id_(device_id){}
5555

56-
void* malloc(size_t size) const {
57-
void* ptr = nullptr;
56+
public:
57+
Allocator(int device_id) : device_id_(device_id) {}
58+
59+
void *malloc(size_t size) const
60+
{
61+
void *ptr = nullptr;
5862
int o_device = 0;
5963
check_cuda_error(get_set_device(device_id_, &o_device));
6064
check_cuda_error(cudaMalloc(&ptr, size));
6165
check_cuda_error(get_set_device(o_device));
6266
return ptr;
6367
}
64-
65-
void free(void* ptr) const {
68+
69+
void free(void *ptr) const
70+
{
6671
int o_device = 0;
6772
check_cuda_error(get_set_device(device_id_, &o_device));
6873
check_cuda_error(cudaFree(ptr));
@@ -71,34 +76,51 @@ class Allocator<AllocatorType::CUDA> : public IAllocator{
7176
}
7277
};
7378

74-
7579
//TODO: allocator of TensorFlow
7680
// You can add context to constructor
7781
#ifdef GOOGLE_CUDA
7882
using namespace tensorflow;
79-
template<>
80-
class Allocator<AllocatorType::TF> : public IAllocator{
83+
template <>
84+
class Allocator<AllocatorType::TF> : public IAllocator
85+
{
8186
OpKernelContext *context_;
82-
public:
83-
Allocator(OpKernelContext *context): context_(context){}
87+
std::vector<Tensor> *allocated_tensor_vector;
88+
89+
public:
90+
Allocator(OpKernelContext *context) : context_(context)
91+
{
92+
allocated_tensor_vector = new std::vector<Tensor>;
93+
}
8494

85-
void* malloc(size_t size) const {
95+
void *malloc(size_t size) const
96+
{
8697
Tensor buf;
8798
long long int buf_size = (long long int)size;
8899
tensorflow::Status status = context_->allocate_temp(DT_UINT8, TensorShape{buf_size}, &buf);
100+
allocated_tensor_vector->push_back(buf);
89101

90-
if(status != tensorflow::Status::OK())
102+
if (status != tensorflow::Status::OK())
91103
throw std::runtime_error("TF error: context->allocate_temp failed");
92104

93105
auto flat = buf.flat<uint8>();
94-
void* ptr = (void*)flat.data();
106+
void *ptr = (void *)flat.data();
107+
cudaMemset(ptr, 0, buf_size);
95108
return ptr;
96109
}
97-
98-
void free(void* ptr) const {
110+
111+
void free(void *ptr) const
112+
{
113+
#ifndef NDEBUG
99114
printf("call from allocator free\n");
115+
#endif
100116
return;
101117
}
118+
119+
~Allocator()
120+
{
121+
allocated_tensor_vector->clear();
122+
delete allocated_tensor_vector;
123+
}
102124
};
103125
#endif
104-
}//namespace fastertransformer
126+
} //namespace fastertransformer

‎fastertransformer/beamsearch_opennmt.h‎

Lines changed: 121 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,121 @@
1+
/*
2+
* Copyright (c) 2020, NVIDIA CORPORATION. All rights reserved.
3+
*
4+
* Licensed under the Apache License, Version 2.0 (the "License");
5+
* you may not use this file except in compliance with the License.
6+
* You may obtain a copy of the License at
7+
*
8+
* http://www.apache.org/licenses/LICENSE-2.0
9+
*
10+
* Unless required by applicable law or agreed to in writing, software
11+
* distributed under the License is distributed on an "AS IS" BASIS,
12+
* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
13+
* See the License for the specific language governing permissions and
14+
* limitations under the License.
15+
*/
16+
17+
/**
18+
* BeamSearch OpenNMT
19+
**/
20+
21+
#pragma once
22+
23+
#include <cuda_runtime.h>
24+
#include "fastertransformer/allocator.h"
25+
#include "fastertransformer/cuda/cuda_kernels.h"
26+
#include "fastertransformer/cuda/open_attention.h"
27+
#include "fastertransformer/cuda/decoding_kernel_check.h"
28+
29+
namespace fastertransformer
30+
{
31+
32+
template <typename T>
33+
void BeamSearch_OpenNMT(
34+
float *log_probs, float *cum_log_probs, bool *finished,
35+
T **key_cache, T **value_cache,
36+
int *parent_ids,
37+
int *sequence_length,
38+
int *word_ids,
39+
int *ids,
40+
int *output_ids,
41+
const int batch_size, const int beam_width,
42+
const int vocab_size, const int hidden_dim, const int step,
43+
const int cache_size, const int decoder_layers, cudaStream_t stream,
44+
const int end_id,
45+
int *finished_count)
46+
{
47+
#ifdef NDEBUG
48+
/* adding cum_log_probs to log_probs */
49+
broadcast_kernelLauncher(log_probs, cum_log_probs, batch_size, beam_width, vocab_size, stream);
50+
#else
51+
broadcast_kernelLauncher(log_probs, cum_log_probs, batch_size, beam_width, vocab_size, stream);
52+
cudaDeviceSynchronize();
53+
check_cuda_error(cudaGetLastError());
54+
55+
/*
56+
User can check the broadcast_kernel by broadcast_kernel_check.
57+
broadcast_kernel_check will compare the results of GPU and CPU.
58+
Note that broadcast_kernel_check contains broadcast_kernelLauncher and uses do not need to call it again.
59+
*/
60+
// broadcast_kernel_check(log_probs, cum_log_probs, batch_size, beam_width, vocab_size, stream);
61+
#endif
62+
63+
#ifdef NDEBUG
64+
/*Use two round kernels to pick the topK values for each batch */
65+
topK(log_probs, ids, batch_size, beam_width, vocab_size, stream);
66+
#else
67+
topK(log_probs, ids, batch_size, beam_width, vocab_size, stream);
68+
cudaDeviceSynchronize();
69+
check_cuda_error(cudaGetLastError());
70+
71+
/*
72+
User can check the topK by topK_check.
73+
topK_check will compare the results of GPU and CPU.
74+
Note that topK_check contains topK and uses do not need to call it again.
75+
*/
76+
// topK_kernel_check(log_probs, ids, batch_size, beam_width, vocab_size, stream);
77+
#endif
78+
79+
#ifdef NDEBUG
80+
update(log_probs, cum_log_probs, ids, finished,
81+
parent_ids, sequence_length, word_ids, output_ids,
82+
batch_size, beam_width, vocab_size, stream,
83+
end_id, finished_count);
84+
#else
85+
update(log_probs, cum_log_probs, ids, finished,
86+
parent_ids, sequence_length, word_ids, output_ids,
87+
batch_size, beam_width, vocab_size, stream,
88+
end_id, finished_count);
89+
cudaDeviceSynchronize();
90+
check_cuda_error(cudaGetLastError());
91+
92+
/*
93+
User can check the update by update_kernel_check.
94+
update_kernel_check will compare the results of GPU and CPU.
95+
Note that update_kernel_check contains update and uses do not need to call it again.
96+
*/
97+
// update_kernel_check(log_probs, cum_log_probs, ids, finished, parent_ids, sequence_length, word_ids, output_ids,
98+
// batch_size, beam_width, vocab_size, stream, end_id, finished_count);
99+
#endif
100+
101+
#ifdef NDEBUG
102+
update_KV_cache<T>(key_cache, value_cache, parent_ids, batch_size,
103+
beam_width, hidden_dim, step, cache_size,
104+
decoder_layers, stream);
105+
#else
106+
update_KV_cache<T>(key_cache, value_cache, parent_ids, batch_size,
107+
beam_width, hidden_dim, step, cache_size,
108+
decoder_layers, stream);
109+
cudaDeviceSynchronize();
110+
check_cuda_error(cudaGetLastError());
111+
112+
/*
113+
User can check the update_KV_cache by update_KV_cache_kernel_check.
114+
update_KV_cache_kernel_check will compare the results of GPU and CPU.
115+
Note that update_KV_cache_kernel_check contains update_KV_cache and uses do not need to call it again.
116+
*/
117+
// update_KV_cache_kernel_check(key_cache, value_cache, parent_ids, batch_size, beam_width, hidden_dim, step, cache_size, decoder_layers, stream);
118+
#endif
119+
}
120+
121+
} // namespace fastertransformer

0 commit comments

Comments
 (0)