saymrwulf/onnxruntime
1
0
Fork 0
mirror of https://github.com/saymrwulf/onnxruntime.git synced 2026-07-29 20:14:01 +00:00
Code Issues Projects Releases Packages Wiki Activity Actions

change half gemm to use compute_32f as default (#7253)

Browse source 
change half gemm to use compute_32f as default; add env variable for configuration
This commit is contained in:
Tianlei Wu 2021-04-08 20:54:37 -07:00 committed by GitHub
parent a4fdb4dbd9
commit 274e2fea0c
No known key found for this signature in database
GPG key ID: 4AEE18F83AFDEB23
7 changed files with 344 additions and 160 deletions
Show stats Download patch file Download diff file Expand all files Collapse all files

21
onnxruntime/contrib_ops/cuda/bert/attention_impl.cu
View file

@ -24,6 +24,7 @@ limitations under the License.
#include <cuda_fp16.h>
#include "core/providers/cuda/cu_inc/common.cuh"
#include "core/providers/cuda/cuda_common.h"
#include "core/providers/cuda/shared_inc/fpgeneric.h"
#include "attention_impl.h"
#include "attention_softmax.h"
@ -88,7 +89,6 @@ bool QkvToContext(
const T* v = k + total_size;
cublasSetStream(cublas, stream);
CublasMathModeSetter helper(prop, cublas, CUBLAS_TENSOR_OP_MATH);
// Concat past (2xBxNxS'xH) to present (2xBxNxS*xH):
// past_k (BxNxS'xH) + k (BxNxSxH) => present_k (BxNxS*xH)
@ -111,11 +111,16 @@ bool QkvToContext(
// Q: BxNxSxH, K (present_k): BxNxS*xH, Q*K': BxNxSxS*
const float rsqrt_head_size = 1.f / sqrt(static_cast<float>(head_size));
const int temp_matrix_size = sequence_length * all_sequence_length;
T one = (T)(1.0f);
T zero = (T)(0.f);
// For raw attention mask, the scalar if 1/sqrt(H) is moved to softmax computation.
T alpha = (T)(use_raw_attention_mask ? 1.0f : rsqrt_head_size);
if (!CUBLAS_CALL(CublasGemmStridedBatched(
cublas, CUBLAS_OP_T, CUBLAS_OP_N, all_sequence_length, sequence_length, head_size, alpha, k, head_size, present_size_per_batch,
q, head_size, size_per_batch, 0.f, scratch1, all_sequence_length, temp_matrix_size, batches))) {
// TODO: move scalar to softmax computation since converting 1/Sqrt(H) to half might have loss in precision.
T alpha = use_raw_attention_mask ? one : (T)(rsqrt_head_size);
if (!CUBLAS_CALL(cublasGemmStridedBatchedHelper(
cublas, CUBLAS_OP_T, CUBLAS_OP_N, all_sequence_length, sequence_length, head_size, &alpha, k, head_size, present_size_per_batch,
q, head_size, size_per_batch, &zero, scratch1, all_sequence_length, temp_matrix_size, batches, prop))) {
return false;
}
@ -138,9 +143,9 @@ bool QkvToContext(
}
// compute P*V (as V*P), and store in scratch3: BxNxSxH
if (!CUBLAS_CALL(CublasGemmStridedBatched(
cublas, CUBLAS_OP_N, CUBLAS_OP_N, head_size, sequence_length, all_sequence_length, 1.f, v, head_size, present_size_per_batch,
scratch2, all_sequence_length, temp_matrix_size, 0.f, scratch3, head_size, size_per_batch, batches))) {
if (!CUBLAS_CALL(cublasGemmStridedBatchedHelper(
cublas, CUBLAS_OP_N, CUBLAS_OP_N, head_size, sequence_length, all_sequence_length, &one, v, head_size, present_size_per_batch,
scratch2, all_sequence_length, temp_matrix_size, &zero, scratch3, head_size, size_per_batch, batches, prop))) {
return false;
}

18
onnxruntime/contrib_ops/cuda/bert/attention_impl.h
View file

@ -38,24 +38,6 @@ bool LaunchAttentionKernel(
void* present // Present state output
);
cublasStatus_t inline CublasGemmStridedBatched(
cublasHandle_t handle, cublasOperation_t transa, cublasOperation_t transb,
int m, int n, int k, const float alpha,
const float* A, int lda, long long int strideA, const float* B, int ldb, long long int strideB,
const float beta, float* C, int ldc, long long int strideC, int batchCount) {
return cublasSgemmStridedBatched(
handle, transa, transb, m, n, k, &alpha, A, lda, strideA, B, ldb, strideB, &beta, C, ldc, strideC, batchCount);
}
cublasStatus_t inline CublasGemmStridedBatched(
cublasHandle_t handle, cublasOperation_t transa, cublasOperation_t transb,
int m, int n, int k, const half alpha,
const half* A, int lda, long long int strideA, const half* B, int ldb, long long int strideB,
const half beta, half* C, int ldc, long long int strideC, int batchCount) {
return cublasHgemmStridedBatched(
handle, transa, transb, m, n, k, &alpha, A, lda, strideA, B, ldb, strideB, &beta, C, ldc, strideC, batchCount);
}
bool LaunchTransCtx(cudaStream_t stream,
const int sequence_length, const int batch_size, const int head_size, const int num_heads,
const int max_threads_per_block, const float* input, float* output);

48
onnxruntime/core/providers/cuda/cuda_common.cc Normal file
View file

@ -0,0 +1,48 @@
// Copyright (c) Microsoft Corporation. All rights reserved.
// Licensed under the MIT License.
#include "core/providers/cuda/cuda_common.h"
#include "core/common/logging/logging.h"
#include "core/common/logging/severity.h"
#include "core/platform/env_var_utils.h"
namespace onnxruntime {
namespace cuda {
// The environment variable is for testing purpose only, and it might be removed in the future.
// The value is an integer, and its bits have the following meaning:
// 0x01 - aggregate in fp16
// 0x02 - disallow reduced precision reduction. No effect when aggregate in fp16.
// 0x04 - pedantic
constexpr const char* kCudaGemmOptions = "ORT_CUDA_GEMM_OPTIONS";
// Initialize the singleton instance
HalfGemmOptions HalfGemmOptions::instance;
const HalfGemmOptions* HalfGemmOptions::GetInstance() {
if (!instance.initialized_) {
// We do not use critical section here since it is fine to initialize multiple times by different threads.
int value = ParseEnvironmentVariableWithDefault<int>(kCudaGemmOptions, 0);
instance.Initialize(value);
}
return &instance;
}
void HalfGemmOptions::Initialize(int value)
{
compute_16f_ = (value & 0x01) > 0;
#if defined(CUDA_VERSION) && CUDA_VERSION >= 11000
disallow_reduced_precision_reduction_ = (value & 0x02) > 0;
pedantic_ = (value & 0x04) > 0;
LOGS_DEFAULT(INFO) << "ORT_CUDA_GEMM_OPTIONS: compute_16f=" << instance.compute_16f_
<< " disallow_reduced_precision_reduction=" << instance.disallow_reduced_precision_reduction_
<< " pedantic=" << instance.pedantic_;
#else
LOGS_DEFAULT(INFO) << "ORT_CUDA_GEMM_OPTIONS: compute_16f=" << instance.compute_16f_;
#endif
initialized_ = true;
}
} // namespace cuda
} // namespace onnxruntime

77
onnxruntime/core/providers/cuda/cuda_common.h
View file

@ -96,27 +96,84 @@ inline bool CalculateFdmStrides(gsl::span<fast_divmod> p, const std::vector<int6
class CublasMathModeSetter {
public:
CublasMathModeSetter(const cudaDeviceProp& prop, cublasHandle_t handle, cublasMath_t mode) : prop_(prop), handle_(handle) {
cublasGetMathMode(handle, &mode_);
#if defined(CUDA_VERSION) && CUDA_VERSION < 11000
if (prop.major >= 7 && mode == CUBLAS_TENSOR_OP_MATH) {
cublasSetMathMode(handle, mode);
}
#if defined(CUDA_VERSION) && CUDA_VERSION < 11000
enable_ = (mode == CUBLAS_TENSOR_OP_MATH ? prop.major >= 7 : true );
#else
enable_ = (mode == CUBLAS_TF32_TENSOR_OP_MATH ? prop.major >= 8 : true);
#endif
#if defined(CUDA_VERSION) && CUDA_VERSION >= 11000
if (prop.major >= 8 && mode == CUBLAS_TF32_TENSOR_OP_MATH) {
cublasSetMathMode(handle, mode);
if (enable_) {
cublasGetMathMode(handle, &mode_);
enable_ = (mode_ != mode);
if (enable_) {
cublasSetMathMode(handle, mode);
}
}
#endif
}
~CublasMathModeSetter() {
cublasSetMathMode(handle_, mode_);
if (enable_) {
cublasSetMathMode(handle_, mode_);
}
}
private:
const cudaDeviceProp& prop_;
cublasHandle_t handle_;
cublasMath_t mode_;
bool enable_;
};
// Cublas Gemm options for half data type
class HalfGemmOptions {
public:
#if defined(CUDA_VERSION) && CUDA_VERSION >= 11000
cublasMath_t GetMathMode() const {
if (pedantic_) {
return CUBLAS_PEDANTIC_MATH;
}
return disallow_reduced_precision_reduction_ && !compute_16f_ ? CUBLAS_MATH_DISALLOW_REDUCED_PRECISION_REDUCTION : CUBLAS_DEFAULT_MATH;
}
cublasComputeType_t GetComputeType() const {
if (compute_16f_) {
return pedantic_ ? CUBLAS_COMPUTE_16F_PEDANTIC : CUBLAS_COMPUTE_16F;
} else {
return pedantic_ ? CUBLAS_COMPUTE_32F_PEDANTIC : CUBLAS_COMPUTE_32F;
}
}
#else
cublasMath_t GetMathMode() const {
// CublasMathModeSetter will check whether device has tensor cores later.
return CUBLAS_TENSOR_OP_MATH;
}
cudaDataType GetComputeType() const {
return compute_16f_ ? CUDA_R_16F : CUDA_R_32F;
}
#endif
static const HalfGemmOptions* GetInstance();
bool IsCompute16F() const { return compute_16f_; }
void Initialize(int value);
private:
// Default is FP32. Aggregate in FP16 might be faster but the cost is loss in precision.
bool compute_16f_{false};
#if defined(CUDA_VERSION) && CUDA_VERSION >= 11000
// Avoid intermediate overflows in accumulation. When compute type is FP32, it will not use FP16 in reduction.
bool disallow_reduced_precision_reduction_{false};
// For numerical robustness studies only. It is much slower.
bool pedantic_{false};
#endif
bool initialized_{false};
static HalfGemmOptions instance;
};
} // namespace cuda

245
onnxruntime/core/providers/cuda/shared_inc/fpgeneric.h
View file

@ -15,25 +15,24 @@
#include "core/providers/cuda/cuda_common.h"
// Generalize library calls to be use in template functions
using namespace onnxruntime::cuda;
// gemm
inline cublasStatus_t cublasGemmHelper(cublasHandle_t handle,
cublasOperation_t transa,
cublasOperation_t transb,
int m, int n, int k,
const float* alpha,
const float* A, int lda,
const float* B, int ldb,
const float* beta,
float* C, int ldc,
const cudaDeviceProp& prop) {
#ifdef ENABLE_TRAINING
// Generalize library calls to be use in template functions
inline cublasStatus_t
cublasGemmHelper(cublasHandle_t handle,
cublasOperation_t transa,
cublasOperation_t transb,
int m, int n, int k,
const float* alpha,
const float* A, int lda,
const float* B, int ldb,
const float* beta,
float* C, int ldc,
const cudaDeviceProp& prop) {
#if defined(CUDA_VERSION) && CUDA_VERSION >= 11000
// TF32 uses 10 bit mantissa which has sufficient margin of precision for most use cases. It gets 8x throughput than FP32 in A100.
// It can be overrided by setting environment variable NVIDIA_TF32_OVERRIDE = 0 to disable TF32
onnxruntime::cuda::CublasMathModeSetter math_mode_setter(prop, handle, CUBLAS_TF32_TENSOR_OP_MATH);
#else
ORT_UNUSED_PARAMETER(prop);
#endif
#else
ORT_UNUSED_PARAMETER(prop);
#endif
@ -48,6 +47,7 @@ inline cublasStatus_t cublasGemmHelper(cublasHandle_t handle,
beta,
C, ldc);
}
inline cublasStatus_t cublasGemmHelper(cublasHandle_t handle,
cublasOperation_t transa,
cublasOperation_t transb,
@ -68,6 +68,7 @@ inline cublasStatus_t cublasGemmHelper(cublasHandle_t handle,
beta,
C, ldc);
}
inline cublasStatus_t cublasGemmHelper(cublasHandle_t handle,
cublasOperation_t transa,
cublasOperation_t transb,
@ -78,36 +79,36 @@ inline cublasStatus_t cublasGemmHelper(cublasHandle_t handle,
const half* beta,
half* C, int ldc,
const cudaDeviceProp& prop) {
onnxruntime::cuda::CublasMathModeSetter math_mode_setter(prop, handle, CUBLAS_TENSOR_OP_MATH);
#ifdef ENABLE_TRAINING
float h_a = onnxruntime::math::halfToFloat(*reinterpret_cast<const uint16_t*>(alpha));
float h_b = onnxruntime::math::halfToFloat(*reinterpret_cast<const uint16_t*>(beta));
// accumulating in FP32
return cublasGemmEx(handle,
transa,
transb,
m, n, k,
&h_a,
A, CUDA_R_16F, lda,
B, CUDA_R_16F, ldb,
&h_b,
C, CUDA_R_16F, ldc,
CUDA_R_32F,
CUBLAS_GEMM_DEFAULT);
#else
// accumulating in FP16
return cublasHgemm(handle,
transa,
transb,
m, n, k,
alpha,
A, lda,
B, ldb,
beta,
C, ldc);
#endif
const HalfGemmOptions* half_options = HalfGemmOptions::GetInstance();
onnxruntime::cuda::CublasMathModeSetter math_mode_setter(prop, handle, half_options->GetMathMode());
if (half_options->IsCompute16F()) {
return cublasGemmEx(handle,
transa,
transb,
m, n, k,
alpha,
A, CUDA_R_16F, lda,
B, CUDA_R_16F, ldb,
beta,
C, CUDA_R_16F, ldc,
half_options->GetComputeType(),
CUBLAS_GEMM_DEFAULT);
} else {
// The alpha and beta shall have same precision as compute type.
float h_a = onnxruntime::math::halfToFloat(*reinterpret_cast<const uint16_t*>(alpha));
float h_b = onnxruntime::math::halfToFloat(*reinterpret_cast<const uint16_t*>(beta));
return cublasGemmEx(handle,
transa,
transb,
m, n, k,
&h_a,
A, CUDA_R_16F, lda,
B, CUDA_R_16F, ldb,
&h_b,
C, CUDA_R_16F, ldc,
half_options->GetComputeType(),
CUBLAS_GEMM_DEFAULT);
}
}
#if defined(CUDA_VERSION) && CUDA_VERSION >= 11000
@ -120,9 +121,7 @@ inline cublasStatus_t cublasGemmHelper(cublasHandle_t handle,
const nv_bfloat16* B, int ldb,
const nv_bfloat16* beta,
nv_bfloat16* C, int ldc,
const cudaDeviceProp& prop) {
onnxruntime::cuda::CublasMathModeSetter math_mode_setter(prop, handle, CUBLAS_DEFAULT_MATH);
const cudaDeviceProp& /*prop*/) {
float h_a = onnxruntime::BFloat16(*reinterpret_cast<const uint16_t*>(alpha)).ToFloat();
float h_b = onnxruntime::BFloat16(*reinterpret_cast<const uint16_t*>(beta)).ToFloat();
@ -153,15 +152,12 @@ inline cublasStatus_t cublasGemmBatchedHelper(cublasHandle_t handle,
float* Carray[], int ldc,
int batch_count,
const cudaDeviceProp& prop) {
#ifdef ENABLE_TRAINING
#if defined(CUDA_VERSION) && CUDA_VERSION >= 11000
onnxruntime::cuda::CublasMathModeSetter math_mode_setter(prop, handle, CUBLAS_TF32_TENSOR_OP_MATH);
#else
ORT_UNUSED_PARAMETER(prop);
#endif
#else
ORT_UNUSED_PARAMETER(prop);
#endif
return cublasSgemmBatched(handle,
transa,
transb,
@ -173,6 +169,7 @@ inline cublasStatus_t cublasGemmBatchedHelper(cublasHandle_t handle,
Carray, ldc,
batch_count);
}
inline cublasStatus_t cublasGemmBatchedHelper(cublasHandle_t handle,
cublasOperation_t transa,
cublasOperation_t transb,
@ -195,6 +192,7 @@ inline cublasStatus_t cublasGemmBatchedHelper(cublasHandle_t handle,
Carray, ldc,
batch_count);
}
inline cublasStatus_t cublasGemmBatchedHelper(cublasHandle_t handle,
cublasOperation_t transa,
cublasOperation_t transb,
@ -206,38 +204,37 @@ inline cublasStatus_t cublasGemmBatchedHelper(cublasHandle_t handle,
half* Carray[], int ldc,
int batch_count,
const cudaDeviceProp& prop) {
onnxruntime::cuda::CublasMathModeSetter math_mode_setter(prop, handle, CUBLAS_TENSOR_OP_MATH);
#ifdef ENABLE_TRAINING
float h_a = onnxruntime::math::halfToFloat(*reinterpret_cast<const uint16_t*>(alpha));
float h_b = onnxruntime::math::halfToFloat(*reinterpret_cast<const uint16_t*>(beta));
// accumulating in FP32
return cublasGemmBatchedEx(handle,
transa,
transb,
m, n, k,
&h_a,
(const void**)Aarray, CUDA_R_16F, lda,
(const void**)Barray, CUDA_R_16F, ldb,
&h_b,
(void**)Carray, CUDA_R_16F, ldc,
batch_count,
CUDA_R_32F,
CUBLAS_GEMM_DEFAULT);
#else
// accumulating in FP16
return cublasHgemmBatched(handle,
transa,
transb,
m, n, k,
alpha,
(const __half**)Aarray, lda,
(const __half**)Barray, ldb,
beta,
(__half**)Carray, ldc,
batch_count);
#endif
const HalfGemmOptions* half_options = HalfGemmOptions::GetInstance();
onnxruntime::cuda::CublasMathModeSetter math_mode_setter(prop, handle, half_options->GetMathMode());
if (half_options->IsCompute16F()) {
return cublasGemmBatchedEx(handle,
transa,
transb,
m, n, k,
alpha,
(const void**)Aarray, CUDA_R_16F, lda,
(const void**)Barray, CUDA_R_16F, ldb,
beta,
(void**)Carray, CUDA_R_16F, ldc,
batch_count,
half_options->GetComputeType(),
CUBLAS_GEMM_DEFAULT);
} else {
float h_a = onnxruntime::math::halfToFloat(*reinterpret_cast<const uint16_t*>(alpha));
float h_b = onnxruntime::math::halfToFloat(*reinterpret_cast<const uint16_t*>(beta));
return cublasGemmBatchedEx(handle,
transa,
transb,
m, n, k,
&h_a,
(const void**)Aarray, CUDA_R_16F, lda,
(const void**)Barray, CUDA_R_16F, ldb,
&h_b,
(void**)Carray, CUDA_R_16F, ldc,
batch_count,
half_options->GetComputeType(),
CUBLAS_GEMM_DEFAULT);
}
}
#if defined(CUDA_VERSION) && CUDA_VERSION >= 11000
@ -251,8 +248,7 @@ inline cublasStatus_t cublasGemmBatchedHelper(cublasHandle_t handle,
const nv_bfloat16* beta,
nv_bfloat16* Carray[], int ldc,
int batch_count,
const cudaDeviceProp& prop) {
onnxruntime::cuda::CublasMathModeSetter math_mode_setter(prop, handle, CUBLAS_TENSOR_OP_MATH);
const cudaDeviceProp& /*prop*/) {
float h_a = onnxruntime::BFloat16(*reinterpret_cast<const uint16_t*>(alpha)).ToFloat();
float h_b = onnxruntime::BFloat16(*reinterpret_cast<const uint16_t*>(beta)).ToFloat();
@ -349,37 +345,38 @@ inline cublasStatus_t cublasGemmStridedBatchedHelper(cublasHandle_t handle,
long long int strideC,
int batch_count,
const cudaDeviceProp& prop) {
onnxruntime::cuda::CublasMathModeSetter math_mode_setter(prop, handle, CUBLAS_TENSOR_OP_MATH);
#ifdef ENABLE_TRAINING
float h_a = onnxruntime::math::halfToFloat(*reinterpret_cast<const uint16_t*>(alpha));
float h_b = onnxruntime::math::halfToFloat(*reinterpret_cast<const uint16_t*>(beta));
// accumulating in FP32
return cublasGemmStridedBatchedEx(handle,
transa,
transb,
m, n, k,
&h_a,
A, CUDA_R_16F, lda, strideA,
B, CUDA_R_16F, ldb, strideB,
&h_b,
C, CUDA_R_16F, ldc, strideC,
batch_count,
CUDA_R_32F,
CUBLAS_GEMM_DEFAULT);
#else
// accumulating in FP16
return cublasHgemmStridedBatched(handle,
transa,
transb,
m, n, k,
alpha,
A, lda, strideA,
B, ldb, strideB,
beta,
C, ldc, strideC,
batch_count);
#endif
const HalfGemmOptions* half_options = HalfGemmOptions::GetInstance();
onnxruntime::cuda::CublasMathModeSetter math_mode_setter(prop, handle, half_options->GetMathMode());
if (half_options->IsCompute16F()) {
return cublasGemmStridedBatchedEx(handle,
transa,
transb,
m, n, k,
alpha,
A, CUDA_R_16F, lda, strideA,
B, CUDA_R_16F, ldb, strideB,
beta,
C, CUDA_R_16F, ldc, strideC,
batch_count,
half_options->GetComputeType(),
CUBLAS_GEMM_DEFAULT);
} else {
// The alpha and beta shall have same precision as compute type.
float h_a = onnxruntime::math::halfToFloat(*reinterpret_cast<const uint16_t*>(alpha));
float h_b = onnxruntime::math::halfToFloat(*reinterpret_cast<const uint16_t*>(beta));
return cublasGemmStridedBatchedEx(handle,
transa,
transb,
m, n, k,
&h_a,
A, CUDA_R_16F, lda, strideA,
B, CUDA_R_16F, ldb, strideB,
&h_b,
C, CUDA_R_16F, ldc, strideC,
batch_count,
half_options->GetComputeType(),
CUBLAS_GEMM_DEFAULT);
}
}
#if defined(CUDA_VERSION) && CUDA_VERSION >= 11000
@ -396,8 +393,7 @@ inline cublasStatus_t cublasGemmStridedBatchedHelper(cublasHandle_t handle,
nv_bfloat16* C, int ldc,
long long int strideC,
int batch_count,
const cudaDeviceProp& prop) {
onnxruntime::cuda::CublasMathModeSetter math_mode_setter(prop, handle, CUBLAS_TENSOR_OP_MATH);
const cudaDeviceProp& /*prop*/) {
float h_a = onnxruntime::BFloat16(*reinterpret_cast<const uint16_t*>(alpha)).ToFloat();
float h_b = onnxruntime::BFloat16(*reinterpret_cast<const uint16_t*>(beta)).ToFloat();
// accumulating in FP32
@ -420,23 +416,24 @@ inline cublasStatus_t cublasGemmStridedBatchedHelper(cublasHandle_t handle,
inline cublasStatus_t cublasTransposeHelper(cudaStream_t, cublasHandle_t handle, cublasOperation_t transa, cublasOperation_t transb, int m, int n, const float* alpha, const float* A, int lda, const float* beta, const float* B, int ldb, float* C, int ldc) {
return cublasSgeam(handle, transa, transb, m, n, alpha, A, lda, beta, B, ldb, C, ldc);
}
inline cublasStatus_t cublasTransposeHelper(cudaStream_t, cublasHandle_t handle, cublasOperation_t transa, cublasOperation_t transb, int m, int n, const double* alpha, const double* A, int lda, const double* beta, const double* B, int ldb, double* C, int ldc) {
return cublasDgeam(handle, transa, transb, m, n, alpha, A, lda, beta, B, ldb, C, ldc);
}
cublasStatus_t cublasTransposeHelper(cudaStream_t, cublasHandle_t, cublasOperation_t, cublasOperation_t, int m, int n, const half*, const half* A, int, const half*, const half*, int, half* C, int);
// copy
inline cublasStatus_t cublasCopyHelper(cudaStream_t, cublasHandle_t handle, int n, const float* x, int incx, float* y, int incy) {
return cublasScopy(handle, n, x, incx, y, incy);
}
inline cublasStatus_t cublasCopyHelper(cudaStream_t, cublasHandle_t handle, int n, const double* x, int incx, double* y, int incy) {
return cublasDcopy(handle, n, x, incx, y, incy);
}
cublasStatus_t cublasCopyHelper(cudaStream_t stream, cublasHandle_t handle, int n, const half* x, int incx, half* y, int incy);
#if defined(CUDA_VERSION) && CUDA_VERSION >= 11000
cublasStatus_t cublasCopyHelper(cudaStream_t stream, cublasHandle_t handle, int n, const nv_bfloat16* x, int incx, nv_bfloat16* y, int incy);
#endif

94
onnxruntime/test/providers/cuda/gemm_options_test.cc Normal file
View file

@ -0,0 +1,94 @@
// Copyright (c) Microsoft Corporation. All rights reserved.
// Licensed under the MIT License.
#ifdef USE_CUDA
#include "gtest/gtest.h"
#include "core/providers/cuda/cuda_common.h"
namespace onnxruntime {
namespace cuda {
namespace test {
TEST(CudaGemmOptionsTest, DefaultOptions) {
HalfGemmOptions gemm_options;
ASSERT_FALSE(gemm_options.IsCompute16F());
#if defined(CUDA_VERSION) && CUDA_VERSION >= 11000
EXPECT_EQ(gemm_options.GetMathMode(), CUBLAS_DEFAULT_MATH);
EXPECT_EQ(gemm_options.GetComputeType(), CUBLAS_COMPUTE_32F);
#else
EXPECT_EQ(gemm_options.GetMathMode(), CUBLAS_TENSOR_OP_MATH);
EXPECT_EQ(gemm_options.GetComputeType(), CUDA_R_32F);
#endif
}
TEST(CudaGemmOptionsTest, Compute16F) {
HalfGemmOptions gemm_options;
gemm_options.Initialize(1);
ASSERT_TRUE(gemm_options.IsCompute16F());
#if defined(CUDA_VERSION) && CUDA_VERSION >= 11000
EXPECT_EQ(gemm_options.GetMathMode(), CUBLAS_DEFAULT_MATH);
EXPECT_EQ(gemm_options.GetComputeType(), CUBLAS_COMPUTE_16F);
#else
EXPECT_EQ(gemm_options.GetMathMode(), CUBLAS_TENSOR_OP_MATH);
EXPECT_EQ(gemm_options.GetComputeType(), CUDA_R_16F);
#endif
}
TEST(CudaGemmOptionsTest, NoReducedPrecision) {
HalfGemmOptions gemm_options;
gemm_options.Initialize(2);
ASSERT_FALSE(gemm_options.IsCompute16F());
#if defined(CUDA_VERSION) && CUDA_VERSION >= 11000
EXPECT_EQ(gemm_options.GetMathMode(), CUBLAS_MATH_DISALLOW_REDUCED_PRECISION_REDUCTION);
EXPECT_EQ(gemm_options.GetComputeType(), CUBLAS_COMPUTE_32F);
#else
EXPECT_EQ(gemm_options.GetMathMode(), CUBLAS_TENSOR_OP_MATH);
EXPECT_EQ(gemm_options.GetComputeType(), CUDA_R_32F);
#endif
}
TEST(CudaGemmOptionsTest, Pedantic) {
HalfGemmOptions gemm_options;
gemm_options.Initialize(4);
ASSERT_FALSE(gemm_options.IsCompute16F());
#if defined(CUDA_VERSION) && CUDA_VERSION >= 11000
EXPECT_EQ(gemm_options.GetMathMode(), CUBLAS_PEDANTIC_MATH);
EXPECT_EQ(gemm_options.GetComputeType(), CUBLAS_COMPUTE_32F_PEDANTIC);
#else
EXPECT_EQ(gemm_options.GetMathMode(), CUBLAS_TENSOR_OP_MATH);
EXPECT_EQ(gemm_options.GetComputeType(), CUDA_R_32F);
#endif
}
TEST(CudaGemmOptionsTest, Compute16F_Pedantic) {
HalfGemmOptions gemm_options;
gemm_options.Initialize(5);
ASSERT_TRUE(gemm_options.IsCompute16F());
#if defined(CUDA_VERSION) && CUDA_VERSION >= 11000
EXPECT_EQ(gemm_options.GetMathMode(), CUBLAS_PEDANTIC_MATH);
EXPECT_EQ(gemm_options.GetComputeType(), CUBLAS_COMPUTE_16F_PEDANTIC);
#else
EXPECT_EQ(gemm_options.GetMathMode(), CUBLAS_TENSOR_OP_MATH);
EXPECT_EQ(gemm_options.GetComputeType(), CUDA_R_16F);
#endif
}
TEST(CudaGemmOptionsTest, Compute16F_NoReducedPrecision) {
HalfGemmOptions gemm_options;
gemm_options.Initialize(3);
ASSERT_TRUE(gemm_options.IsCompute16F());
#if defined(CUDA_VERSION) && CUDA_VERSION >= 11000
EXPECT_EQ(gemm_options.GetMathMode(), CUBLAS_DEFAULT_MATH);
EXPECT_EQ(gemm_options.GetComputeType(), CUBLAS_COMPUTE_16F);
#else
EXPECT_EQ(gemm_options.GetMathMode(), CUBLAS_TENSOR_OP_MATH);
EXPECT_EQ(gemm_options.GetComputeType(), CUDA_R_16F);
#endif
}
} // namespace test
} // namespace cuda
} // namespace onnxruntime
#endif

1
tools/ci_build/amd_hipify.py
View file

@ -151,6 +151,7 @@ provider_excluded_files = [
'cuda_allocator.cc',
'cuda_allocator.h',
'cuda_call.cc',
'cuda_common.cc',
'cuda_common.h',
'cuda_execution_provider_info.cc',
'cuda_execution_provider_info.h',