Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Task01 Даниил Ушков ITMO #13

Closed
wants to merge 7 commits into from

Conversation

daniil-ushkov
Copy link

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

$ ./build/aplusb 
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 float *as, __global 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
    size_t id = get_global_id(0);

    // P.S. В общем случае количество элементов для сложения может быть некратно размеру WorkGroup, тогда размер рабочего пространства округлен вверх от числа элементов до кратности на размер WorkGroup
    // и в таком случае, если сделать обращение к массиву просто по индексу=get_global_id(0), будет undefined behaviour (вплоть до повисания ОС)
    // поэтому нужно либо дополнить массив данных длиной до кратности размеру рабочей группы,
    // либо сделать return в кернеле до обращения к данным в тех WorkItems, где 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.0340718+-0.000353536 s
GFlops: 2.93498
VRAM bandwidth: 32.8009 GB/s
Result data transfer time: 0.0210688+-0.000253849 s
VRAM -> RAM bandwidth: 17.6816 GB/s

Вывод Github CI

TODO

@simiyutin
Copy link
Collaborator

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

@simiyutin simiyutin closed this Sep 20, 2024
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Labels
None yet
Projects
None yet
Development

Successfully merging this pull request may close these issues.

2 participants