1 /*
2 * Copyright (c) 2024, Oracle and/or its affiliates. All rights reserved.
3 * DO NOT ALTER OR REMOVE COPYRIGHT NOTICES OR THIS FILE HEADER.
4 *
5 * This code is free software; you can redistribute it and/or modify it
6 * under the terms of the GNU General Public License version 2 only, as
7 * published by the Free Software Foundation. Oracle designates this
8 * particular file as subject to the "Classpath" exception as provided
9 * by Oracle in the LICENSE file that accompanied this code.
10 *
11 * This code is distributed in the hope that it will be useful, but WITHOUT
12 * ANY WARRANTY; without even the implied warranty of MERCHANTABILITY or
13 * FITNESS FOR A PARTICULAR PURPOSE. See the GNU General Public License
14 * version 2 for more details (a copy is included in the LICENSE file that
15 * accompanied this code).
16 *
17 * You should have received a copy of the GNU General Public License version
18 * 2 along with this work; if not, write to the Free Software Foundation,
19 * Inc., 51 Franklin St, Fifth Floor, Boston, MA 02110-1301 USA.
20 *
21 * Please contact Oracle, 500 Oracle Parkway, Redwood Shores, CA 94065 USA
22 * or visit www.oracle.com if you need additional information or have any
23 * questions.
24 */
25
26 #include "opencl_backend.h"
27 class KernelContextWithBufferState {
28 public:
29 int x;
30 int maxX;
31 BufferState bufferState;
32 };
33 struct ArgArray_2 {
34 int argc;
35 u8_t pad12[12];
36 KernelArg argv[2];
37 };
38
39
40 struct S32Array1024WithBufferState {
41 int length;
42 int array[1024];
43 BufferState bufferState;
44 };
45 int main(int argc, char **argv) {
46 OpenCLBackend openclBackend(0
47 | Backend::Config::INFO_BIT
48 | Backend::Config::Config::TRACE_CALLS_BIT
49 | Backend::Config::Config::TRACE_COPIES_BIT
50 );
51
52 //std::string cudaPath = "/home/gfrost/github/grfrost/babylon-grfrost-fork/hat/squares.cuda";
53 OpenCLSource openclSource((char *) R"(
54 #define NDRANGE_OPENCL
55 #pragma OPENCL EXTENSION cl_khr_global_int32_base_atomics : enable
56 #pragma OPENCL EXTENSION cl_khr_local_int32_base_atomics : enable
57 #ifndef NULL
58 #define NULL 0
59 #endif
60 #pragma OPENCL EXTENSION cl_khr_global_int32_base_atomics : enable
61 #pragma OPENCL EXTENSION cl_khr_local_int32_base_atomics : enable
62 typedef char s8_t;
63 typedef char byte;
64 typedef char boolean;
65 typedef unsigned char u8_t;
66 typedef short s16_t;
67 typedef unsigned short u16_t;
68 typedef unsigned int u32_t;
69 typedef int s32_t;
70 typedef float f32_t;
71 typedef long s64_t;
72 typedef unsigned long u64_t;
73 typedef struct KernelContext_s{
74 int x;
75 int maxX;
76 }KernelContext_t;
77 typedef struct S32Array_s{
78 int length;
79 int array[1];
80 }S32Array_t;
81
82
83
84 inline int squareit(
85 int v
86 ){
87 return v*v;
88 }
89
90
91 __kernel void squareKernel(
92 __global KernelContext_t *global_kc, __global S32Array_t* s32Array
93 ){
94 KernelContext_t mine;
95 KernelContext_t* kc=&mine;
96 kc->x=get_global_id(0);
97 kc->maxX=global_kc->maxX;
98 if(kc->x<kc->maxX){
99 int value = s32Array->array[(long)kc->x];
100 s32Array->array[(long)kc->x]=squareit(value);
101 }
102 return;
103 }
104 )");
105
106 auto *program =openclBackend.compileProgram(openclSource);
107 int maxX = 32;
108 auto *kernelContextWithBufferState = bufferOf<KernelContextWithBufferState>("kernelContext");
109 kernelContextWithBufferState->x=0;
110 kernelContextWithBufferState->maxX=maxX;
111
112 auto *s32Array1024WithBufferState = bufferOf<S32Array1024WithBufferState>("s32ArrayX1024");
113
114 s32Array1024WithBufferState->length=maxX;
115
116 for (int i=0; i < s32Array1024WithBufferState->length; i++){
117 s32Array1024WithBufferState->array[i]=i;
118 }
119
120 ArgArray_2 args2Array{.argc = 2, .argv={
121 {.idx = 0, .variant = '&',.value = {.buffer ={.memorySegment = (void *) kernelContextWithBufferState, .sizeInBytes = sizeof(KernelContextWithBufferState), .access = RO_BYTE}}},
122 {.idx = 1, .variant = '&',.value = {.buffer ={.memorySegment = (void *) s32Array1024WithBufferState, .sizeInBytes = sizeof(S32Array1024WithBufferState), .access = RW_BYTE}}}
123 }};
124 auto kernel = program->getOpenCLKernel((char*)"squareKernel");
125 kernel->ndrange( reinterpret_cast<ArgArray_s *>(&args2Array));
126 for (int i=0; i < s32Array1024WithBufferState->length; i++){
127 std::cout << i << " array[" << i << "]=" << s32Array1024WithBufferState->array[i] << std::endl;
128 }
129 }
130