Stripped down to relevant bits

bPhysics.cpp

void bPhysics::FixedUpdate_CL(int _numWidgets, widget*& _widgets)
{
	const float currentTime = glfwGetTime();

	double frameTime = currentTime - frameStart;

	// Store the time elapsed since the last frame began
	accumulator += frameTime;

	// Record the starting of this frame
	frameStart = currentTime;

	// Avoid spiral of death and clamp dt, thus clamping
	// how many times the UpdatePhysics can be called in
	// a single game loop.
	if (accumulator > 0.2f)
		accumulator = 0.2f;

	while (accumulator >= dt)
	{
		Step_CL(_numWidgets, _widgets);
		accumulator -= dt;
	}

	const float alpha = accumulator / dt;
}

void bPhysics::Step_CL(int numWidgets, widget*& _widgets)
{
	int numContacts = 0;
	manifold_cl* contacts_cl = (manifold_cl*)malloc(numWidgets * numWidgets * sizeof(manifold_cl));
	contacts_cl = d_contacts;
	contacts.clear();
	updateBuffers_CL(_widgets, numWidgets, numContacts, contacts_cl);
	runGeneratePairsCL(_widgets, numWidgets, numContacts, contacts_cl);
	this->ComputePhysics_CL(numWidgets, numContacts, _widgets, contacts_cl);
}

bool bPhysics::bCL_Init(widget*& _widgets, int sizeWidgets)
{
	/* Get Platform and Device Info */
	ret = clGetPlatformIDs(1, &platform_id, &ret_num_platforms);
	if (ret < 0)
		return false;

	ret = clGetDeviceIDs(platform_id, CL_DEVICE_TYPE_GPU, 0, NULL, &ret_num_devices);
	//ret = clGetDeviceIDs(platform_id, CL_DEVICE_TYPE_DEFAULT, 1, &device_id, &ret_num_devices);
	if (ret < 0)
		return false;

	std::cout << " # of devices = " << ret_num_devices << std::endl;
	cdDevices = (cl_device_id*)malloc(ret_num_devices * sizeof(cl_device_id));
	ret = clGetDeviceIDs(platform_id, CL_DEVICE_TYPE_GPU, ret_num_devices, cdDevices, NULL);
	if (ret < 0)
		return false;
	uiTargetDevice = glm::clamp((int)uiTargetDevice, (int)0, (int)(ret_num_devices - 1));

	std::cout << "Using device #: " << uiTargetDevice << std::endl;
	clGetDeviceInfo(cdDevices[uiTargetDevice], CL_DEVICE_MAX_COMPUTE_UNITS, sizeof(uiNumComputeUnits), &uiNumComputeUnits, NULL);
	std::cout << " # of Compute Units = " << uiNumComputeUnits << std::endl;


	/* Create OpenCL context */
	context = clCreateContext(NULL, 1, &cdDevices[uiTargetDevice], NULL, NULL, &ret);
	if (ret < 0)
		return false;

	/* Create Command Queue */
	command_queue = clCreateCommandQueue(context, cdDevices[uiTargetDevice], 0, &ret);
	if (ret < 0)
		return false;

	d_widgets = (widget*)malloc(sizeWidgets * sizeof(widget));
	d_contacts = (manifold_cl*)malloc(sizeWidgets * sizeWidgets * sizeof(manifold_cl));

  // initialize contacts array
	int max = (sizeWidgets*sizeWidgets);
	for (int i = 0; i < max; i++)
	{
		manifold_cl _m;
		_m.index = -1;
		_m.isColliding_cl = false;
		d_contacts[i] = _m;
	}

	d_contacts_size = 0;

	char string[MEM_SIZE];
	fileName = "broadphase.cl";
	/* Load the source code containing the kernel*/
	fopen_s(&fp, fileName, "r");
	if (!fp) {
		fprintf(stderr, "Failed to load kernel.\n");
		exit(1);
	}
	source_str = (char*)malloc(MAX_SOURCE_SIZE);
	source_size = fread(source_str, 1, MAX_SOURCE_SIZE, fp);
	fclose(fp);


	/* Create Kernel Program from the source */
	program = clCreateProgramWithSource(context, 1, (const char **)&source_str,
		(const size_t *)&source_size, &ret);

	if (ret < 0)
		return false;

	/* Build Kernel Program */
	ret = clBuildProgram(program, 1, &cdDevices[uiTargetDevice], NULL, NULL, NULL);
	if (ret < 0)
		return false;

	/* Create OpenCL Kernel */
	kernel = clCreateKernel(program, "broadphase_kernel", &ret);
	//kernel = clCreateKernel(program, "test_kernel", &ret);
	if (ret < 0)
		return false;

	manifold_cl* _contacts = (manifold_cl*)malloc((sizeWidgets * sizeWidgets) * sizeof(manifold_cl));
	_contacts = d_contacts;
	int sizeContacts = 0;

	/* Create Memory Buffer */
	memobj_widgets = clCreateBuffer(context, CL_MEM_READ_WRITE, sizeWidgets * sizeof(widget), NULL, &ret);
	if (ret < 0)
		return false;

	memobj_contacts = clCreateBuffer(context, CL_MEM_READ_WRITE, (sizeWidgets * sizeWidgets) * sizeof(manifold_cl), NULL, &ret);
	if (ret < 0)
		return false;

	memobj_contacts_size = clCreateBuffer(context, CL_MEM_READ_WRITE, 1 * sizeof(int), NULL, &ret);
	if (ret < 0)
		return false;
	memobj_widgets_size = clCreateBuffer(context, CL_MEM_READ_WRITE, 1 * sizeof(int), NULL, &ret);
	if (ret < 0)
		return false;
}

