OpenCL on FPGA — Kernel·Channel·Burst Memory 분석

1
__kernel void vec_add(__global float *a, __global float *b, __global float *c) {
2
    int i = get_global_id(0);
3
    c[i] = a[i] + b[i];
4
}
5
// host: clEnqueueNDRangeKernel(...) global_size = 1024

1
__kernel void vec_add(__global float *a, __global float *b,
2
                      __global float *c, int n) {
3
    for (int i = 0; i < n; i++) {
4
        c[i] = a[i] + b[i];
5
    }
6
}

1
#pragma OPENCL EXTENSION cl_intel_channels : enable
2
channel float c1 __attribute__((depth(1024)));
3
channel float c2 __attribute__((depth(1024)));
4

5
__kernel void producer(__global float *in, int n) {
6
    for (int i = 0; i < n; i++) {
7
        float x = in[i];
8
        write_channel_intel(c1, x);
9
    }
10
}
11

12
__kernel void filter(int n) {
13
    for (int i = 0; i < n; i++) {
14
        float x = read_channel_intel(c1);
15
        write_channel_intel(c2, x * 0.5f);
16
    }
17
}
18

19
__kernel void consumer(__global float *out, int n) {
20
    for (int i = 0; i < n; i++) {
21
        out[i] = read_channel_intel(c2);
22
    }
23
}

1
__kernel void vec_add(__global const float * restrict a,
2
                      __global const float * restrict b,
3
                      __global       float * restrict c,
4
                      int n) {
5
    for (int i = 0; i < n; i++) {
6
        c[i] = a[i] + b[i];
7
    }
8
}

1
#pragma ivdep
2
for (int i = 0; i < n; i++) {
3
    out[i] = process(in[i]);
4
}

1
__kernel void conv2d(__global const float *in, __global float *out,
2
                     int h, int w) {
3
    __local float buf[3][MAX_W];     // BRAM에 매핑
4
    // ... line buffer 패턴
5
}

1
__kernel void mac(__global float *a, __global float *b, float *out) {
2
    float sum = 0;
3
    #pragma unroll
4
    for (int i = 0; i < 8; i++) {
5
        sum += a[i] * b[i];
6
    }
7
    *out = sum;
8
}

1
// NDRange (GPU style)
2
__kernel void vec_add_ndr(__global float *a, __global float *b, __global float *c) {
3
    int i = get_global_id(0);
4
    c[i] = a[i] + b[i];
5
}
6

7
// Single work-item (FPGA style)
8
__kernel void vec_add_swi(__global float *a, __global float *b, __global float *c, int n) {
9
    for (int i = 0; i < n; i++)
10
        c[i] = a[i] + b[i];
11
}

1
#include <CL/cl.h>
2

