Posts Tagged ‘xcode’

Friday, January 6th, 2017

Like clockwork, my matlab updates have gotten out of sync with Xcode updates. It seems like fixing this SDK error always requires a different hack each year. This year I got the error:

No supported compiler or SDK was found. For options, visit
http://www.mathworks.com/support/compilers/current_release/. 

To fix it, I replaced all occurrences of 10.9 with 10.11 in /Applications/MATLAB_R2017a.app/bin/maci64/mexopts/clang{++,}_maci64.xml

I’m still getting linker warnings:

ld: warning: object file was built for newer OSX version (10.11) than being linked (10.9)

For now, I’m assuming that I can ignore them. We’ll see how far that gets me.

Fast Automatic Skinning Transformations prototype program on github

Tuesday, February 17th, 2015

I’ve put up the code for my old prototype program corresponding to the paper “Fast Automatic Skinning Transformations” on github. Enjoy.

Linking against static library using Eigen produces many direct access ... to global weak symbol warnings

Friday, April 13th, 2012

Our group builds a static library of useful functions implemented using the Eigen matrix library in C++. We build the library with a standard Makefile and gcc. But when I try to link against this library in an Xcode project (which also used Eigen) I got many warnings from the linker of the form:


ld: warning: direct access in void Eigen::internal::computeProductBlockingSizes<float, float, 4>(long&, long&, long&)to global weak symbol Eigen::internal::manage_caching_sizes(Eigen::Action, long*, long*)::m_l2CacheSizemeans the weak symbol cannot be overridden at runtime. This was likely caused by different translation units being compiled with different visibility settings.

There seem to be no problems at runtime, but nonetheless it’s annoying to have so many warnings.

To make the warnings go away, I went to my project settings in Xcode and made sure that Symbols Hidden by Default was set to No.

Xcode symbols hidden by default

Basic OpenGL Cocoa App using C

Thursday, September 22nd, 2011

I’m rewriting this tutorial for Xcode 4 on Mac OS X 10.6. For whatever reason when I tried it on my computer it gave compilation errors. The only thing I change is the code inside glView.mm and glView.h

– create a new Cocoa Application
– ceate glView.h/glView.m
– copy and paste in this code and save files
– add frameworks: QuartzCore and OpenGL
– open up MainMenu.xib
– drag Custom View into main window and resize to fill window
– set the autosizing of view to have the two inner arrows on so it resizes automatically with the window
– set the class of the Custom View to be glView
– save nib
– go back to Xcode and run project

Should show a yellow window. Here’s the code:

glView.h

#import <Cocoa/Cocoa.h>

// for display link
#import <QuartzCore/QuartzCore.h>

@interface glView : NSOpenGLView
{
  CVDisplayLinkRef displayLink;
  
  double    deltaTime;
  double    outputTime;
  float    viewWidth;
  float    viewHeight;
}

@end

glView.mm


#import "glView.h"

@interface glView (InternalMethods)

- (CVReturn)getFrameForTime:(const CVTimeStamp *)outputTime;
- (void)drawFrame;

@end

@implementation glView

#pragma mark -
#pragma mark Display Link

static CVReturn MyDisplayLinkCallback(CVDisplayLinkRef displayLink, const CVTimeStamp *now,
                                      const CVTimeStamp *outputTime, CVOptionFlags flagsIn,
                                      CVOptionFlags *flagsOut, void *displayLinkContext)
{
  // go back to Obj-C for easy access to instance variables
  CVReturn result = [(glView *)displayLinkContext getFrameForTime:outputTime];
  return result;
}

- (CVReturn)getFrameForTime:(const CVTimeStamp *)outputTime
{
  // deltaTime is unused in this bare bones demo, but here's how to calculate it using display link info
  deltaTime = 1.0 / (outputTime->rateScalar * (double)outputTime->videoTimeScale / (double)outputTime->videoRefreshPeriod);
  
  [self drawFrame];
  
  return kCVReturnSuccess;
}

- (void)dealloc
{
  CVDisplayLinkRelease(displayLink);
  
  [super dealloc];
}

- (id)initWithFrame:(NSRect)frameRect
{
  // context setup
  NSOpenGLPixelFormat        *windowedPixelFormat;
  NSOpenGLPixelFormatAttribute    attribs[] = {
    NSOpenGLPFAWindow,
    NSOpenGLPFAColorSize, 32,
    NSOpenGLPFAAccelerated,
    NSOpenGLPFADoubleBuffer,
    NSOpenGLPFASingleRenderer,
    0 };
  
  windowedPixelFormat = [[NSOpenGLPixelFormat alloc] initWithAttributes:attribs];
  if (windowedPixelFormat == nil)
  {
    NSLog(@"Unable to create windowed pixel format.");
    exit(0);
  }
  self = [super initWithFrame:frameRect pixelFormat:windowedPixelFormat];
  if (self == nil)
  {
    NSLog(@"Unable to create a windowed OpenGL context.");
    exit(0);
  }
  [windowedPixelFormat release];
  
  // set synch to VBL to eliminate tearing
  GLint    vblSynch = 1;
  [[self openGLContext] setValues:&vblSynch forParameter:NSOpenGLCPSwapInterval];
  
  // set up the display link
  CVDisplayLinkCreateWithActiveCGDisplays(&displayLink);
  CVDisplayLinkSetOutputCallback(displayLink, MyDisplayLinkCallback, self);
  CGLContextObj cglContext = (CGLContextObj)[[self openGLContext] CGLContextObj];
  CGLPixelFormatObj cglPixelFormat = (CGLPixelFormatObj)[[self pixelFormat] CGLPixelFormatObj];
  CVDisplayLinkSetCurrentCGDisplayFromOpenGLContext(displayLink, cglContext, cglPixelFormat);
  
  return self;
}

- (void)awakeFromNib
{
  NSSize    viewBounds = [self bounds].size;
  viewWidth = viewBounds.width;
  viewHeight = viewBounds.height;
  
  // activate the display link
  CVDisplayLinkStart(displayLink);
}

- (void)reshape
{
  NSSize    viewBounds = [self bounds].size;
  viewWidth = viewBounds.width;
  viewHeight = viewBounds.height;
  
  NSOpenGLContext    *currentContext = [self openGLContext];
  [currentContext makeCurrentContext];
  
  // remember to lock the context before we touch it since display link is threaded
  CGLLockContext((CGLContextObj)[currentContext CGLContextObj]);
  
  // let the context know we've changed size
  [[self openGLContext] update];
  
  CGLUnlockContext((CGLContextObj)[currentContext CGLContextObj]);
}

- (void)drawRect:(NSRect)rect
{
  [self drawFrame];
}

- (void)drawFrame
{
  NSOpenGLContext    *currentContext = [self openGLContext];
  [currentContext makeCurrentContext];
  
  // must lock GL context because display link is threaded
  CGLLockContext((CGLContextObj)[currentContext CGLContextObj]);
  
  glViewport(0, 0, viewWidth, viewHeight);
  
  // Draw something that changes over time to prove to yourself that it's really updating in a tight loop
  glClearColor(
    sin(CFAbsoluteTimeGetCurrent()),
    sin(7.0*CFAbsoluteTimeGetCurrent()),
    sin(CFAbsoluteTimeGetCurrent()/3.0),0);
  glClear(GL_COLOR_BUFFER_BIT);
  
  // draw here
  
  [currentContext flushBuffer];
  
  CGLUnlockContext((CGLContextObj)[currentContext CGLContextObj]);
}

@end

Using Cocoa and Xcode 4 to set up a simple OpenGL-based app

Tuesday, September 13th, 2011

Recently I wanted to start a new OpenGL prototyping application. I imagined that it might grow into something bigger so instead of just using GLUT I wanted to start off using a full-on Cocoa app. This will be handy in the end, because developing with Xcode is tailored for Cocoa apps and since I’m usually the only one running my code having native hooks like the correct keyboard modifiers and file dialogs will be convenient. I’m new to Xcode 4, however, and found it extremely frustrating to get a simple Cocoa app with an OpenGL context up and running. In the end here’s what I’ve done to make a basic OpenGL app in Cocoa with TODO: comments left in all the places where one might want to have application specific code.

Make a new Xcode project.

Choose File > New > New Project …
Then choose Mac OS X > Application > Cocoa Application …
Give your “Product” a name, from here on out I will use “BasicOpenGL” as my product/project name.
Leave “Create Document Based Application” unchecked
Uncheck “Create a git repository” (optional, of course)
At this point you should be able to build and run and see a boring blank window pop up.

Create the BasicOpenGLView class

The layout of xcode 4 is rather annoying. Depending on what “Navigator” your looking at you see different GUI elements and have different access to different things. Maybe this is handy once you’re used to it. I am not used to it and it is not handy.
Set up the Navigator to view the project file hierarchy: View > Navigators > Project
Create a file called BasicOpenGLView.h with the following contents:


#import <OpenGL/gl.h>
#import <Cocoa/Cocoa.h>


///////////////////////////////////////////////////////////////////////////////
// OpenGL error handling
///////////////////////////////////////////////////////////////////////////////
// error reporting as both window message and debugger string
void reportError (char * strError);
// if error dump gl errors to debugger string, return error
GLenum glReportError (void);

@interface BasicOpenGLView : NSOpenGLView
{
  // TODO: add object (pointer) for your OpenGL-based app
  
  // Timer object that can be attached to animationTimer method for time-based
  // animations
  NSTimer* timer;
  // Current time at construction
  CFAbsoluteTime start_time;
  // True only if a redraw is necessary, i.e. scene has changed 
  bool damage;
  // This is a cheesy global variable that allows me to to call certain code
  // only once per OpenGL session, I'm not sure how necessary this is
  bool openGL_initialized;
}

///////////////////////////////////////////////////////////////////////////////
// File IO
///////////////////////////////////////////////////////////////////////////////
// Default open function, called when using File > Open dialog, but also when
// files (of correct type) are dragged to the dock icon 
- (IBAction) openDocument: (id) sender;
// Helper method for openDocument that takes the file name of the given file
// directly so that it can be hooked up to the event of double clicking a file
// in Finder
// Inputs:
//   file_name  string containing path to file to be opened
// Returns:
//   true only on success
- (BOOL) openDocumentFromFileName: (NSString *)file_name;
// Default save function, called when using File > Save as ... dialog
- (IBAction) saveDocumentAs: (id) sender;


