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;
}
static VALUE
rcl_mem_create_image_3d(VALUE mod, VALUE context, VALUE flags,
VALUE image_format,
VALUE width, VALUE height, VALUE depth,
VALUE row_pitch, VALUE slice_pitch,
VALUE host_ptr)
{
EXPECT_RCL_TYPE(context, Context);
EXPECT_FIXNUM(flags);
if (CLASS_OF(image_format) != rcl_cImageFormat) {
rb_raise(rb_eTypeError, "expected argument 3 is a ImageFormat.");
}
cl_mem_flags mf = FIX2INT(flags);
EXTRACT_SIZE(width, w);
EXTRACT_SIZE(height, h);
EXTRACT_SIZE(depth, d);
EXTRACT_SIZE(row_pitch, rp);
EXTRACT_SIZE(slice_pitch, sp);
cl_context cxt = ContextPtr(context);
EXTRACT_IMAGE_FORMAT(image_format, imgfmt);
EXTRACT_POINTER(host_ptr, hp);
cl_int res;
cl_mem img = clCreateImage3D(cxt, mf, &imgfmt, w, h, d, rp, sp, hp, &res);
CHECK_AND_RAISE(res);
return RMemory(img);
}
/**
* \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 depth Depth 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, size_t depth, 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 = clCreateImage3D(this->_context->id(), flags, &format, width, height, depth, 0, 0, NULL, &status);
#else
_cl_image_desc desc;
desc.image_type = CL_MEM_OBJECT_IMAGE3D;
desc.image_height = height;
desc.image_width = width;
desc.image_depth = depth;
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 3D image.");
}
cl_mem WINAPI wine_clCreateImage3D(cl_context context, cl_mem_flags flags, cl_image_format * image_format,
size_t image_width, size_t image_height, size_t image_depth, size_t image_row_pitch, size_t image_slice_pitch,
void * host_ptr, cl_int * errcode_ret)
{
cl_mem ret;
TRACE("\n");
ret = clCreateImage3D(context, flags, image_format, image_width, image_height, image_depth, image_row_pitch, image_slice_pitch, host_ptr, errcode_ret);
return ret;
}
Image2D createImage(size3 size, cl_mem_flags flags = CL_MEM_READ_WRITE, typename ChannelType<T>::Type* host_ptr = NULL, size_t row_pitch = 0, size_t slice_pitch = 0)
{
cl_image_format format;
format.image_channel_order = O;
format.image_channel_data_type = T;
cl_int err;
cl_mem mem = clCreateImage3D(id(), flags, &format, size.s[0], size.s[1], size.s[2], row_pitch, slice_pitch, host_ptr, err);
CLPP_CHECK_ERROR(err);
return Image3D(mem);
}
/*!
Creates a 3D OpenCL image object with the specified \a format,
\a width, \a height, \a depth, and \a access mode.
The image memory is created on the device and will not be accessible
to the host via a direct pointer. Use createImage3DHost() to
create a host-accessible image.
Returns the new 3D OpenCL image object, or a null object
if the image could not be created.
\sa createImage3DHost(), createImage3DCopy()
*/
QCLImage3D QCLContext::createImage3DDevice
(const QCLImageFormat &format, int width, int height, int depth,
QCLMemoryObject::Access access)
{
Q_D(QCLContext);
cl_int error = CL_INVALID_CONTEXT;
cl_mem_flags flags = cl_mem_flags(access);
cl_mem mem = clCreateImage3D
(d->id, flags, &(format.m_format), width, height, depth, 0, 0,
0, &error);
reportError("QCLContext::createImage3DDevice:", error);
if (mem)
return QCLImage3D(this, mem);
else
return QCLImage3D();
}
void OpenCLImage::initWithoutTexture(int w,
int h,
int d,
cl_channel_order imageChannelOrder,
cl_channel_type imageChannelDataType,
cl_mem_flags memFlags,
void *dataPtr,
bool blockingWrite)
{
ofLog(OF_LOG_VERBOSE, "OpenCLImage::initWithoutTexture");
init(w, h, d);
cl_int err;
cl_image_format imageFormat;
imageFormat.image_channel_order = imageChannelOrder;
imageFormat.image_channel_data_type = imageChannelDataType;
int image_row_pitch = 0; // TODO
int image_slice_pitch = 0;
if(clMemObject) clReleaseMemObject(clMemObject);
if(depth == 1) {
clMemObject = clCreateImage2D(pOpenCL->getContext(), memFlags, &imageFormat, width, height, image_row_pitch, memFlags & CL_MEM_USE_HOST_PTR ? dataPtr : NULL, &err);
} else {
clMemObject = clCreateImage3D(pOpenCL->getContext(), memFlags, &imageFormat, width, height, depth, image_row_pitch, image_slice_pitch, memFlags & CL_MEM_USE_HOST_PTR ? dataPtr : NULL, &err);
}
assert(err != CL_INVALID_CONTEXT);
assert(err != CL_INVALID_VALUE);
assert(err != CL_INVALID_IMAGE_FORMAT_DESCRIPTOR);
assert(err != CL_INVALID_IMAGE_SIZE );
assert(err != CL_INVALID_HOST_PTR);
assert(err != CL_IMAGE_FORMAT_NOT_SUPPORTED);
assert(err != CL_MEM_OBJECT_ALLOCATION_FAILURE);
assert(err != CL_INVALID_OPERATION);
assert(err != CL_OUT_OF_HOST_MEMORY );
assert(err == CL_SUCCESS);
assert(clMemObject);
if(dataPtr) {
write(dataPtr, blockingWrite);
}
if(texture) delete texture;
texture = NULL;
}
/*!
Creates a 3D OpenCL image object with the specified \a format,
\a width, \a height, \a depth, and \a access mode.
If \a bytesPerLine is not zero, it indicates the number of bytes
between lines in \a data. If \a bytesPerSlice is not zero,
it indicates the number of bytes between slices in \a data.
The image is initialized with a copy of the contents of \a data.
The application's \a data can be discarded after the image
is created.
Returns the new 3D OpenCL image object, or a null object
if the image could not be created.
\sa createImage3DDevice(), createImage3DHost()
*/
QCLImage3D QCLContext::createImage3DCopy
(const QCLImageFormat &format, const void *data,
int width, int height, int depth,
QCLMemoryObject::Access access, int bytesPerLine, int bytesPerSlice)
{
Q_D(QCLContext);
cl_int error = CL_INVALID_CONTEXT;
cl_mem_flags flags = cl_mem_flags(access) | CL_MEM_COPY_HOST_PTR;
cl_mem mem = clCreateImage3D
(d->id, flags, &(format.m_format),
width, height, depth, bytesPerLine, bytesPerSlice,
const_cast<void *>(data), &error);
reportError("QCLContext::createImage3DCopy:", error);
if (mem)
return QCLImage3D(this, mem);
else
return QCLImage3D();
}
void OpenCLImage3D :: 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;
bool at_least_3d = array->header.dimcount >= 3;
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);
size_t depth = at_least_3d ? array->header.dim[2] : 1;
size_t planestride = at_least_3d ? array->header.stride[2] : allo_array_size(array);
cl_image_format format = OpenCLImageFormat::format_from_array(array);
cl_int res = CL_SUCCESS;
cl_mem mem = clCreateImage3D(
ctx.get_context(),
usage,
&format,
width,
height,
depth,
rowstride,
planestride,
array->data.ptr,
&res
);
if(opencl_error(res, "clCreateImage3D error creating buffer")) {
return;
}
mMem = mem;
ctx.attach_resource(this);
}
//--------------------------------------------------------------
bool RayTracingKernel::updateVolume(cl_context context, unsigned char* h_volume, int w, int h, int d)
{
// create 3D array and copy data to device
cl_image_format volume_format;
volume_format.image_channel_order = CL_R;
volume_format.image_channel_data_type = CL_UNORM_INT8;
printf("Load image of size: (%d, %d, %d)\n", w, h, d);
if(d_volumeArray != NULL) clReleaseMemObject(d_volumeArray);
d_volumeArray = clCreateImage3D(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, &volume_format,
w,h,d,
w, w*h,
h_volume, &clErr);
if (!checkOpenClError(clErr, "clCreateImage3D")) return false;
clErr = clSetKernelArg(kernel, 8, sizeof(cl_mem), (void *) &d_volumeArray);
return checkOpenClError(clErr, "clSetKernelArg 8");
}
cl_mem
clCreateImage(cl_context d_context,
cl_mem_flags flags,
const cl_image_format * image_format,
const cl_image_desc * image_desc,
void * host_ptr,
cl_int * errcode_ret)
{
cl_int dummy_errcode;
cl_mem image;
auto context = pobj(d_context);
if (!errcode_ret)
errcode_ret = &dummy_errcode;
if (!image_desc) {
*errcode_ret = CL_INVALID_IMAGE_DESCRIPTOR;
return 0;
}
size_t image_width = image_desc->image_width;
size_t image_height = image_desc->image_height;
size_t image_row_pitch = image_desc->image_row_pitch;
size_t image_depth = image_desc->image_depth;
size_t image_slice_pitch = image_desc->image_slice_pitch;
/* Just pass on to corresponding clCreateImage2D or clCreateImage3D functions: */
if (image_depth == 1) {
image = clCreateImage2D(context, flags, image_format, image_width, image_height,
image_row_pitch, host_ptr, errcode_ret);
}
else {
image = clCreateImage3D(context, flags, image_format, image_width, image_height,
image_depth, image_row_pitch, image_slice_pitch, host_ptr, errcode_ret);
}
return image;
}
void OpenCLImage3D :: create(
OpenCLContext &ctx,
cl_mem_flags usage,
const cl_image_format *format,
size_t width,
size_t height,
size_t depth,
size_t rowstride,
size_t planestride,
void *ptr
) {
destroy();
detach();
usage = OpenCLMemoryBuffer::check_memory_flags(usage, ptr);
cl_int res = CL_SUCCESS;
cl_mem mem = clCreateImage3D(
ctx.get_context(),
usage,
format,
width,
height,
depth,
rowstride,
planestride,
ptr,
&res
);
if(opencl_error(res, "clCreateImage3D error creating buffer")) {
return;
}
mMem = mem;
ctx.attach_resource(this);
}
cl_mem AllocateImage(clLabviewDevice *d, size_t Width, size_t Height, size_t Depth,
int ChannelOrder, int DataType, int *Error){
cl_mem Image = 0;
#ifndef NO_OPENCL
*Error = clLabviewDevice::Error(clLabviewDevice::SanitizeDevice(d));
if(*Error != 0)
return NULL;
cl_channel_order channel_order;
cl_channel_type data_type;
switch(DataType){
case 0:
data_type = CL_FLOAT;
break;
case 1:
data_type = CL_UNSIGNED_INT8;
break;
case 3:
data_type = CL_UNSIGNED_INT16;
break;
case 5:
data_type = CL_UNSIGNED_INT32;
break;
case 2:
data_type = CL_SIGNED_INT8;
break;
case 4:
data_type = CL_SIGNED_INT16;
break;
case 6:
data_type = CL_SIGNED_INT32;
break;
}
switch(ChannelOrder){
case 0:
channel_order = CL_LUMINANCE;
break;
case 1:
channel_order = CL_INTENSITY;
break;
case 2:
channel_order = CL_RGBA;
break;
case 3:
channel_order = CL_ARGB;
break;
case 4:
channel_order = CL_BGRA;
break;
case 5:
channel_order = CL_RG;
break;
case 6:
channel_order = CL_RA;
break;
case 7:
channel_order = CL_R;
break;
case 8:
channel_order = CL_A;
break;
}
cl_image_format image_format;
image_format.image_channel_data_type = data_type;
image_format.image_channel_order = channel_order;
//#ifdef CL_VERSION_1_2
// cl_image_desc image_disc;
// image_disc.num_mip_levels = 0;
// image_disc.num_samples = 0;
// image_disc.image_row_pitch = 0;
// image_disc.image_slice_pitch = 0;
// image_disc.image_depth = Depth;
// image_disc.image_height = Height;
// image_disc.image_width = Width;
// image_disc.image_array_size = Width*Height*Depth;
// image_disc.buffer = NULL;
//
// if(Depth == 1){
// image_disc.image_type = CL_MEM_OBJECT_IMAGE2D;
// }else{
// image_disc.image_type = CL_MEM_OBJECT_IMAGE3D;
// }
//
// cl_mem Image = clCreateImage(d->GetContext(), CL_MEM_ALLOC_HOST_PTR, &image_format, &image_disc, NULL, Error);
//#endif
//#ifdef CL_VERSION_1_1
if(Depth == 1){
Image = clCreateImage2D(d->GetContext(), CL_MEM_ALLOC_HOST_PTR, &image_format, Width, Height, NULL, NULL, Error);
}else{
Image = clCreateImage3D(d->GetContext(), CL_MEM_ALLOC_HOST_PTR, &image_format, Width, Height, Depth, NULL, NULL, NULL, Error);
}
//#endif
*Error = clLabviewDevice::Error(*Error);
#endif
if(*Error == 0){
return Image;
}else{
return NULL;
}
}
bool CL_Image3D::Create(cl_uint uWidth, cl_uint uHeight, cl_uint uDepth, cl_uint uRowPitch, void* pImgInput, CL_ImageOrder OrderType, CL_ImageChannel ChannelType, CL_MemAccess AccessType, CL_MemStorage StorageType)
{
CL_CPP_CONDITIONAL_RETURN_FALSE(m_Image);
CL_CPP_CONDITIONAL_RETURN_FALSE(!m_pContextRef);
CL_CPP_CONDITIONAL_RETURN_FALSE(!m_pContextRef->IsValid());
// CL_CPP_CONDITIONAL_RETURN_FALSE((StorageType == CL_MemStorage_UseHostInput || StorageType == CL_MemStorage_CopyInputToDevice) && !pImgInput);
cl_mem_flags uMemFlags = 0;
cl_image_format ImgFormat;
// Determine the access flags.
switch(AccessType)
{
case CL_MemAccess_ReadOnly: uMemFlags = CL_MEM_READ_ONLY; break;
case CL_MemAccess_WriteOnly: uMemFlags = CL_MEM_WRITE_ONLY; break;
case CL_MemAccess_ReadWrite: uMemFlags = CL_MEM_READ_WRITE; break;
default:
return false;
}
// Determine the storage flags.
switch(StorageType)
{
case CL_MemStorage_AllocateOnDevice: /* default setting, do nothing */ break;
case CL_MemStorage_AllocateOnHost: uMemFlags |= CL_MEM_ALLOC_HOST_PTR; break;
case CL_MemStorage_UseHostInput: uMemFlags |= CL_MEM_USE_HOST_PTR; break;
case CL_MemStorage_CopyInputToDevice: uMemFlags |= CL_MEM_COPY_HOST_PTR; break;
default:
return false;
}
// Determine the image channel order.
switch(OrderType)
{
case CL_ImageOrder_R: ImgFormat.image_channel_order = CL_R; break;
case CL_ImageOrder_A: ImgFormat.image_channel_order = CL_A; break;
case CL_ImageOrder_RG: ImgFormat.image_channel_order = CL_RG; break;
case CL_ImageOrder_RA: ImgFormat.image_channel_order = CL_RA; break;
case CL_ImageOrder_RGB: ImgFormat.image_channel_order = CL_RGB; break;
case CL_ImageOrder_RGBA: ImgFormat.image_channel_order = CL_RGBA; break;
case CL_ImageOrder_BGRA: ImgFormat.image_channel_order = CL_BGRA; break;
case CL_ImageOrder_ARGB: ImgFormat.image_channel_order = CL_ARGB; break;
case CL_ImageOrder_Intensity: ImgFormat.image_channel_order = CL_INTENSITY; break;
case CL_ImageOrder_Luminance: ImgFormat.image_channel_order = CL_LUMINANCE; break;
#ifdef CL_VERSION_1_1
case CL_ImageOrder_Rx: ImgFormat.image_channel_order = CL_Rx; break;
case CL_ImageOrder_RGx: ImgFormat.image_channel_order = CL_RGx; break;
case CL_ImageOrder_RGBx: ImgFormat.image_channel_order = CL_RGBx; break;
#endif
}
// Determine the image channel data type.
switch(ChannelType)
{
case CL_ImageChannel_Norm_Int8: ImgFormat.image_channel_data_type = CL_SNORM_INT8; break;
case CL_ImageChannel_Norm_Int16: ImgFormat.image_channel_data_type = CL_SNORM_INT16; break;
case CL_ImageChannel_Norm_UInt8: ImgFormat.image_channel_data_type = CL_UNORM_INT8; break;
case CL_ImageChannel_Norm_UInt16: ImgFormat.image_channel_data_type = CL_UNORM_INT16; break;
case CL_ImageChannel_Norm_UShort_555: ImgFormat.image_channel_data_type = CL_UNORM_SHORT_555; break;
case CL_ImageChannel_Norm_UShort_565: ImgFormat.image_channel_data_type = CL_UNORM_SHORT_565; break;
case CL_ImageChannel_Norm_UInt_101010: ImgFormat.image_channel_data_type = CL_UNORM_INT_101010; break;
case CL_ImageChannel_Int8: ImgFormat.image_channel_data_type = CL_SIGNED_INT8; break;
case CL_ImageChannel_Int16: ImgFormat.image_channel_data_type = CL_SIGNED_INT16; break;
case CL_ImageChannel_Int32: ImgFormat.image_channel_data_type = CL_SIGNED_INT32; break;
case CL_ImageChannel_UInt8: ImgFormat.image_channel_data_type = CL_UNSIGNED_INT8; break;
case CL_ImageChannel_UInt16: ImgFormat.image_channel_data_type = CL_UNSIGNED_INT16; break;
case CL_ImageChannel_UInt32: ImgFormat.image_channel_data_type = CL_UNSIGNED_INT32; break;
case CL_ImageChannel_Float16: ImgFormat.image_channel_data_type = CL_HALF_FLOAT; break;
case CL_ImageChannel_Float32: ImgFormat.image_channel_data_type = CL_FLOAT; break;
}
cl_context Context = m_pContextRef->GetContext();
cl_int iErrorCode = CL_SUCCESS;
cl_uint uSlicePitch = uRowPitch * uHeight;
// Create the image object.
#if defined(CL_VERSION_1_2)
cl_image_desc ImgDesc;
ImgDesc.image_width = uWidth;
ImgDesc.image_height = uHeight;
ImgDesc.image_depth = uDepth;
ImgDesc.image_array_size = 1;
ImgDesc.image_row_pitch = (pImgInput) ? uRowPitch : 0;
ImgDesc.image_slice_pitch = (pImgInput) ? uSlicePitch : 0;
ImgDesc.num_mip_levels = 0;
ImgDesc.num_samples = 0;
ImgDesc.buffer = NULL;
m_Image = clCreateImage(Context, uMemFlags, &ImgFormat, &ImgDesc, pImgInput, &iErrorCode);
#else
m_Image = clCreateImage3D(Context, uMemFlags, &ImgFormat, uWidth, uHeight, uDepth, (pImgInput) ? uRowPitch : 0, (pImgInput) ? uSlicePitch : 0, pImgInput, &iErrorCode);
#endif
CL_CPP_CATCH_ERROR(iErrorCode);
CL_CPP_ON_ERROR_RETURN_FALSE(iErrorCode);
m_uWidth = uWidth;
m_uHeight = uHeight;
m_uDepth = uDepth;
m_uRowPitch = uRowPitch;
m_uSlicePitch = uSlicePitch;
m_uTotalSize = m_uSlicePitch * m_uHeight;
// Ask OpenCL for the sizeo of each individual element in this image.
clGetImageInfo(m_Image, CL_IMAGE_ELEMENT_SIZE, sizeof(size_t), &m_uElementSize, NULL);
m_ImageOrder = OrderType;
m_ImageChannel = ChannelType;
m_MemAccess = AccessType;
m_MemStorage = StorageType;
return true;
}
int
SimpleImage::setupCL()
{
cl_int status = CL_SUCCESS;
#if 0
cl_device_type dType;
if(deviceType.compare("cpu") == 0)
{
dType = CL_DEVICE_TYPE_CPU;
}
else //deviceType = "gpu"
{
dType = CL_DEVICE_TYPE_GPU;
}
size_t deviceListSize;
/*
* Have a look at the available platforms and pick either
* the AMD one if available or a reasonable default.
*/
cl_uint numPlatforms;
cl_platform_id platform = NULL;
status = clGetPlatformIDs(0, NULL, &numPlatforms);
if(!sampleCommon->checkVal(status,
CL_SUCCESS,
"clGetPlatformIDs failed."))
{
return SDK_FAILURE;
}
if (0 < numPlatforms)
{
cl_platform_id* platforms = new cl_platform_id[numPlatforms];
status = clGetPlatformIDs(numPlatforms, platforms, NULL);
if(!sampleCommon->checkVal(status,
CL_SUCCESS,
"clGetPlatformIDs failed."))
{
return SDK_FAILURE;
}
for (unsigned i = 0; i < numPlatforms; ++i)
{
char pbuf[100];
status = clGetPlatformInfo(platforms[i],
CL_PLATFORM_VENDOR,
sizeof(pbuf),
pbuf,
NULL);
if(!sampleCommon->checkVal(status,
CL_SUCCESS,
"clGetPlatformInfo failed."))
{
return SDK_FAILURE;
}
platform = platforms[i];
if (!strcmp(pbuf, "Advanced Micro Devices, Inc."))
{
break;
}
}
delete[] platforms;
}
if(NULL == platform)
{
sampleCommon->error("NULL platform found so Exiting Application.");
return SDK_FAILURE;
}
// Display available devices.
if(!sampleCommon->displayDevices(platform, dType))
{
sampleCommon->error("sampleCommon::displayDevices() failed");
return SDK_FAILURE;
}
/*
* If we could find our platform, use it. Otherwise use just available platform.
*/
cl_context_properties cps[3] =
{
CL_CONTEXT_PLATFORM,
(cl_context_properties)platform,
0
};
context = clCreateContextFromType(
cps,
dType,
NULL,
NULL,
&status);
if(!sampleCommon->checkVal(status,
CL_SUCCESS,
"clCreateContextFromType failed."))
{
return SDK_FAILURE;
}
/* First, get the size of device list data */
status = clGetContextInfo(
context,
CL_CONTEXT_DEVICES,
0,
NULL,
&deviceListSize);
if(!sampleCommon->checkVal(
status,
CL_SUCCESS,
"clGetContextInfo failed."))
return SDK_FAILURE;
int deviceCount = (int)(deviceListSize / sizeof(cl_device_id));
if(!sampleCommon->validateDeviceId(deviceId, deviceCount))
{
sampleCommon->error("sampleCommon::validateDeviceId() failed");
return SDK_FAILURE;
}
/* Now allocate memory for device list based on the size we got earlier */
devices = (cl_device_id*)malloc(deviceListSize);
if(devices == NULL)
{
sampleCommon->error("Failed to allocate memory (devices).");
return SDK_FAILURE;
}
/* Now, get the device list data */
status = clGetContextInfo(
context,
CL_CONTEXT_DEVICES,
deviceListSize,
devices,
NULL);
if(!sampleCommon->checkVal(
status,
CL_SUCCESS,
"clGetContextInfo failed."))
return SDK_FAILURE;
/* Check for image support */
status = clGetDeviceInfo(devices[deviceId],
CL_DEVICE_IMAGE_SUPPORT,
sizeof(cl_bool),
&imageSupport,
0);
if(!sampleCommon->checkVal(
status,
CL_SUCCESS,
"clGetDeviceInfo failed."))
return SDK_FAILURE;
if(!imageSupport)
{
std::cout << "Error : Images are not supported on this device!\n";
return SDK_EXPECTED_FAILURE;
}
/* Create command queue */
cl_command_queue_properties prop = 0;
if(timing)
prop |= CL_QUEUE_PROFILING_ENABLE;
commandQueue = clCreateCommandQueue(
context,
devices[deviceId],
prop,
&status);
if(!sampleCommon->checkVal(
status,
CL_SUCCESS,
"clCreateCommandQueue failed."))
{
return SDK_FAILURE;
}
/*
* Create and initialize image objects
*/
/* Create 2D input image */
inputImage2D = clCreateImage2D(context,
CL_MEM_READ_ONLY | CL_MEM_USE_HOST_PTR,
&imageFormat,
width,
height,
0,
inputImageData,
&status);
if(!sampleCommon->checkVal(
status,
CL_SUCCESS,
"clCreateBuffer failed. (inputImageBuffer)"))
{
return SDK_FAILURE;
}
/* Create 2D output image */
outputImage2D = clCreateImage2D(context,
CL_MEM_WRITE_ONLY,
&imageFormat,
width,
height,
0,
0,
&status);
if(!sampleCommon->checkVal(
status,
CL_SUCCESS,
"clCreateBuffer failed. (inputImageBuffer)"))
{
return SDK_FAILURE;
}
/* Create 3D input image */
inputImage3D = clCreateImage3D(context,
CL_MEM_READ_ONLY | CL_MEM_USE_HOST_PTR,
&imageFormat,
width,
height / 2,
2, //2 slices
0,
0,
inputImageData,
&status);
if(!sampleCommon->checkVal(
status,
CL_SUCCESS,
"clCreateBuffer failed. (inputImageBuffer)"))
{
return SDK_FAILURE;
}
/* Writes to 3D images not allowed in spec currently */
outputImage3D = clCreateImage2D(context,
CL_MEM_WRITE_ONLY,
&imageFormat,
width,
height,
0,
0,
&status);
if(!sampleCommon->checkVal(
status,
CL_SUCCESS,
"clCreateBuffer failed. (inputImageBuffer)"))
{
return SDK_FAILURE;
}
if(!sampleCommon->checkVal(status,
CL_SUCCESS,
"clCreateBuffer failed. (outputImageBuffer)"))
{
return SDK_FAILURE;
}
/* create a CL program using the kernel source */
streamsdk::SDKFile kernelFile;
std::string kernelPath = sampleCommon->getPath();
if(isLoadBinaryEnabled())
{
kernelPath.append(loadBinary.c_str());
if(!kernelFile.readBinaryFromFile(kernelPath.c_str()))
{
std::cout << "Failed to load kernel file : " << kernelPath << std::endl;
return SDK_FAILURE;
}
const char * binary = kernelFile.source().c_str();
size_t binarySize = kernelFile.source().size();
program = clCreateProgramWithBinary(context,
1,
&devices[deviceId],
(const size_t *)&binarySize,
(const unsigned char**)&binary,
NULL,
&status);
if(!sampleCommon->checkVal(status,
CL_SUCCESS,
"clCreateProgramWithBinary failed."))
{
return SDK_FAILURE;
}
}
else
{
kernelPath.append("SimpleImage_Kernels.cl");
if(!kernelFile.open(kernelPath.c_str()))
{
std::cout << "Failed to load kernel file : "<< kernelPath << std::endl;
return SDK_FAILURE;
}
const char *source = kernelFile.source().c_str();
size_t sourceSize[] = {strlen(source)};
program = clCreateProgramWithSource(context,
1,
&source,
sourceSize,
&status);
if(!sampleCommon->checkVal(
status,
CL_SUCCESS,
"clCreateProgramWithSource failed."))
return SDK_FAILURE;
}
/* create a cl program executable for all the devices specified */
status = clBuildProgram(
program,
1,
&devices[deviceId],
NULL,
NULL,
NULL);
if(status != CL_SUCCESS)
{
if(status == CL_BUILD_PROGRAM_FAILURE)
{
cl_int logStatus;
char *buildLog = NULL;
size_t buildLogSize = 0;
logStatus = clGetProgramBuildInfo (program,
devices[deviceId],
CL_PROGRAM_BUILD_LOG,
buildLogSize,
buildLog,
&buildLogSize);
if(!sampleCommon->checkVal(
logStatus,
CL_SUCCESS,
"clGetProgramBuildInfo failed."))
return SDK_FAILURE;
buildLog = (char*)malloc(buildLogSize);
if(buildLog == NULL)
{
sampleCommon->error("Failed to allocate host memory. (buildLog)");
return SDK_FAILURE;
}
memset(buildLog, 0, buildLogSize);
logStatus = clGetProgramBuildInfo (program,
devices[deviceId],
CL_PROGRAM_BUILD_LOG,
buildLogSize,
buildLog,
NULL);
if(!sampleCommon->checkVal(
logStatus,
CL_SUCCESS,
"clGetProgramBuildInfo failed."))
{
free(buildLog);
return SDK_FAILURE;
}
std::cout << " \n\t\t\tBUILD LOG\n";
std::cout << " ************************************************\n";
std::cout << buildLog << std::endl;
std::cout << " ************************************************\n";
free(buildLog);
}
if(!sampleCommon->checkVal(
status,
CL_SUCCESS,
"clBuildProgram failed."))
return SDK_FAILURE;
}
/* get a kernel object handle for a kernel with the given name */
kernel2D = clCreateKernel(program, "image2dCopy", &status);
if(!sampleCommon->checkVal(
status,
CL_SUCCESS,
"clCreateKernel failed."))
{
return SDK_FAILURE;
}
kernel3D = clCreateKernel(program, "image3dCopy", &status);
if(!sampleCommon->checkVal(
status,
CL_SUCCESS,
"clCreateKernel failed."))
{
return SDK_FAILURE;
}
/* Check group size against group size returned by kernel */
status = clGetKernelWorkGroupInfo(kernel2D,
devices[deviceId],
CL_KERNEL_WORK_GROUP_SIZE,
sizeof(size_t),
&kernel2DWorkGroupSize,
0);
if(!sampleCommon->checkVal(
status,
CL_SUCCESS,
"clGetKernelWorkGroupInfo failed."))
{
return SDK_FAILURE;
}
/* Check group size against group size returned by kernel */
status = clGetKernelWorkGroupInfo(kernel3D,
devices[deviceId],
CL_KERNEL_WORK_GROUP_SIZE,
sizeof(size_t),
&kernel3DWorkGroupSize,
0);
if(!sampleCommon->checkVal(
status,
CL_SUCCESS,
"clGetKernelWorkGroupInfo failed."))
{
return SDK_FAILURE;
}
cl_uint temp = (cl_uint)min(kernel2DWorkGroupSize, kernel3DWorkGroupSize);
if((blockSizeX * blockSizeY) > temp)
{
if(!quiet)
{
std::cout << "Out of Resources!" << std::endl;
std::cout << "Group Size specified : "
<< blockSizeX * blockSizeY << std::endl;
std::cout << "Max Group Size supported on the kernel(s) : "
<< temp << std::endl;
std::cout << "Falling back to " << temp << std::endl;
}
if(blockSizeX > temp)
{
blockSizeX = temp;
blockSizeY = 1;
}
}
#endif
return SDK_SUCCESS;
}
/** vglClUpload branch3d
*/
void vglClUpload(VglImage* img)
{
if (Interop && img->nChannels > 1)
{
vglClUploadInterop(img);
}
else
{
if (img->nChannels == 3)
{
fprintf(stderr, "%s: %s: Error: image with 3 channels not supported. Please convert to 4 channels.\n", __FILE__, __FUNCTION__);
exit(1);
}
cl_int err;
if ( !vglIsInContext(img, VGL_RAM_CONTEXT) &&
!vglIsInContext(img, VGL_BLANK_CONTEXT) )
{
fprintf(stderr, "vglClUpload: Error: image context = %d not in VGL_RAM_CONTEXT or VGL_BLANK_CONTEXT\n", img->inContext);
return;
}
if (img->oclPtr == NULL)
{
/*if (img->fbo != -1)
{
img->oclPtr = clCreateFromGLTexture2D(cl.context,CL_MEM_READ_WRITE,GL_TEXTURE_2D,0,img->fbo,&err);
vglClCheckError( err, (char*) "clCreateFromGlTexture2D interop" );
clEnqueueAcquireGLObjects(cl.commandQueue, 1, &img->oclPtr, 0,0,0);
}
else
{*/
cl_image_format format;
if (img->nChannels == 1)
{
format.image_channel_order = CL_R;
}
else if (img->nChannels == 4)
{
format.image_channel_order = CL_RGBA;
}
if (img->depth == IPL_DEPTH_8U)
{
format.image_channel_data_type = CL_UNORM_INT8;
}
else if (img->depth == IPL_DEPTH_16U)
{
format.image_channel_data_type = CL_UNORM_INT16;
}
else if (img->depth == IPL_DEPTH_32S)
{
format.image_channel_data_type = CL_SIGNED_INT32;
}
else
{
fprintf(stderr, "%s: %s: Error: Unsupported image depth = %d.\n", __FILE__, __FUNCTION__, img->depth);
format.image_channel_data_type = CL_UNORM_INT8;
}
if (img->ndim == 2)
{
img->oclPtr = clCreateImage2D(cl.context, CL_MEM_READ_WRITE, &format, img->shape[VGL_WIDTH], img->shape[VGL_HEIGHT], 0, NULL, &err);
vglClCheckError( err, (char*) "clCreateImage2D" );
}
else if(img->ndim == 3)
{
img->oclPtr = clCreateImage3D(cl.context, CL_MEM_READ_WRITE, &format, img->shape[VGL_WIDTH], img->shape[VGL_HEIGHT], img->shape[VGL_LENGTH], 0, 0, NULL, &err);
vglClCheckError( err, (char*) "clCreateImage3D" );
}
else
{
img->oclPtr = clCreateBuffer(cl.context, CL_MEM_READ_WRITE, img->getTotalSizeInBytes(), NULL, &err);
vglClCheckError( err, (char*) "clCreateNDImage" );
}
/*
cl_image_desc desc;
if (img->ndim == 2)
{
desc.image_type = CL_MEM_OBJECT_IMAGE2D;
desc.image_width = img->shape[VGL_WIDTH];
desc.image_height = img->shape[VGL_HEIGHT];
desc.image_depth = 0;
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;
}
else
{
desc.image_type = CL_MEM_OBJECT_IMAGE3D;
desc.image_width = img->shape[VGL_WIDTH];
desc.image_height = img->shape[VGL_HEIGHT];
desc.image_depth = img->shape[VGL_LENGTH];
desc.image_array_size = 0;
desc.image_row_pitch = 0;
desc.image_slice_pitch = 0;
desc.num_mip_levels = 0;
desc.num_samples = 0;
desc.buffer = NULL;
}
img->oclPtr = clCreateImage(cl.context,CL_MEM_READ_WRITE, &format, &desc,NULL,&err);
vglClCheckError(err, (char*) "clCreateImage");
*/
}
if (vglIsInContext(img, VGL_RAM_CONTEXT))
{
size_t Origin[3] = { 0, 0, 0};
int nFrames = 1;
if(img->ndim == 3)
{
nFrames = img->shape[VGL_LENGTH];
}
void* imageData = img->getImageData();
if (!imageData)
{
fprintf(stderr, "%s: %s: Error: both ipl and ndarray are NULL.\n", __FILE__, __FUNCTION__);
exit(1);
}
if ( (img->ndim == 2) || (img->ndim == 3) )
{
size_t Size3d[3] = {img->shape[VGL_WIDTH], img->shape[VGL_HEIGHT], nFrames};
err = clEnqueueWriteImage( cl.commandQueue, img->oclPtr, CL_TRUE, Origin, Size3d, 0, 0, (char*)imageData, 0, NULL, NULL );
vglClCheckError( err, (char*) "clEnqueueWriteImage" );
clFinish(cl.commandQueue);
}
else
{
err = clEnqueueWriteBuffer(cl.commandQueue, img->oclPtr, CL_TRUE, 0, img->getTotalSizeInBytes(), imageData, 0, NULL, NULL);
vglClCheckError( err, (char*) "clEnqueueWriteBuffer" );
clFinish(cl.commandQueue);
}
}
vglAddContext(img, VGL_CL_CONTEXT);
}
}
void initCLVolume(float *h_Volume, int DATA_W, int DATA_H, int DATA_D)
{
ciErrNum = CL_SUCCESS;
printf("b image support is %i \n",g_bImageSupport);
if (g_bImageSupport)
//if (true)
{
// create 3D array and copy data to device
cl_image_format volume_format;
volume_format.image_channel_order = CL_RGBA;
volume_format.image_channel_data_type = CL_UNORM_INT8;
unsigned char* h_tempVolume = (unsigned char*)malloc(DATA_W * DATA_H * DATA_D * sizeof(unsigned char) * 4);
for(int i = 0; i <(int)(DATA_W * DATA_H * DATA_D); i++)
{
h_tempVolume[4 * i] = (unsigned char)(h_Volume[i] / 10.0f);
//h_tempVolume[4 * i] = h_Volume[i]/10.0f;
//h_tempVolume[4 * i + 0] = 10;
//h_tempVolume[4 * i + 1] = 10;
//h_tempVolume[4 * i + 2] = 10;
//h_tempVolume[4 * i + 3] = 10;
}
int error;
d_volumeArray = clCreateImage3D(cxGPUContext, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, &volume_format,
DATA_W, DATA_H, DATA_D,
(DATA_W * 4), (DATA_W * DATA_H * 4),
//0, 0,
h_tempVolume, &error);
printf("Create image 3D error is %i \n",error);
//oclCheckErrorEX(ciErrNum, CL_SUCCESS, pCleanup);
free (h_tempVolume);
// create transfer function texture
float transferFunc[] = {
0.0, 0.0, 0.0, 0.0,
1.0, 0.0, 0.0, 1.0,
1.0, 0.5, 0.0, 1.0,
1.0, 1.0, 0.0, 1.0,
0.0, 1.0, 0.0, 1.0,
0.0, 1.0, 1.0, 1.0,
0.0, 0.0, 1.0, 1.0,
1.0, 0.0, 1.0, 1.0,
0.0, 0.0, 0.0, 0.0,
};
cl_image_format transferFunc_format;
transferFunc_format.image_channel_order = CL_RGBA;
transferFunc_format.image_channel_data_type = CL_FLOAT;
d_transferFuncArray = clCreateImage2D(cxGPUContext, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, &transferFunc_format,
9, 1, sizeof(float) * 9 * 4,
transferFunc, &ciErrNum);
printf("Error 8 is %i \n",ciErrNum);
//oclCheckErrorEX(ciErrNum, CL_SUCCESS, pCleanup);
// Create samplers for transfer function, linear interpolation and nearest interpolation
transferFuncSampler = clCreateSampler(cxGPUContext, true, CL_ADDRESS_CLAMP_TO_EDGE, CL_FILTER_LINEAR, &ciErrNum);
printf("Error 9 is %i \n",ciErrNum);
//oclCheckErrorEX(ciErrNum, CL_SUCCESS, pCleanup);
volumeSamplerLinear = clCreateSampler(cxGPUContext, true, CL_ADDRESS_REPEAT, CL_FILTER_LINEAR, &ciErrNum);
printf("Error 10 is %i \n",ciErrNum);
//oclCheckErrorEX(ciErrNum, CL_SUCCESS, pCleanup);
volumeSamplerNearest = clCreateSampler(cxGPUContext, true, CL_ADDRESS_REPEAT, CL_FILTER_NEAREST, &ciErrNum);
printf("Error 11 is %i \n",ciErrNum);
//oclCheckErrorEX(ciErrNum, CL_SUCCESS, pCleanup);
// set image and sampler args
ciErrNum = clSetKernelArg(ckKernel, 8, sizeof(cl_mem), (void *) &d_volumeArray);
ciErrNum |= clSetKernelArg(ckKernel, 9, sizeof(cl_mem), (void *) &d_transferFuncArray);
ciErrNum |= clSetKernelArg(ckKernel, 10, sizeof(cl_sampler), linearFiltering ? &volumeSamplerLinear : &volumeSamplerNearest);
ciErrNum |= clSetKernelArg(ckKernel, 11, sizeof(cl_sampler), &transferFuncSampler);
//oclCheckErrorEX(ciErrNum, CL_SUCCESS, pCleanup);
printf("Error 12 is %i \n",ciErrNum);
}
// init invViewMatrix
d_invViewMatrix = clCreateBuffer(cxGPUContext, CL_MEM_READ_ONLY, 12 * sizeof(float), 0, &ciErrNum);
ciErrNum |= clSetKernelArg(ckKernel, 7, sizeof(cl_mem), (void *) &d_invViewMatrix);
printf("Error 13 is %i \n",ciErrNum);
//oclCheckErrorEX(ciErrNum, CL_SUCCESS, pCleanup);
}
void obj_mem_manager::g_arrange_mem()
{
std::vector<int>().swap(obj_mem_manager::tdescrip.texture_nums);
std::vector<int>().swap(obj_mem_manager::tdescrip.texture_sizes);
std::vector<int>().swap(obj_mem_manager::obj_sub_nums);
///int maxnum = return_max_num(size);
int wt = which_temp_object;
cl_uint trianglecount=0;
std::vector<obj_g_descriptor> desc;
unsigned int n=0;
std::vector<int> newtexid;
std::vector<int> mtexids; ///mipmaps
///process loaded objects
for(unsigned int i=0; i<objects_container::obj_container_list.size(); i++)
{
objects_container *obj = &objects_container::obj_container_list[i];
if(obj->isloaded == false)
{
obj->call_load_func(&objects_container::obj_container_list[i]);
obj->set_active_subobjs(true);
}
}
///what to do with bumpmaps?
///process textures in active texture list
for(unsigned int i=0; i<texture::active_textures.size(); i++)
{
texture *tex = &texture::texturelist[texture::active_textures[i]];
if(tex->loaded == false)
{
tex->call_load_func(tex);
}
}
std::vector<std::pair<int, int> > unique_sizes;
///obj_mem_manager::c_texture_array
for(unsigned int i=0; i<texture::active_textures.size(); i++)
{
texture *T = &texture::texturelist[texture::active_textures[i]];
int s = T->get_largest_dimension();
bool iswithin = false;
for(unsigned int j=0; j<unique_sizes.size(); j++)
{
if(unique_sizes[j].first == s)
{
unique_sizes[j].second++;
iswithin = true;
}
}
if(!iswithin)
{
unique_sizes.push_back(std::make_pair(s, 1));
}
}
unsigned int final_memory_size = 0; ///doesn't do mipmaps, eh
for(unsigned int i=0; i<unique_sizes.size(); i++)
{
int size = unique_sizes[i].first;
int num = unique_sizes[i].second;
int num_pages = calc_num_slices(size, num);
for(int i=0; i<num_pages; i++)
{
obj_mem_manager::tdescrip.texture_sizes.push_back(size);
obj_mem_manager::tdescrip.texture_nums.push_back(0);
}
final_memory_size+=num_pages;
}
obj_mem_manager::c_texture_array = new cl_uchar4[max_tex_size*max_tex_size*final_memory_size];
std::vector<cl_uint> tex_num_ids;
std::vector<cl_uint> tex_active_ids;
int b = 0;
for(unsigned int i=0; i<texture::active_textures.size(); i++)
{
if(texture::texturelist[texture::active_textures[i]].type==0)
{
int t=0;
int mipmaps[MIP_LEVELS];
add_texture_and_mipmaps(texture::texturelist[texture::active_textures[i]], mipmaps, t);
newtexid.push_back(t);
tex_num_ids.push_back(b);
tex_active_ids.push_back(texture::active_textures[i]);
for(int n=0; n<MIP_LEVELS; n++)
{
mtexids.push_back(mipmaps[n]);
}
b++;
}
}
int mipbegin=newtexid.size();
for(unsigned int i=0; i<mtexids.size(); i++)
{
newtexid.push_back(mtexids[i]);
}
b = 0;
///if active texture increment and then do the check
///fill in obj_g_descriptors for all the subobjects of the objects in the scene
cl_uint cumulative_bump = 0;
for(unsigned int i=0; i<objects_container::obj_container_list.size(); i++)
{
objects_container* obj = &objects_container::obj_container_list[i];
obj_sub_nums.push_back(obj->objs.size());
obj->arrange_id = i;
for(std::vector<object>::iterator it=obj->objs.begin(); it!=obj->objs.end(); it++) ///if you call this more than once, it will break. Need to store how much it has already done, and start it again from there to prevent issues with mipmaps
{
it->object_g_id = n;
obj_g_descriptor g;
desc.push_back(g);
desc[n].tri_num=(it)->tri_num;
desc[n].start=trianglecount;
cl_uint num_id = 0;
for(unsigned int i=0; i<tex_active_ids.size(); i++)
{
if(tex_active_ids[i] == it->tid && texture::texturelist[it->tid].type == 0)
{
num_id = tex_num_ids[i];
}
}
desc[n].tid = num_id;
for(int i=0; i<MIP_LEVELS; i++)
{
desc[n].mip_level_ids[i]=mipbegin + desc[n].tid*MIP_LEVELS + i;
}
desc[n].world_pos=(it)->pos;
desc[n].world_rot=(it)->rot;
desc[n].has_bump = it->has_bump;
desc[n].cumulative_bump = cumulative_bump;
cumulative_bump+=it->has_bump;
trianglecount+=(it)->tri_num;
n++;
}
}
///allocate memory on the gpu
temporary_objects[wt].g_texture_sizes = clCreateBuffer(cl::context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, sizeof(int)*obj_mem_manager::tdescrip.texture_sizes.size(), obj_mem_manager::tdescrip.texture_sizes.data(), &cl::error);
temporary_objects[wt].g_texture_nums = clCreateBuffer(cl::context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, sizeof(int)*newtexid.size(), newtexid.data(), &cl::error);
cl_image_format fermat;
fermat.image_channel_order=CL_RGBA;
fermat.image_channel_data_type=CL_UNSIGNED_INT8;
///2048*4 2048*2048*4 are row pitch and row size
temporary_objects[wt].g_texture_array=clCreateImage3D(cl::context, CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR, &fermat, 2048, 2048, obj_mem_manager::tdescrip.texture_sizes.size(), 2048*4, (2048*2048*4), obj_mem_manager::c_texture_array, &cl::error);
///now, we need to lump texture sizes into catagories
temporary_objects[wt].g_obj_desc = clCreateBuffer(cl::context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, sizeof(obj_g_descriptor)*n, desc.data(), &cl::error);
temporary_objects[wt].g_obj_num = clCreateBuffer(cl::context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, sizeof(cl_uint), &n, &cl::error);
temporary_objects[wt].g_tri_mem = clCreateBuffer(cl::context, CL_MEM_READ_ONLY, sizeof(triangle)*trianglecount, NULL, &cl::error);
temporary_objects[wt].g_cut_tri_mem= clCreateBuffer(cl::context, CL_MEM_READ_WRITE, sizeof(cl_float4)*trianglecount*3, NULL, &cl::error);
if(cl::error!=0)
{
std::cout << "g_tri_mem create" << std::endl;
exit(cl::error);
}
temporary_objects[wt].g_tri_num = clCreateBuffer(cl::context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR , sizeof(cl_uint), &trianglecount, &cl::error);
temporary_objects[wt].g_cut_tri_num = clCreateBuffer(cl::context, CL_MEM_READ_WRITE, sizeof(cl_uint), NULL, &cl::error);
if(cl::error!=0)
{
std::cout << "g_tri_num create" << std::endl;
exit(cl::error);
}
cl_uint running=0;
int obj_id=0;
///write triangle data to gpu
int p=0;
for(std::vector<objects_container>::iterator it2 = objects_container::obj_container_list.begin(); it2!=objects_container::obj_container_list.end(); it2++)
{
objects_container* obj = &(*it2);
for(std::vector<object>::iterator it=obj->objs.begin(); it!=obj->objs.end(); it++)
{
for(int i=0; i<(*it).tri_num; i++)
{
(*it).tri_list[i].vertices[0].pad[1]=obj_id;
p++;
}
clEnqueueWriteBuffer(cl::cqueue, temporary_objects[wt].g_tri_mem, CL_TRUE, sizeof(triangle)*running, sizeof(triangle)*(*it).tri_num, (*it).tri_list.data(), 0, NULL, NULL);
running+=(*it).tri_num;
obj_id++;
}
}
///opencl is a tad stupid, to force it to actually write the memory (rather than just stupidly keep it allocated somewhere), you have call a trivial kernel to force it to be written
clSetKernelArg(cl::trivial_kernel, 0, sizeof(cl_mem), &temporary_objects[wt].g_tri_mem);
clSetKernelArg(cl::trivial_kernel, 1, sizeof(cl_mem), &temporary_objects[wt].g_texture_array);
clSetKernelArg(cl::trivial_kernel, 2, sizeof(cl_mem), &temporary_objects[wt].g_cut_tri_num);
size_t num = 100;
size_t local = 1;
clEnqueueNDRangeKernel(cl::cqueue, cl::trivial_kernel, 1, NULL, &num, &local, 0, NULL, NULL);
temporary_objects[wt].tri_num=trianglecount;
clFinish(cl::cqueue);
delete [] obj_mem_manager::c_texture_array; ///instead of reallocating this entire thing, keep it in memory and simply add bits on?
obj_mem_manager::c_texture_array = NULL;
}
