Постоянное использование памяти в коде CUDA

Я сам не могу понять, как лучше всего обеспечить постоянство памяти, используемой в моем ядре. Аналогичный вопрос есть на http: //stackoverflow...r-pleasant-way . Я работаю с GTX580 и компилирую только для возможностей 2.0. Мое ядро ​​выглядит так:

__global__ Foo(const int *src, float *result) {...}

Я выполняю следующий код на хосте:

cudaMalloc(src, size);
cudaMemcpy(src, hostSrc, size, cudaMemcpyHostToDevice);
Foo<<<...>>>(src, result);

альтернативный способ - добавить

__constant__ src[size];

в файл .cu, удалить указатель src из ядра и выполнить

cudaMemcpyToSymbol("src", hostSrc, size, 0, cudaMemcpyHostToDevice);
Foo<<<...>>>(result);

Are эти два способа эквивалентны или первый не гарантирует использование постоянной памяти вместо глобальной памяти? размер изменяется динамически, поэтому второй способ в моем случае не подходит.