Результаты кода Cudafy отличаются, когда BlockSize 7 и BlockSize 8

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.

ПРИМЕР 1: Когда CudafyModes.Target = eGPUType. OpenCL и BlockSize - 1,2,3,4,5,6 и 7. Результаты исправьте.

CASE 2: CudafyModes.Target = eGPUType. OpenCL и BlockSize 8,9,10,11, .... и больше. Полученные результаты: Не верно.

CASE 3: CudafyModes.Target = eGPUType. Эмулятор и BlockSize - 1,2,3,4,5,6,7,8,9,10,11, .... и более. Результаты: правильный.

Пример кода показан ниже. Инициализация переменных:

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 Launch Код:

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; 
    } 
} 

Правильный (CASE 1 и СЛУЧАЙ 3) Результаты:

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 

Неверные (случай 2) Результаты:

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 в 2d массив, первый столбец является 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; 
    } 
}