bias — Row/Col-wise Bias Add Kernels

본 모듈은 per‑sample 행렬 표현(rows × cols)에 대해 행(열) 방향 브로드캐스트 bias add를 제공합니다. add_bias_rowwise.cuh에는 커널과 런처가 함께 정의되어 있습니다

 

1. 목적/책임

• Dense/Activation 직후 등에서 열 방향(row‑wise) bias를 더하는 연산.
• CNN 등에서 채널 단위로 행 방향(col‑wise) bias를 브로드캐스트하는 연산.
• 단순 합 연산이므로 in‑place 사용 가능(out==in).

 

2. Shape/브로드캐스트 규약

• Per‑Sample Shape 규칙과 정합:
  • 실행기에서 B = batch_size * rows_per_sample을 계산해 rows로 전달합니다.
• row‑wise: (B, C) + (1, C) → (B, C)
• col‑wise: (B, C) + (rows_per_sample, 1) → (B, C) (실제로는 길이 rows_per_sample 벡터를 사용)
• CNN에서 채널 편향을 col‑wise로 쓰는 경우, rows_per_sample = channels(또는 내부 표현상의 행 수).

 

3. 메모리/레이아웃/성능

• 포인터는 모두 float32, C‑contiguous 가정.
• x축=열(col)로 매핑하여 coalesced read/write.
• In‑place 지원: out을 in과 동일 포인터로 넘겨도 안전.
• 튜닝 아이디어:
  • float2/float4 벡터화(열 길이가 2/4 배수일 때)
  • 큰 cols에서는 blockDim.x=128/256로 조정 후 occupancy 측정
  • L2 hit 개선을 위해 bias를 __ldg 또는 shared에 캐시(현대 GPU에선 자동 캐시로 충분한 경우 多)