57 lines
2.6 KiB
Plaintext
57 lines
2.6 KiB
Plaintext
// CelerisLab – streaming/pull_double_buffer.cuh
|
||
// Double-buffer PULL streaming.
|
||
// f_local[i] = f_in[ neighbor(n, -c_i), i ]
|
||
// f_out[n, i] = f_collided[i] (after collision, written to second buffer)
|
||
//
|
||
// This replaces the old PUSH scheme where each thread writes to neighbors.
|
||
// Pull streaming provides better memory coalescing on GPU.
|
||
// ============================================================================
|
||
|
||
#ifndef CELERIS_STREAMING_PULL_DOUBLE_BUFFER_CUH
|
||
#define CELERIS_STREAMING_PULL_DOUBLE_BUFFER_CUH
|
||
|
||
// ---------------------------------------------------------------------------
|
||
// stream_pull_load: gather DDF from neighboring nodes into local f[NQ]
|
||
// n – current node
|
||
// f – output: local distribution (compute precision, float)
|
||
// fi – input: source DDF array (storage precision, fpxx)
|
||
// j – neighbor indices (j[i] = neighbor in direction i)
|
||
//
|
||
// Pull semantic: f[i] at node n = fi[ j[opp(i)], i ]
|
||
// i.e. the population traveling in direction i at node n
|
||
// arrived from the node at n - c_i = neighbor in direction opp(i)
|
||
// ---------------------------------------------------------------------------
|
||
__device__ __forceinline__ void stream_pull_load(unsigned long n,
|
||
float* __restrict__ f,
|
||
const fpxx* __restrict__ fi,
|
||
const unsigned long* __restrict__ j)
|
||
{
|
||
// i=0 (rest): always from self
|
||
f[0] = load_ddf(fi, index_f(n, 0u));
|
||
|
||
// For paired ordering, opp(i) swaps within each pair:
|
||
// opp(1)=2, opp(2)=1, opp(3)=4, opp(4)=3, ...
|
||
#pragma unroll
|
||
for (int i = 1; i < NQ; i += 2) {
|
||
// population i arrives from direction opp(i) = i+1
|
||
f[i] = load_ddf(fi, index_f(j[i+1], (unsigned int)i));
|
||
// population i+1 arrives from direction opp(i+1) = i
|
||
f[i + 1] = load_ddf(fi, index_f(j[i], (unsigned int)(i + 1)));
|
||
}
|
||
}
|
||
|
||
// ---------------------------------------------------------------------------
|
||
// stream_pull_store: write collided f[NQ] to the OUTPUT buffer at node n
|
||
// ---------------------------------------------------------------------------
|
||
__device__ __forceinline__ void stream_pull_store(unsigned long n,
|
||
const float* __restrict__ f,
|
||
fpxx* __restrict__ fi_out)
|
||
{
|
||
#pragma unroll
|
||
for (int i = 0; i < NQ; i++) {
|
||
store_ddf(fi_out, index_f(n, (unsigned int)i), f[i]);
|
||
}
|
||
}
|
||
|
||
#endif // CELERIS_STREAMING_PULL_DOUBLE_BUFFER_CUH
|