diff --git a/onnxruntime/contrib_ops/cuda/bert/attention_impl.cu b/onnxruntime/contrib_ops/cuda/bert/attention_impl.cu index cfe61d9df4..8f7616c588 100644 --- a/onnxruntime/contrib_ops/cuda/bert/attention_impl.cu +++ b/onnxruntime/contrib_ops/cuda/bert/attention_impl.cu @@ -24,6 +24,7 @@ limitations under the License. #include #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(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; } diff --git a/onnxruntime/contrib_ops/cuda/bert/attention_impl.h b/onnxruntime/contrib_ops/cuda/bert/attention_impl.h index ea2ba8b22f..c5064c3744 100644 --- a/onnxruntime/contrib_ops/cuda/bert/attention_impl.h +++ b/onnxruntime/contrib_ops/cuda/bert/attention_impl.h @@ -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); diff --git a/onnxruntime/core/providers/cuda/cuda_common.cc b/onnxruntime/core/providers/cuda/cuda_common.cc new file mode 100644 index 0000000000..771c0338f1 --- /dev/null +++ b/onnxruntime/core/providers/cuda/cuda_common.cc @@ -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(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 diff --git a/onnxruntime/core/providers/cuda/cuda_common.h b/onnxruntime/core/providers/cuda/cuda_common.h index f3994bf43f..6a6eca8184 100644 --- a/onnxruntime/core/providers/cuda/cuda_common.h +++ b/onnxruntime/core/providers/cuda/cuda_common.h @@ -96,27 +96,84 @@ inline bool CalculateFdmStrides(gsl::span p, const std::vector= 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 diff --git a/onnxruntime/core/providers/cuda/shared_inc/fpgeneric.h b/onnxruntime/core/providers/cuda/shared_inc/fpgeneric.h index 6b46550fae..6752316ef5 100644 --- a/onnxruntime/core/providers/cuda/shared_inc/fpgeneric.h +++ b/onnxruntime/core/providers/cuda/shared_inc/fpgeneric.h @@ -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(alpha)); - float h_b = onnxruntime::math::halfToFloat(*reinterpret_cast(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(alpha)); + float h_b = onnxruntime::math::halfToFloat(*reinterpret_cast(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(alpha)).ToFloat(); float h_b = onnxruntime::BFloat16(*reinterpret_cast(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(alpha)); - float h_b = onnxruntime::math::halfToFloat(*reinterpret_cast(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(alpha)); + float h_b = onnxruntime::math::halfToFloat(*reinterpret_cast(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(alpha)).ToFloat(); float h_b = onnxruntime::BFloat16(*reinterpret_cast(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(alpha)); - float h_b = onnxruntime::math::halfToFloat(*reinterpret_cast(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(alpha)); + float h_b = onnxruntime::math::halfToFloat(*reinterpret_cast(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(alpha)).ToFloat(); float h_b = onnxruntime::BFloat16(*reinterpret_cast(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 - - - - diff --git a/onnxruntime/test/providers/cuda/gemm_options_test.cc b/onnxruntime/test/providers/cuda/gemm_options_test.cc new file mode 100644 index 0000000000..3dec06da25 --- /dev/null +++ b/onnxruntime/test/providers/cuda/gemm_options_test.cc @@ -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 diff --git a/tools/ci_build/amd_hipify.py b/tools/ci_build/amd_hipify.py index 4d3bea4c81..3818602c42 100644 --- a/tools/ci_build/amd_hipify.py +++ b/tools/ci_build/amd_hipify.py @@ -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',