2 * This program is free software: you can redistribute it and/or modify
3 * it under the terms of the GNU General Public License as published by
4 * the Free Software Foundation, either version 3 of the License, or
5 * (at your option) any later version.
7 * This program is distributed in the hope that it will be useful,
8 * but WITHOUT ANY WARRANTY; without even the implied warranty of
9 * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the
10 * GNU General Public License for more details.
12 * You should have received a copy of the GNU General Public License
13 * along with this program. If not, see <http://www.gnu.org/licenses/>.
15 * Copyright (C) 2010 Pavel Herrmann (morpheus.ibis@gmail.com)
27 #define L(cstbuf) sizeof(cstbuf)-1
28 #define scpy(target,source) strncpy(target+target##idx,source, L(source));target##idx+=L(source)
31 char* createShadeKernel(int mtlnum
, struct mtl_def_s
* materials
, char** mtllist
)
33 const char types
[]="#include \"types.h\"\n#include \"cl/algs.h\"\n#include \"mat/background.cl\"\n";
34 const char inc
[]="#include \"mat/";
35 const char sep
[]=".cl\"\n";
36 const char shade_head
[]="__kernel void shade_kernel(__global struct hit_s * hits, __const int raycount,\n"
37 "__global const int * rayindex, __global struct ray_s * rays, __global struct pixel_s * image,\n"
38 "__global const struct light_sample_s * light_samples, __global const struct light_s * lights, __const int lightcount,\n"
39 "__global const struct triangle_s * triangles, __global const struct accel_s * accel";
40 const char count_head
[]="__kernel void shade_count_kernel(__global struct hit_s * hits, __const int raycount, __global int * rayindex,\n"
41 "__global int * newcount, __global const struct triangle_s * triangles";
42 const char prepare
[]="__private int hitnum = get_global_id(0)+get_global_id(1)*get_global_size(0);\n"
43 "if (hitnum >= raycount) return;\n"
44 "__private struct hit_s hit = hits[hitnum];\n";
45 const char shade_bkgnd
[]="if (islive(&(hit.reconstructInfo))==0) return;\nif(hit.triIndex<0){\n"
46 "background(&hit, image);\n}else{\n";
47 const char count_bkgnd
[]="if (islive(&(hit.reconstructInfo))==0) hit.triIndex=-1;\nif(hit.triIndex>=0){\n";
48 const char prepare_valid
[]="__global const struct triangle_s * triangle = &(triangles[hit.triIndex]);\n"
49 "__private int materialindex = triangle->objectIndex;\n";
50 const char count_begin
[]="__private int cnt=0;\n";
51 const char begin
[]="switch(materialindex) {\n";
52 const char params
[]="(&hit, triangle";
53 const char shade_params
[]=", &(rays[rayindex[hitnum]]), image, &(light_samples[hitnum*lightcount]), lights, lightcount";
54 const char count_work
[]="cnt=";
55 const char ret
[]="return;\n";
56 const char brk
[]="break;\n";
57 const char block_begin
[]=")\n{\n";
58 const char shade_end
[]="}}\n}\n";
59 const char count_postfix
[]="_count";
60 const char count_end
[]="}}\n__local volatile int lcnt;\n__local int base;\n__private int offset=0;\nlcnt=0;\n"
61 "barrier(CLK_LOCAL_MEM_FENCE);\noffset=atomic_add(&lcnt,cnt);\nbarrier(CLK_LOCAL_MEM_FENCE);\n"
62 "if ((get_local_id(0)==0)&&(get_local_id(1)==0))\n base=atomic_add(newcount,lcnt);\nbarrier(CLK_LOCAL_MEM_FENCE);\n"
63 "rayindex[hitnum]=base+offset;\n}\n";
64 const char work_end
[]=");\n";
66 const char csf
[]="case %3d:\n"; //10 chars
67 const char txf
[]=", __read_only image%dd_t t%03d"; //28 chars
68 const char tf
[]=", t%03d"; //6 chars
85 shadelen
= L(shade_head
) + L(block_begin
) + L(prepare
) + L(shade_bkgnd
) + L(prepare_valid
) + L(begin
) + L(shade_end
);
86 countlen
= L(count_head
) + L(block_begin
) + L(prepare
) + L(count_bkgnd
) + L(prepare_valid
) + L(count_begin
) + L(begin
) + L(count_end
);
87 for (int i
=0;i
<mtlnum
;i
++)
90 headlen
+=L(inc
) + t
+ L(sep
);
91 countlen
+=10/*L(csf)*/ + L(count_work
) + t
+ L(count_postfix
) + L(params
) + L(work_end
) + L(brk
);
92 shadelen
+=10/*L(csf)*/ + t
+ L(params
) + L(shade_params
) + L(work_end
) + L(ret
);
93 for (int j
=0;j
<materials
[i
].texcount
;j
++)
96 if(texcount
<materials
[i
].texidx
[j
]) texcount
=materials
[i
].texidx
[j
];
99 shadelen
+=28/*L(txf)*/ * (texcount
+1);
101 head
= malloc(sizeof(char)* (headlen
+1));
102 shade
= malloc(sizeof(char)* (shadelen
+1));
103 count
= malloc(sizeof(char)* (countlen
+1));
115 scpy(count
,count_head
);
116 scpy(count
,block_begin
);
117 scpy(count
,count_begin
);
119 scpy(count
,count_bkgnd
);
120 scpy(count
,prepare_valid
);
123 scpy(shade
,shade_head
);
125 for (int i
=0;i
<texcount
+1;i
++)
127 sprintf(shade
+shadeidx
,txf
,2,i
);
130 scpy(shade
,block_begin
);
132 scpy(shade
,shade_bkgnd
);
133 scpy(shade
,prepare_valid
);
137 for (int i
=0;i
<mtlnum
;i
++)
139 t
=strlen(mtllist
[i
]);
141 strncpy(head
+headidx
,mtllist
[i
],t
);
145 sprintf(count
+countidx
,csf
,i
);
147 scpy(count
,count_work
);
148 strncpy(count
+countidx
,mtllist
[i
],t
);
150 scpy(count
,count_postfix
);
152 scpy(count
,work_end
);
155 sprintf(shade
+shadeidx
,csf
,i
);
157 strncpy(shade
+shadeidx
,mtllist
[i
],t
);
160 scpy(shade
,shade_params
);
162 for (int j
=0;j
<materials
[i
].texcount
;j
++)
164 sprintf(shade
+shadeidx
,tf
,materials
[i
].texidx
[j
]);
167 scpy(shade
,work_end
);
171 scpy(count
,count_end
);
172 scpy(shade
,shade_end
);
175 retval
= malloc(sizeof(char)*(headlen
+shadelen
+countlen
+1));
176 strncpy(retval
,head
,headlen
);
177 strncpy(retval
+headlen
,shade
,shadelen
);
178 strncpy(retval
+headlen
+shadelen
,count
,countlen
);
179 retval
[headlen
+shadelen
+countlen
]=0;
181 //printf("%s\n",retval);
186 int loadTextures(int mtlnum
, struct texinfo_s
** textures
, struct mtl_def_s
* materials
)
195 for (int i
=0;i
<mtlnum
;i
++)
197 materials
[i
].texidx
= malloc(sizeof(int)*materials
[i
].texcount
);
198 for (int j
=0;j
<materials
[i
].texcount
;j
++)
199 materials
[i
].texidx
[j
]=texnum
+j
;
200 texnum
+=materials
[i
].texcount
;
202 *textures
= malloc(sizeof(struct texinfo_s
)*texnum
);
204 for (int i
=0;i
<mtlnum
;i
++)
205 for (int j
=0;j
<materials
[i
].texcount
;j
++)
207 ilLoadImage(materials
[i
].texnames
[j
]);
209 (*textures
)[texnum
].bpp
=ilGetInteger(IL_IMAGE_BITS_PER_PIXEL
);
210 (*textures
)[texnum
].dimensions
.x
=ilGetInteger(IL_IMAGE_WIDTH
);
211 (*textures
)[texnum
].dimensions
.y
=ilGetInteger(IL_IMAGE_HEIGHT
);
212 (*textures
)[texnum
].dimensions
.z
=ilGetInteger(IL_IMAGE_DEPTH
);
213 imglength
=((*textures
)[texnum
].dimensions
.x
* (*textures
)[texnum
].dimensions
.y
* (*textures
)[texnum
].dimensions
.z
* 4);
214 (*textures
)[texnum
].data
=malloc(sizeof(unsigned char)*imglength
);
216 switch (ilGetInteger(IL_IMAGE_BITS_PER_PIXEL
))
219 for (int k
=0;k
<(imglength
/4);k
++)
221 (*textures
)[texnum
].data
[k
*4+0] = imgdata
[k
*3+0];
222 (*textures
)[texnum
].data
[k
*4+1] = imgdata
[k
*3+1];
223 (*textures
)[texnum
].data
[k
*4+2] = imgdata
[k
*3+2];
224 (*textures
)[texnum
].data
[k
*4+3] = 255;
228 memcpy((*textures
)[texnum
].data
,imgdata
,imglength
);
233 (*textures
)[texnum
].bpp
=32;