CUDA — большое число и выделение памяти

У меня очень странный баг в программе. Я потратил много часов на это, но я не нашел решения. Я написал простую программу, чтобы воспроизвести мою проблему. Может быть, кто-нибудь поможет мне. Я пробовал cuda-memcheck и Каков канонический способ проверки ошибок с помощью API среды выполнения CUDA? но я не получаю никаких ошибок.

Подробности:

версия nvcc — V6.0.1

версия gcc - 4.8.1

Полный код:

#include <stdio.h>

__constant__ unsigned long long int bigNumber = 83934243334343;
__device__ bool isFound = false;
__global__ void kernel(int *dev_number) {

    unsigned long long int id = threadIdx.x + (blockIdx.x * blockDim.x);
    while (id < bigNumber && isFound==false) {

        if(id == 10) {
            *dev_number = 4;
            isFound=true;
        }
        id++;
    }
}

int main(int argc, char *argv[]) {
    int number = 0;
    int *dev_number;

    printf("Number: %d\n", number);

    return 0;
}

Компиляция и запуск:

nvcc myprogram.cu
./myprogram

Когда я запускаю эту программу, я не получаю никакого возвращаемого значения. Но когда переменная - bigNumber имеет меньшее значение или я не использую cudaMalloc и cudaMemcpy, это работает (это означает, что вызывается возврат 0). Какое соединение должно выделять память для другой переменной с константой bigNumber? В чем проблема?


person Bakus123    schedule 13.04.2015    source источник
comment
разве 839299365868340224 не должно быть 839299365868340224ULL? в противном случае это int.   -  person mch    schedule 13.04.2015
comment
@mch, я уже пробовал...   -  person Bakus123    schedule 13.04.2015


Ответы (1)


Теперь, когда вы изменили код на что-то более разумное, я сразу же получаю следующий результат:

__device__ volatile bool isFound = false;

Спецификатор volatile заставляет компилятор пропускать любые оптимизации, которые не позволяют каждому потоку читать глобальную копию переменной.

Из документации.

Компилятор может оптимизировать чтение и запись в глобальную или разделяемую память (например, кэшируя глобальное чтение в регистры или кэш L1), если он соблюдает семантику упорядочения памяти функций ограничения памяти (функции ограничения памяти) и семантику видимости памяти. функций синхронизации (Synchronization Functions).

Эти оптимизации можно отключить с помощью ключевого слова volatile: если переменная, расположенная в глобальной или разделяемой памяти, объявлена ​​как volatile, компилятор предполагает, что ее значение может быть изменено или использовано в любое время другим потоком, и поэтому любая ссылка на эту переменную компилируется в фактическая инструкция чтения или записи памяти.

Если вы не используете квалификатор volatile, то только один поток принимает условие раннего выхода (isFound), а все остальные должны выполнять цикл в течение очень долгого времени, пока их значение id не превысит bigNumber.

person Robert Crovella    schedule 13.04.2015