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

Jump to bottom

Add Tensor Core support for convolutions #5

Merged
merged 2 commits into from Jul 15, 2020
Merged

Add Tensor Core support for convolutions #5

Changes from 1 commit
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
80f4bb7
Add Tensor Core support for convolutions
Jul 14, 2020
a98e443
Rename ETL_TENSOR_CORES define to ETL_EXPERIMENTAL_TENSOR_CORES
Jul 15, 2020
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Jump to
Jump to file
Failed to load files.
Diff view
Diff view
20 changes: 20 additions & 0 deletions include/etl/impl/cudnn/conv.hpp
View file
Open in desktop
Expand Up @@ -12,6 +12,8 @@

#pragma once

#define ETL_TENSOR_CORES
wichtounet marked this conversation as resolved.
Show resolved Hide resolved

#ifdef ETL_CUDNN_MODE

#include "etl/impl/cublas/cuda.hpp"
Expand Down Expand Up @@ -74,6 +76,9 @@ void conv2_valid_set(I&& input, K&& kernel, C&& conv, size_t s1, size_t s2, size
cudnnConvolutionDescriptor_t convolution;
cudnn_check(cudnnCreateConvolutionDescriptor(&convolution));
cudnn_check(cudnnSetConvolution2dDescriptor(convolution, p1, p2, s1, s2, 1, 1, mode, data_type));
#ifdef ETL_TENSOR_CORES
cudnn_check(cudnnSetConvolutionMathType(convolution, CUDNN_TENSOR_OP_MATH_ALLOW_CONVERSION));
#endif

// Find the algorithm to use
cudnnConvolutionFwdAlgo_t conv_algo;
Expand Down Expand Up @@ -185,6 +190,9 @@ void conv4_forward_set(I&& input, K&& kernel, C&& conv, size_t s1, size_t s2, si
cudnnConvolutionDescriptor_t convolution;
cudnn_check(cudnnCreateConvolutionDescriptor(&convolution));
cudnn_check(cudnnSetConvolution2dDescriptor(convolution, p1, p2, s1, s2, 1, 1, mode, data_type));
#ifdef ETL_TENSOR_CORES
cudnn_check(cudnnSetConvolutionMathType(convolution, CUDNN_TENSOR_OP_MATH_ALLOW_CONVERSION));
#endif

// Find the algorithm to use
cudnnConvolutionFwdAlgo_t conv_algo;
Expand Down Expand Up @@ -289,6 +297,9 @@ void conv4_backward_filter_set(I&& input, K&& kernel, C&& conv, size_t s1, size_
cudnnConvolutionDescriptor_t convolution;
cudnn_check(cudnnCreateConvolutionDescriptor(&convolution));
cudnn_check(cudnnSetConvolution2dDescriptor(convolution, p1, p2, s1, s2, 1, 1, mode, data_type));
#ifdef ETL_TENSOR_CORES
cudnn_check(cudnnSetConvolutionMathType(convolution, CUDNN_TENSOR_OP_MATH_ALLOW_CONVERSION));
#endif

// Find the algorithm to use
cudnnConvolutionBwdFilterAlgo_t conv_algo;
Expand Down Expand Up @@ -395,6 +406,9 @@ void conv2_full_set(I&& input, K&& kernel, C&& conv, cudnnConvolutionMode_t mode
cudnnConvolutionDescriptor_t convolution;
cudnn_check(cudnnCreateConvolutionDescriptor(&convolution));
cudnn_check(cudnnSetConvolution2dDescriptor(convolution, 0, 0, 1, 1, 1, 1, mode, data_type));
#ifdef ETL_TENSOR_CORES
cudnn_check(cudnnSetConvolutionMathType(convolution, CUDNN_TENSOR_OP_MATH_ALLOW_CONVERSION));
#endif

// Find the algorithm to use
cudnnConvolutionBwdDataAlgo_t conv_algo;
Expand Down Expand Up @@ -495,6 +509,9 @@ void conv2_valid_multi_set(I& input, K&& kernel, C&& conv, size_t s1, size_t s2,
cudnnConvolutionDescriptor_t convolution;
cudnn_check(cudnnCreateConvolutionDescriptor(&convolution));
cudnn_check(cudnnSetConvolution2dDescriptor(convolution, p1, p2, s1, s2, 1, 1, mode, data_type));
#ifdef ETL_TENSOR_CORES
cudnn_check(cudnnSetConvolutionMathType(convolution, CUDNN_TENSOR_OP_MATH_ALLOW_CONVERSION));
#endif

// Find the algorithm to use
cudnnConvolutionFwdAlgo_t conv_algo;
Expand Down Expand Up @@ -600,6 +617,9 @@ void conv4_backward_data_set(I&& input, K&& kernel, C&& conv, cudnnConvolutionMo
cudnnConvolutionDescriptor_t convolution;
cudnn_check(cudnnCreateConvolutionDescriptor(&convolution));
cudnn_check(cudnnSetConvolution2dDescriptor(convolution, p1, p2, s1, s2, 1, 1, mode, data_type));
#ifdef ETL_TENSOR_CORES
cudnn_check(cudnnSetConvolutionMathType(convolution, CUDNN_TENSOR_OP_MATH_ALLOW_CONVERSION));
#endif

// Find the algorithm to use
cudnnConvolutionBwdDataAlgo_t conv_algo;
Expand Down