3
int main() {
4
    cl_platform_id plat;
5
    cl_device_id dev;
6
    clGetPlatformIDs(1, &plat, NULL);
7
    clGetDeviceIDs(plat, CL_DEVICE_TYPE_ACCELERATOR, 1, &dev, NULL);
8

9
    cl_context ctx = clCreateContext(NULL, 1, &dev, NULL, NULL, NULL);
10
    cl_command_queue q = clCreateCommandQueue(ctx, dev, 0, NULL);
11

12
    // Load aocx
13
    FILE *f = fopen("vec_add.aocx", "rb");
14
    fseek(f, 0, SEEK_END);
15
    size_t sz = ftell(f);
16
    fseek(f, 0, SEEK_SET);
17
    uint8_t *bin = malloc(sz);
18
    fread(bin, sz, 1, f);
19
    fclose(f);
20

21
    cl_program prog = clCreateProgramWithBinary(ctx, 1, &dev, &sz,
22
                                                 (const uint8_t**)&bin, NULL, NULL);
23
    clBuildProgram(prog, 1, &dev, "", NULL, NULL);
24
    cl_kernel k = clCreateKernel(prog, "vec_add", NULL);
25

26
    // Buffers
27
    cl_mem ba = clCreateBuffer(ctx, CL_MEM_READ_ONLY,  N*sizeof(float), NULL, NULL);
28
    cl_mem bb = clCreateBuffer(ctx, CL_MEM_READ_ONLY,  N*sizeof(float), NULL, NULL);
29
    cl_mem bc = clCreateBuffer(ctx, CL_MEM_WRITE_ONLY, N*sizeof(float), NULL, NULL);
30

31
    clEnqueueWriteBuffer(q, ba, CL_TRUE, 0, N*sizeof(float), host_a, 0, NULL, NULL);
32
    clEnqueueWriteBuffer(q, bb, CL_TRUE, 0, N*sizeof(float), host_b, 0, NULL, NULL);
33

34
    clSetKernelArg(k, 0, sizeof(cl_mem), &ba);
35
    clSetKernelArg(k, 1, sizeof(cl_mem), &bb);
36
    clSetKernelArg(k, 2, sizeof(cl_mem), &bc);
37
    int n = N;
38
    clSetKernelArg(k, 3, sizeof(int), &n);
39

40
    size_t one = 1;
41
    clEnqueueTask(q, k, 0, NULL, NULL);    /* single work-item */
42
    clFinish(q);
43

44
    clEnqueueReadBuffer(q, bc, CL_TRUE, 0, N*sizeof(float), host_c, 0, NULL, NULL);
45
}

1
#include <sycl/sycl.hpp>
2
#include <sycl/ext/intel/fpga_extensions.hpp>
3

4
using namespace sycl;
5

6
int main() {
7
    queue q{ext::intel::fpga_emulator_selector{}};   // emulator
8
    // or fpga_selector{} for real hardware
9

10
    constexpr int N = 1024;
11
    std::vector<float> a(N, 1.0f), b(N, 2.0f), c(N);
12

13
    {
14
        buffer ba{a}, bb{b}, bc{c};
15
        q.submit([&](handler &h) {
16
            accessor aa{ba, h, read_only};
17
            accessor ab{bb, h, read_only};
18
            accessor ac{bc, h, write_only, no_init};
19
            h.single_task<class vadd>([=]() {
20
                for (int i = 0; i < N; i++)
21
                    ac[i] = aa[i] + ab[i];
22
            });
23
        });
24
    }
25

26
    for (int i = 0; i < 4; i++) std::cout << c[i] << "\n";
27
}

1
__kernel void fir(__global const float *in, __global float *out, int n) {
2
    float buf[8] = {0};
3
    const float coef[8] = {0.1f, 0.2f, 0.3f, 0.4f, 0.4f, 0.3f, 0.2f, 0.1f};
4

5
    for (int i = 0; i < n; i++) {
6
        #pragma unroll
7
        for (int j = 7; j > 0; j--) buf[j] = buf[j-1];
8
        buf[0] = in[i];
9

10
        float sum = 0;
11
        #pragma unroll
12
        for (int j = 0; j < 8; j++) sum += coef[j] * buf[j];
13

14
        out[i] = sum;
15
    }
16
}

1
Throughput:    100 MSPS
2
II:            1
3
Latency:       12 cycles
4
Resource:
5
  ALM:         3500 / 245440  (1.4%)
6
  DSP:         8
7
  M20K:        2
8
  Logic util:  2.1%

1
// host
2
auto kernel = xrt::kernel(device, uuid, "vec_add");
3
auto run = kernel(a_buf, b_buf, c_buf, N);
4
run.wait();

1
__kernel void foo(__global float *a) {
2
    int i = get_global_id(0);
3
    a[i] = a[i] * 2.0f;
4
}
5
// global_size = 1024

1
__kernel void f(__global float *in, __global float *out, int n) {
2
    for (int i = 0; i < n; i++) out[i] = in[i];
3
}
4
// → in과 out이 alias 가능 → cycle마다 1 access

1
channel float c __attribute__((depth(4)));