| rocm_available | 1 if a HIP GPU is usable (initializes the backend on first call). |
| rocm_device_name | Name of the selected GPU ("AMD Instinct MI100"), or "" if none. |
| rocm_write | |
| rocm_read | |
| rocm_sync | wait for all queued work |
| rocm_gemm | c[m,n] = op(a)[m,k] * op(b)[k,n] (rocBLAS sgemm) |
| rocm_add | out = a + b |
| rocm_sub | |
| rocm_mul | |
| rocm_div | |
| rocm_pow | |
| rocm_scale | out = a * k |
| rocm_relu | |
| rocm_exp | |
| rocm_log | |
| rocm_sqrt | |
| rocm_gelu | tanh approximation (gelu_pytorch_tanh) |
| rocm_rmsnorm | RMSNorm per row: y[r] = x[r] / sqrt(mean(x[r]^2) + eps) * w, over 'rows'
rows of 'n' values; w has n values, or dev == 0 for no weight. y may be x. |
| rocm_relu_grad | out = z > 0 ? g : 0 |
| rocm_sgd | w -= lr * g, in place |
| rocm_sum | out1[0] = sum(a) |
| rocm_mse_grad | Fused MSE gradient: dy = 2 (y - t) and loss1[0] = sum((y - t)^2), in one pass. |
| rocm_copy | dst[dstOff .. dstOff+n) = src[srcOff .. srcOff+n) (device-to-device copy) |
| rocm_concat | out[b] = (a[b], c[b]) for 'batch' rows of na and nc elements, one launch
(the channel concat of NCHW tensors: na = aChannels * height * width). |
| rocm_upsample2x_nearest | Nearest-neighbour 2x upsample of 'planes' row-major h x w planes into
planes x 2h x 2w: out[p][y][x] = in[p][y/2][x/2] (cnn.h's FeatureMap::upsample2x_nearest). |
| rocm_upsample2x_concat | out[b] = (x[b].upsample2x_nearest(), skip[b]) along channels, in one pass:
x is [batch, xChannels, h, w], skip [batch, skipChannels, 2h, 2w], out
[batch, xChannels + skipChannels, 2h, 2w] (a U-Net decoder step). |
| rocm_bias_channels | NCHW per-channel bias, in place: y[b][c][i] += bias[c] for every one of the
'plane' (= height*width) elements, then max(0, .) if relu != 0. The
epilogue gpu_miopen.h runs after a convolution. |
| rocm_half_write | |
| rocm_half_read | |
| rocm_to_half | Queued conversions over dst.len elements. |
| rocm_from_half | exact |
| rocm_half_gemm | c[m,n] = op(a)[m,k] * op(b)[k,n], 16-bit a and b of the same dtype, fp32
accumulation; c in the same 16-bit dtype, or fp32 with _f32. |
| rocm_half_gemm_f32 | |
| rocm_half_sgd | w = round(w - lr * g), computed in fp32, in place. |
| rocm_half_gemm_relu | c = relu(op(a) op(b)), all in a's dtype (the GEMM, then ReLU in place). |
| rocm_half_gemv | Matrix-vector product for decoding one token at a time: y = w x with w an
[n][k] fp16/bf16 matrix (row-major, e.g. a linear layer's weight) and x,
y fp32 (x is not rounded to 16 bits). Memory-bound: 0.8-1.07 TB/s of
weights on an MI100, 1.04-1.8x faster than a one-row rocm_half_gemm_f32
(most on short outputs, e.g. 1536 x 2048: 7.8 vs 13.7 us). Fastest when k
is a multiple of 8. |
| rocm_half_relu | Training ops, as their RocmBuf versions (loss stays fp32). |
| rocm_half_relu_grad | |
| rocm_half_mse_grad | |
| rocm_half_sgd2 | rocm_half_sgd on two weight/gradient pairs in one launch. |
| rocm_half_convert | The usual layout: fp32 master weights and optimizer state, a 16-bit copy
of the weights for the forward / backward passes, 16-bit or fp32
gradients with the loss scale in them. One step: |
| rocm_to_half_many | Every src[i] -> dst[i] (lengths dst[i].len / src[i].len as in
rocm_to_half / rocm_from_half), 32 buffers per launch: a model's
parameters in one or a few launches instead of one per tensor. |
| rocm_from_half_many | |
| rocm_half_unscale_many | Gradient unscaling: dst[i] = src[i] * scale in fp32 (16-bit or fp32
sources; the fp32 form may work in place, dst == src), and flag[0] = 1
when any result is inf or NaN. The flag is only ever set: zero it
(rocm_zero) before the step. |
| rocm_unscale_many | |
| rocm_zero | b = 0 (all of b.len) |
| rocm_mp_sgd | Optimizer steps on fp32 master weights w with gradient g, as nn.h's
Optimizer::step (OPT_SGD / OPT_MOMENTUM when vel is valid, OPT_ADAM /
OPT_ADAMW when decoupled; 'step' is the 1-based step count for the bias
correction). The updated weights are also written, rounded, to wHalf when
it is valid. Skipped entirely when skip is valid and skip[0] != 0. Pass an
invalid buffer (dev == 0) for any optional argument. |
| rocm_mp_adam | |
| rocm_half_add | 16-bit elementwise ops (fp32 arithmetic, one rounding per value; a, b and
out of one dtype and a.len elements). |
| rocm_half_sub | |
| rocm_half_mul | |
| rocm_half_scale | |
| rocm_half_axpy | a + k b |
| rocm_half_bias | y[r][c][p] += bias[c] (rows x channels x plane, row-major), then ReLU when
relu: a linear layer's bias (plane 1) or an NCHW convolution's. |
| rocm_half_maxpool2d | 16-bit max / avg pooling on NCHW (x: [batch, channels, height, width]),
kernel x kernel window, given stride, no padding: the same windows as
gpu_miopen.h's rocm_maxpool2d. fp32 accumulation, one rounding; max
pooling copies the largest value's bits. |
| rocm_half_avgpool2d | |
| rocm_capture_begin | Ops issued between begin and end are recorded instead of run. Don't call
rocm_read/rocm_write/rocm_sync while recording. |
| rocm_capture_end | exec == 0 on failure |
| rocm_graph_launch | queued like any other op |
| rocm_stream_handle | |
| rocm_flash_attention | fp32, batched over heads: q and out are [heads][seqLen][headDim], k and v
[heads][kvLen][headDim], all row-major; headDim <= 128. 1 on success, 0 on
a bad shape or failure. Forward only, queued on the stream like every op. |
| rocm_linear_attention | Linear attention with phi(x) = elu(x) + 1, as attention.h's
linear_attention_forward: out_i = phi(q_i) S / phi(q_i) z with S, z the
sums of phi(k_j)^T v_j and phi(k_j). Chunked (128 rows on matrix cores,
32 otherwise): per-chunk states, a prefix over chunks, then each chunk's
rows -- so causal (j <= i, which needs kvLen == seqLen) runs in parallel
too. Cost O(n d^2) instead of
O(n^2 d). The chunk states live in a scratch buffer this library keeps
and grows; during graph capture a call that would need to grow it fails. |
| rocm_attention_use_mfma | on = 0 makes both ops use the CUDA-core kernels even on matrix-core GPUs
(for comparison and testing); on = 1 restores the default. Returns 1 if the
matrix-core kernels are in use after the call. |
| rocm_add_f32 | |
| rocm_sub_f32 | |
| rocm_mul_f32 | |
| rocm_div_f32 | |
| rocm_pow_f32 | |
| rocm_relu_f32 | |
| rocm_log_f32 | |
| rocm_exp_f32 | |
| rocm_sqrt_f32 | |
| rocm_scale_f32 | |
| rocm_sum_f32 | |
| rocm_flash_attention_f32 | The attention ops above on host arrays (same layouts). |
| rocm_linear_attention_f32 | |
| rocm_matmul_f32 | out[M,N] = a[M,K] * b[K,N] with this library's own tiled HIP kernel
(64x64 tiles in LDS, 4x4 per thread). Kept as a dependency-free
reference; rocm_matmul_f32_blas is several times faster at every size. |
| rocm_matmul_f32_blas | Same product through rocBLAS. |