Source code
extern __shared__ cudafloat iw[];
iw[threadIdx.x] = CUDA_VALUE(0.0);
for(int i = threadIdx.x; i <= numInputs; i += blockDim.x) {
cudafloat i_w = weights[NEURON * (numInputs + 1) + i];
if (i > BIAS) i_w *= inputs[PATTERN * numInputs + (i - 1)];
iw[threadIdx.x] += i_w;
}
__syncthreads();
SumBeforeWarp<blockSize>(iw);
if (threadIdx.x < 32) {
SumWarp<blockSize>(iw);
if (threadIdx.x == 0) {
int n = PATTERN * NUM_NEURONS + NEURON;
int nSelAct = PATTERN * totalNeuronsWithSelectiveActivation + NEURON + mOffset;
cudafloat output = CUDA_SIGMOID(iw[0]);
cudafloat M = (m != nullptr) ? m[nSelAct] : CUDA_VALUE(1.0);
cudafloat outn = output * M;
cudafloat error = (desiredOutputs[n] - outn);
if (m != nullptr) localGradientSpaceNet[nSelAct] = error * output * CUDA_SIGMOID_DERIVATE(M);
outputs[n] = outn;
localGradient[n] = error * M * CUDA_SIGMOID_DERIVATE(output);
Run 1
The code has a clear overall structure (shared memory initialization, parallel reduction, single-thread output computation) but suffers from several readability issues. The logic is reasonably organized, yet variable names like `iw`, `M`, `n`, and `nSelAct` are cryptic. Nesting is moderate with only one level of control flow inside the final conditional. The main problem is action density: multiple operations are packed into single lines (e.g., `iw[threadIdx.x] += i_w;` combined with conditional multiplication). Reading flow is disrupted by inconsistent spacing, multi-statement lines, and magic constants like `32` and `blockSize`. The intent is inferable but not obvious—especially the selective activation logic and gradient computations. Overall, an experienced CUDA programmer could parse this, but it requires effort and domain knowledge.
Run 2
The code has a clear high-level structure: initialization, parallel reduction, and final output computation. However, readability suffers due to several issues. Variable names like 'iw', 'n', 'M', and 'outn' are cryptic and don't convey intent. The logic for handling BIAS, weights, and selective activation is compressed into a few dense lines, making the mathematical relationships hard to follow without domain knowledge. The reduction pattern is split across two functions, which adds clarity, but the conditional 'if (m != nullptr)' checks appear twice, increasing cognitive load. Nesting is shallow but the code within the 'if (threadIdx.x == 0)' block packs many operations together, weakening action granularity. The code is concise but at the cost of clarity, with multiple operations per line and unclear data flow. Reading flow is somewhat natural, following the execution order, but the lack of comments and descriptive names obscures the purpose of each computation.
Run 3
The code has a fairly logical structure and avoids excessive nesting, but it suffers from poor clarity of intent due to cryptic variable names (e.g., 'iw', 'nSelAct', 'mOffset'), unexplained numerical constants (like 32 and 1.0), and heavy reliance on macros (CUDA_VALUE, CUDA_SIGMOID_DERIVATE). The reading flow is disrupted by implicit domain-specific logic (bias handling, selective activation) that is not self-evident. Action granularity is mostly good—one meaningful operation per line—but the dense arithmetic in the final block reduces readability. The shared memory reduction pattern is recognizable but could benefit from comments or clearer naming. Overall, the code is reasonably concise and structured, but ambiguous naming and domain assumptions lower its readability.
Run 1
Code Structure: The sequence of accumulation, synchronization, reduction, activation, and gradient calculation is logically organized, but the snippet ends without closing braces. Nesting: Moderate nesting reflects CUDA execution constraints, though the nested thread checks add indentation. Clarity of intent: Familiar operations and descriptive array names help, but abbreviations such as iw, i_w, M, and outn, unexplained indexing offsets, and undocumented reduction helpers require domain knowledge. Code length: Compact, with little obvious redundancy. Action granularity: Most lines perform one action; inline conditional assignments slightly reduce scanability. Reading flow: Blank lines separate stages naturally, but inconsistent spacing, missing explanatory comments, and the incomplete ending interrupt comprehension.
Run 2
Code Structure: Organized into weighted accumulation, reduction, activation, and gradient calculation, but the snippet ends without closing braces. Nesting: Moderate nesting through a loop and two thread-selection conditions; manageable but less easy to scan. Clarity of intent: CUDA operations and descriptive array names indicate the purpose, while terse identifiers such as iw, M, n, and nSelAct and unexplained indexing conventions require domain knowledge. Code length: Compact and mostly necessary, though long indexing expressions add cognitive load. Action granularity: Generally one action per line, but inline conditional assignments compress conditions and operations together. Reading flow: Execution proceeds naturally from initialization to outputs and gradients, with whitespace separating stages; absent comments, mixed naming styles, and the incomplete ending reduce readability.
Run 3
Code Structure: Organized into accumulation, reduction, activation, and gradient computation, but the missing closing braces make the snippet incomplete. Nesting: Moderate nesting combines a loop, warp condition, and thread-zero condition without becoming deeply tangled. Clarity of intent: CUDA operations and output calculations suggest neural-network processing, but terse names such as iw, M, n, and nSelAct and unexplained indexing reduce clarity. Code length: Compact and mostly purposeful, though complex index expressions require careful reading. Action granularity: Most lines perform one action; inline conditional assignments slightly reduce scanability. Reading flow: Execution proceeds naturally from weighted inputs to outputs and gradients, with whitespace separating stages, but macro-heavy expressions and CUDA-specific reduction helpers require domain knowledge.