diff --git a/package.json b/package.json index 80212f96..19067167 100644 --- a/package.json +++ b/package.json @@ -9,7 +9,7 @@ }, "license": "MIT", "devDependencies": { - "grunt": "0.4.5", + "grunt": "~0.4.5", "grunt-browserify": "^3.2.0", "grunt-closure-compiler": "0.0.21", "grunt-contrib-clean": "0.6.0", diff --git a/src/utils/webcl.js b/src/utils/webcl.js index 0e674ec1..85f4e6dc 100644 --- a/src/utils/webcl.js +++ b/src/utils/webcl.js @@ -169,7 +169,7 @@ */ function hasWebCLNamespace() { - WebCLNamespaceAvailable = window.WebCL && WebCL.getPlatforms; + WebCLNamespaceAvailable = window.webcl && webcl.getPlatforms; return WebCLNamespaceAvailable; @@ -187,7 +187,7 @@ OpenCLDriversAvailable = true; try { - platArr = WebCL.getPlatforms(); + platArr = webcl.getPlatforms(); } catch (e) { OpenCLDriversAvailable = false; } @@ -239,7 +239,7 @@ devices = getDevicesByType(type || DEFAULT_DEVICE); // Creating default context - ctx = createContext({devices: devices}); + ctx = createContext(devices); return true; } @@ -253,7 +253,7 @@ function getPlatforms() { if(platforms.length === 0) { - platforms = WebCL.getPlatforms(); + platforms = webcl.getPlatforms(); } return platforms; @@ -279,11 +279,11 @@ try { if (type === "CPU") { - deviceArr = platform.getDevices(WebCL.DEVICE_TYPE_CPU); + deviceArr = platform.getDevices(webcl.DEVICE_TYPE_CPU); } else if (type === "GPU") { - deviceArr = platform.getDevices(WebCL.DEVICE_TYPE_GPU); + deviceArr = platform.getDevices(webcl.DEVICE_TYPE_GPU); } else if (type === "ALL") { - deviceArr = platform.getDevices(WebCL.DEVICE_TYPE_ALL); + deviceArr = platform.getDevices(webcl.DEVICE_TYPE_ALL); } } catch (e) { @@ -344,7 +344,7 @@ } try { - platform = device.getInfo(WebCL.DEVICE_PLATFORM); + platform = device.getInfo(webcl.DEVICE_PLATFORM); } catch (e) { var errCode = getErrorCodeFromCLError(e); @@ -373,7 +373,7 @@ } try { - deviceArr = clCtx.getInfo(WebCL.CONTEXT_DEVICES); + deviceArr = clCtx.getInfo(webcl.CONTEXT_DEVICES); }catch(e) { errCode = getErrorCodeFromCLError(e); @@ -397,14 +397,15 @@ */ function createContext(properties) { + /* var props = { devices: getDevicesByType(DEFAULT_DEVICE) }, context; - XML3D.extend(props, properties); + XML3D.extend(props, properties);*/ try { - context = WebCL.createContext(props); + var context = webcl.createContext(properties); } catch (e) { var errCode = getErrorCodeFromCLError(e); @@ -494,7 +495,7 @@ throw e; } throw new WebCLError(getCLErrorName(errCode), - program.getBuildInfo(deviceArr[0], WebCL.CL_PROGRAM_BUILD_LOG)); + program.getBuildInfo(deviceArr[0], WebCL.PROGRAM_BUILD_LOG)); } return program; @@ -602,11 +603,11 @@ try { if (type === "r") { - return clCtx.createBuffer(WebCL.CL_MEM_READ_ONLY, size); + return clCtx.createBuffer(webcl.MEM_READ_ONLY, size); } else if (type === "w") { - return clCtx.createBuffer(WebCL.CL_MEM_WRITE_ONLY, size); + return clCtx.createBuffer(webcl.MEM_WRITE_ONLY, size); } else if (type === "rw") { - return clCtx.createBuffer(WebCL.CL_MEM_READ_WRITE, size); + return clCtx.createBuffer(webcl.MEM_READ_WRITE, size); } else { XML3D.debug.logError("WebCL API: createBuffer(): Unknown buffer type:", type); return false; @@ -752,7 +753,7 @@ return false; } - nKernelArgs = kernel.getInfo(WebCL.CL_KERNEL_NUM_ARGS); + nKernelArgs = kernel.getInfo(webcl.KERNEL_NUM_ARGS); if (inputArgs.length > nKernelArgs) { XML3D.debug.logWarning("WebCL: setArgs: Input args amount > kernel program args amount! Ignoring extra arguments."); @@ -761,7 +762,7 @@ return false; } - XML3D.debug.logDebug("Args for kernel:", kernel.getInfo(WebCL.KERNEL_FUNCTION_NAME)); + XML3D.debug.logDebug("Args for kernel:", kernel.getInfo(webcl.KERNEL_FUNCTION_NAME)); i = nKernelArgs; diff --git a/src/xflow/interface/graph.js b/src/xflow/interface/graph.js index 4855906e..bfa26199 100644 --- a/src/xflow/interface/graph.js +++ b/src/xflow/interface/graph.js @@ -61,7 +61,7 @@ function initWebCLPlatform(graph) { // Creating a new WebCL context try { - clCtx = webcl.createContext({devices: clDevices}); + clCtx = webcl.createContext(clDevices); } catch (e) { return false; } diff --git a/src/xflow/operator/cl-program.js b/src/xflow/operator/cl-program.js index c9420a6d..21961088 100644 --- a/src/xflow/operator/cl-program.js +++ b/src/xflow/operator/cl-program.js @@ -103,11 +103,51 @@ this.type = "int"; } break; + case xflowDataTypes.INT4: + { + helperMap = helperParamMap.buffer; + this.type = "int4*"; + addressSpace = "__global"; + this.needsMemObject = true; + } + break; case xflowDataTypes.FLOAT: { this.type = "float"; } break; + case xflowDataTypes.FLOAT2: + { + helperMap = helperParamMap.buffer; + this.type = "float2*"; + addressSpace = "__global"; + this.needsMemObject = true; + } + break; + case xflowDataTypes.FLOAT3: + { + helperMap = helperParamMap.buffer; + this.type = "float*"; + addressSpace = "__global"; + this.needsMemObject = true; + } + break; + case xflowDataTypes.FLOAT4: + { + helperMap = helperParamMap.buffer; + this.type = "float4*"; + addressSpace = "__global"; + this.needsMemObject = true; + } + break; + case xflowDataTypes.FLOAT4X4: + { + helperMap = helperParamMap.buffer; + this.type = "float16*"; + addressSpace = "__global"; + this.needsMemObject = true; + } + break; default: return; } @@ -152,6 +192,32 @@ helperVal = self.entryValue.height; } else if (p === "length") { helperVal = self.entryValue.length; + var type = self.xflowType; + switch (type) { + case Xflow.DATA_TYPE.FLOAT2: + { + helperVal = helperVal / 2; + } + break; + case Xflow.DATA_TYPE.FLOAT3: + { + helperVal = helperVal / 3; + } + break; + case Xflow.DATA_TYPE.FLOAT4: + case Xflow.DATA_TYPE.INT4: + { + helperVal = helperVal / 4; + } + break; + case Xflow.DATA_TYPE.FLOAT4X4: + { + helperVal = helperVal / 16; + } + break; + default: + break; + } } self.helpers.push(new KernelParam(self.program, pName, null, helperMap.type, new Uint32Array([helperVal]))); }); @@ -184,12 +250,30 @@ this.arg.release(); } - this.byteSize = byteSize; + this.byteSize = isNaN(byteSize) ? 1 : byteSize; if (this.xflowType === Xflow.DATA_TYPE.TEXTURE) { // Texture is a special case memObjectSize = entryValue.width * entryValue.height * byteSize; + this.byteSize = 4; } else { - memObjectSize = entryValue.length * byteSize; + switch (this.xflowType) { + case Xflow.DATA_TYPE.INT4: + { + memObjectSize = entryValue.length * Int32Array.BYTES_PER_ELEMENT; + } + break; + case Xflow.DATA_TYPE.FLOAT2: + case Xflow.DATA_TYPE.FLOAT3: + case Xflow.DATA_TYPE.FLOAT4: + case Xflow.DATA_TYPE.FLOAT4x4: + { + memObjectSize = entryValue.length * Float32Array.BYTES_PER_ELEMENT; + } + break; + default: + memObjectSize = entryValue.length * 4; + break; + } } memObject = clAPI.createBuffer(memObjectSize, memObjectMode, clCtx); @@ -204,7 +288,7 @@ updateValue: function (entry) { if (this.hasMemObject) { - this.val = entry.data; + this.val = entry.data === undefined ? entry : entry.data; this.entryValue = entry; this.checkEntrySize(); } else { @@ -222,7 +306,14 @@ } else if (name.indexOf("height") !== -1) { p.updateValue(new Uint32Array([self.entryValue.height])); } else if (name.indexOf("length") !== -1) { - p.updateValue(new Uint32Array([self.entryValue.length])); + var len = self.entryValue.length; + if (self.xflowType === Xflow.DATA_TYPE.FLOAT4 || + self.xflowType === Xflow.DATA_TYPE.INT4) { + len = len / 4; + } else if (self.xflowType === Xflow.DATA_TYPE.FLOAT4X4) { + len = len / 16; + } + p.updateValue(new Uint32Array([len])); } }); }, @@ -232,6 +323,13 @@ if(this.xflowType === Xflow.DATA_TYPE.TEXTURE) { newSize = entryVal.width * entryVal.height * this.byteSize; + } else if (this.xflowType === Xflow.DATA_TYPE.FLOAT2 || + this.xflowType === Xflow.DATA_TYPE.FLOAT3 || + this.xflowType === Xflow.DATA_TYPE.FLOAT4 || + this.xflowType === Xflow.DATA_TYPE.FLOAT4X4) { + newSize = entryVal.length * Float32Array.BYTES_PER_ELEMENT; + } else if (this.xflowType === Xflow.DATA_TYPE.INT4) { + newSize = entryVal.length * Int32Array.BYTES_PER_ELEMENT; } else { newSize = entryVal.length * this.byteSize; } @@ -488,7 +586,7 @@ codeLines.push("int " + firstInput.name + "_i = get_global_id(0);"); - codeLines.push("if (int " + firstInput.name + "_i >= " + firstInput.name + "_length) return;"); + codeLines.push("if (" + firstInput.name + "_i >= " + firstInput.name + "_length) return;"); } return codeLines.join('\n'); @@ -535,17 +633,17 @@ len = inputMemObjs.length; for (i = 0; i < len; i++) { memObj = inputMemObjs[i]; - cmdQueue.enqueueWriteBuffer(memObj.arg, false, 0, memObj.arg.getInfo(WebCL.MEM_SIZE), memObj.val, []); + cmdQueue.enqueueWriteBuffer(memObj.arg, false, 0, /*memObj.arg.getInfo(WebCL.MEM_SIZE)*/memObj.memObjectSize, memObj.val, []); } // Execute (enqueue) kernel - cmdQueue.enqueueNDRangeKernel(kernel, WSSizes[1].length, [], WSSizes[1], WSSizes[0], []); + cmdQueue.enqueueNDRangeKernel(kernel, WSSizes[1].length, [], WSSizes[1], WSSizes[0]); // Read the result buffer from OpenCL device len = outputMemObjs.length; for (i = 0; i < len; i++) { memObj = outputMemObjs[i]; - cmdQueue.enqueueReadBuffer(memObj.arg, false, 0, memObj.arg.getInfo(WebCL.MEM_SIZE), memObj.val, []); + cmdQueue.enqueueReadBuffer(memObj.arg, false, 0, /*memObj.arg.getInfo(WebCL.MEM_SIZE)*/memObj.memObjectSize, memObj.val, []); } cmdQueue.finish(); //Finish all the operations @@ -613,12 +711,31 @@ globalWS = [Math.ceil(entryVal.width / localWS[0]) * localWS[0], Math.ceil(entryVal.height / localWS[1]) * localWS[1]]; } else { + var k = 1; + switch (targetInput.xflowType) { + case Xflow.DATA_TYPE.INT4: + case Xflow.DATA_TYPE.FLOAT4: + { + k = 4; + } + break; + case Xflow.DATA_TYPE.FLOAT3: + { + k = 3; + } + break; + case Xflow.DATA_TYPE.FLOAT4X4: + { + k = 16; + } + break; + } localWS = [16]; - globalWS = [Math.ceil(entryVal.length / localWS[0]) * localWS[0]]; + globalWS = [Math.ceil(entryVal.length / (localWS[0] * k)) * localWS[0]]; } return [localWS, globalWS]; } -}()); \ No newline at end of file +}());