#312

Committed 06 May 2025 08:41PM UTC coverage: 67.925% (+0.4%) from 67.512%

Build # #312

Build Type

push

Committed by

Commit Message

Refactor OpenCL I/O to use unified chunk layout and consistent endianness handling

Run Details

0 of 10 new or added lines in 3 files covered. (0.0%)

2 existing lines in 2 files now uncovered.

1224 of 1802 relevant lines covered (67.92%)

0.68 hits per line

Source File
Press 'n' to go to next uncovered line, 'b' for previous

0.0

/src/main/java/net/ladenthin/bitcoinaddressfinder/OpenClTask.java

// @formatter:off
/**
 * Copyright 2020 Bernard Ladenthin bernard.ladenthin@gmail.com
 *
 * Licensed under the Apache License, Version 2.0 (the "License");
 * you may not use this file except in compliance with the License.
 * You may obtain a copy of the License at
 *
 *    http://www.apache.org/licenses/LICENSE-2.0
 *
 * Unless required by applicable law or agreed to in writing, software
 * distributed under the License is distributed on an "AS IS" BASIS,
 * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
 * See the License for the specific language governing permissions and
 * limitations under the License.
 *
 */
// @formatter:on
package net.ladenthin.bitcoinaddressfinder;

import com.google.common.annotations.VisibleForTesting;
import java.math.BigInteger;
import java.nio.ByteBuffer;
import java.nio.ByteOrder;
import net.ladenthin.bitcoinaddressfinder.configuration.CProducer;
import static org.jocl.CL.CL_MEM_READ_ONLY;
import static org.jocl.CL.CL_MEM_USE_HOST_PTR;
import static org.jocl.CL.CL_MEM_WRITE_ONLY;
import static org.jocl.CL.CL_TRUE;
import static org.jocl.CL.clCreateBuffer;
import static org.jocl.CL.clEnqueueNDRangeKernel;
import static org.jocl.CL.clEnqueueReadBuffer;
import static org.jocl.CL.clEnqueueWriteBuffer;
import static org.jocl.CL.clFinish;
import static org.jocl.CL.clReleaseMemObject;
import static org.jocl.CL.clSetKernelArg;
import org.jocl.Pointer;
import org.jocl.Sizeof;
import org.jocl.cl_command_queue;
import org.jocl.cl_context;
import org.jocl.cl_kernel;
import org.jocl.cl_mem;
import org.slf4j.Logger;
import org.slf4j.LoggerFactory;

public class OpenClTask {

    protected Logger logger = LoggerFactory.getLogger(this.getClass());
    
    /**
     * I din't know which is better.
     */
    private static final boolean USE_HOST_PTR = false;
    
    private final CProducer cProducer;

    private final cl_context context;
    private final ByteBuffer srcByteBuffer;
    private final Pointer srcPointer;

    private final cl_mem srcMem;
    private final BitHelper bitHelper;
    private final ByteBufferUtility byteBufferUtility;
    private final BigInteger maxPrivateKeyForBatchSize;

    // Only available after init
    public OpenClTask(cl_context context, CProducer cProducer, BitHelper bitHelper, ByteBufferUtility byteBufferUtility) {
        this.context = context;
        this.cProducer = cProducer;
        this.bitHelper = bitHelper;
        this.byteBufferUtility = byteBufferUtility;

        int srcSizeInBytes = getSrcSizeInBytes();
        maxPrivateKeyForBatchSize = KeyUtility.getMaxPrivateKeyForBatchSize(cProducer.batchSizeInBits);
        srcByteBuffer = ByteBuffer.allocateDirect(srcSizeInBytes);
        srcPointer = Pointer.to(srcByteBuffer);
        srcMem = clCreateBuffer(
                context,
                CL_MEM_READ_ONLY | CL_MEM_USE_HOST_PTR,
                srcSizeInBytes,
                srcPointer,
                null
        );
    }

    public int getSrcSizeInBytes() {
        return PublicKeyBytes.PRIVATE_KEY_MAX_NUM_BYTES;
    }

    public long getDstSizeInBytes() {
        return (long) PublicKeyBytes.CHUNK_SIZE_NUM_BYTES * bitHelper.convertBitsToSize(cProducer.batchSizeInBits);
    }

