pulsatrix::HIPBackend Class Reference

HIP-resident DeviceBackend implementation, targeting AMD GPUs via ROCm. More...

#include <hip_backend.hpp>

Inheritance diagram for pulsatrix::HIPBackend:

Collaboration diagram for pulsatrix::HIPBackend:

Public Member Functions
	HIPBackend ()
	~HIPBackend () override
	HIPBackend (const HIPBackend &)=delete
HIPBackend &	operator= (const HIPBackend &)=delete
DeviceType	device () const noexcept override
	Which device this backend's buffers reside on.
void *	allocate (size_t bytes) override
	Allocates a buffer of the given size.
void	free (void *ptr) noexcept override
	Frees a buffer previously returned by allocate(). Safe to call with nullptr.
void	copy (void *dst, const void *src, size_t bytes, CopyDirection dir) override
	Copies bytes between buffers.
void	fill (void *ptr, float value, size_t n) override
	Fills every element of a float buffer with a constant value.
void	gemm (const float *a, const float *b, float *out, size_t m, size_t k, size_t n) override
	Row-major matrix multiply: out = a * b.
void	elementwise (ElementwiseOp op, const float *in, float *out, size_t n) override
	Applies a unary elementwise operation to every element of a buffer.
void	add (const float *a, const float *b, float *out, size_t n) override
	Elementwise binary addition: out[i] = a[i] + b[i] for i in [0, n).
void	mul (const float *a, const float *b, float *out, size_t n) override
	Elementwise binary (Hadamard) multiplication: out[i] = a[i] * b[i] for i in [0, n).
void	gemm_ex (const float *a, bool transpose_a, const float *b, bool transpose_b, float *out, size_t m, size_t k, size_t n, float beta) override
	General row-major matrix multiply: out = op(A) * op(B) + beta * out.
void	column_sums (const float *in, float *out, size_t rows, size_t cols, float beta) override
	Per-column sum of a (rows x cols) row-major matrix: out[j] = beta*out[j] + sum_i in[i][j].
void	add_row_vector (const float *in, const float *row, float *out, size_t rows, size_t cols) override
	Broadcast row add: out[i][j] = in[i][j] + row[j] for a (rows x cols) matrix.
void	elementwise_backward (ElementwiseOp op, const float *x, const float *grad_out, float *grad_in, size_t n) override
	Activation backward: grad_in[i] = grad_out[i] * f'(x[i]), f = op, x = the forward input.
void	axpby (float alpha, const float *x, float beta, const float *y, float *out, size_t n) override
	out[i] = alpha * x[i] + beta * y[i]. out may alias x or y.
float	dot (const float *a, const float *b, size_t n) override
	Dot product sum_i a[i]*b[i], returned to the host.
void	softmax_rows (const float *in, float *out, size_t rows, size_t cols) override
	Row-wise softmax of a (rows x cols) matrix, max-subtracted. out may alias in.
void	softmax_rows_backward (const float *y, const float *dy, float *dx, size_t rows, size_t cols) override
	Softmax backward from its output y: dx[i][j] = y[i][j] * (dy[i][j] - sum_k y[i][k]*dy[i][k]).
void	logsumexp_rows (const float *in, float *out, size_t rows, size_t cols) override
	Per-row log-sum-exp, max-subtracted: out[i] = log sum_j exp(in[i][j]).
void	adam_step (float *param, const float *grad, float *m, float *v, size_t n, float lr, float beta1, float beta2, float eps, float bias_correction1, float bias_correction2) override
	One fused Adam update over n parameters.
float	sum (const float *in, size_t n) override
	sum_i in[i], returned to the host. Same reduction order as dot(). Synchronizes.
void	dropout_forward (const float *in, float *out, float *mask, size_t n, float p, float scale, uint64_t seed, uint64_t offset) override
	Inverted dropout with a counter-based RNG: element i is dropped iff uniform(seed, offset + i) < p; kept elements are scaled by scale.
void	bce_with_logits (const float *logits, const float *target, float *out, size_t n) override
	Per-element binary cross-entropy with logits: out[i] = max(x, 0) - x*y + log1p(exp(-|x|)), x = logits[i], y = target[i].
void	bce_with_logits_grad (const float *logits, const float *target, float *grad, size_t n, float scale) override
	BCE-with-logits gradient: grad[i] = (sigmoid(x) - y) * scale, using the overflow-free sigmoid (exp(x) / (1 + exp(x)) for x < 0).
void	layer_norm_forward (const float *in, const float *gamma, const float *beta, float *xhat, float *out, float *row_std, size_t rows, size_t cols, float eps) override
	LayerNorm forward per row: xhat = (x - mean)/sqrt(var + eps), out = gamma*xhat + beta.
void	layer_norm_backward (const float *grad_out, const float *gamma, const float *xhat, const float *row_std, float *grad_in, size_t rows, size_t cols) override
	LayerNorm input gradient per row, from the cached xhat and per-row std.
void	rms_norm_forward (const float *in, const float *gamma, float *out, float *row_rms, size_t rows, size_t cols, float eps) override
	RMSNorm forward per row: out = gamma * x / sqrt(mean(x^2) + eps).
void	rms_norm_backward (const float *grad_out, const float *gamma, const float *in, const float *row_rms, float *grad_in, float *gamma_terms, size_t rows, size_t cols) override
	RMSNorm input gradient per row, plus gamma_terms[r][i] = grad_out * x / rms – the per-row contributions column_sums then reduces into gamma's gradient.
void	rope_rotate (const float *in, const float *cos_table, const float *sin_table, float *out, size_t num_slices, size_t seq_len, size_t head_dim, bool inverse) override
	Rotary position embedding over (num_slices, seq_len, head_dim) data.
void	permute_0213 (const float *in, float *out, size_t d0, size_t d1, size_t d2, size_t d3) override
	Swaps the middle two axes: in (d0, d1, d2, d3) -> out (d0, d2, d1, d3).
void	gather_rows (const float *table, const float *indices, float *out, size_t count, size_t dim) override
	out[i][:] = table[indices[i]][:] for count rows of width dim.
