private OpenCLKernel InitKernel(BigInteger elementMontgomery, BigInteger generatorMontgomery, out OpenCLBuffer <int> gpuCounterBuffer, out BigInteger?answer)
{
// Make all inputs GPU ready, by converting them to uint arrays.
// Note: each number (e.g. special points, the modulus etc.) will be represented with
// wordsPerNumber uints (wordsPerNumber * 32 bits).
uint[] gpuModulus = modulus.ToUintArray().PadWithDefaultForLength(wordsPerNumber); // TODO: Remove padding?
uint[] gpuModulusPrime = modulusPrime.ToUintArray().PadWithDefaultForLength(wordsPerNumber);
uint[] gpuElement = elementMontgomery.ToUintArray().PadWithDefaultForLength(wordsPerNumber);
uint[] gpuGenerator = generatorMontgomery.ToUintArray().PadWithDefaultForLength(wordsPerNumber);
// Input buffers.
OpenCLBuffer <uint> gpuModulusBuffer = new OpenCLBuffer <uint>(program, gpuModulus);
OpenCLBuffer <uint> gpuGeneratorBuffer = new OpenCLBuffer <uint>(program, gpuGenerator);
OpenCLBuffer <uint> gpuElementBuffer = new OpenCLBuffer <uint>(program, gpuElement);
// Buffers for local memory. There is room for an additional 2 numbers, which will be used to store
// the generator and element in local memory.
OpenCLBuffer <uint> gpuNumbersBuffer = new OpenCLBuffer <uint>(program, new uint[wordsPerNumber * (2 + 32)]);
// Counter buffer.
gpuCounterBuffer = new OpenCLBuffer <int>(program, new int[1]);
// Buffers for saving numbers between kernel executions.
uint[] startingPointsArray = startingPointGenerator.GetVerticalStartingPointsArray(NUM_GPU_THREADS, out answer);
OpenCLBuffer <uint> gpuSavedNumbersBuffer = new OpenCLBuffer <uint>(program, startingPointsArray);
OpenCLBuffer <uint> gpuUsedStartingPointBuffer = new OpenCLBuffer <uint>(program, startingPointsArray);
OpenCLBuffer <long> gpuIterationCounts = new OpenCLBuffer <long>(program, new long[NUM_GPU_THREADS]);
// Fill the gpuStartingPointBuffer.
answer = startingPointGenerator.FillStartingPointsBuffer(4 * NUM_GPU_THREADS);
OpenCLKernel kernel = new OpenCLKernel(program, "generate_chain");
// Set the kernelarguments.
kernel.SetArgument(0, gpuStartingPointsBuffer);
kernel.SetArgument(1, gpuSavedNumbersBuffer);
kernel.SetArgument(2, gpuUsedStartingPointBuffer);
kernel.SetLocalArgument(3, gpuNumbersBuffer);
kernel.SetArgument(4, gpuModulusBuffer);
kernel.SetArgument <uint>(5, gpuModulusPrime[0]);
kernel.SetArgument(6, gpuGeneratorBuffer);
kernel.SetArgument(7, gpuElementBuffer);
kernel.SetArgument(8, gpuSpecialPointsBuffer);
kernel.SetArgument <int>(9, wordsPerNumber);
kernel.SetArgument(10, gpuCounterBuffer);
kernel.SetArgument(11, gpuIterationCounts);
kernel.SetArgument <long>(12, 1 << (Program.K + 4)); // Maximum chain length is 16 * 2^k.
kernel.SetArgument <int>(13, Program.K / 32); // Value of k, in words.
kernel.SetArgument <int>(14, Program.K % 32); // Remaining value of k.
return(kernel);
}
public StartingPointGenerator(InputTuple input, int rAsPower, int wordsPerNumber, Pollard_Rho pRho, OpenCLProgram program, OpenCLBuffer <uint> startingPointsBuffer)
{
this.modulus = input.Modulus;
this.generator = input.Generator;
this.order = input.Order;
this.element = input.Element;
this.rAsPower = rAsPower;
this.wordsPerNumber = wordsPerNumber;
this.pRho = pRho;
this.kernel = new OpenCLKernel(program, "add_new_starting_points");
this.newStartingPointsBuffer = new OpenCLBuffer <uint>(program, new uint[4 * DLPSolver.NUM_GPU_THREADS * wordsPerNumber]);
this.startingPointPool = new uint[4 * DLPSolver.NUM_GPU_THREADS * wordsPerNumber];
kernel.SetArgument(0, startingPointsBuffer);
kernel.SetArgument(1, newStartingPointsBuffer);
}
/// <summary>
/// Enqueues a command to execute a single <see cref="OpenCLKernel"/>.
/// </summary>
/// <param name="kernel"> The <see cref="OpenCLKernel"/> to execute. </param>
/// <param name="events"> A collection of events that need to complete before this particular command can be executed. If <paramref name="events"/> is not <c>null</c> or read-only a new <see cref="OpenCLEvent"/> identifying this command is created and attached to the end of the collection. </param>
public void ExecuteTask(OpenCLKernel kernel, IReadOnlyList<OpenCLEventBase> events = null, IList<OpenCLEventBase> newEvents = null)
{
int eventWaitListSize;
CLEventHandle[] eventHandles = OpenCLTools.ExtractHandles(events, out eventWaitListSize);
CLEventHandle[] newEventHandle = (newEvents != null) ? new CLEventHandle[1] : null;
OpenCLErrorCode error = CL10.EnqueueTask(Handle, kernel.Handle, eventWaitListSize, eventHandles, newEventHandle);
OpenCLException.ThrowOnError(error);
if (newEvents != null)
{
lock (newEvents)
{
newEvents.Add(new OpenCLEvent(newEventHandle[0], this));
}
}
}
public static Error clSetKernelArg(OpenCLKernel kernel, Int32 arg_index, IntPtr arg_size, ref OpenCLMem arg_value)
{
Console.WriteLine("Calling Error clSetKernelArg(OpenCLKernel kernel, Int32 arg_index, IntPtr arg_size, ref OpenCLMem arg_value)");
return default(Error);
}
public OpenCLTask(OpenCLKernel kernel, AutoResetEvent taskCompletedEvent)
{
Kernel = kernel;
TaskCompletedEvent = taskCompletedEvent;
}
public static Error clGetKernelWorkGroupInfo(OpenCLKernel kernel, OpenCLDevice device, KernelWorkGroupInfo param_name, IntPtr param_value_size, IntPtr param_value, IntPtr param_value_size_ret)
{
Console.WriteLine("Calling Error clGetKernelWorkGroupInfo(OpenCLKernel kernel, OpenCLDevice device, KernelWorkGroupInfo param_name, IntPtr param_value_size, IntPtr param_value, IntPtr param_value_size_ret)");
return default(Error);
}
public static Error clRetainKernel(OpenCLKernel kernel)
{
Console.WriteLine("Calling Error clRetainKernel(OpenCLKernel kernel)");
return default(Error);
}
public static extern Error clEnqueueTask(OpenCLCommandQueue command_queue, OpenCLKernel kernel, Int32 num_events_in_wait_list, [In] OpenCLEvent[] event_wait_list, OpenCLEvent e);
public static Error clGetKernelInfo(OpenCLKernel kernel, KernelInfo param_name, IntPtr param_value_size, IntPtr param_value, out IntPtr param_value_size_ret)
{
if (param_name == KernelInfo.NumArgs)
{
Marshal.WriteIntPtr(param_value, new IntPtr(3));
param_value_size_ret = new IntPtr(IntPtr.Size);
}
else if (param_name == KernelInfo.FunctionName)
{
String name = "VectorAdd";
for (int i = 0; i < name.Length; i++)
{
Marshal.WriteByte(param_value, i, (byte)name[i]);
}
param_value_size_ret = new IntPtr(name.Length);
}
else
{
param_value_size_ret = default(IntPtr);
}
Console.WriteLine("Calling Error clGetKernelInfo(OpenCLKernel kernel, KernelInfo param_name, IntPtr param_value_size, IntPtr param_value, out IntPtr param_value_size_ret)");
return default(Error);
}
public static Error clCreateKernelsInProgram(OpenCLProgram program, Int32 num_kernels, [Out] OpenCLKernel[] kernels, out Int32 num_kernels_ret)
{
num_kernels_ret = 1;
kernels = new OpenCLKernel[] { new OpenCLKernel() };
Console.WriteLine("Calling Error clCreateKernelsInProgram(OpenCLProgram program, Int32 num_kernels, [Out] OpenCLKernel[] kernels, out Int32 num_kernels_ret)");
return default(Error);
}
/// <summary>
/// Tries to solve the current instance of the DLP.
/// </summary>
public BigInteger?Solve()
{
// Fill the startingPointPool with starting points.
BigInteger?answer = startingPointGenerator.FillStartingPointPool(4 * NUM_GPU_THREADS);
if (answer != null)
{
return(answer);
}
modulusPrime = -Utility.ExtendedEuclidesGCD1(BigInteger.One << 32, modulus).Item2;
if (modulusPrime < 0)
{
modulusPrime += (BigInteger.One << 32);
}
BigInteger generatorMontgomery = ToMontgomery(generator);
BigInteger elementMontgomery = ToMontgomery(element);
// Initialize the kernel and buffers.
OpenCLBuffer <int> countersBuffer;
OpenCLKernel kernel = InitKernel(elementMontgomery, generatorMontgomery, out countersBuffer, out answer);
if (answer != null)
{
return(answer);
}
// Let the GPU do its job.
int counter = 0;
while (true)
{
Console.WriteLine($"Starting kernel {counter}.");
kernel.Execute(new long[] { NUM_GPU_THREADS }, new long[] { 32 });
kernel.WaitTillQueueFinish();
Console.WriteLine($"Kernel {counter} finished.");
countersBuffer.CopyFromDevice();
Console.WriteLine($"Found special points: {countersBuffer[0]}.");
// The number of used starting points has exceeded the threshold.
if (countersBuffer[0] > NUM_GPU_THREADS)
{
Console.WriteLine("Generating new starting points.");
answer = startingPointGenerator.FillStartingPointsBuffer(countersBuffer[0]);
if (answer != null)
{
return(answer);
}
Console.WriteLine("Retrieving special points.");
specialPointCollector.CollectSpecialPoints(countersBuffer[0]);
answer = specialPointCollector.FindCollision();
if (answer != null)
{
return(answer);
}
// Reset the counter.
countersBuffer[0] = 0;
countersBuffer.CopyToDevice();
}
counter++;
}
}
public static extern Error clRetainKernel(OpenCLKernel kernel);
public static extern Error clSetKernelArg(OpenCLKernel kernel, Int32 arg_index, IntPtr arg_size, ref OpenCLMem arg_value);
public static extern Error clReleaseKernel(OpenCLKernel kernel);
public static extern Error clGetKernelWorkGroupInfo(OpenCLKernel kernel, OpenCLDevice device, KernelWorkGroupInfo param_name, IntPtr param_value_size, IntPtr param_value, IntPtr param_value_size_ret);
public static extern Error clGetKernelInfo(OpenCLKernel kernel, KernelInfo param_name, IntPtr param_value_size, IntPtr param_value, IntPtr param_value_size_ret);
/// <summary>
/// Enqueues a command to execute a range of <see cref="OpenCLKernel"/>s in parallel.
/// </summary>
/// <param name="kernel"> The <see cref="OpenCLKernel"/> to execute. </param>
/// <param name="globalWorkOffset"> An array of values that describe the offset used to calculate the global ID of a work-item instead of having the global IDs always start at offset (0, 0,... 0). </param>
/// <param name="globalWorkSize"> An array of values that describe the number of global work-items in dimensions that will execute the kernel function. The total number of global work-items is computed as global_work_size[0] *...* global_work_size[work_dim - 1]. </param>
/// <param name="localWorkSize"> An array of values that describe the number of work-items that make up a work-group (also referred to as the size of the work-group) that will execute the <paramref name="kernel"/>. The total number of work-items in a work-group is computed as local_work_size[0] *... * local_work_size[work_dim - 1]. </param>
/// <param name="events"> A collection of events that need to complete before this particular command can be executed. If <paramref name="events"/> is not <c>null</c> or read-only a new <see cref="OpenCLEvent"/> identifying this command is created and attached to the end of the collection. </param>
public void Execute(OpenCLKernel kernel, long[] globalWorkOffset, long[] globalWorkSize, long[] localWorkSize, IReadOnlyList<OpenCLEventBase> events = null, IList<OpenCLEventBase> newEvents = null)
{
int eventWaitListSize;
CLEventHandle[] eventHandles = OpenCLTools.ExtractHandles(events, out eventWaitListSize);
CLEventHandle[] newEventHandle = (newEvents != null) ? new CLEventHandle[1] : null;
OpenCLErrorCode error = CL10.EnqueueNDRangeKernel(Handle, kernel.Handle, globalWorkSize.Length, OpenCLTools.ConvertArray(globalWorkOffset), OpenCLTools.ConvertArray(globalWorkSize), OpenCLTools.ConvertArray(localWorkSize), eventWaitListSize, eventHandles, newEventHandle);
OpenCLException.ThrowOnError(error);
if (newEvents != null)
{
lock (newEvents)
{
newEvents.Add(new OpenCLEvent(newEventHandle[0], this));
}
}
}
public static Error clEnqueueTask(OpenCLCommandQueue command_queue, OpenCLKernel kernel, Int32 num_events_in_wait_list, [In] OpenCLEvent[] event_wait_list, OpenCLEvent e)
{
Console.WriteLine("Calling Error clEnqueueTask(OpenCLCommandQueue command_queue, OpenCLKernel kernel, Int32 num_events_in_wait_list, [In] OpenCLEvent[] event_wait_list, OpenCLEvent e)");
return default(Error);
}
public static Error clEnqueueNDRangeKernel(OpenCLCommandQueue command_queue, OpenCLKernel kernel, Int32 work_dim, [In] IntPtr[] global_work_offset, [In] IntPtr[] global_work_size, [In] IntPtr[] local_work_size, Int32 num_events_in_wait_list, [In] OpenCLEvent[] event_wait_list, OpenCLEvent e)
{
Console.WriteLine("Calling Error clEnqueueNDRangeKernel(OpenCLCommandQueue command_queue, OpenCLKernel kernel, Int32 work_dim, [In] IntPtr[] global_work_offset, [In] IntPtr[] global_work_size, [In] IntPtr[] local_work_size, Int32 num_events_in_wait_list, [In] OpenCLEvent[] event_wait_list, OpenCLEvent e)");
return default(Error);
}
public static extern Error clEnqueueNDRangeKernel(OpenCLCommandQueue command_queue, OpenCLKernel kernel, Int32 work_dim, [In] IntPtr[] global_work_offset, [In] IntPtr[] global_work_size, [In] IntPtr[] local_work_size, Int32 num_events_in_wait_list, [In] OpenCLEvent[] event_wait_list, OpenCLEvent e);
