template<> unsigned CData::initCData<Uint8ClampedArray>(cl_command_queue aQueue, cl_mem aMemObj, ArrayBufferView::ViewType aType, unsigned aLength, unsigned aSize, PassRefPtr<Uint8ClampedArray> anArray)
{
cl_int err_code;
m_type = aType;
m_length = aLength;
m_size = aSize;
m_memObj = aMemObj;
if (anArray.get())
m_theUint8ClampedArray = anArray;
else
m_theUint8ClampedArray.clear();
DEBUG_LOG_STATUS("initCData", "queue is " << aQueue << " buffer is " << aMemObj);
err_code = clRetainCommandQueue(m_queue);
if (err_code != CL_SUCCESS) {
DEBUG_LOG_ERROR("initCData", err_code);
// We should really fail here but a bug in the whatif OpenCL
// makes the above retain operation always fail.
m_isRetained = false;
} else
m_isRetained = true;
m_queue = aQueue;
return RT_OK;
}
/// Creates a new command queue object as a copy of \p other.
command_queue(const command_queue &other)
: m_queue(other.m_queue)
{
if(m_queue){
clRetainCommandQueue(m_queue);
}
}
cl_int WINAPI wine_clRetainCommandQueue(cl_command_queue command_queue)
{
cl_int ret;
TRACE("(%p)\n", command_queue);
ret = clRetainCommandQueue(command_queue);
TRACE("(%p)=%d\n", command_queue, ret);
return ret;
}
UfoNode *
ufo_gpu_node_new (gpointer cmd_queue)
{
UfoGpuNode *node;
g_return_val_if_fail (cmd_queue != NULL, NULL);
node = UFO_GPU_NODE (g_object_new (UFO_TYPE_GPU_NODE, NULL));
node->priv->cmd_queue = cmd_queue;
clRetainCommandQueue (cmd_queue);
return UFO_NODE (node);
}
/*!
Requests that this kernel instance be run on globalWorkSize() items,
optionally subdivided into work groups of localWorkSize() items.
The kernel will be enqueued and executed in a background thread.
Returns a QFuture object that can be used to wait for the kernel
to finish execution. The request is executed on the active
command queue for context().
Usually runInThread() is called implicitly via QtConcurrent::run():
\code
kernel.setGlobalWorkSize(100, 100);
QFuture<void> future = QtConcurrent::run(kernel, a1, b1);
future.waitForFinished();
\endcode
The kernel object must not be reused until the background
thread finishes execution of the kernel. Thus, the following
code will have unexpected effects:
\code
QFuture<void> future1 = QtConcurrent::run(kernel, a1, b1);
QFuture<void> future2 = QtConcurrent::run(kernel, a2, b2);
future1.waitForFinished();
future2.waitForFinished();
\endcode
The recommended method to run the same kernel multiple times in a
background thread is as follows:
\code
void runKernelTwice(QCLKernel &kernel)
{
kernel(a1, b1);
kernel(a2, b2).waitForFinished();
}
QFuture<void> future = QtConcurrent::run(runKernelTwice, kernel);
\endcode
\sa run()
*/
QFuture<void> QCLKernel::runInThread()
{
Q_D(const QCLKernel);
cl_kernel kernel = m_kernelId;
cl_command_queue queue = d->context->activeQueue();
if (!kernel || !queue)
return QFuture<void>();
clRetainKernel(kernel);
clRetainCommandQueue(queue);
return QtConcurrent::run
(qt_run_kernel, kernel, queue, d->globalWorkSize, d->localWorkSize);
}
/// Copies the command queue object from \p other to \c *this.
command_queue& operator=(const command_queue &other)
{
if(this != &other){
if(m_queue){
clReleaseCommandQueue(m_queue);
}
m_queue = other.m_queue;
if(m_queue){
clRetainCommandQueue(m_queue);
}
}
return *this;
}
/**
* gopencl_commandqueue_ref:
* @self: an instance of GopenclCommandqueue.
*
* Increments the object reference count.
*
* Returns: the instance itself.
*
* Since: 0.1
*/
gpointer
gopencl_commandqueue_ref (GopenclCommandqueue *self)
{
cl_command_queue cl_queue = NULL;
cl_int err = CL_SUCCESS;
g_object_get(self, "id", &cl_queue, NULL);
err = clRetainCommandQueue(cl_queue);
if (err != CL_SUCCESS) {
gopencl_format_error(err, NULL);
return NULL;
}
return g_object_ref(self);
}
nsresult dpoCKernel::InitKernel(cl_command_queue aCmdQueue, cl_kernel aKernel, cl_mem aFailureMem)
{
cl_int err_code;
kernel = aKernel;
err_code = clRetainCommandQueue( aCmdQueue);
if (err_code != CL_SUCCESS) {
DEBUG_LOG_ERROR("initCData", err_code);
return NS_ERROR_NOT_AVAILABLE;
}
cmdQueue = aCmdQueue;
failureMem = aFailureMem;
err_code = clSetKernelArg(kernel, 0, sizeof(cl_mem), &failureMem);
if (err_code != CL_SUCCESS) {
DEBUG_LOG_ERROR("initCData", err_code);
return NS_ERROR_NOT_AVAILABLE;
}
return NS_OK;
}
WEAK void halide_init_kernels(const char* src, int size) {
int err;
cl_device_id dev;
// Initialize one shared context for all Halide compiled instances
if (!cl_ctx) {
const cl_uint maxPlatforms = 4;
cl_platform_id platforms[maxPlatforms];
cl_uint platformCount = 0;
err = clGetPlatformIDs( maxPlatforms, platforms, &platformCount );
CHECK_ERR( err, "clGetPlatformIDs" );
cl_platform_id platform = NULL;
const char * name = get_opencl_platform();
if (name != NULL) {
for (cl_uint i = 0; i < platformCount; ++i) {
const cl_uint maxPlatformName = 256;
char platformName[maxPlatformName];
err = clGetPlatformInfo( platforms[i], CL_PLATFORM_NAME, maxPlatformName, platformName, NULL );
if (err != CL_SUCCESS) continue;
if (strstr(platformName, name))
{
platform = platforms[i];
break;
}
}
} else if (platformCount > 0) {
platform = platforms[0];
}
if (platform == NULL){
halide_printf("Failed to find OpenCL platform\n");
return;
}
#ifdef DEBUG
const cl_uint maxPlatformName = 256;
char platformName[maxPlatformName];
err = clGetPlatformInfo( platform, CL_PLATFORM_NAME, maxPlatformName, platformName, NULL );
CHECK_ERR( err, "clGetPlatformInfo" );
halide_printf("Got platform '%s', about to create context (t=%lld)\n",
platformName, (long long)halide_current_time_ns());
#endif
// Make sure we have a device
const cl_uint maxDevices = 4;
cl_device_id devices[maxDevices];
cl_uint deviceCount = 0;
err = clGetDeviceIDs( platform, CL_DEVICE_TYPE_ALL, maxDevices, devices, &deviceCount );
CHECK_ERR( err, "clGetDeviceIDs" );
if (deviceCount == 0) {
halide_printf("Failed to get device\n");
return;
}
dev = devices[deviceCount-1];
#ifdef DEBUG
const cl_uint maxDeviceName = 256;
char deviceName[maxDeviceName];
err = clGetDeviceInfo( dev, CL_DEVICE_NAME, maxDeviceName, deviceName, NULL );
CHECK_ERR( err, "clGetDeviceInfo" );
halide_printf("Got device '%s', about to create context (t=%lld)\n",
deviceName, (long long)halide_current_time_ns());
#endif
// Create context
cl_context_properties properties[] = { CL_CONTEXT_PLATFORM, (cl_context_properties)platform, 0 };
cl_ctx = clCreateContext(properties, 1, &dev, NULL, NULL, &err);
CHECK_ERR( err, "clCreateContext" );
// cuEventCreate(&__start, 0);
// cuEventCreate(&__end, 0);
halide_assert(!cl_q);
cl_q = clCreateCommandQueue(cl_ctx, dev, 0, &err);
CHECK_ERR( err, "clCreateCommandQueue" );
} else {
// Maintain ref count of context.
clRetainContext(cl_ctx);
clRetainCommandQueue(cl_q);
}
// Initialize a module for just this Halide module
if ((!__mod) && (size > 1)) {
#ifdef DEBUG
halide_printf("Compiling kernel (%i bytes)\n", size);
#endif
// Create module
cl_device_id devices[] = { dev };
size_t lengths[] = { size };
if (strstr(src, "/*OpenCL C*/")) {
// Program is OpenCL C.
const char * sources[] = { src };
__mod = clCreateProgramWithSource(cl_ctx, 1, &sources[0], NULL, &err );
CHECK_ERR( err, "clCreateProgramWithSource" );
} else {
// Program is SPIR binary.
const unsigned char * binaries[] = { (unsigned char *)src };
__mod = clCreateProgramWithBinary(cl_ctx, 1, devices, lengths, &binaries[0], NULL, &err );
CHECK_ERR( err, "clCreateProgramWithBinary" );
}
err = clBuildProgram( __mod, 1, &dev, NULL, NULL, NULL );
if (err != CL_SUCCESS) {
size_t len;
char buffer[2048];
halide_printf("Error: Failed to build program executable! err = %d\n", err);
if (clGetProgramBuildInfo(__mod, dev, CL_PROGRAM_BUILD_LOG, sizeof(buffer), buffer, &len) == CL_SUCCESS)
halide_printf("%s\n", buffer);
else
halide_printf("clGetProgramBuildInfo failed to get build log!\n");
halide_assert(err == CL_SUCCESS);
}
}
}
static void inc(cl_command_queue & something)
{
cl_int err = clRetainCommandQueue(something);
VIENNACL_ERR_CHECK(err);
}
void*
pic2bs(int h, int v, int* pic, float* res){
//ocl init
cl_platform_id * platform = NULL;
cl_device_id * device_list = NULL;
cl_context context;
cl_command_queue queue;
cl_uint n_platforms;
cl_uint n_devices;
cl_program prog;
cl_kernel kern;
cl_int status;
struct timeval tvb, tva;
//int* buf = new int [h];
//p i
status = clGetPlatformIDs(0, NULL, &n_platforms);
platform = new cl_platform_id[n_platforms];
status = clGetPlatformIDs(n_platforms, platform, NULL);
//d i
status = clGetDeviceIDs(platform[1], CL_DEVICE_TYPE_GPU, 0, NULL, &n_devices);
device_list = new cl_device_id [n_devices];
status = clGetDeviceIDs(platform[1], CL_DEVICE_TYPE_GPU, n_devices, device_list, NULL);
//c i
context = clCreateContext(NULL, n_devices, device_list, NULL, NULL, &status);
queue = clCreateCommandQueue(context, *device_list, 0, &status);
//m i
cl_mem clm_pic = clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_USE_HOST_PTR, v*h*sizeof(cl_int), pic, &status);
cl_mem clm_res = clCreateBuffer(context, CL_MEM_READ_WRITE | CL_MEM_USE_HOST_PTR, 2*sizeof(cl_float), res, &status);
//cping
status = clFinish(queue);
// status = clEnqueueWriteBuffer(queue, clm_pic, CL_TRUE, 0, h*v*sizeof(cl_int), pic, 0, NULL, NULL);
// status = clEnqueueWriteBuffer(queue, clm_href, CL_TRUE, 0, h*sizeof(cl_float), href, 0, NULL, NULL);
//
prog = clCreateProgramWithSource(context, 1, (const char**)&s_pic2bs, NULL, &status);
status |= clBuildProgram(prog, 1, device_list, NULL, NULL, NULL);
if (status){
cl_int logStatus;
char * buildLog = NULL;
size_t buildLogSize = 0;
logStatus = clGetProgramBuildInfo(prog, *device_list, CL_PROGRAM_BUILD_LOG, buildLogSize, buildLog, &buildLogSize);
buildLog = new char [buildLogSize];
logStatus = clGetProgramBuildInfo(prog, *device_list, CL_PROGRAM_BUILD_LOG, buildLogSize, buildLog, NULL);
std::cout << "Here comes the compilation log\n";
std::cout << buildLog;
exit(0);
}
kern = clCreateKernel(prog, "pic2hprof", &status);
//s a
gettimeofday(&tvb,NULL);
status = clSetKernelArg(kern, 0, sizeof(cl_mem), (void*)&clm_pic);
status = clSetKernelArg(kern, 1, sizeof(cl_mem), (void*)&clm_res);
status = clSetKernelArg(kern, 2, sizeof(cl_int), (void*)&h);
status = clSetKernelArg(kern, 3, sizeof(cl_int), (void*)&v);
//
size_t globalThreads = h/4;
size_t localThreads = THREADS;
//
status = clEnqueueNDRangeKernel(queue, kern, 1, NULL, &globalThreads, NULL, 0, NULL, NULL);
// result
status = clEnqueueReadBuffer(queue, clm_res, CL_TRUE, 0, 2*sizeof(float), res, 0, NULL, NULL);
// flush
status = clFlush(queue);
status = clFinish(queue);
gettimeofday(&tva,NULL);
std::cout << "Kernel t " << tva.tv_usec - tvb.tv_usec << " us\n";
//
status = clReleaseKernel(kern);
status = clReleaseProgram(prog);
status = clReleaseMemObject(clm_res);
status = clReleaseMemObject(clm_pic);
status = clRetainCommandQueue(queue);
status = clReleaseContext(context);
delete[] platform;
delete[] device_list;
return 0;
}
PIGLIT_CL_API_TEST_CONFIG_END
enum piglit_result
piglit_cl_test(const int argc,
const char** argv,
const struct piglit_cl_api_test_config* config,
const struct piglit_cl_api_test_env* env)
{
int ref_count = 0;
const int max_ref_count = 10;
cl_int errNo;
cl_uint* ref_count_ptr;
/*** Normal usage ***/
cl_command_queue command_queue = clCreateCommandQueue(env->context->cl_ctx,
env->device_id,
0,
&errNo);
if(!piglit_cl_check_error(errNo, CL_SUCCESS)){
fprintf(stderr,
"Failed (error code: %s): Create a command queue.\n",
piglit_cl_get_error_name(errNo));
return PIGLIT_FAIL;
}
ref_count_ptr = piglit_cl_get_command_queue_info(command_queue,
CL_QUEUE_REFERENCE_COUNT);
if(*ref_count_ptr != 1) {
free(ref_count_ptr);
fprintf(stderr,
"CL_QUEUE_REFERENCE_COUNT should be 1 after creating command queue.\n");
return PIGLIT_FAIL;
}
free(ref_count_ptr);
/* increase by two and decrease by one on each itreation */
for(ref_count = 1; ref_count < max_ref_count; ref_count++) {
errNo = clRetainCommandQueue(command_queue);
if(!piglit_cl_check_error(errNo, CL_SUCCESS)) {
fprintf(stderr,
"clRetainCommandQueue: Failed (error code: %s): Retain command queue.\n",
piglit_cl_get_error_name(errNo));
return PIGLIT_FAIL;
}
errNo = clReleaseCommandQueue(command_queue);
if(!piglit_cl_check_error(errNo, CL_SUCCESS)){
fprintf(stderr,
"clReleaseCommandQueue: Failed (error code: %s): Release command queue.\n",
piglit_cl_get_error_name(errNo));
return PIGLIT_FAIL;
}
errNo = clRetainCommandQueue(command_queue);
if(!piglit_cl_check_error(errNo, CL_SUCCESS)){
fprintf(stderr,
"clRetainCommandQueue: Failed (error code: %s): Retain command queue.\n",
piglit_cl_get_error_name(errNo));
return PIGLIT_FAIL;
}
/* check internal value of reference count */
ref_count_ptr =
piglit_cl_get_command_queue_info(command_queue,
CL_QUEUE_REFERENCE_COUNT);
if(*ref_count_ptr != (ref_count+1)) {
free(ref_count_ptr);
fprintf(stderr,
"CL_QUEUE_REFERENCE_COUNT is not changing accordingly.\n");
return PIGLIT_FAIL;
}
free(ref_count_ptr);
}
/* Decrease reference count to 0 */
for(ref_count = max_ref_count; ref_count > 0; ref_count--) {
errNo = clReleaseCommandQueue(command_queue);
if(!piglit_cl_check_error(errNo, CL_SUCCESS)){
fprintf(stderr,
"clReleaseCommandQueue: Failed (error code: %s): Release command queue.\n",
piglit_cl_get_error_name(errNo));
return PIGLIT_FAIL;
}
/* check internal value of reference count */
if(ref_count > 1) {
ref_count_ptr =
piglit_cl_get_command_queue_info(command_queue,
CL_QUEUE_REFERENCE_COUNT);
if(*ref_count_ptr != (ref_count-1)) {
free(ref_count_ptr);
fprintf(stderr,
"CL_QUEUE_REFERENCE_COUNT is not changing accordingly.\n");
return PIGLIT_FAIL;
}
free(ref_count_ptr);
}
}
return PIGLIT_PASS;
}
WEAK void halide_init_kernels(void *user_context, const char* src, int size) {
int err;
cl_device_id dev;
// Initialize one shared context for all Halide compiled instances
if (!(*cl_ctx)) {
const cl_uint maxPlatforms = 4;
cl_platform_id platforms[maxPlatforms];
cl_uint platformCount = 0;
err = clGetPlatformIDs( maxPlatforms, platforms, &platformCount );
CHECK_ERR( err, "clGetPlatformIDs" );
cl_platform_id platform = NULL;
// Find the requested platform, or the first if none specified.
const char * name = getenv("HL_OCL_PLATFORM");
if (name != NULL) {
for (cl_uint i = 0; i < platformCount; ++i) {
const cl_uint maxPlatformName = 256;
char platformName[maxPlatformName];
err = clGetPlatformInfo( platforms[i], CL_PLATFORM_NAME, maxPlatformName, platformName, NULL );
if (err != CL_SUCCESS) continue;
if (strstr(platformName, name))
{
platform = platforms[i];
break;
}
}
} else if (platformCount > 0) {
platform = platforms[0];
}
if (platform == NULL){
halide_printf(user_context, "Failed to find OpenCL platform\n");
return;
}
#ifdef DEBUG
const cl_uint maxPlatformName = 256;
char platformName[maxPlatformName];
err = clGetPlatformInfo( platform, CL_PLATFORM_NAME, maxPlatformName, platformName, NULL );
CHECK_ERR( err, "clGetPlatformInfo" );
halide_printf(user_context, "Got platform '%s', about to create context (t=%lld)\n",
platformName, (long long)halide_current_time_ns(user_context));
#endif
cl_device_type device_type = 0;
// Find the device types requested.
const char * dev_type = getenv("HL_OCL_DEVICE");
if (dev_type != NULL) {
if (strstr("cpu", dev_type))
device_type |= CL_DEVICE_TYPE_CPU;
if (strstr("gpu", dev_type))
device_type |= CL_DEVICE_TYPE_GPU;
}
// If no devices are specified yet, just use all.
if (device_type == 0)
device_type = CL_DEVICE_TYPE_ALL;
// Make sure we have a device
const cl_uint maxDevices = 4;
cl_device_id devices[maxDevices];
cl_uint deviceCount = 0;
err = clGetDeviceIDs( platform, device_type, maxDevices, devices, &deviceCount );
CHECK_ERR( err, "clGetDeviceIDs" );
if (deviceCount == 0) {
halide_printf(user_context, "Failed to get device\n");
return;
}
dev = devices[deviceCount-1];
#ifdef DEBUG
const cl_uint maxDeviceName = 256;
char deviceName[maxDeviceName];
err = clGetDeviceInfo( dev, CL_DEVICE_NAME, maxDeviceName, deviceName, NULL );
CHECK_ERR( err, "clGetDeviceInfo" );
halide_printf(user_context, "Got device '%s', about to create context (t=%lld)\n",
deviceName, (long long)halide_current_time_ns(user_context));
#endif
// Create context
cl_context_properties properties[] = { CL_CONTEXT_PLATFORM, (cl_context_properties)platform, 0 };
*cl_ctx = clCreateContext(properties, 1, &dev, NULL, NULL, &err);
CHECK_ERR( err, "clCreateContext" );
// cuEventCreate(&__start, 0);
// cuEventCreate(&__end, 0);
halide_assert(user_context, !(*cl_q));
*cl_q = clCreateCommandQueue(*cl_ctx, dev, 0, &err);
CHECK_ERR( err, "clCreateCommandQueue" );
} else {
#ifdef DEBUG
halide_printf(user_context, "Already had context %p\n", *cl_ctx);
#endif
// Maintain ref count of context.
CHECK_CALL( clRetainContext(*cl_ctx), "clRetainContext" );
CHECK_CALL( clRetainCommandQueue(*cl_q), "clRetainCommandQueue" );
CHECK_CALL( clGetContextInfo(*cl_ctx, CL_CONTEXT_DEVICES, sizeof(dev), &dev, NULL), "clGetContextInfo" );
}
// Initialize a module for just this Halide module
if ((!__mod) && (size > 1)) {
// Create module
cl_device_id devices[] = { dev };
size_t lengths[] = { size };
if (strstr(src, "/*OpenCL C*/")) {
// Program is OpenCL C.
#ifdef DEBUG
halide_printf(user_context, "Compiling OpenCL C kernel: %s\n\n", src);
#endif
const char * sources[] = { src };
__mod = clCreateProgramWithSource(*cl_ctx, 1, &sources[0], NULL, &err );
CHECK_ERR( err, "clCreateProgramWithSource" );
} else {
// Program is SPIR binary.
#ifdef DEBUG
halide_printf(user_context, "Compiling SPIR kernel (%i bytes)\n", size);
#endif
const unsigned char * binaries[] = { (unsigned char *)src };
__mod = clCreateProgramWithBinary(*cl_ctx, 1, devices, lengths, &binaries[0], NULL, &err );
CHECK_ERR( err, "clCreateProgramWithBinary" );
}
err = clBuildProgram( __mod, 1, &dev, NULL, NULL, NULL );
if (err != CL_SUCCESS) {
size_t len;
char buffer[2048];
halide_printf(user_context, "Error: Failed to build program executable! err = %d\n", err);
if (clGetProgramBuildInfo(__mod, dev, CL_PROGRAM_BUILD_LOG, sizeof(buffer), buffer, &len) == CL_SUCCESS)
halide_printf(user_context, "Build Log:\n %s\n-----\n", buffer);
else
halide_printf(user_context, "clGetProgramBuildInfo failed to get build log!\n");
halide_assert(user_context, err == CL_SUCCESS);
}
}
}
