Warrior_EA/DirectML/opencl_seq_syntax_check.cpp

14 lines
605 B
C++
Raw Permalink Normal View History

feat(ai): sequence-LSTM kernels for the OpenCL tier Closes the gap left by 7a08197, which refused sequence mode under OpenCL. That was defensible for a private build and not for a shipped one: the release path includes an OpenCL laptop, and a customer with a GPU would have found LSTM and HYBRID simply unavailable. One launch PER TIMESTEP rather than a single kernel looping with barrier(). Every hidden unit's gates read all of h_{t-1}, OpenCL barriers only span a work-group, and nothing here constrains how the runtime partitions the global size - so an in-kernel loop would be correct only by luck of the partitioning. Host-driven launches make each step an implicit global barrier: more enqueues, correct on every device. Backward reuses the buffers the single-timestep path leaves idle in sequence mode - ConcatenatedGradient (4H) for gate gradients, HiddenCache (H) for dh, Memory (2H) for dc - so BPTT costs no extra allocations. dW is zeroed once and accumulated across steps, matching the fused DLL kernel. Verification available on this machine has limits worth recording. The math is the same as CPU_LSTMSeqForward/Backward, which is gradient-checked to 2.3e-10; the kernels are syntax/type-checked offline (DirectML\opencl_seq_syntax_check.cpp, compiled as C++ with OpenCL shims) because there is no OpenCL device or ICD here. That check exists because a typo in Network.cl fails the WHOLE program build, which would take the dense and conv kernels down with it - not just the new ones. KernelCreate results are now checked and reported for these four for the same reason; a build failure degrades to "LSTM/HYBRID unavailable on this device" instead of an Execute error mid-training. STILL NEEDS A RUN ON REAL OPENCL HARDWARE before release. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
2026-07-30 18:40:56 -04:00
// Offline syntax/type check for the sequence-LSTM OpenCL kernels.
// Not a semantic test - it verifies the kernel bodies are well-formed C so a
// typo cannot fail the whole Network.cl program build on a customer's GPU
// (which would take the dense and conv kernels down with it, not just LSTM).
#include <cmath>
#define __kernel
#define __global
#define MIN_ACTIVATION_DERIVATIVE 1.0e-3f
static int g_id = 0;
static inline int get_global_id(int) { return g_id; }
static inline float fmaxf_(float a, float b) { return a > b ? a : b; }
#define fmax fmaxf_
#include "kernels.inc"
int main() { return 0; }