Тест OpenCL - советы по изменяющимся параметрам

Question

Я хотел бы выполнить контрольный показатель времени выполнения двухэтапного сокращения суммы с помощью OpenCL (от этого AMD link) на Radeon HD 7970 Tahiti XT.

Первоначально я взял первую версию кода, в которой я не использовал первый цикл, который выполняет сокращение от входного массива размером N до выходного массива размером NworkItems. Вот этот первый цикл в код ядра:

int global_index = get_global_id(0); 
    float accumulator = 0; 
    // Loop sequentially over chunks of input vector 
    while (global_index < length) { 
    float element = buffer[global_index]; 
    accumulator += element; 
    global_index += get_global_size(0); 
    } 

Так что с этой первой версии, я измерил время работы в зависимости от размера входного массива (которое равно общему числу нитей) и для различных размеров работы группа. Вот результаты:

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

От this link говорится, что AMD рекомендует несколько 64 для размера рабочей группы (32 для NVIDIA).

Кроме того, из последнего комментария на this other link, рекомендуется установить размер рабочей группы как: WorkGroup size = (Number of total threads)/(Compute Units). На моей плате GPU у меня есть 32 вычислительных устройства.

Поэтому я хотел бы получить советы, зная, какие параметры будут интересны для сравнения, чтобы сравнить время автономной работы в этой второй версии (с первым циклом восстановления). Например, я могу принимать различные значения для отношения (N size of input array)/(total NworkItems) и фиксированного значения для WorkGroup size (см выражения выше),

или же, наоборот, т.е. я должен изменять значение для WorkGroup size и зафиксировать отношение (N size of input array)/(total NworkItems)?

Любая идея приветствуется, спасибо

Answer 1

Вы должны суммировать данные локальных данных вместо рассредоточены, чтобы помочь передаче памяти (соединялись доступа к данным). Поэтому используйте это вместо:

int chunk_size = length/get_global_size(0)+(length%get_global_size(0) > 0); //Will give how many items each work item needs to process 
    int global_index = get_group_id(0)*get_local_size(0)*chunk_size + get_local_id(0); //Start at this address for this work item 
    float accumulator = 0; 

    for(int i=0; i<chunk_size; i++) 
    // Loop sequentially over chunks of input vector 
    if (global_index < length) { 
     float element = buffer[global_index]; 
     accumulator += element; 
     global_index += get_local_size(0); 
    } 
    } 

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