OpenCL и Tesla M1060

Question

Я использую Tesla m1060 для вычисления GPGPU. Она имеет following specs:

# of Tesla GPUs 1 
# of Streaming Processor Cores (XXX per processor) 240 
Memory Interface (512-bit per GPU) 512-bit 

Когда я использую OpenCL, можно отобразить следующую информацию о плате:

available platform OpenCL 1.1 CUDA 6.5.14 
device Tesla M1060 type:CL_DEVICE_TYPE_GPU 
max compute units:30 
max work item dimensions:3 
max work item sizes (dim:0):512 
max work item sizes (dim:1):512 
max work item sizes (dim:2):64 
global mem size(bytes):4294770688 local mem size:16383 

Как я могу соотносить информацию GPU карт информаций памяти OpenCL?

Например:

• Что делает средства "Memory Interace"? Связано ли это с рабочим элементом?
• Как я могу связать «240 ядер» графического процессора с рабочими группами/элементами?
• Как я могу сопоставить рабочие группы с ним (каково будет количество рабочих групп)?

Благодаря

EDIT:

После следующих ответов, есть вещь, которая до сих пор неясно мне:

Значение CL_KERNEL_PREFERRED_WORK_GROUP_SIZE_MULTIPLE 32 для ядра я использую.

Однако мое устройство имеет CL_DEVICE_MAX_COMPUTE_UNITS значение 30.

В OpenCL 1.1 Апи, написано (стр. 15):

Compute Unit: An OpenCL device has one or more compute units. A work-group executes on a single compute unit

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

Как отмечалось ранее, когда я установил число рабочих групп до 32, программы выдает следующее сообщение об ошибке:

Entry function uses too much shared data (0x4020 bytes, 0x4000 max).

Значение 16 работ.

Добавление

Вот моя Kernel подпись:

// enable double precision (not enabled by default) 
#ifdef cl_khr_fp64 
    #pragma OPENCL EXTENSION cl_khr_fp64 : enable 
#else 
    #error "IEEE-754 double precision not supported by OpenCL implementation." 
#endif 

#define BLOCK_SIZE 16 // --> this is what defines the WG size to me 

__kernel __attribute__((reqd_work_group_size(BLOCK_SIZE, BLOCK_SIZE, 1))) 
    void mmult(__global double * A, __global double * B, __global double * C, const unsigned int q) 
{ 
    __local double A_sub[BLOCK_SIZE][BLOCK_SIZE]; 
    __local double B_sub[BLOCK_SIZE][BLOCK_SIZE]; 
    // stuff that does matrix multiplication with __local 
} 

В коде принимающей стороны:

#define BLOCK_SIZE 16 
... 
const size_t local_work_size[2] = {BLOCK_SIZE, BLOCK_SIZE}; 
... 
status = clEnqueueNDRangeKernel(command_queue, kernel, 2, NULL, global_work_size, local_work_size, 0, NULL, NULL); 

Answer 1

Интерфейс памяти ничего к приложению OpenCL не значит. Это количество бит, которое контроллер памяти имеет для чтения/записи в память (часть ddr5 в современном gpus). Формула максимальной глобальной скорости памяти примерно равна: pipeWidth * memoryClockSpeed, но поскольку opencl предназначен для кросс-платформенной, вам не нужно будет знать это значение, если вы не попытаетесь определить верхнюю границу производительности памяти. Знание о 512-битном интерфейсе несколько полезно, когда вы имеете дело с объединением памяти.wiki: Coalescing (computer science)

Максимальные размеры рабочего элемента связаны с 1) тем, как аппаратные средства планируют вычисления, и 2) количеством низкоуровневой памяти на устройстве - например. частной памяти и локальной памяти.

Цифра 240 тоже не имеет значения для opencl. Вы можете определить, что каждый из 30 вычислительных блоков состоит из 8 потоковых процессорных ядер для этой архитектуры gpu (потому что 240/30 = 8). Если вы запросите CL_KERNEL_PREFERRED_WORK_GROUP_SIZE_MULTIPLE, это будет очень похоже на кратное 8 для этого устройства. см.: clGetKernelWorkGroupInfo

Я ответил на аналогичные вопросы по размеру рабочей группы. see here, and here

В конечном счете, вам необходимо настроить приложение и ядра на основе собственных результатов настольной маркировки. Я считаю, что стоит потратить время, чтобы написать много тестов с различными размерами рабочих групп и, в конечном итоге, с жестким кодом оптимального размера.

Answer 2

Добавление другого ответа для решения проблемы с локальной памятью.

Entry function uses too much shared data (0x4020 bytes, 0x4000 max)

Поскольку вы распределения A_sub и B_sub, каждый из которых имеет 32 * 32 * SizeOf (двойной), вы бежите из локальной памяти. Устройство должно позволить вам выделять 16kb или 0x4000 байт локальной памяти без проблем.

0x4020 - это 32 байта или 4 удваивается больше, чем позволяет ваше устройство. Есть только две вещи, которые я могу придумать, которые могут вызвать ошибку: 1) может быть ошибка с вашим устройством или драйверами, препятствующая вам выделить полный 16kb, или 2) вы выделяете память где-то еще в вашем ядре.

Чтобы обойти это, вам нужно использовать значение BLOCK_SIZE меньше 32.

Однако есть хорошие новости. Если вы только хотите набрать несколько CL_KERNEL_PREFERRED_WORK_GROUP_SIZE_MULTIPLE в качестве размера рабочей группы, BLOCK_SIZE = 16 уже делает это за вас. (16 * 16 = 256 = 32 * 8). Чтобы лучше воспользоваться локальной памятью, попробуйте BLOCK_SIZE = 24. (576 = 32 * 18)