bool bPhysics::updateBuffers_CL(widget*& _widgets, int sizeWidgets, int& sizeContacts, manifold_cl*& _contacts)
{
	int* d_size = (int*)malloc(1 * sizeof(int));
	d_size[0] = sizeWidgets;

	int* d_size_contacts = (int*)malloc(1 * sizeof(int));
	d_size_contacts[0] = sizeContacts;
	/* Copy input data to the memory buffer */
	ret = clEnqueueWriteBuffer(command_queue, memobj_widgets, CL_TRUE, 0, sizeWidgets * sizeof(widget), _widgets, 0, NULL, NULL);
	if (ret < 0)
		return false;
	ret = clEnqueueWriteBuffer(command_queue, memobj_contacts, CL_TRUE, 0, (sizeWidgets * sizeWidgets) * sizeof(manifold_cl), _contacts, 0, NULL, NULL);
	if (ret < 0)
		return false;
	ret = clEnqueueWriteBuffer(command_queue, memobj_widgets_size, CL_TRUE, 0, 1 * sizeof(int), d_size, 0, NULL, NULL);
	if (ret < 0)
		return false;
	ret = clEnqueueWriteBuffer(command_queue, memobj_contacts_size, CL_TRUE, 0, 1 * sizeof(int), d_size_contacts, 0, NULL, NULL);
	if (ret < 0)
		return false;

	/* Set OpenCL Kernel Parameters */
	ret = clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *)&memobj_widgets);
	if (ret < 0)
		return false;

	ret = clSetKernelArg(kernel, 1, sizeof(cl_mem), (void *)&memobj_contacts);
	if (ret < 0)
		return false;
	ret = clSetKernelArg(kernel, 2, sizeof(int), &memobj_widgets_size);
	if (ret < 0)
		return false;
	ret = clSetKernelArg(kernel, 3, sizeof(cl_mem), (void *)&memobj_contacts_size);
	if (ret < 0)
		return false;
}

