PyCUDA - передача матрицы по ссылке из python в C ++ CUDA-код - PullRequest
Новая
       35

PyCUDA - передача матрицы по ссылке из python в C ++ CUDA-код

0 голосов
/ 07 августа 2011

Мне нужно написать функцию PyCUDA, которая получает две матрицы Nx3 и Mx3 и возвращает матрицу NxM, но я не могу понять, как передать по ссылке матрицу, не зная количества столбцов.

Мой код в основном такой: 

#kernel declaration
mod = SourceModule("""
__global__ void distance(int N, int M, float d1[][3], float d2[][3], float res[][M])
{
    int i = threadIdx.x;
    int j = threadIdx.y;
    float x, y, z;
    x = d2[j][0]-d1[i][0];
    y = d2[j][1]-d1[i][1];
    z = d2[j][2]-d1[i][2];
    res[i][j] = x*x + y*y + z*z;
}
""")

#load data
data1 = numpy.loadtxt("data1.txt").astype(numpy.float32) # Nx3 matrix
data2 = numpy.loadtxt("data2.txt").astype(numpy.float32) # Mx3 matrix
N=data1.shape[0]
M=data2.shape[0]
res = numpy.zeros([N,M]).astype(numpy.float32) # NxM matrix

#invoke kernel
dist_gpu = mod.get_function("distance")
dist_gpu(cuda.In(numpy.int32(N)), cuda.In(numpy.int32(M)), cuda.In(data1), cuda.In(data2), cuda.Out(res), block=(N,M,1))

#save data
numpy.savetxt("results.txt", res)

При компиляции я получаю сообщение об ошибке: 

kernel.cu(3): error: a parameter is not allowed

то есть я не могу использовать M в качестве числа столбцов для res [] [] в объявлении функции. Я не могу оставить количество столбцов необъявленным ...

Мне нужна матрица NxM в качестве выхода, но я не могу понять, как это сделать. Вы можете мне помочь?

1 Ответ

1 голос
/ 07 августа 2011

Вы должны использовать тональный линейный доступ к памяти внутри ядра, то есть ndarray и gpuarray хранят данные внутри, а PyCUDA передаст указатель на данные в памяти gpu, выделенные для gpuarray, когда они будут предоставлены.в качестве аргумента для ядра PyCUDA.Поэтому (если я понимаю, что вы пытаетесь сделать), ваше ядро ​​должно быть записано примерно так: 

__device__ unsigned int idx2d(int i, int j, int lda)
{
    return j + i*lda;
}

__global__ void distance(int N, int M, float *d1, float *d2, float *res)
{
    int i = threadIdx.x + blockDim.x * blockIdx.x;
    int j = threadIdx.y + blockDim.y * blockIdx.y;
    float x, y, z;
    x = d2[idx2d(j,0,3)]-d1[idx2d(i,0,3)];
    y = d2[idx2d(j,1,3)]-d1[idx2d(i,1,3)];
    z = d2[idx2d(j,2,3)]-d1[idx2d(i,2,3)];

    res[idx2d(i,j,N)] = x*x + y*y + z*z;
}

Здесь я принял основной порядок строк по умолчанию numpy при определении вспомогательной функции idx2d,Все еще есть проблемы со стороной Python кода, который вы разместили, но я думаю, вы уже это знаете.

РЕДАКТИРОВАТЬ : Вот полный рабочий сценарий воспроизведения, основанный на кодеопубликовано в вашем вопросе.Обратите внимание, что в нем используется только один блок (как и в оригинале), поэтому помните о размерах блоков и сетки, пытаясь запустить его для чего-либо, кроме как в небольших случаях. 

import numpy as np
from pycuda import compiler, driver
from pycuda import autoinit

#kernel declaration
mod = compiler.SourceModule("""
__device__ unsigned int idx2d(int i, int j, int lda)
{
    return j + i*lda;
}

__global__ void distance(int N, int M, float *d1, float *d2, float *res)
{
    int i = threadIdx.x + blockDim.x * blockIdx.x;
    int j = threadIdx.y + blockDim.y * blockIdx.y;
    float x, y, z;
    x = d2[idx2d(j,0,3)]-d1[idx2d(i,0,3)];
    y = d2[idx2d(j,1,3)]-d1[idx2d(i,1,3)];
    z = d2[idx2d(j,2,3)]-d1[idx2d(i,2,3)];

    res[idx2d(i,j,N)] = x*x + y*y + z*z;
}
""")

#make data
data1 = np.random.uniform(size=18).astype(np.float32).reshape(-1,3)
data2 = np.random.uniform(size=12).astype(np.float32).reshape(-1,3)
N=data1.shape[0]
M=data2.shape[0]
res = np.zeros([N,M]).astype(np.float32) # NxM matrix

#invoke kernel
dist_gpu = mod.get_function("distance")
dist_gpu(np.int32(N), np.int32(M), driver.In(data1), driver.In(data2), \
        driver.Out(res), block=(N,M,1), grid=(1,1))

print res
...