Условия гонки, несмотря на атомные функции (CUDA)?

Question

У меня проблема, которая параллельна на двух уровнях: у меня есть тонна наборов пар координат (x0, x1, y0, y1), которые превращаются в переменные vdx, vdy, vyy, и для каждого из этих множеств я пытаюсь вычислить значения всех «мономов», составленных из них до степени n (т.е. все возможные комбинации их различных степеней, такие как vdx^3*vdy*vyy^2 или vdx*1*vyy^4). Эти значения затем складываются по всем наборам.

Моя стратегия (и на данный момент я бы просто хотел, чтобы она работала, ее не нужно оптимизировать с помощью нескольких ядер или сложных сокращений, если это действительно необходимо) заключается в том, чтобы каждая нить имела дело с одним набор пар координат и вычислить значения всех их соответствующих одночленов. В общей памяти каждого блока хранятся все мономиальные суммы, и когда блок выполнен, первый поток в блоке добавляет результат к глобальной сумме. Поскольку доступ к общей памяти каждого блока осуществляется всеми потоками во всех местах, я использую atomicAdd; то же самое с блоками и глобальной памятью.

К сожалению, до сих пор, кажется, что-то есть условие гонки, так как у меня разные результаты каждый раз, когда я запускаю ядро.

Если это поможет, я в настоящее время с помощью degree = 3 и опуская одну из переменных, а это означает, что в коде ниже, сокровенное цикл (более evbl) ничего не делать и просто повторяется 4 раза. Действительно, вывод ядра выглядит следующим образом: 51502,55043.1,55043.1,51502,47868.5,47868.5,48440.5,48440.6,46284.7,46284.7,46284.7,46284.7,46034.3,46034.3,46034.3,46034.3,44972.8,44972.8,44972.8,44972.8,43607.6,43607.6,43607.6,43607.6,43011,43011,43011,43011,42747.8,42747.8,42747.8,42747.8,45937.8,45937.8,46509.9,46509.9,..., и это замечательно, что есть (грубая) схема из 4-х кортежей. Но каждый раз, когда я запускаю его, все значения очень разные.

Все в поплавках, но я на 2.1 GPU, и это не должно быть проблемой. cuda-memcheck также не сообщает об ошибках.

Может ли кто-нибудь, у кого больше опыта в CUDA, дать мне несколько указателей, как отслеживать состояние гонки здесь?

__global__ void kernel(...) { 

    extern __shared__ float s_data[]; 

    // just use global memory for now 
    // get threadID: 
    int idx = blockIdx.x * blockDim.x + threadIdx.x; 
    if(idx >= nPairs) return; 

    // ... do some calculations to get x/y... 

    // calculate vdx, vdy and vyy 
    float vdx = (x1 - x0)/(float)xheight; 
    float vdy = (y1 - y0)/(float)xheight; 
    float vyy = 0.5*(y0 + y1)/(float)xheight; 


    const int offs1 = degree + 1; 
    const int offs2 = offs1 * offs1; 
    const int offs3 = offs2 * offs1; 
    float sol = 1.0; 

    // now calculate monomial results and store in shared memory 

    for(int evdx = 0; evdx <= degree; evdx++) { 
    for(int evdy = 0; evdy <= degree; evdy++) { 
     for(int evyy = 0; evyy <= degree; evyy++) { 
     for(int evbl = 0; evbl <= degree; evbl++) { 
      s = powf(vdx, evdx) + powf(vdy, evdy) + powf(vyy, evyy); 
      atomicAdd(&(s_data[evbl + offs1*evyy + offs2*evdy + 
       offs3*evdx]), sol/1000.0); 

     } 
     } 
    } 
    } 

    // now copy shared memory to global 
    __syncthreads(); 
    if(threadIdx.x == 0) { 
    for(int i = 0; i < nMonomials; i++) { 
     atomicAdd(&outmD[i], s_data[i]); 
    } 
    } 
} 

Answer 1

Вы используете общую память, но вы никогда не инициализируете ее.