Use x64 numbers

raub · raub · commit 365b43d1828a · 2025-04-03T20:26:35.000+04:00
diff --git a/README.md b/README.md
@@ -37,11 +37,8 @@ See [TypeScript declarations](/index.d.ts) for more details.
 	```
 2. Fetch the CL control objects:
 	```ts
-	const platform = cl.getPlatformIDs()[0];
-	const devices = cl.getDeviceIDs(platform, cl.DEVICE_TYPE_ALL);
-	const context = cl.createContext([cl.CONTEXT_PLATFORM, platform], devices);
-	const device = cl.getContextInfo(context, cl.CONTEXT_DEVICES)[0];
-	const queue = cl.createCommandQueue(context, device, null);
+	const { context, device } = cl.quickStart(); // see /index.js
+	const queue = cl.createCommandQueue(context, device);
 	```
 3. Prepare the data input/output buffers:
 	```ts
@@ -67,7 +64,8 @@ See [TypeScript declarations](/index.d.ts) for more details.
 4. Create a valid CL program, e.g. from source:
 	```ts
 	const program = cl.createProgramWithSource(context, `
-		__kernel void vadd(__global int *a, __global int *b, __global int *c, uint num) {
+		__kernel
+		void vadd(__global int *a, __global int *b, __global int *c, uint num) {
 			size_t i = get_global_id(0);
 			if (i < num) {
 				c[i] = a[i] + b[i];
@@ -92,7 +90,7 @@ See [TypeScript declarations](/index.d.ts) for more details.
 	// Do the work
 	cl.enqueueWriteBuffer(queue, bufferA, true, 0, BYTE_SIZE, arrayA);
 	cl.enqueueWriteBuffer(queue, bufferB, true, 0, BYTE_SIZE, arrayB);
-	cl.enqueueNDRangeKernel(queue, kernel, 1, null, [BUFFER_SIZE], null);
+	cl.enqueueNDRangeKernel(queue, kernel, 1, null, [BUFFER_SIZE]);
 	cl.enqueueReadBuffer(queue, bufferC, true, 0, BYTE_SIZE, arrayC);
 	```
 7. See if it worked:
diff --git a/examples/saxpy.ts b/examples/saxpy.ts
@@ -1,12 +1,9 @@
 import cl from '../index.js';
 
 const getEventExecTime = (event: cl.TClEvent): number => {
-	// times are 64-bit values in nanoseconds. They are returned as [hi, lo] a 2-integer array
-	// here we use the lo parts since this example is unlikely to go beyond 2^31 nanseconds per event.
 	const startTime = cl.getEventProfilingInfo(event, cl.PROFILING_COMMAND_START);
 	const endTime = cl.getEventProfilingInfo(event, cl.PROFILING_COMMAND_END);
-
-	return (endTime[1] - startTime[1]) * 1e-6; // report in millisecond (from nanoseconds)
+	return (endTime - startTime) * 1e-6; // nanoseconds to milliseconds
 };
 
 const VECTOR_SIZE = 512 * 1024;
diff --git a/examples/simple.ts b/examples/simple.ts
@@ -46,7 +46,7 @@ cl.setKernelArg(kernel, 3, 'uint', BUFFER_SIZE);
 // Do the work
 cl.enqueueWriteBuffer(queue, bufferA, true, 0, BYTE_SIZE, arrayA);
 cl.enqueueWriteBuffer(queue, bufferB, true, 0, BYTE_SIZE, arrayB);
-cl.enqueueNDRangeKernel(queue, kernel, 1, null, [BUFFER_SIZE], null);
+cl.enqueueNDRangeKernel(queue, kernel, 1, null, [BUFFER_SIZE]);
 cl.enqueueReadBuffer(queue, bufferC, true, 0, BYTE_SIZE, arrayC);
 
 // print results
diff --git a/examples/square-events.ts b/examples/square-events.ts
@@ -27,9 +27,9 @@ cl.setKernelArg(kern, 0, 'uint*', inputsMem);
 cl.setKernelArg(kern, 1, 'uint*', outputsMem);
 cl.setKernelArg(kern, 2, 'uint', NVALUES);
 
-const cq = cl.createCommandQueue(context, device, null);
+const cq = cl.createCommandQueue(context, device);
 
-cl.enqueueNDRangeKernel(cq, kern, 1, null, [NVALUES], null);
+cl.enqueueNDRangeKernel(cq, kern, 1, null, [NVALUES]);
 
 // here we use the returned user event to associate a callback that will be called from OpenCL
 // once read buffer is complete.
diff --git a/examples/square.ts b/examples/square.ts
@@ -36,7 +36,7 @@ cl.setKernelArg(kern2, 0, 'uint*', inputsMem);
 cl.setKernelArg(kern2, 1, 'uint*', outputsMem);
 cl.setKernelArg(kern2, 2, 'uint', NVALUES);
 
-cl.enqueueNDRangeKernel(cq, kern2, 1, null, [NVALUES], null);
+cl.enqueueNDRangeKernel(cq, kern2, 1, null, [NVALUES]);
 cl.enqueueReadBuffer(cq, outputsMem, true, 0, NVALUES * BYTES_PER_ELEMENT, outputs);
 cl.finish(cq);
 
diff --git a/examples/vector-add.ts b/examples/vector-add.ts
@@ -58,9 +58,7 @@ cl.enqueueWriteBuffer (queue, aBuffer, true, 0, A.length * Uint32Array.BYTES_PER
 cl.enqueueWriteBuffer (queue, bBuffer, true, 0, B.length * Uint32Array.BYTES_PER_ELEMENT, B);
 
 // Execute (enqueue) kernel
-cl.enqueueNDRangeKernel(
-	queue, kernel, 1, null, [BUFFER_SIZE], null,
-);
+cl.enqueueNDRangeKernel(queue, kernel, 1, null, [BUFFER_SIZE]);
 
 // get results and block while getting them
 const C = new Uint32Array(BUFFER_SIZE);
diff --git a/index.d.ts b/index.d.ts
@@ -68,6 +68,9 @@ declare namespace cl {
 	const INVALID_COMPILER_OPTIONS: Error;
 	const INVALID_LINKER_OPTIONS: Error;
 	const INVALID_DEVICE_PARTITION_COUNT: Error;
+	// Additional Error Codes
+	const PLATFORM_NOT_FOUND_KHR: Error;
+	const INVALID_GL_SHAREGROUP_REFERENCE_KHR: Error;
 	
 	// OpenCL Version
 	const VERSION_1_0: number;
@@ -457,13 +460,13 @@ declare namespace cl {
 	const WGL_HDC_KHR: number;
 	const GLX_DISPLAY_KHR: number;
 	const CGL_SHAREGROUP_KHR: number;
+	const CURRENT_DEVICE_FOR_GL_CONTEXT_KHR: number;
+	const DEVICES_FOR_GL_CONTEXT_KHR: number;
 	
 	// ---- Not available on Apple:
 	// cl_khr_icd extension:
 	// cl_platform_info
 	const PLATFORM_ICD_SUFFIX_KHR: number;
-	// Additional Error Codes
-	const PLATFORM_NOT_FOUND_KHR: number;
 	
 	// cl_nv_device_attribute_query extension:
 	// no extension exports since it has no functions
@@ -567,7 +570,7 @@ declare namespace cl {
 		kernel: TClKernel,
 		device: TClDevice,
 		param_name: number,
-	) => TClStatus;
+	) => (number | number[]);
 	
 	
 	/**
@@ -1068,9 +1071,9 @@ declare namespace cl {
 		queue: TClQueue,
 		kernel: TClKernel,
 		work_dim: number,
-		work_offset: number[] | null,
-		work_global: number[] | null,
-		work_local: number[] | null,
+		work_offset?: number[] | null,
+		work_global?: number[] | null,
+		work_local?: number[] | null,
 		event_wait_list?: TClEvent[] | null,
 		hasEvent?: boolean,
 	) => TClEventOrVoid;
@@ -1300,8 +1303,14 @@ declare namespace cl {
 	const getEventProfilingInfo: (
 		event: TClEvent,
 		param_name: number,
-	) => [number, number];
+	) => number;
 	
+	/**
+	 * A small helper to quickly grab a good CL context.
+	 * 
+	 * Multiple calls to this function will return the same `context`.
+	 * Pass `true` to see the intermediate device list (from where we choose).
+	 */
 	const quickStart: (isLoggingDevices?: boolean) => Readonly<{
 		platform: TClPlatform,
 		device: TClDevice,
diff --git a/src/cpp/bindings.cpp b/src/cpp/bindings.cpp
@@ -206,6 +206,7 @@ Napi::Object initModule(Napi::Env env, Napi::Object exports) {
 	JS_CL_ERROR(INVALID_DEVICE_PARTITION_COUNT);
 	// Additional Error Codes
 	JS_CL_ERROR(PLATFORM_NOT_FOUND_KHR);
+	JS_CL_ERROR(INVALID_GL_SHAREGROUP_REFERENCE_KHR);
 	
 	JS_CL_CONSTANT(SUCCESS);
 	
@@ -597,6 +598,8 @@ Napi::Object initModule(Napi::Env env, Napi::Object exports) {
 	JS_CL_CONSTANT(WGL_HDC_KHR);
 	JS_CL_CONSTANT(GLX_DISPLAY_KHR);
 	JS_CL_CONSTANT(CGL_SHAREGROUP_KHR);
+	JS_CL_CONSTANT(CURRENT_DEVICE_FOR_GL_CONTEXT_KHR);
+	JS_CL_CONSTANT(DEVICES_FOR_GL_CONTEXT_KHR);
 	
 	// cl_platform_info
 	JS_CL_CONSTANT(PLATFORM_ICD_SUFFIX_KHR);
diff --git a/src/cpp/common.cpp b/src/cpp/common.cpp
@@ -140,6 +140,8 @@ const char* getExceptionMessage(const cl_int code) {
 			return "Invalid device partition count";
 		case CL_PLATFORM_NOT_FOUND_KHR:
 			return "Platform not found (ICD)";
+		case CL_INVALID_GL_SHAREGROUP_REFERENCE_KHR:
+			return "Invalid GL sharegroup reference";
 		default:
 			printf("OpenCL Unknown error: %d\n", code);
 			return "Unknown error";
diff --git a/src/cpp/common.hpp b/src/cpp/common.hpp
@@ -17,13 +17,6 @@
 	#define strncasecmp _strnicmp
 #endif
 
-#ifndef CL_CURRENT_DEVICE_FOR_GL_CONTEXT_KHR
-	#define CL_CURRENT_DEVICE_FOR_GL_CONTEXT_KHR 0x2006
-#endif
-#ifndef CL_DEVICES_FOR_GL_CONTEXT_KHR
-	#define CL_DEVICES_FOR_GL_CONTEXT_KHR 0x2007
-#endif
-
 #define CHECK_ERR(code) {                                        \
 	cl_int _err = (code);                                        \
 	if (_err != CL_SUCCESS) {                                    \
@@ -41,6 +34,15 @@ namespace opencl {
 void getPtrAndLen(Napi::Object obj, void** ptr, size_t *len);
 const char* getExceptionMessage(cl_int code);
 
+inline Napi::Number NewInt64(napi_env env, int64_t val) {
+	napi_value value;
+	napi_status status = napi_create_int64(env, val, &value);
+	NAPI_THROW_IF_FAILED(env, status, Napi::Number());
+	return Napi::Number(env, value);
+}
+
+#define RET_X64(VAL) return NewInt64(env, static_cast<int64_t>(VAL))
+
 JS_METHOD(createKernel);
 JS_METHOD(createKernelsInProgram);
 JS_METHOD(retainKernel);
diff --git a/src/cpp/device.cpp b/src/cpp/device.cpp
@@ -197,22 +197,7 @@ inline Napi::Value getDeviceInfoUlong(Napi::Env env, cl_device_id device_id, uin
 		&param_value,
 		nullptr
 	));
-	
-	/**
-		JS Compatibility
-		
-		As JS does not support 64 bits integer, we return a 2-integer array with
-			output_values[0] = (input_value >> 32) & 0xffffffff;
-			output_values[1] = input_value & 0xffffffff;
-		
-		and reconstruction as
-			input_value = ((int64_t) output_values[0]) << 32) | output_values[1];
-	*/
-	// TODO: maybe bollox, try modern x64 values
-	Napi::Array arr = Napi::Array::New(env);
-	arr.Set(0u, JS_NUM(param_value>>32)); // hi
-	arr.Set(1u, JS_NUM(param_value & 0xffffffff)); // lo
-	RET_VALUE(arr);
+	RET_X64(param_value);
 }
 
 inline Napi::Value getDeviceInfoSize(Napi::Env env, cl_device_id device_id, uint32_t param_name) {
diff --git a/src/cpp/event.cpp b/src/cpp/event.cpp
@@ -1,6 +1,7 @@
 #include <uv.h>
 
 #include "types.hpp"
+#include "common.hpp"
 
 
 namespace opencl {
@@ -115,16 +116,6 @@ JS_METHOD(getEventProfilingInfo) { NAPI_ENV;
 		case CL_PROFILING_COMMAND_SUBMIT:
 		case CL_PROFILING_COMMAND_START:
 		case CL_PROFILING_COMMAND_END: {
-			/**
-				JS Compatibility
-				
-				As JS does not support 64 bits integer, we return a 2-integer array with
-					output_values[0] = (input_value >> 32) & 0xffffffff;
-					output_values[1] = input_value & 0xffffffff;
-				
-				and reconstruction as
-					input_value = ((int64_t) output_values[0]) << 32) | output_values[1];
-			*/
 			cl_ulong val;
 			CHECK_ERR(clGetEventProfilingInfo(
 				ev,
@@ -134,10 +125,7 @@ JS_METHOD(getEventProfilingInfo) { NAPI_ENV;
 				nullptr
 			));
 			
-			Napi::Array arr = Napi::Array::New(env);
-			arr.Set(0u, JS_NUM(val>>32)); // hi
-			arr.Set(1u, JS_NUM(val & 0xffffffff)); // lo
-			RET_VALUE(arr);
+			RET_X64(val);
 		}
 	}
 	
diff --git a/src/cpp/kernel.cpp b/src/cpp/kernel.cpp
@@ -478,20 +478,7 @@ JS_METHOD(getKernelWorkGroupInfo) { NAPI_ENV;
 				&sz,
 				nullptr
 			));
-			/**
-				JS Compatibility
-
-				As JS does not support 64 bits integer, we return a 2-integer array with
-					output_values[0] = (input_value >> 32) & 0xffffffff;
-					output_values[1] = input_value & 0xffffffff;
-
-				and reconstruction as
-					input_value = ((int64_t) output_values[0]) << 32) | output_values[1];
-			*/
-			Napi::Array arr = Napi::Array::New(env);
-			arr.Set(0u, sz>>32); // hi
-			arr.Set(1u, sz & 0xffffffff); // lo
-			RET_VALUE(arr);
+			RET_X64(sz);
 		}
 	}
 	
diff --git a/test/device.test.ts b/test/device.test.ts
@@ -50,16 +50,6 @@ describe('Device', () => {
 		});
 	};
 	
-	const test64Array = (name: keyof typeof cl) => {
-		it(name + ' returns a 2 integers array', (t, done) => {
-			const val = cl.getDeviceInfo(device, cl[name] as number) as [number, number];
-			U.assertType(val, 'array');
-			U.assertType(val[0], 'number');
-			U.assertType(val[1], 'number');
-			done();
-		});
-	};
-	
 	describe('#getDeviceInfo()', () => {
 		testString('DEVICE_NAME');
 		testString('DEVICE_VENDOR');
@@ -111,11 +101,11 @@ describe('Device', () => {
 		testInteger('DEVICE_REFERENCE_COUNT');
 		testInteger('DEVICE_PARTITION_MAX_SUB_DEVICES');
 		
-		test64Array('DEVICE_GLOBAL_MEM_CACHE_SIZE');
-		test64Array('DEVICE_GLOBAL_MEM_SIZE');
-		test64Array('DEVICE_LOCAL_MEM_SIZE');
-		test64Array('DEVICE_MAX_CONSTANT_BUFFER_SIZE');
-		test64Array('DEVICE_MAX_MEM_ALLOC_SIZE');
+		testInteger('DEVICE_GLOBAL_MEM_CACHE_SIZE');
+		testInteger('DEVICE_GLOBAL_MEM_SIZE');
+		testInteger('DEVICE_LOCAL_MEM_SIZE');
+		testInteger('DEVICE_MAX_CONSTANT_BUFFER_SIZE');
+		testInteger('DEVICE_MAX_MEM_ALLOC_SIZE');
 		testInteger('DEVICE_IMAGE2D_MAX_HEIGHT');
 		testInteger('DEVICE_IMAGE2D_MAX_WIDTH');
 		testInteger('DEVICE_IMAGE3D_MAX_DEPTH');
diff --git a/test/exports.test.ts b/test/exports.test.ts
@@ -24,6 +24,7 @@ const constants: ReadonlyArray<keyof typeof cl> = [
 	'INVALID_MIP_LEVEL', 'INVALID_GLOBAL_WORK_SIZE', 'INVALID_PROPERTY',
 	'INVALID_IMAGE_DESCRIPTOR', 'INVALID_COMPILER_OPTIONS', 'INVALID_LINKER_OPTIONS',
 	'INVALID_DEVICE_PARTITION_COUNT', 'PLATFORM_NOT_FOUND_KHR',
+	'INVALID_GL_SHAREGROUP_REFERENCE_KHR',
 	'SUCCESS', 'VERSION_1_0', 'VERSION_1_1', 'VERSION_1_2',
 	'FALSE', 'TRUE', 'BLOCKING', 'NON_BLOCKING',
 	'PLATFORM_PROFILE', 'PLATFORM_VERSION', 'PLATFORM_NAME', 'PLATFORM_VENDOR',
@@ -126,6 +127,7 @@ const constants: ReadonlyArray<keyof typeof cl> = [
 	'DEVICE_COMPUTE_CAPABILITY_MAJOR_NV', 'DEVICE_COMPUTE_CAPABILITY_MINOR_NV',
 	'DEVICE_REGISTERS_PER_BLOCK_NV', 'DEVICE_WARP_SIZE_NV', 'DEVICE_GPU_OVERLAP_NV',
 	'DEVICE_KERNEL_EXEC_TIMEOUT_NV', 'DEVICE_INTEGRATED_MEMORY_NV',
+	'CURRENT_DEVICE_FOR_GL_CONTEXT_KHR', 'DEVICES_FOR_GL_CONTEXT_KHR',
 ];
 
 const methods: ReadonlyArray<keyof typeof cl> = [
diff --git a/test/kernel.test.ts b/test/kernel.test.ts
@@ -323,8 +323,8 @@ describe('Kernel', () => {
 		testForType('KERNEL_COMPILE_WORK_GROUP_SIZE', (v) => U.assertType(v, 'array'));
 		testForType('KERNEL_PREFERRED_WORK_GROUP_SIZE_MULTIPLE', (v) => U.assertType(v, 'number'));
 		testForType('KERNEL_WORK_GROUP_SIZE', (v) => U.assertType(v, 'number'));
-		testForType('KERNEL_LOCAL_MEM_SIZE', (v) => U.assertType(v, 'array'));
-		testForType('KERNEL_PRIVATE_MEM_SIZE', (v) => U.assertType(v, 'array'));
+		testForType('KERNEL_LOCAL_MEM_SIZE', (v) => U.assertType(v, 'number'));
+		testForType('KERNEL_PRIVATE_MEM_SIZE', (v) => U.assertType(v, 'number'));
 		
 		it('throws INVALID_VALUE when looking for KERNEL_GLOBAL_WORK_SIZE', () => {
 			U.withProgram(context, squareKern, (prg) => {
diff --git a/test/queues.test.ts b/test/queues.test.ts
@@ -125,7 +125,7 @@ describe('CommandQueue - Common', () => {
 				cl.setKernelArg(kern, 1, 'uint*', outputsMem);
 				cl.setKernelArg(kern, 2, 'uint', 10000);
 				
-				cl.enqueueNDRangeKernel(cq, kern, 1, null, [100], null);
+				cl.enqueueNDRangeKernel(cq, kern, 1, null, [100]);
 				
 				cl.releaseMemObject(inputsMem);
 				cl.releaseMemObject(outputsMem);
@@ -142,7 +142,6 @@ describe('CommandQueue - Common', () => {
 						1,
 						null,
 						[100],
-						null,
 					),
 					new Error('Argument 1 must be of type `Object`'),
 				);
@@ -171,7 +170,6 @@ describe('CommandQueue - Common', () => {
 						1,
 						null,
 						[100, 200],
-						null,
 					),
 					cl.INVALID_GLOBAL_WORK_SIZE,
 				);