void	scatter_add_rows (const float *src, const float *indices, float *table, size_t count, size_t dim) override
	table[indices[i]][:] += src[i][:] for i = 0..count-1, in increasing i.
void	tanh_gaussian_forward (const float *mean, const float *log_std, const float *eps, float *action, float *std_cache, float *log_prob, size_t rows, size_t cols, float stabilizer, double half_log_two_pi) override
	TanhGaussianPolicy sampling per (rows, cols) row: action = tanh(mean + exp(log_std)*eps), std_cache = exp(log_std), log_prob[r] accumulated in double.
void	tanh_gaussian_backward (const float *action, const float *std_cache, const float *eps, const float *grad_action, const float *grad_log_prob, float *grad_mean, float *grad_log_std, size_t n, float stabilizer) override
	TanhGaussianPolicy gradients w.r.t. mean and log_std, per element.
void	lrp_linear (const float *x, const float *w, const float *z, const float *r, float *r_in, size_t rows, size_t in_features, size_t out_features, float eps) override
	LinearModule epsilon rule: r_in[n][i] = sum_j (x[n][i] w[i][j] / stab(z[n][j])) r[n][j].
void	lrp_residual_split (const float *a, const float *b, const float *r, float *r_a, float *r_b, size_t n, float eps) override
	Epsilon split of a residual sum y = a + b: r_a = (a / stab(y)) r, r_b = (b / stab(y)) r.
void	lrp_bilinear_elementwise (const float *a, const float *b, const float *r, float *r_out, size_t n, float eps) override
	Eq. 15 for an elementwise product c = a*b: r_out = (a b / (2c + eps sign c)) r (same for both).
void	lrp_bilinear_matmul (const float *a, const float *b, const float *o, const float *r_o, float *r_a, float *r_b, size_t slices, size_t m, size_t p, size_t q, float eps, bool b_transposed) override
	Eq. 15 for slices independent matmuls O = A @ B (A (M x P), B (P x Q), O and r_o (M x Q)).
void	lrp_softmax_rows (const float *x, const float *y, const float *r, float *r_in, size_t rows, size_t cols) override
	SoftmaxModule rule per row: r_in = x * (r - y * sum(r)).
void	lrp_rope (const float *x, const float *y, const float *r, const float *cos_table, const float *sin_table, float *r_in, size_t slices, size_t seq_len, size_t head_dim, float eps) override
	RoPEModule epsilon rule over (slices, seq_len, head_dim), tables as for rope_rotate.
void	logic_pointwise (LogicOp op, int norm, const float *a, const float *b, const float *g_or_r, const float *y, float *out_a, float *out_b, size_t n, float eps) override
	One elementwise pass of Conjunction/Disjunction for operands a, b.
void	aggregator_forward (const float *x, float *mean_pow, float *out, size_t n, size_t cols, float p) override
	AggregatorModule power mean over the leading axis of an (n, cols) input, per column.
void	aggregator_backward (const float *x, const float *mean_pow, const float *grad_out, float *grad_in, size_t n, size_t cols, float p) override
	AggregatorModule input gradient, per column.
void	aggregator_lrp (const float *x, const float *mean_pow, const float *r_out, float *r_in, size_t n, size_t cols, float p, float eps) override
	AggregatorModule epsilon rule, per column.
void	im2col (const float *in, float *col, size_t n, size_t c, size_t h, size_t w, const ConvGeometry &geometry) override
	Unfolds (n, c, h, w) into (n, c*kh*kw, out_h*out_w) patches, with out_h = (h + 2*pad_h - kh) / stride_h + 1 (likewise out_w). Taps that fall in the zero padding read 0.
void	col2im_add (const float *col, float *out, size_t n, size_t c, size_t h, size_t w, const ConvGeometry &geometry) override
	Folds (n, c*kh*kw, out_h*out_w) patches back, adding into out (n, c, h, w).
void	add_channel_vector (const float *in, const float *vec, float *out, size_t n, size_t c, size_t inner) override
	out[i][ch][k] = in[i][ch][k] + vec[ch] over (n, c, inner). out may alias in.
void	lrp_conv (const float *col, const float *kernel, const float *pre_bias, const float *r, float *r_col, size_t n, size_t out_channels, size_t p, size_t q, float eps) override
	Conv2D epsilon rule in patch space: r_col (n, p, q) from the cached patches, the kernel (out_channels, p) and the pre-bias outputs (n, out_channels, q).
void	lrp_stabilized_divide (const float *r, const float *denom, const float *gate, float *out, size_t n, float eps, LrpGate gate_mode) override
	Gated stabilized division, the one non-gemm step of the affine LRP rules: out[i] = passes(gate[i]) ? r[i] / (denom[i] + eps sign(denom[i])) : 0, sign(0) = +1.
void	max_pool_forward (const float *in, float *out, float *argmax, size_t planes, size_t h, size_t w, size_t kh, size_t kw) override
	Non-overlapping max pool over planes of (h, w); argmax = flat in-plane index (first max wins).
void	max_unpool (const float *src, const float *argmax, float *dst, size_t planes, size_t h, size_t w, size_t kh, size_t kw) override
	dst[plane][argmax] = src for every pooled element; dst must be zeroed by the caller.
void	avg_pool_forward (const float *in, float *out, size_t planes, size_t h, size_t w, size_t kh, size_t kw) override
	Non-overlapping average pool over planes of (h, w).
void	avg_pool_backward (const float *grad_out, float *grad_in, size_t planes, size_t h, size_t w, size_t kh, size_t kw) override
	Spreads grad_out / (kh*kw) over each window; grad_in must be zeroed by the caller.
void	lrp_avg_pool (const float *x, const float *r, float *r_in, size_t planes, size_t h, size_t w, size_t kh, size_t kw, float eps) override
	AvgPool2D epsilon rule; r_in must be zeroed by the caller.
void	batch_norm_forward (const float *in, const float *gamma, const float *beta, float *xhat, float *out, float *channel_std, size_t n, size_t c, size_t spatial, float eps) override
	BatchNorm over (n, spatial) per channel of (n, c, spatial) data.
