Results 1 to 3 of 3

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

  1. #1
    Junior Member
    Join Date
    Apr 2012
    Posts
    4

    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:
    Code :
    glGenTextures(1, &voxelTexture);
    glActiveTexture(GL_TEXTURE0);
    glBindTexture(GL_TEXTURE_3D, voxelTexture);
    glTexStorage3D(GL_TEXTURE_3D, numMipMapLevels, GL_RGBA8, sideLength, sideLength, sideLength);
    During initialization:

    Code :
     
    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...

    Code :
     
    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:

    Code :
     
    #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?

  2. #2
    Junior Member
    Join Date
    Apr 2012
    Posts
    4

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

    If anyone is interested in reading the complete version of the code, here is is:

    Code :
    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;
    }

    Code :
    // 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

    Code :
    #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

    Code :
    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

  3. #3
    Junior Member
    Join Date
    Apr 2012
    Posts
    4

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

    I think I have pinpointed the problem further.

    The problem starts with

    Code :
    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?

Similar Threads

  1. glTexSubImage2D and cube map texture
    By HappyZappy in forum OpenGL ES general technical discussions
    Replies: 1
    Last Post: 11-09-2010, 08:09 AM
  2. Writing to an OpenGL Texture
    By pettersson in forum OpenCL
    Replies: 2
    Last Post: 11-29-2009, 01:53 PM

Posting Permissions

  • You may not post new threads
  • You may not post replies
  • You may not post attachments
  • You may not edit your posts
  •