    /**
    * Writes the base private key to the source buffer in the format expected by the OpenCL kernel.
    * <p>
    * The method ensures that the provided private key is valid for the current batch size. If it exceeds
    * the allowed range, a {@link PrivateKeyTooLargeException} is thrown.
    * <p>
    * Internally, the private key is first converted to a byte array in Big-Endian format (as returned
    * by {@link BigInteger#toByteArray()}). Because the OpenCL kernel expects the private key as a
    * {@code __global const u32 *k} array in <strong>Little-Endian</strong> word order, the byte array
    * is then converted from Big-Endian to Little-Endian before being written to the OpenCL input buffer.
    * <p>
    * This matches the behavior of the OpenCL kernel {@code generateKeysKernel_grid}, which reads the key
    * using {@code copy_u32_array(k_littleEndian_local, k, ...)} assuming Little-Endian input and applies
    * the work-item ID to the least-significant word.
    *
    * @param privateKeyBase the base private key used as input to the OpenCL kernel
    * @throws PrivateKeyTooLargeException if the key is too large for the current batch size
    */
    public void setSrcPrivateKeyChunk(BigInteger privateKeyBase) {
        if (KeyUtility.isInvalidWithBatchSize(privateKeyBase, maxPrivateKeyForBatchSize)) {
            throw new PrivateKeyTooLargeException(privateKeyBase, maxPrivateKeyForBatchSize, cProducer.batchSizeInBits);
        }

        // BigInteger.toByteArray() always returns a big-endian (MSB-first) representation, 
        // meaning the most significant byte (MSB) comes first.
        // Therefore, the source format is always Big Endian.
        byte[] byteArray = byteBufferUtility.bigIntegerToBytes(privateKeyBase);
        EndiannessConverter endiannessConverter = new EndiannessConverter(ByteOrder.BIG_ENDIAN, ByteOrder.LITTLE_ENDIAN, byteBufferUtility);
        endiannessConverter.convertEndian(byteArray);
        byteBufferUtility.putToByteBuffer(srcByteBuffer, byteArray);
    }
    
    @VisibleForTesting
    public ByteBuffer getSrcByteBuffer() {
        return srcByteBuffer;
    }

    public ByteBuffer executeKernel(cl_kernel kernel, cl_command_queue commandQueue) {
        final long dstSizeInBytes = getDstSizeInBytes();
        // allocate a new dst buffer that a clone afterwards is not necessary
        final ByteBuffer dstByteBuffer = ByteBuffer.allocateDirect(ByteBufferUtility.ensureByteBufferCapacityFitsInt(dstSizeInBytes));
        final Pointer dstPointer = Pointer.to(dstByteBuffer);
        final cl_mem dstMem;
        if (USE_HOST_PTR) {
            dstMem = clCreateBuffer(context,
                    CL_MEM_USE_HOST_PTR, dstSizeInBytes,
                    dstPointer,
                    null
            );
        } else {
            dstMem = clCreateBuffer(context,
                    CL_MEM_WRITE_ONLY, dstSizeInBytes,
                    null,
                    null
            );
        }

        // Set the arguments for the kernel
        clSetKernelArg(kernel, 0, Sizeof.cl_mem, Pointer.to(dstMem));
        clSetKernelArg(kernel, 1, Sizeof.cl_mem, Pointer.to(srcMem));

        // Set the work-item dimensions
        long global_work_size[] = new long[]{bitHelper.convertBitsToSize(cProducer.batchSizeInBits)};
        long localWorkSize[] = null; // new long[]{1}; // enabling the system to choose the work-group size.
        int workDim = 1;

        {
            // write src buffer
            clEnqueueWriteBuffer(
                    commandQueue,
                    srcMem,
                    CL_TRUE,
                    0,
                    getSrcSizeInBytes(),
                    srcPointer,
                    0,
                    null,
                    null
            );
            clFinish(commandQueue);
        }
        {
            // execute the kernel
            long beforeExecute = System.currentTimeMillis();
            clEnqueueNDRangeKernel(
                    commandQueue,
                    kernel,
                    workDim,
                    null,
                    global_work_size,
                    localWorkSize,
                    0,
                    null,
                    null
            );
            clFinish(commandQueue);

            long afterExecute = System.currentTimeMillis();
            
            if (logger.isTraceEnabled()) {
                logger.trace("Executed OpenCL kernel in " + (afterExecute - beforeExecute) + "ms");
            }
        }
        {
            // read the dst buffer
            long beforeRead = System.currentTimeMillis();

            clEnqueueReadBuffer(commandQueue,
                    dstMem,
                    CL_TRUE,
                    0,
                    dstSizeInBytes,
                    dstPointer,
                    0,
                    null,
                    null
            );
            clFinish(commandQueue);
            clReleaseMemObject(dstMem);

            long afterRead = System.currentTimeMillis();
            if (logger.isTraceEnabled()) {
                logger.trace("Read OpenCL data "+((dstSizeInBytes / 1024) / 1024) + "Mb in " + (afterRead - beforeRead) + "ms");
            }
        }
        return dstByteBuffer;
    }

