Код CUDA не работает, почему? - PullRequest
Новая
       45

Код CUDA не работает, почему?

1 голос
/ 21 февраля 2012

У меня есть этот код, но иногда он работает, иногда НЕТ (напишите printf ("ERR:% d \ n", id)).Я работаю под CUDA 4.1 и у меня есть GTS450, который имеет вычислительные возможности 2.1.

Код не имеет более высокой цели, я просто пытаюсь выяснить, почему он не работает, потому что Мой разум говорит мне, что правильно:]

Если вы хотите запустить его, может быть, вам нужно выполнить несколько раз, когда появляется «ошибка» или изменить размер сетки!

PS: здесь вы можете скачать exe-файл для win64 - вам нужноиметь драйвер cuda4.1 

class MAN
{
public:
    int m_id;
    int m_use;

    __device__
    MAN()
    {
        m_id = -1;
        m_use = 0;
    }
};

__device__ int* d_ids = NULL;
__device__ int d_last_ids = 0;

__device__ MAN* d_mans = NULL;


__global__ void init()
{
    d_mans = new MAN[500];  //note: 500 is more than enough!
    d_ids = new int[500];

    for(int i=0; i < 500; i++)
        d_ids[i] = 0;
}


__device__ int getMAN() //every block get unique number, so at one moment all running blocks has different id
{
    while(true)
    {
        for(int i=0; i < 500; i++)
            if(atomicCAS(&(d_mans[i].m_use), 0, 1)==0)
                return i;
    }
}
__device__ void returnMAN(int id)
{
    int s = atomicExch(&(d_mans[id].m_use), 0);
}



__global__ void testIt()
{
    if(threadIdx.x==0)
    {
        int man = getMAN();

        int id = d_mans[man].m_id;
        if(id == -1)    //If It never works with this "id", its creating new
        {
            id = atomicAdd(&d_last_ids, 2);

            d_ids[id] = 10; //set to non-zero
            d_mans[man].m_id = id;  //save new id for next time

            printf("ADD:%d\n", id);
        }

        if(d_ids[id]==0)
            printf("ERR:%d\n", id); //THIS SHOULD NEVER HAPPEN, BUT BECOMES !!!

        returnMAN(man);
    }
}



int main()
{
    init<<<1, 1>>>();
    printf("init() err: %d\n", cudaDeviceSynchronize());

    testIt<<<20000, 512>>>();
    printf("testIt() err: %d\n", cudaDeviceSynchronize());

    getchar();
    return 0;
}

Ответы [ 2 ]

1 голос
/ 21 февраля 2012

Это, кажется, происходит, потому что этот код 

    int id = d_mans[man].m_id;
    if(id == -1)    //If It never works with this "id", its creating new
    {
        id = atomicAdd(&d_last_ids, 2);

        d_ids[id] = 10; //set to non-zero
        d_mans[man].m_id = id;  //save new id for next time

        printf("ADD:%d\n", id);
    }

    if(d_ids[id]==0)
        printf("ERR:%d\n", id); //THIS SHOULD NEVER HAPPEN, BUT BECOMES !!!

Содержит состояние гонки, если какой-то блок записал в d_mans [man] .m_id, но еще не записал в d_ids [id]. Вероятно, компилятор обменивается инструкцией «установлено на ненулевое значение» и «сохранить новый идентификатор для следующего раза» или кэш просто не обновляется вовремя.

На самом деле проблема в вашем распределителе - лучше запомнить индекс последнего использованного 'man', чем искать его.

0 голосов
/ 23 февраля 2012

Я изменил это: 

__device__ int* d_ids = NULL;

на это: 

__device__ volatile int* d_ids = NULL;

и все работает нормально !!!

И даже Ему не нужно __threadfence ();

...