> -----Original Message----- > From: Beignet [mailto:[email protected]] On Behalf Of > Guo Yejun > Sent: Wednesday, December 16, 2015 9:00 > To: [email protected] > Cc: Guo, Yejun > Subject: [Beignet] [PATCH V2] output warning message if the > global/local_work_size is not good in Debug mode > > the known issue is that utest will output lots of warning messages since the > total number of work-items is less than 64. > > V2: use DEBUGP to wrap the code, and also refine the macro > Signed-off-by: Guo Yejun <[email protected]> > --- > src/cl_api.c | 11 ++++++++++- > src/cl_command_queue_gen7.c | 6 +++--- > src/cl_utils.h | 12 ++++++++++-- > 3 files changed, 23 insertions(+), 6 deletions(-) > > diff --git a/src/cl_api.c b/src/cl_api.c index 510941e..05ca889 100644 > --- a/src/cl_api.c > +++ b/src/cl_api.c > @@ -2999,6 +2999,7 @@ clEnqueueNDRangeKernel(cl_command_queue > command_queue, > fixed_local_sz[1] = 1; > } else { > uint j, maxDimSize = 64 /* from 64? */, maxGroupSize = 256; > //MAX_WORK_GROUP_SIZE may too large > + size_t realGroupSize = 1; > for (i = 0; i< work_dim; i++) { > for (j = maxDimSize; j > 1; j--) { > if (global_work_size[i] % j == 0 && j <= maxGroupSize) { @@ -3008,7 > +3009,10 @@ clEnqueueNDRangeKernel(cl_command_queue > command_queue, > break; //choose next work_dim > } > } > + realGroupSize *= fixed_local_sz[i]; > } > + if (realGroupSize < 8 || realGroupSize % 8 != 0) > + DEBUGP(DL_WARNING, "unable to find good values for > + local_work_size[i], please provide local_work_size[] explicitly, you > + can find good values with trial-and-error method."); > } > } > > @@ -3016,8 +3020,13 @@ clEnqueueNDRangeKernel(cl_command_queue > command_queue, > fixed_global_sz[0] = (global_work_size[0]+15) / 16 * 16; > fixed_global_sz[1] = (global_work_size[1]+15) / 16; > } else { > - for (i = 0; i < work_dim; ++i) > + size_t total_work_items = 1; > + for (i = 0; i < work_dim; ++i) { > fixed_global_sz[i] = global_work_size[i]; > + total_work_items *= fixed_global_sz[i]; > + } > + if (total_work_items < 64) > + DEBUGP(DL_WARNING, "too small work-items (see values in > + global_work_size[]) might result in bad performance."); > } > if (global_work_offset != NULL) > for (i = 0; i < work_dim; ++i) > diff --git a/src/cl_command_queue_gen7.c > b/src/cl_command_queue_gen7.c index 44db7ed..791a7ca 100644 > --- a/src/cl_command_queue_gen7.c > +++ b/src/cl_command_queue_gen7.c > @@ -369,21 +369,21 @@ > cl_command_queue_ND_range_gen7(cl_command_queue queue, > > /* Compute the number of HW threads we need */ > if(UNLIKELY(err = cl_kernel_work_group_sz(ker, local_wk_sz, 3, > &local_sz) != CL_SUCCESS)) { > - DEBUGP("Beignet: Work group size exceed Kernel's work group size.\n"); > + DEBUGP(DL_ERROR, "Work group size exceed Kernel's work group > + size."); > return err; > } > kernel.thread_n = thread_n = (local_sz + simd_sz - 1) / simd_sz; > kernel.curbe_sz = cst_sz; > > if (scratch_sz > ker->program->ctx->device->scratch_mem_size) { > - DEBUGP("Beignet: Out of scratch memory %d.\n", scratch_sz); > + DEBUGP(DL_ERROR, "Out of scratch memory %d.", scratch_sz); > return CL_OUT_OF_RESOURCES; > } > /* Curbe step 1: fill the constant urb buffer data shared by all threads */ > if (ker->curbe) { > kernel.slm_sz = cl_curbe_fill(ker, work_dim, global_wk_off, global_wk_sz, > local_wk_sz, thread_n); > if (kernel.slm_sz > ker->program->ctx->device->local_mem_size) { > - DEBUGP("Beignet: Out of shared local memory %d.\n", kernel.slm_sz); > + DEBUGP(DL_ERROR, "Out of shared local memory %d.", > + kernel.slm_sz); > return CL_OUT_OF_RESOURCES; > } > } > diff --git a/src/cl_utils.h b/src/cl_utils.h index 7595158..a0e37b2 100644 > --- a/src/cl_utils.h > +++ b/src/cl_utils.h > @@ -31,11 +31,19 @@ > #define JOIN(X, Y) _DO_JOIN(X, Y) > #define _DO_JOIN(X, Y) _DO_JOIN2(X, Y) > #define _DO_JOIN2(X, Y) X##Y > +enum DEBUGP_LEVEL > +{ > + DL_INFO, > + DL_WARNING, > + DL_ERROR > +}; > #ifdef NDEBUG > #define DEBUGP(...) > #else > - #define DEBUGP(fmt, ...) \ > - fprintf(stderr, fmt, ##__VA_ARGS__) > + //TODO: print or not with the value of level (DEBUGP_LEVEL) > + #define DEBUGP(level, fmt, ...) \ > + fprintf(stderr, "Beignet: "#fmt, ##__VA_ARGS__); \ fprintf(stderr, > + "\n"); Define level but always used stderrr? typo here?
> #endif > > /* Check compile time errors */ > -- > 1.9.1 > > _______________________________________________ > Beignet mailing list > [email protected] > http://lists.freedesktop.org/mailman/listinfo/beignet _______________________________________________ Beignet mailing list [email protected] http://lists.freedesktop.org/mailman/listinfo/beignet