void	batch_norm_backward (const float *grad_out, const float *gamma, const float *xhat, const float *channel_std, float *grad_in, float *gamma_grad, float *beta_grad, size_t n, size_t c, size_t spatial) override
	BatchNorm input gradient plus this call's gamma/beta gradients (overwritten, per channel).
void	batch_norm_update_running (const float *in, float *running_mean, float *running_var, size_t n, size_t c, size_t spatial, float momentum) override
	Folds this batch's per-channel mean and unbiased variance into the running ones (PyTorch's momentum rule; the variance is kept when a channel has one value). FND-5.
void	batch_norm_eval_forward (const float *in, const float *gamma, const float *beta, const float *running_mean, const float *running_var, float *xhat, float *out, float *channel_std, size_t n, size_t c, size_t spatial, float eps) override
	Eval-mode BatchNorm from the running statistics: a per-channel affine map. FND-5.
void	batch_norm_eval_backward (const float *grad_out, const float *gamma, const float *xhat, const float *channel_std, float *grad_in, float *gamma_grad, float *beta_grad, size_t n, size_t c, size_t spatial) override
	Eval-mode BatchNorm gradient: grad_out * gamma / std, plus gamma/beta gradients (overwritten, per channel). FND-5.
void	group_norm_forward (const float *in, const float *gamma, const float *beta, float *xhat, float *out, float *group_std, size_t n, size_t c, size_t spatial, size_t num_groups, float eps) override
	GroupNorm per (example, group) of (n, c, spatial) data; group_std is (n, num_groups).
void	group_norm_backward (const float *grad_out, const float *gamma, const float *xhat, const float *group_std, float *grad_in, float *gamma_grad, float *beta_grad, size_t n, size_t c, size_t spatial, size_t num_groups) override
	GroupNorm input gradient plus this call's gamma/beta gradients (overwritten, per channel).
void	copy_2d (float *dst, size_t dst_stride, const float *src, size_t src_stride, size_t rows, size_t cols) override
	Strided 2-D copy: rows of cols floats from src (row stride src_stride) to dst (row stride dst_stride) – e.g. one timestep of an (N, L, D) sequence.
void	accumulate_rows (const float *in, float *out, size_t rows, size_t cols) override
	out[j] += in[i][j] for i = 0..rows-1 in order, accumulating straight into out.
void	recurrent_cell (RecurrentCellOp op, const RecurrentCellArgs &args, size_t n) override
	One fused recurrent-cell pass over n elements (see RecurrentCellOp for slots).
void	gru_lrp_hprev (const float *h_prev, const float *w_hn, const float *hn, const float *r_term_b, const float *direct, float *r_hprev, size_t rows, size_t hidden, float eps) override
	GRU's R_hprev for (rows, hidden): the recurrent epsilon-rule sum through W_hn, with each element's direct term inserted where the original loop added it.
void	ssm_pass (SsmPassOp op, const SsmPassArgs &args) override
	One fused Mamba / RWKV / RetNet pass (see SsmPassOp for lanes and slots).
void	rl_rows (RlRowOp op, const RlRowArgs &args) override
	One fused RL loss / target / Polyak pass (see RlRowOp for lanes and slots).
void	top_k_rows (const float *in, float *values, float *indices, size_t rows, size_t cols, size_t k, bool largest) override
	Per row of a row-major (rows, cols) matrix: the k largest (or smallest) values in rank order into values (rows, k), and their column indices into indices (rows, k) as whole-number floats.
Public Member Functions inherited from pulsatrix::DeviceBackend
virtual	~DeviceBackend ()=default

Detailed Description

HIP-resident DeviceBackend implementation, targeting AMD GPUs via ROCm.

Note

gemm uses hipBLAS (hipblasSgemm) with the column-major swap-and-transpose trick, since hipBLAS inherits cuBLAS's column-major assumption and this project's Tensor is row-major – see gpu_backend_programming/context_gpu_cublas_cudnn_integration.md. elementwise/add are hand-written kernels, one thread per element.

Structurally a mirror of CUDABackend, deliberately. Phase 1.6's HIPIFY triage found 39 of 40 references auto-convertible with no warnings in the kernel bodies at all (<<<>>>, __global__, threadIdx/blockIdx are identical in HIP), so a divergent design would add risk without adding value and would make Mission 2's three-way equivalence suite harder to read.

