From 9233f7fdc849c8249e50d45ba435890c0c035083 Mon Sep 17 00:00:00 2001
From: Hayodea Hekol <hayodea@protonmail.com>
Date: Sat, 8 Nov 2025 10:26:17 -0400
Subject: [PATCH] livoxG1: Add OpenCl kernels for collation

---
 stimBuffApis/livoxGen1/CMakeLists.txt     | 11 +++++++++++
 stimBuffApis/livoxGen1/collateDgrams.cl   |  3 +++
 stimBuffApis/livoxGen1/openClKernels.cl.S | 17 +++++++++++++++++
 stimBuffApis/livoxGen1/openClKernels.h    | 17 +++++++++++++++++
 4 files changed, 48 insertions(+)
 create mode 100644 stimBuffApis/livoxGen1/collateDgrams.cl
 create mode 100644 stimBuffApis/livoxGen1/openClKernels.cl.S
 create mode 100644 stimBuffApis/livoxGen1/openClKernels.h

diff --git a/stimBuffApis/livoxGen1/CMakeLists.txt b/stimBuffApis/livoxGen1/CMakeLists.txt
index 93a56b6..40e732a 100644
--- a/stimBuffApis/livoxGen1/CMakeLists.txt
+++ b/stimBuffApis/livoxGen1/CMakeLists.txt
@@ -37,12 +37,23 @@ if(ENABLE_STIMBUFFAPI_livoxGen1)
 		message(STATUS "Found OpenCL using pkg-config")
 	endif()
 
+	# Enable assembly language
+	enable_language(ASM)
+
 	add_library(livoxGen1 SHARED
 		livoxGen1.cpp
 		stagingBuffer.cpp
 		pcloudStimulusBuffer.cpp
 		ioUringAssemblyEngine.cpp
 		openClCollatingAndMeshingEngine.cpp
+		openClKernels.cl.S
+	)
+
+	# Set assembler working directory so .incbin can find the .cl file
+	# Also declare dependency on collateDgrams.cl
+	set_source_files_properties(openClKernels.cl.S PROPERTIES
+		COMPILE_FLAGS "-I${CMAKE_CURRENT_SOURCE_DIR}"
+		OBJECT_DEPENDS "${CMAKE_CURRENT_SOURCE_DIR}/collateDgrams.cl"
 	)
 
 	target_include_directories(livoxGen1 PUBLIC
diff --git a/stimBuffApis/livoxGen1/collateDgrams.cl b/stimBuffApis/livoxGen1/collateDgrams.cl
new file mode 100644
index 0000000..36a83f6
--- /dev/null
+++ b/stimBuffApis/livoxGen1/collateDgrams.cl
@@ -0,0 +1,3 @@
+__kernel void collate(__global uchar* assembly, __global uchar* collation) {
+	// Placeholder kernel - will be replaced with actual collation logic
+}
diff --git a/stimBuffApis/livoxGen1/openClKernels.cl.S b/stimBuffApis/livoxGen1/openClKernels.cl.S
new file mode 100644
index 0000000..87d50ee
--- /dev/null
+++ b/stimBuffApis/livoxGen1/openClKernels.cl.S
@@ -0,0 +1,17 @@
+	.section .rodata
+	.global collateKernelStart
+	.global collateKernelEnd
+	.global collateKernelNBytes
+	.type collateKernelStart, @object
+	.type collateKernelEnd, @object
+	.type collateKernelNBytes, @object
+
+collateKernelStart:
+	.incbin "collateDgrams.cl"
+collateKernelEnd:
+
+	.section .data
+collateKernelNBytes:
+	.quad collateKernelEnd - collateKernelStart
+
+	.section .note.GNU-stack,"",@progbits
diff --git a/stimBuffApis/livoxGen1/openClKernels.h b/stimBuffApis/livoxGen1/openClKernels.h
new file mode 100644
index 0000000..9dce67c
--- /dev/null
+++ b/stimBuffApis/livoxGen1/openClKernels.h
@@ -0,0 +1,17 @@
+#ifndef _LIVOX_GEN1_OPENCL_KERNELS_H
+#define _LIVOX_GEN1_OPENCL_KERNELS_H
+
+#include <cstddef>
+
+namespace smo {
+namespace stim_buff {
+
+// External symbols for collate kernel
+extern const char collateKernelStart[];
+extern const char collateKernelEnd[];
+extern const size_t collateKernelNBytes;
+
+} // namespace stim_buff
+} // namespace smo
+
+#endif // _LIVOX_GEN1_OPENCL_KERNELS_H