C++ clCreateImage2D函数代码示例

bool Sobel::runOpenCL(Image input, Image output, const Params& params)
{
    if (!initCL(params, sobel_kernel, "-cl-fast-relaxed-math"))
    {
        return false;
    }

    cl_int err;
    cl_kernel kernel;
    cl_mem d_input, d_output;
    cl_image_format format = {CL_RGBA, CL_UNORM_INT8};

    kernel = clCreateKernel(m_program, "sobel", &err);
    CHECK_ERROR_OCL(err, "creating kernel", return false);

    d_input = clCreateImage2D(
                  m_context, CL_MEM_READ_ONLY, &format,
                  input.width, input.height, 0, NULL, &err);
    CHECK_ERROR_OCL(err, "creating input image", return false);

    d_output = clCreateImage2D(
                   m_context, CL_MEM_WRITE_ONLY, &format,
                   input.width, input.height, 0, NULL, &err);
    CHECK_ERROR_OCL(err, "creating output image", return false);

    size_t origin[3] = {0, 0, 0};
    size_t region[3] = {input.width, input.height, 1};
    err = clEnqueueWriteImage(
              m_queue, d_input, CL_TRUE,
              origin, region, 0, 0, input.data, 0, NULL, NULL);
    CHECK_ERROR_OCL(err, "writing image data", return false);

    err  = clSetKernelArg(kernel, 0, sizeof(cl_mem), &d_input);
    err |= clSetKernelArg(kernel, 1, sizeof(cl_mem), &d_output);
    CHECK_ERROR_OCL(err, "setting kernel arguments", return false);

    reportStatus("Running OpenCL kernel");

    const size_t global[2] = {output.width, output.height};
    const size_t *local = NULL;
    if (params.wgsize[0] && params.wgsize[1])
    {
        local = params.wgsize;
    }

    // Timed runs
    for (int i = 0; i < params.iterations + 1; i++)
    {
        err = clEnqueueNDRangeKernel(
                  m_queue, kernel, 2, NULL, global, local, 0, NULL, NULL);
        CHECK_ERROR_OCL(err, "enqueuing kernel", return false);

        // Start timing after warm-up run
        if (i == 0)
        {
            err = clFinish(m_queue);
            CHECK_ERROR_OCL(err, "running kernel", return false);
            startTiming();
        }
    }

PassRefPtr<WebCLImage> WebCLContext::createImage2DBase(unsigned flags, unsigned width, unsigned height, unsigned rowPitch, unsigned channelOrder, unsigned channelType, void* data, ExceptionState& es)
{
    if (!width || !height) {
        es.throwWebCLException(WebCLException::InvalidImageFormatDescriptor, WebCLException::invalidImageFormatDescriptorMessage);
        return nullptr;
    }

    if (!WebCLInputChecker::isValidMemoryObjectFlag(flags)) {
        es.throwWebCLException(WebCLException::InvalidValue, WebCLException::invalidValueMessage);
        return nullptr;
    }

    ASSERT(data);
    flags |= CL_MEM_COPY_HOST_PTR;
    WebCLImageDescriptor imageDescriptor;
    imageDescriptor.setWidth(width);
    imageDescriptor.setHeight(height);
    imageDescriptor.setRowPitch(rowPitch);
    imageDescriptor.setChannelOrder(channelOrder);
    imageDescriptor.setChannelType(channelType);
    cl_image_format imageFormat = {channelOrder, channelType};

    cl_int err = CL_SUCCESS;
    cl_mem clMemId = clCreateImage2D(m_clContext, flags, &imageFormat, width, height, rowPitch, data, &err);
    if (err != CL_SUCCESS) {
        WebCLException::throwException(err, es);
        return nullptr;
    }

    RefPtr<WebCLImage> image = WebCLImage::create(clMemId, imageDescriptor, this);
    return image.release();
}

cl_mem *cl_allocTexture(int width, int height, void *data, size_t elementSize, cl_channel_type type) {
	cl_int status;

	cl_mem *mem;
	mem = (cl_mem *) malloc(sizeof(cl_mem));

	cl_image_format image_format;
	image_format.image_channel_order = CL_R;
	image_format.image_channel_data_type = type;

	*mem = clCreateImage2D(clGPUContext,
		CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,
		&image_format,
		width,
		height,
		width*elementSize,  // TODO make sure this is good
		data,
		&status);

	if(cl_errChk(status, "creating texture")) {
		exit(1);
	}

	return mem;
}

cl_mem
piglit_cl_create_image(piglit_cl_context context, cl_mem_flags flags,
                       const cl_image_format *format,
                       const piglit_image_desc *desc)
{
	cl_int errNo;
	cl_mem image = NULL;

#ifdef CL_VERSION_1_2
	if (piglit_cl_get_platform_version(context->platform_id) >= 12) {
		image = clCreateImage(context->cl_ctx, flags, format, desc, NULL, &errNo);
	} else
#endif
	if (desc->image_type == CL_MEM_OBJECT_IMAGE2D) {
		image = clCreateImage2D(context->cl_ctx, flags, format,
		                        desc->image_width, desc->image_height, 0,
		                        NULL, &errNo);
	} else if (desc->image_type == CL_MEM_OBJECT_IMAGE3D) {
		image = clCreateImage3D(context->cl_ctx, flags, format,
		                        desc->image_width, desc->image_height,
		                        desc->image_depth, 0, 0,
		                        NULL, &errNo);
	} else {
		fprintf(stderr,
		        "Invalid image mem object type: %s\n",
		        piglit_cl_get_enum_name(desc->image_type));
	}
	if(!piglit_cl_check_error(errNo, CL_SUCCESS)) {
		fprintf(stderr,
		        "Could not create image: %s\n",
		        piglit_cl_get_error_name(errNo));
	}

	return image;
}

/**
 * \brief ocl::Image::create Creates cl_mem for this Image.
 *
 * Note that no Memory is allocated. Allocation takes place when data is transfered.
 * It is assumed that an active Queue exists.
 *
 * \param width Width of the image.
 * \param height Height of the image.
 * \param type Channeltype of the image.
 * \param order Channelorder of the image.
 */
void ocl::Image::create(size_t width, size_t height, ChannelType type, ChannelOrder order, Access access)
{
    TRUE_ASSERT(this->_context != 0, "Context not valid - cannot create Image");
    cl_mem_flags flags = access;
    cl_image_format format;
    format.image_channel_order = order;
    format.image_channel_data_type = type;
    cl_int status;

#if defined(OPENCL_V1_0) || defined(OPENCL_V1_1)
    this->_id = clCreateImage2D(this->_context->id(), flags, &format, width, height, 0, NULL, &status);
#else
    _cl_image_desc desc;
    desc.image_type = CL_MEM_OBJECT_IMAGE2D;
    desc.image_height = height;
    desc.image_width = width;
    desc.image_depth = 1;
    desc.image_array_size = 1;
    desc.image_row_pitch = 0;
    desc.image_slice_pitch = 0;
    desc.num_mip_levels = 0;
    desc.num_samples = 0;
    desc.buffer = NULL;

    this->_id = clCreateImage(this->_context->id(), flags, &format, &desc, NULL, &status);
#endif
    OPENCL_SAFE_CALL(status);
    TRUE_ASSERT(this->_id != 0, "Could not create 2D image.");
}

/*!
    Creates a 2D OpenCL image object from \a image with the
    specified \a access mode.

    The OpenCL image is initialized with a copy of the contents of
    \a image.  The application's \a image can be discarded after the
    OpenCL image is created.

    Returns the new 2D OpenCL image object, or a null object
    if the image could not be created.  If \a image has a zero size,
    this function will return a null QCLImage2D object.

    \sa createImage2DDevice(), createImage2DHost()
*/
QCLImage2D QCLContext::createImage2DCopy
    (const QImage &image, QCLMemoryObject::Access access)
{
    Q_D(QCLContext);

    // Validate the image.
    if (image.width() < 1 || image.height() < 1)
        return QCLImage2D();
    QCLImageFormat format(image.format());
    if (format.isNull()) {
        qWarning("QCLContext::createImage2DCopy: QImage format %d "
                 "does not have an OpenCL equivalent", int(image.format()));
        return QCLImage2D();
    }

    // Create the image object.
    cl_int error = CL_INVALID_CONTEXT;
    cl_mem_flags flags = cl_mem_flags(access) | CL_MEM_COPY_HOST_PTR;
    cl_mem mem = clCreateImage2D
        (d->id, flags, &(format.m_format),
         image.width(), image.height(), image.bytesPerLine(),
         const_cast<uchar *>(image.bits()), &error);
    reportError("QCLContext::createImage2DCopy:", error);
    if (mem)
        return QCLImage2D(this, mem);
    else
        return QCLImage2D();
}

static VALUE
rcl_mem_create_image_2d(VALUE mod, VALUE context,
                                   VALUE flags,
                                   VALUE image_format,
                                   VALUE width, VALUE height, VALUE row_pitch,
                                   VALUE host_ptr)
{
    EXPECT_RCL_TYPE(context, Context);
    EXPECT_FIXNUM(flags);

    EXTRACT_SIZE(width, w);
    EXTRACT_SIZE(height, h);
    EXTRACT_SIZE(row_pitch, rp);

    cl_context cxt = ContextPtr(context);
    cl_mem_flags mf = FIX2INT(flags);
    EXTRACT_IMAGE_FORMAT(image_format, imgfmt);

    EXTRACT_POINTER(host_ptr, hp);

    cl_int res;
    cl_mem img = clCreateImage2D(cxt, mf, &imgfmt, w, h, rp, hp, &res);
    CHECK_AND_RAISE(res);

    return RMemory(img);
}

//(__read_only image2d_t  input, __read_only image2d_t  input2,__write_only image2d_t output, const int x1, const int x2, const int y1, const int y2, const int op)
void overlay(Image * I, Image * other, int x1, int y1, int op)
{
	setFormat(I,1);
	setFormat(other,1);
	cl_int result;
	size_t gsize[2];
	gsize[0] = I->width;
	gsize[1] = I->height;
	cl_mem tmp = clCreateImage2D(context, CL_MEM_READ_WRITE | CL_MEM_ALLOC_HOST_PTR, &format, I->width, I->height, 0,NULL,&result);
	
	size_t lsize[2];
	lsize[0] = getLocalSize(I->width);
	lsize[1] = getLocalSize(I->height);	
	
	check(result);
	check(clSetKernelArg(overlay_kernel,0,sizeof(cl_mem),&I->GPU_Image));
	check(clSetKernelArg(overlay_kernel,1,sizeof(cl_mem),&other->GPU_Image));
	check(clSetKernelArg(overlay_kernel,2,sizeof(cl_mem),&tmp));
	check(clSetKernelArg(overlay_kernel,3,sizeof(int),&x1));
	check(clSetKernelArg(overlay_kernel,4,sizeof(int),&other->width));
	check(clSetKernelArg(overlay_kernel,5,sizeof(int),&y1));
	check(clSetKernelArg(overlay_kernel,6,sizeof(int),&other->height));
	check(clSetKernelArg(overlay_kernel,7,sizeof(int),&op));
	check(clEnqueueNDRangeKernel(queue, overlay_kernel, 2, 0, gsize, lsize, 0 ,NULL ,NULL));
	
	clReleaseMemObject(I->GPU_Image);
	I->GPU_Image = tmp;
}

void convolve(Image * I, ConvMatrix * matrix)
{
	matrix_transfer(matrix,1);
	transfer(I,1);
	cl_int result;
	size_t gsize[2];
	gsize[0] = I->width;
	gsize[1] = I->height;
	
	size_t lsize[2];
	lsize[0] = getLocalSize(I->width);
	lsize[1] = 1;
		
	cl_mem tmp = clCreateImage2D(context,CL_MEM_READ_WRITE | CL_MEM_ALLOC_HOST_PTR, &format, I->width, I->height, 0,NULL,&result);
	
	check(clSetKernelArg(convolution_kernel,0,sizeof(cl_mem),&I->GPU_Image));
	check(clSetKernelArg(convolution_kernel,1,sizeof(cl_mem),&tmp));	
	check(clSetKernelArg(convolution_kernel,2,sizeof(cl_mem),&matrix->GPU_Matrix));	
	check(clSetKernelArg(convolution_kernel,3,sizeof(int),&matrix->dim));
	check(clSetKernelArg(convolution_kernel,4,sizeof(float),&matrix->div));
	check(clEnqueueNDRangeKernel(queue, convolution_kernel, 2, 0, gsize, NULL, 0, NULL, NULL));
	
	clReleaseMemObject(I->GPU_Image);
	I->GPU_Image = tmp;	
}

void OpenCLImage2D :: create(
	OpenCLContext &ctx, 
	cl_mem_flags usage, 
	AlloArray *array
) {
	destroy();
	detach();

	usage = OpenCLMemoryBuffer::check_memory_flags(usage, array->data.ptr);
	
	bool at_least_2d = array->header.dimcount >= 2;
	size_t width = array->header.dim[0];
	size_t height = at_least_2d ? array->header.dim[1] : 1;
	size_t rowstride = at_least_2d ? array->header.stride[1] : allo_array_size(array);
	cl_image_format format = OpenCLImageFormat::format_from_array(array);
	
	cl_int res = CL_SUCCESS;
	cl_mem mem = clCreateImage2D(
		ctx.get_context(),
		usage,
		&format,
		width,
		height,
		rowstride,
		array->data.ptr,
		&res
	);
	
	if(opencl_error(res, "clCreateImage2D error creating buffer")) {
		return;
	}
	
	mMem = mem;
	ctx.attach_resource(this);
}

void avgCase(Image * I, Image * other, convMatrix * matrix)
{
	
	matrix_transfer(matrix,1);
	transfer(I,1);
	transfer(other,1);
	cl_int result;
	
	size_t gsize[2];
	gsize[0] = other->width;
	gsize[1] = other->height;
	I->width = other->width;
	I->height = other->height;
	
	size_t lsize[2];
	lsize[0] = getLocalSize(I->width);
	lsize[1] = getLocalSize(I->height);
	
	cl_mem tmp = clCreateImage2D(context, CL_MEM_READ_WRITE | CL_MEM_ALLOC_HOST_PTR, &format, I->width, I->height, 0, NULL, &result);
	check(result);

	check(clSetKernelArg(avgCase_kernel,0,sizeof(cl_mem),&I->GPU_Image));
	check(clSetKernelArg(avgCase_kernel,1,sizeof(cl_mem),&other->GPU_Image));
	check(clSetKernelArg(avgCase_kernel,2,sizeof(cl_mem),&tmp));
	check(clSetKernelArg(avgCase_kernel,3,sizeof(cl_mem),&matrix->GPU_Matrix));
	check(clSetKernelArg(avgCase_kernel,4,sizeof(int),&matrix->dim));
	check(clSetKernelArg(avgCase_kernel,5,sizeof(float),&matrix->div));
	check(clEnqueueNDRangeKernel(queue, avgCase_kernel , 2, 0, gsize, NULL, 0, NULL, NULL));
	clReleaseMemObject(I->GPU_Image);
	free(I->CPU_Image);
	
	I->CPU_Image = NULL;
	I->CPU_Image = (TYPE*) malloc(sizeof(TYPE)*I->width*I->height*DEPTH);
	I->GPU_Image = tmp;
}

Image * new_ImageFromFile(char * file)
{
	Image * tmp = (Image*) malloc(sizeof(Image));
	unsigned char * data = NULL;
	LodePNG_decode32_file(&data, &tmp->width, &tmp->height, file);
	tmp->CPU_Image = (TYPE*) malloc(sizeof(TYPE)*tmp->width*tmp->height*DEPTH);
	for(unsigned int i = 0; i < tmp->height; i++)
	{
		for(unsigned int j = 0; j < tmp->width; j++)
		{
			for(unsigned int k = 0; k < DEPTH; k++)
			{
				 tmp->CPU_Image[(i*tmp->width + j)*DEPTH + k] = (data[(i*tmp->width + j)*DEPTH + k]);	
			}			
		}
	}	
	free(data);
	
	int result;
	tmp->GPU_Image = clCreateImage2D(context,CL_MEM_READ_WRITE | CL_MEM_ALLOC_HOST_PTR, &format, tmp->width, tmp->height, 0,NULL,&result);
	check(result);
	tmp->locality = 0;
	tmp->type = RGBTYPE;
	return tmp;
}

cl_mem OpenCLDevice::COM_clAttachMemoryBufferToKernelParameter(cl_kernel kernel, int parameterIndex, int offsetIndex,
                                                               list<cl_mem> *cleanup, MemoryBuffer **inputMemoryBuffers,
                                                               ReadBufferOperation *reader)
{
	cl_int error;
	
	MemoryBuffer *result = reader->getInputMemoryBuffer(inputMemoryBuffers);

	const cl_image_format imageFormat = {
		CL_RGBA,
		CL_FLOAT
	};

	cl_mem clBuffer = clCreateImage2D(this->m_context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, &imageFormat, result->getWidth(),
	                                  result->getHeight(), 0, result->getBuffer(), &error);

	if (error != CL_SUCCESS) { printf("CLERROR[%d]: %s\n", error, clewErrorString(error));  }
	if (error == CL_SUCCESS) cleanup->push_back(clBuffer);

	error = clSetKernelArg(kernel, parameterIndex, sizeof(cl_mem), &clBuffer);
	if (error != CL_SUCCESS) { printf("CLERROR[%d]: %s\n", error, clewErrorString(error));  }

	COM_clAttachMemoryBufferOffsetToKernelParameter(kernel, offsetIndex, result);
	return clBuffer;
}

OpenCLHandle FilterEffect::createOpenCLImageResult(uint8_t* source)
{
    FilterContextOpenCL* context = FilterContextOpenCL::context();
    ASSERT(context);

    if (context->inError())
        return 0;

    ASSERT(!hasResult());
    cl_image_format clImageFormat;
    clImageFormat.image_channel_order = CL_RGBA;
    clImageFormat.image_channel_data_type = CL_UNORM_INT8;

    int errorCode = 0;
#ifdef CL_API_SUFFIX__VERSION_1_2
    cl_image_desc imageDescriptor = { CL_MEM_OBJECT_IMAGE2D, m_absolutePaintRect.width(), m_absolutePaintRect.height(), 0, 0, 0, 0, 0, 0, 0};
    m_openCLImageResult = clCreateImage(context->deviceContext(), CL_MEM_READ_WRITE | (source ? CL_MEM_COPY_HOST_PTR : 0),
        &clImageFormat, &imageDescriptor, source, &errorCode);
#else
    m_openCLImageResult = clCreateImage2D(context->deviceContext(), CL_MEM_READ_WRITE | (source ? CL_MEM_COPY_HOST_PTR : 0),
        &clImageFormat, m_absolutePaintRect.width(), m_absolutePaintRect.height(), 0, source, &errorCode);
#endif
    if (context->isFailed(errorCode))
        return 0;

    return m_openCLImageResult;
}

bool AutoExposure::init(cl_context context, cl_device_id device, QSize computeSize, int updatePeriod)
{
    assert(!_initialized);

    _updatePeriod= updatePeriod;
    _lumaSize= computeSize;

    // Compile kernel and set arguments
    CLUtils::KernelDefines downDefines;
    downDefines["GAMMA_CORRECT"]= "2.2f";
    _downKernel= CLUtils::loadKernelPath(context, device, ":/kernels/lumaDownsample.cl",
            "lumaDownsample", downDefines, QStringList("../res/kernels/"));
    if(!_downKernel) {
        debugWarning("Could not compile kernel.");
        return false;
    }

    _lumaData= (uchar*)malloc(lumaDataBytes());
    if(!_lumaData) {
        debugWarning("Could not allocate data.");
        return false;
    }

    cl_int error;
    // This image could be WRITE_ONLY
    _lumaImage= clCreateImage2D(context, CL_MEM_READ_WRITE, clFormatGL(GL_R),
            _lumaSize.width(), _lumaSize.height(), 0, 0, &error);
    if(clCheckError(error, "clCreateImage2D")) {
        free(_lumaData);
        return false;
    }

    _initialized= true;
    return true;
}