BlockSize 7 和 BlockSize 8 时 Cudafy 代码结果不同

Question

我正在使用 Cudafy.NET，我对 BlockSize 有一些困难。它在某些情况下会产生不同的结果。很快的区别就在这里：

//correct results when using this line
gpu.Launch(1, 7, "kernelfx_alldata", 10, devdata, devnmin, devnmax, devgmin, devgmax, devtest);

//incorrect results when using this line
gpu.Launch(1, 8, "kernelfx_alldata", 10, devdata, devnmin, devnmax, devgmin, devgmax, devtest);

问题的详细解释：

我有 10 个项目要循环。 GridSize 为 1.

CASE 1: When CudafyModes.Target = eGPUType.OpenCL and the BlockSize is 1,2,3,4,5,6 and 7. The results are correct.

CASE 2: CudafyModes.Target = eGPUType.OpenCL and the BlockSize is 8,9,10,11, .... and more. The results are incorrect.

CASE 3: CudafyModes.Target = eGPUType.Emulator and the BlockSize is 1,2,3,4,5,6,7,8,9,10,11, .... and more. The results are correct.

示例代码如下所示。初始化变量：

double[,] data;
double[] nmin, nmax, gmin, gmax;

void initializeVars()
{
    data = new double[10, 10];
    for (int i = 0; i < 10; i++)
        {
            data[i, 0] = 100 + i;
            data[i, 1] = 32 + i;
            data[i, 2] = 22 + i;
            data[i, 3] = -20 - i;
            data[i, 4] = 5522 + 10 * i;
            data[i, 5] = 40 + i;
            data[i, 6] = 14 - i;
            data[i, 7] = 12 + i;
            data[i, 8] = -10 + i;
            data[i, 9] = 10 + 10 * i;
        }
    nmin = new double[10];
    nmax= new double[10];
    gmin = new double[10];
    gmax = new double[10];
    for (int i = 0; i < 10; i++)
    {
        nmin[i] = -1;
        nmax[i] = 1;
        gmin[i] = i;
        gmax[i] = 11 * i*i+1;
    }
}

GPU 启动代码：

private void button1_Click(object sender, EventArgs e)
{
    CudafyModes.Target = eGPUType.OpenCL;
    CudafyModes.DeviceId = 0;
    CudafyTranslator.Language = eLanguage.OpenCL;
    CudafyModule km = CudafyTranslator.Cudafy();
    Cudafy.Host.GPGPU gpu = Cudafy.Host.CudafyHost.GetDevice(CudafyModes.Target, CudafyModes.DeviceId);
    gpu.LoadModule(km);
    initializeVars();
    double[,] devdata = gpu.Allocate<double>(data); gpu.CopyToDevice(data, devdata);
    double[] devnmin = gpu.Allocate<double>(nmin); gpu.CopyToDevice(nmin, devnmin);
    double[] devnmax = gpu.Allocate<double>(nmax); gpu.CopyToDevice(nmax, devnmax);
    double[] devgmin = gpu.Allocate<double>(gmin); gpu.CopyToDevice(gmin, devgmin);
    double[] devgmax = gpu.Allocate<double>(gmax); gpu.CopyToDevice(gmax, devgmax);
    double[] test = new double[10];
    double[] devtest = gpu.Allocate<double>(test);
    gpu.Launch(1, 8, "kernelfx_alldata", 10, devdata, devnmin,
           devnmax, devgmin, devgmax,  devtest);
    gpu.CopyFromDevice(devtest, test);
    gpu.FreeAll();
}

Cudafy 内核

[Cudafy]
public static void kernelfx_alldata(GThread thread, int N, double[,] data, double[] nmin, double[] nmax, double[] gmin, double[] gmax, double[] test)
{
    int tid = thread.threadIdx.x + thread.blockIdx.x * thread.blockDim.x;
    while (tid < N)
    {
        double[] tmp = thread.AllocateShared<double>("tmp", 10);
        tmp[0] = 1; 
        for (int i = 1; i < 10; i++)
        {
            tmp[i] = data[tid, i - 1];
        }
        for (int i = 1; i < 10; i++)
        {
            tmp[i] = (nmax[i - 1] - nmin[i - 1]) / (gmax[i - 1] - gmin[i - 1]) * (tmp[i] - gmin[i - 1]) + nmin[i - 1];
        }
        test[tid] = tmp[1];

        tid = tid + thread.blockDim.x * thread.gridDim.x;
    }
}

The Correct (CASE 1 and CASE 3) Results are:

test[0]=199.0
test[1]=201.0
test[2]=203.0
test[3]=205.0
test[4]=207.0
test[5]=209.0
test[6]=211.0
test[7]=213.0
test[8]=215.0
test[9]=217.0

Incorrect (CASE 2) results are:

test[0]=213.0
test[1]=213.0
test[2]=213.0
test[3]=213.0
test[4]=213.0
test[5]=213.0
test[6]=213.0
test[7]=213.0
test[8]=217.0
test[9]=217.0

当BlockSize小于8时，结果正确。但是当 BlockSize 大于 8 时，结果不正确。为了有效地使用 gpu，blockSize 必须大于 8。

这段代码有什么问题？

最好的问候...

Answer 1

将 tmp 声明为二维数组，第一列是 threadId 解决了问题。工作代码如下：

[Cudafy]
public static void kernelfx_alldata(GThread thread, int N, double[,] data, double[] nmin,
                                double[] nmax, double[] gmin, double[] gmax, double[] test)
{
    int tid = thread.threadIdx.x + thread.blockIdx.x * thread.blockDim.x;
    double[,] tmp = thread.AllocateShared<double>("tmp", 10, 10);
    while (tid < N)
    {

        tmp[tid, 0] = 1;
        for (int i = 1; i < 10; i++)
        {
            tmp[tid, i] = data[tid, i - 1];
        }
        for (int i = 1; i < 10; i++)
        {
            tmp[tid, i] = (nmax[i - 1] - nmin[i - 1]) / (gmax[i - 1] - gmin[i - 1]) * (tmp[tid, i] - gmin[i - 1]) + nmin[i - 1];
        }
        test[tid] = tmp[tid, 1];

        tid = tid + thread.blockDim.x * thread.gridDim.x;
    }
}

BlockSize 7 和 BlockSize 8 时 Cudafy 代码结果不同

Cudafy code results different when BlockSize 7 and BlockSize 8

c#

opencl

gpgpu

gpu

cudafy.net