На этом шаге мы рассмотрим глобальную память.
Основная часть DRAM GPU доступна приложениям именно как глобальная память. Доступный ее объем фактически определяется объемом установленной DRAM-памяти на устройстве. Глобальная память является основным местом для размещения и хранения большого объема данных для обработки ядрами, она сохраняет свои значения между вызовами ядер, что позволяет через нее передавать данные между ядрами.
Функции cudaMalloc и cudaFree являются стандартными функциями выделения и освобождения глобальной памяти. Поскольку для эффективного доступа к глобальной памяти важным требованием является выравнивание данных в памяти, то для выделения памяти под двумерные массивы лучше использовать функцию cudaMallocPitch, которая при выделении памяти может увеличить объем памяти под каждую строку, чтобы гарантировать выравнивание всех строк. При этом дополнительный объем памяти в байтах, служащий для выравнивания строки, возвращается через параметр pitch.
Приведем ниже функцию копирования данных.
cudaError_t cudaMemcpy(void *dst, const void *src, size_t size, enum cudaMemcpyKind kind); cudaFree(void *dst);
Параметр dst задает место в памяти, в которую копируются данные.
Параметр src задает место в памяти, из которой копируются данные.
Параметр size задает сколько байтов памяти, на которую указывает src, нужно скопировать в dst.
Параметр kind задает направление копирования и может принимать только одно из следующих значений:
- cudaMemcpyHostToHost - скопировать из памяти CPU в память CPU;
- cudaMemcpyHostToDevice - скопировать из памяти CPU в память GPU;
- cudaMemcpyDeviceToHost - скопировать из памяти GPU в память CPU;
- cudaMemcpyDeviceToDevice - скопировать из памяти GPU в память GPU.
Обратите внимание, что хотя функции выделения глобальной памяти и возвращают указатель на память, но это указатель на память GPU, поэтому доступ по данному указателю может осуществлять только код, выполняемый на GPU.
Обратим внимание, что копирование памяти между CPU и GPU является очень дорогостоящей операцией, и число подобных операций должно быть минимизировано. Это связано с тем, что скорость передачи данных ограничивается скоростью передачи данных по шине PCI Express.
В то же время передача данных в пределах DRAM оказывается намного выше. Например, у GPU GeForce GTX 280 она составляет 141 Гбайт/сек.
Рассмотрим простейший пример - построение таблицы значений заданной функции с заданным шагом.
Для этого необходимо выделить два массива одинакового размера для хранения результата: один - в памяти CPU, другой - в памяти GPU. После этого запускается ядро, заполняющее массив в глобальной памяти заданными значениями. После завершении ядра необходимо скопировать результаты вычислений из памяти GPU в память CPU и освободить выделенную глобальную память.
Приведем ниже часть программы, в которой используется технология CUDA.
__global__ void tableKernel(float *devPtr, float step){ int index=blockIdx.x*blockDim.x+threadIdx.x; float x=step*index; devPtr[index]=__sinf(__sqrtf(x)); } void buildTable(float *res, int n, float step){ float *devPtr; assert(n%256==0); cudaMalloc(&devPtr, n*sizeof(float)); tableKernel<<<dim3(n/256),dim3(256)>>>(devPtr,step); cudaMemcpy(res,devPtr,n*sizeof(float),cudaMemcpyDeviceToHost); cudaFree(devPtr); }
На следующем шаге мы рассмотрим оптимизацию использования глобальной памяти.
