TF32 (TensorFloat-32) = FP32 와 같은 지수 + 더 짧은 가수 : FP 32 와 동일한 다이내믹 레인지 정밀도는 FP 16 과 비슷
연산 : Tensor Core 에서 TF32 로 곱셈, 누산은 FP32 로 함
장점 : 기존 FP32 대비 GEMM 속도 크게 향상 (배치, 행렬이 클수록)
static inline void gemm_rm_strided_batched_tf32(
cublasHandle_t h,
bool transA, bool transB,
int M, int N, int K,
const float* A, int lda, long long strideA,
const float* B, int ldb, long long strideB,
float* C, int ldc, long long strideC,
int batch,
float alpha=1.f, float beta=0.f)
{
cublasOperation_t opA = transA ? CUBLAS_OP_T : CUBLAS_OP_N;
cublasOperation_t opB = transB ? CUBLAS_OP_T : CUBLAS_OP_N;
CUBLAS_CHECK(
cublasGemmStridedBatchedEx(
h,
/*opB,opA*/ opB, opA,
/*m,n,k*/ N, M, K,
&alpha,
/*B*/ B, CUDA_R_32F, ldb, strideB,
/*A*/ A, CUDA_R_32F, lda, strideA,
&beta,
/*C*/ C, CUDA_R_32F, ldc, strideC,
/*batch*/ batch,
/*computeType*/ CUBLAS_COMPUTE_32F_FAST_TF32,
/*algo*/ CUBLAS_GEMM_DEFAULT_TENSOR_OP
)
);
}
MATMUL, ADD 의 호출 함수 변경으로 수정,
추후 옵션별 호출 가능
'dev_AI_framework' 카테고리의 다른 글
| Regression 에서 오류 발생, XOR 은 잘 되는데 왜...? - 해결완료 (0) | 2025.08.19 |
|---|---|
| 후에 구현할 기능 고민해보자... (5) | 2025.08.15 |
| CUDA 커널 호출 구조 개선 - 한 번의 호출로 학습 완료! (2) | 2025.08.15 |
| 배치 루프 (여러 번 런치) vs 배치 결합 (한 번에 런치) (4) | 2025.08.15 |
| 연산 방식의 변경 - 임시 버퍼/전치 커널 , 커스텀 matmul backward 제거 (0) | 2025.08.15 |