From 065319efdb2ffdfcb0f599452570ee3fcaaf4cde Mon Sep 17 00:00:00 2001 From: Benjamin Segovia Date: Thu, 17 May 2012 17:41:49 +0000 Subject: Added more tests for function arguments --- kernels/compiler_function_argument.cl | 7 +++++ kernels/compiler_function_argument0.cl | 7 +++++ kernels/compiler_function_argument1.cl | 7 +++++ src/cl_command_queue_gen7.c | 28 +++++++++---------- utests/CMakeLists.txt | 5 +++- utests/compiler_function_argument.cpp | 46 +++++++++++++++++++++++++++++++ utests/compiler_function_argument0.cpp | 45 ++++++++++++++++++++++++++++++ utests/compiler_function_argument1.cpp | 50 ++++++++++++++++++++++++++++++++++ 8 files changed, 180 insertions(+), 15 deletions(-) create mode 100644 kernels/compiler_function_argument.cl create mode 100644 kernels/compiler_function_argument0.cl create mode 100644 kernels/compiler_function_argument1.cl create mode 100644 utests/compiler_function_argument.cpp create mode 100644 utests/compiler_function_argument0.cpp create mode 100644 utests/compiler_function_argument1.cpp diff --git a/kernels/compiler_function_argument.cl b/kernels/compiler_function_argument.cl new file mode 100644 index 00000000..fe6de28f --- /dev/null +++ b/kernels/compiler_function_argument.cl @@ -0,0 +1,7 @@ +__kernel void +compiler_function_argument(__global int *dst, int value) +{ + int id = (int)get_global_id(0); + dst[id] = value; +} + diff --git a/kernels/compiler_function_argument0.cl b/kernels/compiler_function_argument0.cl new file mode 100644 index 00000000..6bc2e921 --- /dev/null +++ b/kernels/compiler_function_argument0.cl @@ -0,0 +1,7 @@ +__kernel void +compiler_function_argument0(__global int *dst, short value) +{ + int id = (int)get_global_id(0); + dst[id] = value; +} + diff --git a/kernels/compiler_function_argument1.cl b/kernels/compiler_function_argument1.cl new file mode 100644 index 00000000..8842b0b4 --- /dev/null +++ b/kernels/compiler_function_argument1.cl @@ -0,0 +1,7 @@ +__kernel void +compiler_function_argument1(__global int *dst, char value, short value0, int value1) +{ + int id = (int)get_global_id(0); + dst[id] = value + value0 + value1; +} + diff --git a/src/cl_command_queue_gen7.c b/src/cl_command_queue_gen7.c index 5da418cf..0ed0fe7a 100644 --- a/src/cl_command_queue_gen7.c +++ b/src/cl_command_queue_gen7.c @@ -102,7 +102,6 @@ error: static void cl_curbe_fill(cl_kernel ker, - char *curbe, const size_t *global_wk_off, const size_t *global_wk_sz, const size_t *local_wk_sz) @@ -110,7 +109,7 @@ cl_curbe_fill(cl_kernel ker, int32_t offset; #define UPLOAD(ENUM, VALUE) \ if ((offset = gbe_kernel_get_curbe_offset(ker->opaque, ENUM, 0)) >= 0) \ - *((uint32_t *) (curbe + offset)) = VALUE; + *((uint32_t *) (ker->curbe + offset)) = VALUE; UPLOAD(GBE_CURBE_LOCAL_SIZE_X, local_wk_sz[0]); UPLOAD(GBE_CURBE_LOCAL_SIZE_Y, local_wk_sz[1]); UPLOAD(GBE_CURBE_LOCAL_SIZE_Z, local_wk_sz[2]); @@ -125,12 +124,12 @@ cl_curbe_fill(cl_kernel ker, UPLOAD(GBE_CURBE_GROUP_NUM_Z, global_wk_sz[2]/local_wk_sz[2]); #undef UPLOAD - /* Write identity for the stack pointer. This will make the stack pointer - * computations faster in the kernel + /* Write identity for the stack pointer. This is required by the stack pointer + * computation in the kernel */ if ((offset = gbe_kernel_get_curbe_offset(ker->opaque, GBE_CURBE_STACK_POINTER, 0)) >= 0) { const uint32_t simd_sz = gbe_kernel_get_simd_width(ker->opaque); - uint32_t *stackptr = (uint32_t *) (curbe + offset); + uint32_t *stackptr = (uint32_t *) (ker->curbe + offset); int32_t i; for (i = 0; i < simd_sz; ++i) stackptr[i] = i; } @@ -170,7 +169,6 @@ cl_command_queue_ND_range_gen7(cl_command_queue queue, { cl_context ctx = queue->ctx; cl_gpgpu gpgpu = queue->gpgpu; - char *curbe = NULL; /* Does not include per-thread local IDs */ char *final_curbe = NULL; /* Includes them and one sub-buffer per group */ cl_gpgpu_kernel kernel; const uint32_t simd_sz = cl_kernel_get_simd_width(ker); @@ -187,8 +185,8 @@ cl_command_queue_ND_range_gen7(cl_command_queue queue, kernel.slm_sz = 0; /* Curbe step 1: fill the constant buffer data shared by all threads */ - curbe = alloca(ker->curbe_sz); - cl_curbe_fill(ker, curbe, global_wk_off, global_wk_sz, local_wk_sz); + if (ker->curbe) + cl_curbe_fill(ker, global_wk_off, global_wk_sz, local_wk_sz); /* Compute the number of HW threads we need */ TRY (cl_kernel_work_group_sz, ker, local_wk_sz, 3, &local_sz); @@ -209,12 +207,14 @@ cl_command_queue_ND_range_gen7(cl_command_queue queue, cl_gpgpu_states_setup(gpgpu, &kernel); /* Curbe step 2. Give the localID and upload it to video memory */ - TRY_ALLOC (final_curbe, (char*) alloca(thread_n * cst_sz)); - if (curbe) - for (i = 0; i < thread_n; ++i) - memcpy(final_curbe + cst_sz * i, curbe, cst_sz); - TRY (cl_set_varying_payload, ker, final_curbe, local_wk_sz, simd_sz, cst_sz, thread_n); - cl_gpgpu_upload_constants(gpgpu, final_curbe, thread_n*cst_sz); + if (ker->curbe) { + assert(cst_sz > 0); + TRY_ALLOC (final_curbe, (char*) alloca(thread_n * cst_sz)); + for (i = 0; i < thread_n; ++i) + memcpy(final_curbe + cst_sz * i, ker->curbe, cst_sz); + TRY (cl_set_varying_payload, ker, final_curbe, local_wk_sz, simd_sz, cst_sz, thread_n); + cl_gpgpu_upload_constants(gpgpu, final_curbe, thread_n*cst_sz); + } /* Start a new batch buffer */ batch_sz = cl_kernel_compute_batch_sz(ker); diff --git a/utests/CMakeLists.txt b/utests/CMakeLists.txt index 0fd5e823..f8b3805d 100644 --- a/utests/CMakeLists.txt +++ b/utests/CMakeLists.txt @@ -29,7 +29,10 @@ ADD_LIBRARY(utests SHARED compiler_lower_return1.cpp compiler_lower_return2.cpp compiler_array.cpp - compiler_array0.cpp) + compiler_array0.cpp + compiler_function_argument.cpp + compiler_function_argument0.cpp + compiler_function_argument1.cpp) TARGET_LINK_LIBRARIES(utests cl m) ADD_EXECUTABLE(run utest_run.cpp) diff --git a/utests/compiler_function_argument.cpp b/utests/compiler_function_argument.cpp new file mode 100644 index 00000000..29775ceb --- /dev/null +++ b/utests/compiler_function_argument.cpp @@ -0,0 +1,46 @@ +/* + * Copyright © 2012 Intel Corporation + * + * This library is free software; you can redistribute it and/or + * modify it under the terms of the GNU Lesser General Public + * License as published by the Free Software Foundation; either + * version 2 of the License, or (at your option) any later version. + * + * This library is distributed in the hope that it will be useful, + * but WITHOUT ANY WARRANTY; without even the implied warranty of + * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the GNU + * Lesser General Public License for more details. + * + * You should have received a copy of the GNU Lesser General Public + * License along with this library. If not, see . + * + * Author: Benjamin Segovia + */ + +#include "utest_helper.hpp" + +void compiler_function_argument(void) +{ + const size_t n = 2048; + const int value = 34; + + // Setup kernel and buffers + OCL_CREATE_KERNEL("compiler_function_argument"); + OCL_CREATE_BUFFER(buf[0], 0, n * sizeof(uint32_t), NULL); + OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]); + OCL_SET_ARG(1, sizeof(int), &value); + + // Run the kernel + globals[0] = n; + locals[0] = 16; + OCL_NDRANGE(1); + OCL_MAP_BUFFER(0); + + // Check results + for (uint32_t i = 0; i < n; ++i) + OCL_ASSERT(((int*)buf_data[0])[i] == value); +} + +MAKE_UTEST_FROM_FUNCTION(compiler_function_argument); + + diff --git a/utests/compiler_function_argument0.cpp b/utests/compiler_function_argument0.cpp new file mode 100644 index 00000000..3aac9cb7 --- /dev/null +++ b/utests/compiler_function_argument0.cpp @@ -0,0 +1,45 @@ +/* + * Copyright © 2012 Intel Corporation + * + * This library is free software; you can redistribute it and/or + * modify it under the terms of the GNU Lesser General Public + * License as published by the Free Software Foundation; either + * version 2 of the License, or (at your option) any later version. + * + * This library is distributed in the hope that it will be useful, + * but WITHOUT ANY WARRANTY; without even the implied warranty of + * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the GNU + * Lesser General Public License for more details. + * + * You should have received a copy of the GNU Lesser General Public + * License along with this library. If not, see . + * + * Author: Benjamin Segovia + */ + +#include "utest_helper.hpp" + +void compiler_function_argument0(void) +{ + const size_t n = 2048; + const short value = 34; + + // Setup kernel and buffers + OCL_CREATE_KERNEL("compiler_function_argument0"); + OCL_CREATE_BUFFER(buf[0], 0, n * sizeof(uint32_t), NULL); + OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]); + OCL_SET_ARG(1, sizeof(short), &value); + + // Run the kernel + globals[0] = n; + locals[0] = 16; + OCL_NDRANGE(1); + OCL_MAP_BUFFER(0); + + // Check results + for (uint32_t i = 0; i < n; ++i) + OCL_ASSERT(((int*)buf_data[0])[i] == value); +} + +MAKE_UTEST_FROM_FUNCTION(compiler_function_argument0); + diff --git a/utests/compiler_function_argument1.cpp b/utests/compiler_function_argument1.cpp new file mode 100644 index 00000000..dff46a81 --- /dev/null +++ b/utests/compiler_function_argument1.cpp @@ -0,0 +1,50 @@ +/* + * Copyright © 2012 Intel Corporation + * + * This library is free software; you can redistribute it and/or + * modify it under the terms of the GNU Lesser General Public + * License as published by the Free Software Foundation; either + * version 2 of the License, or (at your option) any later version. + * + * This library is distributed in the hope that it will be useful, + * but WITHOUT ANY WARRANTY; without even the implied warranty of + * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the GNU + * Lesser General Public License for more details. + * + * You should have received a copy of the GNU Lesser General Public + * License along with this library. If not, see . + * + * Author: Benjamin Segovia + */ + +#include "utest_helper.hpp" + +void compiler_function_argument1(void) +{ + const size_t n = 2048; + const char value = 34; + const short value0 = 31; + const int value1 = 3; + + // Setup kernel and buffers + OCL_CREATE_KERNEL("compiler_function_argument1"); + OCL_CREATE_BUFFER(buf[0], 0, n * sizeof(uint32_t), NULL); + OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]); + OCL_SET_ARG(1, sizeof(char), &value); + OCL_SET_ARG(2, sizeof(short), &value0); + OCL_SET_ARG(3, sizeof(int), &value1); + + // Run the kernel + globals[0] = n; + locals[0] = 16; + OCL_NDRANGE(1); + OCL_MAP_BUFFER(0); + + // Check results + for (uint32_t i = 0; i < n; ++i) + OCL_ASSERT(((int*)buf_data[0])[i] == value + value0 + value1); +} + +MAKE_UTEST_FROM_FUNCTION(compiler_function_argument1); + + -- cgit v1.2.1