Gahow Wang
6035ffdc0b
- Batched GEMM via cublasGemmStridedBatchedEx - Causal mask CUDA kernel (F32 + BF16) - Element-wise scale CUDA kernel (F32 + BF16) - attention() composing: batched_matmul + scale + causal_mask + softmax - Fixed to_device/contiguous infinite recursion (GPU contiguous via CPU round-trip) - 5 attention tests passing (max_err < 3e-7 F32) - Total: 61 tests passing across all crates Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
91 lines
3.1 KiB
Plaintext
91 lines
3.1 KiB
Plaintext
#include <cuda_bf16.h>
|
|
#include <math.h>
|
|
|
|
// GELU (tanh approximation):
|
|
// gelu(x) = 0.5 * x * (1 + tanh(sqrt(2/pi) * (x + 0.044715 * x^3)))
|
|
__device__ __forceinline__ float gelu_f(float x) {
|
|
const float SQRT_2_OVER_PI = 0.7978845608f;
|
|
float cube = x * x * x;
|
|
float inner = SQRT_2_OVER_PI * (x + 0.044715f * cube);
|
|
return 0.5f * x * (1.0f + tanhf(inner));
|
|
}
|
|
|
|
// SiLU (Swish): silu(x) = x * sigmoid(x) = x / (1 + exp(-x))
|
|
__device__ __forceinline__ float silu_f(float x) {
|
|
return x / (1.0f + expf(-x));
|
|
}
|
|
|
|
__global__ void gelu_f32(const float* x, float* out, int n) {
|
|
int idx = blockIdx.x * blockDim.x + threadIdx.x;
|
|
if (idx < n) out[idx] = gelu_f(x[idx]);
|
|
}
|
|
|
|
__global__ void gelu_bf16(const __nv_bfloat16* x, __nv_bfloat16* out, int n) {
|
|
int idx = blockIdx.x * blockDim.x + threadIdx.x;
|
|
if (idx < n) out[idx] = __float2bfloat16(gelu_f(__bfloat162float(x[idx])));
|
|
}
|
|
|
|
__global__ void silu_f32(const float* x, float* out, int n) {
|
|
int idx = blockIdx.x * blockDim.x + threadIdx.x;
|
|
if (idx < n) out[idx] = silu_f(x[idx]);
|
|
}
|
|
|
|
__global__ void silu_bf16(const __nv_bfloat16* x, __nv_bfloat16* out, int n) {
|
|
int idx = blockIdx.x * blockDim.x + threadIdx.x;
|
|
if (idx < n) out[idx] = __float2bfloat16(silu_f(__bfloat162float(x[idx])));
|
|
}
|
|
|
|
__global__ void scale_f32_kernel(const float* x, float* out, float scale, int n) {
|
|
int idx = blockIdx.x * blockDim.x + threadIdx.x;
|
|
if (idx < n) out[idx] = x[idx] * scale;
|
|
}
|
|
|
|
__global__ void scale_bf16_kernel(const __nv_bfloat16* x, __nv_bfloat16* out, float scale, int n) {
|
|
int idx = blockIdx.x * blockDim.x + threadIdx.x;
|
|
if (idx < n) out[idx] = __float2bfloat16(__bfloat162float(x[idx]) * scale);
|
|
}
|
|
|
|
extern "C" {
|
|
|
|
void launch_gelu_f32(const void* x, void* out, int n, void* stream) {
|
|
int block = 256;
|
|
int grid = (n + block - 1) / block;
|
|
gelu_f32<<<grid, block, 0, (cudaStream_t)stream>>>((const float*)x, (float*)out, n);
|
|
}
|
|
|
|
void launch_gelu_bf16(const void* x, void* out, int n, void* stream) {
|
|
int block = 256;
|
|
int grid = (n + block - 1) / block;
|
|
gelu_bf16<<<grid, block, 0, (cudaStream_t)stream>>>(
|
|
(const __nv_bfloat16*)x, (__nv_bfloat16*)out, n);
|
|
}
|
|
|
|
void launch_silu_f32(const void* x, void* out, int n, void* stream) {
|
|
int block = 256;
|
|
int grid = (n + block - 1) / block;
|
|
silu_f32<<<grid, block, 0, (cudaStream_t)stream>>>((const float*)x, (float*)out, n);
|
|
}
|
|
|
|
void launch_silu_bf16(const void* x, void* out, int n, void* stream) {
|
|
int block = 256;
|
|
int grid = (n + block - 1) / block;
|
|
silu_bf16<<<grid, block, 0, (cudaStream_t)stream>>>(
|
|
(const __nv_bfloat16*)x, (__nv_bfloat16*)out, n);
|
|
}
|
|
|
|
void launch_scale_f32(const void* x, void* out, float scale, int n, void* stream) {
|
|
int block = 256;
|
|
int grid = (n + block - 1) / block;
|
|
scale_f32_kernel<<<grid, block, 0, (cudaStream_t)stream>>>(
|
|
(const float*)x, (float*)out, scale, n);
|
|
}
|
|
|
|
void launch_scale_bf16(const void* x, void* out, float scale, int n, void* stream) {
|
|
int block = 256;
|
|
int grid = (n + block - 1) / block;
|
|
scale_bf16_kernel<<<grid, block, 0, (cudaStream_t)stream>>>(
|
|
(const __nv_bfloat16*)x, (__nv_bfloat16*)out, scale, n);
|
|
}
|
|
|
|
}
|