0

2つの行列Nx3とMx3を取得し、行列NxMを返すPyCUDA関数を作成する必要がありますが、列の数がわからないと、参照によって行列を渡す方法がわかりません。

私のコードは基本的に次のようなものです。

#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[][]の列数としてMを使用することはできません。列数を宣言しないままにすることもできません...

出力として行列NxMが必要ですが、これを行う方法がわかりません。手伝って頂けますか？

4

1 に答える 1

1

カーネル内でピッチ線形メモリアクセスを使用する必要があります。つまり、データを内部に格納する方法ですndarray。PyCUDAは、PyCUDAカーネルに引数として提供されるときに、にgpuarray割り当てられたgpuメモリ内のデータへのポインタを渡します。gpuarrayしたがって（私があなたがやろうとしていることを理解しているなら）あなたのカーネルは次のように書かれるべきです：

__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
于 2011-08-07T07:31:53.100 に答える