OpenCL to OpenGL texture problems

Question

I'm trying to use OpenCL to draw to a cl_image that I got from a OpenGL texture and then render that texture. The problem is when I run my code on CL_DEVICE_TYPE_CPU it works fine however when I run on CL_DEVICE_TYPE_GPU the texture appears to be some random pixels. I'm new to OpenCL and not sure what's goin on so I'll post code below, also using OpenCL on OSX.

Host Code:

#import "GLView.h"
#import <GLKit/GLKit.h>
#import <OpenCL/OpenCL.h>
#import "kernel.cl.h"

#define WIDTH 500
#define HEIGHT 500

static GLfloat squareVertexData[] =
{
    -0.5f, -0.5f, 0.0f,     0.0f, 0.0f,
    0.5f, -0.5f, 0.0f,      1.0f, 0.0f,
    0.5f, 0.5f, 0.0f,       1.0f, 1.0f,
    -0.5f, -0.5f, 0.0f,     0.0f, 0.0f,
    0.5f, 0.5f, 0.0f,       1.0f, 1.0f,
    -0.5f, 0.5f, 0.0f,      0.0f, 1.0f
};

@interface GLView ()

@property (strong, nonatomic) GLKBaseEffect *effect;
@property (nonatomic) GLuint vertexArray;
@property (nonatomic) GLuint vertexBuffer;
@property (nonatomic) GLuint tex;
@property (nonatomic) cl_image clImage;
@property (nonatomic) dispatch_queue_t queue;

@end

@implementation GLView

#pragma mark - Lifecycle methods

- (void)awakeFromNib
{
    NSOpenGLPixelFormatAttribute attrs[] =
    {
        NSOpenGLPFADoubleBuffer,
        NSOpenGLPFAOpenGLProfile,
        NSOpenGLProfileVersion3_2Core,
        NSOpenGLPFAColorSize, 8,
        NSOpenGLPFAAlphaSize, 8,
        0
    };

    NSOpenGLPixelFormat *pixelFormat = [[NSOpenGLPixelFormat alloc] initWithAttributes:attrs];
    [self setPixelFormat:pixelFormat];

    NSOpenGLContext *context = [[NSOpenGLContext alloc] initWithFormat:pixelFormat shareContext:nil];
    CGLEnable([context CGLContextObj], kCGLCECrashOnRemovedFunctions);
    [self setOpenGLContext:context];
}

- (void)dealloc
{
    [self.openGLContext makeCurrentContext];

    glDeleteBuffers(1, &_vertexBuffer);
    glDeleteVertexArrays(1, &_vertexArray);

    glDeleteTextures(1, &_tex);
    gcl_release_image(self.clImage);
}

#pragma mark - NSOpenGLView methods

- (void)prepareOpenGL
{
    GLint swapInterval = 1;
    [self.openGLContext setValues:&swapInterval forParameter:NSOpenGLCPSwapInterval];

    glGenVertexArrays(1, &_vertexArray);
    glBindVertexArray(self.vertexArray);

    glGenBuffers(1, &_vertexBuffer);
    glBindBuffer(GL_ARRAY_BUFFER, _vertexBuffer);
    glBufferData(GL_ARRAY_BUFFER, sizeof(squareVertexData), squareVertexData, GL_STATIC_DRAW);

    glEnableVertexAttribArray(GLKVertexAttribPosition);
    glVertexAttribPointer(GLKVertexAttribPosition, 3, GL_FLOAT, GL_FALSE, 20, (char *)NULL + (0));

    glEnableVertexAttribArray(GLKVertexAttribTexCoord0);
    glVertexAttribPointer(GLKVertexAttribTexCoord0, 2, GL_FLOAT, GL_FALSE, 20, (char *)NULL + (12));

    glBindVertexArray(0);

    self.effect = [[GLKBaseEffect alloc] init];
    self.effect.transform.projectionMatrix = GLKMatrix4MakeOrtho(-1.0f, 1.0f, -1.0f, 1.0f, -1.0f, 1.0f);

    glGenTextures(1, &_tex);
    glBindTexture(GL_TEXTURE_2D, self.tex);
    glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_MAG_FILTER, GL_LINEAR);
    glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_MIN_FILTER, GL_LINEAR);
    glTexImage2D(GL_TEXTURE_2D, 0, GL_RGBA, WIDTH, HEIGHT, 0, GL_RGBA, GL_UNSIGNED_BYTE, NULL);

    self.effect.texture2d0.name = self.tex;

    CGLShareGroupObj sharegroup = CGLGetShareGroup(self.openGLContext.CGLContextObj);
    gcl_gl_set_sharegroup(sharegroup);

    self.clImage = gcl_gl_create_image_from_texture(GL_TEXTURE_2D, 0, self.tex);

    self.queue = gcl_create_dispatch_queue(CL_DEVICE_TYPE_GPU, NULL);

    dispatch_sync(self.queue, ^{

        cl_ndrange range = {
            2,
            { 0 },
            { WIDTH, HEIGHT },
            { 0 }
        };

        mykernel_kernel(&range, self.clImage);
    });

    glClearColor(0.65f, 0.65f, 0.65f, 1.0f);
    glClear(GL_COLOR_BUFFER_BIT);

    glBindVertexArray(self.vertexArray);

    [self.effect prepareToDraw];

    glDrawArrays(GL_TRIANGLES, 0, 6);

    glBindVertexArray(0);

    [self.openGLContext flushBuffer];
}

@end

Kernel Code:

kernel void mykernel(write_only image2d_t image)
{
    size_t x = get_global_id(0);
    size_t y = get_global_id(1);

    uint4 color;
    color.x = 255;
    color.w = 255;

    write_imageui(image, (int2) (x,y), color);
}

Where you assign the values to 'color' in the kernel, is there a reason why you only set the 'x' and 'w' components? — brm
no reason, just testing to see if it works, i see a red texture when on cpu, but weird random looking noise on gpu — marchinram
Do you also get the random noise if you do set all components? Perhaps not setting two of them can be the source of some randomness? The kernel itself looks fine otherwise, very similar to one that I've used for testing. — brm
I set them all now to 255, and when I run on CPU I see white texture, however on GPU it's always black, really weird what's happening — marchinram
The only other thing I can suggest is trying to use GL_NEAREST instead of GL_LINEAR — brm

Wil Shipley Wil Shipley · Accepted Answer · 2014-04-29T09:42:13

I’m not clear why you’re asking for:

    NSOpenGLPFAColorSize, 8,

This seems like not a lot of color. Possibly it’d work better if you ask for 32-bits (8 per pixel, plus alpha at end?):

    NSOpenGLPFAColorSize, 32,

That’s the default created by Scene Kit, anyhow.

OpenCL to OpenGL texture problems

1 Answers