///////////////////////////////////////////////////////////////////////////////
// Mouse and Keyboard Input
///////////////////////////////////////////////////////////////////////////////
- (void) keyDown:(NSEvent *)theEvent;
- (void) keyUp:(NSEvent *)theEvent;
- (void) mouseDown:(NSEvent *)theEvent;
- (void) rightMouseDown:(NSEvent *)theEvent;
- (void) otherMouseDown:(NSEvent *)theEvent;
- (void) mouseUp:(NSEvent *)theEvent;
- (void) rightMouseUp:(NSEvent *)theEvent;
- (void) otherMouseUp:(NSEvent *)theEvent;
- (void) mouseMoved:(NSEvent *)theEvent;
- (void) mouseDragged:(NSEvent *)theEvent;
- (void) rightMouseDragged:(NSEvent *)theEvent;
- (void) otherMouseDragged:(NSEvent *)theEvent;
- (void) scrollWheel:(NSEvent *)theEvent;
- (void) viewDidMoveToWindow;
// OpenGL apps like to think of (0,0) being the top left corner, cocoa apps
// think of (0,0) as the bottom left corner. This simple flips the y coordinate
// according to the current height
// Inputs:
//   location  point of click according to Cocoa
// Returns
//   point of click according to with y coordinate flipped
- (NSPoint) flip_y:(NSPoint)location;


///////////////////////////////////////////////////////////////////////////////
// OpenGL
///////////////////////////////////////////////////////////////////////////////
// Called whenever openGL context changes size
- (void) reshape;
// Main display or draw function, called when redrawn
- (void) drawRect:(NSRect)rect;
// Set initial OpenGL state (current context is set)
// called after context is created
- (void) prepareOpenGL;
// this can be a troublesome call to do anything heavyweight, as it is called
// on window moves, resizes, and display config changes.  So be careful of
// doing too much here.  window resizes, moves and display changes (resize,
// depth and display config change)
- (void) update;


///////////////////////////////////////////////////////////////////////////////
// Animation
///////////////////////////////////////////////////////////////////////////////

// per-window timer function, basic time based animation preformed here
- (void) animationTimer:(NSTimer *)timer;
// Set an instance variable to the current time, so as to keep track of the time
// at which the app was started
- (void) setStartTime;
// Get the time since the start time: elapsed time since app was started
- (CFAbsoluteTime) getElapsedTime;


///////////////////////////////////////////////////////////////////////////////
// Cocoa
///////////////////////////////////////////////////////////////////////////////
- (BOOL) acceptsFirstResponder;
- (BOOL) becomeFirstResponder;
- (BOOL) resignFirstResponder;
- (void) awakeFromNib;
- (void) terminate:(NSNotification *)aNotification;

@end

And also a file called BasicOpenGLView.h with the following contents:


#import "BasicOpenGLView.h"
// For functions like gluErrorString()
#import <OpenGL/glu.h>
#ifdef __APPLE__
#define _MACOSX
#endif


void reportError (char * strError)
{
  // Set up a fancy font/display for error messages
  NSMutableDictionary *attribs = [NSMutableDictionary dictionary];
  [attribs setObject: [NSFont fontWithName: @"Monaco" size: 9.0f] 
    forKey: NSFontAttributeName];
  [attribs setObject: [NSColor whiteColor] 
    forKey: NSForegroundColorAttributeName];
  // Build the error message string
  NSString * errString = [NSString stringWithFormat:@"Error: %s.", strError];
  // Display to log
  NSLog (@"%@\n", errString);
}

GLenum glReportError (void)
{
  // Get current OpenGL error flag
  GLenum err = glGetError();
  // If there's an error report it
  if (GL_NO_ERROR != err)
  {
    reportError ((char *) gluErrorString (err));
  }
  return err;
}

@implementation BasicOpenGLView

  -(IBAction) openDocument: (id) sender
  {
    NSOpenPanel *tvarNSOpenPanelObj  = [NSOpenPanel openPanel];
    // TODO: Add a item to this list corresponding to each file type extension
    // this app supports opening
    // Create an array of strings specifying valid extensions and HFS file types.
    NSArray *fileTypes = [NSArray arrayWithObjects:
      @"obj",
      @"OBJ",
      NSFileTypeForHFSTypeCode('TEXT'),
      nil];
    // Create an Open file... dialog
    NSInteger tvarNSInteger = [tvarNSOpenPanelObj runModalForTypes:fileTypes];
    // If the user selected OK then load the file
    if(tvarNSInteger == NSOKButton)
    {
      // Pass on file name to opener helper
      [self openDocumentFromFileName:[tvarNSOpenPanelObj filename]];
    }
  }

  - (BOOL)openDocumentFromFileName:(NSString *) file_name
  {
    // convert cocoa string to c string
    const char * c_file_name = [file_name UTF8String];
    // TODO: handle loading a file from filename
    NSLog(@"Opening file: %s", c_file_name);
    damage = true;
    return false;
  }

  -(IBAction) saveDocumentAs: (id) sender
  {
    NSSavePanel *savePanel = [NSSavePanel savePanel]; 
    [savePanel setTitle:@"Save as (.obj by default)"];
    // TODO: Add a item to this list corresponding to each file type extension
    // this app supports opening
    // Create an array of strings specifying valid extensions and HFS file types.
    NSArray *fileTypes = [NSArray arrayWithObjects:
      @"obj",
      @"OBJ",
      NSFileTypeForHFSTypeCode('TEXT'),
      nil];
    // Only allow these file types
    [savePanel setAllowedFileTypes:fileTypes]; 
    [savePanel setTreatsFilePackagesAsDirectories:NO]; 
    // Allow user to save file as he likes
    [savePanel setAllowsOtherFileTypes:YES];
    // Create save as... dialog
    NSInteger user_choice =  
      [savePanel runModalForDirectory:NSHomeDirectory() file:@""];
    // If user selected OK then save the file
    if(NSOKButton == user_choice)
    {
      // convert cocoa string to c string
      const char * file_name = [[savePanel filename] UTF8String];
      // TODO: handle saving default file
      NSLog(@"Saving file to %s", file_name);
    } 
  }

  -(void)keyDown:(NSEvent *)theEvent
  {
    // NOTE: holding a key on the keyboard starts to signal multiple down
    // events (the only one final up event)
    NSString *characters = [theEvent characters];
    if ([characters length])
    {
      // convert characters to single char
      char character = [characters characterAtIndex:0];
      // TODO: Handle key down event
      NSLog(@"Keyboard down: %c\n",character);
    }
    damage = true;
  }

  -(void)keyUp:(NSEvent *)theEvent
  {
    NSString *characters = [theEvent characters];
    if ([characters length])
    {
      // convert characters to single char
      char character = [characters characterAtIndex:0];
      // TODO: Handle key up event
      NSLog(@"Keyboard up: %c\n",character);
    }
    damage = true;
  }

  - (void)mouseDown:(NSEvent *)theEvent
  {
    // Get location of the click
    NSPoint location = 
      [self flip_y:
        [self convertPoint:[theEvent locationInWindow] fromView:nil]];
    // TODO: Handle mouse up event
    NSLog(@"Mouse down at (%g,%g)\n",location.x,location.y);
    damage = true;
  }

  - (void)rightMouseDown:(NSEvent *)theEvent
  {
    // TODO: Handle right mouse button down event
    // For now just treat as left mouse button down event
    [self mouseDown: theEvent];
  }

  - (void)otherMouseDown:(NSEvent *)theEvent
  {
    // TODO: Handle other strange mouse button bown events
    // For now just treat as left mouse button down event
    [self mouseDown: theEvent];
  }

  - (void)mouseUp:(NSEvent *)theEvent
  {
    // Get location of the click
    NSPoint location = 
      [self flip_y:
        [self convertPoint:[theEvent locationInWindow] fromView:nil]];
    // TODO: Handle mouse up event
    NSLog(@"Mouse up at (%g,%g)\n",location.x,location.y);
    damage = true;
  }

  - (void)rightMouseUp:(NSEvent *)theEvent
  {  
    // TODO: Handle right mouse button up event
    // For now just treat as left mouse button up event
    [self mouseUp: theEvent];
  }

  - (void)otherMouseUp:(NSEvent *)theEvent
  {
    // TODO: Handle other strange mouse button up events  
    // For now just treat as left mouse button up event
    [self mouseUp: theEvent];
  }

  - (void)mouseMoved:(NSEvent *)theEvent
  {
    NSPoint location = 
      [self flip_y:
        [self convertPoint:[theEvent locationInWindow] fromView:nil]];
    // TODO: Handle mouse move event
    NSLog(@"Mouse moved to (%g,%g)\n",location.x,location.y);
    damage = true;
  }

  - (void)mouseDragged:(NSEvent *)theEvent
  {
    
    NSPoint location = 
      [self flip_y:
        [self convertPoint:[theEvent locationInWindow] fromView:nil]];
    // TODO: Handle mouse drag event
    NSLog(@"Mouse dragged to (%g,%g)\n",location.x,location.y);
    damage = true;
  }

  - (void)rightMouseDragged:(NSEvent *)theEvent
  { 
    // TODO: Handle right mouse button drag event
    // For now just treat as left mouse button drag event
    [self mouseDragged: theEvent];
  }

  - (void)otherMouseDragged:(NSEvent *)theEvent
  {
    // TODO: Handle other strange mouse button drag event
    // For now just treat as left mouse button drag event
    [self mouseDragged: theEvent];
  }


  - (void)scrollWheel:(NSEvent *)theEvent
  {

    NSPoint location = 
      [self flip_y:
        [self convertPoint:[theEvent locationInWindow] fromView:nil]];
    // TODO: Handle mouse scroll event
    NSLog(@"Mouse scroll wheel at (%g,%g) by (%g,%g)\n",
      location.x,location.y,[theEvent deltaX],[theEvent deltaY]);
    damage = true;
  }

  - (void) viewDidMoveToWindow
  {
    // Listen to all mouse move events (not just dragging)
    [[self window] setAcceptsMouseMovedEvents:YES];
    // When view changes to this window then be sure that we start responding
    // to mouse events
    [[self window] makeFirstResponder:self];
  }

  - (NSPoint) flip_y:(NSPoint) location
  {
    // Get openGL context size
    NSRect rectView = [self bounds];
    // Cocoa gives opposite of OpenGL y direction, flip y direction
    location.y = rectView.size.height - location.y;
    return location;
  }

  - (void) reshape
  {
    NSRect rectView = [self bounds];
    // TODO: Handle resize window using the following
    NSLog(@"New context size: %g %g\n",
      rectView.size.width,rectView.size.height);
  }

  - (void) drawRect:(NSRect)rect
  {
    // TODO: handle draw event
    // For now just clear the screen with a time dependent color
    glClearColor(
      fabs(sin([self getElapsedTime])),
      fabs(sin([self getElapsedTime]/3)),
      fabs(sin([self getElapsedTime]/7)),
      0);
    glClear(GL_COLOR_BUFFER_BIT | GL_DEPTH_BUFFER_BIT);
    // Elapsed time in seconds: getElapsedTime()
    // Report any OpenGL errors
    glReportError ();
    // Flush all OpenGL calls
    glFlush();
    // Flush OpenGL context
    [[self openGLContext] flushBuffer];
  }

  - (void) prepareOpenGL
  {
    const GLint swapInt = 1;
    // set to vbl sync
    [[self openGLContext] setValues:&swapInt 
      forParameter:NSOpenGLCPSwapInterval];
    if(!openGL_initialized)
    {
      // Get command line arguments and find whether stealFocus is set to YES
      NSUserDefaults *args = [NSUserDefaults standardUserDefaults];
      // also find out if app should steal focus
      bool stealFocus = [args boolForKey:@"stealFocus"];
      if(stealFocus)
      {
        // Steal focus means that the apps window will appear in front of all
        // other programs when it launches even in front of the calling
        // application (e.g. a terminal)
        [[NSApplication sharedApplication] activateIgnoringOtherApps:YES];
      }
      // TODO: Initialize OpenGL app, do anything here that you need to *after*
      // the OpenGL context is initialized (load textures, shaders, etc.)
      openGL_initialized = true;
    }
    NSLog(@"prepareOpenGL\n");
  }

  - (void) update 
  {
    [super update];  
  }

  - (void)animationTimer:(NSTimer *)timer
  { 
    // TODO: handle timer based redraw (animation) here
    bool your_app_says_to_redraw = true;
    if(your_app_says_to_redraw || damage)
    {
      damage = false;
      [self drawRect:[self bounds]];
    }
  }

  - (void) setStartTime
  {   
    start_time = CFAbsoluteTimeGetCurrent ();
  }

  - (CFAbsoluteTime) getElapsedTime
  {   
    return CFAbsoluteTimeGetCurrent () - start_time;
  }

  - (BOOL)acceptsFirstResponder
  {
    return YES;
  }

  - (BOOL)becomeFirstResponder
  {
    return  YES;
  }

  - (BOOL)resignFirstResponder
  {
    return YES;
  }

  - (void) awakeFromNib
  {
    openGL_initialized = false;
    // keep track of start/launch time
    [self setStartTime];
    // start animation timer
    timer = [NSTimer timerWithTimeInterval:(1.0f/60.0f) target:self 
      selector:@selector(animationTimer:) userInfo:nil repeats:YES];
    [[NSRunLoop currentRunLoop] addTimer:timer forMode:NSDefaultRunLoopMode];
    // ensure timer fires during resize
    [[NSRunLoop currentRunLoop] addTimer:timer 
      forMode:NSEventTrackingRunLoopMode]; 
  }

  - (void) terminate:(NSNotification *)aNotification
  {
    // TODO: delete your app's object
    NSLog(@"Terminating");
  }

@end

This class, BasicOpenGLView, will contain the OpenGL context and hooks (marked with TODO:s and temp code for application specific function calls.

For the code to build again we will need to link to OpenGL

Link to OpenGL framework

Set up the Navigator to view the project file hierarchy: View > Navigators > Project
Select the Project, and make sure the BasicOpenGL Target is selected.
Under Linked Frameworks and Libraries, click the little plus sign and add OpenGL.framework

Now the project should build and run again (although it’s not actually using our new class)

Hook up BasicOpenGLView class to the application

Change the contents of BasicOpenGLAppDelegate.h to the following:


//
//  BasicOpenGLDelegate.h
//  BasicOpenGL
//
//  Created by Alec Jacobson on 9/13/11.
//  Copyright 2011 New York University. All rights reserved.
//

#import <Cocoa/Cocoa.h>
#import "BasicOpenGLView.h"

@interface BasicOpenGLAppDelegate : NSObject
{
    IBOutlet BasicOpenGLView * basic_opengl_view;
}
// Set the application to terminate when all windows (there is only one) are 
// closed
- (BOOL)applicationShouldTerminateAfterLastWindowClosed:
  (NSApplication *)theApplication;
// Allows the application to open a file on launch and when files are dragged
// to the dock icon of the app
- (BOOL)application:(NSApplication *)theApplication openFile:
  (NSString *)filename;
// Triggered when application is about to terminate
- (void)applicationWillTerminate:(NSNotification *)aNotification;
@end

And the contents of BasicOpenGLAppDelegate.m to:


//
//  BasicOpenGLAppDelegate.m
//  BasicOpenGL
//
//  Created by Alec Jacobson on 9/13/11.
//  Copyright 2011 New York University. All rights reserved.
//

#import "BasicOpenGLAppDelegate.h"

@implementation BasicOpenGLAppDelegate

- (BOOL)applicationShouldTerminateAfterLastWindowClosed:(NSApplication *)theApplication
{
	return YES;
}
- (BOOL)application:(NSApplication *)theApplication openFile:(NSString *)filename
{ 
  [basic_opengl_view openDocumentFromFileName:filename];
  return YES;
}  
- (void)applicationWillTerminate:(NSNotification *)aNotification
{
  [basic_opengl_view terminate:aNotification];
}

@end

This allows your app to hook up the our class via Cocoa outlets and sets up launching and exiting the program.

The code again should build and run, but we still don’t see any effect because the class hasn’t been added to the GUI window.

Hook up BasicOpenGLView class to GUI

Open up MainMenu.xib

  1. Open View > Utilities > Identity Inspector
  2. Click Window on the right column
  3. Click little cube in bottom right (Object Library)
  4. Drag OpenGL View onto your window
  5. In the View > Utilities > Identity Inspector change NsOpenGLView to the class name to BasicOpenGLView

In the View > Utilities > Connections Inspector, drag a new referencing Outlet to the little blue square in the left column that says “Basic OpenGL App Delegate”

At this point you should be able to build and run and see now our window has a rectangle in it changing color over time (this is a result of the default code in BasicOpenGLView.mm actually being executed now). You’ll also see that logs are printed to the output concerning keyboard and mouse interactions. However, this colored rectangle, our OpenGL context, is not well placed in the window and does not resize automatically.

Set OpenGL display options and auto-resizing

Still in MainMenu.xib with the openGL context selected.
In the View > Utilities > Attributes Inspector change the following:
Color: Default
Depth: 32 bit
Sampling: 16 Samples
Anti-aliasing: Default anti-aliasing
(Anti aliasing will not show up in the preview unless you reopen the x-code project)

In the View > Utilities > Size Inspector you may want to set up default sizes for your window, view and openglview objects, also automatic resizing

Select the containing window and in View > Utilities > Attributes Inspector uncheck Memory > “One Shot” (otherwise you’ll get a warning at compile time)

Now if you build and run you’ll see a full-window, opengl display of a changing color.

Recognize certain file types

Currently with the files set up as above, you can “open” and “save” .obj files. Opening and saving does nothing more than print Log statements, but this gives you the idea of what part of the code to edit to implement something interesting. The current application can save through the File &gt: Save as … menu item. It can open files through File > Open. But we can also set the app up to open files by dragging .obj files to the application’s dock icon or right clicking on a .obj file and selecting BasicOpenGL.app. Actually if you don’t have any Mac applications already claiming .obj files you’ll see that by just building and running your app all .obj files in finder have been associated with BasicOpenGL.app

In the file BasicOpenGLView, replace the default file extension I’ve used, .obj, in the relevant places with your application specific file type.

To add a certain file type to be recognized by your app (so you can drag the file to the dock icon or “Open with …”) then View > Navigators > Project. Select your project, then your target then Info > Document types. In the bottom right choose Add > Add document type. Then give your type a name (3D Object File), an extension (obj) and perhaps an icon (.icns file).

Download a working copy of all this as a zip

Yet another macports/qt/mac nightmare…

Friday, May 6th, 2011

Today I wasted countless hours recompiling all sorts of junk on my computer. The culprits once again are macports, Qt and Apple’s 32-bit/64-bit androgyny.

I recently upgraded from mac os x 10.5 (where 32-bits is default) to mac os x 10.6 (where 64-bits is default). Macports naturally pretends to work just fine after I make the switch. Low and behold it’s little house on the sand is about to wash away. Today I wanted to upgrade to the latest version of pdflatex. I tried to do so using macports. No go. Eventually I found the “migration” guide for “migrating” you macports after you upgrade from mac os x 10.5 to mac os x 10.6. Their idea of “migration” is uninstalling and reinstalling everything! All of my 300 some ports needed to be recompiled. Their instructions by the way did not work as it didn’t account for external dependencies still hanging around in i386, 32-bit mode.

I’ve finally reinstalled pdflatex. But now my Qt installation doesn’t work.

I tried:


sudo port install qt4-mac

and about a year later everything seemed to have worked fine. Kudos to macports for having the bright idea to restrain from distributing precompiled binaries. I’m learning patience!

But everything did not work fine. I re-qmake-ed my current project into an xcode project which went fine, but when I tried to build my project in xcode I got the following mysterious error:


pbxcp: warning: couldn't strip: /absolute/path/to/my/app: No such file or directory

Searching around the only ideas I found were to delete the build directory, clean and rebuild which did not work.

Solution:


sudo port deactivate qt4-mac

Install Qt SDK directly from Qt site (hours faster than macports by the way).

Anti-aliasing in Cocoa OpenGL app using Xcode

Tuesday, March 29th, 2011

It was incredibly difficult to figure this out. I finally found a solution to allowing my OpenGL app to be anti-aliased in an OpenGL context within a Cocoa App’s window using Xcode.

In the end the steps are quite simple. Open you .xib that contains the OpenGL context. Select that container and open View > Utilities > Attributes Inspector

xcode utilities attributes inspector

Then find the Sampling and Anti-Aliasing parameters and make sure they’re both set to something reasonable.

xcode cocoa anti-aliasing

And here’s a screen capture comparing without anti-aliasing (left) and with anti-aliasing (right).
with and without anti-aliasing

Combining CUDA, Qt, and Xcode

Monday, November 29th, 2010

As a proof of concept and a skeleton for some more intense code, I wanted to be sure that I could get a simple example program that used CUDA, Qt and Xcode. A priori there is no reason that this shouldn’t work.

The simplest way to do this is to separate CUDA from the main program entirely. I found a discussion of how to do this on the nvidia site. I will base my example on the final snipets on that thread.

Building a CUDA library

The first step will be to bake my gpu code into a static library. The idea is then to call that library from my Qt main program.

This involved five files.

HelloWorld.cu:


#include "HelloWorld.cuh"
#include <stdio.h>

// Kernel functions must be inlined (?)
#include "cuPrintf.cu"
__global__ void HelloFromDevice(void)
{
  cuPrintf("Hello, world from block %d thread %d!\n",blockIdx.x,threadIdx.x);
}

int HelloWorld()
{
  // greet from the host
  printf("Hello, world from the host!\n");

  // initialize cuPrintf
  cudaPrintfInit();

  // launch a kernel with a single thread to greet from the device
  HelloFromDevice<<<10,64>>>();

  // display the device's greeting
  cudaPrintfDisplay();
  
  // clean up after cuPrintf
  cudaPrintfEnd();

  return 0;
}

HelloWorld.cuh


int HelloWorld();

cuPrintf.cu (by Nvidia)


/*
	Copyright 2009 NVIDIA Corporation.  All rights reserved.

	NOTICE TO LICENSEE:   

	This source code and/or documentation ("Licensed Deliverables") are subject 
	to NVIDIA intellectual property rights under U.S. and international Copyright 
	laws.  

	These Licensed Deliverables contained herein is PROPRIETARY and CONFIDENTIAL 
	to NVIDIA and is being provided under the terms and conditions of a form of 
	NVIDIA software license agreement by and between NVIDIA and Licensee ("License 
	Agreement") or electronically accepted by Licensee.  Notwithstanding any terms 
	or conditions to the contrary in the License Agreement, reproduction or 
	disclosure of the Licensed Deliverables to any third party without the express 
	written consent of NVIDIA is prohibited.     

	NOTWITHSTANDING ANY TERMS OR CONDITIONS TO THE CONTRARY IN THE LICENSE AGREEMENT, 
	NVIDIA MAKES NO REPRESENTATION ABOUT THE SUITABILITY OF THESE LICENSED 
	DELIVERABLES FOR ANY PURPOSE.  IT IS PROVIDED "AS IS" WITHOUT EXPRESS OR IMPLIED 
	WARRANTY OF ANY KIND. NVIDIA DISCLAIMS ALL WARRANTIES WITH REGARD TO THESE 
	LICENSED DELIVERABLES, INCLUDING ALL IMPLIED WARRANTIES OF MERCHANTABILITY, 
	NONINFRINGEMENT, AND FITNESS FOR A PARTICULAR PURPOSE.   NOTWITHSTANDING ANY 
	TERMS OR CONDITIONS TO THE CONTRARY IN THE LICENSE AGREEMENT, IN NO EVENT SHALL 
	NVIDIA BE LIABLE FOR ANY SPECIAL, INDIRECT, INCIDENTAL, OR CONSEQUENTIAL DAMAGES, 
	OR ANY DAMAGES WHATSOEVER RESULTING FROM LOSS OF USE, DATA OR PROFITS,	WHETHER 
	IN AN ACTION OF CONTRACT, NEGLIGENCE OR OTHER TORTIOUS ACTION,  ARISING OUT OF 
	OR IN CONNECTION WITH THE USE OR PERFORMANCE OF THESE LICENSED DELIVERABLES.  

	U.S. Government End Users. These Licensed Deliverables are a "commercial item" 
	as that term is defined at  48 C.F.R. 2.101 (OCT 1995), consisting  of 
	"commercial computer  software"  and "commercial computer software documentation" 
	as such terms are  used in 48 C.F.R. 12.212 (SEPT 1995) and is provided to the 
	U.S. Government only as a commercial end item.  Consistent with 48 C.F.R.12.212 
	and 48 C.F.R. 227.7202-1 through 227.7202-4 (JUNE 1995), all U.S. Government 
	End Users acquire the Licensed Deliverables with only those rights set forth 
	herein. 

	Any use of the Licensed Deliverables in individual and commercial software must 
	include, in the user documentation and internal comments to the code, the above 
	Disclaimer and U.S. Government End Users Notice.
 */

/*
 *	cuPrintf.cu
 *
 *	This is a printf command callable from within a kernel. It is set
 *	up so that output is sent to a memory buffer, which is emptied from
 *	the host side - but only after a cudaThreadSynchronize() on the host.
 *
 *	Currently, there is a limitation of around 200 characters of output
 *	and no more than 10 arguments to a single cuPrintf() call. Issue
 *	multiple calls if longer format strings are required.
 *
 *	It requires minimal setup, and is *NOT* optimised for performance.
 *	For example, writes are not coalesced - this is because there is an
 *	assumption that people will not want to printf from every single one
 *	of thousands of threads, but only from individual threads at a time.
 *
 *	Using this is simple - it requires one host-side call to initialise
 *	everything, and then kernels can call cuPrintf at will. Sample code
 *	is the easiest way to demonstrate:
 *
	#include "cuPrintf.cu"
 	
	__global__ void testKernel(int val)
	{
		cuPrintf("Value is: %d\n", val);
	}

	int main()
	{
		cudaPrintfInit();
		testKernel<<< 2, 3 >>>(10);
		cudaPrintfDisplay(stdout, true);
		cudaPrintfEnd();
        return 0;
	}
 *
 *	See the header file, "cuPrintf.cuh" for more info, especially
 *	arguments to cudaPrintfInit() and cudaPrintfDisplay();
 */

#ifndef CUPRINTF_CU
#define CUPRINTF_CU

#include "cuPrintf.cuh"
#if __CUDA_ARCH__ > 100      // Atomics only used with > sm_10 architecture
#include <sm_11_atomic_functions.h>
#endif

// This is the smallest amount of memory, per-thread, which is allowed.
// It is also the largest amount of space a single printf() can take up
const static int CUPRINTF_MAX_LEN = 256;

// This structure is used internally to track block/thread output restrictions.
typedef struct __align__(8) {
	int threadid;				// CUPRINTF_UNRESTRICTED for unrestricted
	int blockid;				// CUPRINTF_UNRESTRICTED for unrestricted
} cuPrintfRestriction;

// The main storage is in a global print buffer, which has a known
// start/end/length. These are atomically updated so it works as a
// circular buffer.
// Since the only control primitive that can be used is atomicAdd(),
// we cannot wrap the pointer as such. The actual address must be
// calculated from printfBufferPtr by mod-ing with printfBufferLength.
// For sm_10 architecture, we must subdivide the buffer per-thread
// since we do not even have an atomic primitive.
__constant__ static char *globalPrintfBuffer = NULL;         // Start of circular buffer (set up by host)
__constant__ static int printfBufferLength = 0;              // Size of circular buffer (set up by host)
__device__ static cuPrintfRestriction restrictRules;         // Output restrictions
__device__ volatile static char *printfBufferPtr = NULL;     // Current atomically-incremented non-wrapped offset

// This is the header preceeding all printf entries.
// NOTE: It *must* be size-aligned to the maximum entity size (size_t)
typedef struct __align__(8) {
    unsigned short magic;                   // Magic number says we're valid
    unsigned short fmtoffset;               // Offset of fmt string into buffer
    unsigned short blockid;                 // Block ID of author
    unsigned short threadid;                // Thread ID of author
} cuPrintfHeader;

// Special header for sm_10 architecture
#define CUPRINTF_SM10_MAGIC   0xC810        // Not a valid ascii character
typedef struct __align__(16) {
    unsigned short magic;                   // sm_10 specific magic number
    unsigned short unused;
    unsigned int thread_index;              // thread ID for this buffer
    unsigned int thread_buf_len;            // per-thread buffer length
    unsigned int offset;                    // most recent printf's offset
} cuPrintfHeaderSM10;


// Because we can't write an element which is not aligned to its bit-size,
// we have to align all sizes and variables on maximum-size boundaries.
// That means sizeof(double) in this case, but we'll use (long long) for
// better arch<1.3 support
#define CUPRINTF_ALIGN_SIZE      sizeof(long long)

// All our headers are prefixed with a magic number so we know they're ready
#define CUPRINTF_SM11_MAGIC  (unsigned short)0xC811        // Not a valid ascii character


//
//  getNextPrintfBufPtr
//
//  Grabs a block of space in the general circular buffer, using an
//  atomic function to ensure that it's ours. We handle wrapping
//  around the circular buffer and return a pointer to a place which
//  can be written to.
//
//  Important notes:
//      1. We always grab CUPRINTF_MAX_LEN bytes
//      2. Because of 1, we never worry about wrapping around the end
//      3. Because of 1, printfBufferLength *must* be a factor of CUPRINTF_MAX_LEN
//
//  This returns a pointer to the place where we own.
//
__device__ static char *getNextPrintfBufPtr()
{
    // Initialisation check
    if(!printfBufferPtr)
        return NULL;

	// Thread/block restriction check
	if((restrictRules.blockid != CUPRINTF_UNRESTRICTED) && (restrictRules.blockid != (blockIdx.x + gridDim.x*blockIdx.y)))
		return NULL;
	if((restrictRules.threadid != CUPRINTF_UNRESTRICTED) && (restrictRules.threadid != (threadIdx.x + blockDim.x*threadIdx.y + blockDim.x*blockDim.y*threadIdx.z)))
		return NULL;

	// Conditional section, dependent on architecture
#if __CUDA_ARCH__ == 100
    // For sm_10 architectures, we have no atomic add - this means we must split the
    // entire available buffer into per-thread blocks. Inefficient, but what can you do.
    int thread_count = (gridDim.x * gridDim.y) * (blockDim.x * blockDim.y * blockDim.z);
    int thread_index = threadIdx.x + blockDim.x*threadIdx.y + blockDim.x*blockDim.y*threadIdx.z +
                       (blockIdx.x + gridDim.x*blockIdx.y) * (blockDim.x * blockDim.y * blockDim.z);
    
    // Find our own block of data and go to it. Make sure the per-thread length
	// is a precise multiple of CUPRINTF_MAX_LEN, otherwise we risk size and
	// alignment issues! We must round down, of course.
    unsigned int thread_buf_len = printfBufferLength / thread_count;
	thread_buf_len &= ~(CUPRINTF_MAX_LEN-1);

	// We *must* have a thread buffer length able to fit at least two printfs (one header, one real)
	if(thread_buf_len < (CUPRINTF_MAX_LEN * 2))
		return NULL;

	// Now address our section of the buffer. The first item is a header.
    char *myPrintfBuffer = globalPrintfBuffer + (thread_buf_len * thread_index);
    cuPrintfHeaderSM10 hdr = *(cuPrintfHeaderSM10 *)(void *)myPrintfBuffer;
    if(hdr.magic != CUPRINTF_SM10_MAGIC)
    {
        // If our header is not set up, initialise it
        hdr.magic = CUPRINTF_SM10_MAGIC;
        hdr.thread_index = thread_index;
        hdr.thread_buf_len = thread_buf_len;
        hdr.offset = 0;         // Note we start at 0! We pre-increment below.
        *(cuPrintfHeaderSM10 *)(void *)myPrintfBuffer = hdr;       // Write back the header

        // For initial setup purposes, we might need to init thread0's header too
        // (so that cudaPrintfDisplay() below will work). This is only run once.
        cuPrintfHeaderSM10 *tophdr = (cuPrintfHeaderSM10 *)(void *)globalPrintfBuffer;
        tophdr->thread_buf_len = thread_buf_len;
    }

    // Adjust the offset by the right amount, and wrap it if need be
    unsigned int offset = hdr.offset + CUPRINTF_MAX_LEN;
    if(offset >= hdr.thread_buf_len)
        offset = CUPRINTF_MAX_LEN;

    // Write back the new offset for next time and return a pointer to it
    ((cuPrintfHeaderSM10 *)(void *)myPrintfBuffer)->offset = offset;
    return myPrintfBuffer + offset;
#else
    // Much easier with an atomic operation!
    size_t offset = atomicAdd((unsigned int *)&printfBufferPtr, CUPRINTF_MAX_LEN) - (size_t)globalPrintfBuffer;
    offset %= printfBufferLength;
    return globalPrintfBuffer + offset;
#endif
}


//
//  writePrintfHeader
//
//  Inserts the header for containing our UID, fmt position and
//  block/thread number. We generate it dynamically to avoid
//	issues arising from requiring pre-initialisation.
//
__device__ static void writePrintfHeader(char *ptr, char *fmtptr)
{
    if(ptr)
    {
        cuPrintfHeader header;
        header.magic = CUPRINTF_SM11_MAGIC;
        header.fmtoffset = (unsigned short)(fmtptr - ptr);
        header.blockid = blockIdx.x + gridDim.x*blockIdx.y;
        header.threadid = threadIdx.x + blockDim.x*threadIdx.y + blockDim.x*blockDim.y*threadIdx.z;
        *(cuPrintfHeader *)(void *)ptr = header;
    }
}


//
//  cuPrintfStrncpy
//
//  This special strncpy outputs an aligned length value, followed by the
//  string. It then zero-pads the rest of the string until a 64-aligned
//  boundary. The length *includes* the padding. A pointer to the byte
//  just after the \0 is returned.
//
//  This function could overflow CUPRINTF_MAX_LEN characters in our buffer.
//  To avoid it, we must count as we output and truncate where necessary.
//
__device__ static char *cuPrintfStrncpy(char *dest, const char *src, int n, char *end)
{
    // Initialisation and overflow check
    if(!dest || !src || (dest >= end))
        return NULL;

    // Prepare to write the length specifier. We're guaranteed to have
    // at least "CUPRINTF_ALIGN_SIZE" bytes left because we only write out in
    // chunks that size, and CUPRINTF_MAX_LEN is aligned with CUPRINTF_ALIGN_SIZE.
    int *lenptr = (int *)(void *)dest;
    int len = 0;
    dest += CUPRINTF_ALIGN_SIZE;

    // Now copy the string
    while(n--)
    {
        if(dest >= end)     // Overflow check
            break;

        len++;
        *dest++ = *src;
        if(*src++ == '\0')
            break;
    }

    // Now write out the padding bytes, and we have our length.
    while((dest < end) && (((long)dest & (CUPRINTF_ALIGN_SIZE-1)) != 0))
    {
        len++;
        *dest++ = 0;
    }
    *lenptr = len;
    return (dest < end) ? dest : NULL;        // Overflow means return NULL
}


//
//  copyArg
//
//  This copies a length specifier and then the argument out to the
//  data buffer. Templates let the compiler figure all this out at
//  compile-time, making life much simpler from the programming
//  point of view. I'm assuimg all (const char *) is a string, and
//  everything else is the variable it points at. I'd love to see
//  a better way of doing it, but aside from parsing the format
//  string I can't think of one.
//
//  The length of the data type is inserted at the beginning (so that
//  the display can distinguish between float and double), and the
//  pointer to the end of the entry is returned.
//
__device__ static char *copyArg(char *ptr, const char *arg, char *end)
{
    // Initialisation check
    if(!ptr || !arg)
        return NULL;

    // strncpy does all our work. We just terminate.
    if((ptr = cuPrintfStrncpy(ptr, arg, CUPRINTF_MAX_LEN, end)) != NULL)
        *ptr = 0;

    return ptr;
}

template <typename T>
__device__ static char *copyArg(char *ptr, T &arg, char *end)
{
    // Initisalisation and overflow check. Alignment rules mean that
    // we're at least CUPRINTF_ALIGN_SIZE away from "end", so we only need
    // to check that one offset.
    if(!ptr || ((ptr+CUPRINTF_ALIGN_SIZE) >= end))
        return NULL;

    // Write the length and argument
    *(int *)(void *)ptr = sizeof(arg);
    ptr += CUPRINTF_ALIGN_SIZE;
    *(T *)(void *)ptr = arg;
    ptr += CUPRINTF_ALIGN_SIZE;
    *ptr = 0;

    return ptr;
}


//
//  cuPrintf
//
//  Templated printf functions to handle multiple arguments.
//  Note we return the total amount of data copied, not the number
//  of characters output. But then again, who ever looks at the
//  return from printf() anyway?
//
//  The format is to grab a block of circular buffer space, the
//  start of which will hold a header and a pointer to the format
//  string. We then write in all the arguments, and finally the
//  format string itself. This is to make it easy to prevent
//  overflow of our buffer (we support up to 10 arguments, each of
//  which can be 12 bytes in length - that means that only the
//  format string (or a %s) can actually overflow; so the overflow
//  check need only be in the strcpy function.
//
//  The header is written at the very last because that's what
//  makes it look like we're done.
//
//  Errors, which are basically lack-of-initialisation, are ignored
//  in the called functions because NULL pointers are passed around
//

// All printf variants basically do the same thing, setting up the
// buffer, writing all arguments, then finalising the header. For
// clarity, we'll pack the code into some big macros.
#define CUPRINTF_PREAMBLE \
    char *start, *end, *bufptr, *fmtstart; \
    if((start = getNextPrintfBufPtr()) == NULL) return 0; \
    end = start + CUPRINTF_MAX_LEN; \
    bufptr = start + sizeof(cuPrintfHeader);

// Posting an argument is easy
#define CUPRINTF_ARG(argname) \
	bufptr = copyArg(bufptr, argname, end);

// After args are done, record start-of-fmt and write the fmt and header
#define CUPRINTF_POSTAMBLE \
    fmtstart = bufptr; \
    end = cuPrintfStrncpy(bufptr, fmt, CUPRINTF_MAX_LEN, end); \
    writePrintfHeader(start, end ? fmtstart : NULL); \
    return end ? (int)(end - start) : 0;

__device__ int cuPrintf(const char *fmt)
{
	CUPRINTF_PREAMBLE;

	CUPRINTF_POSTAMBLE;
}
template <typename T1> __device__ int cuPrintf(const char *fmt, T1 arg1)
{
	CUPRINTF_PREAMBLE;
	    
	CUPRINTF_ARG(arg1);

	CUPRINTF_POSTAMBLE;
}
template <typename T1, typename T2> __device__ int cuPrintf(const char *fmt, T1 arg1, T2 arg2)
{
	CUPRINTF_PREAMBLE;
	    
	CUPRINTF_ARG(arg1);
	CUPRINTF_ARG(arg2);

	CUPRINTF_POSTAMBLE;
}
template <typename T1, typename T2, typename T3> __device__ int cuPrintf(const char *fmt, T1 arg1, T2 arg2, T3 arg3)
{
	CUPRINTF_PREAMBLE;
	    
	CUPRINTF_ARG(arg1);
	CUPRINTF_ARG(arg2);
	CUPRINTF_ARG(arg3);

	CUPRINTF_POSTAMBLE;
}
template <typename T1, typename T2, typename T3, typename T4> __device__ int cuPrintf(const char *fmt, T1 arg1, T2 arg2, T3 arg3, T4 arg4)
{
	CUPRINTF_PREAMBLE;
	    
	CUPRINTF_ARG(arg1);
	CUPRINTF_ARG(arg2);
	CUPRINTF_ARG(arg3);
	CUPRINTF_ARG(arg4);

	CUPRINTF_POSTAMBLE;
}
template <typename T1, typename T2, typename T3, typename T4, typename T5> __device__ int cuPrintf(const char *fmt, T1 arg1, T2 arg2, T3 arg3, T4 arg4, T5 arg5)
{
	CUPRINTF_PREAMBLE;
	    
	CUPRINTF_ARG(arg1);
	CUPRINTF_ARG(arg2);
	CUPRINTF_ARG(arg3);
	CUPRINTF_ARG(arg4);
	CUPRINTF_ARG(arg5);

	CUPRINTF_POSTAMBLE;
}
template <typename T1, typename T2, typename T3, typename T4, typename T5, typename T6> __device__ int cuPrintf(const char *fmt, T1 arg1, T2 arg2, T3 arg3, T4 arg4, T5 arg5, T6 arg6)
{
	CUPRINTF_PREAMBLE;
	    
	CUPRINTF_ARG(arg1);
	CUPRINTF_ARG(arg2);
	CUPRINTF_ARG(arg3);
	CUPRINTF_ARG(arg4);
	CUPRINTF_ARG(arg5);
	CUPRINTF_ARG(arg6);
	CUPRINTF_POSTAMBLE;
}
template <typename T1, typename T2, typename T3, typename T4, typename T5, typename T6, typename T7> __device__ int cuPrintf(const char *fmt, T1 arg1, T2 arg2, T3 arg3, T4 arg4, T5 arg5, T6 arg6, T7 arg7)
{
	CUPRINTF_PREAMBLE;
	    
	CUPRINTF_ARG(arg1);
	CUPRINTF_ARG(arg2);
	CUPRINTF_ARG(arg3);
	CUPRINTF_ARG(arg4);
	CUPRINTF_ARG(arg5);
	CUPRINTF_ARG(arg6);
	CUPRINTF_ARG(arg7);

	CUPRINTF_POSTAMBLE;
}
template <typename T1, typename T2, typename T3, typename T4, typename T5, typename T6, typename T7, typename T8> __device__ int cuPrintf(const char *fmt, T1 arg1, T2 arg2, T3 arg3, T4 arg4, T5 arg5, T6 arg6, T7 arg7, T8 arg8)
{
	CUPRINTF_PREAMBLE;

	CUPRINTF_ARG(arg1);
	CUPRINTF_ARG(arg2);
	CUPRINTF_ARG(arg3);
	CUPRINTF_ARG(arg4);
	CUPRINTF_ARG(arg5);
	CUPRINTF_ARG(arg6);
	CUPRINTF_ARG(arg7);
	CUPRINTF_ARG(arg8);

	CUPRINTF_POSTAMBLE;
}
template <typename T1, typename T2, typename T3, typename T4, typename T5, typename T6, typename T7, typename T8, typename T9> __device__ int cuPrintf(const char *fmt, T1 arg1, T2 arg2, T3 arg3, T4 arg4, T5 arg5, T6 arg6, T7 arg7, T8 arg8, T9 arg9)
{
	CUPRINTF_PREAMBLE;
	    
	CUPRINTF_ARG(arg1);
	CUPRINTF_ARG(arg2);
	CUPRINTF_ARG(arg3);
	CUPRINTF_ARG(arg4);
	CUPRINTF_ARG(arg5);
	CUPRINTF_ARG(arg6);
	CUPRINTF_ARG(arg7);
	CUPRINTF_ARG(arg8);
	CUPRINTF_ARG(arg9);

	CUPRINTF_POSTAMBLE;
}
template <typename T1, typename T2, typename T3, typename T4, typename T5, typename T6, typename T7, typename T8, typename T9, typename T10> __device__ int cuPrintf(const char *fmt, T1 arg1, T2 arg2, T3 arg3, T4 arg4, T5 arg5, T6 arg6, T7 arg7, T8 arg8, T9 arg9, T10 arg10)
{
	CUPRINTF_PREAMBLE;
	    
	CUPRINTF_ARG(arg1);
	CUPRINTF_ARG(arg2);
	CUPRINTF_ARG(arg3);
	CUPRINTF_ARG(arg4);
	CUPRINTF_ARG(arg5);
	CUPRINTF_ARG(arg6);
	CUPRINTF_ARG(arg7);
	CUPRINTF_ARG(arg8);
	CUPRINTF_ARG(arg9);
	CUPRINTF_ARG(arg10);

	CUPRINTF_POSTAMBLE;
}
#undef CUPRINTF_PREAMBLE
#undef CUPRINTF_ARG
#undef CUPRINTF_POSTAMBLE


//
//	cuPrintfRestrict
//
//	Called to restrict output to a given thread/block.
//	We store the info in "restrictRules", which is set up at
//	init time by the host. It's not the cleanest way to do this
//	because it means restrictions will last between
//	invocations, but given the output-pointer continuity,
//	I feel this is reasonable.
//
__device__ void cuPrintfRestrict(int threadid, int blockid)
{
    int thread_count = blockDim.x * blockDim.y * blockDim.z;
	if(((threadid < thread_count) && (threadid >= 0)) || (threadid == CUPRINTF_UNRESTRICTED))
		restrictRules.threadid = threadid;

	int block_count = gridDim.x * gridDim.y;
	if(((blockid < block_count) && (blockid >= 0)) || (blockid == CUPRINTF_UNRESTRICTED))
		restrictRules.blockid = blockid;
}


///////////////////////////////////////////////////////////////////////////////
// HOST SIDE

#include <stdio.h>
static FILE *printf_fp;

static char *printfbuf_start=NULL;
static char *printfbuf_device=NULL;
static int printfbuf_len=0;


//
//  outputPrintfData
//
//  Our own internal function, which takes a pointer to a data buffer
//  and passes it through libc's printf for output.
//
//  We receive the formate string and a pointer to where the data is
//  held. We then run through and print it out.
//
//  Returns 0 on failure, 1 on success
//
static int outputPrintfData(char *fmt, char *data)
{
    // Format string is prefixed by a length that we don't need
    fmt += CUPRINTF_ALIGN_SIZE;

    // Now run through it, printing everything we can. We must
    // run to every % character, extract only that, and use printf
    // to format it.
    char *p = strchr(fmt, '%');
    while(p != NULL)
    {
        // Print up to the % character
        *p = '\0';
        fputs(fmt, printf_fp);
        *p = '%';           // Put back the %

        // Now handle the format specifier
        char *format = p++;         // Points to the '%'
        p += strcspn(p, "%cdiouxXeEfgGaAnps");
        if(*p == '\0')              // If no format specifier, print the whole thing
        {
            fmt = format;
            break;
        }

        // Cut out the format bit and use printf to print it. It's prefixed
        // by its length.
        int arglen = *(int *)data;
        if(arglen > CUPRINTF_MAX_LEN)
        {
            fputs("Corrupt printf buffer data - aborting\n", printf_fp);
            return 0;
        }

        data += CUPRINTF_ALIGN_SIZE;
        
        char specifier = *p++;
        char c = *p;        // Store for later
        *p = '\0';
        switch(specifier)
        {
            // These all take integer arguments
            case 'c':
            case 'd':
            case 'i':
            case 'o':
            case 'u':
            case 'x':
            case 'X':
            case 'p':
                fprintf(printf_fp, format, *((int *)data));
                break;

            // These all take double arguments
            case 'e':
            case 'E':
            case 'f':
            case 'g':
            case 'G':
            case 'a':
            case 'A':
                if(arglen == 4)     // Float vs. Double thing
                    fprintf(printf_fp, format, *((float *)data));
                else
                    fprintf(printf_fp, format, *((double *)data));
                break;

            // Strings are handled in a special way
            case 's':
                fprintf(printf_fp, format, (char *)data);
                break;

            // % is special
            case '%':
                fprintf(printf_fp, "%%");
                break;

            // Everything else is just printed out as-is
            default:
                fprintf(printf_fp, format);
                break;
        }
        data += CUPRINTF_ALIGN_SIZE;         // Move on to next argument
        *p = c;                     // Restore what we removed
        fmt = p;                    // Adjust fmt string to be past the specifier
        p = strchr(fmt, '%');       // and get the next specifier
    }

    // Print out the last of the string
    fputs(fmt, printf_fp);
    return 1;
}


//
//  doPrintfDisplay
//
//  This runs through the blocks of CUPRINTF_MAX_LEN-sized data, calling the
//  print function above to display them. We've got this separate from
//  cudaPrintfDisplay() below so we can handle the SM_10 architecture
//  partitioning.
//
static int doPrintfDisplay(int headings, int clear, char *bufstart, char *bufend, char *bufptr, char *endptr)
{
    // Grab, piece-by-piece, each output element until we catch
    // up with the circular buffer end pointer
    int printf_count=0;
    char printfbuf_local[CUPRINTF_MAX_LEN+1];
    printfbuf_local[CUPRINTF_MAX_LEN] = '\0';

    while(bufptr != endptr)
    {
        // Wrap ourselves at the end-of-buffer
        if(bufptr == bufend)
            bufptr = bufstart;

        // Adjust our start pointer to within the circular buffer and copy a block.
        cudaMemcpy(printfbuf_local, bufptr, CUPRINTF_MAX_LEN, cudaMemcpyDeviceToHost);

        // If the magic number isn't valid, then this write hasn't gone through
        // yet and we'll wait until it does (or we're past the end for non-async printfs).
        cuPrintfHeader *hdr = (cuPrintfHeader *)printfbuf_local;
        if((hdr->magic != CUPRINTF_SM11_MAGIC) || (hdr->fmtoffset >= CUPRINTF_MAX_LEN))
        {
            //fprintf(printf_fp, "Bad magic number in printf header\n");
            break;
        }

        // Extract all the info and get this printf done
        if(headings)
            fprintf(printf_fp, "[%d, %d]: ", hdr->blockid, hdr->threadid);
        if(hdr->fmtoffset == 0)
            fprintf(printf_fp, "printf buffer overflow\n");
        else if(!outputPrintfData(printfbuf_local+hdr->fmtoffset, printfbuf_local+sizeof(cuPrintfHeader)))
            break;
        printf_count++;

        // Clear if asked
        if(clear)
            cudaMemset(bufptr, 0, CUPRINTF_MAX_LEN);

        // Now advance our start location, because we're done, and keep copying
        bufptr += CUPRINTF_MAX_LEN;
    }

    return printf_count;
}


//
//  cudaPrintfInit
//
//  Takes a buffer length to allocate, creates the memory on the device and
//  returns a pointer to it for when a kernel is called. It's up to the caller
//  to free it.
//
extern "C" cudaError_t cudaPrintfInit(size_t bufferLen)
{
    // Fix up bufferlen to be a multiple of CUPRINTF_MAX_LEN
    bufferLen = (bufferLen < CUPRINTF_MAX_LEN) ? CUPRINTF_MAX_LEN : bufferLen;
    if((bufferLen % CUPRINTF_MAX_LEN) > 0)
        bufferLen += (CUPRINTF_MAX_LEN - (bufferLen % CUPRINTF_MAX_LEN));
    printfbuf_len = (int)bufferLen;

    // Allocate a print buffer on the device and zero it
    if(cudaMalloc((void **)&printfbuf_device, printfbuf_len) != cudaSuccess)
		return cudaErrorInitializationError;
    cudaMemset(printfbuf_device, 0, printfbuf_len);
    printfbuf_start = printfbuf_device;         // Where we start reading from

	// No restrictions to begin with
	cuPrintfRestriction restrict;
	restrict.threadid = restrict.blockid = CUPRINTF_UNRESTRICTED;
	cudaMemcpyToSymbol(restrictRules, &restrict, sizeof(restrict));

    // Initialise the buffer and the respective lengths/pointers.
    cudaMemcpyToSymbol(globalPrintfBuffer, &printfbuf_device, sizeof(char *));
    cudaMemcpyToSymbol(printfBufferPtr, &printfbuf_device, sizeof(char *));
    cudaMemcpyToSymbol(printfBufferLength, &printfbuf_len, sizeof(printfbuf_len));

    return cudaSuccess;
}


//
//  cudaPrintfEnd
//
//  Frees up the memory which we allocated
//
extern "C" void cudaPrintfEnd()
{
    if(!printfbuf_start || !printfbuf_device)
        return;

    cudaFree(printfbuf_device);
    printfbuf_start = printfbuf_device = NULL;
}


//
//  cudaPrintfDisplay
//
//  Each call to this function dumps the entire current contents
//	of the printf buffer to the pre-specified FILE pointer. The
//	circular "start" pointer is advanced so that subsequent calls
//	dumps only new stuff.
//
//  In the case of async memory access (via streams), call this
//  repeatedly to keep trying to empty the buffer. If it's a sync
//  access, then the whole buffer should empty in one go.
//
//	Arguments:
//		outputFP     - File descriptor to output to (NULL => stdout)
//		showThreadID - If true, prints [block,thread] before each line
//
extern "C" cudaError_t cudaPrintfDisplay(void *outputFP, bool showThreadID)
{
	printf_fp = (FILE *)((outputFP == NULL) ? stdout : outputFP);

    // For now, we force "synchronous" mode which means we're not concurrent
	// with kernel execution. This also means we don't need clearOnPrint.
	// If you're patching it for async operation, here's where you want it.
    bool sync_printfs = true;
	bool clearOnPrint = false;

    // Initialisation check
    if(!printfbuf_start || !printfbuf_device || !printf_fp)
        return cudaErrorMissingConfiguration;

    // To determine which architecture we're using, we read the
    // first short from the buffer - it'll be the magic number
    // relating to the version.
    unsigned short magic;
    cudaMemcpy(&magic, printfbuf_device, sizeof(unsigned short), cudaMemcpyDeviceToHost);

    // For SM_10 architecture, we've split our buffer into one-per-thread.
    // That means we must do each thread block separately. It'll require
    // extra reading. We also, for now, don't support async printfs because
    // that requires tracking one start pointer per thread.
    if(magic == CUPRINTF_SM10_MAGIC)
    {
        sync_printfs = true;
	    clearOnPrint = false;
        int blocklen = 0;
        char *blockptr = printfbuf_device;
        while(blockptr < (printfbuf_device + printfbuf_len))
        {
            cuPrintfHeaderSM10 hdr;
            cudaMemcpy(&hdr, blockptr, sizeof(hdr), cudaMemcpyDeviceToHost);

            // We get our block-size-step from the very first header
            if(hdr.thread_buf_len != 0)
                blocklen = hdr.thread_buf_len;

            // No magic number means no printfs from this thread
            if(hdr.magic != CUPRINTF_SM10_MAGIC)
            {
                if(blocklen == 0)
                {
                    fprintf(printf_fp, "No printf headers found at all!\n");
                    break;                              // No valid headers!
                }
                blockptr += blocklen;
                continue;
            }

            // "offset" is non-zero then we can print the block contents
            if(hdr.offset > 0)
            {
                // For synchronous printfs, we must print from endptr->bufend, then from start->end
                if(sync_printfs)
                    doPrintfDisplay(showThreadID, clearOnPrint, blockptr+CUPRINTF_MAX_LEN, blockptr+hdr.thread_buf_len, blockptr+hdr.offset+CUPRINTF_MAX_LEN, blockptr+hdr.thread_buf_len);
                doPrintfDisplay(showThreadID, clearOnPrint, blockptr+CUPRINTF_MAX_LEN, blockptr+hdr.thread_buf_len, blockptr+CUPRINTF_MAX_LEN, blockptr+hdr.offset+CUPRINTF_MAX_LEN);
            }

            // Move on to the next block and loop again
            blockptr += hdr.thread_buf_len;
        }
    }
    // For SM_11 and up, everything is a single buffer and it's simple
    else if(magic == CUPRINTF_SM11_MAGIC)
    {
	    // Grab the current "end of circular buffer" pointer.
        char *printfbuf_end = NULL;
        cudaMemcpyFromSymbol(&printfbuf_end, printfBufferPtr, sizeof(char *));

        // Adjust our starting and ending pointers to within the block
        char *bufptr = ((printfbuf_start - printfbuf_device) % printfbuf_len) + printfbuf_device;
        char *endptr = ((printfbuf_end - printfbuf_device) % printfbuf_len) + printfbuf_device;

        // For synchronous (i.e. after-kernel-exit) printf display, we have to handle circular
        // buffer wrap carefully because we could miss those past "end".
        if(sync_printfs)
            doPrintfDisplay(showThreadID, clearOnPrint, printfbuf_device, printfbuf_device+printfbuf_len, endptr, printfbuf_device+printfbuf_len);
        doPrintfDisplay(showThreadID, clearOnPrint, printfbuf_device, printfbuf_device+printfbuf_len, bufptr, endptr);

        printfbuf_start = printfbuf_end;
    }
    else
        ;//printf("Bad magic number in cuPrintf buffer header\n");

    // If we were synchronous, then we must ensure that the memory is cleared on exit
    // otherwise another kernel launch with a different grid size could conflict.
    if(sync_printfs)
        cudaMemset(printfbuf_device, 0, printfbuf_len);

    return cudaSuccess;
}

// Cleanup
#undef CUPRINTF_MAX_LEN
#undef CUPRINTF_ALIGN_SIZE
#undef CUPRINTF_SM10_MAGIC
#undef CUPRINTF_SM11_MAGIC

#endif

cuPrintf.cuh (also by Nvidia):


/*
	Copyright 2009 NVIDIA Corporation.  All rights reserved.

	NOTICE TO LICENSEE:   

	This source code and/or documentation ("Licensed Deliverables") are subject 
	to NVIDIA intellectual property rights under U.S. and international Copyright 
	laws.  

	These Licensed Deliverables contained herein is PROPRIETARY and CONFIDENTIAL 
	to NVIDIA and is being provided under the terms and conditions of a form of 
	NVIDIA software license agreement by and between NVIDIA and Licensee ("License 
	Agreement") or electronically accepted by Licensee.  Notwithstanding any terms 
	or conditions to the contrary in the License Agreement, reproduction or 
	disclosure of the Licensed Deliverables to any third party without the express 
	written consent of NVIDIA is prohibited.     

	NOTWITHSTANDING ANY TERMS OR CONDITIONS TO THE CONTRARY IN THE LICENSE AGREEMENT, 
	NVIDIA MAKES NO REPRESENTATION ABOUT THE SUITABILITY OF THESE LICENSED 
	DELIVERABLES FOR ANY PURPOSE.  IT IS PROVIDED "AS IS" WITHOUT EXPRESS OR IMPLIED 
	WARRANTY OF ANY KIND. NVIDIA DISCLAIMS ALL WARRANTIES WITH REGARD TO THESE 
	LICENSED DELIVERABLES, INCLUDING ALL IMPLIED WARRANTIES OF MERCHANTABILITY, 
	NONINFRINGEMENT, AND FITNESS FOR A PARTICULAR PURPOSE.   NOTWITHSTANDING ANY 
	TERMS OR CONDITIONS TO THE CONTRARY IN THE LICENSE AGREEMENT, IN NO EVENT SHALL 
	NVIDIA BE LIABLE FOR ANY SPECIAL, INDIRECT, INCIDENTAL, OR CONSEQUENTIAL DAMAGES, 
	OR ANY DAMAGES WHATSOEVER RESULTING FROM LOSS OF USE, DATA OR PROFITS,	WHETHER 
	IN AN ACTION OF CONTRACT, NEGLIGENCE OR OTHER TORTIOUS ACTION,  ARISING OUT OF 
	OR IN CONNECTION WITH THE USE OR PERFORMANCE OF THESE LICENSED DELIVERABLES.  

	U.S. Government End Users. These Licensed Deliverables are a "commercial item" 
	as that term is defined at  48 C.F.R. 2.101 (OCT 1995), consisting  of 
	"commercial computer  software"  and "commercial computer software documentation" 
	as such terms are  used in 48 C.F.R. 12.212 (SEPT 1995) and is provided to the 
	U.S. Government only as a commercial end item.  Consistent with 48 C.F.R.12.212 
	and 48 C.F.R. 227.7202-1 through 227.7202-4 (JUNE 1995), all U.S. Government 
	End Users acquire the Licensed Deliverables with only those rights set forth 
	herein. 

	Any use of the Licensed Deliverables in individual and commercial software must 
	include, in the user documentation and internal comments to the code, the above 
	Disclaimer and U.S. Government End Users Notice.
 */

#ifndef CUPRINTF_H
#define CUPRINTF_H

/*
 *	This is the header file supporting cuPrintf.cu and defining both
 *	the host and device-side interfaces. See that file for some more
 *	explanation and sample use code. See also below for details of the
 *	host-side interfaces.
 *
 *  Quick sample code:
 *
	#include "cuPrintf.cu"
 	
	__global__ void testKernel(int val)
	{
		cuPrintf("Value is: %d\n", val);
	}

	int main()
	{
		cudaPrintfInit();
		testKernel<<< 2, 3 >>>(10);
		cudaPrintfDisplay(stdout, true);
		cudaPrintfEnd();
        return 0;
	}
 */

///////////////////////////////////////////////////////////////////////////////
// DEVICE SIDE
// External function definitions for device-side code

// Abuse of templates to simulate varargs
__device__ int cuPrintf(const char *fmt);
template <typename T1> __device__ int cuPrintf(const char *fmt, T1 arg1);
template <typename T1, typename T2> __device__ int cuPrintf(const char *fmt, T1 arg1, T2 arg2);
template <typename T1, typename T2, typename T3> __device__ int cuPrintf(const char *fmt, T1 arg1, T2 arg2, T3 arg3);
template <typename T1, typename T2, typename T3, typename T4> __device__ int cuPrintf(const char *fmt, T1 arg1, T2 arg2, T3 arg3, T4 arg4);
template <typename T1, typename T2, typename T3, typename T4, typename T5> __device__ int cuPrintf(const char *fmt, T1 arg1, T2 arg2, T3 arg3, T4 arg4, T5 arg5);
template <typename T1, typename T2, typename T3, typename T4, typename T5, typename T6> __device__ int cuPrintf(const char *fmt, T1 arg1, T2 arg2, T3 arg3, T4 arg4, T5 arg5, T6 arg6);
template <typename T1, typename T2, typename T3, typename T4, typename T5, typename T6, typename T7> __device__ int cuPrintf(const char *fmt, T1 arg1, T2 arg2, T3 arg3, T4 arg4, T5 arg5, T6 arg6, T7 arg7);
template <typename T1, typename T2, typename T3, typename T4, typename T5, typename T6, typename T7, typename T8> __device__ int cuPrintf(const char *fmt, T1 arg1, T2 arg2, T3 arg3, T4 arg4, T5 arg5, T6 arg6, T7 arg7, T8 arg8);
template <typename T1, typename T2, typename T3, typename T4, typename T5, typename T6, typename T7, typename T8, typename T9> __device__ int cuPrintf(const char *fmt, T1 arg1, T2 arg2, T3 arg3, T4 arg4, T5 arg5, T6 arg6, T7 arg7, T8 arg8, T9 arg9);
template <typename T1, typename T2, typename T3, typename T4, typename T5, typename T6, typename T7, typename T8, typename T9, typename T10> __device__ int cuPrintf(const char *fmt, T1 arg1, T2 arg2, T3 arg3, T4 arg4, T5 arg5, T6 arg6, T7 arg7, T8 arg8, T9 arg9, T10 arg10);


//
//	cuPrintfRestrict
//
//	Called to restrict output to a given thread/block. Pass
//	the constant CUPRINTF_UNRESTRICTED to unrestrict output
//	for thread/block IDs. Note you can therefore allow
//	"all printfs from block 3" or "printfs from thread 2
//	on all blocks", or "printfs only from block 1, thread 5".
//
//	Arguments:
//		threadid - Thread ID to allow printfs from
//		blockid - Block ID to allow printfs from
//
//	NOTE: Restrictions last between invocations of
//	kernels unless cudaPrintfInit() is called again.
//
#define CUPRINTF_UNRESTRICTED	-1
__device__ void cuPrintfRestrict(int threadid, int blockid);



///////////////////////////////////////////////////////////////////////////////
// HOST SIDE
// External function definitions for host-side code

//
//	cudaPrintfInit
//
//	Call this once to initialise the printf system. If the output
//	file or buffer size needs to be changed, call cudaPrintfEnd()
//	before re-calling cudaPrintfInit().
//
//	The default size for the buffer is 1 megabyte. For CUDA
//	architecture 1.1 and above, the buffer is filled linearly and
//	is completely used;	however for architecture 1.0, the buffer
//	is divided into as many segments are there are threads, even
//	if some threads do not call cuPrintf().
//
//	Arguments:
//		bufferLen - Length, in bytes, of total space to reserve
//		            (in device global memory) for output.
//
//	Returns:
//		cudaSuccess if all is well.
//
extern "C" cudaError_t cudaPrintfInit(size_t bufferLen=1048576);   // 1-meg - that's enough for 4096 printfs by all threads put together

//
//	cudaPrintfEnd
//
//	Cleans up all memories allocated by cudaPrintfInit().
//	Call this at exit, or before calling cudaPrintfInit() again.
//
extern "C" void cudaPrintfEnd();

//
//	cudaPrintfDisplay
//
//	Dumps the contents of the output buffer to the specified
//	file pointer. If the output pointer is not specified,
//	the default "stdout" is used.
//
//	Arguments:
//		outputFP     - A file pointer to an output stream.
//		showThreadID - If "true", output strings are prefixed
//		               by "[blockid, threadid] " at output.
//
//	Returns:
//		cudaSuccess if all is well.
//
extern "C" cudaError_t cudaPrintfDisplay(void *outputFP=NULL, bool showThreadID=false);

#endif  // CUPRINTF_H

Makefile


HELLOWORLDLIB := libHelloWorld.a

all : $(HELLOWORLDLIB)

CUDA_INSTALL_PATH ?= /usr/local/cuda

NVCC       := $(CUDA_INSTALL_PATH)/bin/nvcc 
CXX	     := g++
ARCHIVER   := ar cqs

TARGETDIR := ../lib
TARGET := $(TARGETDIR)/$(HELLOWORLDLIB)

VERBOSE :=

CUDAINCLUDES  += -I$(CUDA_INSTALL_PATH)/include
COMMONFLAGS += -DUNIX

CXXFLAGS := \
	-W -Wall \
	-Wimplicit \
	-Wswitch \
	-Wformat \
	-Wchar-subscripts \
	-Wparentheses \
	-Wmultichar \
	-Wtrigraphs \
	-Wpointer-arith \
	-Wcast-align \
	-Wreturn-type \
	-Wno-unused-function \
	$(SPACE)
	
NVCCFLAGS := \
	-c -Xopencc \
	-OPT:unroll_size=200000 

# Debug/release configuration
ifeq ($(dbg),1)
	COMMONFLAGS += -g
	NVCCFLAGS   += -D_DEBUG
	CXXFLAGS	+= -D_DEBUG
	CFLAGS    += -D_DEBUG
	OBJDIR   := debug
	LIBSUFFIX   := D
else 
	COMMONFLAGS += -O2 
	OBJDIR   := release
	LIBSUFFIX   :=
	NVCCFLAGS   += --compiler-options -fno-strict-aliasing
	CXXFLAGS	+= -fno-strict-aliasing
	CFLAGS    += -fno-strict-aliasing
endif
	
CUDALIB := -L$(CUDA_INSTALL_PATH)/lib
CUDALIB += -lcudart -lcutil

# Add common flags
NVCCFLAGS += $(COMMONFLAGS) $(COMMONINCLUDES) $(CUDAINCLUDES)
CFLAGS  += $(COMMONFLAGS) $(COMMONINCLUDES) 
CXXFLAGS  += $(COMMONFLAGS) $(COMMONINCLUDES) 

CUDAOBJS := \
	$(OBJDIR)/HelloWorld.cu.o 

$(HELLOWORLDLIB): directories $(CUDAOBJS)
	$(ARCHIVER) $(TARGET) $(CUDAOBJS)

$(OBJDIR)/HelloWorld.cu.o : HelloWorld.cu $(CU_DEPS)
	$(VERBOSE)$(NVCC) $(NVCCFLAGS) -I. -o $(OBJDIR)/HelloWorld.cu.o -c HelloWorld.cu

directories:
	$(VERBOSE)mkdir -p $(OBJDIR)
	$(VERBOSE)mkdir -p $(TARGETDIR)

clean:
	$(VERBOSE)rm -r $(OBJDIR)
	$(VERBOSE)rm -r $(TARGET)

Be sure to change any relevant paths in the makefile. I save all of these files in a folder called cuda/src/, then if I issue:


cd cuda/src
make

I build the static library cuda/lib/libHelloWorld.a

Building a qt app

My qt app consists of four files located in the same directory as the cuda subdirectory mentioned above:

HelloWorldQt.pro


INCLUDEPATH += cuda/src
        
CUDA_LIBDIR = /usr/local/cuda/lib
CUDALIB = -L$$CUDA_LIBDIR -lcudart

SOURCES += main.cpp \
	   HelloButton.cpp
HEADERS += HelloButton.h

LIBS += -Lcuda/lib -lHelloWorld $$CUDALIB

HelloButton.cpp


#include <HelloWorld.cuh>
#include <HelloButton.h>

HelloButton::HelloButton(const QString & text, QWidget * parent) : 
  QPushButton(text, parent)
{
}

void HelloButton::on_clicked()
{
  HelloWorld();
}

HelloButton.h


#include <QPushButton>
class HelloButton : public QPushButton
{
  Q_OBJECT
  public:
  public:
    HelloButton(const QString & text, QWidget * parent = 0);
    virtual ~HelloButton(){};
  public slots:
    void on_clicked();
};

main.cpp


#include <QApplication>
#include <HelloButton.h>
int main(int argc, char * argv[])
{
  QApplication app(argc, argv);
  HelloButton hello_button("Hello, GPU!");
  QObject::connect(
    &hello_button,
    SIGNAL(clicked()), 
    &hello_button, 
    SLOT(on_clicked()));
  hello_button.show();
  return app.exec();
}

Now you can generate an Xcode project using Qmake:


qmake-mac -spec macx-xcode HelloWorldQt.pro

Building and running with Xcode

There’s some trickiness getting executables to run correctly when linking to cuda libraries. If you just build and run the project generate by the above you may see errors like:


dyld: Library not loaded: @rpath/libcudart.dylib
  Referenced from: /Users/ajx/Code/Cuda/HelloWorldQt/build/Debug/HelloWorldQt.app/Contents/MacOS/HelloWorldQt
  Reason: image not found

There are a few ways to fix this. I prefer this simple one, but the down side is that the final app must be run from Xcode.

Then open HelloWorldQt.xcodeproj, in the side bar open Executables > Right click on HelloWorldQt and select Get Info. Then click the Arguments tab, add a new variable “to be set in the environment”:
Name: DYLD_LIBRARY_PATH
Value: /usr/local/lib/cuda

This will let you build and run from Xcode, to run the app NOT via Xcode you will have to do fancy stuff with otool that I’m not bothering with as of yet.

Download project tree source code

AntTweakBar GLUT example app using Xcode

Monday, October 11th, 2010

Here is a short tutorial to create a simple GLUT xcode project using the prototyping UI library AntTweakBar. I expect that you have compiled AntTweakBar as a static libary.

Create an xcode project

Open Xcode. File > New Project. Then choose Other > Empty Project. (The dialog looks different for different versions of Xcode, so just find “Empty Project” some place).
anttweak bar glut example tutorial

Call your project: AntTweakBarGLUTStarter
anttweak bar glut example tutorial

Add a target executable

Right-click on Targets then select Add > New Target… (or go to Project > New Target…)
anttweak bar glut example tutorial

Choose Cocoa > Application as the template of your new target (Don’t worry we’re not really making a cocoa app, it’s still going to be in pure GLUT. This just helps make the bundle).
anttweak bar glut example tutorial

Call your target: AntTweakBarGLUTStarter
anttweak bar glut example tutorial

Important:
Under the Build settings of your new target (Right-click on the AntTweakBarGLUTStarter target and select Get Info, then choose the Build tab), you need to remove the GCC_PREFIX_HEADER entry.
anttweak bar glut example tutorial

anttweak bar glut example tutorial

Link to libraries

Right-click on your target and select Add > Existing Frameworks.
anttweak bar glut example tutorial

This will open up your target’s general info tab. At the bottom left click the “+” sign to add linked libraries. Select GLUT.framework and OpenGL.framework.
anttweak bar glut example tutorial

Click the “+” sign again and then “Add Other…”. Add your libAntTweakBar.a file. I always copy my external dependencies into a folder called [my xcode project]/external. So mine is located at AntTweakBarGLUTStarter/external/AntTweakBar/lib/libAntTweakBar.a
anttweak bar glut example tutorial

This will prompt you with some options. Just choose “add”.
anttweak bar glut example tutorial

Add AntTweakBar header

I’m not sure if there is a more pleasant way to do this. I add the AntTweakBar.h file directly to my xcode project. Right-click on AntTweakBarGLUTStarter > Add > Existing Files… .
anttweak bar glut example tutorial

Then add AntTweakBar.h (mine is located at AntTweakBarGLUTStarter/external/AntTweakBar/include/AntTweakBar.h).
anttweak bar glut example tutorial

Again, just click “add” at the prompt.
anttweak bar glut example tutorial

Create a main program

Right-click on AntTweakBarGLUTStarter > Add > New File.
anttweak bar glut example tutorial

Choose C++ file.
anttweak bar glut example tutorial

Call it main.cpp
anttweak bar glut example tutorial

Now you are ready to put whatever you want in your program. I copied the contents from the AntTweakBar GLUT example (TwSimpleGLUT.c) into main.cpp. This works almost as is. Just add at the very top you:


#ifdef __APPLE__
#define _MACOSX
#endif

Build and go

anttweak bar glut example tutorial

Adding an icon

You’ll see in the above screen shot that I have a beautiful dock icon for my app. This is very easy in the current setup.

Make some icon using Photoshop/Gimp whatever you like. I just save the image as a .png file and use http://iconverticons.com/ to convert it to a .icns file.

Once you have the .icns file add it to your project. I made a new folder AntTweakBarGLUTStarter/data and put my icon there: AntTweakBarGLUTStarter/data/AntTweakBarGLUTStarter.icns. So Right-click on AntTweakBarGLUTStarter > Add > Existing File.
anttweak bar glut example tutorial

Choose your icon file.
anttweak bar glut example tutorial

Again, just click “add” at the prompt.
anttweak bar glut example tutorial

To set this file as your icon Right-click on your target and select Get Info…
anttweak bar glut example tutorial

Then choose the Properties tab. Under “Icon file:” type the name of your icon file.
anttweak bar glut example tutorial

Download

The source and .xcodeproj file in a zipped directory
The application binary (10.5 32-bits)

Source: The above is largely based off of: http://blog.onesadcookie.com/2007/12/xcodeglut-tutorial.html

Build and run current xcode project from command line

Friday, June 18th, 2010

I edit my source code via the command line using vim. Then I need to go into xcode and push build/run to debug the code I’m working on. I’d rather just issue a command from vim. Here’s a cheap way to build and run via and applescript which can then be called by vim.
Save the following applescript as xcode_build_and_launch.scpt


tell application "Xcode"
  activate
  set targetProject to project of active project document
	
  if (build targetProject) starts with "Build succeeded" then
    launch targetProject
  end if
end tell

Then from vim you can issue:


:!osascript path/to/xcode_build_and_launch.scpt

This will activate xcode and build the current project, and if that succeeds then it will launch the current executable target.

Note: I stripped the above out of code from a question on stackoverflow. Apparently this will not work for iPhone simulator projects.

Update: I like to call it from vim with the bash time command and then I have an idea of how long the compile took.


:!time osascript path/to/xcode_build_and_launch.scpt