Task01 Василий Можаев SPbU

[mozhaa]

Локальный вывод

Data generated for n=100000000!
#ifdef __CLION_IDE__
    // Этот include виден только для CLion парсера, это позволяет IDE "знать" ключевые слова вроде __kernel, __global
    // а также уметь подсказывать OpenCL методы, описанные в данном инклюде (такие как get_global_id(...) и get_local_id(...))
    #include "clion_defines.cl"
#endif

#line 8// Седьмая строчка теперь восьмая (при ошибках компиляции в логе компиляции будут указаны корректные строчки благодаря этой директиве)

// TODO 5 реализуйте кернел:
// - От обычной функции кернел отличается модификатором __kernel и тем, что возвращаемый тип всегда void
// - На вход дано три массива float чисел; единственное, чем они отличаются от обычных указателей - модификатором __global, т.к. это глобальная память устройства (видеопамять)
// - Четвертым и последним аргументом должно быть передано количество элементов в каждом массиве (unsigned int, главное, чтобы тип был согласован с типом в соответствующем clSetKernelArg в T0D0 10)

__kernel void aplusb(__global const float as[], __global const float bs[], __global float cs[], unsigned int n) {
    // Узнать, какой workItem выполняется в этом потоке поможет функция get_global_id
    // см. в документации https://www.khronos.org/registry/OpenCL/sdk/1.2/docs/man/xhtml/
    // OpenCL Compiler -> Built-in Functions -> Work-Item Functions

    // P.S. В общем случае количество элементов для сложения может быть некратно размеру WorkGroup, тогда размер рабочего пространства округлен вверх от числа элементов до кратности на размер WorkGroup
    // и в таком случае, если сделать обращение к массиву просто по индексу=get_global_id(0), будет undefined behaviour (вплоть до повисания ОС)
    // поэтому нужно либо дополнить массив данных длиной до кратности размеру рабочей группы,
    // либо сделать return в кернеле до обращения к данным в тех WorkItems, где get_global_id(0) выходит за границы данных (явной проверкой)

    int id = get_global_id(0);
    if (id >= n)
        return;
    cs[id] = as[id] + bs[id];
}
Log:


Kernel average time: 0.0344513+-3.63146e-06 s
GFlops: 2.90265
VRAM bandwidth: 32.4397 GB/s
Result data transfer time: 0.11914+-0.000158638 s
VRAM -> RAM bandwidth: 3.12681 GB/s

Вывод Github CI

Data generated for n=100000000!
#ifdef __CLION_IDE__
    // Этот include виден только для CLion парсера, это позволяет IDE "знать" ключевые слова вроде __kernel, __global
    // а также уметь подсказывать OpenCL методы, описанные в данном инклюде (такие как get_global_id(...) и get_local_id(...))
    #include "clion_defines.cl"
#endif

#line 8// Седьмая строчка теперь восьмая (при ошибках компиляции в логе компиляции будут указаны корректные строчки благодаря этой директиве)

// TODO 5 реализуйте кернел:
// - От обычной функции кернел отличается модификатором __kernel и тем, что возвращаемый тип всегда void
// - На вход дано три массива float чисел; единственное, чем они отличаются от обычных указателей - модификатором __global, т.к. это глобальная память устройства (видеопамять)
// - Четвертым и последним аргументом должно быть передано количество элементов в каждом массиве (unsigned int, главное, чтобы тип был согласован с типом в соответствующем clSetKernelArg в T0D0 10)

__kernel void aplusb(__global const float as[], __global const float bs[], __global float cs[], unsigned int n) {
    // Узнать, какой workItem выполняется в этом потоке поможет функция get_global_id
    // см. в документации https://www.khronos.org/registry/OpenCL/sdk/1.2/docs/man/xhtml/
    // OpenCL Compiler -> Built-in Functions -> Work-Item Functions

    // P.S. В общем случае количество элементов для сложения может быть некратно размеру WorkGroup, тогда размер рабочего пространства округлен вверх от числа элементов до кратности на размер WorkGroup
    // и в таком случае, если сделать обращение к массиву просто по индексу=get_global_id(0), будет undefined behaviour (вплоть до повисания ОС)
    // поэтому нужно либо дополнить массив данных длиной до кратности размеру рабочей группы,
    // либо сделать return в кернеле до обращения к данным в тех WorkItems, где get_global_id(0) выходит за границы данных (явной проверкой)

    int id = get_global_id(0);
    if (id >= n)
        return;
    cs[id] = as[id] + bs[id];
}
Log:
Compilation started
Compilation done
Linking started
Linking done
Device build started
Device build done
Kernel  was successfully vectorized (8)
Done.
Kernel average time: 0.0313088+-0.000122142 s
GFlops: 3.194
VRAM bandwidth: 35.6957 GB/s
Result data transfer time: 6.25e-06+-7.21688e-07 s
VRAM -> RAM bandwidth: 59604.6 GB/s

[mozhaa]

Обновил вывод github CI

[simiyutin]

Все хорошо, задача зачтена, 5/5 баллов 👍