Created
October 3, 2011 21:57
-
-
Save chadbrewbaker/1260367 to your computer and use it in GitHub Desktop.
Simple WebCL benchmark on an array using the JavaScript timer
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
<!DOCTYPE html PUBLIC "-//W3C//DTD XHTML 1.0 Strict//EN" "http://www.w3.org/TR/xhtml1/DTD/xhtml1-strict.dtd"> | |
<html xmlns="http://www.w3.org/1999/xhtml"> | |
<head> | |
<meta http-equiv="content-type" content="text/html; charset=utf-8" /> | |
<title>Learning WebCL, lesson 3</title> | |
<meta name="keywords" content="webcl, webgl, gpu, opencl, opengl" /> | |
<meta name="description" content="WebCL extension for Firefox, providing direct GPU access from JavaScript" /> | |
<link href="/default.css" rel="stylesheet" type="text/css" /> | |
</head> | |
<body id="tutorial"> | |
<div id="header"> | |
<div id="logo"> | |
<h1>WebCL</h1> | |
</div> <!-- end div#logo --> | |
</div> <!-- end div#header --> | |
<div id="page"> | |
<div id="page-bgtop"> | |
<div id="content"> | |
<div class="post box"> | |
<h2 class="title">Vector Add Benchmark</h2> | |
<div class="entry"> | |
<script id="clProgramVectorAdd" type="text/x-opencl"> | |
__kernel void ckVectorAdd(__global unsigned int* vectorIn1, __global unsigned int* vectorIn2, | |
__global unsigned int* vectorOut, unsigned int uiVectorWidth) | |
{ | |
unsigned int x = get_global_id(0); | |
if (x >= (uiVectorWidth)) | |
{ | |
return; | |
} | |
// add the vector elements | |
vectorOut[x] = vectorIn1[x] + vectorIn2[x]; | |
} | |
</script> | |
<script type="text/javascript"> | |
function loadKernel(id){ | |
var kernelElement = document.getElementById(id); | |
var kernelSource = kernelElement.text; | |
if (kernelElement.src != "") { | |
var mHttpReq = new XMLHttpRequest(); | |
mHttpReq.open("GET", kernelElement.src, false); | |
mHttpReq.send(null); | |
kernelSource = mHttpReq.responseText; | |
} | |
return kernelSource; | |
} | |
</script> | |
<script type="text/javascript"> | |
function CL_vectorAdd () { | |
var timer = { | |
startTimes : [], | |
start : function(id) { | |
this.startTimes[id] = new Date().getTime(); | |
}, | |
elapsed : function(id, log) { | |
var elapsed = new Date().getTime() - this.startTimes[id]; | |
if (log==true) { console.log(id, "took", elapsed, "ms") }; | |
return elapsed; | |
} | |
}; | |
timer.start("everything"); | |
// All output is written to element by id "output" | |
var output = document.getElementById("output"); | |
output.innerHTML = ""; | |
try { | |
// First check if the WebCL extension is installed at all | |
if (window.WebCL == undefined) { | |
alert("Unfortunately your system does not support WebCL. " + | |
"Make sure that you have both the OpenCL driver " + | |
"and the WebCL browser extension installed."); | |
return false; | |
} | |
// Generate input vectors | |
var vectorLength = 30; | |
var UIvector1 = new Uint32Array(vectorLength); | |
var UIvector2 = new Uint32Array(vectorLength); | |
for ( var i=0; i<vectorLength; i=i+1) { | |
UIvector1[i] = Math.floor(Math.random() * 100); //Random number 0..99 | |
UIvector2[i] = Math.floor(Math.random() * 100); //Random number 0..99 | |
} | |
output.innerHTML += "<br>Vector length = " + vectorLength; | |
// Setup WebCL context using the default device of the first available platform | |
var platforms = WebCL.getPlatformIDs(); | |
var ctx = WebCL.createContextFromType ([WebCL.CL_CONTEXT_PLATFORM, platforms[0]], | |
WebCL.CL_DEVICE_TYPE_DEFAULT); | |
// Reserve buffers | |
var bufSize = vectorLength * 4; // size in bytes | |
output.innerHTML += "<br>Buffer size: " + bufSize + " bytes"; | |
var bufIn1 = ctx.createBuffer (WebCL.CL_MEM_READ_ONLY, bufSize); | |
var bufIn2 = ctx.createBuffer (WebCL.CL_MEM_READ_ONLY, bufSize); | |
var bufOut = ctx.createBuffer (WebCL.CL_MEM_WRITE_ONLY, bufSize); | |
// Create and build program for the first device | |
var kernelSrc = loadKernel("clProgramVectorAdd"); | |
var program = ctx.createProgramWithSource(kernelSrc); | |
var devices = ctx.getContextInfo(WebCL.CL_CONTEXT_DEVICES); | |
try { | |
program.buildProgram ([devices[0]], ""); | |
} catch(e) { | |
alert ("Failed to build WebCL program. Error " | |
+ program.getProgramBuildInfo (devices[0], WebCL.CL_PROGRAM_BUILD_STATUS) | |
+ ": " + program.getProgramBuildInfo (devices[0], WebCL.CL_PROGRAM_BUILD_LOG)); | |
throw e; | |
} | |
// Create kernel and set arguments | |
var kernel = program.createKernel ("ckVectorAdd"); | |
kernel.setKernelArg (0, bufIn1); | |
kernel.setKernelArg (1, bufIn2); | |
kernel.setKernelArg (2, bufOut); | |
kernel.setKernelArg (3, vectorLength, WebCL.types.UINT); | |
// Create command queue using the first available device | |
var cmdQueue = ctx.createCommandQueue (devices[0], 0); | |
// Write the buffer to OpenCL device memory | |
var dataObject1 = WebCL.createDataObject (); | |
dataObject1.allocate(bufSize); | |
dataObject1.set (UIvector1); | |
var dataObject2 = WebCL.createDataObject (); | |
dataObject2.allocate(bufSize); | |
dataObject2.set (UIvector2); | |
cmdQueue.enqueueWriteBuffer (bufIn1, false, 0, dataObject1.length, dataObject1, []); | |
cmdQueue.enqueueWriteBuffer (bufIn2, false, 0, dataObject2.length, dataObject2, []); | |
// Init ND-range | |
var localWS = [8]; | |
var globalWS = [Math.ceil (vectorLength / localWS) * localWS]; | |
output.innerHTML += "<br>Global work item size: " + globalWS; | |
output.innerHTML += "<br>Local work item size: " + localWS; | |
// Execute (enqueue) kernel | |
cmdQueue.enqueueNDRangeKernel(kernel, globalWS.length, [], globalWS, localWS, []); | |
// Read the result buffer from OpenCL device | |
cmdQueue.enqueueReadBuffer (bufOut, false, 0, bufSize, dataObject1, []); | |
cmdQueue.finish (); //Finish all the operations | |
outBuffer = new Uint32Array(vectorLength); | |
var utils = WebCL.getUtils (); | |
utils.writeDataObjectToTypedArray (dataObject1, outBuffer); | |
//Print input vectors and result vector | |
output.innerHTML += "<br>Vector1 = "; | |
for (var i = 0; i < vectorLength; i = i + 1) { | |
output.innerHTML += UIvector1[i] + ", "; | |
} | |
output.innerHTML += "<br>Vector2 = "; | |
for (var i = 0; i < vectorLength; i = i + 1) { | |
output.innerHTML += UIvector2[i] + ", "; | |
} | |
output.innerHTML += "<br>Result = "; | |
for (var i = 0; i < vectorLength; i = i + 1) { | |
output.innerHTML += outBuffer[i] + ", "; | |
} | |
var elapsed = timer.elapsed("everything"); | |
//alert(elapsed+" ticks"); | |
output.innerHTML += "<br> ticks:"+elapsed; | |
} catch(e) { | |
document.getElementById("output").innerHTML += "<h3>ERROR:</h3><pre style=\"color:red;\">" + e.message + "</pre>"; | |
throw e; | |
} | |
} | |
</script> | |
<div><input type="button" value="Run Vector Add Benchmark" onclick="CL_vectorAdd()"></div> | |
<div id="output"></div> | |
</div> | |
</div> | |
</div> | |
<!-- end div#content --> | |
<div id="sidebar"> | |
<div class="box"> | |
<ul> | |
<li> | |
</li> | |
</ul> | |
</div> | |
</div> | |
<!-- end div#sidebar --> | |
<div style="clear: both; height: 1px"></div> | |
</div> | |
</div> <!-- end div#page --> | |
<div id="footer-wrapper"> | |
<div id="footer"> | |
</div> | |
</div> <!-- end div#footer-wrapper --> | |
<!--<script type="text/javascript"> | |
// var _gaq = _gaq || []; | |
_gaq.push(['_setAccount', 'UA-22607782-1']); | |
_gaq.push(['_trackPageview']); | |
(function() { | |
var ga = document.createElement('script'); ga.type = 'text/javascript'; ga.async = true; | |
ga.src = ('https:' == document.location.protocol ? 'https://ssl' : 'http://www') + '.google-analytics.com/ga.js'; | |
var s = document.getElementsByTagName('script')[0]; s.parentNode.insertBefore(ga, s); | |
})(); | |
</script> | |
--> | |
</body> | |
</html> |
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment