Browse Source

(committing very old stuff) Did some stuff with OpenCL once

master
Joshua Moerman 8 years ago
commit
bb39483eee
  1. 3
      .gitignore
  2. 26
      Fractal.fsh
  3. 18
      Fractal.vsh
  4. 8
      Kernel.cl
  5. 38
      NSGLWrapper.hpp
  6. 95
      NSGLWrapper.mm
  7. 41
      NSWrapper.hpp
  8. 216
      NSWrapper.mm
  9. 297
      XcodeOpenCL.xcodeproj/project.pbxproj
  10. 12452
      cl.hpp
  11. 76
      cl2.hpp
  12. 215
      main.cpp

3
.gitignore

@ -0,0 +1,3 @@
*.xcworkspace
xcuserdata

26
Fractal.fsh

@ -0,0 +1,26 @@
#version 330
uniform float rotation;
in vec2 pos;
in vec2 start;
out vec4 fragColor;
void main(){
vec2 z = start;
int i = 0;
for (i = 0; i < 30; i++) {
vec2 zsq = z*z;
if(zsq.x + zsq.y > 16.0) break;
float t = zsq.x - zsq.y + pos.x;
z.y = 2.0*z.x*z.y + pos.y;
z.x = t;
}
fragColor = vec4(float(i) / 30.0);
fragColor.bg = 0.5 * sin(z) + 0.5;
}

18
Fractal.vsh

@ -0,0 +1,18 @@
#version 330
uniform float rotation;
in vec4 position;
in vec4 color;
out vec2 pos;
out vec2 start;
void main(){
pos = position.xy;
pos.x -= 0.5;
start = sqrt(1.0 + 0.01 * rotation) * 0.3 * sin(vec2(0.1, 0.1337) * rotation);
gl_Position = position;
}

8
Kernel.cl

@ -0,0 +1,8 @@
kernel void square(global float* input, size_t width, global float* output){
size_t x = get_global_id(0);
size_t y = get_global_id(1);
float i = 2.0 * (input[x + width*y] - 0.5);
output[x + width*y] = 0.5 + 0.5*sin((1.0 - i) * sin(i*i + 6.2*float(x)/width) + cos(1.0/i + 6.2*float(y)/width));
}

38
NSGLWrapper.hpp

@ -0,0 +1,38 @@
//
// NSGLWrapper.hpp
// XcodeOpenCL
//
// Created by Joshua Moerman on 13/04/14.
//
//
#pragma once
#include <memory>
#include <functional>
#include <OpenGL/OpenGL.h>
#include <CoreVideo/CVDisplayLink.h>
struct GLContext {
GLContext();
~GLContext();
void set_as_current_context() const;
void lock();
void unlock();
CGLContextObj ctx{};
CGLPixelFormatObj pix{};
};
struct CVDisplayLinky {
CVDisplayLinky(GLContext const &);
~CVDisplayLinky();
void start();
void stop();
std::function<void()> callback;
CVDisplayLinkRef displayLink;
};

95
NSGLWrapper.mm

@ -0,0 +1,95 @@
//
// NSGLWrapper.mm
// XcodeOpenCL
//
// Created by Joshua Moerman on 13/04/14.
//
//
#include "NSGLWrapper.hpp"
#include <iostream>
#include <stdexcept>
#import <Cocoa/Cocoa.h>
#import <OpenGL/OpenGL.h>
static void check_cgl(CGLError e) {
if(e != kCGLNoError){
throw std::runtime_error(CGLErrorString(e));
}
}
GLContext::GLContext() {
GLint npix{};
const CGLPixelFormatAttribute attributes[]{
kCGLPFADoubleBuffer,
kCGLPFADepthSize, static_cast<CGLPixelFormatAttribute>(24),
kCGLPFAOpenGLProfile, static_cast<CGLPixelFormatAttribute>(kCGLOGLPVersion_3_2_Core),
static_cast<CGLPixelFormatAttribute>(0)
};
check_cgl(CGLChoosePixelFormat(attributes, &pix, &npix));
check_cgl(CGLCreateContext(pix, 0, &ctx));
GLint swap_interval = 1;
check_cgl(CGLSetParameter(ctx, kCGLCPSwapInterval, &swap_interval));
set_as_current_context();
}
GLContext::~GLContext() {
CGLReleaseContext(ctx);
CGLReleasePixelFormat(pix);
}
void GLContext::set_as_current_context() const {
check_cgl(CGLSetCurrentContext(ctx));
}
void GLContext::lock() {
check_cgl(CGLLockContext(ctx));
}
void GLContext::unlock() {
check_cgl(CGLUnlockContext(ctx));
}
static void check_cv(CVReturn r) {
if(r != kCVReturnSuccess){
throw std::runtime_error("Some CV display link issue");
}
}
static CVReturn CVDisplayLinkyInvokeCallback(CVDisplayLinkRef displayLink, const CVTimeStamp* now, const CVTimeStamp* outputTime, CVOptionFlags flagsIn, CVOptionFlags* flagsOut, void* displayLinkContext) {
if(!displayLinkContext){
throw std::runtime_error("There is no context in the display link");
}
auto& foo = *(CVDisplayLinky*)displayLinkContext;
if(foo.callback){
foo.callback();
}
return kCVReturnSuccess;
}
CVDisplayLinky::CVDisplayLinky(GLContext const & c) {
check_cv(CVDisplayLinkCreateWithActiveCGDisplays(&displayLink));
check_cv(CVDisplayLinkSetOutputCallback(displayLink, &CVDisplayLinkyInvokeCallback, this));
check_cv(CVDisplayLinkSetCurrentCGDisplayFromOpenGLContext(displayLink, c.ctx, c.pix));
}
CVDisplayLinky::~CVDisplayLinky() {
stop();
CVDisplayLinkRelease(displayLink);
}
void CVDisplayLinky::start() {
check_cv(CVDisplayLinkStart(displayLink));
}
void CVDisplayLinky::stop() {
check_cv(CVDisplayLinkStop(displayLink));
}

41
NSWrapper.hpp

@ -0,0 +1,41 @@
//
// NSWrapper.hpp
// XcodeOpenCL
//
// Created by Joshua Moerman on 13/04/14.
//
//
#pragma once
#include <memory>
#include <functional>
#include <CoreGraphics/CoreGraphics.h>
struct GLContext;
struct ContextParameters{
GLContext & context;
};
struct GLViewFunctionality {
// All three will only be called with the current gl context already set
// May be called from different threads
std::function<void(ContextParameters)> initialize_callback;
std::function<void(ContextParameters)> draw_callback;
std::function<void(ContextParameters, CGFloat, CGFloat)> resize_callback;
};
// returns a GLViewFunctionality with only its draw function set.
GLViewFunctionality simple_draw(std::function<void()> f);
struct NSAppWrapper {
NSAppWrapper();
~NSAppWrapper();
void run();
void create_window(GLViewFunctionality const &);
std::unique_ptr<struct NSAppWrapperImpl> impl;
};

216
NSWrapper.mm

@ -0,0 +1,216 @@
//
// NSWrapper.mm
// XcodeOpenCL
//
// Created by Joshua Moerman on 13/04/14.
//
//
#include "NSWrapper.hpp"
#include "NSGLWrapper.hpp"
#include <vector>
#import <Cocoa/Cocoa.h>
#import <GLKit/GLKit.h>
GLViewFunctionality simple_draw(std::function<void()> f){
return {nullptr, [f](ContextParameters){ f(); }, nullptr};
}
@interface GLView : NSOpenGLView
// All three will only be called with the current gl context already set
// May be called from different threads (but the properties are atomic by default)
@property GLViewFunctionality functionality;
@end
struct NSAppWrapperImpl {
NSAppWrapperImpl() {
app = [NSApplication sharedApplication];
app.activationPolicy = NSApplicationActivationPolicyRegular;
appName = [[NSProcessInfo processInfo] processName];
}
void set_up_menu() {
NSMenu* appMenu = [NSMenu new];
[appMenu addItemWithTitle:[@"Quit " stringByAppendingString:appName]
action:@selector(terminate:)
keyEquivalent:@"q"];
NSMenuItem * menuItem = [NSMenuItem new];
menuItem.submenu = appMenu;
app.mainMenu = [NSMenu new];
[app.mainMenu addItem:menuItem];
}
void create_window(GLViewFunctionality const & f) {
NSWindow * window = [[NSWindow alloc] initWithContentRect:NSMakeRect(0, 0, 500, 500)
styleMask:NSTitledWindowMask | NSResizableWindowMask | NSMiniaturizableWindowMask
backing:NSBackingStoreBuffered
defer:YES];
window.title = appName;
window.collectionBehavior |= NSWindowCollectionBehaviorFullScreenPrimary;
[window center];
[window makeKeyAndOrderFront:nil];
NSRect frame = [window contentRectForFrameRect:window.frame];
frame.origin = {0, 0};
GLView * view = [[GLView alloc] initWithFrame:frame];
view.functionality = f;
window.contentView = view;
windows.push_back(window);
}
void run() {
[app run];
}
NSApplication * app;
std::vector<NSWindow *> windows;
NSString * appName;
};
NSAppWrapper::NSAppWrapper()
: impl(new NSAppWrapperImpl){
impl->set_up_menu();
}
NSAppWrapper::~NSAppWrapper() = default;
void NSAppWrapper::run() {
impl->run();
}
void NSAppWrapper::create_window(GLViewFunctionality const & f){
impl->create_window(f);
}
/*
Some general notes about GLView:
- It is in a multithreaded context. The CVDisplayLink will run on a diferent thread than the main thread. However the main thread *will also draw* (eg. when reshaping, to avoid flickering). So there are some locks here and there.
- Hopefully the C++ objects are destructed nicely
- There is syntax for the constructor, so I used unique_ptr...
*/
@implementation GLView {
std::unique_ptr<GLContext> context;
std::unique_ptr<CVDisplayLinky> linky;
}
@synthesize functionality;
- (id) initWithFrame:(NSRect)frameRect {
if(self = [super initWithFrame:frameRect]){
context = std::unique_ptr<GLContext>(new GLContext);
self.pixelFormat = [[NSOpenGLPixelFormat alloc] initWithCGLPixelFormatObj: context->pix];
self.openGLContext = [[NSOpenGLContext alloc] initWithCGLContextObj: context->ctx];
self.wantsBestResolutionOpenGLSurface = YES; // Opt-in for retina
}
return self;
}
- (void) prepareOpenGL{
[super prepareOpenGL];
context->set_as_current_context();
// Synchronize buffer swaps with vertical refresh rate
GLint swapInt = 1;
[[self openGLContext] setValues:&swapInt forParameter:NSOpenGLCPSwapInterval];
if(functionality.initialize_callback){
functionality.initialize_callback({*context});
}
// Create a display link capable of being used with all active displays
linky = std::unique_ptr<CVDisplayLinky>(new CVDisplayLinky(*context));
linky->callback = [=]{
// will be called in a different thread => need new autoreleasepool
@autoreleasepool {
[self drawView];
}
};
linky->start();
// Register to be notified when the window closes so we can stop the displaylink
[[NSNotificationCenter defaultCenter] addObserver:self
selector:@selector(windowWillClose:)
name:NSWindowWillCloseNotification
object:[self window]];
}
- (void) windowWillClose:(NSNotification*)notification{
// The default OpenGL render buffers will be destroyed. If display link
// continues to fire without renderbuffers, OpenGL draw calls will set errors.
linky->stop();
}
- (void) reshape{
[super reshape];
context->lock();
// [NSView convertRectToBacking] converts point sizes to pixel sizes.
// viewRectPixels will be larger (2x) than viewRectPoints for retina displays.
// viewRectPixels will be the same as viewRectPoints for non-retina displays
NSRect viewRectPixels = [self convertRectToBacking:self.bounds];
if(functionality.resize_callback){
functionality.resize_callback({*context}, viewRectPixels.size.width, viewRectPixels.size.height);
}
context->unlock();
}
- (void)renewGState{
// Called whenever graphics state updated (such as window resize)
// OpenGL rendering is not synchronous with other rendering on the OSX.
// Therefore, call disableScreenUpdatesUntilFlush so the window server
// doesn't render non-OpenGL content in the window asynchronously from
// OpenGL content, which could cause flickering. (non-OpenGL content
// includes the title bar and drawing done by the app with other APIs)
[[self window] disableScreenUpdatesUntilFlush];
[super renewGState];
}
- (void) drawRect: (NSRect) theRect{
[self drawView];
}
- (void) drawView{
context->set_as_current_context();
// We draw on a secondary thread through the display link
// When resizing the view, -reshape is called automatically on the main
// thread. Add a mutex around to avoid the threads accessing the context
// simultaneously when resizing
context->lock();
NSRect viewRectPixels = [self convertRectToBacking:self.bounds];
glBindFramebuffer(GL_FRAMEBUFFER, 0);
glViewport(0, 0, (int)viewRectPixels.size.width, (int)viewRectPixels.size.height);
if(functionality.draw_callback){
functionality.draw_callback({*context});
}
CGLFlushDrawable(context->ctx);
context->unlock();
}
@end

297
XcodeOpenCL.xcodeproj/project.pbxproj

@ -0,0 +1,297 @@
// !$*UTF8*$!
{
archiveVersion = 1;
classes = {
};
objectVersion = 46;
objects = {
/* Begin PBXBuildFile section */
4258B19C1C864DE00050E2BE /* libpng.a in Frameworks */ = {isa = PBXBuildFile; fileRef = 4258B19B1C864DE00050E2BE /* libpng.a */; };
428B574618E584F200A6D212 /* main.cpp in Sources */ = {isa = PBXBuildFile; fileRef = 428B573618E584A100A6D212 /* main.cpp */; };
428B574B18E58D8800A6D212 /* OpenCL.framework in Frameworks */ = {isa = PBXBuildFile; fileRef = 428B574A18E58D8800A6D212 /* OpenCL.framework */; };
42F63BB418FA943500F2C81A /* Cocoa.framework in Frameworks */ = {isa = PBXBuildFile; fileRef = 42F63BB318FA943500F2C81A /* Cocoa.framework */; };
42F63BB618FA943900F2C81A /* Foundation.framework in Frameworks */ = {isa = PBXBuildFile; fileRef = 42F63BB518FA943900F2C81A /* Foundation.framework */; };
42F63BB918FAAE4F00F2C81A /* NSWrapper.mm in Sources */ = {isa = PBXBuildFile; fileRef = 42F63BB718FAAE4F00F2C81A /* NSWrapper.mm */; };
42F63BBC18FAB2C500F2C81A /* NSGLWrapper.mm in Sources */ = {isa = PBXBuildFile; fileRef = 42F63BBA18FAB2C500F2C81A /* NSGLWrapper.mm */; };
42F63BBE18FAC32100F2C81A /* CoreGraphics.framework in Frameworks */ = {isa = PBXBuildFile; fileRef = 42F63BBD18FAC32100F2C81A /* CoreGraphics.framework */; };
42F63BC018FAC32700F2C81A /* OpenGL.framework in Frameworks */ = {isa = PBXBuildFile; fileRef = 42F63BBF18FAC32700F2C81A /* OpenGL.framework */; };
42F63BC218FAC37F00F2C81A /* QuartzCore.framework in Frameworks */ = {isa = PBXBuildFile; fileRef = 42F63BC118FAC37F00F2C81A /* QuartzCore.framework */; };
42F63BC718FB158500F2C81A /* Fractal.fsh in CopyFiles */ = {isa = PBXBuildFile; fileRef = 42F63BC318FB13AF00F2C81A /* Fractal.fsh */; };
42F63BC818FB158500F2C81A /* Fractal.vsh in CopyFiles */ = {isa = PBXBuildFile; fileRef = 42F63BC418FB13AF00F2C81A /* Fractal.vsh */; };
42F63BCA18FB17FE00F2C81A /* Kernel.cl in Sources */ = {isa = PBXBuildFile; fileRef = 42F63BC918FB17FE00F2C81A /* Kernel.cl */; };
42F63BCB18FB181F00F2C81A /* Kernel.cl in CopyFiles */ = {isa = PBXBuildFile; fileRef = 42F63BC918FB17FE00F2C81A /* Kernel.cl */; };
/* End PBXBuildFile section */
/* Begin PBXCopyFilesBuildPhase section */
428B573A18E584DB00A6D212 /* CopyFiles */ = {
isa = PBXCopyFilesBuildPhase;
buildActionMask = 12;
dstPath = "";
dstSubfolderSpec = 16;
files = (
42F63BCB18FB181F00F2C81A /* Kernel.cl in CopyFiles */,
42F63BC718FB158500F2C81A /* Fractal.fsh in CopyFiles */,
42F63BC818FB158500F2C81A /* Fractal.vsh in CopyFiles */,
);
runOnlyForDeploymentPostprocessing = 0;
};
/* End PBXCopyFilesBuildPhase section */
/* Begin PBXFileReference section */
4249DFE518F87EB3007C4554 /* libpng15.a */ = {isa = PBXFileReference; lastKnownFileType = archive.ar; name = libpng15.a; path = ../../../../../usr/local/lib/libpng15.a; sourceTree = "<group>"; };
4258B19B1C864DE00050E2BE /* libpng.a */ = {isa = PBXFileReference; lastKnownFileType = archive.ar; name = libpng.a; path = ../../../../../usr/local/lib/libpng.a; sourceTree = "<group>"; };
428B573618E584A100A6D212 /* main.cpp */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.cpp.cpp; path = main.cpp; sourceTree = "<group>"; };
428B573C18E584DB00A6D212 /* main */ = {isa = PBXFileReference; explicitFileType = "compiled.mach-o.executable"; includeInIndex = 0; path = main; sourceTree = BUILT_PRODUCTS_DIR; };
428B574A18E58D8800A6D212 /* OpenCL.framework */ = {isa = PBXFileReference; lastKnownFileType = wrapper.framework; name = OpenCL.framework; path = System/Library/Frameworks/OpenCL.framework; sourceTree = SDKROOT; };
428B574C18E5C2D500A6D212 /* cl.hpp */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.cpp.h; path = cl.hpp; sourceTree = "<group>"; };
428B574D18E5D19600A6D212 /* cl2.hpp */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.cpp.h; path = cl2.hpp; sourceTree = "<group>"; };
42F63BB318FA943500F2C81A /* Cocoa.framework */ = {isa = PBXFileReference; lastKnownFileType = wrapper.framework; name = Cocoa.framework; path = System/Library/Frameworks/Cocoa.framework; sourceTree = SDKROOT; };
42F63BB518FA943900F2C81A /* Foundation.framework */ = {isa = PBXFileReference; lastKnownFileType = wrapper.framework; name = Foundation.framework; path = System/Library/Frameworks/Foundation.framework; sourceTree = SDKROOT; };
42F63BB718FAAE4F00F2C81A /* NSWrapper.mm */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.cpp.objcpp; path = NSWrapper.mm; sourceTree = "<group>"; };
42F63BB818FAAE4F00F2C81A /* NSWrapper.hpp */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.cpp.h; path = NSWrapper.hpp; sourceTree = "<group>"; };
42F63BBA18FAB2C500F2C81A /* NSGLWrapper.mm */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.cpp.objcpp; path = NSGLWrapper.mm; sourceTree = "<group>"; };
42F63BBB18FAB2C500F2C81A /* NSGLWrapper.hpp */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.cpp.h; path = NSGLWrapper.hpp; sourceTree = "<group>"; };
42F63BBD18FAC32100F2C81A /* CoreGraphics.framework */ = {isa = PBXFileReference; lastKnownFileType = wrapper.framework; name = CoreGraphics.framework; path = System/Library/Frameworks/CoreGraphics.framework; sourceTree = SDKROOT; };
42F63BBF18FAC32700F2C81A /* OpenGL.framework */ = {isa = PBXFileReference; lastKnownFileType = wrapper.framework; name = OpenGL.framework; path = System/Library/Frameworks/OpenGL.framework; sourceTree = SDKROOT; };
42F63BC118FAC37F00F2C81A /* QuartzCore.framework */ = {isa = PBXFileReference; lastKnownFileType = wrapper.framework; name = QuartzCore.framework; path = System/Library/Frameworks/QuartzCore.framework; sourceTree = SDKROOT; };
42F63BC318FB13AF00F2C81A /* Fractal.fsh */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.glsl; path = Fractal.fsh; sourceTree = "<group>"; };
42F63BC418FB13AF00F2C81A /* Fractal.vsh */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.glsl; path = Fractal.vsh; sourceTree = "<group>"; };
42F63BC918FB17FE00F2C81A /* Kernel.cl */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.opencl; path = Kernel.cl; sourceTree = "<group>"; };
/* End PBXFileReference section */
/* Begin PBXFrameworksBuildPhase section */
428B573918E584DB00A6D212 /* Frameworks */ = {
isa = PBXFrameworksBuildPhase;
buildActionMask = 2147483647;
files = (
42F63BC218FAC37F00F2C81A /* QuartzCore.framework in Frameworks */,
42F63BC018FAC32700F2C81A /* OpenGL.framework in Frameworks */,
42F63BBE18FAC32100F2C81A /* CoreGraphics.framework in Frameworks */,
42F63BB618FA943900F2C81A /* Foundation.framework in Frameworks */,
42F63BB418FA943500F2C81A /* Cocoa.framework in Frameworks */,
428B574B18E58D8800A6D212 /* OpenCL.framework in Frameworks */,
4258B19C1C864DE00050E2BE /* libpng.a in Frameworks */,
);
runOnlyForDeploymentPostprocessing = 0;
};
/* End PBXFrameworksBuildPhase section */
/* Begin PBXGroup section */
428B572E18E5844100A6D212 = {
isa = PBXGroup;
children = (
4258B19B1C864DE00050E2BE /* libpng.a */,
42F63BC918FB17FE00F2C81A /* Kernel.cl */,
42F63BC318FB13AF00F2C81A /* Fractal.fsh */,
42F63BC418FB13AF00F2C81A /* Fractal.vsh */,
42F63BBA18FAB2C500F2C81A /* NSGLWrapper.mm */,
42F63BBB18FAB2C500F2C81A /* NSGLWrapper.hpp */,
42F63BB718FAAE4F00F2C81A /* NSWrapper.mm */,
42F63BB818FAAE4F00F2C81A /* NSWrapper.hpp */,
428B574D18E5D19600A6D212 /* cl2.hpp */,
428B574C18E5C2D500A6D212 /* cl.hpp */,
428B573618E584A100A6D212 /* main.cpp */,
4249DFE518F87EB3007C4554 /* libpng15.a */,
42F63BC118FAC37F00F2C81A /* QuartzCore.framework */,
42F63BB518FA943900F2C81A /* Foundation.framework */,
42F63BB318FA943500F2C81A /* Cocoa.framework */,
42F63BBD18FAC32100F2C81A /* CoreGraphics.framework */,
42F63BBF18FAC32700F2C81A /* OpenGL.framework */,
428B574A18E58D8800A6D212 /* OpenCL.framework */,
428B573D18E584DB00A6D212 /* Products */,
);
sourceTree = "<group>";
};
428B573D18E584DB00A6D212 /* Products */ = {
isa = PBXGroup;
children = (
428B573C18E584DB00A6D212 /* main */,
);
name = Products;
sourceTree = "<group>";
};
/* End PBXGroup section */
/* Begin PBXNativeTarget section */
428B573B18E584DB00A6D212 /* main */ = {
isa = PBXNativeTarget;
buildConfigurationList = 428B574318E584DB00A6D212 /* Build configuration list for PBXNativeTarget "main" */;
buildPhases = (
428B573818E584DB00A6D212 /* Sources */,
428B573918E584DB00A6D212 /* Frameworks */,
428B573A18E584DB00A6D212 /* CopyFiles */,
);
buildRules = (
);
dependencies = (
);
name = main;
productName = main;
productReference = 428B573C18E584DB00A6D212 /* main */;
productType = "com.apple.product-type.tool";
};
/* End PBXNativeTarget section */
/* Begin PBXProject section */
428B572F18E5844100A6D212 /* Project object */ = {
isa = PBXProject;
attributes = {
LastUpgradeCheck = 0510;
};
buildConfigurationList = 428B573218E5844100A6D212 /* Build configuration list for PBXProject "XcodeOpenCL" */;
compatibilityVersion = "Xcode 3.2";
developmentRegion = English;
hasScannedForEncodings = 0;
knownRegions = (
en,
);
mainGroup = 428B572E18E5844100A6D212;
productRefGroup = 428B573D18E584DB00A6D212 /* Products */;
projectDirPath = "";
projectRoot = "";
targets = (
428B573B18E584DB00A6D212 /* main */,
);
};
/* End PBXProject section */
/* Begin PBXSourcesBuildPhase section */
428B573818E584DB00A6D212 /* Sources */ = {
isa = PBXSourcesBuildPhase;
buildActionMask = 2147483647;
files = (
42F63BBC18FAB2C500F2C81A /* NSGLWrapper.mm in Sources */,
42F63BB918FAAE4F00F2C81A /* NSWrapper.mm in Sources */,
428B574618E584F200A6D212 /* main.cpp in Sources */,
42F63BCA18FB17FE00F2C81A /* Kernel.cl in Sources */,
);
runOnlyForDeploymentPostprocessing = 0;
};
/* End PBXSourcesBuildPhase section */
/* Begin XCBuildConfiguration section */
428B573318E5844100A6D212 /* Debug */ = {
isa = XCBuildConfiguration;
buildSettings = {
CLANG_CXX_LANGUAGE_STANDARD = "c++0x";
CLANG_CXX_LIBRARY = "libc++";
HEADER_SEARCH_PATHS = (
"$(inherited)",
/Applications/Xcode.app/Contents/Developer/Toolchains/XcodeDefault.xctoolchain/usr/include,
/usr/local/include,
);
LIBRARY_SEARCH_PATHS = /usr/local/lib;
OPENCL_FAST_RELAXED_MATH = YES;
OPENCL_MAD_ENABLE = YES;
OPENCL_OPTIMIZATION_LEVEL = 3;
RUN_CLANG_STATIC_ANALYZER = YES;
};
name = Debug;
};
428B573418E5844100A6D212 /* Release */ = {
isa = XCBuildConfiguration;
buildSettings = {
CLANG_CXX_LANGUAGE_STANDARD = "c++0x";
CLANG_CXX_LIBRARY = "libc++";
HEADER_SEARCH_PATHS = (
"$(inherited)",
/Applications/Xcode.app/Contents/Developer/Toolchains/XcodeDefault.xctoolchain/usr/include,
/usr/local/include,
);
LIBRARY_SEARCH_PATHS = /usr/local/lib;
OPENCL_FAST_RELAXED_MATH = YES;
OPENCL_MAD_ENABLE = YES;
OPENCL_OPTIMIZATION_LEVEL = 3;
RUN_CLANG_STATIC_ANALYZER = YES;
};
name = Release;
};
428B574418E584DB00A6D212 /* Debug */ = {
isa = XCBuildConfiguration;
buildSettings = {
CLANG_CXX_LANGUAGE_STANDARD = "gnu++0x";
CLANG_CXX_LIBRARY = "libc++";
CLANG_ENABLE_OBJC_ARC = YES;
CLANG_WARN_BOOL_CONVERSION = YES;
CLANG_WARN_CONSTANT_CONVERSION = YES;
CLANG_WARN_DIRECT_OBJC_ISA_USAGE = YES_ERROR;
CLANG_WARN_EMPTY_BODY = YES;
CLANG_WARN_ENUM_CONVERSION = YES;
CLANG_WARN_INT_CONVERSION = YES;
CLANG_WARN_OBJC_ROOT_CLASS = YES_ERROR;
CLANG_WARN__DUPLICATE_METHOD_MATCH = YES;
COPY_PHASE_STRIP = NO;
GCC_OPTIMIZATION_LEVEL = 0;
GCC_PREPROCESSOR_DEFINITIONS = (
"DEBUG=1",
"$(inherited)",
);
GCC_SYMBOLS_PRIVATE_EXTERN = NO;
GCC_WARN_64_TO_32_BIT_CONVERSION = YES;
GCC_WARN_ABOUT_RETURN_TYPE = YES_ERROR;
GCC_WARN_UNDECLARED_SELECTOR = YES;
GCC_WARN_UNINITIALIZED_AUTOS = YES_AGGRESSIVE;
GCC_WARN_UNUSED_FUNCTION = YES;
GCC_WARN_UNUSED_VARIABLE = YES;
MACOSX_DEPLOYMENT_TARGET = 10.10;
ONLY_ACTIVE_ARCH = YES;
PRODUCT_NAME = "$(TARGET_NAME)";
SDKROOT = macosx;
};
name = Debug;
};
428B574518E584DB00A6D212 /* Release */ = {
isa = XCBuildConfiguration;
buildSettings = {
CLANG_CXX_LANGUAGE_STANDARD = "gnu++0x";
CLANG_CXX_LIBRARY = "libc++";
CLANG_ENABLE_OBJC_ARC = YES;
CLANG_WARN_BOOL_CONVERSION = YES;
CLANG_WARN_CONSTANT_CONVERSION = YES;
CLANG_WARN_DIRECT_OBJC_ISA_USAGE = YES_ERROR;
CLANG_WARN_EMPTY_BODY = YES;
CLANG_WARN_ENUM_CONVERSION = YES;
CLANG_WARN_INT_CONVERSION = YES;
CLANG_WARN_OBJC_ROOT_CLASS = YES_ERROR;
CLANG_WARN__DUPLICATE_METHOD_MATCH = YES;
COPY_PHASE_STRIP = YES;
DEBUG_INFORMATION_FORMAT = "dwarf-with-dsym";
ENABLE_NS_ASSERTIONS = NO;
GCC_WARN_64_TO_32_BIT_CONVERSION = YES;
GCC_WARN_ABOUT_RETURN_TYPE = YES_ERROR;
GCC_WARN_UNDECLARED_SELECTOR = YES;
GCC_WARN_UNINITIALIZED_AUTOS = YES_AGGRESSIVE;
GCC_WARN_UNUSED_FUNCTION = YES;
GCC_WARN_UNUSED_VARIABLE = YES;
MACOSX_DEPLOYMENT_TARGET = 10.10;
PRODUCT_NAME = "$(TARGET_NAME)";
SDKROOT = macosx;
};
name = Release;
};
/* End XCBuildConfiguration section */
/* Begin XCConfigurationList section */
428B573218E5844100A6D212 /* Build configuration list for PBXProject "XcodeOpenCL" */ = {
isa = XCConfigurationList;
buildConfigurations = (
428B573318E5844100A6D212 /* Debug */,
428B573418E5844100A6D212 /* Release */,
);
defaultConfigurationIsVisible = 0;
defaultConfigurationName = Release;
};
428B574318E584DB00A6D212 /* Build configuration list for PBXNativeTarget "main" */ = {
isa = XCConfigurationList;
buildConfigurations = (
428B574418E584DB00A6D212 /* Debug */,
428B574518E584DB00A6D212 /* Release */,
);
defaultConfigurationIsVisible = 0;
defaultConfigurationName = Release;
};
/* End XCConfigurationList section */
};
rootObject = 428B572F18E5844100A6D212 /* Project object */;
}

12452
cl.hpp

File diff suppressed because it is too large

76
cl2.hpp

@ -0,0 +1,76 @@
//
// cl2.hpp
// XcodeOpenCL
//
// Created by Joshua Moerman on 28/03/14.
//
//
#ifndef XcodeOpenCL_cl2_hpp
#define XcodeOpenCL_cl2_hpp
#include "cl.hpp"
#include <cassert>
#include <iostream>
// debugging tool
static void check(cl_int err){
assert(err == CL_SUCCESS);
}
// simple variadic wrapper for kernels/arguments (1D)
struct KernelOp{
KernelOp(const cl::Program& program, const char* name, cl_int* err)
: kernel(program, name, err)
{}
template <typename... Args>
cl_int operator()(cl::CommandQueue & queue, size_t W, size_t H, Args&&... args){
return apply(queue, W, H, 0, std::forward<Args>(args)...);
}
private:
cl::Kernel kernel;
template <typename T, typename... Args>
cl_int apply(cl::CommandQueue & queue, size_t W, size_t H, cl_uint n, T&& t, Args&&... args){
check(kernel.setArg(n, t));
return apply(queue, W, H, n+1, std::forward<Args>(args)...);
}
cl_int apply(cl::CommandQueue & queue, size_t W, size_t H, size_t){
return queue.enqueueNDRangeKernel(kernel, cl::NullRange, cl::NDRange(W, H), cl::NullRange);
}
};
// We need to put the op<<'s in namespace cl (ADL)
namespace cl {
std::ostream& operator<<(std::ostream& out, Platform const & platform){
return out << platform.getInfo<CL_PLATFORM_NAME>() << ", " << platform.getInfo<CL_PLATFORM_VERSION>();
}
std::ostream& operator<<(std::ostream& out, Device const & device){
out << device.getInfo<CL_DEVICE_NAME>() << ", " << device.getInfo<CL_DEVICE_VERSION>() << ", " << device.getInfo<CL_DRIVER_VERSION>();
return out;
}
std::ostream& operator<<(std::ostream& out, Context const & context){
Platform p;
context.getInfo(CL_CONTEXT_PLATFORM, &p);
out << "Platform:\t" << p << '\n';
auto devices = context.getInfo<CL_CONTEXT_DEVICES>();
out << "Number of devices:\t" << devices.size() << '\n';
int i = 0;
for(auto&& d : devices){
out << ++i << "\t" << d << '\n';
}
return out;
}
std::ostream& operator<<(std::ostream& out, Program const & program){
return out << "Kernels in program:\t" << program.getInfo<CL_PROGRAM_KERNEL_NAMES>();
}
}
#endif

