errorCode = dispatchQueue->enqueueWriteBuffer(
        inputBuffer,
        CL_FALSE,//CL_TRUE,
        0,
        sizeof(cl_uchar) * inputBufferSize,
        passwordBuffer,
        NULL,
        &dispatchDelegate);

errorCode = dispatchQueue->enqueueNDRangeKernel(
    *kernel,
    NullRange,
    NDRange(globalWorkgroupSize, 1),
    NullRange,
    NULL,
    NULL);

__kernel void sha1Crack(__global uchar* out, __global uchar* in,
                        __constant int* passLen, __constant int* targetHash,
                        __global bool* collisionFound)
{
    //Kernel Instance Global GPU Mem IO Mapping:
    __private int id = get_global_id(0);
    __private int inputIndexStart = id * passwordLen;

    //Select Password input key space:
    #pragma unroll
    for (i = 0; i < passwordLen; i++)
    {
        inputMem[i] = in[inputIndexStart + i];
    }

    //SHA1 Code omitted for brevity...
}