Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Add device options to enable tensor core math mode. #3066

Merged
merged 7 commits into from
Mar 4, 2019
Merged

Add device options to enable tensor core math mode. #3066

Show file tree
Hide file tree
Changes from 3 commits
Commits
Show all changes
7 commits
Select commit Hold shift + click to select a range
7fc1ca5
Add device options to enable tensor core math mode.
luitjens Mar 1, 2019
7982a20
fixed typo in comment
luitjens Mar 2, 2019
3d111ec
Style fixes
luitjens Mar 2, 2019
e5469b7
Fixed one more style issue
luitjens Mar 2, 2019
919f32c
Remove trailing _ from struct
luitjens Mar 2, 2019
742ec88
Register device options in a few of the nnet3 binaries.
luitjens Mar 2, 2019
cd812a0
Register before read. Updated example.
luitjens Mar 2, 2019
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
10 changes: 10 additions & 0 deletions src/cudamatrix/cu-device.cc
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -110,6 +110,14 @@ void CuDevice::Initialize() {
// Initialize CUBLAS.
CUBLAS_SAFE_CALL(cublasCreate(&cublas_handle_));
CUBLAS_SAFE_CALL(cublasSetStream(cublas_handle_, cudaStreamPerThread));

if (device_options_.use_tensor_cores_) {
// Enable tensor cores in CUBLAS
// Note if the device does not support tensor cores this will fall back to normal math mode
CUBLAS_SAFE_CALL(cublasSetMathMode(cublas_handle_,
CUBLAS_TENSOR_OP_MATH));
}

// Initialize the cuSPARSE library
CUSPARSE_SAFE_CALL(cusparseCreate(&cusparse_handle_));
CUSPARSE_SAFE_CALL(cusparseSetStream(cusparse_handle_, cudaStreamPerThread));
Expand Down Expand Up @@ -525,6 +533,8 @@ CuDevice::~CuDevice() {
// Each thread has its own copy of the CuDevice object.
// Note: this was declared "static".
thread_local CuDevice CuDevice::this_thread_device_;

CuDevice::CuDeviceOptions CuDevice::device_options_;

// define and initialize the static members of the CuDevice object.
int32 CuDevice::device_id_ = -1;
Expand Down
22 changes: 22 additions & 0 deletions src/cudamatrix/cu-device.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -184,8 +184,30 @@ class CuDevice {
/// (i.e. from outside the class), call this only if Enabled() returns true.
bool IsComputeExclusive();

// Register command line options for CUDA device.
// This must be done before calling CuDevice::Initialize()
// Example:
// CuDevice::RegisterDeviceOptions(&po);
// CuDevice::Initialize();
static void RegisterDeviceOptions(OptionsItf *po) {
CuDevice::device_options_.Register(po);
}
~CuDevice();
private:

struct CuDeviceOptions {
bool use_tensor_cores_; // Enable tensor cores
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

should be use_tensor_cores, no trailing underscore (since it's a struct)... also two spaces before the //.

CuDeviceOptions () : use_tensor_cores_(false) {};
void Register(OptionsItf *po) {
po->Register("cuda-use-tensor-cores",&use_tensor_cores_,
"Enable FP16 tensor math. "
"This is higher performance but less accuracy. "
"This is only recommended for inference.");
}
};

static CuDeviceOptions device_options_;

// Default constructor used to initialize this_thread_device_
CuDevice();
CuDevice(CuDevice&); // Disallow.
Expand Down