215
main.cpp

@ -0,0 +1,215 @@
//
// main.cpp
// XcodeOpenCL
//
// Created by Joshua Moerman on 28/03/14.
//
//
#include "cl2.hpp"
#include "NSWrapper.hpp"
#include "../../CodeTA/ImageStreams/include/png.hpp"
#include <OpenGL/gl3.h>
#include <vector>
#include <iostream>
#include <fstream>
#include <sstream>
#include <cassert>
#include <cmath>
using namespace std;
#define BUFFER_OFFSET(i) ((char *)NULL + (i))
const GLfloat quad[] = {
// x y z, r g b
1, -1, 0, 1, 0, 0,
-1, -1, 0, 0, 1, 0,
1, 1, 0, 0, 0, 1,
-1, 1, 0, 1, 1, 1
};
static void check_shader(GLuint s){
GLint logLength;
glGetShaderiv(s, GL_INFO_LOG_LENGTH, &logLength);
if (logLength > 0) {
GLchar *log = (GLchar *)malloc(logLength);
glGetShaderInfoLog(s, logLength, &logLength, log);
std::cout << "Shader compile log:\n%s" << log << std::endl;
free(log);
}
}
string slurp(string filename) {
ifstream in(filename);
stringstream sstr;
sstr << in.rdbuf();
return sstr.str();
}
struct App {
float time;
GLuint program;
GLint uniform;
GLuint vao;
GLuint posBufferName;
void initialize(){
time = 0;
program = glCreateProgram();
auto v_source_str = slurp("Fractal.vsh");
auto v_source = v_source_str.c_str();
GLuint v = glCreateShader(GL_VERTEX_SHADER);
glShaderSource(v, 1, &v_source, NULL);
glCompileShader(v);
check_shader(v);
auto f_source_str = slurp("Fractal.fsh");
auto f_source = f_source_str.c_str();
GLuint f = glCreateShader(GL_FRAGMENT_SHADER);
glShaderSource(f, 1, &f_source, NULL);
glCompileShader(f);
check_shader(f);
glAttachShader(program, v);
glAttachShader(program, f);
glBindAttribLocation(program, 0, "position");
glBindAttribLocation(program, 1, "color");
glLinkProgram(program);
uniform = glGetUniformLocation(program, "rotation");
glDetachShader(program, v);
glDeleteShader(v);
glDetachShader(program, f);
glDeleteShader(f);
glGenVertexArrays(1, &vao);
glBindVertexArray(vao);
glGenBuffers(1, &posBufferName);
glBindBuffer(GL_ARRAY_BUFFER, posBufferName);
glBufferData(GL_ARRAY_BUFFER, 4*6*4, quad, GL_STATIC_DRAW);
glEnableVertexAttribArray(0);
glEnableVertexAttribArray(1);
glVertexAttribPointer(0, 3, GL_FLOAT, GL_FALSE, 6*4, BUFFER_OFFSET(0));
glVertexAttribPointer(1, 3, GL_FLOAT, GL_FALSE, 6*4, BUFFER_OFFSET(3*4));
}
void draw(){
time += 1/60.0f;
glClearColor(0.65f, 0.65f, 0.65f, 1.0f);
glClear(GL_COLOR_BUFFER_BIT | GL_DEPTH_BUFFER_BIT);
glUseProgram(program);
glUniform1f(uniform, time);
glBindVertexArray(vao);
glDrawArrays(GL_TRIANGLE_STRIP, 0, 4);
}
};
int main(){
const bool mandelbrot = false;
if(mandelbrot) {
App a;
NSAppWrapper app;
app.create_window({
[&](ContextParameters){
a.initialize();
},
[&](ContextParameters){
a.draw();
},
nullptr
});
app.run();
return 0;
}
// else, this other weird thing
auto context = cl::Context::getDefault();
cout << context << endl;
cl_int err = 0;
// build the program
auto KernelSource = slurp("Kernel.cl");
cl::Program program(context, {KernelSource.c_str(), KernelSource.size()}, true, &err);
check(err);
cout << program << endl;
// grab the kernel
KernelOp kernel(program, "square", &err);
check(err);
// create a queue
cl::CommandQueue queue(context, context.getInfo<CL_CONTEXT_DEVICES>().front());
// make a lot of data
constexpr size_t W = 1280 * 4;
constexpr size_t H = 800 * 4;
std::vector<cl_float> input_vector(W*H);
for(int y = 0; y < H; ++y){
for(int x = 0; x < W; ++x){
input_vector[x + W*y] = 10 * ((x / double(W) - 0.5) + 1.3371337*(y / double(H) - 0.5));
}
}
// transfer data into buffers
cl::Buffer input(context, input_vector.begin(), input_vector.end(), false, true);
int r = 80, g = 40, b = 20;
// DO IT (in place)
for(int i = 0; i < r; ++i){
check(kernel(queue, W, H, input, W, input));
}
// read back
queue.finish();
auto red = input_vector;
// DO IT (in place)
for(int i = 0; i < g; ++i){
check(kernel(queue, W, H, input, W, input));
}
// read back
queue.finish();
auto green = input_vector;
// DO IT (in place)
for(int i = 0; i < b; ++i){
check(kernel(queue, W, H, input, W, input));
}
// read back
queue.finish();
auto& blue = input_vector;
// test
cout << "opencl is done" << endl;
png::ostream<> image(W, H, "test.png");
for(int i = 0; i < red.size(); ++i){
image << png::ostream<>::pixel(blue[i], red[i], green[i]);
}
cout << "png is saved" << endl;
}