bool bPhysics::runGeneratePairsCL(widget*& _widgets, int sizeWidgets, int& sizeContacts, manifold_cl*& _contacts)
{
	/* Execute OpenCL Kernel */
	//ret = clEnqueueTask(command_queue, kernel, 0, NULL, NULL);

	const size_t dimSize = 2;
	size_t global_item_size[dimSize];
	global_item_size[0] = (sizeWidgets);
	global_item_size[1] = sizeWidgets;
	//size_t local_item_size = sizeWidgets;

	/* Execute OpenCL kernel as data parallel */
	ret = clEnqueueNDRangeKernel(command_queue, kernel, dimSize, NULL, global_item_size, NULL, 0, NULL, NULL);
	//ret = clEnqueueTask(command_queue, kernel, NULL, NULL, NULL);
	if (ret < 0)
		return false;

	/* Copy results from the memory buffer */
	/* Transfer result to host */
	ret = clEnqueueReadBuffer(command_queue, memobj_widgets, CL_TRUE, 0, sizeWidgets * sizeof(widget), _widgets, 0, NULL, NULL);
	ret = clEnqueueReadBuffer(command_queue, memobj_contacts, CL_TRUE, 0, (sizeWidgets * sizeWidgets) * sizeof(manifold_cl), _contacts, 0, NULL, NULL);
	ret = clEnqueueReadBuffer(command_queue, memobj_contacts_size, CL_TRUE, 0, 1 * sizeof(int), &sizeContacts, 0, NULL, NULL);
	if (ret < 0)
		return false;

	ret = clFinish(command_queue);
	if (ret < 0)
		return false;
	/* Display Result */
	//puts(string);
	// Sort
	manifold_cl* temp_contacts = (manifold_cl*)malloc(sizeWidgets * sizeWidgets * sizeof(manifold_cl));
	int counter = 0;
	int inv_counter = (sizeWidgets*sizeWidgets)-1;

	if (sizeContacts > 0)
	{
		std::cout << "Contact! " << std::endl;
	}

	for (int i = 0; i < (sizeWidgets*sizeWidgets); i++)
	{
		std::cout << "Index: " << _contacts[i].index << std::endl;
		// this doesn't work yet because the "valid" entries opencl deletes
		if (_contacts[i].index > -1)
		{
			temp_contacts[counter] = _contacts[i];
			counter++;
		}
		else
		{
			temp_contacts[inv_counter] = _contacts[i];
			inv_counter--;
		}

	}
	free(temp_contacts);

	return true;
}

bPhysics.h

	// OpenCL
	void GeneratePairs_CL(manifold_cl*& _contacts, widget*& _widgets, int _WidgetSize, int& _contactsSize);
	void FixedUpdate_CL(int _numWidgets, widget*& _widgets);
	void ComputePhysics_CL(int numWidgets, int& numContacts, widget*& _widgets, manifold_cl*& _contacts);
	void Step_CL(int numWidgets, widget*& _widgets);

	// more
	cl_device_id device_id = NULL;
	cl_context context = NULL;				// OpenCL Context
	cl_command_queue command_queue = NULL;	// OpenCL Command Queue
	cl_mem memobj_widgets = NULL;
	cl_mem memobj_widgets_size = NULL;
	cl_mem memobj_contacts = NULL;
	cl_mem memobj_contacts_size = NULL;
	cl_program program = NULL;
	cl_kernel kernel = NULL;
	cl_device_id *cdDevices = NULL;     // OpenCL device list
	cl_platform_id platform_id = NULL;		// OpenCL Platform
	cl_uint ret_num_devices;
	cl_uint ret_num_platforms;
	cl_int ret;
	cl_uint uiNumComputeUnits;
	cl_uint uiTargetDevice = 0;	        // OpenCL Device to compute on
	//cl_int err;
	size_t logSize;
	char *programLog;
	cl_build_status status;

	FILE *fp;
	char *fileName;
	char *source_str;
	size_t source_size;

	widget* d_widgets;
	manifold_cl* d_contacts;
	int d_contacts_size;
	int d_widgets_size;

	bool bCL_Init(widget*& _widgets, int sizeWidgets);
	bool cleanupCL();
	bool opencl_ReadKernel(int sizeWidgets);
	bool updateBuffers_CL(widget*& _widgets, int sizeWidgets, int& sizeContacts, manifold_cl*& _contacts);
	bool runGeneratePairsCL(widget*& _widgets, int sizeWidgets, int& sizeContacts, manifold_cl*& _contacts);

