Как я могу создать структуру родных указателей в JCuda
У меня есть ядро CUDA, которое принимает список структур.
kernel<<<blockCount,blockSize>>>(MyStruct *structs);
Каждая структура содержит 3 указателя.
typedef struct __align(16)__ {
float* pointer1;
float* pointer2;
float* pointer3;
}
У меня есть три массива устройств, содержащих числа с плавающей точкой, и каждый указатель в структуре указывает на число с плавающей точкой в одном из трех массивов устройств.
Список структур представляет собой древовидную / графическую структуру, которая позволяет ядру выполнять рекурсивные операции в зависимости от порядка списка структур, отправляемых в ядро. (Этот бит работает в C++, поэтому не связан с моей проблемой)
То, что я хотел бы сделать, это иметь возможность отправлять мою структуру указателей из JCuda. Я понимаю, что это изначально невозможно, если оно не сведено к массиву с дополнениями, как в этом посте.
Я понимаю все проблемы с выравниванием и заполнением, которые могут возникнуть при отправке списка структур, по сути, это повторяющийся массив с дополнениями, с которым я в порядке.
Но я не уверен, как это сделать, это заполнить мой уплощенный структурный буфер указателями, например, я думаю, что я могу сделать что-то вроде этого:
Pointer A = ....(underlying device array1)
Pointer B = ....(underlying device array2)
Pointer C = ....(underlying device array3)
ByteBuffer structListBuffer = ByteBuffer.allocate(16*noSteps);
for(int x = 0; x<noSteps; x++) {
// Get the underlying pointer values
long pointer1 = A.withByteOffset(getStepOffsetA(x)).someGetUnderlyingPointerValueFunction();
long pointer2 = B.withByteOffset(getStepOffsetB(x)).someGetUnderlyingPointerValueFunction();
long pointer3 = C.withByteOffset(getStepOffsetC(x)).someGetUnderlyingPointerValueFunction();
// Build the struct
structListBuffer.asLongBuffer().append(pointer1);
structListBuffer.asLongBuffer().append(pointer2);
structListBuffer.asLongBuffer().append(pointer3);
structListBuffer.asLongBuffer().append(0); //padding
}
structListBuffer будет содержать список структур так, как ядро будет ожидать этого.
Так есть ли способ сделать someGetUnderlyingPointerValueFunction() из байтового буфера?
1 ответ
Если я все правильно понял, главный вопрос в том, есть ли такая магическая функция, как
long address = pointer.someGetUnderlyingPointerValueFunction();
который возвращает адрес родного указателя.
Краткий ответ: нет, такой функции нет.
(Примечание: аналогичная функциональность уже была запрошена довольно давно, но я еще не добавил ее. Главным образом потому, что такая функция не имеет смысла для указателей на массивы Java или (непрямых) байтовых буферов. Кроме того, вручную обработка структур с их дополнениями и выравниваниями, а также указатели разных размеров на 32- и 64-битных машинах, а также буферы с прямым или прямым порядком байтов являются бесконечным источником головной боли. Но я вижу смысл, возможный случай применения и т. д. Я, скорее всего, добавлю что-то вроде getAddress() функция. Может быть, только к CUdeviceptr класс, где это определенно имеет смысл - по крайней мере, больше, чем в Pointer учебный класс. Люди будут использовать этот метод, чтобы делать странные вещи, и они будут делать вещи, которые вызовут неприятные сбои виртуальной машины, но сама JCuda является настолько тонким слоем абстракции, что в любом случае в этом отношении нет системы безопасности...)
Тем не менее, вы можете обойти текущее ограничение, используя такой метод:
private static long getPointerAddress(CUdeviceptr p)
{
// WORKAROUND until a method like CUdeviceptr#getAddress exists
class PointerWithAddress extends Pointer
{
PointerWithAddress(Pointer other)
{
super(other);
}
long getAddress()
{
return getNativePointer() + getByteOffset();
}
}
return new PointerWithAddress(p).getAddress();
}
Конечно, это некрасиво и явно противоречит намерению сделать getNativePointer() а также getByteOffset() методы protected, Но в конечном итоге он может быть заменен каким-то "официальным" методом:
private static long getPointerAddress(CUdeviceptr p)
{
return p.getAddress();
}
и до сих пор это, вероятно, решение, наиболее близкое к тому, что вы можете сделать на стороне C.
Вот пример, который я написал для тестирования этого. Ядро является всего лишь фиктивным ядром, которое заполняет структуру "идентифицируемыми" значениями (чтобы увидеть, оказываются ли они в нужном месте), и предполагается, что оно запускается только с 1 потоком:
typedef struct __declspec(align(16)) {
float* pointer1;
float* pointer2;
float* pointer3;
} MyStruct;
extern "C"
__global__ void kernel(MyStruct *structs)
{
structs[0].pointer1[0] = 1.0f;
structs[0].pointer1[1] = 1.1f;
structs[0].pointer1[2] = 1.2f;
structs[0].pointer2[0] = 2.0f;
structs[0].pointer2[1] = 2.1f;
structs[0].pointer2[2] = 2.2f;
structs[0].pointer3[0] = 3.0f;
structs[0].pointer3[1] = 3.1f;
structs[0].pointer3[2] = 3.2f;
structs[1].pointer1[0] = 11.0f;
structs[1].pointer1[1] = 11.1f;
structs[1].pointer1[2] = 11.2f;
structs[1].pointer2[0] = 12.0f;
structs[1].pointer2[1] = 12.1f;
structs[1].pointer2[2] = 12.2f;
structs[1].pointer3[0] = 13.0f;
structs[1].pointer3[1] = 13.1f;
structs[1].pointer3[2] = 13.2f;
}
Это ядро запускается в следующей программе (Примечание. Компиляция файла PTX выполняется здесь на лету с настройками, которые могут не соответствовать вашему приложению. В случае сомнений, вы можете скомпилировать файл PTX вручную).
pointer1, pointer2 а также pointer3 указатели каждой структуры инициализируются так, что они указывают на последовательные элементы буферов устройства A, B а также C соответственно, каждый со смещением, позволяющим идентифицировать значения, записываемые ядром. (Обратите внимание, что я пытался справиться с двумя возможными случаями запуска этого на 32-битной или 64-битной машине, что предполагает различный размер указателя - хотя в настоящее время я могу протестировать только 32-битную версию)
import static jcuda.driver.JCudaDriver.*;
import java.io.ByteArrayOutputStream;
import java.io.File;
import java.io.IOException;
import java.io.InputStream;
import java.nio.ByteBuffer;
import java.nio.ByteOrder;
import java.nio.IntBuffer;
import java.nio.LongBuffer;
import java.util.Arrays;
import jcuda.Pointer;
import jcuda.Sizeof;
import jcuda.driver.CUcontext;
import jcuda.driver.CUdevice;
import jcuda.driver.CUdeviceptr;
import jcuda.driver.CUfunction;
import jcuda.driver.CUmodule;
import jcuda.driver.JCudaDriver;
public class JCudaPointersInStruct
{
public static void main(String args[]) throws IOException
{
JCudaDriver.setExceptionsEnabled(true);
String ptxFileName = preparePtxFile("JCudaPointersInStructKernel.cu");
cuInit(0);
CUdevice device = new CUdevice();
cuDeviceGet(device, 0);
CUcontext context = new CUcontext();
cuCtxCreate(context, 0, device);
CUmodule module = new CUmodule();
cuModuleLoad(module, ptxFileName);
CUfunction function = new CUfunction();
cuModuleGetFunction(function, module, "kernel");
int numElements = 9;
CUdeviceptr A = new CUdeviceptr();
cuMemAlloc(A, numElements * Sizeof.FLOAT);
cuMemsetD32(A, 0, numElements);
CUdeviceptr B = new CUdeviceptr();
cuMemAlloc(B, numElements * Sizeof.FLOAT);
cuMemsetD32(B, 0, numElements);
CUdeviceptr C = new CUdeviceptr();
cuMemAlloc(C, numElements * Sizeof.FLOAT);
cuMemsetD32(C, 0, numElements);
int numSteps = 2;
int sizeOfStruct = Sizeof.POINTER * 4;
ByteBuffer hostStructsBuffer =
ByteBuffer.allocate(numSteps * sizeOfStruct);
if (Sizeof.POINTER == 4)
{
IntBuffer b = hostStructsBuffer.order(
ByteOrder.nativeOrder()).asIntBuffer();
for(int x = 0; x<numSteps; x++)
{
CUdeviceptr pointer1 = A.withByteOffset(getStepOffsetA(x));
CUdeviceptr pointer2 = B.withByteOffset(getStepOffsetB(x));
CUdeviceptr pointer3 = C.withByteOffset(getStepOffsetC(x));
//System.out.println("Step "+x+" pointer1 is "+pointer1);
//System.out.println("Step "+x+" pointer2 is "+pointer2);
//System.out.println("Step "+x+" pointer3 is "+pointer3);
b.put((int)getPointerAddress(pointer1));
b.put((int)getPointerAddress(pointer2));
b.put((int)getPointerAddress(pointer3));
b.put(0);
}
}
else
{
LongBuffer b = hostStructsBuffer.order(
ByteOrder.nativeOrder()).asLongBuffer();
for(int x = 0; x<numSteps; x++)
{
CUdeviceptr pointer1 = A.withByteOffset(getStepOffsetA(x));
CUdeviceptr pointer2 = B.withByteOffset(getStepOffsetB(x));
CUdeviceptr pointer3 = C.withByteOffset(getStepOffsetC(x));
//System.out.println("Step "+x+" pointer1 is "+pointer1);
//System.out.println("Step "+x+" pointer2 is "+pointer2);
//System.out.println("Step "+x+" pointer3 is "+pointer3);
b.put(getPointerAddress(pointer1));
b.put(getPointerAddress(pointer2));
b.put(getPointerAddress(pointer3));
b.put(0);
}
}
CUdeviceptr structs = new CUdeviceptr();
cuMemAlloc(structs, numSteps * sizeOfStruct);
cuMemcpyHtoD(structs, Pointer.to(hostStructsBuffer),
numSteps * sizeOfStruct);
Pointer kernelParameters = Pointer.to(
Pointer.to(structs)
);
cuLaunchKernel(function,
1, 1, 1,
1, 1, 1,
0, null, kernelParameters, null);
cuCtxSynchronize();
float hostA[] = new float[numElements];
cuMemcpyDtoH(Pointer.to(hostA), A, numElements * Sizeof.FLOAT);
float hostB[] = new float[numElements];
cuMemcpyDtoH(Pointer.to(hostB), B, numElements * Sizeof.FLOAT);
float hostC[] = new float[numElements];
cuMemcpyDtoH(Pointer.to(hostC), C, numElements * Sizeof.FLOAT);
System.out.println("A "+Arrays.toString(hostA));
System.out.println("B "+Arrays.toString(hostB));
System.out.println("C "+Arrays.toString(hostC));
}
private static long getStepOffsetA(int x)
{
return x * Sizeof.FLOAT * 4 + 0 * Sizeof.FLOAT;
}
private static long getStepOffsetB(int x)
{
return x * Sizeof.FLOAT * 4 + 1 * Sizeof.FLOAT;
}
private static long getStepOffsetC(int x)
{
return x * Sizeof.FLOAT * 4 + 2 * Sizeof.FLOAT;
}
private static long getPointerAddress(CUdeviceptr p)
{
// WORKAROUND until a method like CUdeviceptr#getAddress exists
class PointerWithAddress extends Pointer
{
PointerWithAddress(Pointer other)
{
super(other);
}
long getAddress()
{
return getNativePointer() + getByteOffset();
}
}
return new PointerWithAddress(p).getAddress();
}
//-------------------------------------------------------------------------
// Ignore this - in practice, you'll compile the PTX manually
private static String preparePtxFile(String cuFileName) throws IOException
{
int endIndex = cuFileName.lastIndexOf('.');
if (endIndex == -1)
{
endIndex = cuFileName.length()-1;
}
String ptxFileName = cuFileName.substring(0, endIndex+1)+"ptx";
File cuFile = new File(cuFileName);
if (!cuFile.exists())
{
throw new IOException("Input file not found: "+cuFileName);
}
String modelString = "-m"+System.getProperty("sun.arch.data.model");
String command =
"nvcc " + modelString + " -ptx -arch sm_11 -lineinfo "+
cuFile.getPath()+" -o "+ptxFileName;
System.out.println("Executing\n"+command);
Process process = Runtime.getRuntime().exec(command);
String errorMessage =
new String(toByteArray(process.getErrorStream()));
String outputMessage =
new String(toByteArray(process.getInputStream()));
int exitValue = 0;
try
{
exitValue = process.waitFor();
}
catch (InterruptedException e)
{
Thread.currentThread().interrupt();
throw new IOException(
"Interrupted while waiting for nvcc output", e);
}
if (exitValue != 0)
{
System.out.println("nvcc process exitValue "+exitValue);
System.out.println("errorMessage:\n"+errorMessage);
System.out.println("outputMessage:\n"+outputMessage);
throw new IOException(
"Could not create .ptx file: "+errorMessage);
}
System.out.println("Finished creating PTX file");
return ptxFileName;
}
private static byte[] toByteArray(InputStream inputStream)
throws IOException
{
ByteArrayOutputStream baos = new ByteArrayOutputStream();
byte buffer[] = new byte[8192];
while (true)
{
int read = inputStream.read(buffer);
if (read == -1)
{
break;
}
baos.write(buffer, 0, read);
}
return baos.toByteArray();
}
}
Результат, как и ожидалось / желательно:
A [1.0, 1.1, 1.2, 0.0, 11.0, 11.1, 11.2, 0.0, 0.0]
B [0.0, 2.0, 2.1, 2.2, 0.0, 12.0, 12.1, 12.2, 0.0]
C [0.0, 0.0, 3.0, 3.1, 3.2, 0.0, 13.0, 13.1, 13.2]