    public void releaseCl() {
        clReleaseMemObject(srcMem);
    }

    /**
     * https://stackoverflow.com/questions/3366925/deep-copy-duplicate-of-javas-bytebuffer/4074089
     */
    private static ByteBuffer cloneByteBuffer(final ByteBuffer original) {
        // Create clone with same capacity as original.
        final ByteBuffer clone = (original.isDirect())
                ? ByteBuffer.allocateDirect(original.capacity())
                : ByteBuffer.allocate(original.capacity());

        // Create a read-only copy of the original.
        // This allows reading from the original without modifying it.
        final ByteBuffer readOnlyCopy = original.asReadOnlyBuffer();

        // Flip and read from the original.
        readOnlyCopy.flip();
        clone.put(readOnlyCopy);

        return clone;
    }

}

1	// @formatter:off
2	/**
3	* Copyright 2020 Bernard Ladenthin bernard.ladenthin@gmail.com
4	*
5	* Licensed under the Apache License, Version 2.0 (the "License");
6	* you may not use this file except in compliance with the License.
7	* You may obtain a copy of the License at
8	*
9	* http://www.apache.org/licenses/LICENSE-2.0
10	*
11	* Unless required by applicable law or agreed to in writing, software
12	* distributed under the License is distributed on an "AS IS" BASIS,
13	* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
14	* See the License for the specific language governing permissions and
15	* limitations under the License.
16	*
17	*/
18	// @formatter:on
19	package net.ladenthin.bitcoinaddressfinder;
20
21	import com.google.common.annotations.VisibleForTesting;
22	import java.math.BigInteger;
23	import java.nio.ByteBuffer;
24	import java.nio.ByteOrder;
25	import net.ladenthin.bitcoinaddressfinder.configuration.CProducer;
26	import static org.jocl.CL.CL_MEM_READ_ONLY;
27	import static org.jocl.CL.CL_MEM_USE_HOST_PTR;
28	import static org.jocl.CL.CL_MEM_WRITE_ONLY;
29	import static org.jocl.CL.CL_TRUE;
30	import static org.jocl.CL.clCreateBuffer;
31	import static org.jocl.CL.clEnqueueNDRangeKernel;
32	import static org.jocl.CL.clEnqueueReadBuffer;
33	import static org.jocl.CL.clEnqueueWriteBuffer;
34	import static org.jocl.CL.clFinish;
35	import static org.jocl.CL.clReleaseMemObject;
36	import static org.jocl.CL.clSetKernelArg;
37	import org.jocl.Pointer;
38	import org.jocl.Sizeof;
39	import org.jocl.cl_command_queue;
40	import org.jocl.cl_context;
41	import org.jocl.cl_kernel;
42	import org.jocl.cl_mem;
43	import org.slf4j.Logger;
44	import org.slf4j.LoggerFactory;
45
46	public class OpenClTask {
47
48	protected Logger logger = LoggerFactory.getLogger(this.getClass());	×
49
50	/**
51	* I din't know which is better.
52	*/
53	private static final boolean USE_HOST_PTR = false;
54
55	private final CProducer cProducer;
56
57	private final cl_context context;
58	private final ByteBuffer srcByteBuffer;
59	private final Pointer srcPointer;
60
61	private final cl_mem srcMem;
62	private final BitHelper bitHelper;
63	private final ByteBufferUtility byteBufferUtility;
64	private final BigInteger maxPrivateKeyForBatchSize;
65
66	// Only available after init
NEW 67	public OpenClTask(cl_context context, CProducer cProducer, BitHelper bitHelper, ByteBufferUtility byteBufferUtility) {	×
68	this.context = context;	×
69	this.cProducer = cProducer;	×
70	this.bitHelper = bitHelper;	×
71	this.byteBufferUtility = byteBufferUtility;	×
72
UNCOV 73	int srcSizeInBytes = getSrcSizeInBytes();	×
74	maxPrivateKeyForBatchSize = KeyUtility.getMaxPrivateKeyForBatchSize(cProducer.batchSizeInBits);	×
75	srcByteBuffer = ByteBuffer.allocateDirect(srcSizeInBytes);	×
76	srcPointer = Pointer.to(srcByteBuffer);	×
77	srcMem = clCreateBuffer(	×
78	context,
79	CL_MEM_READ_ONLY \| CL_MEM_USE_HOST_PTR,
80	srcSizeInBytes,
81	srcPointer,
82	null
83	);
84	}	×
85
86	public int getSrcSizeInBytes() {
87	return PublicKeyBytes.PRIVATE_KEY_MAX_NUM_BYTES;	×
88	}
89
90	public long getDstSizeInBytes() {
NEW 91	return (long) PublicKeyBytes.CHUNK_SIZE_NUM_BYTES * bitHelper.convertBitsToSize(cProducer.batchSizeInBits);	×
92	}
93
94	/**
95	* Writes the base private key to the source buffer in the format expected by the OpenCL kernel.
96	* <p>
97	* The method ensures that the provided private key is valid for the current batch size. If it exceeds
98	* the allowed range, a {@link PrivateKeyTooLargeException} is thrown.
99	* <p>
100	* Internally, the private key is first converted to a byte array in Big-Endian format (as returned
101	* by {@link BigInteger#toByteArray()}). Because the OpenCL kernel expects the private key as a
102	* {@code __global const u32 *k} array in <strong>Little-Endian</strong> word order, the byte array
103	* is then converted from Big-Endian to Little-Endian before being written to the OpenCL input buffer.
104	* <p>
105	* This matches the behavior of the OpenCL kernel {@code generateKeysKernel_grid}, which reads the key
106	* using {@code copy_u32_array(k_littleEndian_local, k, ...)} assuming Little-Endian input and applies
107	* the work-item ID to the least-significant word.
108	*
109	* @param privateKeyBase the base private key used as input to the OpenCL kernel
110	* @throws PrivateKeyTooLargeException if the key is too large for the current batch size
111	*/
112	public void setSrcPrivateKeyChunk(BigInteger privateKeyBase) {
113	if (KeyUtility.isInvalidWithBatchSize(privateKeyBase, maxPrivateKeyForBatchSize)) {	×
114	throw new PrivateKeyTooLargeException(privateKeyBase, maxPrivateKeyForBatchSize, cProducer.batchSizeInBits);	×
115	}
116
117	// BigInteger.toByteArray() always returns a big-endian (MSB-first) representation,
118	// meaning the most significant byte (MSB) comes first.
119	// Therefore, the source format is always Big Endian.
120	byte[] byteArray = byteBufferUtility.bigIntegerToBytes(privateKeyBase);	×
NEW 121	EndiannessConverter endiannessConverter = new EndiannessConverter(ByteOrder.BIG_ENDIAN, ByteOrder.LITTLE_ENDIAN, byteBufferUtility);	×
122	endiannessConverter.convertEndian(byteArray);	×
123	byteBufferUtility.putToByteBuffer(srcByteBuffer, byteArray);	×
124	}	×
125
126	@VisibleForTesting
127	public ByteBuffer getSrcByteBuffer() {
128	return srcByteBuffer;	×
129	}
130
131	public ByteBuffer executeKernel(cl_kernel kernel, cl_command_queue commandQueue) {
132	final long dstSizeInBytes = getDstSizeInBytes();	×
133	// allocate a new dst buffer that a clone afterwards is not necessary
134	final ByteBuffer dstByteBuffer = ByteBuffer.allocateDirect(ByteBufferUtility.ensureByteBufferCapacityFitsInt(dstSizeInBytes));	×
135	final Pointer dstPointer = Pointer.to(dstByteBuffer);	×
136	final cl_mem dstMem;
137	if (USE_HOST_PTR) {
138	dstMem = clCreateBuffer(context,
139	CL_MEM_USE_HOST_PTR, dstSizeInBytes,
140	dstPointer,
141	null
142	);
143	} else {
144	dstMem = clCreateBuffer(context,	×
145	CL_MEM_WRITE_ONLY, dstSizeInBytes,
146	null,
147	null
148	);
149	}
150
151	// Set the arguments for the kernel
152	clSetKernelArg(kernel, 0, Sizeof.cl_mem, Pointer.to(dstMem));	×
153	clSetKernelArg(kernel, 1, Sizeof.cl_mem, Pointer.to(srcMem));	×
154
155	// Set the work-item dimensions
156	long global_work_size[] = new long[]{bitHelper.convertBitsToSize(cProducer.batchSizeInBits)};	×
157	long localWorkSize[] = null; // new long[]{1}; // enabling the system to choose the work-group size.	×
158	int workDim = 1;	×
159
160	{
161	// write src buffer
162	clEnqueueWriteBuffer(	×
163	commandQueue,
164	srcMem,
165	CL_TRUE,
166	0,
167	getSrcSizeInBytes(),	×
168	srcPointer,
169	0,
170	null,
171	null
172	);
173	clFinish(commandQueue);	×
174	}
175	{
176	// execute the kernel
177	long beforeExecute = System.currentTimeMillis();	×
178	clEnqueueNDRangeKernel(	×
179	commandQueue,
180	kernel,
181	workDim,
182	null,
183	global_work_size,
184	localWorkSize,
185	0,
186	null,
187	null
188	);
189	clFinish(commandQueue);	×
190
191	long afterExecute = System.currentTimeMillis();	×
192
193	if (logger.isTraceEnabled()) {	×
194	logger.trace("Executed OpenCL kernel in " + (afterExecute - beforeExecute) + "ms");	×
195	}
196	}
197	{
198	// read the dst buffer
199	long beforeRead = System.currentTimeMillis();	×
200
201	clEnqueueReadBuffer(commandQueue,	×
202	dstMem,
203	CL_TRUE,
204	0,
205	dstSizeInBytes,
206	dstPointer,
207	0,
208	null,
209	null
210	);
211	clFinish(commandQueue);	×
212	clReleaseMemObject(dstMem);	×
213
214	long afterRead = System.currentTimeMillis();	×
215	if (logger.isTraceEnabled()) {	×
216	logger.trace("Read OpenCL data "+((dstSizeInBytes / 1024) / 1024) + "Mb in " + (afterRead - beforeRead) + "ms");	×
217	}
218	}
219	return dstByteBuffer;	×
220	}
221
222	public void releaseCl() {
223	clReleaseMemObject(srcMem);	×
224	}	×
225
226	/**
227	* https://stackoverflow.com/questions/3366925/deep-copy-duplicate-of-javas-bytebuffer/4074089
228	*/
229	private static ByteBuffer cloneByteBuffer(final ByteBuffer original) {
230	// Create clone with same capacity as original.
231	final ByteBuffer clone = (original.isDirect())	×
232	? ByteBuffer.allocateDirect(original.capacity())	×
233	: ByteBuffer.allocate(original.capacity());	×
234
235	// Create a read-only copy of the original.
236	// This allows reading from the original without modifying it.
237	final ByteBuffer readOnlyCopy = original.asReadOnlyBuffer();	×
238
239	// Flip and read from the original.
240	readOnlyCopy.flip();	×
241	clone.put(readOnlyCopy);	×
242
243	return clone;	×
244	}
245
246	}

bernardladenthin / BitcoinAddressFinder / #312

Source File Press 'n' to go to next uncovered line, 'b' for previous

Source File
Press 'n' to go to next uncovered line, 'b' for previous