No warpSize handling is needed anywhere in this class: every kernel here is one-thread-per-element with no cross-lane operation, so the CDNA-wavefront-64 hazard that context_accel_rocm_hip.md flags for CUDA ports does not arise. (This project's dev device, gfx1151, reports warpSize 32 in any case.)

The same Phase 1.5 scope limit applies here: only Tensor operations that route entirely through DeviceBackend's own primitives are safe against a HIP-backed Tensor. Phase 1's Module backward/LRP/optimizer code is host-loop-only and is guarded by PULSATRIX_REQUIRE_HOST, which covers DeviceType::Hip identically to DeviceType::Cuda – those guards need no change for this backend.

Constructor & Destructor Documentation

HIPBackend() [1/2]

pulsatrix::HIPBackend::HIPBackend

(

)

~HIPBackend()

pulsatrix::HIPBackend::~HIPBackend ( )

override

HIPBackend() [2/2]

pulsatrix::HIPBackend::HIPBackend ( const HIPBackend &  )

delete

Member Function Documentation

accumulate_rows()

void pulsatrix::HIPBackend::accumulate_rows ( const float *  in, float *  out, size_t  rows, size_t  cols  )

overridevirtual

out[j] += in[i][j] for i = 0..rows-1 in order, accumulating straight into out.

Note

Unlike column_sums(beta = 1), which adds a finished column sum to out, this adds row by row into the running value – the association the recurrent modules' bias gradients always used.

Implements pulsatrix::DeviceBackend.

adam_step()

void pulsatrix::HIPBackend::adam_step ( float *  param, const float *  grad, float *  m, float *  v, size_t  n, float  lr, float  beta1, float  beta2, float  eps, float  bias_correction1, float  bias_correction2  )

overridevirtual

One fused Adam update over n parameters.

Parameters

bias_correction1	1 - beta1^t, computed once on the host per step.
bias_correction2	1 - beta2^t, likewise.

Note

Same per-element expression order as the pre-campaign AdamOptimizer host loop: m = b1*m + (1-b1)*g; v = b2*v + (1-b2)*g*g; p -= lr * (m/bc1) / (sqrt(v/bc2) + eps).

Implements pulsatrix::DeviceBackend.

add()

void pulsatrix::HIPBackend::add ( const float *  a, const float *  b, float *  out, size_t  n  )

overridevirtual

Elementwise binary addition: out[i] = a[i] + b[i] for i in [0, n).

Parameters

a	First operand, must hold at least n floats.
b	Second operand, must hold at least n floats.
out	Output buffer, must hold at least n floats. May alias a or b for an in-place accumulation.
n	Number of elements.

Note

A separate method from elementwise() rather than extending ElementwiseOp with a binary case – see ElementwiseOp's own

. Added in Mission 3 (Phase 0 autograd) specifically because gradient accumulation (multiple children contributing to a shared parent's gradient) needs it; this is the resolution of the binary-op design decision deferred since Mission 0.

Implements pulsatrix::DeviceBackend.

add_channel_vector()

void pulsatrix::HIPBackend::add_channel_vector ( const float *  in, const float *  vec, float *  out, size_t  n, size_t  c, size_t  inner  )

overridevirtual

out[i][ch][k] = in[i][ch][k] + vec[ch] over (n, c, inner). out may alias in.

Implements pulsatrix::DeviceBackend.

add_row_vector()

void pulsatrix::HIPBackend::add_row_vector ( const float *  in, const float *  row, float *  out, size_t  rows, size_t  cols  )

overridevirtual

Broadcast row add: out[i][j] = in[i][j] + row[j] for a (rows x cols) matrix.

Note

out may alias in.

Implements pulsatrix::DeviceBackend.

aggregator_backward()

void pulsatrix::HIPBackend::aggregator_backward ( const float *  x, const float *  mean_pow, const float *  grad_out, float *  grad_in, size_t  n, size_t  cols, float  p  )

overridevirtual

AggregatorModule input gradient, per column.

Implements pulsatrix::DeviceBackend.

aggregator_forward()

void pulsatrix::HIPBackend::aggregator_forward ( const float *  x, float *  mean_pow, float *  out, size_t  n, size_t  cols, float  p  )

overridevirtual

AggregatorModule power mean over the leading axis of an (n, cols) input, per column.

Implements pulsatrix::DeviceBackend.

aggregator_lrp()

void pulsatrix::HIPBackend::aggregator_lrp ( const float *  x, const float *  mean_pow, const float *  r_out, float *  r_in, size_t  n, size_t  cols, float  p, float  eps  )

overridevirtual

AggregatorModule epsilon rule, per column.

Implements pulsatrix::DeviceBackend.

allocate()

void * pulsatrix::HIPBackend::allocate ( size_t  bytes )

overridevirtual

Allocates a buffer of the given size.

Parameters

bytes

Number of bytes to allocate. A request of 0 bytes returns nullptr by convention (not an error) — there is nothing to allocate.

Returns

Pointer to the allocated buffer, or nullptr iff bytes == 0.

Exceptions

std::runtime_error

if a nonzero-size allocation fails. Never silently returns nullptr for a nonzero request.

Implements pulsatrix::DeviceBackend.

avg_pool_backward()

void pulsatrix::HIPBackend::avg_pool_backward ( const float *  grad_out, float *  grad_in, size_t  planes, size_t  h, size_t  w, size_t  kh, size_t  kw  )

overridevirtual

Spreads grad_out / (kh*kw) over each window; grad_in must be zeroed by the caller.

Implements pulsatrix::DeviceBackend.

avg_pool_forward()

void pulsatrix::HIPBackend::avg_pool_forward ( const float *  in, float *  out, size_t  planes, size_t  h, size_t  w, size_t  kh, size_t  kw  )

overridevirtual

Non-overlapping average pool over planes of (h, w).

Implements pulsatrix::DeviceBackend.

axpby()

void pulsatrix::HIPBackend::axpby ( float  alpha, const float *  x, float  beta, const float *  y, float *  out, size_t  n  )

overridevirtual

out[i] = alpha * x[i] + beta * y[i]. out may alias x or y.

Note

beta == 0 does not read y (BLAS convention), so out = alpha * x exactly even where y holds inf/NaN – otherwise 0 * inf would turn a scale-only call into NaN.

Implements pulsatrix::DeviceBackend.

batch_norm_backward()

void pulsatrix::HIPBackend::batch_norm_backward ( const float *  grad_out, const float *  gamma, const float *  xhat, const float *  channel_std, float *  grad_in, float *  gamma_grad, float *  beta_grad, size_t  n, size_t  c, size_t  spatial  )

overridevirtual

BatchNorm input gradient plus this call's gamma/beta gradients (overwritten, per channel).

Implements pulsatrix::DeviceBackend.

batch_norm_eval_backward()

void pulsatrix::HIPBackend::batch_norm_eval_backward ( const float *  grad_out, const float *  gamma, const float *  xhat, const float *  channel_std, float *  grad_in, float *  gamma_grad, float *  beta_grad, size_t  n, size_t  c, size_t  spatial  )

overridevirtual

Eval-mode BatchNorm gradient: grad_out * gamma / std, plus gamma/beta gradients (overwritten, per channel). FND-5.

Implements pulsatrix::DeviceBackend.

batch_norm_eval_forward()

void pulsatrix::HIPBackend::batch_norm_eval_forward ( const float *  in, const float *  gamma, const float *  beta, const float *  running_mean, const float *  running_var, float *  xhat, float *  out, float *  channel_std, size_t  n, size_t  c, size_t  spatial, float  eps  )

overridevirtual

Eval-mode BatchNorm from the running statistics: a per-channel affine map. FND-5.

Implements pulsatrix::DeviceBackend.

batch_norm_forward()

void pulsatrix::HIPBackend::batch_norm_forward ( const float *  in, const float *  gamma, const float *  beta, float *  xhat, float *  out, float *  channel_std, size_t  n, size_t  c, size_t  spatial, float  eps  )

overridevirtual

BatchNorm over (n, spatial) per channel of (n, c, spatial) data.

Implements pulsatrix::DeviceBackend.

batch_norm_update_running()

void pulsatrix::HIPBackend::batch_norm_update_running ( const float *  in, float *  running_mean, float *  running_var, size_t  n, size_t  c, size_t  spatial, float  momentum  )

overridevirtual

Folds this batch's per-channel mean and unbiased variance into the running ones (PyTorch's momentum rule; the variance is kept when a channel has one value). FND-5.

Implements pulsatrix::DeviceBackend.

bce_with_logits()

void pulsatrix::HIPBackend::bce_with_logits ( const float *  logits, const float *  target, float *  out, size_t  n  )

overridevirtual

Per-element binary cross-entropy with logits: out[i] = max(x, 0) - x*y + log1p(exp(-|x|)), x = logits[i], y = target[i].

Note

Fused rather than composed from elementwise ops: this exact evaluation order is what BCEWithLogitsLoss has always computed, and no composition reproduces its rounding.

Implements pulsatrix::DeviceBackend.

bce_with_logits_grad()

void pulsatrix::HIPBackend::bce_with_logits_grad ( const float *  logits, const float *  target, float *  grad, size_t  n, float  scale  )

overridevirtual

BCE-with-logits gradient: grad[i] = (sigmoid(x) - y) * scale, using the overflow-free sigmoid (exp(x) / (1 + exp(x)) for x < 0).

Implements pulsatrix::DeviceBackend.

col2im_add()

void pulsatrix::HIPBackend::col2im_add ( const float *  col, float *  out, size_t  n, size_t  c, size_t  h, size_t  w, const ConvGeometry &  geometry  )

overridevirtual

Folds (n, c*kh*kw, out_h*out_w) patches back, adding into out (n, c, h, w).

Note

A gather: one GPU thread per output pixel sums its window contributions in the same order the CPU scatter loop adds them – deterministic, no atomics.

Implements pulsatrix::DeviceBackend.

column_sums()

void pulsatrix::HIPBackend::column_sums ( const float *  in, float *  out, size_t  rows, size_t  cols, float  beta  )

overridevirtual

Per-column sum of a (rows x cols) row-major matrix: out[j] = beta*out[j] + sum_i in[i][j].

Note

Rows are summed in increasing i on every backend. out must not alias in.

Implements pulsatrix::DeviceBackend.

copy()

void pulsatrix::HIPBackend::copy ( void *  dst, const void *  src, size_t  bytes, CopyDirection  dir  )

overridevirtual

Copies bytes between buffers.

Parameters

dst	Destination buffer, must be large enough to hold bytes.
src	Source buffer.
bytes	Number of bytes to copy.
dir	Direction of the copy (informs device-specific implementations which memory space each pointer lives in; CPUBackend ignores it).

Implements pulsatrix::DeviceBackend.

copy_2d()

void pulsatrix::HIPBackend::copy_2d ( float *  dst, size_t  dst_stride, const float *  src, size_t  src_stride, size_t  rows, size_t  cols  )

overridevirtual

Strided 2-D copy: rows of cols floats from src (row stride src_stride) to dst (row stride dst_stride) – e.g. one timestep of an (N, L, D) sequence.

Implements pulsatrix::DeviceBackend.

device()

DeviceType pulsatrix::HIPBackend::device ( ) const

inlineoverridevirtualnoexcept

Which device this backend's buffers reside on.

Note

Tensor's constructors that take no explicit DeviceType tag the Tensor with this, so a temporary allocated through a CUDA/HIP backend can no longer be silently labelled Cpu (the default-tag defect found by the GPU-native-kernels campaign recon: a mislabelled tensor makes Tensor pick HostToHost copies on device memory and lets host-only guards pass).

Implements pulsatrix::DeviceBackend.

dot()

float pulsatrix::HIPBackend::dot ( const float *  a, const float *  b, size_t  n  )

overridevirtual

Dot product sum_i a[i]*b[i], returned to the host.

Note

Synchronizes. Fixed-order reduction on every backend, but GPU order differs from CPU's sequential sum, so CPU and GPU agree to rounding, not bitwise.

Implements pulsatrix::DeviceBackend.

dropout_forward()

void pulsatrix::HIPBackend::dropout_forward ( const float *  in, float *  out, float *  mask, size_t  n, float  p, float  scale, uint64_t  seed, uint64_t  offset  )

overridevirtual

Inverted dropout with a counter-based RNG: element i is dropped iff uniform(seed, offset + i) < p; kept elements are scaled by scale.

Parameters

mask	Receives 1.0 (kept) or 0.0 (dropped) per element, for backward().

Note

uniform(seed, k) is splitmix64(seed + (k + 1) * golden_gamma), top 24 bits as a float in [0, 1). Stateless and identical on every backend, so a GPU mask is bit-identical to the CPU one for the same (seed, offset) – the property that lets Dropout be tested CPU-vs-GPU at all. A dropped element is written as 0 (a select, not in * 0), so an inf/NaN input there still yields 0.

Implements pulsatrix::DeviceBackend.

elementwise()

void pulsatrix::HIPBackend::elementwise ( ElementwiseOp  op, const float *  in, float *  out, size_t  n  )

overridevirtual

Applies a unary elementwise operation to every element of a buffer.

Parameters

op	Operation to apply.
in	Input buffer, must hold at least n floats.
out	Output buffer, must hold at least n floats. May alias in for an in-place application.
n	Number of elements.

Implements pulsatrix::DeviceBackend.

elementwise_backward()

void pulsatrix::HIPBackend::elementwise_backward ( ElementwiseOp  op, const float *  x, const float *  grad_out, float *  grad_in, size_t  n  )

overridevirtual

Activation backward: grad_in[i] = grad_out[i] * f'(x[i]), f = op, x = the forward input.

Note

Derivatives: Relu selects grad_out where x > 0, else 0 (0 at x == 0 and for a non-finite grad_out, matching ReluModule); Neg -1; Tanh 1 - tanh(x)^2; Sigmoid s(1 - s); Silu s + x*s*(1 - s), s = sigmoid(x); Exp exp(x). grad_in may alias grad_out or x.

Implements pulsatrix::DeviceBackend.

fill()

void pulsatrix::HIPBackend::fill ( void *  ptr, float  value, size_t  n  )

overridevirtual

Fills every element of a float buffer with a constant value.

Parameters

ptr	Buffer to fill, must hold at least n floats.
value	Fill value.
n	Number of elements to fill.

Implements pulsatrix::DeviceBackend.

free()

void pulsatrix::HIPBackend::free ( void *  ptr )

overridevirtualnoexcept

Frees a buffer previously returned by allocate(). Safe to call with nullptr.

Implements pulsatrix::DeviceBackend.

gather_rows()

void pulsatrix::HIPBackend::gather_rows ( const float *  table, const float *  indices, float *  out, size_t  count, size_t  dim  )

overridevirtual

out[i][:] = table[indices[i]][:] for count rows of width dim.

Implements pulsatrix::DeviceBackend.

gemm()

void pulsatrix::HIPBackend::gemm ( const float *  a, const float *  b, float *  out, size_t  m, size_t  k, size_t  n  )

overridevirtual

Row-major matrix multiply: out = a * b.

Parameters

a	Pointer to A (m x k), row-major.
b	Pointer to B (k x n), row-major.
out	Pointer to the output buffer (m x n), row-major. Must be pre-allocated by the caller.
m	Rows of A / rows of out.
k	Columns of A / rows of B.
n	Columns of B / columns of out.

Implements pulsatrix::DeviceBackend.

gemm_ex()

void pulsatrix::HIPBackend::gemm_ex ( const float *  a, bool  transpose_a, const float *  b, bool  transpose_b, float *  out, size_t  m, size_t  k, size_t  n, float  beta  )

overridevirtual

General row-major matrix multiply: out = op(A) * op(B) + beta * out.

Parameters

a	A, stored (m x k) row-major, or (k x m) when transpose_a.
transpose_a	Use A^T.
b	B, stored (k x n) row-major, or (n x k) when transpose_b.
transpose_b	Use B^T.
out	(m x n) row-major. Read only when beta != 0 (beta == 0 overwrites, so out may hold garbage). Must not alias a or b.
beta	Scale on the existing out; 1 accumulates (gradient accumulation).

Note

Replaces the explicit host-side transpose() copies modules built to feed gemm().

Implements pulsatrix::DeviceBackend.

group_norm_backward()

void pulsatrix::HIPBackend::group_norm_backward ( const float *  grad_out, const float *  gamma, const float *  xhat, const float *  group_std, float *  grad_in, float *  gamma_grad, float *  beta_grad, size_t  n, size_t  c, size_t  spatial, size_t  num_groups  )

overridevirtual

GroupNorm input gradient plus this call's gamma/beta gradients (overwritten, per channel).

Implements pulsatrix::DeviceBackend.

group_norm_forward()

void pulsatrix::HIPBackend::group_norm_forward ( const float *  in, const float *  gamma, const float *  beta, float *  xhat, float *  out, float *  group_std, size_t  n, size_t  c, size_t  spatial, size_t  num_groups, float  eps  )

overridevirtual

GroupNorm per (example, group) of (n, c, spatial) data; group_std is (n, num_groups).

Implements pulsatrix::DeviceBackend.

gru_lrp_hprev()

void pulsatrix::HIPBackend::gru_lrp_hprev ( const float *  h_prev, const float *  w_hn, const float *  hn, const float *  r_term_b, const float *  direct, float *  r_hprev, size_t  rows, size_t  hidden, float  eps  )

overridevirtual

GRU's R_hprev for (rows, hidden): the recurrent epsilon-rule sum through W_hn, with each element's direct term inserted where the original loop added it.

Implements pulsatrix::DeviceBackend.

im2col()

void pulsatrix::HIPBackend::im2col ( const float *  in, float *  col, size_t  n, size_t  c, size_t  h, size_t  w, const ConvGeometry &  geometry  )

overridevirtual

Unfolds (n, c, h, w) into (n, c*kh*kw, out_h*out_w) patches, with out_h = (h + 2*pad_h - kh) / stride_h + 1 (likewise out_w). Taps that fall in the zero padding read 0.

Implements pulsatrix::DeviceBackend.

layer_norm_backward()

void pulsatrix::HIPBackend::layer_norm_backward ( const float *  grad_out, const float *  gamma, const float *  xhat, const float *  row_std, float *  grad_in, size_t  rows, size_t  cols  )

overridevirtual

LayerNorm input gradient per row, from the cached xhat and per-row std.

Implements pulsatrix::DeviceBackend.

layer_norm_forward()

void pulsatrix::HIPBackend::layer_norm_forward ( const float *  in, const float *  gamma, const float *  beta, float *  xhat, float *  out, float *  row_std, size_t  rows, size_t  cols, float  eps  )

overridevirtual

LayerNorm forward per row: xhat = (x - mean)/sqrt(var + eps), out = gamma*xhat + beta.

Implements pulsatrix::DeviceBackend.

logic_pointwise()

void pulsatrix::HIPBackend::logic_pointwise ( LogicOp  op, int  norm, const float *  a, const float *  b, const float *  g_or_r, const float *  y, float *  out_a, float *  out_b, size_t  n, float  eps  )

overridevirtual

One elementwise pass of Conjunction/Disjunction for operands a, b.

Parameters

g_or_r	Upstream gradient (Backward) or relevance (Lrp); unused by Forward.
y	Cached forward output (Lrp only).
out_a	Forward: the output. Backward/Lrp: a's gradient/relevance.
out_b	Backward/Lrp: b's gradient/relevance; unused by Forward.

Implements pulsatrix::DeviceBackend.

logsumexp_rows()

void pulsatrix::HIPBackend::logsumexp_rows ( const float *  in, float *  out, size_t  rows, size_t  cols  )

overridevirtual

Per-row log-sum-exp, max-subtracted: out[i] = log sum_j exp(in[i][j]).

Implements pulsatrix::DeviceBackend.

lrp_avg_pool()

void pulsatrix::HIPBackend::lrp_avg_pool ( const float *  x, const float *  r, float *  r_in, size_t  planes, size_t  h, size_t  w, size_t  kh, size_t  kw, float  eps  )

overridevirtual

AvgPool2D epsilon rule; r_in must be zeroed by the caller.

Implements pulsatrix::DeviceBackend.

lrp_bilinear_elementwise()

void pulsatrix::HIPBackend::lrp_bilinear_elementwise ( const float *  a, const float *  b, const float *  r, float *  r_out, size_t  n, float  eps  )

overridevirtual

Eq. 15 for an elementwise product c = a*b: r_out = (a b / (2c + eps sign c)) r (same for both).

Implements pulsatrix::DeviceBackend.

lrp_bilinear_matmul()

void pulsatrix::HIPBackend::lrp_bilinear_matmul ( const float *  a, const float *  b, const float *  o, const float *  r_o, float *  r_a, float *  r_b, size_t  slices, size_t  m, size_t  p, size_t  q, float  eps, bool  b_transposed  )

overridevirtual

Eq. 15 for slices independent matmuls O = A @ B (A (M x P), B (P x Q), O and r_o (M x Q)).

Parameters

b_transposed

B is stored as B^T (Q x P); r_b is then written in that same layout.

Note

r_a and r_b are overwritten (not accumulated into).

Implements pulsatrix::DeviceBackend.

lrp_conv()

void pulsatrix::HIPBackend::lrp_conv ( const float *  col, const float *  kernel, const float *  pre_bias, const float *  r, float *  r_col, size_t  n, size_t  out_channels, size_t  p, size_t  q, float  eps  )

overridevirtual

Conv2D epsilon rule in patch space: r_col (n, p, q) from the cached patches, the kernel (out_channels, p) and the pre-bias outputs (n, out_channels, q).

Implements pulsatrix::DeviceBackend.

lrp_linear()

void pulsatrix::HIPBackend::lrp_linear ( const float *  x, const float *  w, const float *  z, const float *  r, float *  r_in, size_t  rows, size_t  in_features, size_t  out_features, float  eps  )

overridevirtual

LinearModule epsilon rule: r_in[n][i] = sum_j (x[n][i] w[i][j] / stab(z[n][j])) r[n][j].

Parameters

z	(N, out) pre-bias outputs; w (in, out); x, r_in (N, in); r (N, out).

Implements pulsatrix::DeviceBackend.

lrp_residual_split()

void pulsatrix::HIPBackend::lrp_residual_split ( const float *  a, const float *  b, const float *  r, float *  r_a, float *  r_b, size_t  n, float  eps  )

overridevirtual

Epsilon split of a residual sum y = a + b: r_a = (a / stab(y)) r, r_b = (b / stab(y)) r.

Implements pulsatrix::DeviceBackend.

lrp_rope()

void pulsatrix::HIPBackend::lrp_rope ( const float *  x, const float *  y, const float *  r, const float *  cos_table, const float *  sin_table, float *  r_in, size_t  slices, size_t  seq_len, size_t  head_dim, float  eps  )

overridevirtual

RoPEModule epsilon rule over (slices, seq_len, head_dim), tables as for rope_rotate.

Implements pulsatrix::DeviceBackend.

lrp_softmax_rows()

void pulsatrix::HIPBackend::lrp_softmax_rows ( const float *  x, const float *  y, const float *  r, float *  r_in, size_t  rows, size_t  cols  )

overridevirtual

SoftmaxModule rule per row: r_in = x * (r - y * sum(r)).

Implements pulsatrix::DeviceBackend.

lrp_stabilized_divide()

void pulsatrix::HIPBackend::lrp_stabilized_divide ( const float *  r, const float *  denom, const float *  gate, float *  out, size_t  n, float  eps, LrpGate  gate_mode  )

overridevirtual

Gated stabilized division, the one non-gemm step of the affine LRP rules: out[i] = passes(gate[i]) ? r[i] / (denom[i] + eps sign(denom[i])) : 0, sign(0) = +1.

Parameters

gate	Read only when gate_mode != LrpGate::None (may then be nullptr).

Note

out may alias r or denom.

Implements pulsatrix::DeviceBackend.

max_pool_forward()

void pulsatrix::HIPBackend::max_pool_forward ( const float *  in, float *  out, float *  argmax, size_t  planes, size_t  h, size_t  w, size_t  kh, size_t  kw  )

overridevirtual

Non-overlapping max pool over planes of (h, w); argmax = flat in-plane index (first max wins).

Implements pulsatrix::DeviceBackend.

max_unpool()

void pulsatrix::HIPBackend::max_unpool ( const float *  src, const float *  argmax, float *  dst, size_t  planes, size_t  h, size_t  w, size_t  kh, size_t  kw  )

overridevirtual

dst[plane][argmax] = src for every pooled element; dst must be zeroed by the caller.

Implements pulsatrix::DeviceBackend.

mul()

void pulsatrix::HIPBackend::mul ( const float *  a, const float *  b, float *  out, size_t  n  )

overridevirtual

Elementwise binary (Hadamard) multiplication: out[i] = a[i] * b[i] for i in [0, n).

Parameters

a	First operand, must hold at least n floats.
b	Second operand, must hold at least n floats.
out	Output buffer, must hold at least n floats. May alias a or b for an in-place application.
n	Number of elements.

Note

Added in Phase 6's SwiGLU mission (mission_swiglu.md) – the third consumer needing a Hadamard product via a raw host loop (after LSTMModule's/GRUModule's gate arithmetic), the trigger point the Risk Register named for finally adding this primitive properly instead of a fourth raw loop. LSTMModule/GRUModule's existing raw loops are not retrofitted to use it – no behavior change needed there, logged as a low-priority future cleanup only.

Implements pulsatrix::DeviceBackend.

operator=()

HIPBackend & pulsatrix::HIPBackend::operator= ( const HIPBackend &  )

delete

permute_0213()

void pulsatrix::HIPBackend::permute_0213 ( const float *  in, float *  out, size_t  d0, size_t  d1, size_t  d2, size_t  d3  )

overridevirtual

Swaps the middle two axes: in (d0, d1, d2, d3) -> out (d0, d2, d1, d3).

Note

Attention's head split ((N, L, H, D) -> (N, H, L, D)) and merge (the reverse) are both this permutation. out must not alias in.

Implements pulsatrix::DeviceBackend.

recurrent_cell()

void pulsatrix::HIPBackend::recurrent_cell ( RecurrentCellOp  op, const RecurrentCellArgs &  args, size_t  n  )

overridevirtual

One fused recurrent-cell pass over n elements (see RecurrentCellOp for slots).

Implements pulsatrix::DeviceBackend.

rl_rows()

void pulsatrix::HIPBackend::rl_rows ( RlRowOp  op, const RlRowArgs &  args  )

overridevirtual

One fused RL loss / target / Polyak pass (see RlRowOp for lanes and slots).

Note

One lane per row: a row's softmax, argmax and gradient writes are owned by its lane and walk the columns in the original loop order. Per-row loss terms are reduced by the caller with column_sums, which adds rows in increasing order like the original loop.

Implements pulsatrix::DeviceBackend.

rms_norm_backward()

void pulsatrix::HIPBackend::rms_norm_backward ( const float *  grad_out, const float *  gamma, const float *  in, const float *  row_rms, float *  grad_in, float *  gamma_terms, size_t  rows, size_t  cols  )

overridevirtual

RMSNorm input gradient per row, plus gamma_terms[r][i] = grad_out * x / rms – the per-row contributions column_sums then reduces into gamma's gradient.

Implements pulsatrix::DeviceBackend.

rms_norm_forward()

void pulsatrix::HIPBackend::rms_norm_forward ( const float *  in, const float *  gamma, float *  out, float *  row_rms, size_t  rows, size_t  cols, float  eps  )

overridevirtual

RMSNorm forward per row: out = gamma * x / sqrt(mean(x^2) + eps).

Implements pulsatrix::DeviceBackend.

rope_rotate()

void pulsatrix::HIPBackend::rope_rotate ( const float *  in, const float *  cos_table, const float *  sin_table, float *  out, size_t  num_slices, size_t  seq_len, size_t  head_dim, bool  inverse  )

overridevirtual

Rotary position embedding over (num_slices, seq_len, head_dim) data.

Parameters

cos_table,sin_table	(seq_len, head_dim/2) per-position rotation tables.
inverse	false: forward rotation; true: its transpose (RoPE's backward).

Note

out must not alias in.

Implements pulsatrix::DeviceBackend.

scatter_add_rows()

void pulsatrix::HIPBackend::scatter_add_rows ( const float *  src, const float *  indices, float *  table, size_t  count, size_t  dim  )

overridevirtual

table[indices[i]][:] += src[i][:] for i = 0..count-1, in increasing i.

Note

Deterministic: the GPU kernel runs one thread per column and walks i in order – repeated indices accumulate in the same order as the CPU, no atomics.

Implements pulsatrix::DeviceBackend.

softmax_rows()

void pulsatrix::HIPBackend::softmax_rows ( const float *  in, float *  out, size_t  rows, size_t  cols  )

overridevirtual

Row-wise softmax of a (rows x cols) matrix, max-subtracted. out may alias in.

Implements pulsatrix::DeviceBackend.

softmax_rows_backward()

void pulsatrix::HIPBackend::softmax_rows_backward ( const float *  y, const float *  dy, float *  dx, size_t  rows, size_t  cols  )

overridevirtual

Softmax backward from its output y: dx[i][j] = y[i][j] * (dy[i][j] - sum_k y[i][k]*dy[i][k]).

Note

dx must not alias y or dy.

Implements pulsatrix::DeviceBackend.

ssm_pass()

void pulsatrix::HIPBackend::ssm_pass ( SsmPassOp  op, const SsmPassArgs &  args  )

overridevirtual

One fused Mamba / RWKV / RetNet pass (see SsmPassOp for lanes and slots).

Note

Recurrences run one lane per independent (batch, channel[, state]) sequence, walking time in order inside the lane – parallel over lanes, sequential over time, the CPU's order. Every reduction is owned by one lane and summed in the original loop order.

Implements pulsatrix::DeviceBackend.

sum()

float pulsatrix::HIPBackend::sum ( const float *  in, size_t  n  )

overridevirtual

sum_i in[i], returned to the host. Same reduction order as dot(). Synchronizes.

Implements pulsatrix::DeviceBackend.

tanh_gaussian_backward()

void pulsatrix::HIPBackend::tanh_gaussian_backward ( const float *  action, const float *  std_cache, const float *  eps, const float *  grad_action, const float *  grad_log_prob, float *  grad_mean, float *  grad_log_std, size_t  n, float  stabilizer  )

overridevirtual

TanhGaussianPolicy gradients w.r.t. mean and log_std, per element.

Implements pulsatrix::DeviceBackend.

tanh_gaussian_forward()

void pulsatrix::HIPBackend::tanh_gaussian_forward ( const float *  mean, const float *  log_std, const float *  eps, float *  action, float *  std_cache, float *  log_prob, size_t  rows, size_t  cols, float  stabilizer, double  half_log_two_pi  )

overridevirtual

TanhGaussianPolicy sampling per (rows, cols) row: action = tanh(mean + exp(log_std)*eps), std_cache = exp(log_std), log_prob[r] accumulated in double.

Implements pulsatrix::DeviceBackend.

top_k_rows()

void pulsatrix::HIPBackend::top_k_rows ( const float *  in, float *  values, float *  indices, size_t  rows, size_t  cols, size_t  k, bool  largest  )

overridevirtual

Per row of a row-major (rows, cols) matrix: the k largest (or smallest) values in rank order into values (rows, k), and their column indices into indices (rows, k) as whole-number floats.

Note

Order: NaN ranks above every number; equal values keep the lower column index first. One lane per row, pure selection: every backend's output is bit-identical.

Preconditions, validated by the caller (top_k()): 1 <= k <= cols, cols <= 2^24.

Implements pulsatrix::DeviceBackend.

---

The documentation for this class was generated from the following file:

• include/pulsatrix/hip_backend.hpp