If anyone is interested in reading the complete version of the code, here is is:
bool begin()
{
//------ Variables ------//
GLuint texture3d;
const unsigned int sideLength = 8;
// OpenCL stuff
cl_platform_id clPlatform;
cl_context clGPUContext;
cl_device_id clDevice;
cl_command_queue clCommandQueue;
cl_program clProgram;
cl_kernel clTexture3dWriteKernel;
cl_int clError;
// CL memory that interlinks with GL memory
cl_mem clTexture3dMip0;
cl_mem clTexture3dMip1;
//------ Initialize OpenCL ------//
// Get an OpenCL platform
cl_platform_id clPlatforms[10];
cl_uint numPlatforms;
clError = clGetPlatformIDs(10, clPlatforms, &numPlatforms);
if (clError != CL_SUCCESS)
printf("could not create platform");
// Chose the platform that contains the AMD card
clPlatform = clPlatforms[0];
// Get the device - for now just assume that the device supports sharing with OpenGL
clError = clGetDeviceIDs(clPlatform, CL_DEVICE_TYPE_GPU, 1, &clDevice, NULL);
if (clError != CL_SUCCESS)
printf("could not get a GPU device on the platform");
// Create the context, with support for sharing with OpenGL
cl_context_properties props[] =
{
CL_GL_CONTEXT_KHR, (cl_context_properties)wglGetCurrentContext(),
CL_WGL_HDC_KHR, (cl_context_properties)wglGetCurrentDC(),
CL_CONTEXT_PLATFORM, (cl_context_properties)clPlatform,
0
};
clGPUContext = clCreateContext(props, 1, &clDevice, NULL, NULL, &clError);
if (clError != CL_SUCCESS)
printf("could not create a context");
// Create a command-queue
clCommandQueue = clCreateCommandQueue(clGPUContext, clDevice, 0, &clError);
if (clError != CL_SUCCESS)
printf("could not create command queue");
// Load program source code
size_t programLength;
char* cSourceCL = loadProgramSource("src/texture3dWrite.cl", &programLength);
if(cSourceCL == NULL)
printf("could not load program source");
// Create the program
clProgram = clCreateProgramWithSource(clGPUContext, 1, (const char **) &cSourceCL, &programLength, &clError);
if (clError != CL_SUCCESS)
printf("could not create program");
// Build the program
clError = clBuildProgram(clProgram, 0, NULL, "-cl-fast-relaxed-math", NULL, NULL);
if (clError != CL_SUCCESS)
{
printf("could not build program");
char cBuildLog[10240];
clGetProgramBuildInfo(clProgram, clDevice, CL_PROGRAM_BUILD_LOG, sizeof(cBuildLog), cBuildLog, NULL);
printf(cBuildLog);
}
// Create the texture 3d write kernel
clTexture3dWriteKernel = clCreateKernel(clProgram, "texture3dWrite", &clError);
if (clError != CL_SUCCESS)
printf("could not create the texture 3d write kernel");
//------ Create OpenGL 3D texture ------//
// Create a 3D texture with 2 mipmap levels
glGenTextures(1, &texture3d);
glBindTexture(GL_TEXTURE_3D, texture3d);
glTexParameteri(GL_TEXTURE_3D, GL_TEXTURE_BASE_LEVEL, 0);
glTexParameteri(GL_TEXTURE_3D, GL_TEXTURE_MAX_LEVEL, 1);
glTexStorage3D(GL_TEXTURE_3D, 2, GL_RGBA8, sideLength, sideLength, sideLength);
//------ Create OpenCL objects from the 2 texture mipmap layers ------//
// Create CL versions of the first and second mip map level of the 3D voxel texture
clTexture3dMip0 = clCreateFromGLTexture(clGPUContext, CL_MEM_WRITE_ONLY, GL_TEXTURE_3D, 0, texture3d, &clError);
if (clError != CL_SUCCESS)
printf("could not create CL texture3D mip level 0 from OpenGL texture3D");
clTexture3dMip1 = clCreateFromGLTexture(clGPUContext, CL_MEM_WRITE_ONLY, GL_TEXTURE_3D, 1, texture3d, &clError);
if (clError != CL_SUCCESS)
printf("could not create CL texture3D mip level 1 from OpenGL texture3D");
//------ Prepare and invoke kernel ------//
glFinish();
// Acquire GL memory
clError = clEnqueueAcquireGLObjects(clCommandQueue, 1, &clTexture3dMip0, 0,0,0);
clError |= clEnqueueAcquireGLObjects(clCommandQueue, 1, &clTexture3dMip1, 0,0,0);
if (clError != CL_SUCCESS)
printf("could not acquire OpenGL memory objects");
// Set parameters of the mip map generator kernel
clError = clSetKernelArg(clTexture3dWriteKernel, 0, sizeof(cl_mem), &clTexture3dMip0);
clError |= clSetKernelArg(clTexture3dWriteKernel, 1, sizeof(cl_mem), &clTexture3dMip1);
if (clError != CL_SUCCESS)
printf("could not set kernel arguments");
// Perpare to call the kernel
const unsigned int globalWorkOffset = 0;
const unsigned int globalWorkSize = 1;
const unsigned int localWorkSize = 1;
// Call the kernel
clError = clEnqueueNDRangeKernel(clCommandQueue, clTexture3dWriteKernel, 1, &globalWorkOffset, &globalWorkSize, &localWorkSize, 0,0,0);
if (clError != CL_SUCCESS)
printf("could not call the kernel");
// Release GL memory
clError = clEnqueueReleaseGLObjects(clCommandQueue, 1, &clTexture3dMip0, 0,0,0);
clError |= clEnqueueReleaseGLObjects(clCommandQueue, 1, &clTexture3dMip1, 0,0,0);
if (clError != CL_SUCCESS)
printf("could not release OpenGL memory objects");
clFinish(clCommandQueue);
//------ Read textures and see if the kernel has worked ------//
std::vector<glm::u8vec4> imageData0(sideLength*sideLength*sideLength);
std::vector<glm::u8vec4> imageData1(sideLength/2*sideLength/2*sideLength/2);
glGetTexImage(GL_TEXTURE_3D, 0, GL_RGBA, GL_UNSIGNED_BYTE, &imageData0[0]);
glGetTexImage(GL_TEXTURE_3D, 1, GL_RGBA, GL_UNSIGNED_BYTE, &imageData1[0]);
if(imageData0[0] == glm::u8vec4(255,255,255,255))
printf("the kernel correctly wrote the color white to the first mipmap image
");
else
printf("the kernel failed to write the color white to the first mipmap image
");
if(imageData1[0] == glm::u8vec4(255,255,255,255))
printf("the kernel correctly wrote the color white to the second mipmap image
");
else
printf("the kernel failed to write the color white to the second mipmap image
");
return true;
}
// From the Nvidia OpenCL utils
char* loadProgramSource(const char* cFilename, size_t* szFinalLength)
{
// locals
FILE* pFileStream = NULL;
size_t szSourceLength;
if(fopen_s(&pFileStream, cFilename, "rb") != 0)
{
return NULL;
}
// get the length of the source code
fseek(pFileStream, 0, SEEK_END);
szSourceLength = ftell(pFileStream);
fseek(pFileStream, 0, SEEK_SET);
// allocate a buffer for the source code string and read it in
char* cSourceString = (char *)malloc(szSourceLength + 1);
if (fread((cSourceString), szSourceLength, 1, pFileStream) != 1)
{
fclose(pFileStream);
free(cSourceString);
return 0;
}
// close the file and return the total length of the string
fclose(pFileStream);
if(szFinalLength != 0)
{
*szFinalLength = szSourceLength;
}
cSourceString[szSourceLength] = '\0';
return cSourceString;
}
texture3dWrite.cl
#pragma OPENCL EXTENSION cl_khr_3d_image_writes : enable
__kernel void texture3dWrite(write_only image3d_t image0, write_only image3d_t image1)
{
int4 destination = (int4)(0,0,0,0);
float4 white = (float4)(1,1,1,1);
write_imagef(image0, destination, white);
write_imagef(image1, destination, white);
};
output
the kernel correctly wrote the color white to the first mipmap image
the kernel failed to write the color white to the second mipmap image