1. 1
2. 2
3. 3
4. 4
5. 5
6. 6
7. 7
8. 8
9. 9
10. 10
11. 11
12. 12
13. 13
14. 14
15. 15
16. 16
17. 17
18. 18
19. 19
20. 20
21. 21
22. 22
23. 23
24. 24
25. 25
26. 26
27. 27
28. 28
29. 29
30. 30
31. 31
32. 32
33. 33
34. 34
35. 35
36. 36
37. 37
38. 38
39. 39
40. 40
41. 41
42. 42
43. 43
44. 44
45. 45
46. 46
47. 47
48. 48
49. 49
50. 50
51. 51
52. 52
53. 53
54. 54
55. 55
56. 56
57. 57
58. 58
59. 59
60. 60
61. 61
62. 62
63. 63
64. 64
65. 65
66. 66
67. 67
68. 68
69. 69
70. 70
71. 71
72. 72
73. 73
74. 74
75. 75
76. 76
77. 77
78. 78
79. 79
80. 80
81. 81
82. 82
83. 83
84. 84
85. 85
86. 86
87. 87
88. 88
89. 89
90. 90
91. 91
92. 92

let src1 = r#"
        __kernel void add1(__global float* A, __global float* BBB, __global float* B, int m, int n) 
        {
            __local float Blo[64];
            int x = get_local_id(0);
            int y = get_local_id(1);
            int i = get_global_id(0);
            int j = get_global_id(1);
            int k = get_global_id(2);
            i += k / 8;
            j += k % 8;

            if (i >= n || j >= m) return;


            Blo[x * 8 + y] = A[i * m + j];




            barrier(CLK_LOCAL_MEM_FENCE);

            float BB = 0;

            for (int xx = 0; xx < 8; ++xx)
                for (int yy = 0; yy < 8; ++yy)
                {
    
                        float c = (2 * xx + 1) * x * 3.1415926535 / 16;
                        float cc = (2 * yy + 1) * y * 3.1415926535 / 16;
                        c = cos(c);
                        cc = cos(cc);
                        BB += Blo[xx * 8 + yy] * c * cc;

                }

            float Ci, Cj;
            if (x == 0)
                Ci = 1 / 1.4142135623;
            else
                Ci = 1;

            if (y == 0)
                Cj = 1 / 1.4142135623;
            else
                Cj = 1;
            B[k * m * n + i * m + j] = Ci * Cj / 4 * BB;

            barrier(CLK_LOCAL_MEM_FENCE);

            i = get_global_id(0);
            j = get_global_id(1);

            float summ = 0;
            for (int ii = 0; ii < 64; ++ii)
                summ += B[ii * m * n + i * m + j];
            BBB[i * m + j] = summ / 64;
            
        }
    "#;


let pro_que = ProQue::builder().src(src1).dims((hi, wi, 64)).build().unwrap();


   let matr11 = Buffer::builder()
        .queue(pro_que.queue().clone())
        .flags(MemFlags::new().read_only().use_host_ptr())
        .dims((hi, wi))
        .host_data(&Resr)
        .build().unwrap();

        let matg11 = Buffer::builder()
        .queue(pro_que.queue().clone())
        .flags(MemFlags::new().read_only().use_host_ptr())
        .dims((hi, wi))
        .host_data(&Resg)
        .build().unwrap();

...

    let mut kernel;
    {
            let wi = wi as i32;
            let hi = hi as i32;
            kernel = pro_que.create_kernel("add1").unwrap().arg_buf(&matr11).arg_buf(&resr11).arg_buf(&bor1).arg_scl(wi).arg_scl(hi);
            kernel.lws((8, 8)).enq().unwrap();
            kernel = pro_que.create_kernel("add1").unwrap().arg_buf(&matg11).arg_buf(&resg11).arg_buf(&bog1).arg_scl(wi).arg_scl(hi);
            kernel.lws((8, 8)).enq().unwrap();
            kernel = pro_que.create_kernel("add1").unwrap().arg_buf(&matb11).arg_buf(&resb11).arg_buf(&bob1).arg_scl(wi).arg_scl(hi);
            kernel.lws((8, 8)).enq().unwrap();
    }