diff options
| author | Homer Hsing <homer.xing@intel.com> | 2013-05-02 09:00:31 +0800 |
|---|---|---|
| committer | Zhigang Gong <zhigang.gong@linux.intel.com> | 2013-05-02 10:46:58 +0800 |
| commit | d0f10c38843b1f8f1e3a688b86a4139661816583 (patch) | |
| tree | 6b8c8d65754630758c67250b57eb7eef6441e835 /src | |
| parent | 29e29dcfe7060be1fbe75c0b5bdb2978e0a59128 (diff) | |
Support global constant arrays
Version 3.
Support global constant arrays defined outside any kernel.
Example:
constant int h[] = {71,72,73,74,75,76,77};
kernel void k(global int *dst) {
int i = get_global_id(0);
dst[i] = h[i % 7];
}
Signed-off-by: Homer Hsing <homer.xing@intel.com>
Reviewed-by: Zhigang Gong <zhigang.gong@linux.intel.com>
Diffstat (limited to 'src')
| -rw-r--r-- | src/cl_command_queue_gen7.c | 8 |
1 files changed, 8 insertions, 0 deletions
diff --git a/src/cl_command_queue_gen7.c b/src/cl_command_queue_gen7.c index 9402549c..108684f0 100644 --- a/src/cl_command_queue_gen7.c +++ b/src/cl_command_queue_gen7.c @@ -120,6 +120,7 @@ cl_curbe_fill(cl_kernel ker, UPLOAD(GBE_CURBE_GROUP_NUM_Y, global_wk_sz[1]/local_wk_sz[1]); UPLOAD(GBE_CURBE_GROUP_NUM_Z, global_wk_sz[2]/local_wk_sz[2]); UPLOAD(GBE_CURBE_THREAD_NUM, thread_n); + UPLOAD(GBE_CURBE_GLOBAL_CONSTANT_OFFSET, gbe_kernel_get_curbe_offset(ker->opaque, GBE_CURBE_GLOBAL_CONSTANT_DATA, 0) + 32); #undef UPLOAD /* Write identity for the stack pointer. This is required by the stack pointer @@ -132,6 +133,13 @@ cl_curbe_fill(cl_kernel ker, for (i = 0; i < (int32_t) simd_sz; ++i) stackptr[i] = i; } + /* Write global constant arrays */ + if ((offset = gbe_kernel_get_curbe_offset(ker->opaque, GBE_CURBE_GLOBAL_CONSTANT_DATA, 0)) >= 0) { + /* Write the global constant arrays */ + gbe_program prog = ker->program->opaque; + gbe_program_get_global_constant_data(prog, ker->curbe + offset); + } + /* Handle the various offsets to SLM */ const int32_t arg_n = gbe_kernel_get_arg_num(ker->opaque); int32_t arg, slm_offset = 0; |