epilogue 동작과정 이해

지금 코드는 가능하면 cuBLASLt 에필로그(BIAS / RELU_BIAS)를 시도하고, 실패하면 GEMM만 수행 → 후처리 커널(bias/relu)을 실행하는 폴백으로 설계

 

1) 에필로그 동작이 “시도”되는 지점

src/my_kernels.cu 내 ge2_launch_gemm_bias_act_tc_f16()에서 첫 번째로 run_matmul(/*try_epilogue_bias=*/ p->has_bias, st)를 호출합니다. 여기서 try_epilogue_bias==true이면 에필로그를 설정

// try_epilogue_bias=true 이면 에필로그(BIAS 또는 RELU_BIAS) 설정
if (try_epilogue_bias && p->has_bias) {
  cublasLtEpilogue_t epi =
      (p->act == 1) ? CUBLASLT_EPILOGUE_RELU_BIAS : CUBLASLT_EPILOGUE_BIAS;
  cublasLtMatmulDescSetAttribute(opDesc, CUBLASLT_MATMUL_DESC_EPILOGUE,
                                 &epi, sizeof(epi));
  const void* biasDev = reinterpret_cast<const void*>(bias_f32);
  cublasLtMatmulDescSetAttribute(opDesc, CUBLASLT_MATMUL_DESC_BIAS_POINTER,
                                 &biasDev, sizeof(biasDev));
}

• CUBLASLT_MATMUL_DESC_EPILOGUE = BIAS 또는 RELU_BIAS
• CUBLASLT_MATMUL_DESC_BIAS_POINTER = 디바이스 메모리의 bias 벡터 포인터(우린 FP32)

그 뒤 cublasLtMatmul(...)을 수행하고 성공하면 에필로그까지 Fuse된 경로가 끝

cublasLtMatmulDescSetAttribute 함수

 

Matmul 연산 설명자 (cublasLtMatmulDesc_t) 의 속성 Attribute 을 설정하는 함수  

cublasStatus_t cublasLtMatmulDescSetAttribute(
    cublasLtMatmulDesc_t matmulDesc,
    cublasLtMatmulDescAttributes_t attr,
    const void* buf,
    size_t sizeInBytes
);

주요 파라미터

• matmulDesc
  • cublasLtMatmulDesc_t 타입객체
  • 행렬 곱 연산에 대한 설명자 ( Descriptor) - 어떤 연산을 할지에 대한 메타데이터 
• attr
  • 어떤 속성을 세팅할지를 나타냄, 열거형
  • (cublasLtMatmulDescAttributes_t) 값으로 지정
• buf
  • 속성에 넣을 값이 담긴 버퍼 포인터.
• sizeInBytes
  • buf 의 크기 

 

자주 쓰이는 Attribute 예시

• CUBLASLT_MATMUL_DESC_TRANSA
  A 행렬을 전치할지 여부 (cublasOperation_t: CUBLAS_OP_N, CUBLAS_OP_T 등).
• CUBLASLT_MATMUL_DESC_TRANSB
  B 행렬을 전치할지 여부.
• CUBLASLT_MATMUL_DESC_EPILOGUE
  Epilogue 연산 지정 (예: Bias 추가, ReLU 활성화, GELU 등).
• CUBLASLT_MATMUL_DESC_EPILOGUE_AUX_POINTER
  Bias 같은 추가 데이터 포인터 지정.
• CUBLASLT_MATMUL_DESC_SCALE_TYPE
  스케일링 팩터 타입 지정 (float, half, etc.).

 

간단 예시

cublasLtMatmulDesc_t operationDesc;
cublasLtMatmulDescCreate(&operationDesc, CUDA_R_32F);

// A 행렬 전치 여부 설정
cublasOperation_t transa = CUBLAS_OP_N;
cublasLtMatmulDescSetAttribute(
    operationDesc,
    CUBLASLT_MATMUL_DESC_TRANSA,
    &transa,
    sizeof(transa)
);

// Epilogue: Bias 추가 + ReLU
cublasLtEpilogue_t epi = CUBLASLT_EPILOGUE_BIAS_RELU;
cublasLtMatmulDescSetAttribute(
    operationDesc,
    CUBLASLT_MATMUL_DESC_EPILOGUE,
    &epi,
    sizeof(epi)
);

 

✅ 정리
cublasLtMatmulDescSetAttribute = 행렬 곱 연산자(MatmulDesc)에 “속성(Attribute)”를 세팅하는 함수
→ 어떤 연산을 할지(전치 여부, 스케일, epilogue 동작 등)를 세부적으로 제어 가능

 

2) 에필로그가 “실패”하면?

같은 함수에서 즉시 폴백

// 1) 에필로그 시도
bool ok = run_matmul(p->has_bias, st);

// 2) 실패 시 폴백: 에필로그 OFF + 후처리 커널
if (!ok) {
  ok = run_matmul(false, st);   // 에필로그 없이 다시 GEMM
  // 이후 add_bias_fp16 / add_bias_relu_fp16 / relu_only_fp16 실행
}

• run_matmul(false, st) 는 에필로그를 설정하지 않는다.
• 그 다음, 후처리 커널 중 하나를 실행해 Bias/Relu 를 적용

성공하면 에필로그 실행, 실패 시 후처리 커널로 적용

 

3) 에필로그의 실제 수학적 의미

현재 다음과 같이 설정

• 입력/출력 : FP16 (Row-major)
• Compute/Accum : FP32  (cublasLtMatmulDescCreate(... CUBLAS_COMPUTE_32F, CUDA_R_32F))
• alpha = 1.0f, beta = 0.0f 

그 결과, 에필로그로 계산되는 형태는

// 1) 에필로그 시도
bool ok = run_matmul(p->has_bias, st);

// 2) 실패 시 폴백: 에필로그 OFF + 후처리 커널
if (!ok) {
  ok = run_matmul(false, st);   // 에필로그 없이 다시 GEMM
  // 이후 add_bias_fp16 / add_bias_relu_fp16 / relu_only_fp16 실행
}