source: CIVL/examples/translation/openclversion2/template.cvl@ 764d472

1.23 2.0 acw/focus-triggers main test-branch
Last change on this file since 764d472 was 1c47cb2, checked in by Jacob Trieu <fuufusuu@…>, 12 years ago

Slight changes to next example, planning for 2.1

git-svn-id: svn://vsl.cis.udel.edu/civl/trunk@1198 fb995dde-84ed-4084-dfe6-e5aef3e2452c

  • Property mode set to 100644
File size: 3.4 KB
RevLine 
[d221246c]1#include <stdio.h>
2#include <stdlib.h>
3#include <string.h>
4
[1c47cb2]5//RANDOM IS UNINTERPRETED
[aec3070]6//THIS ASSUMES THE PROGRAMMER HAS ALREADY RAN THE ORIGINAL OPENCL CODE AND IS NOT RESPONSIBLE FOR CHECKING CORRECT USAGE OF METHODS AS OF NOW
7
[d221246c]8
9$input int NUM_DEVICES;
10$input int MAX_NUM_DEVICES;
[aec3070]11$input int CL_DEVICE_MAX_WORK_GROUP_SIZE;
12$input int LOCAL;
13$assume 0 < NUM_DEVICES && NUM_DEVICES < MAX_NUM_DEVICES;
[d221246c]14$gbarrier gbarrier = $gbarrier_create($here, NUM_DEVICES);
15//struct goes here
[aec3070]16
17
[d221246c]18typedef struct
19{
20 int device_id;
[aec3070]21 int workgroup;
22 int global_id;
23 int local_id;
24
25 //kernel variables
26}kernel;
[d221246c]27
28//kernel goes here
29
[aec3070]30void workfunc(size_t local, size_t global, kernel param)
31{
32 for(int i = local * param.workgroup; i < local * param.workgroup + local; i++)
33 {
34 param.local_id = i % local;
35 param.global_id = i;
36 //printf("My workgroup id is %d, my global id is %d, my local id is %d\n", param.workgroup, param.global_id, param.local_id);
37 func(param.workgroup, param.global_id, param.local_id, );
38 }
39}
40
[d221246c]41/*
42Note that the original lines were "__kernel void square( \n" \
[aec3070]43" __global int* input, \n" \
44" __global int* output, \n" \
[d221246c]45" const unsigned int count) \n" \
46
47Any parser must take note of and don't input \n, "", or \ as is
[aec3070]48__global int * input, __global int * output, int count;
[d221246c]49*/
50
51
52int main(int argc, char** argv)
53{
54 //get the number from clGetDeviceIDs 3rd parameter
[aec3070]55 $assert (LOCAL < CL_DEVICE_MAX_WORK_GROUP_SIZE);
[d221246c]56 //variables from __kernel come here
[aec3070]57
[d221246c]58
59 //from the code before
[aec3070]60
61 //handle the definitions being put in different places
62
[d221246c]63
64 //comes from clCreateBuffer
[aec3070]65
[d221246c]66
67 //Possibly keep a list of variables, with a flag for whether they are init or not
68 //Not init, malloc one from what is found in
[aec3070]69 //output = clCreateBuffer(context, CL_MEM_WRITE_ONLY, sizeof(int) * count, NULL, NULL);
70
[d221246c]71
72
73 //came from clEnqueueWriteBuffer rather than the start of code
74
75
76 //Phase after this is the definitions
[aec3070]77 int device_id[NUM_DEVICES];
[d221246c]78 //put device_ids
79 for(int i = 0; i < NUM_DEVICES; i++)
80 {
[aec3070]81 device_id[i] = i;
[d221246c]82 }
83
[aec3070]84 //from clCreateContext, uses the device input
85 //but there may be a loop, take it into account in the next example
86
87 //
88 //"Get" local size from clEnqueueNDRangeKernel, but is really an input
89
90 //Creates an array of the struct according to clEnqueueNDRangeKernel
91 //Have to split array into parts using local and global, and those are a workgroup
92 //For now, assume local is 1, or else inputting the arrays will be odd, for now
93 kernel param[global/local];
94 for(int i = 0; i < global/local; i++)
[d221246c]95 {
[aec3070]96 //Also picks the device to be used
97 param[i].device_id = device_id[0];
98 //other parts of the struct
[d221246c]99 }
100
[aec3070]101
102 //spawns processes according to parameters in clEnqueueNDRangeKernel
103//spawns processes according to parameters in clEnqueueNDRangeKernel
104 $proc procs[global/local];
105 for(int i = 0; i < global/local; i++)
106 {
107 param[i].workgroup = i;
108 procs[i] = $spawn workfunc(local, global, param[i]);
109 }
110
111 for(int i = 0; i < global/local; i++)
[d221246c]112 {
113 $wait(procs[i]);
114 }
115
116 //$barrier barrier = $barrier_create($here, gbarrier, now[i].device_id);
117 //$barrier_call(barrier);
118 //$barrier_destroy(barrier);
119
120 $gbarrier_destroy(gbarrier);
121
122 //use the information from clEnqueueReadBuffer
123 //may have to alter later
124
125 return 0;
126}
Note: See TracBrowser for help on using the repository browser.