broadphase.cl

	#pragma OPENCL EXTENSION cl_khr_global_int32_base_atomics : enable
	#pragma OPENCL EXTENSION cl_khr_local_int32_base_atomics : enable
	__kernel void broadphase_kernel(__global widget *_widgets, __global manifold_cl *_contacts, int _WidgetSize, __global int* _contactsSize)
	{

		int i = get_global_id(0);
		int j = get_global_id(1);
		int index;
		index = (_WidgetSize * get_global_id(1)) + get_global_id(0);

		manifold_cl _m;
		bVec3 temp;
		temp.x = 0;
		temp.y = 0;
		temp.z = 0;
		_m.normal_cl = temp;
		_m.penetration_cl = 0;
		_m.A_cl = i;
		_m.B_cl = j;
		_m.isColliding_cl = false;

		_m.index = -1;
		_m.gid = index;
				_contacts[index] = _m;
		atom_inc(&_contactsSize[0]);
	}

bWidget.h

typedef struct tag_bVec3 {
	cl_float x;
	cl_float y;
	cl_float z;
} bVec3;

typedef struct tag_transform_cl {
	bVec3 position_cl;
	cl_float orientation_in_radians_cl;
	//glm::mat4 translation_cl;
	//glm::mat4 rotation_cl;
	//glm::mat4 scale_cl;
} transform_cl;

typedef struct tag_widget {
	cl_int id_cl;

	cl_int shape_cl;
	bool isStatic_cl;
	cl_float density_cl;
	cl_float inv_intertia_cl;
	cl_float inertia_cl;
	cl_float inv_mass_cl;
	cl_float mass_cl;
	bVec3 force_cl;
	bVec3 velocity_cl;
	//glm::vec3 force;
	//glm::vec3 velocity;
	cl_float angularVelocity_cl;
	cl_float torque_cl;
	cl_float impulse_cl;
	cl_float restitution_cl;
	cl_float staticFriction_cl;
	cl_float dynamicFriction_cl;
	cl_float width_cl;
	cl_float height_cl;
	transform_cl t_cl;
} widget;

// manifold
typedef struct tag_manifold_cl {
	cl_int A_cl; // widget_A_index;
	cl_int B_cl; // widget_B_index;
	int gid;
	int index;
	cl_float sf_cl;
	cl_float df_cl;
	cl_float penetration_cl;
	cl_float e_cl;
	bVec3 normal_cl;      // From A to B
	//glm::vec3 normal;
	bool isColliding_cl;
} manifold_cl;

do 
{
    double x, y;
    glfwGetCursorPos(window, &x, &y);
    glfwGetWindowSize(window, &new_width, &new_height);     
    float newAspect = (float)new_width / (float)new_height;
    main_scene.aspect = 1;

    glm::vec4 updatedScalePair = ComputeAspectRatio(new_width, new_height, originalAspect);

    glViewport(updatedScalePair.z, updatedScalePair.w, WIDTH * updatedScalePair.x, HEIGHT * updatedScalePair.y);
    glScissor(updatedScalePair.z, updatedScalePair.w, WIDTH * updatedScalePair.x, HEIGHT * updatedScalePair.y);
    glEnable(GL_SCISSOR_TEST);
    //glViewport(0, 0, new_width, new_height);
    main_scene.ProjectionMatrix = glm::ortho(-((float)WIDTH / 2) / originalAspect, ((float)WIDTH / 2) / originalAspect, -((float)HEIGHT / 2) / originalAspect, ((float)HEIGHT / 2) / originalAspect, 0.0f, 100.0f);

    tempWidgets = main_scene.getWidgetList();
    tempBodies = main_scene.getBodyList();

    // Clear the screen
    glClear(GL_COLOR_BUFFER_BIT | GL_DEPTH_BUFFER_BIT);
    mPhysicsController.bodies = tempBodies;
    main_scene.setWidgetList(tempWidgets);
    main_scene.setBodiesList(tempBodies);

    mPhysicsController.FixedUpdate_CL(objID + 1, widgets);
    main_scene.renderScene_CL(objID + 1, widgets);
    // Bind our texture in Texture Unit 0
    glActiveTexture(GL_TEXTURE0);
    glBindTexture(GL_TEXTURE_2D, Texture);
    // Set our "myTextureSampler" sampler to user Texture Unit 0
    glUniform1i(TextureID, 0);

    // Swap buffers
    glfwSwapBuffers(window);
    glfwPollEvents();

} // Check if the ESC key was pressed or the window was closed
while (glfwGetKey(window, GLFW_KEY_ESCAPE) != GLFW_PRESS &&
    glfwWindowShouldClose(window) == 0);