Writing to the different mip-map levels of a 3d texture

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?

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

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?