FIXME: Report correct return value - Needs change in KPRINTF
| 203 | // FIXME: Report correct return value - Needs change in KPRINTF |
| 204 | // |
| 205 | static ssize_t |
| 206 | generator( |
| 207 | char *buf, |
| 208 | size_t buflen, |
| 209 | const struct SubproblemDim *subdims, |
| 210 | const struct PGranularity *pgran, |
| 211 | void *extra) |
| 212 | { |
| 213 | |
| 214 | DUMMY_ARGS_USAGE_2(pgran, subdims); |
| 215 | char tempTemplate[32*1024]; |
| 216 | |
| 217 | if ( buf == NULL) // return buffer size |
| 218 | { |
| 219 | buflen = (32 * 1024 * sizeof(char)); |
| 220 | return (ssize_t)buflen; |
| 221 | } |
| 222 | CLBLASKernExtra *extraFlags = ( CLBLASKernExtra *)extra; |
| 223 | |
| 224 | #ifdef DEBUG_AXPY |
| 225 | printf("AXPY GENERATOR called....\n"); |
| 226 | printf("dataType : %c\n", Prefix[extraFlags->dtype]); |
| 227 | #endif |
| 228 | |
| 229 | unsigned int vecLenA = extraFlags->vecLenA; |
| 230 | |
| 231 | #ifdef DEBUG_AXPY |
| 232 | printf("Vector length used : %d\n\n", vecLenA); |
| 233 | #endif |
| 234 | |
| 235 | bool doVLOAD = false; |
| 236 | if( extraFlags->flags & KEXTRA_NO_COPY_VEC_A ) |
| 237 | { |
| 238 | doVLOAD = true; |
| 239 | #ifdef DEBUG_AXPY |
| 240 | printf("DOing VLOAD as Aligned Data Pointer not Availabe\n"); |
| 241 | #endif |
| 242 | } |
| 243 | else |
| 244 | { |
| 245 | #ifdef DEBUG_AXPY |
| 246 | printf("Using Aligned Data Pointer .......\n"); |
| 247 | #endif |
| 248 | } |
| 249 | strcpy( tempTemplate, (char*)axpy_kernel ); |
| 250 | kprintf kobj( Prefix[extraFlags->dtype], vecLenA, doVLOAD, doVLOAD); |
| 251 | kobj.spit((char*)buf, tempTemplate); |
| 252 | |
| 253 | return (32 * 1024 * sizeof(char)); |
| 254 | } |
| 255 | |
| 256 | /* |
| 257 | __kernel void %PREFIXaxpy_kernel( %TYPE _alpha, __global %TYPE *_X, __global %TYPE *_Y, |