PDA

View Full Version : Writing to the different mip-map levels of a 3d texture



sebby_man
09-21-2012, 11:04 PM
My kernel is able to write to the first mip map level of a 3D texture but not the second. As background for this problem, I am trying to create an OpenCL mip map generator.

I created an OpenGL 3d texture that has two mip-map levels and am sharing it with OpenCL through two calls to clCreateFromGLTexture3D (with miplevel being 0 and 1). This gives me two cl images which I send to my kernel. No errors are generated so far.

In the kernel I am using the cl_khr_image_writes extension. Now I write the color white to the first texel of the first image and the first texel of the second image. I see white for the first mip but not the second.

The way I verify whether the write works or not is by using a separate 3D texture viewer, where each texel is represented by a colored cube. The viewer also allows switching between the mipmap levels, so I can clearly see that writing to the first mip works but writing to the second doesn't.

I am using a Radeon 7750 with Catalyst version 12.8

Now here most of the code. I omitted some lines where I do error checking

Creation of 3D texture:


glGenTextures(1, &voxelTexture);
glActiveTexture(GL_TEXTURE0);
glBindTexture(GL_TEXTURE_3D, voxelTexture);
glTexStorage3D(GL_TEXTURE_3D, numMipMapLevels, GL_RGBA8, sideLength, sideLength, sideLength);

During initialization:




glFinish();
clVoxelTextureMip0 = clCreateFromGLTexture3D(clGPUContext, CL_MEM_WRITE_ONLY, GL_TEXTURE_3D, 0, voxelTexture, &clError);
clVoxelTextureMip1 = clCreateFromGLTexture3D(clGPUContext, CL_MEM_WRITE_ONLY, GL_TEXTURE_3D, 1, voxelTexture, &clError);
clError = clSetKernelArg(clMipMapGeneratorKernel, 0, sizeof(cl_mem), &clVoxelTextureMip0);
clError |= clSetKernelArg(clMipMapGeneratorKernel, 1, sizeof(cl_mem), &clVoxelTextureMip1);
clFinish(clCommandQueue);



Then...




glFinish();
clError = clEnqueueAcquireGLObjects(clCommandQueue, 1, &clVoxelTextureMip0, 0,0,0);
clError |= clEnqueueAcquireGLObjects(clCommandQueue, 1, &clVoxelTextureMip1, 0,0,0);

const unsigned int globalWorkOffset[3] = {0,0,0};
const unsigned int globalWorkSize[3] = {textureSideLength/2, textureSideLength/2, textureSideLength/2};
const unsigned int localWorkSize[3] = {4,4,4};

clError = clEnqueueNDRangeKernel(clCommandQueue, clMipMapGeneratorKernel, 3, globalWorkOffset, globalWorkSize, localWorkSize, 0,0,0);

clError = clEnqueueReleaseGLObjects(clCommandQueue, 1, &clVoxelTextureMip0, 0,0,0);
clError |= clEnqueueReleaseGLObjects(clCommandQueue, 1, &clVoxelTextureMip1, 0,0,0);

clFinish(clCommandQueue);


And the kernel:




#pragma OPENCL EXTENSION cl_khr_3d_image_writes : enable

__kernel void mipMapGenerator(write_only image3d_t mipLevel0, write_only image3d_t mipLevel1)
{
float4 white = (float4)(1,1,1,1);
int4 first = (int4)(0,0,0,0);
write_imagef(mipLevel0, first, white);
write_imagef(mipLevel1, first, white);
};



I hope this gives a good idea of my problem. Any thoughts?

sebby_man
09-22-2012, 02:24 PM
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\n");
else
printf("the kernel failed to write the color white to the first mipmap image\n");

if(imageData1[0] == glm::u8vec4(255,255,255,255))
printf("the kernel correctly wrote the color white to the second mipmap image\n");
else
printf("the kernel failed to write the color white to the second mipmap image\n");

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

sebby_man
09-22-2012, 05:20 PM
I think I have pinpointed the problem further.

The problem starts with



clTexture3dMip1 = clCreateFromGLTexture(clGPUContext, CL_MEM_WRITE_ONLY, GL_TEXTURE_3D, 1, texture3d, &clError);


When I actually read the image with clEnqueueReadImage, the result is all 0's. However if I read from the 0th mipmap layer instead, the results are correct. Yet at no point do i get a CL error. I tried this with 2D textures as well and got the same outcome.


Something seems to be wrong with creating CL images from GL texture mipmaps levels above 0. Is this a problem that others have had?