FIXME: Report correct return value - Needs change in KPRINTF
| 218 | // FIXME: Report correct return value - Needs change in KPRINTF |
| 219 | // |
| 220 | static ssize_t |
| 221 | generator( |
| 222 | char *buf, |
| 223 | size_t buflen, |
| 224 | const struct SubproblemDim *subdims, |
| 225 | const struct PGranularity *pgran, |
| 226 | void *extra) |
| 227 | { |
| 228 | DUMMY_ARG_USAGE(subdims); |
| 229 | size_t BLOCKSIZE = pgran->wgSize[0]; |
| 230 | char tempTemplate[32*1024]; |
| 231 | |
| 232 | if ( buf == NULL) // return buffer size |
| 233 | { |
| 234 | buflen = (32 * 1024 * sizeof(char)); |
| 235 | return (ssize_t)buflen; |
| 236 | } |
| 237 | CLBLASKernExtra *extraFlags = ( CLBLASKernExtra *)extra; |
| 238 | SolutionStep *step = container_of( pgran , pgran, SolutionStep); |
| 239 | CLBlasKargs* kargs = (CLBlasKargs*) &(step->args); |
| 240 | char const *kernName; |
| 241 | |
| 242 | if(kargs->redctnType == REDUCE_BY_SUM) { |
| 243 | kernName = red_sum_kernel; |
| 244 | } else if(kargs->redctnType == REDUCE_BY_MAX) { |
| 245 | kernName = red_max_kernel; |
| 246 | } else if(kargs->redctnType == REDUCE_BY_MIN) { |
| 247 | kernName = red_min_kernel; |
| 248 | } else if(kargs->redctnType == REDUCE_MAX_WITH_INDEX) { |
| 249 | kernName = red_with_index_kernel; |
| 250 | } else if(kargs->redctnType == REDUCE_BY_HYPOT) { |
| 251 | kernName = red_hypot_kernel; |
| 252 | } else if(kargs->redctnType == REDUCE_BY_SSQ) { |
| 253 | kernName = red_ssq_kernel; |
| 254 | } |
| 255 | |
| 256 | #ifdef DEBUG_REDUCTION |
| 257 | printf("REDUCTION GENERATOR called....\n"); |
| 258 | printf("dataType : %c\n", Prefix[extraFlags->dtype]); |
| 259 | printf("Vector length used : %d\n\n", vecLenA); |
| 260 | #endif |
| 261 | |
| 262 | unsigned int vecLenA = extraFlags->vecLenA; |
| 263 | bool doVLOAD = false; |
| 264 | if( extraFlags->flags & KEXTRA_NO_COPY_VEC_A ) |
| 265 | { |
| 266 | doVLOAD = true; |
| 267 | } |
| 268 | strcpy( tempTemplate, kernName ); |
| 269 | kprintf kobj( Prefix[extraFlags->dtype], vecLenA, doVLOAD, doVLOAD, BLOCKSIZE); |
| 270 | kobj.spit((char*)buf, tempTemplate); |
| 271 | |
| 272 | return (32 * 1024 * sizeof(char)); |
| 273 | } |
| 274 | |
| 275 | /* |
| 276 | __kernel void %PREFIXred_sum_kernel( __global %TYPE *_X, __global %TYPE *_res, |