PyCUDA - 传递一个矩阵引用从Python到C++ CUDA代码

Question

我必须写一个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 

是，我不能通过并购作为列的RES [] [数字]在函数的声明中。我不能留下未申报的列数...

我需要一个矩阵NxM作为输出，但我不知道如何做到这一点。你可以帮我吗？

Answer 1

您应该使用内核中投线性存储器访问，那是多么ndarray和gpuarray数据存储在内部，并且当它作为一个参数传递给供应PyCUDA将传递一个指针在分配给gpuarray在GPU mempoy数据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