“use of undeclared identifier” using MSAOpenGL

#1

Hi everyone,
I found the 1M Particle example in memo’s ofxMSAOpenGL(https://github.com/memo/ofxMSAOpenCL) and successfully compiled and ran it.

Then, in Particle.cl, I comment out some variables(DAMP, CENTER_FORCE, MOUSE_FORCE, MIN_SPEED) and add them inside the struct, because I want to access these value in OF and use ofxGui to control these value.

But I got these errors:

[verbose] OpenCL::setupFromOpenGL 
[verbose] 1 devices found, on 1 platforms

[verbose] 

*********
OpenCL Device information for device #0
 cl_device_id................0x1024500
 vendorName..................Intel Inc.
 deviceName..................Intel(R) Iris(TM) Graphics 6100
 driverVersion...............1.2(Mar 15 2018 22:04:25)
 deviceVersion...............OpenCL 1.2 
 maxComputeUnits.............48
 maxWorkItemDimensions.......3
 maxWorkItemSizes[0].........256
 maxWorkGroupSize............256
 maxClockFrequency...........1100
 maxMemAllocSize.............384.000 MB
 imageSupport................YES
 maxReadImageArgs............128
 maxWriteImageArgs...........8
 image2dMaxWidth.............16384
 image2dMaxHeight............16384
 image3dMaxWidth.............2048
 image3dMaxHeight............2048
 image3dMaxDepth.............2048
 maxSamplers.................16
 maxParameterSize............1024
 globalMemCacheSize..........0.000 MB
 globalMemSize...............1536.000 MB
 maxConstantBufferSize.......64.000 KB
 maxConstantArgs.............8
 localMemSize................64.000 KB
 errorCorrectionSupport......NO
 profilingTimerResolution....80
 endianLittle................1
 profile.....................FULL_PROFILE
 extensions..................cl_APPLE_SetMemObjectDestructor cl_APPLE_ContextLoggingFunctions cl_APPLE_clut cl_APPLE_query_kernel_names cl_APPLE_gl_sharing cl_khr_gl_event cl_khr_global_int32_base_atomics cl_khr_global_int32_extended_atomics cl_khr_local_int32_base_atomics cl_khr_local_int32_extended_atomics cl_khr_byte_addressable_store cl_khr_image2d_from_buffer cl_khr_gl_depth_images cl_khr_depth_images cl_khr_3d_image_writes 
*********


[verbose] OpenCLBuffer::initBuffer
[verbose] OpenCLMemoryObject::memoryObjectInit
[verbose] OpenCLBuffer::initFromGLObject
[verbose] OpenCLMemoryObject::memoryObjectInit
[verbose] OpenCL::loadProgramFromFile
[verbose] OpenCLProgram::OpenCLProgram
[verbose] OpenCLProgram::loadFromFile MSAOpenCL/Particle.cl, isBinary: 0
[verbose] OpenCLProgram::loadFromSource 
[verbose] OpenCLProgram::build options: -I "../../../data/" 
[ error ] 

 ***** Error building program. ***** 
 ***********************************


[ error ] <program source>:29:10: error: use of undeclared identifier 'mouseForce'
        diff *= mouseForce * invDistSQ;
         ^
<program source>:31:51: error: use of undeclared identifier 'centerForce'
    p->vel += (dimensions*0.5f - posBuffer[id]) * centerForce - diff* p->mass;
                                                  ^
<program source>:36:12: error: use of undeclared identifier 'minSpeed'
        if(speed2<minSpeed) posBuffer[id] = mousePos + diff * (1.0f + p->mass);
           ^
<program source>:39:12: error: use of undeclared identifier 'damp'
        p->vel *= damp;
           ^

[verbose] OpenCL::loadKernel updateParticle, 0x0
[verbose] OpenCLProgram::loadKernel updateParticle
[verbose] OpenCLKernel::OpenCLKernel 0x100a559c0, updateParticle
[ error ] Error creating kernel: updateParticle
[verbose] OpenCLMemoryObject::~OpenCLMemoryObject
[verbose] OpenCLMemoryObject::~OpenCLMemoryObject
[verbose] OpenCLKernel::~OpenCLKernel updateParticle
2019-05-25 09:59:04.664448-0500 190524_OpenCL_ParticleDebug[53962:23935931] [CL_INVALID_KERNEL] : OpenCL Error : Failed to release kernel! Kernel given is invalid!
[verbose] OpenCLProgram::~OpenCLProgram
[verbose] OpenCL::~OpenCL


(Particle.cl)

//#define DAMP            0.95f
//#define CENTER_FORCE    0.007f  //0.007f
//#define MOUSE_FORCE        300.0f    //300
//#define MIN_SPEED        0.1f


typedef struct{
	float2 vel;
	float mass;
    float damp;
    
    float centerForce;
    float mouseForce;
    float minSpeed;
    float dummy0;        // need this to make sure the float2 vel is aligned to a 16 byte boundary

    
    
} Particle;



__kernel void updateParticle(__global Particle* particles, __global float2* posBuffer, const float2 mousePos, const float2 dimensions){
	int id = get_global_id(0);
	__global Particle *p = &particles[id];
	
	float2 diff = mousePos - posBuffer[id];
	float invDistSQ = 1.0f / dot(diff, diff);
	diff *= mouseForce * invDistSQ;

    p->vel += (dimensions*0.5f - posBuffer[id]) * centerForce - diff* p->mass;
//    p->vel += posBuffer[id] * CENTER_FORCE - diff * p->mass;  // amo made changes this line

	
	float speed2 = dot(p->vel, p->vel);
	if(speed2<minSpeed) posBuffer[id] = mousePos + diff * (1.0f + p->mass);

	posBuffer[id] += p->vel;
	p->vel *= damp;
}


my ofApp.cpp

/******
 
 This example updates 1M particles on the GPU using OpenCL
 The OpenCL kernel writes position data directly to a VBO stored in the OpenGL device memory
 so now data transfer between host and device during runtime
 
 
 Kernel based on Rui's ofxOpenCL particle example opencl particles 001b.zip 
 at http://code.google.com/p/ruisource/ 
 *****/


#include "ofApp.h"
#include "MSAOpenCL.h"

#define NUM_PARTICLES (10000)


typedef struct{
    float2 vel;
    float mass;
    float damp;
    
    
    float centerForce;
    float mouseForce;
    float minSpeed;
    float dummy0;        // need this to make sure the float2 vel is aligned to a 16 byte boundary
    
} Particle;


float2				mousePos;
float2				dimensions;


msa::OpenCL			opencl;

msa::OpenCLBufferManagedT<Particle>	particles; // vector of Particles on host and corresponding clBuffer on device

float2				particlesPos[NUM_PARTICLES];
msa::OpenCLBufferManagedT<float2> particlePos; // vector of particle positions on host and corresponding clBuffer, and vbo on device

GLuint				vbo;


//--------------------------------------------------------------
void ofApp::setup(){
    gui.setup();
    gui.add(mass.setup("mass", 1, 0.1, 10));
    gui.add(damp.setup("damp", 0.95, 0, 1));
    gui.add(centerForce.setup("centerForce", 0.007, 0, 5));
    gui.add(mouseForce.setup("mouseForce", 300, 0, 1000));
    gui.add(minSpeed.setup("minSpeed", 0.1, 0, 5));


    
	ofBackground(0, 0, 0);
	ofSetLogLevel(OF_LOG_VERBOSE);
	ofSetVerticalSync(false);
	
	opencl.setupFromOpenGL();
    
    // create vbo
    glGenBuffersARB(1, &vbo);
	glBindBufferARB(GL_ARRAY_BUFFER_ARB, vbo);
	glBufferDataARB(GL_ARRAY_BUFFER_ARB, sizeof(float4) * NUM_PARTICLES, 0, GL_DYNAMIC_COPY_ARB);
	glBindBufferARB(GL_ARRAY_BUFFER_ARB, 0);

    // init host and CL buffers
    particles.initBuffer(NUM_PARTICLES);
    particlePos.initFromGLObject(vbo, NUM_PARTICLES);

    // init data
	
	for(int i=0; i<NUM_PARTICLES; i++) {
		Particle &p = particles[i];
		p.vel.set(0, 0);
        p.mass = ofRandom(0.5, 1);
    
        p.damp = 0.95;
        p.centerForce = 0.007;
        p.mouseForce = 300;
        p.minSpeed = 0.1;
        
		particlePos[i].set(ofRandomWidth(), ofRandomHeight());
	}
	
    particles.writeToDevice();
    particlePos.writeToDevice(); ///< uploads buffer data to shared CL/GL memory, so the vbo and the cl buffer are written in one go, since they occupy the same memory locations.
	
	opencl.loadProgramFromFile("MSAOpenCL/Particle.cl");
	opencl.loadKernel("updateParticle");

	opencl.kernel("updateParticle")->setArg(0, particles);
	opencl.kernel("updateParticle")->setArg(1, particlePos);
    opencl.kernel("updateParticle")->setArg(2, mousePos);//.getPtr(), sizeof(float2));
    opencl.kernel("updateParticle")->setArg(3, dimensions);//.getPtr(), sizeof(float2));
	
	glPointSize(1);
}




//--------------------------------------------------------------
void ofApp::update(){
	
	mousePos.x = ofGetMouseX();
	mousePos.y = ofGetMouseY();
	dimensions.x = ofGetWidth();
	dimensions.y = ofGetHeight();
    
    float DAMP;
    float CENTER_FORCE;
    float MOUSE_FORCE;
    float MIN_SPEED;
    
    for(int i=0; i<NUM_PARTICLES; i++) {
        Particle &p = particles[i];
        p.damp = damp;
        p.centerForce = centerForce;
        p.mouseForce = mouseForce;
        p.minSpeed = minSpeed;
    }
	
    opencl.kernel("updateParticle")->setArg(2, mousePos);//.getPtr(), sizeof(float2));
    opencl.kernel("updateParticle")->setArg(3, dimensions);//.getPtr(), sizeof(float2) );
	glFlush();
	
	opencl.kernel("updateParticle")->run1D(NUM_PARTICLES);
}


//--------------------------------------------------------------
void ofApp::draw(){
	opencl.finish();

    glColor4d(1.0f, 1.0f, 1.0f, 0.5f);
//    glColor3f(1.0f, 1.0f, 1.0f);
	glBindBufferARB(GL_ARRAY_BUFFER_ARB, vbo);

	opencl.finish();

	glEnableClientState(GL_VERTEX_ARRAY);
	glVertexPointer(2, GL_FLOAT, 0, 0);
	glDrawArrays(GL_POINTS, 0, NUM_PARTICLES);
	glBindBufferARB(GL_ARRAY_BUFFER_ARB, 0);
	
	
	glColor3f(1, 1, 1);
	string info = "fps: " + ofToString(ofGetFrameRate()) + "\nnumber of particles: " + ofToString(NUM_PARTICLES);
	ofDrawBitmapString(info, 20, 20);
    
    gui.draw();
}

//--------------------------------------------------------------
void ofApp::keyPressed(int key){
    ofToggleFullscreen();
}

ofApp.h

/******
 
 This example updates 1M particles on the GPU using OpenCL
 The OpenCL kernel writes position data directly to a VBO stored in the OpenGL device memory
 so now data transfer between host and device during runtime
 
 
 Kernel based on Rui's ofxOpenCL particle example opencl particles 001b.zip
 at http://code.google.com/p/ruisource/
 *****/






#pragma once


#include "ofMain.h"
#include "ofxGui.h"


class ofApp : public ofBaseApp{
	
public:
	void setup();
	void update();
	void draw();
    void keyPressed(int key);
    
    
    ofxPanel gui;
    ofxFloatSlider mass, damp, centerForce, mouseForce, minSpeed;

};

Is anyone know why I got these error?
Thank you so much!

#2

Oh I figure it out, i need -> to access the variable in the struct, my bad, lol

so for example, this is the wrong version:

should change to this:

p->vel *= p->damp;
1 Like