Глобальные (как в C) динамические массивы CUDA, выделенные в памяти устройства.
https://stackoverflow.com/questions/79121
italiano
english
français
española
中国
日本の
العربية
Deutsch
한국어
Português
RussianВопрос
Итак, я пытаюсь написать код, использующий архитектуру CUDA Nvidia.Я заметил, что копирование на устройство и с него серьезно снижает мою общую производительность, поэтому теперь я пытаюсь переместить на устройство большой объем данных.
Поскольку эти данные используются во многих функциях, я бы хотел, чтобы они были глобальными.Да, я могу передавать указатели, но мне бы очень хотелось знать, как работать с глобальными переменными в этом случае.
Итак, у меня есть функции устройства, которые хотят получить доступ к массиву, выделенному устройству.
В идеале я мог бы сделать что-то вроде:
__device__ float* global_data;
main()
{
cudaMalloc(global_data);
kernel1<<<blah>>>(blah); //access global data
kernel2<<<blah>>>(blah); //access global data again
}
Однако я не понял, как создать динамический массив.Я нашел обходной путь, объявив массив следующим образом:
__device__ float global_data[REALLY_LARGE_NUMBER];
И хотя для этого не требуется вызов cudaMalloc, я бы предпочел подход динамического выделения.
Решение
Что-то вроде этого, вероятно, должно сработать.
#include <algorithm>
#define NDEBUG
#define CUT_CHECK_ERROR(errorMessage) do { \
cudaThreadSynchronize(); \
cudaError_t err = cudaGetLastError(); \
if( cudaSuccess != err) { \
fprintf(stderr, "Cuda error: %s in file '%s' in line %i : %s.\n", \
errorMessage, __FILE__, __LINE__, cudaGetErrorString( err) );\
exit(EXIT_FAILURE); \
} } while (0)
__device__ float *devPtr;
__global__
void kernel1(float *some_neat_data)
{
devPtr = some_neat_data;
}
__global__
void kernel2(void)
{
devPtr[threadIdx.x] *= .3f;
}
int main(int argc, char *argv[])
{
float* otherDevPtr;
cudaMalloc((void**)&otherDevPtr, 256 * sizeof(*otherDevPtr));
cudaMemset(otherDevPtr, 0, 256 * sizeof(*otherDevPtr));
kernel1<<<1,128>>>(otherDevPtr);
CUT_CHECK_ERROR("kernel1");
kernel2<<<1,128>>>();
CUT_CHECK_ERROR("kernel2");
return 0;
}
Дайте ему крутиться.
Другие советы
Потратьте некоторое время на изучение обширной документации, предлагаемой NVIDIA.
Из Руководства по программированию:
float* devPtr;
cudaMalloc((void**)&devPtr, 256 * sizeof(*devPtr));
cudaMemset(devPtr, 0, 256 * sizeof(*devPtr));
Это простой пример того, как распределять память.Теперь в ваших ядрах вы должны принять указатель на число с плавающей запятой следующим образом:
__global__
void kernel1(float *some_neat_data)
{
some_neat_data[threadIdx.x]++;
}
__global__
void kernel2(float *potentially_that_same_neat_data)
{
potentially_that_same_neat_data[threadIdx.x] *= 0.3f;
}
Итак, теперь вы можете вызывать их следующим образом:
float* devPtr;
cudaMalloc((void**)&devPtr, 256 * sizeof(*devPtr));
cudaMemset(devPtr, 0, 256 * sizeof(*devPtr));
kernel1<<<1,128>>>(devPtr);
kernel2<<<1,128>>>(devPtr);
Поскольку эти данные используются в многочисленных функциях, я бы хотел, чтобы они были глобальными.
Есть несколько веских причин использовать глобальные переменные.Это определенно не тот.Я оставлю это в качестве упражнения, чтобы расширить этот пример, включив в него перемещение «devPtr» в глобальную область видимости.
РЕДАКТИРОВАТЬ:
Хорошо, основная проблема заключается в следующем:ваши ядра могут обращаться только к памяти устройства, и единственные указатели глобальной области, которые они могут использовать, — это указатели графического процессора.При вызове ядра из вашего ЦП за кулисами происходит следующее: указатели и примитивы копируются в регистры графического процессора и/или общую память до того, как ядро будет выполнено.
Итак, самое близкое, что я могу предложить, это:используйте cudaMemcpyToSymbol() для достижения своих целей.Но в целом учтите, что другой подход может быть правильным.
#include <algorithm>
__constant__ float devPtr[1024];
__global__
void kernel1(float *some_neat_data)
{
some_neat_data[threadIdx.x] = devPtr[0] * devPtr[1];
}
__global__
void kernel2(float *potentially_that_same_neat_data)
{
potentially_that_same_neat_data[threadIdx.x] *= devPtr[2];
}
int main(int argc, char *argv[])
{
float some_data[256];
for (int i = 0; i < sizeof(some_data) / sizeof(some_data[0]); i++)
{
some_data[i] = i * 2;
}
cudaMemcpyToSymbol(devPtr, some_data, std::min(sizeof(some_data), sizeof(devPtr) ));
float* otherDevPtr;
cudaMalloc((void**)&otherDevPtr, 256 * sizeof(*otherDevPtr));
cudaMemset(otherDevPtr, 0, 256 * sizeof(*otherDevPtr));
kernel1<<<1,128>>>(otherDevPtr);
kernel2<<<1,128>>>(otherDevPtr);
return 0;
}
Не забудьте для этого примера '--host-compilation=c++'.
Я пошел дальше и попробовал решение: выделить временный указатель и передать его простой глобальной функции, похожей на kernel1.
Хорошая новость в том, что это работает :)
Однако я думаю, что это сбивает с толку компилятор, поскольку теперь я получаю сообщение «Рекомендация:Не могу сказать, на что указывает указатель, предполагая пространство глобальной памяти всякий раз, когда я пытаюсь получить доступ к глобальным данным.К счастью, предположение оказывается верным, но предупреждения раздражают.
В любом случае, для справки: я просмотрел множество примеров и пробежался по упражнениям nvidia, целью которых было заставить вывод сказать «Правильно!».Однако я не смотрел все из них.Если кто-нибудь знает пример SDK, в котором они выполняют динамическое глобальное распределение памяти устройства, мне все равно хотелось бы знать.
Хм, именно эта проблема с перемещением devPtr в глобальную область была моей проблемой.
У меня есть реализация, которая делает именно это: два ядра имеют указатель на передаваемые данные.Я явно не хочу передавать эти указатели.
Я довольно внимательно прочитал документацию и заглянул на форумы nvidia (и Google искал около часа), но я не нашел реализации глобального динамического массива устройств, который действительно работает (я пробовал несколько, которые компилируются и затем потерпите неудачу новыми и интересными способами).
ознакомьтесь с образцами, включенными в SDK.Многие из этих примеров проектов — достойный способ учиться на собственном примере.
Поскольку эти данные используются во многих функциях, я бы хотел, чтобы они были глобальными.
-
Есть несколько веских причин использовать глобальные переменные.Это определенно не тот.Я оставлю это в качестве упражнения, чтобы расширить этот пример, включив перемещение «devptr» в глобальную область.
Что, если ядро работает с большой константной структурой, состоящей из массивов?Использование так называемой постоянной памяти не является вариантом, поскольку ее размер очень ограничен.так что тогда вам придется поместить его в глобальную память..?
