Spatchcock

On to WordPress

2011-12-13T12:26:00.002-05:00

As a demonstration, I followed WordPress's five-minute installation process and the results are fairly gratifying. Thus, I plan on ceasing future updates to this blog (I write about 30 posts a year and publish 20 of them; it's not like they're that frequent).

Look for future updates at http://akerr.net/andrew/blog/.

See you there.

Software Rasterization

2011-09-01T23:03:00.002-04:00

The previous post awakened the part of me that is somewhat passionate about graphics programming. In the last three nights, I threw together a software rasterizer that is now capable of perspective-correct texture mapping in addition to point lights.

Here is a screenshot:

And here is a video that VLC should have no trouble playing.

Some details:

The rasterization algorithm is very primitive. Triangle setup first ensures clockwise vertex ordering (cross product of the two edge vectors followed by a sign comparison) and culls backfaces. Next, a transformation to Barycentric coordinates is computed, and blocks of pixels in the render target are transformed and compared. If their Barycentric coordinates are found to lie within the triangle, they are used to interpolate vertex color and texture samples.

It is not yet GPU accelerated, though expect a CUDA implementation to follow.

ASCII Torus

2011-07-22T11:59:00.003-04:00

Inspired by this guy's 3D torus rendered to an ASCII palette, I decided to spend an hour last night implementing something similar. Here's the source, and here's the result. It was a fun exercise:


                             @@@@@@@                            
                    @@@@$##*!!!===!!!*##$@@@@                   
               @@@$$#*!=;;:~~-~~-~~~:~:;==!*#$$@@@              
            $@$$$#*!=;:~~~-~-----~~~~~~:::;=!!*#$$$@$           
          $$$$$##*=;~~~-----~~~~:::::::::~::;=!**#$$$$$         
        *$$$$$#*!=:~~------~~~~~::::::::::::::;!**#$$$$$#       
       *#$$$$$#*=:~--,,,--             ~:::::::=!*##$$$$#*      
      !*#$$$$$#*=~-,,,                     ~:~::!*#$$$$$#*!     
     ;!*#$$$$$#*=,.                           ~~!*#$$$$$#*!;    
    -;!*#$$$$$$#!,                             ~*#$$$$$$#*!;,   
    ~:;==**#$$$@@@$                           #@@@$$$#**!=;~,   
    ~~~:;=!*###$$@@@@                       @@@@$$###*!=;:~-,   
     ~~~::;=!***#$$$@@@@@@             @@@@@@$$$#***!=;:~~--    
     :~~~~~::;=!!**###$$$$@@@@@@@@@@@@@$$$$###**!!=;;:~~---,    
      ::~~~~~~~:;;==!!****#############****!!==;;;:~~-----,     
        ::~~~~-~~~~::;;=====!!!!!!!!!=====;;:::~---------       
         :::~~~~~~~-----~~:::::::::::::~~~-------------,        
           :::::~~~~~-~------------------------------,          
              :::::~~~~~~~~-----------------------,             
                 :::::::~~~~~~~~~~~~~-----------                
                      :::::::~~~~~~~~~~~~~-

And here it is rendered by the same utility to a higher resolution.

Screencapture Utility

2011-07-20T12:03:00.006-04:00

I wrote a quick'n'dirty alternative to Skitch for Linux because Shutter was always so crappy. It's a simple Python script that enables on-the-fly cropping and immediate upload to the configured server. Took about ten minutes to research, write, and debug.

Here's a sample:

Alpine Ascents Successful Climb of Mount Everest

2011-06-08T12:55:00.003-04:00

In brighter news, Seattle-based Alpine Ascents is currently wrapping up a successful summit assault on both Everest and nearby Lhotse a day later. A topological map illustrates the incredible vertical and horizontal distances covered.

This shows the route from base camp, traversal to Camp I across the Khumbu Icefall, Advanced Base Camp II at 21,000 ft, Camp III at 24,000 ft, Lhotse High Camp, and finally Camp IV at 26,000ft. From Camp IV, the route to the summit ascends 3,000 ft during a grueling 12 hour climb requiring a 9pm departure time to reach the summit by the next morning. The elevation profile shows the several returns to base camp from Camp II during the acclimatization period. It also shows the ascent of Lhotse by three of the climbers from the Everest group prior to the group's descent.

Alpine Ascents does a pretty thorough job listing their itinerary and equipment recommendations. Check it out, particularly the group electronics gear.

Into Thin Air

2011-06-08T11:39:00.010-04:00

I recently finished reading Into Thin Air, Jon Krakauer's incredible description of his summit attempt on Mount Everest during what became the 1996 Mount Everest Disaster. The narrative describes Krakauer's experience during the heyday of commercial mountaineering guide services. As a journalist writing for Outside magazine, Krakauer described what was hoped to be a routine climb along Mount Everest's South Col summit route used by Hillary and Norgay during their record-setting first summit in 1953, now the standard route.

During summit day, slow progress by inexperienced climbers and judgment errors by guides set the stage for tragedy when a storm quickly materialized along the southern ridge. Blowing 40 knots and producing incredible wind chills, many tardy climbers including the leader of Krakauer's expedition were trapped high on the upper mountain. By this time, Krakauer himself had nearly made it to high camp in the col and narrowly escaped incredible amounts of exposure in the high winds. Lack of visibility thwarted rescue attempts, and by the next day eight members from two expeditions had perished while others suffered extreme frost bite.

An expedition filming the IMAX documentary Everest offered assistance and supplies such as additional bottled oxygen in addition to arranging a helicopter evacuation of two frost bitten climbers. This film hauntingly documents their perspective of the disaster.

One of the most striking components of this story is how altitude clouded the judgments of the people involved. Several times, someone attempting to adjust their oxygen regulators would get it wrong and either turn off the flow of O2 entirely or drain the tank in minutes. Tragically, one of the guides who had reached high camp during the storm incorrectly stated that all of the oxygen bottles were empty when, in fact, at least two were not. Krakauer speculates this incorrect information changed the plans of Rob Hall, the lead climber, who elected to stay with a client near Hillary Step, just below the summit, where both he and his client perished.

The 1996 Everest Disaster demonstrates the consequences of hubris high on a mountain. One of the guides from an American expedition wasn't climbing with oxygen. While this might be a personal challenge to overcome at other times, doing so while employed to take care of others is exceedingly dangerous. When disaster struck, he was unable to assist; he couldn't even wait around with a distressed climber, for the lack of oxygen slows human metabolism such that the only substantial source of internal heat is produced from muscles as they are operating. Additionally, the many climbers from multiple commercial expeditions created bottlenecks that delayed summit times and increased exposure to weather-related hazards; if the storm had struck even an hour later, many more people are likely to have made it to [limited] shelter in the South Col. Finally, turn around times were ignored at great peril.

One of the other lessons here is no climb is guaranteed, and achieving the summit is only the half-way point. Mountaineering is inherently quite dangerous, particularly above 8000m where oxygen levels are so low that it's virtually impossible to carry someone. Anyone considering undertaking such a venture should come to terms with the fact that rescue is nearly impossible; doubts in one's abilities, fitness levels, and weather conditions should be headed. Underscoring this point, most of the climbing along the upper mountain is done without ropes linking the climbers; falls are so perilous on the steep slopes that doing so would simply increase the scope of one person's tragedy.

Coquette

2011-05-11T10:25:00.003-04:00

WindFinder is a coquette and a tease. Spectacular sailing conditions are always forecast four days out. When time moves forward, things always settle down to calm conditions that aren't worth the drive up.

Summer must be upon us.

DEC is Dead

2011-04-10T22:08:00.003-04:00

I was perusing Amazon and came upon this title: DEC is Dead. Long Live DEC. This book describes the rise and fall of Digital Equipment Corporation, a legend in the computing industry that introduced many innovations in processor and system design. Many of the architects went on to other companies such as AMD and Transmetta and continued to advance the state of the art. It's an appealing topic, to be sure.

I wanted to point that out.

Additionally, the title seems to confuse the etymological implications of "The x is dead. Long live y." In this case, even if x and y appear to be the same improper noun, lexically, they never refer to the same individual or entity; that'd be a contradiction. The phrase actually refers to a transitional event, where the life of an old entity, x, is over and we should now bestow praise and hope on the life of a new entity, y. See this for a foothold on more information.

I've seen a lot of smart people get this wrong, and evidently someone solidified their mistake as the title of their book. In this case, I suggest DEC is dead. Long live Intel, AMD, and Transmetta.

You're welcome.

Sailing Adventure

2011-04-04T13:27:00.003-04:00

Saturday, I went sailing 420s on Lake Lanier with a labmate. Of five boats that went out, we were one of two that didn't both capsize and require a rescue. The other boat that avoided capsize spent much of the morning with only the jib flying and left early.

Winds were intense, from 15 to 20 mph, and waves were high. Our 420 achieved planing speed several times; the sudden boost in speed and rudder authority was reminiscent of a jetski.

Windfinder provides a great and detailed wind forecast.

Exploring R.M.S. Queen Mary

2011-03-10T15:02:00.016-05:00

RMS Queen Mary is an ocean liner from the Cunard Line in service from 1936 to 1967. Propelled by four screws, each moved by a collection of high-speed steam turbines, Queen Mary offered the fastest civilian Atlantic crossings available. Her record setting voyage of 3 days, 21 hours earned her the Blue Riband which she held for 14 years. Unlike many of the other liners of the day, Queen Mary survived her service life and found a permanent home along the coast of Long Beach, California, where she now accommodates hotel guests, numerous restaurants, tourists, and a community center built into her forward engine room and aft boiler rooms.

On Sunday morning, I drove up to Long Beach along the Pacific Coast Highway to spend as much of the day as possible exploring Queen Mary's decks and interior spaces, hoping to see as much of the ship as possible. Displacing 81,000 tons and with a length of 1,020 ft, I would surely run out of daylight long before I ran out of things to see.

Hopefully this blog post captures some of the highlights of my visit.

Forward well deck, superstructure, and bridge of RMS Queen Mary

Port side showing bridge wing, hull, and superstructure.

Various controls and telegraphs in the aft engine room. The forward engine room housing propulsion machinery for the outer screws has been emptied of contents, and a public area has been constructed in its place. I think they may have removed part of the bulkhead separating it from the aft-most boiler room.

Controls and telegraphs in the aft engine room. The telegraph is how the bridge suggests a speed and direction to the engineers who actually adjust the steam valves accordingly. Turbines aren't typically reversible, so a separate set of high-pressure and intermediate pressure turbines are provided for moving astern. Note, the "astern" valves are open in spite of the request to move "ahead slow;" this is an unrealistic configuration.

The engine room is large, and structural elements are clearly visible. Several of the turbines and gearboxes are in place and exposed revealing large cast iron gearing and the lowest-pressure turbine rotor. The author appreciates seeing exposed machinery.

Looking forward along the port side at the Promenade (A) Deck level. Many of the internal passage ways inside the ship are blocked, so getting from the engine room to the promenade deck can only be completed by this external catwalk system. Lame.

It's a long way down from the boat deck to the water line. Ripples indicate 1 on the Beaufort scale.

Looking forward, view of B deck, forward well deck, and the forecastle. The rifle battery is an homage to her WWII service. Queen Mary holds the record for largest number of people transported on a ship at any given time (16,000 or so).

I never look happy about anything. I am, though.

Queen Mary's outer port propeller is still on its shaft. A tank built around the ship and illumination sources make it visible. It's an incredible sight. In viewing this, I realized QM is actually grounded and not afloat making her immune to the tides and quite stationary.

I was clearly here all day.

RMS Queen Mary's bridge.

More images are visible in this directory.

Ferroequinology: More Running Boards

2011-02-21T10:31:00.006-05:00

Saturday, I revisited SERM to spend some time on A&WP 290. The previous visit, we cleared several brackets and attached steam and air pipes to expose a nearly clear boiler shell. In preparation for Ultrasonic Thickness testing, a grid over each course of the shell must be arranged and test points ground to expose bare metal. To access the entire shell, running boards on the fireman's (left) side of the locomotive needed to be installed.

Installing the fireman's side running boards used all of the mounting hardware remaining and, to our chagrin, exposed mistakes in mounting the engineer's side running boards. This image depicts the two large castings that support both running boards and compressed air tanks. Each of the four castings is distinct, not correctly labeled, and not easily positioned even with a crane (of which I have become a precise and efficient operator).

Ultimately, we had to remove the engineer's side running boards, remove what we guessed was the incorrectly placed bracket, place a different bracket, reinstall the running board, install that bracket on the other side in both possible locations, and hopefully achieve the best fit. Moreover, studs can be bent so it's a matter of luck and judgment to determine whether a part doesn't fit because it's not the correct part or because you're not using a large enough hammer.

By end of day, with many possible combinations attempted, all air tank brackets were soundly mounted, running boards placed, and the boiler now accessible from all sides. During the next disassembly, we'll surely stamp each casting and avoid confusion in the future... ha.

Up next, ultrasonic thickness measurement over every square foot of a 37ft pressure vessel with working pressure 201 psi and operational temperature of 450 degF.

Orthographic Meet 2011

2011-02-20T14:58:00.010-05:00

Last night, Emma and I attended the Atlanta Open Orthographic Meet, an annual spelling competition held at Manuel's Tavern. My score was 19 total, although I feel bad pronunciation robbed me of of 'pharoah.' These are the words from each of four rounds with my best effort to recall the particular definition provided during the competition.

Round 1

pharoah - a tyrant
shellacked - to defeat decisively
corpuscle - an unattached body cell, such as a blood or lymph cell
artisan - a skilled manual worker; a craftsperson
ululate - to howl or wail, as with grief
metallurgy - the study of metals and their properties in bulk and at the atomic level
bulwark - wall or embankment raised as a defensive fortification
privilege - a special advantage, immunity, permission
artesian - a well from which water flows spontaneously due to internal pressure
mischief - behavior that causes discomfiture or annoyance in anothe
sauerbraten - pot roast of beef marinated in vinegar, water, wine, and spices before being cooked
poltergeist - a ghost that manifests itself by noises, rappings, and the creation of disorder
deuce - a tied score in tennis
lien - right to take and hold the property of a debtor as security for a debt
apropos - being at once opportune and to the point
rendezvous - meeting at a prearranged time and place
vaudeville - theatrical performance of this kind; a variety show
pharisaic - hypocritically self-righteous and condemnatory
apiary - place where bees are kept
bamboozle - to take in by elaborate methods of deceit

Round 2

grok - to understand profoundly through intuition
uxorious - excessively submissive or devoted to one's wife
apocopate - to omit the final sound or sounds of (a word)
diurnal - occurring or active during the daytime rather than at night
succedaneum - a substitute
poinciana - having large orange or red flowers
snollygoster - one, especially a politician, who is guided by personal advantage rather than by consistent, respectable principles
involucrar - ??
viand - a very choice or delicious dish
onomatopoeia - the formation or use of words such as buzz or murmur that imitate the sounds associated with the objects or actions they refer to
??? -
antipyretic - a medication that reduces fever
trepanation - the medical procedure of making a hole in the skull
anneal - to subject glass or metal to a process of heating and slow cooling in order to toughen and reduce brittleness
coccygodynia - pain in the coccyx or tailbone area

Round 3

kuemmel - a sweet, colorless liqueur flavored with caraway seed
zugzwang - a situation where one player is put at a disadvantage because he has to make a movel the player would prefer to pass and make no move
fissiparous - inclined to cause or undergo division into separate parts or groups
psammite - a metamorphosed rock unit with a dominantly sandstone
smaragd - having the color of emeralds
coryphee - a ballet dancer
ratamacue -
majuscule - capital letters
mirliton - a pear-shaped vegetable or its vine
claquer - an organized body of professional applauders

Round 4

cimicine -
rijstaffel - a dish originating in Indonesia; a wide variety of foods and sauces are served with rice
mbaqanga - a style of Black popular music of urban South Africa
jobbernowl - a blockhead
tchick - a slight sound such as that made by pressing the tongue against the roof of the mouth and explosively sucking out the air at one side, as in urging on a horse

Install GPU Ocelot on Ubuntu 10.04

2011-02-07T15:01:00.004-05:00

We've recently acquired a new machine to test Advanced Vector Extensions with GPU Ocelot's experimental vectorizing execution manager and code generator. I decided to document the complete set of procedures needed to build Ocelot on a fresh installation of Ubuntu 10.04.

1.) Install Ubuntu updates

2.) Install prerequisites:


sudo apt-get install g++ subversion libboost-all-dev 
sudo apt-get install glutg3-dev libglew1.5-dev glew-utils 
sudo apt-get install bison flex automake autoconf libtool

3.) Install NVIDIA CUDA 3.2 for Ubuntu 10.04. This includes an updated video driver and CUDA 3.2 Toolkit.

4.) Download, build, and install LLVM


svn co http://llvm.org/svn/llvm-project/llvm/trunk llvm
cd llvm
./configure && make ENABLE_OPTIMIZED=1 -j 8
sudo make install ENABLE_OPTIMIZED=1

5.) Download, build, and install GPU Ocelot


svn checkout http://gpuocelot.googlecode.com/svn/trunk/ gpuocelot-read-only
cd gpuocelot-read-only/ocelot
libtoolize && aclocal && autoconf && automake && ./configure
make check -j 8
sudo make install

6.) Run built in unit tests:


make test

Attention all Personnel

2011-01-31T11:20:00.004-05:00

Update your blogs, folks! I know you've been up to something; what is it? What did you learn from it?

A short observation: I discovered the level of accuracy in my technique at pool vastly improved if I decreased the distance between cue and cue ball. Short stroke trumps long stroke.

Another short observation: Under sail, when preparing for a gybe, be sure you are running directly down wind. A gybe should not involve significant rudder/tiller movements, and one's heading should not change dramatically. In this video, the nose of the boat does not turn during the gybe; the main is simply sheeted in until the wind carries it across.

Strolls through Piedmont Park are best completed with Emma and a smoothie.

Taste of AVX

2011-01-18T16:23:00.004-05:00

One can use Intel's Software Development Emulator to execute programs utilizing unsupported ISA extensions such as Advanced Vector Extensions. This tool, based on Pin, instruments binaries and emulates instructions not supported in hardware. It also provides detailed instruction opcode histograms so you can, for instance, figure out whether some video game or library was compiled with SSE support.

Here's an example x86-64 function implementing element-wise multiply-accumulate of two single-precision floating-point vectors of arbitrary length.


; void vectorFMA(float *R, float *A, float *B, int N)
;
; R[i] += A[i] * B[i] for i = 0 .. N-1
;
; uses AVX to multiply two vectors, elementwise, and add
; results to a third vector
;
; rdi - R
; rsi - A
; rdx - B
; rcx - N
;
vectorFMA:
 
.L1:
 ; load words from R, A, and B
 vmovups ymm1, [rdi]
 vmovups ymm2, [rsi]
 vmovups ymm3, [rdx]
 
 ; *R += *A * *B
 vmulps ymm4, ymm2, ymm3
 vaddps ymm5, ymm1, ymm4
 vmovups [rdi], ymm5
 
 ; R+=8, A+=8, B+=8
 add rdi, 32
 add rsi, 32
 add rdx, 32
 
 ; if (--rcx == 0) break;
 dec rcx
 jnz .L1
.L5:
 ret

When our SandyBridge machine arrives next week, I should have a handful of interesting microbenchmarks to run on it.

Ferroequinology: Running Boards

2011-01-18T10:55:00.007-05:00

We hoisted engineer's side running boards onto the brackets mounted the previous week. Now, it is possible to access the top of A&WP 290's boiler shell. We used the opportunity to heat and remove a handful of brackets to clear several sand and steam pipes.

Protip: when applying a socket wrench to a red-hot steel nut, make sure your gloves are dry to prevent any unwanted heat transfer...

Update: water soaked into gloves acts as a thermal conductor resulting in scalds if a gloved hand is applied to steel whose temperature is several hundred degrees above the boiling point of water.

Wanted

2011-01-14T14:04:00.001-05:00

My notebook. Last seen December 6, 2010.

Constant-time Set Membership

2011-01-14T10:23:00.009-05:00

While scouring the internet yesterday, I came upon an article written by Preston Briggs, currently at NVIDIA and a former office mate in Bellevue. Efficient representation for sparse sets.

This paper considers set representations of objects for which a bijection to the natural numbers exists and provides a novel alternative to bit vectors.

Sets represented by bit vectors are one of the more common approaches [what we use in Ocelot]. Given a universe of size u, bit vectors indicate whether element i is a member of a set by setting bit i providing access times on the order O(u). This representation also requires O(u) time to clear the bit vector initially. If short-lived sets with large u are needed frequently, overheads can be prohibitive. A sparse representation may be desirable, but it should have O(1) access times.

Preston's approach represents a sparse set with two arrays of size u, denoted dense and sparse, and an integer k representing the number of elements in the set. The array dense stores a packed sequence of elements that have been added to the set, and valid indices are 0 ... k-1. The array sparse provides a mapping of natural numbers to possible indices into dense. For example, assume j belongs to the set and is stored in location dense[p]; the following is true: sparse[j] = p.

With two array access, both O(1) we can efficiently determine whether element j is in the set:


def isIn(j, S):
  return (S.sparse[j] < S.k) and (S.dense[S.sparse[j]] == j)

That is, if sparse[j] is a valid index (from 0 to k-1) and the object at that index is indeed the element, then it belongs to the set. This technique has the advantage that neither dense nor sparse needs to be cleared initially. To clear the set, simply set k = 0.

Elements are added by:


def insert(S, j):
  S.dense[S.k] = j
  S.sparse[j] = S.k
  S.k += 1

Storage requirements are 2*(u+1) for a set of u integers. I thought this was interesting.

Human Knowledge

2011-01-12T19:53:00.000-05:00

PhD in pictures

Delegates in C++

2011-01-11T11:34:00.006-05:00

I reimplemented the Baker valve gear simulator in C++ a few weeks ago. This involved adding a nonlinear equation solver to libkerr's matrix library, my custom-made high-performance linear algebra library that's been reused over the years in dozens of small projects and a few large ones.

The algorithm itself is nothing special: Newton's method. Implementing it in a general and yet efficient manner in C++, however, presents a design choice.

I was tempted to accept a template argument and call that class's operator() member function.


template < typename Real, typename Functor >
int fsolve(matrix &x, const Functor &F, Real precision, int maxIterations) {
  ..
  while () {
    matrix fx = F(x);  // operator()
  }
  ..
}

This, however, would require a separate class for each constraint function. In my case of coupled systems, I need to determine roots for two distinct constraint functions that shared both constant state and analytically computed state. Additionally, the zeros of the first function would be needed as input to the second function.

Instead, I opted for the delegate design pattern. The functor includes both the object as well as references to member functions.


template <typename Real, typename Functor>
int fsolve(matrix<Real> &x, const Functor &F, matrix<Real> (Functor::*ptr)(matrix<Real>), Real precision, int maxIterations) {
  ..
  while () {
    matrix<Real> fx = (F.*ptr)(x);  // delegate
  }
  ..
}
..
BakerValve instance;
fsolve<double, BakerValve>(x1, instance, &BakerValve::constraintF1);
..
fsolve<double, BakerValve>(x2, instance, &BakerValve::constraintF2);

Some additional features of my solver are the Jacobian can either be estimated by a default solver or computed analytically by a user-supplied function.

The implementation seems to work well and always converge. I haven't implemented an aesthetically pleasing visualization of its output, as this was mostly scratching an itch for implementing the constraint solver.

Ferroequinology: clearing A&WP 290's Boiler Shell

2011-01-09T13:19:00.010-05:00

Yesterday, I visited the Southeastern Railway Museum and worked on A&WP 290 with Dale Grice.

In continuing with the goal of removing everything from the boiler shell to prepare for an eventual blast with baking soda, we removed a steam pipe carrying a "signal" that actuates cylinder cocks. A valve in the engineer's side of the cab raises steam pressure in this pipe which actuates valves at the bottoms of the main cylinders at the front of the locomotive that permit condensation and steam to blow out when the locomotive starts. This is why steam engines seem to start in a cloud of steam.

This required removing two unions coupling a network of steel pipes. I tried without luck during a previous visit to wrench them apart. Yesterday, we used a "rosebud" oxy-acetylene torch to heat the union coupling to dull cherry red then quenched it with a cup of water. The first coupling I tried this on came loose readily and after one revolution with the wrench [and plenty of time to cool], I removed it the rest of the way with torque applied by [gloved] hand.

The second union would not budge in spite of this treatment. Between heatings, I hit it with a hammer but it remained fast. Consequently, we used a cutting torch which emits a jet of oxygen gas through a flame and oxidizes steel. A short segment of pipe was cut between the union and a T-coupling thus freeing the main pipe along the boiler.

With unions apart, the cylinder cock pipe would come loose with the removal of a single bracket fastened to the boiler shell via a threaded stud. Unfortunately, the nut holding the bracket was corroded and could not be made to turn without heat. After heating to cherry red, a wrench with torque magnified by a breaker bar [steel pipe] eventually made some progress. Or we thought it was progress. After about half a revolution the nut fell off revealing we had literally torqued the 1/2" stud apart and never moved the nut. One more thing to fix on the boiler...

I manned the portable crane and we pulled the pipes off the boiler then applied a set of labels so they can be replaced or new copies made.

To reach the top of the boiler, the next step was to replace air cylinder brackets. These are large castings [weighing probably 50 or 60 pounds each] that suspend two large air cylinders along the locomotive, one on each side, and support running boards.

Again, I manned the crane lifting the castings up to the boiler while Dale aligned them with 1.5" diameter threaded studs. These studs were placed in the boiler from within when it was built and cannot be easily replaced without potentially compromising the boiler's integrity [not to mention incurring a laborious and detailed inspection]. The first bracket went on without much trouble but we discovered it was mis-labeled and didn't sit flush with the boiler shell. The next bracket didn't slide onto the studs at all, and we discovered the reason was due to the forward stud bent up and forward. The studs needed to be parallel for the bracket to slide on, and this clearly wouldn't do.

To straighten it, Dale fitted a large 12' steel pipe over the stud, and I climbed up on a scaffold. The plan was for me to bend the stud using the pipe as a lever. Doing this proved problematic; the pipe was not adequately stiff, and all force I applied deformed the pipe like a spring. I had visions of losing my grip and the pipe springing back; the last thing I would see is a 2.5" steel pipe swinging back to hit me in the face...

We abandoned this idea pretty quickly. When Dale attached a nut to the end of the stud and hit it with a sledgehammer without any success, he decided heat was the only solution. Again, the rosebud torch came out and in no time, the base of the stud was glowing cherry red. Each gentle stroke of the sledgehammer moved the end of it by about 1/4", and in no time we had straightened it.

Both brackets slid on without difficulty after this, and now we'll be able to add running boards and access the top of the boiler. Progress!

My Struggles in Trigonometry Class

2011-01-05T16:37:00.007-05:00

During a period of restlessness last night, while lying in bed I recalled a project I completed during 11th grade trigonometry class in high school. Today, I don't remember the minimum scope requirements, but I chose to satisfy them by building a game that simulated artillery. The game environment consisted of a simulated island (a 3D heightfield) and two textured stone bunkers corresponding to each player's location.

In another window with GUI controls, the player submitted elevation, azimuth, and velocity of an artillery shell. Then, they would click FIRE and a shell would launch from their location, arc over the simulated island, and impact somewhere. If it impacted the other player's bunker (or their own), the player who's bunker survived received a point and the game restarted with new random locations. The camera was attached to the shell so the firing player would have a chance to observe the island, reconnoiter their opponents location, and update their trajectory. Differences in elevation between the two players' bunkers and physical obstructions such as the island itself meant each player needed to determine a unique trajectory to the other. I believe there was also simulated wind.

I spent plenty of time working on it and made it finally ready to present the day all projects were due. No one else did anything related to computing, let alone implement a real-time 3D renderer with a game built around it. The morning before class, I installed it on the antiquated classroom PC to test it out. To my horror, the low max-resolution of the video card prevented the GUI controls from appearing; the input window was obscured and not sizable. The controls for entering artillery parameters and firing the shell were out of view. Our teacher was perpetually suspicious of technology and told me I would have to submit it as is; no deadlines would be granted. I think she may have even been slightly pleased that smarty-pants Andykerr would finally be hoist with his own petard (when taken literally, 'sent into the air by one's own powder keg' -- a remarkably delightful metaphor given the context of my game).

I had until sixth period, classtime, to resolve the issue and make the game playable. Unfortunately, I did not have the source code, and even if I did, the machines I had access to with Visual Studio lacked the DirectX software development kit. Recompiling it was impossible.

So I edited the binary instead.

Windows applications store layouts of controls as resource files embedded in the executable. These are binary also, but I knew enough to expect the X and Y coordinates of the critical controls to be stored as 16-bit unsigned integers. Using MSPaint, I took a screenshot of the application to determine what the locations of the buttons and textboxes actually were. I visited the electronics lab where computers with Visual Studio were installed, and I used Visual Studio's hex editor to search and edit the executable. Fortunately for me, the coordinates made unique binary sequences for that particular application, so finding them was straightforward. I made up new values, overwrote the old ones, and hoped.

I brought it Trigonometry during sixth period, installed it on the classroom PC, and was met with pleased astonishment. The controls were in new positions, visible, and the game could be played as designed. The textures were crappy, and the slow video card made gameplay uninteresting, but it ran and I received full credit.

3D games rely heavily on trigonometric relationships to transform and project mathematical representations of a virtual world onto a raster display, but the math was by *far* not the most difficult part of that project. I don't think anyone else appreciated what I went through for that damn grade...

A Key Insight on Heterogeneity in Computing

2010-12-28T09:34:00.002-05:00

Reading a PhD dissertation on parallel computing [link PDF].

"The Cell Broadband Engine is a heterogeneous processor by itself,.."

I've been arguing that today's modern multicore CPUs are heterogeneous processors and some folks at Intel seem to agree. Processors have multiple functional units with differing capabilities, a scheduler, and wide vector units with their own Instruction Set Architecture. Each revision of Intel and AMD x86 processors seem to include a few additional instructions to power some specialized unit such as video processing or floating point.

AndyKerr insight: every modern superscalar processor is a heterogeneous platform by itself, and the extent of heterogeneity is rising.

Fall 2010 Complete

2010-12-14T16:20:00.003-05:00

In one week, I represented GT during Micro-43, wrote a paper submitted to GPGPU-4, and completed two extensive meetings with industry sponsors. Winter 2010 will hopefully provide a much-needed reprieve.

Anyone interested in a trip to the Tennessee Valley Railroad Museum to ride #610 before her FRA boiler certificate expires? I should call up the museum to see if they plan on any steam-ups of ex Southern #630, a splendid 2-8-0 steam locomotive in the final stages of a complete restoration. Chattanooga isn't too far, but we combine trips when we can to make the most of things.

I would like to perform one sailing adventure. Things are cold, to be sure, but the winds should be stellar. If we can keep the 420 from capsizing, I'm sure we'll have one hell of a time. Robust constitutions only need apply.

A range visit is scheduled for this Friday.

Back to Christmas shopping...

Happy Thanksgiving

2010-11-24T14:20:00.004-05:00

While preparing your Thanksgiving meal, remember the following, dear readers.

Happy Thanksgiving!

B-Tree

2010-11-20T14:42:00.007-05:00

Saturday morning algorithms practice: I implemented insert() and find() members std::map<> as a B-tree.

Each node in the B-tree has at most 2*t-1 keys and at most 2*t children. The best-case depth of a B-tree is log_{2*t}(N) and the worst case is log_{t}(N). During a search, O(2t) comparisons are performed per node, and O(height) nodes are visited.

For t = 3, my implementation shows the following complexity for searching:


N = 1000
Average nodes visited: 4.68
log_3(N) = 6.28771  (worst)
log_6(N) = 3.85529  (best)
tree depth = 5

For N = 25000, you may compare to an earlier implementation of skiplists.


N = 25000
Average nodes visited: 7.68968
log_3(N) = 9.21766  (worst)
log_6(N) = 5.65178  (best)
depth = 8

LLVM Floating-point Performance Bug

2010-11-03T13:55:00.002-04:00

I've noticed an unusual behavior of the LLVM x86 code generator that results in
nearly a 4x slow-down in floating-point throughput that I would like to report.

I've written a compute-intensive microbenchmark to approach theoretical peak
throughput of the target processor by issuing a large number of independent
floating-point multiplies. The distance between dependent instructions is at
least four instructions to match the latency of the floating-point functional
unit in my Intel Core2 Quad (Q9550 at 2.83 GHz).

The microbenchmark itself replicates the following block 512 times:


 .
 .
 {
  p1 = p1 * a; 
  p2 = p2 * b;
  p3 = p3 * c; 
  p4 = p4 * d;
 }
 .
 .

Compiling with NVCC, Ocelot, and LLVM, I can confirm the interleaved instruction
schedule with a four-instruction reuse distance. An excerpt follows:


  .
  .
  %r1500 = fmul float %r1496, %r24 ; compute %1500
  %r1501 = fmul float %r1497, %r23
  %r1502 = fmul float %r1498, %r22
  %r1503 = fmul float %r1499, %r21

  %r1504 = fmul float %r1500, %r24  ; first use of %1500
  %r1505 = fmul float %r1501, %r23
  %r1506 = fmul float %r1502, %r22
  %r1507 = fmul float %r1503, %r21

  %r1508 = fmul float %r1504, %r24  ; first use of %1504
  .
  .

The JIT compiler, however, seems to break the interleaving of independent instructions
and issues a long sequence of instructions with back-to-back dependencies. It is as if
all p1 = .. expressions are collected at once followed by all p2 = .. expressions and so
forth.


 p1 = p1 * a
 p1 = p1 * a
 .
 .

 p2 = p2 * b
 p2 = p2 * b
 .
 .

 p3 = p3 * c
 p3 = p3 * c
 .
 .

An actual excerpt of the generated x86 assembly follows:


  mulss %xmm8, %xmm10
  mulss %xmm8, %xmm10
  .
  . repeated 512 times 
  .

  mulss %xmm7, %xmm9
  mulss %xmm7, %xmm9
  .
  . repeated 512 times 
  .
  
  mulss %xmm6, %xmm3
  mulss %xmm6, %xmm3
  .
  . repeated 512 times 
  .

Since p1, p2, p3, and p4 are all independent, this reordering is correct. This would have
the possible advantage of reducing live ranges of values. However, in this microbenchmark,
the number of live values is eight single-precision floats - well within SSE's sixteen
architectural registers.

Moreover, the pipeline latency is four cycles meaning back-to-back instructions with true
dependencies cannot issue immediately. The benchmark was intentionally written to avoid this
hazard but LLVM's code generator seems to ignore that when it schedules instructions.

When I run this benchmark on my 2.83 GHz CPU [45.3 GFLOP/s stated single-precision throughput],
I observe the following performance results:


  1 threads       0.648891 GFLOP/s
  2 threads       1.489049 GFLOP/s
  3 threads       2.209838 GFLOP/s
  4 threads       2.940443 GFLOP/s

When I rewrite the generated assembly to exhibit the same interleaving as in the LLVM IR form


  .
  .
  mulss %xmm8, %xmm10
  mulss %xmm7, %xmm9
  mulss %xmm6, %xmm3
  mulss %xmm5, %xmm11

  mulss %xmm8, %xmm10
  mulss %xmm7, %xmm9
  mulss %xmm6, %xmm3
  mulss %xmm5, %xmm11

  mulss %xmm8, %xmm10
  mulss %xmm7, %xmm9
  .
  .

observed performance increases by nearly a factor of four:


  1 threads       2.067118 GFLOP/s
  2 threads       5.569419 GFLOP/s
  3 threads       8.285519 GFLOP/s
  4 threads       10.81742 GFLOP/s

I show similar results with packed single-precision floating point instructions. The LLVM-generated
machine code is nearly 4x slower than the interleaved schedule suggested by the LLVM IR input.


  Vectorized - No instruction interleaving - back-to-back dependencies

  1 threads        1.540621 GFLOP/s
  2 threads        5.900833 GFLOP/s
  3 threads        8.755953 GFLOP/s
  4 threads        11.257122 GFLOP/s

  Vectorized - Interleaved - stride-4 reuse distance

  1 threads       3.157255 GFLOP/s
  2 threads       22.104369 GFLOP/s
  3 threads       32.300111 GFLOP/s
  4 threads       39.112162 GFLOP/s

It is worth noting that 39.1 GFLOP/s is approaching the theoretical limits of the processor
(stated to be 45.25 GFLOP/s single-precision).

Skiplists

2010-09-29T15:33:00.007-04:00

Skiplists are probabilistic data structures with many of the same properties as balanced binary search trees. They are analogous to express and local passenger trains and have O(log N) insert and search complexities.

I implemented a skiplist container class in C++ to have the same semantics as STL's std::map<>. Here are some results for insertion of unique integers in random order:


iteration and search tests pass.

N = 25000
avg insert: 13.1927
log_2(N): 14.6096

N = 25000
avg search: 14.503
log_2(N): 14.6096

Another neat property is there isn't an expensive re-balancing operation as seen in red-black and AVL trees. The insert() and delete() operations simply need to update at most log_2(N) pointers. Haven't looked at actual performance of my implementation but I thought it was an interesting exercise.

Here is an example skiplist after 11 inserts (after the jump).


[header]
 |         |         |        
 -> [0]   -> [1]   -> [2]  

[0] = 0
 |        
 -> [1]  

[1] = 10
 |         |        
 -> [2]   -> [2]  

[2] = 20
 |         |         |        
 -> [3]   -> [6]   -> [6]  

[3] = 30
 |        
 -> [4]  

[4] = 40
 |        
 -> [5]  

[5] = 50
 |        
 -> [6]  

[6] = 60
 |         |         |        
 -> [7]   -> null   -> null  

[7] = 70
 |        
 -> [8]  

[8] = 80
 |        
 -> [9]  

[9] = 90
 |        
 -> [10]  

[10] = 100
 |        
 -> null

Searching for key [9] would consist of visiting the following sequence of nodes: (header, 2, 6, 7, 8, 9). Searching for key [3] would consist of visiting (header, 2, 3).

Ascent of Mount Rainier

2010-08-16T13:39:00.007-04:00

On August 14, I set out with three companion to summit Mount Rainier. At 9:15 am yesterday, August 15, we reached the summit crater.

The climb was well-documented with a digital camera, but here are some snapshots I took with my iPhone.

Mount Rainier seen from Muir Snowfield

Summit crater of Mount Rainier

Me, Ryan, and Tim (left to right) on Columbia Crest looking northwest toward Seattle.

Me on Columbia Crest with Mount Adams in background

Comprehensive post to follow.

Mount Rainier Weather Forecast

2010-08-13T14:00:00.005-04:00

I've been refreshing this page continuously for the past few days. If the forecast is correct, the summit should be warm and relatively calm. I'll have my four thermal layers nevertheless, but I feel the need to post the current forecast (below the break).

Mount Rainier, August 22, 2009


MOUNT RAINIER RECREATIONAL FORECAST
NATIONAL WEATHER SERVICE SEATTLE WA
410 AM PDT FRI AUG 13 2010

.SYNOPSIS...A STRONG UPPER LEVEL HIGH PRESSURE AREA OFFSHORE WILL 
BRING A PERIOD OF SUNNY DAYS...WITH TEMPERATURES WARMING THROUGH THE 
WEEKEND AS THE LOW LEVEL FLOW TURNS OFFSHORE. CLEAR WEATHER WILL
CONTINUE EARLY NEXT WEEK WITH TEMPERATURES LIKELY TO MODERATE
GRADUALLY.

&&

.FRIDAY...SUNNY. FREEZING LEVEL ABOVE 15000 FEET. 
.FRIDAY NIGHT...CLEAR. FREEZING LEVEL ABOVE 15000 FEET. 
.SATURDAY...SUNNY. FREEZING LEVEL ABOVE 15000 FEET. 
.SATURDAY NIGHT...CLEAR. FREEZING LEVEL ABOVE 15000 FEET. 
.SUNDAY...SUNNY. FREEZING LEVEL ABOVE 15000 FEET. 

&&
TEMPERATURE AND WIND FORECASTS FOR SELECTED LOCATIONS.

                       FRI    FRI    SAT    SAT    SUN  
                            NIGHT         NIGHT        

SUMMIT   (14411 FT)     34     33     37     34     38
                      N 25  NE 14   E 10  NE  5   N  5 

CAMP MUIR(10188 FT)     55     43     55     46     55
                      E 10   E 15   E 15   E  5  SE  5 

PARADISE  (5420 FT)     75     58     78     59     80
                      E 10   E 10   E 10   E 10   E 10 

LONGMIRE  (2700 FT)     81     58     84     58     88
                      CALM   CALM   CALM   CALM   CALM

++ TEMPERATURES AND WIND FOR THE SUMMIT AND CAMP MUIR ARE AVERAGE
   CONDITIONS EXPECTED IN THE FREE AIR AT THOSE ELEVATIONS.
++ TEMPERATURES FOR PARADISE AND LONGMIRE ARE THE EXPECTED HIGHS AND
   LOWS. WIND IS THE AVERAGE WIND EXPECTED DURING THAT PERIOD.

.EXTENDED FORECAST...

.SUNDAY NIGHT THROUGH WEDNESDAY NIGHT...CLEAR. FREEZING LEVEL ABOVE
15000 FEET. 
.THURSDAY...PARTLY SUNNY. A CHANCE OF SHOWERS IN THE AFTERNOON. 
FREEZING LEVEL 15000 FEET. 

$$

WEATHER.GOV/SEATTLE

Mount Rainier Climbing Information

2010-08-12T20:52:00.004-04:00

Here are some Mount Rainier climbing resources if you're interested.

Climbing information

Climbing routes (look for Paradise Routes - this is the best topographic map of the route to summit)

Route information

Equipment checklist

Weather forecast

To Climb Mount Rainier

2010-08-12T00:28:00.005-04:00

The stars have aligned such that I might be a part of a three or four person rope team to attempt a summit of Mount Rainier on the weekend of August 13 - August 15. If all goes according to plan, we will reach Mount Rainier's 14,400 ft summit crater by approximately 10am Sunday morning.

A detail plan will come out of a meeting between the climbers tomorrow evening. We will follow the Disappointment Cleaver route. This is a two-day climb from the trailhead at Paradise to high camp at Camp Muir (10,100 ft) on a ridge between Cowlitz and Nisqually Glacier on Rainier's southern face. Early the next morning, we'll rope up and cross Cowlitz glacier and Ingraham flats toward Disappointment Cleaver, a rock formation dividing Emmons and Ingraham Glaciers. Once we reach the top of this formation, we'll proceed over snow up toward the crater rim, and from there to Columbia Crest (summit ridge) if weather permits.

I hope it permits.

Update:

Follow Mount Rainier's weather forecast http://www.atmos.washington.edu/data/rainier_report.html.


.THURSDAY...SUNNY. FREEZING LEVEL 14000 FEET. 
.THURSDAY NIGHT...MOSTLY CLEAR. FREEZING LEVEL 15000 FEET.
.FRIDAY...SUNNY. FREEZING LEVEL ABOVE 15000 FEET. 
.FRIDAY NIGHT...CLEAR. FREEZING LEVEL ABOVE 15000 FEET.
.SATURDAY...SUNNY. FREEZING LEVEL ABOVE 15000 FEET.

.SATURDAY NIGHT THROUGH WEDNESDAY...CLEAR NIGHTS AND SUNNY DAYS. FREEZING
LEVEL ABOVE 15000 FEET.

That appears to be ideal climbing weather.

Glacier Training: Prusik

2010-07-28T17:06:00.009-04:00

The essence of human-powered ascent, by any means, is a lift followed by a rest. When climbing stairs, you lift yourself then stand while hoisting your free leg to the next tread. When climbing a rope, it's quite beneficial to have similar options. It is ideal if your legs can bear most of the lifting load, yet at any point relax while muscles regain strength.

Enter the prusik hitch.

Three-wrap prusik hitch:

This is a simple knot that may be tied from a loop of line that grips a climbing rope with great force when loaded, but when unloaded may be moved upwards or downwards. A pair of these hitches may be used to ascend a climbing rope as follows. A short loop makes the "waist prusik" and includes a three-wrap prusik hitch around the climbing rope as illustrated above and links to the climber's harness via a locking carabiner. A longer length of rope makes up the leg prusik; stirrups at the ends tie around boots.

To ascend, start from sitting, suspended by the waist prusik with legs in stirrups of unloaded leg prusik. Employ the following procedure.

Stretch out legs and attempt to stand in leg prusik, using arms to maintain vertical posture

Use free hand to move unloaded waist prusik hitch as high up on the rope as possible

Sit in (now raised) waist prusik which should immediately grip climbing rope. You are now ~8 inches higher than you were before.

Fold legs inward, unloading leg prusik

Use free hand to raise leg prusik upwards, just below waist prusik

Repeat until at top of climbing rope

To descend, the same process is used but in reverse, but you must govern yourself and not be too greedy. Only attempt ~8 inches of down-climb per iteration.

While sitting, lower leg prusik 8" from below waist prusik hitch

Stand in stirrups, unloading waist prusik

Lower waist prusik to top of leg prusik hitch

Sit in (now lowered) waist prusik to unload leg prusik hitch

Repeat until on ground or bottom of climbing rope

After every eight to ten feet of elevation change, it would be nice to save progress so, in the event of slippage or other accident, you do not wind up at the bottom again. To accomplish this, make a "bight" in the climbing rope just below the prusik hitches and tie a figure eight knot. Clip this into your climbing harness with a free carabiner [NEVER open the gate of a carabiner that is providing protection to you or your gear]. Once this is done, a fall will at most set you back to the highest safety knot.

If carrying a heavy backpack, it's very difficult to execute the above without quickly fatiguing. The pack moves your center of gravity outward putting more load on your arms. The only viable solution we were able to work out was to attach a third prusik hitch between the bag and the climbing rope below the leg prusik hitch.

The sequence is modified by performing the following during every iteration:

While seated, after raising leg prusik, lift bag hanging below.

Use free hand to raise unloaded bag prusik hitch to below leg prusik.

Resume basic prusiking process, but execute this procedure after sitting every iteration.

Other options beyond a bag prusik were attempted, but none of them yielded satisfactory results when clipping into knots in the climbing rope. One method was to carabiner the bag to a loop that forms in the climbing rope, but this adds weight and pulls downward on your harness reducing the distance change per iteration. Another method ties the bag to the end of the climbing rope, but the first safety knot you make will transfer the bag's load to you. Neither of these were particularly satisfactory or as elegant as the third prusik.

I took this camping:

Coming soon: glacier traversal

Glacier Training on Mount Baker

2010-07-26T13:37:00.011-04:00

The climb of Mount Adams two weeks ago was exhilarating and exhausting, but an important element of mountaineering was [thankfully] lacking: ice. Glaciers present a new set of hazards and require specialized techniques to overcome them. It behooves every mountaineer to obtain training and plenty of practice such that, in an emergency, the mountaineer can immediately swing into action and apply rescue procedures quickly and correctly. We react how we train.

I was fortunate enough to spend this weekend camping with Brooke and Ryan next to a glacier on Mount Baker. Ryan had taken a glacier class from a mountaineering guide service the previous weekend and explained to us as much of the material as we could get through. On the glacier, we put these concepts into action. Here is a list of the glacier traversal and rescue procedures we covered.

prusiking

glacier traversal

crevasse identification

anchoring in snow

anchoring in ice

roping up

self-arrest

crevasse rescue

ice climbing

The two-day trip consisted of the following events.

On Day 1, we: drove to Mount Baker-Snoqualmie National Forest, we hiked 2.5 miles from the trailhead to the camp site crossing numerous streams carrying glacier melt water down the volcano, setup camp, discussed rope climbing fundamentals, prusiked up a tree, had camp supper, enjoyed a view of the heavily glaciated face of Mount Baker, enjoyed discussion with some other climbers camping there, and I slept in a bivouac sack+thermarest+sleeping bag.

On Day 2, we: carried gear onto the glacier, practiced traversal techniques, climbed over some ice mounds, made all sorts of anchors, setup ropes, practiced rescue techniques, practiced self-arrests, ice climbed, broke camp, hiked out, and enjoyed a nice supper at a brewery in Bellingham.

What will eventually follow is a brief description of each of the hazards and techniques we trained for. The entire two-day trip was well documented, for the sake of both posterity and so I can review material. It will take several days to coalesce photos from two cameras and write things down, but I'll try to complete at least one section before the end of the week.

Coming soon: Prusiking

Cimbing Mount Adams

2010-07-20T03:07:00.001-04:00

Mount Adams.

Climbing Mount Adams

2010-07-14T13:37:00.016-04:00

Mount Adams Climb Report. What follows is a somewhat meandering description of my climb of Mount Adams. I wrote most of this several days after the climb but didn't wrap it up until just now. One could conceivably annotate each paragraph with photos from the next blog post.

Mount Adams is a prominent 12,280 ft stratovolcano in southern Washington state approximately 25 miles north of the Hood River and 20 miles east of Mount St. Helens. In early July, its beautiful slopes are snow-covered and offers an 'easy' climbing route along its southern face that is free from significant mountaineering hazards - crevasses to fall into, cliffs from which rocks fall, and avalanche-prone slopes - making it a popular climb for residents of Portland and Seattle. On July 11, 2010, I accompanied my roommate Tim and his friend Rumi to the summit of Mount Adams in a single day's climb.

Getting to Mount Adams wilderness requires approximately four hours of driving time from Bellevue, so we left at 11pm the night before and drove all night. Delayed slightly in acquiring our third climber, we didn't really depart Federal Way until close to midnight. During the drive, I managed to get perhaps 45 minutes of sleep but was otherwise conscious for the journey. By the time morning twilight approached, we were still beyond the trailhead wondering through backroads north of the Hood River and kicking ourselves for getting such a late start. Other climbers would be just waking up now, having camped at the trailhead or - better - at 9,000 ft and would easily summit during mid-morning when the snow is still frozen solid from the previous night's cold. Every second that passed as we blazed over the narrow dirt roads made the summit seem more and more distant.

By about 6am, we reached the trailhead. Mt. Rainier had left me with false expectations, as Paradise, the trailhead for the skyline trail below the Muir Snowfields, is the location of an enormous resort hotel. Cold Springs on Mount Adams, in contrast, was a sequence of camp sites at the end of a labyrinth of dirt roads. On this day, it was full of vacant tents and quite calm. We recorded our location via a Garmin hand-held GPS and noted our starting elevation of approximately 5,500 ft above sea level.

We donned our boots, grabbed our packs, and started up a dirt road consisting of four or five switchbacks. Starting time was 6:30am. In approximately 0.75 miles, we reached the first patches of snow. The sun was up, the sky bright blue, and filtering through the evergreen trees. This being the first time I wore my mountaineering boots outside of my bedroom and office, I noted they were remarkably comfortable. They demonstrated good traction on the snow, and nylon gaiters around my ankles prevented snow from entering above the laces.

Thirty minutes in, we stopped to reconfigure our packs and apply sun screen. In spite of the snow on the ground, temperatures were warm, and our bodies' heat output was high. I wore a Patagonia capalene long-sleeved shirt whose loose weave enabled moisture to escape but kept me at a good temperature. Unfortunately, we knew that this day would be a warm one for southern Washington, and that would mean the snowfields we would be traversing would be soft and weak by the time we reached higher elevations. At best, this would mean more energy spent with each step. Worst case could mean avalanche danger in which entire slabs of snow pull apart and slide down the mountain. Things were calm and serene at this point, so we didn't dwell on that.

We progressed quickly, and in spite of being the least experienced climber of this group, I kept up a good pace. In no time, our elevation gains became apparent. Mount Hood, directly to our south, loomed above Oregon. Occasionally, we'd obtain glimpses of the snow-covered Mount St. Helens, though her crater was not visible. Looking north, the 11,500 ft false summit of Mount Adams appeared as merely a point off in the distance. Insects buzzed around in spite of the snow, and the density of timber began to thin. What had once been small ridges lying distantly in our path had grown to be towering massifs. At some point, the hike had become a climb.

Ascending a snow-covered slope is an iterative process in which the toe of one's boot is kicked into the snow with sufficient force to crush down the fluffy snow into a hard, foot-shaped step. The lead climber "breaks a trail" leaving a foot path the others follow. Kicking is hard work, but at this point no one was tired enough to feel a real need to take advantage of existing paths. That kind of bravado probably isn't a good idea, and we should have walked single file over the existing trail from earlier climbers. Below 8,000 ft, we maintained a fast aerobic pace and took breaks only for water, my 96 ounces of which were beginning to weigh upon me.

The suns rays began making an impact. Snow is an excellent reflector, and we made frequent stops to apply sunscreen to our exposed parts. We got hotter, and there was frequent talk of being "cooked like a lobster." I donned glacier goggles, and hats came out.

Just after noon, we reached the "Lunch Counter," a comparatively flat area at approximately 9,000 ft with exposed rocks. Sensible climbers attempting to summit Mount Adams in two-day climbs typically set up tents here using natural landforms to reduce exposure to winds. Brooke and Ryan camped here the weekend before my climb, for example. Climbers attempting to summit in a single day eat lunch here, as it is the last reprieve before a steep 2,000 ft ascent up to the false summit. We sat down on some warm volcanic rocks and ate a much-needed lunch. I ate a banana, peanut butter and jelly sandwich, drank some Gatorade, and consumed a Gu packet. The sweets were beginning to make me nauseated, but I didn't bring other options. We looked out past the beautiful river valley to Mount Hood and admired the view. Sparse clouds slowly rolled in from the east but never engulfed the mountain. Other climbers glissaded down from the summit, having sensibly gotten a far earlier start than we did.

Setting off from the lunch counter, the climb quickly grew more challenging. We were now on an exposed face with a strong wind blowing down the mountain. The slope quickly steepened to the maximum of 30 degrees, and our ever increasing altitude made aerobic respiration difficult. Now afternoon, the overhead sun had warmed the snowfield such that each step sloshed through nearly a foot of snow, and the ice axe in my hand penetrated the entire length of its shaft as I dug in for additional support. Suspecting that at any moment my feet might slip and I would fall, I prepared to dig in with the ice axe for a self arrest with every step.

I had trouble obtaining a satisfying pace and frequently climbed too fast then ran out of oxygen and had to pause to catch my breath. This quickly wore me out and didn't make this part of the climb too pleasant. A steady stream of climbers glissaded by, as our trail took us up between well-established glissade paths. Skiiers and snowboarders also whizzed by every so often, so we shouted out to them. I can't imagine a better way to give back nearly 7,000 ft of hard-won elevation than by skiing down, and these slopes were pretty ideal for that.

After ages went by, Tim and I reached the false summit. The altitude was beginning to get to me, and I had to breath hard for nearly ten minutes to finally catch my breath. My throat was parched, but the Gatorade I had on hand left me nauseated. I highly recommend bringing water over Gatorade. We sat just below the peak of the false summit and admired the view from 11,500 ft. Just above our heads, a fast wind whipped over the peak and motivated better insulation. I added long underwear under my cargo pants, donned a $10.00 fleece jacket purchased from a Pier 39 giftshop in San Francisco, and put on a borrowed nylon shell jacket. This kept the wind out, me insulated, and ventilated to avoid perspiration building up inside.

Peaking our heads over the false summit, an intense and continuous headwind blasted us. We had over 800 ft more to ascend and a daunting slope with plenty of wind exposure. We moved swiftly along a shallow basin toward the next slope and looked upwards as tiny and distant climbers provided us a sense of scale. Tim, suffering slightly from the altitude and from the intentionally heavy pack he was carrying to train with, suggested I take the lead.

With no one in front of me, I settled into a good pace that made what might otherwise have been a dreadful climb at the end of a long haul much more pleasant. I employed the "rest step" in which, after each step, a complete breath was inhaled and exhaled. If I put my foot down more than eight or nine inches above the previous step, I took a second breath. It took a deliberate effort to keep things slow, but consequently I could sustain over 200 steps between short rest breaks and I never reached a point in which anaerobic respiration got the better of me. If any part of the slope above looked like it might give way, I would have enough energy to attempt try to outrun it.

In what felt like no time at all, I had climbed the last major stretch and stopped on a short crest awaiting my two other companions to reach me. I used the time to take panorama photographs with my iPhone whose LCD was completely invisible in the burning sunlight. Three or four individuals came down from the now-visible summit a mere hundred yards or so beyond me. I felt like racing up to it but instead elected to wait for my friends to arrive.

They did, and we marched up to a small mound of snow covering the remains of an old mining shack. Volcanos are good sources of sulfur, evidently. This was the true summit, 12,280 ft above sea level. Sparse cloud cover at 8,000 ft added to the incredible view. Taking our first glimpses to the north, Mount Rainier loomed over the horizon, its three peaks easily resolvable. Mount Hood had grown less captivating with each exhausting push up the route, but suddenly it was much more charming, particularly since we were looking "down" on its 11,500 ft peak. Mount St. Helens was consumed by its cloud cover and barely visible. Several features on the north side of the mountain came into view. I never could make out the rounded concavity of a volcanic crater, but I'm told one exists and would probably be visible at the time of this writing due to melting.

After plenty of summit photos, we took shelter behind the true summit. Winds had shifted, and it no longer concealed us from the blast, so we darted off back down the trail as quickly as possible. With five hours of daylight remaining, we weren't technically rushed, but "summit fever" had worn off.

To enable a productive glissade, I duct-taped a plastic bag around the seat of my pants hoping it would reduce friction and keep me dry. Predictably, it didn't do either of those things and ripped off immediately. Nevertheless, on each of the steep slopes, I made an attempt to glissade down the well-established paths in the snow. And each time, I barely got moving at all. My pants are not slippery, and treking poles attached to my pack dug in. These factors combined with wet snow made the experience somewhat demeaning.

In approximately three hours from the summit, we made it back to the trail and not long after the car. Energized by completion of an intense and impressive climb, we chattered on in the car for twenty minutes as we drove back to civilization. This lasted for twenty minutes and then I hit a wall. Over 42 hours of being awake hit me hard, and I became irritable and passed in and out of consciousness. The microsleeping lasted the entire way back; I tried to avoid actually going to sleep, as the driver was in the same sleep-exhausted condition as I was and, toward the end, was bouncing between rumble strips like photons in a fiber.

Survived that, made it home at 3am. Arrived to work at 10:00am. Best climb to date!

Blog Slacking

2010-07-13T14:01:00.010-04:00

Since May, I've driven across the country, lived in Seattle for a week, visited San Francisco for a week, attended Emma's graduation ceremony, moved to Bellevue, written a whole lot of code, climbed several hills, biked all around, gone sailing, visited Illinois, and climbed Mount Adams. Shockingly, the only thing I've blogged during all of this is a simple quine I wrote in CUDA, and even then it was only to document having done it first.

I'm a slacker.

In related news, I've implemented a photo gallery for FanBolt that I also plan on installing on akerr.net. It should satisfy the purposes of this blog quite well. As there is a value cliff coinciding with ComiCon for getting the photo gallery into production, it better be working by the end of this weekend. At that point in time, I'll upload photos from each of the above adventures to get some kind of documentation to you [few, eternally optimistic] readers!

Blogger iframe test:

Your browser does not support iframes.

First CUDA Quine

2010-06-29T19:29:00.007-04:00

As you may know, a quine is a short program that prints its own source code using common facilities of the language. With CUDA 3.0 and Fermi's support of printf(), this isn't very difficult to do with CUDA. Here is an example, the first I've seen to run on a GPU.


$ cat quine.cu 
#include <stdio.h>
__device__ const char *f="#include <stdio.h>%c__device__ const char *f=%c%s%c; __global__ void k(){ printf(f,10,34,f,34,10); } int main() { k<<< dim3(1,1),dim3(1,1)>>>(); cudaThreadExit(); return 0; }%c"; __global__ void k(){ printf(f,10,34,f,34,10); } int main() { k<<< dim3(1,1),dim3(1,1)>>>(); cudaThreadExit(); return 0; }

$ nvcc --run quine.cu -arch sm_20 > quine2.cu
$ diff quine.cu quine2.cu
$ nvcc --run quine2.cu -arch sm_20 > quine3.cu
$ diff quine.cu quine3.cu
$

Ocelot Overview

2010-03-30T15:14:00.005-04:00

With 2010 being the year of the Ocelot, I've made a slight modification to the Ocelot Overview image that has found its way into so many presentations [not all of which have been mine]. See if you can spot it.

Baker Valve Gear Simulator in Python

2010-03-06T19:03:00.021-05:00

What started off as a compile-time project has finally yielded some results.

As some of you are aware, my interests in engineering exceed computing and from time to time take the front seat for brief periods. In this case, I've spent two weekends volunteering at the Southeastern Railway Museum shops with the hope of working on the preservation and restoration of the exquisite A&WP 290. The neat thing about steam locomotives is they're incredibly overt exhibitions of thermodynamics and mechanics. The boiler is a giant heat exchanger that sends steam to cylinders where work is performed. A mechanical linkage outside the frame transmits this power to the wheels and uses the wheel position to configure valves that admit and exhaust the working fluid.

One of the things I was interested in was a mathematical model of this linkage. This sort of modelling has applications in robotics and kinematics simulations, so it's not too bizarre of a topic for a computer engineer to dabble in. To simulate the valve gear, I constructed a set of polynomial equations modelling each rod as a constant length between two points capable of 2D motion in a plane. The collection of pin joints and rigid rods determines a finite set of possible static equilibria for the system. Enough points can be computed analytically so a numerical non-linear equation solver can compute the others.

I wrote a Python application [distributed as source] using SciPy's fsolve() to simulate the running gear. Plots generated with matplotlib provide visual output.

Here are some of the results. This wikipedia article explains the Baker valve gear and this animated gif shows it in operation [generated from an existing valve gear simulator which I've not yet used to validate mine]. Other valve gears are explained on this page if you're interested.

The following plot illustrates the state of the running gear for several wheel rotations. Not every part is a two-joint rod; there is a bellcrank and an eccentric crank mounted to the main crank pin. See the Wikipedia article for a better illustration.

The following plot illustrates the offsets of the crosshead and cylinder valves for two different cutoffs and two different laps. Cutoff is the angle between vertical and the valve gear's yoke. 0 degrees is effectively neutral, and +25 degrees is "well forward." Setting it to a negative angle reverses the locomotive. Lap is, in a manner of speaking, the amount of motion derived from the crosshead; the valves can move this much without actually admitting any steam. It's worth noting that lap is built into the valve gear, while cutoff is a setting that is adjusted during operation.

To achieve smooth operation, it's desirable to admit steam into the cylinder before the piston has reached front or back dead center. This "lead" is effectively a phase offset that depends on the direction of motion. The amount of lap affects the amount of lead by mixing the 90-degree offset motion of the valve gear with the crosshead via the combination lever. The following plot illustrates the displacement of the cylinder valves as zero-mean curves to enable a direct comparison of 0" and 3.5" of lap. In practice, it's easy to add a constant linear offset to the cylinder valve spindle, and this is one of the variables that must be tuned when the locomotive is overhauled.

The rod sizes and geometry used in this simulation were estimated from photographs of A&WP 290 as well as direct measurements. I didn't get every measurement I wanted, as I was taking them by hand with a tape measure that falls apart, and real locomotives are 3 dimensional and quite large. Nevertheless, this was an interesting project and I'm very gratified by the results. Maybe one day I'll build a 7.5" gauge live steamer and dust off this code for design purposes...

NP

2010-02-17T13:02:00.003-05:00

I thought of this XKCD comic last night as I was showing problems were in NP. Resurrected for your amusement.

Read up on Karp's NP-Complete Problems.

Upwelling Within

2010-02-15T16:17:00.003-05:00

I have the urge to build something in ways that do not involve typing. Something mechanical with moving parts.

In potentially related news, this looked neat: Crower six-stroke diesel.

Motherboard Death Throws

2010-02-14T19:22:00.011-05:00

My desktop is back. Evidently, one [of three] PCI-Express x16 lanes is dead as is one of the memory channels. The system refuses to POST with something occupying any of the damaged connectors. This leaves two [though one inaccessible] PCI-Express x16 slots and one memory channel still viable. I can make do with this configuration for the time being.

The experience has prompted the question of what I want in a home desktop. While there was a time when GPU VSIPL development and benchmarking demanded a high-end system, now I do most work on other resources purchased by the research budget. I don't game, and I do most home development on my laptop which is typically booted into Linux. I wouldn't have a need for Windows at all except to synchronize my iPhone, though it's nice to have that kind of compatibility for special interests and the occasional foray into 3D graphics programming. That said, a low-end multicore CPU and medium-range GPU could satisfy those needs quite well.

With a medium-performance laptop dual booting Windows and Linux, do I need the desktop at all?

The strongest reason I can see for wanting one is to have a stable place for an array of harddrives to ensure reliable access to a decade or so of accumulated data. That's a task easily serviced by network-attached storage [1].

This has disturbing implications for the value of my research interests at the consumer level...

[1] BUT NOT IF IT'S MADE BY LACIE - you're better off keeping your data on floppy disks under your mattress than shoveling it into that black hole of a storage medium; it's a place where entropy goes to die; fail fail fail.

GPU Ocelot

2010-02-10T16:37:00.003-05:00

What follows are a set of ideas I've been babbling to people over the past few days:

A recent call for papers to contribute to NVIDIA's GPU Computing Gems has prompted the brainstorming of potential uses for GPU Ocelot. It occurs to me that I've not done a post that addresses an important segment of the GPU Ocelot user space: CUDA developers.

While Ocelot makes a good infrastructure for researching GPU architecture, programming models, compiler transformations, and operating systems, it's also incredibly useful for CUDA application development and performance tuning. Ocelot consists of an implementation of the CUDA Runtime API with three backends for executing CUDA kernels on: a full-featured PTX emulator with instrumentation, an efficient translator to multicore CPU, and a CUDA-capable GPU. I'll be focusing on the capabilities enabled by Ocelot's PTX emulator for most of this post.

PTX is the virtual instruction set that the CUDA compiler targets. Ocelot includes an instruction-by-instruction emulation of PTX kernels in a manner that is faithful to the PTX execution model and very similar to how GPUs might execute the kernel. Since Ocelot explicitly computes the state of a PTX kernel for every instruction and every thread, detailed instrumentation and analysis is possible. Here are some of the features that greatly assist in developing CUDA applications and tuning them for efficient execution on GPUs.

Ocelot's PTX emulator is capable of checking for out-of-bounds memory accesses. Without Ocelot, your application might simply seg-fault on the GPU and possibly crash your video driver. There aren't really any tools currently available to help you find these problems, and padding added to memory allocations by the CUDA driver may conceal problems on the system you test with only to explode on your users whose systems may be different.

Beyond global memory accesses, Ocelot's emulator detects race conditions in shared memory that necessitate use of synchronization as well as deadlocks in which not all threads reach a synchronization point. To achieve maximum performance from the memory system, Ocelot includes a way of determining whether the application is taking advantage of spatial locality of data whose accesses can be coalesced in time. At present, we do not detect bank conflicts in which multiple concurrent threads attempt to access the same port of shared memory; on a GPU, these must be serialized greatly impacting performance. I plan to implement that soon.

Finally, Ocelot supports user-contributed trace generators. You could write, for example, a trace generator that watches a particular location in memory and identifies when that location was updated and by which thread(s). I believe I will write this and then, in a blog post, demonstrate how to extend it.

NVIDIA's CUDA does include an emulation target, but this is implemented completely differently using one host thread per logical CUDA thread. This MIMD approach to kernel execution is terribly inefficient, as most of the application runtime is spent by the OS switching contexts between threads. Moreover, it constitutes a radical departure from the way kernels are executed by GPUs to the extent that CUDA kernels that are invalid for subtle reasons may execute correctly on their emulator but incorrectly on a GPU. Ocelot's PTX emulator presents a more similar execution environment that captures the subtleties of the GPU execution model.

GPU Ocelot is fairly stable as of this writing. It has been validated with over a hundred applications from the CUDA SDK, third party benchmark suites, CUDA-based libraries such as GPU VSIPL and Thrust, and custom CUDA applications. Ocelot has been used to capture performance characteristics of these applications as well as find nuanced bugs in existing CUDA programs. It constitutes a great improvement over NVIDIA's emulation mode CUDA runtime in terms of both accuracy, observability of potential bugs, and emulation performance.

JSON

2010-02-02T11:08:00.004-05:00

Last October, I wrote a JSON parser in C++ that constructs a document object structure. A visitor class provides convenient access to the document object, and a number of tests validate the parser. I've also relaxed the specification to permit strings without quotation marks assuming they meet the same rules as variables in a C-like programming language. Multi-byte encodings aren't as well-tested or covered.

I'm preparing to make it available under a BSD license. Here's a sample usage scenario from the test suite.


//
// Example JSON usage
//
try {
  std::stringstream ss;
  json::Parser parser;

  // without quotes, and with weird spacing
  ss << "{objects:[{id:1, location : [ 0, 0,0]},{id:2,location:[1.25,0.25,4]}]}";
  
  json::Object *object = parser.parse_object(ss);
  json::Visitor visitor(object);
  
  json::Visitor obj0 = visitor["objects"][0];
  json::Visitor obj1 = visitor["objects"][1];
  
  if ((int)(obj0["id"]) != 1 || (double)(obj0["location"][0]) != 0) {
    if (verbose) {
      std::cout << "visitor test 0 - failed for obj0\n";
    }
    passes = false;
  }
  
  if ((int)(obj1["id"]) != 2 || (double)(obj1["location"][0]) != 1.25 ||
    (double)(obj1["location"][1]) != 0.25 || (double)(obj1["location"][2]) != 4) {
    if (verbose) {
      std::cout << "visitor test 0 - failed for obj1\n";
    }
    passes = false;     
  }
  
  delete object;
}
catch (common::exception &exp) {
  passes = false;
  std::cout << "parse failed for visitor test 0\n";
  std::cout << exp.message << std::endl;
}

Turkey Albondigas

2010-01-30T18:51:00.000-05:00

Emma and I are making this tonight: Turkey albondigas soup. It is spectacular.

Posted here for posterity.

Spring 2010

2010-01-17T15:08:00.009-05:00

I'm enrolled in two Computer Science courses: CS 6505 for an unneeded grade, and CS 6520 for fun.

In CS6505 on Friday, I was asked to show the rationals are countable. Having seen this problem before, I quickly mapped the rationals onto discrete points in a 2D plane and, while waving hands, showed how they could be counted by spiralling outward from the origin. "There are different classes of infinity," I struggled to explain. The three gentlemen I was sitting with were not necessarily convinced I had done any math or that my explanation was adequate, so I cast about for a better method.

Here it is. Recall that N is the set of natural numbers: {1, 2, 3, ..}, integers are the set Z = {.., -2, -1, 0, 1, 2, ..}, and the rational numbers are numbers that are the ratio of an integer to a natural number. A set is countable iff there exists a mapping to the natural numbers.

Finite sets are of course countable, but showing that infinite sets are countable requires a bit more thought. The integers, for example, may be mapped to natural numbers as follows: let "0" be integer number 1, +1 be integer number 2, -1 be integer number 3, +2 be integer number 4, -2 be integer number 5, and so on. Easy.

Now for the rationals. Consider an integer a represented by the bits (a0 a1 a2 .. an), where a0 is 0 if a is positive and 1 if a is negative, while bits (a1 a2 .. an) denote the absolute value of a. Assume b is an n-bit natural number formed by bits (b1 b2 .. bn). The pair (a, b) represents some rational number. We may now form a string from the bit representation of (a, b) as follows:

(1 a0 a1 b1 a2 b2 a3 b3 .. an bn)

This binary string maps to a unique natural number. Moreover, we see that the set ZxN is a superset of the rationals (-10/5 and -2/1 are the same rational number but distinct in the superset ZxN). Because each element of this superset maps to a unique natural number, it is countable, and therefore the rationals are countable. qed.

My personal take-aways from this experience: (1) writing things down like this rather than drawing visual aids and pointing seems to make people more comfortable, and (2) trying to show something in two slightly different ways is a good exercise.

Southeastern Railway Museum

2009-12-28T15:13:00.016-05:00

The day after Christmas, my parents, Emma, and I spent the afternoon at the Southeastern Railway Museum. Their large and growing collection includes a number of cosmetically restored rolling stock from railroads in the Southeast including several large steam locomotives. This post covers some of the photographs from that visit.

Savannah and Atlant #750 - This 4-6-2 "Light Pacific" built in 1910 pulled excursions all around the Southeastern U.S. including one in 1989 that I rode on. It was last operated that year and currently sits on static display, having been cosmetically restored.

Click on photos to expand them to their full version.

My darling Emma at the throttle. Her right hand is on the power reverse lever, and her legs are crossed because she is a lady.

Me in the cab. The smaller numbers [expand the photo] refer to dimensions of the running gear: 69" driver wheels, 22" diameter cylinders, 26" stroke.

This is a direct shot of the Walschaerts valve gear. A valve gear provides a linear reciprocating motion that is either 90 or 270 degrees out of phase with the main crank pin plus a small amount of lead thereby determining which side of the piston steam is admitted and for how long during the stroke. To reverse the steam locomotive, the radius rod pivot is moved to the top of the expansion link.

Attention all Boiler Makers!

S&A #750 - This photograph of #750's firebox illustrates why it is very unlikely she will ever run under steam again. The crownsheet was worn considerably and panels were cut out along the top sides of the firebox. Needless to say, a lot would have to be repaired here for #750's boiler to be certified.

Also visible in several of the large flues are the ends of superheater steam pipes. To improve the efficiency of the boiler, steam is piped down and back through the large fire tubes before it reaches the throttle valve. This increases the stagnation temperature of the working fluid yielding what is known as "dry steam" and increases the amount of work it can do in the cylinders. In effect, this applies the concept of a water tube boiler to the fire tube design used in nearly all steam locomotives in the U.S..

A&WP 290

Later in the day and being the trouble maker that I am, I gave myself a self-guided tour of the SERM's machine shop which houses an assortment of locomotives and rolling stop in various stages of disassembly.

A&WP #290 - Here, we see Atlanta and West Point #290, the [formerly] beautiful 4-6-2 "heavy Pacific" that ran with #750 in the New Georgia Railroad until steam excursions stopped entirely. #290 was featured in the film Fried Green Tomatoes [killing Buddy and severing Buddy Jr.'s arm] and has been covered in several Pentrex videos: link [YouTube].

I didn't have my tripod in the shop, and exposure times were typically between 2 and 3 seconds, so the photos are slightly blurry. Last time I was here [May 2005], #290 was in the midst of restoration to operational status. The jacketing that insulates the boiler has been removed, as seen here, and preperations were underway to lift the boiler and firebox off the chassis to have it shipped to Alabama.

There, it was planned that her boiler be thoroughly inspected and analyzed for repair and recertification, something every steam boiler must undergo every ten years. Unfortunately, funding never crystalized, and Norfolk Southern has since refused to permit steam operations on its mainlines. Lacking cash and a place to operate her, SERM has since stopped hard work on #290 though they hope to have her on static display at some point. Talking with one of the shop foremen, it seems her boiler "should" be viable and her running gear is in excellent shape. "If 4501 can cross the country, why not #290?" Why not indeed!

This is what a "good" firebox looks like.

Built in 1926, A&WP #290's valve gear is of the more modern Baker type which was adopted by all USRA "Heavy Pacifics." I find it exotic and interesting. Apparently this design has joints that are easier to lubricate and less susceptible to wear.

More photos of #290 in operation may be found on the Internets.

Georgia Power #97

SERM is more concerned with restoring their 0-6-0 Georgia Power #97. This is its boiler, primered and ready to be shipped off for analysis. Its smaller, cheaper to operate, and they can run it within the grounds of the museum to give rides and do work making it much more practical.

This is its chassis.

The End

Feel free to visit the disorganized dump of photos taken that day. Most were recorded as RAW that I hastily and automatically converted and downsampled to JPEG.

Selection sort, 64-bit

2009-12-23T00:21:00.004-05:00

Having haphazardly embarked on Agner Fog's informative C++ optimization guide [pdf], I decided it was time to write some x86-64 assembly to gain familiarity with the x86-64 instruction set.

Selection sort is an easy hello-world type application for assembly programming. On a Windows platform, getting Visual Studio and NASM to build a project together demanded a moment's tinkering, so I didn't want to be too ambitious.

The code listing follows. The best part is there is no stack frame to prepare or registers to spill. This doesn't do anything sophisticated like prefetch data, and of course isn't the best algorithm for sorting. It is pretty easy to spot the nested loops.

Posted for your amusement.


section .text code align=16

global selsort

; void selsort(int *A, int N)
;
; A: rcx
; N: rdx
;
selsort:
 xor rax, rax
 
L1:
 mov r8, rax
 inc r8
 mov r10, rax
 
L2:
 cmp r8, rdx
 jz L3
 
 mov r11d, dword[rcx+r8*4]
 mov r9d, dword[rcx+r10*4]
 
 cmp r9d, r11d
 cmova r10, r8

 inc r8
 jmp L2
 
L3:
 mov r11d, dword[rcx+rax*4]
 mov r9d, dword[rcx+r10*4]
 
 cmp r9d, r11d
 jge L4
 
 mov dword[rcx+r10*4], r11d
 mov dword[rcx+rax*4], r9d
 
L4:
 inc rax
 cmp rax, rdx
 jnz L1
 
 ret

netherp - share your naughty bits!

2009-12-14T23:07:00.005-05:00

At Josh's request, I threw the source code to netherp v2 up on a public Github repository. Here is the URL:

http://github.com/kerrmudgeon/kerrutils/blob/master/netherp

netherp was originally intended to be a minimal-configuration HTTP server for sharing files among two machines of arbitrary make and model. One day a few weeks ago I became enthusiastic about asynchronous I/O and non-threaded server architectures, so I wrote a second version from scratch. It doesn't support HTTP PUT or POST or any method other than GET. There is no authentication, and I suspect there may be ways of compromising it.

Nevertheless, there it is. It's single-threaded, platform-independent [within reason], and relatively fast. On my quad-core 64-bit Linux platform in the lab, I ran http_load requesting a variety of files from my local GPU Ocelot repository ranging in size from a 1kB to 100s of kBs. It handled http_load's maximum of 1,000 requests per second without trouble; in this case, the network was the bottleneck. It has also served a 400 MB file over a residential cable connection, this time running on a Windows machine.

So, it's vetted. It works.

There are several barbaric components that ought to be hastily buried underneath a pile of complexity theory [i.e. the MIME type selector], but this does demonstrate working use of select() which lends some much-needed pedagogical value to this exercise.

Additional comments are certainly welcome.

One other facet of interest is I wanted this to be distributable as source, so I wrote a program to embed binary image files used as logos and icons into the source code of the application. It expresses a binary file as a C++ style array declaration. i.e.:


// 'images/file.png'
const unsigned int image_file[] = {
 0x474e5089, 0xa1a0a0d, 0xd000000, 0x52444849, 
 0x30000000, 0x28000000, 0x608, 0x7987b800, 
 ...
};

May not be the best way, but at least it's out the door.

Frigid Digit

2009-11-16T10:59:00.008-05:00

Georgia Tech Sailing Club hosted its intra-club regatta on Saturday. Four 420s sailed out into Lake Lanier and completed a few races. It was very informal, and the light winds made things maddeningly slow at times, although things were never so calm as to prevent deliberate maneuvering.

I didn't win any of the races, but my crew and I came in second place a few times [which doesn't technically matter - sailing rewards winners only]. There are subtle tactical and strategic differences that distinguish experts that I began to pick up on. A strong start is key to dominating a pack of boats - you receive "puffs" of stronger winds earlier and can speed up to stronger positions. Additionally, good positioning and tacking strategies tend to grant a smart skipper right of way during those moments when two boats come upon each other.

There were some tactical matters as well. One boat just seemed to "go faster" in identical wind conditions. Since this regatta consisted of boats of the same class and make, presumably the fastest were those trimmed the best. After we stopped racing, I requested the experienced crews to sail a straight course up wind so that I could follow along and trim sails and crew locations to improve my speed relative to them. That helped considerably to identify the optimal sail positions.

Ralph's Picassa gallery of the day.

Here I am sailing downwind. This is the boring part of the race.

Image

The Nacra catameran that capsized after flying one of its hulls:

Image

Update

While getting my haircut today, I thought about what it might take to determine the optimal strategy for a particular boat under a given set of wind conditions in a regatta. Simplifying assumptions would be:

uniform wind direction and speed across the entire area

no other boats to create local calm areas with bad air or force suboptimal maneuvering

uniform sea state

ideal sail and boat trim

This hypothetical systematic approach would attempt to minimize total time around a pair of marks. Specifically, this would determine:

what course to make that minimizes time to reach the upwind mark

which tack to start on

when to tack

whether to "run" or "broad reach" during the downwind leg

Since being at the right place at the right time with the wind on the starboard side is advantageous when encountering other boats, the model might favor a starboard tack at certain points in the course.

It seems like this could be computed given

an accurate speed model of the Club 420 under typical loading conditions for each point of sail

performance model of a well-executed tack

performance model of a well-executed gybe

This would attack the problem that, I suppose, gets answered through intuition and experience by 'good' skippers. Of course, it completely ignores the influence of other boats whose presence could easily compromise a good single-boat strategy. It also assumes the skipper and crew are good enough to maneuver in ways that satisfy the performance models; this depends on training and skill. Nevertheless, it seems like having a good theory here would be a great way to quickly gain a practical intuition that could be deployed out on the water.

Someone get me a performance model!

Git solves your problems

2009-11-12T10:20:00.010-05:00

So last night, I wanted to install a Subversion server on Emma's darling quad-core behemoth hosting akerr.net among *other things*. Inspecting the installation, it seems mod_dav was not compiled into Apache2.x, and no one after Apache 1.3 seems to suggest loading it as a dynamic shared library. I successfully compiled several test builds of Apache2.2 in user space, but matching Emma's configuration in a maintainable way would have been a nightmare inside of doubtfulness wrapped in the promise of sleepless nights down the road if things go wrong.

Then I went looking for alternatives. It seems the only administrative prerequisites for hosting a remote Git repository are that the host runs sshd and you to have an account on said machine. After installing the lovely msysgit git client for Windows on my home desktop, I performed the following.

On the remote side:


$ ssh akerr@akerr.net:
$ mkdir akerr.git && cd akerr.git
$ git --bare init
Initialized empty Git repository in /home/akerr/akerr.git/
$

On the local side:


$ cd C:/research/akerr
$ git init
$ git remote add origin "ssh://akerr@akerr.net:<ssh port>/~/akerr.git"
$ git add infrastructure/
 ...
$ git commit -a -m "initial revision of libkerr"
$ git push origin master
 ...
$

This is all that is needed. To pull from the remote host, simply perform the following:


$ git pull origin master
 ...
$

The elegant components of this solution are:
- doesn't require superuser access or configuration changes on the host
- security is a function of SSH
- git is fast and functional in an isolated environment

Thank you, Emma, for the free hosting. This helps you too. : )

Wheel of Reincarnation

2009-11-05T17:00:00.002-05:00

I read an interesting paper describing the Wheel of Reincarnation - the phenomenon in which certain specialized computations are handled by coprocessors but then these computations are shifted back to the main processor as CPUs gain in performance.

General computing power, whatever its purpose, should come from the central resources of the system. If these resources should prove inadequate, then it is the system, not the display, that needs more computing power.

[Myer and Sutherland. "On the design of display processors"]

I guess this phenomenon is satisfied by general-purpose computing on GPUs that brings the graphics processor back into the 'system' instead of being a specialized co-processor for accelerating DirectX.

With GPU Ocelot, the system is getting more power.

Emma's Birthday!

2009-10-09T17:20:00.003-04:00

Let us also observe that my darling Emma is twenty six years awesome as of yesterday. We celebrated with supper, exquisite company, Maker's Mark, and cookie cake. Presumably, celebratory ice skating is upcoming.

Happy birthday, Emma!

Atlanta Calling

2009-10-09T17:15:00.004-04:00

By now, you're surely tired of seeing the evidence of my awesome summer. To keep you interested in lieu of a complete blog post chronicling my travels to HPEC'09 and then IISWC, here is a Google Image Search that should elicit a chuckle among some of you.

Example:

Real post coming soon.

Post Summer 2009

2009-09-08T16:54:00.005-04:00

I am back from Seattle. Here is a taste of what I did this summer. Special thanks go to all of the people I met along the way.

Flying a Cessna 172S:

link

Sailing on Lake Union:

link

Bike ride to Lake Washington:

link

Helicopter flight over Seattle:

link

10,000 ft climb of Mt. Rainier:

Things I want to accomplish personally this Fall:

* nail down research agenda working toward thesis
* paper at PACT
* sail considerably, taking you and others with me
* bicycle over a notable distance
* hike somewhere notable
* undertake a slaying expedition
* gain upper body strength for rock climbing options class next Spring

Let's see if a more productive Andykerr can accomplish all of this. I'll have to put considerably less whiskey in my coffee the next few mornings.

VSIPL Design Thoughts

2009-05-13T18:59:00.002-04:00

As an implementer of the VSIPL API, I've thought of several small criticisms of the library. By posting them here, Google will record them for posterity.

* VSIPL Core Lite should require implementations to include vsip_complete() -
This could easily be a no-op on any implementation that doesn't perform some sort of lazy evaluation. Applications could call it without fear that things may break if they move to another implementation.

* VSIPL API should include vsip_usermalloc() -
This would allocate user-accessible data blocks in a manner that could be used efficiently by the library. The alternative is let the user perform its own memory allocation; this has no regard to alignment or pinned paging.

* VSIPL should include more robust error reporting methods -
Reports might be implementation-dependent, but it would enable much better robust composition.

I am confident none of these requirements would break existing implementations or negatively impact users.

Inexpensive Performance

2009-04-13T23:05:00.004-04:00

Taking a break from MemoryTraceGenerator.cpp to mention:

I'm replacing the motherboard and CPU in my dad's home desktop with an Athlon 64 x2 Kuma at 2.7 GHz and an $80 EVGA motherboard with a GeForce 8200 integrated video chipset. The total here is roughly $200 if you also purchase several GB of DDR2 RAM. This wouldn't be a terrible CUDA development platform, particularly if someone loaned you a higher-end GPU. Even without, its integrated GPU is still capable of accelerating typical applications such as video decoding.

I only mention this because what one might consider today's low-end desktop is astonishingly powerful.

It's also my first time touching AMD.

Winds

2009-04-06T15:59:00.000-04:00

Design fail: high winds bristling down Peachtree St at the 130 ft level tend to pass over and around
the balcony railings of our condo creating quite a bit of commotion. I would estimate the power spectrum has peaks at the 10 kHz frequency as well as the 5 Hz and ~100 Hz frequencies. The noise was so loud and so stationary that, at Emma's suggestion, I slept wearing ear plugs from my shooting kit. That worked for roughly two hours.

Can't we just remove the railings from the balcony during high winds?

Webcam

2009-04-04T15:59:00.003-04:00

Emma's webcam is now operational. I am using webcamXP 5 [free!] to make HTTP POST calls to upload images. At present, the server will store only the most recent image, so don't expect a log. It is updated once every hour. I may enable serving the most recent 24 hours, but that functionality will not come online today.

I used a PHP script on the server to receive the files and authenticate the poster. If you guessed the passphrase and spoofed our IP address, I guess you could upload an image of a giant eyeball and it would appear on our webcam page.

Here it is. Bookmark this:

webcam

Note the billboard between the Viewpoint and the T-Mobile sales office. The nude male ad has been replaced for an ad for The Mighty Boosh on Adult Swim.

GPU VSIPL Press

2009-04-02T11:23:00.002-04:00

GPU VSIPL is now featured on NVIDIA's CUDA Zone: GPU VSIPL. NVIDIA lists it as an application, although we only wrote the application to demonstrate the library (and achieved a 75x speedup over the world-famous TASP VSIPL distribution).

In unrelated news, Emma and I have completed our move. I need to set up her webcam.

GPGPU'09

2009-03-09T19:04:00.004-04:00

The presentation at GPGPU'09 went well yesterday. I lasted several minutes over budget [as usual], but I felt good about it. When I get home and find my credentials to the GPU VSIPL website, I'll upload my presentation and cite my paper.

Today was the first day of ASPLOS. I wound up sleeping in and missed the keynote, but I sat at the keynote speaker's table during lunch. He's a Google guy and regaled the group with clever ideas, charm, and wit.

The trouble with conferences is they fill you with a billion great ideas and no time to implement them. The wild and crazy idea session included:

* bubble wrap cores - disposable CPU cores so you can run several hotter and faster and burn them out quickly
* on-chip power - nuclear or piezoelectric
* neuro-implants - CPUs in your brain
* purely speculative cores - run parts of your program long before you ever get to them

2 of 3 beer vouchers remain.

$QR$ Presentation

2009-03-05T01:16:00.004-05:00

I finally finished my GPGPU '09 presentation. I need to script it and make sure I can present all 45 slides in 20 minutes. Some of them are closer to animations than "view graphs" so the actual number of slides I'd spend more than a few sentences on is lower.

How do you guys prepare for presentations?

Word-for-word scripting, bullet points you remember to mention in the context of a larger discussion you have with your audience, mix Adderall and a vodka tonic shortly before the presentation and hope for the best? Comment!

GPU VSIPL

2009-02-27T12:53:00.001-05:00

Released a new GPU VSIPL. Our range-doppler map application achieved 75x speedup on my machine at work.

Productivity

2009-02-19T10:06:00.001-05:00

Installing Wubi, Ubuntu for Windows. I rather wish I had a 64-bit operating system installed on this machine.

Traveling

2009-02-16T21:39:00.003-05:00

I've been fairly remiss in keeping up this blog and in keeping up appearances in real life. I'll try to at least provide an overview in this post.

I'm typing this from a hotel room in Raleigh, NC. I'm attending High-Performance Computer Architecture conference going on until Wednesday. Three other gentleman of acclaim from my lab are here, though none of us are presenting here. We're here to attend talks, ask questions, converse with presenters in the hallways, and cash in our liquor vouchers at the appropriate times.

Today's keynote was delivered by a fellow at HP and summarized the elements of HP Labs's research thrusts: intelligent infrastructure, content transformations, information management, analytics, cloud computing, and sustainability.

The first three "award" paper presentations covered (1) prefetching for linked data structures, (2) recovery and prediction from voltage emergencies, and (3) long-wire topologies for low-latency 3D networks on chip.

This afternoon, I attended several papers on multicore caching, coherence protocols, a software-managed cache implementation for IBM Cell [looking shockingly similar and not materially superior to either a software-managed cache I implemented for our PTX->Cell runtime or the SMC that ships with IBM's SDK; I didn't pipe up and notify him about this, as I don't think anyone expected it to be].

Georgia Tech is well represented here, with several full profs in attendance [not mine] and a number of students. The panel session today featured the legendary William "Bill" Dally, now of NVIDIA. As Mark Richards described, he is a brilliant and charismatic engineer who portrayed GPUs in very positive light, having lead a number of streaming architecture projects at Stanford in the past. They all agreed that multicore is important, computer architecture needs to inform computer science systems projects at various levels of abstraction [i.e. compiler writers, OS designers, and CPU architects need to have lunch at the same table every day], research projects need to be more aggressive and less product-like [i.e. make crazy and wild assumptions about the future target platforms and build from that; if your results show significant gains even though you couldn't actually build it today, maybe that's how future systems will actually be developed]. And apparently the difference between parallel computing and distributed computing is in distributed computing, you don't assume all machines are working.

Overall, this conference is far more academic than HPEC, which is very much targeted for DoD types. HPCA is very much on the academic side with some very deep and nuanced computer architecture topics occupying many of the papers. The award paper covering prefetching for linked data structures, for instance, discusses a method by which pointers may be identified on cache misses and automatically prefetched with the assumption that they point to nodes in a data structure likely to be fetched shortly anyway. The complete proceedings are presented to attendees in both print form and cleverly stored on a complimentary 1 GB USB drive / laser pointer / ballpoint pen.

That's all for now. I'll be back at some point before Thursday at 10am. I've just submitted a revised edition of my paper to GPGPU 2009. I'll fly to DC on March 7 and come back Mar 12. For future reference, make sure your ACM submission has the copyright notice on the first page; they get mad if you don't.

See you soon, dear reader(s).

New Work

2009-01-21T14:18:00.002-05:00

Just built the new PC at work.

andrew@:~$ cat /proc/cpuinfo
processor : 0
vendor_id : GenuineIntel
cpu family : 6
model : 23
model name : Intel(R) Core(TM)2 Quad CPU Q9550 @ 2.83GHz
stepping : 10
cpu MHz : 2833.268
cache size : 6144 KB
physical id : 0
siblings : 4
core id : 0
cpu cores : 4
apicid : 0
initial apicid : 0
fpu : yes
fpu_exception : yes
cpuid level : 13
wp : yes
flags : fpu vme de pse tsc msr pae mce cx8 apic sep mtrr pge mca cmov pat pse36 clflush dts acpi mmx fxsr sse sse2 ss ht tm pbe syscall nx lm constant_tsc arch_perfmon pebs bts rep_good nopl pni monitor ds_cpl vmx smx est tm2 ssse3 cx16 xtpr sse4_1 lahf_lm
bogomips : 5666.53
clflush size : 64
cache_alignment : 64
address sizes : 36 bits physical, 48 bits virtual
power management:

It says the same for the other three cores.

Defeated

2008-12-13T22:56:00.005-05:00

This is the gas valve from a Yugoslavian SKS I J.B.Welded into the gas block. When I got tired of banging my head against my desk trying to make my code run fast, I went into the living room and banged a hammer against a punch [and against my fingers] to remove this reluctant part so I could replace it.

Strangely, the difficulty was not that the J.B.Weld refused to give; I achieved small motion early on. Rather, the simple matter of non-zero volume wedging the part against the gas block walls made extracting it difficult.

Advice: attempting to join parts that were intended to be separate limits your options in the future.

Advice for those who ignored first bit of advice: buy a good selection of punches now because you'll need them. Drill a small pit in the part you're trying to separate, then use the correctly sized punch so you don't damage the surrounding material. Avoid stressing things like barrels by properly supporting the job with wooden blocks or substitute.

Native Client and Alchemy: Software as a Service?

2008-12-13T22:36:00.004-05:00

Not so long ago, Google announced Native Client which describes a way in which a browser may execute native instructions in secure sandbox. The natively executing code would run at full speed, and the sandbox would provide a platform-independent layer of abstraction so applications targeting whatever framework was distributed with Native Client The Product (tm) could run on Mac, Linux, or Windows. The paper linked from the "NaCl" page describes the security measures taken to ensure natively executing code can do nothing harmful. Presumably, just running a NaCl browser on a Linux, Windows, or Mac system as a user with minimal privileges to the file system and other resources is already sufficiently secure, but the sandbox protects the browser process as well.

Previously, Adobe announced Alchemy leveraging the Low-Level Virtual Machine bundled with the Flash player as a way to run C++ applications at near native speed with the safety of executing in a virtual machine and dynamic compilation framework. LLVM has the added benefit of being a mature project with fairly robust and well-designed development tools and runtime features.

The commonly cited application for these frameworks is "to run Quake in a browser!" While interesting in its own right, a conversation with Uniquely Joe sparked the idea that perhaps this offers a new business model for software development. Instead of buying shrink-wrapped software in a box, or downloading a torrent from your favorite pirate search engine, you simply log into your software account provided by Adobe, Google, Microsoft, or other, and use the applications you've purchased a service plan for. This benefits you because you don't have to worry about maintaining all of these software installations on the several computers you may work from, software upgrades are more seamless [perhaps less expensive for each subsequent version but now there are many more versions though only your checking account really notices], and perhaps there is a security benefit too if most of your applications run in some sort of sandbox.

This benefits software publishers because distribution costs diminish as there are no shrink-wrap boxes to produce, revenues are more constant due to the service-oriented usage model, you can push advertisements to your users and probably achieve better product integration, and you can curtail piracy since presumably you could build in a cryptographic certificate to prevent execution of your application except where authorized by your license server. With everything running in a sandbox with a uniform abstraction layer on all architectures, testing and support costs should also diminish. As we learned in CS 8803, dynamic compilation frameworks are able to perform many optimizations at runtime that enable an application to [eventually] achieve high performance tuned to the actual system on which it is executed. 'Application + dynamic optimizer' is clearly parallelizable so additional processor cores should reduce the impact of dynamic translation overheads.

Will consumers like it? While some won't even notice a difference, many might be attracted for the reasons previously listed. If Adobe can count on users having LLVM installed, or users willing to install it, it may take this route exclusively. How many people would be willing to eschew Adobe's impressive productivity suite in favor of the disjoint FOSS projects out there?

Comments welcome.

GPU QR

2008-12-11T16:59:00.002-05:00

My GPU QR algorithm achieves a 10.8x speedup over the native MATLAB QR for large matrices. This is faster than I have seen anyone report by far.

I'm working on an Intel MKL testbed application, as that's probably the fastest QR available for x86 and is the ultimate reference implementation.

Organization of Knowledge

2008-12-04T18:04:00.002-05:00

I clicked through to this page on Wikipedia. Clearly, pages like this are useful. I'm pretty sure any good PhD student, by the time they're ready to propose, should be able to sit down at Rocky Mountain, drink several pints, and rattle off something similar for their field.

News from the GPU VSIPL Front

2008-11-22T13:22:00.003-05:00

We're planning to undergo talks for licensing GPU VSIPL. Presently, we only distribute the binary and insist it not be redistributed. This lets early adopters test it out but minimizes the scope of what we're committed to supporting. When the implementation is stable and complete, we're hoping to make it available in some additional capacity. Ideally, we would partner with someone and include it in a larger VSIPL distribution.

Additionally, I'm responsible for implementing a sample application to publish at CUDA Zone. I need a team of undergraduates to finish the Test Suite for me.

Wii on PS3

2008-11-21T21:56:00.004-05:00

Someone should implement a Wii virtual machine for PS3.

This should include a hardware receiver (and drivers) for the Wii controllers. GameCube compatibility could be accomplished by translating the inputs from the PS3 controllers into what appear to be GameCube controller signals.

The Cell B.E. processor includes a 64-bit "Power Processing Unit." The Wii's CPU is also a POWER implementation. They both use GPUs developed from the same technology that goes into PC GPUs, so I would expect the PS3 to implement whatever primitive shader model is offered in the Wii. Moreover, I would expect the performance advantage of the PS3 to be sufficient to emulate any hardware-provided functionality in the Wii (use the idle SPUs for dynamic translation!).

Nintendo could continue to sell Wii games, and they'd penetrate formerly PS3-only households. I'd wager the technical obstacles to this idea aren't insurmountable.

Cells and PS3

2008-11-20T19:29:00.002-05:00

Guess the PlayStation3 was simply ahead of its time.

We're in the process of building a system that, if talks go well, will have several high-end (9800GX2) GPUs, one Cell processor, and one quad-core x86. I have no plans for tossing in a Blu-Ray player though that would be amusing.

GPU VSIPL

2008-11-19T10:26:00.002-05:00

GPU VSIPL is getting some press:

http://www.marketwatch.com/news/story/Mercury-Computer-Systems-Unveils-Multi/story.aspx?guid={A3EA2E33-A633-4021-A6C5-CFB83FBAC70B}

Here is a link to the main GPU VSIPL page, where we claim speedup of two orders of magnitude for applications well-suited to GPUs without ever writing any GPU code.

http://gpu-vsipl.gtri.gatech.edu

We just released a new version on Sunday with a number of enhancements from VSIPL Core, though we do not yet claim VSIPL Core compliance so I will not enumerate them yet. That will come by Christmas.

Super Obama World

2008-11-17T22:49:00.001-05:00

Watch out for soccer moms.

http://superobamaworld.com/

SSE3 and Manhattans

2008-11-09T23:42:00.008-05:00

Now that the semester's coding for the PTX to SPU translator is complete, I spent the weekend researching some areas that I've been thinking about but haven't had much time to investigate. So, it was a weekend spent mostly* hacking.

1.)
Streaming SIMD Extensions are a set of instructions added to the x86 instruction set beginning with the Pentium II. These instructions apply the same operator, typically floating-point {*,+, -, /}, to corresponding elements of 128-bit 4-element vector registers. Since they are parallel, you can typically perform more operations in a given number of clock cycles than with scalar floating-point code.

SSE2 and SSE3 are revisions that have added additional instructions as programmers demanded them. If you have a Pentium 4 Prescott or better, you have SSE3. If you have a 2.2 GHz P4 Northwood, you only support SSE2, and you miss out on the faster-better-cheaper possibilities concomitant with SSE3.

I spent a few hours updating my hand-rolled matrix class with compiler instrinsics (statements with the semantics of C functions but direct correspondence to CPU instructions; portable too), and I achieved 2.5x speedup for matrix multiply. SSE3 provides support for horizontal operations - operators apply to elements within the same 128-bit register. This permits the implementation of dot products and complex arithmetic without shuffle instructions and makes the code a lot faster. If your CPU doesn't support SSE3, you should probably build a new system (and use the existing system as a dedicated build machine).

2.)
CUDA is interoperable with OpenGL and Direct3D9. I spent a few hours tonight writing a quick DirectX application that renders a textured quad then performs post processing (separable 2D convolution) with a CUDA kernel. The immediate application of this would be to produce efficient visualizations for GPU-based simulations. Other ideas are to perform 3D rendering with DirectX and post-processing image-space operations with CUDA though Cg/HLSL is still probably the right way to implement that.

Also, the fragmented nature of OpenGL distributions across versions and driver providers made it more of a debugging hassle to get working than CUDA-DirectX interoperability.

3.)
Identified the need for a new power supply. Apparently, a GeForce GTX 280 has been purchased for me. I'd like to use it along with the GeForce 9800 GX2 giving me a grand total of 3 GPUs and 2 GB of GDDR3 memory. I'm working on ways to leverage all three at once, so this isn't a fool's errand. Unfortunately, my power supply cannot source enough current on enough lines to power both cards. During Christmas break, I'll make the transition.

*
During a trip to Harry and Sons, I decided to modify my usual order of Chicken Larb (Thai salad). I still ordered it, but I augmented it with a Manhattan. For those of you who don't know, a Manhattan is a cocktail of whiskey and sweet vermouth. Typically, I avoid cocktails because (1) I'd only really had bad examples and (2) cola + {rum, whiskey} is difficult to beat. The stigma of cocktails being girly drinks may have originated during my freshman year's first experiences with alcohol; vodka, grenadine, and orange juice are simply not something I'll ever combine again.

A well-made Manhattan, on the other hand, is quite strong yet simple enough to order from a busy bartender. In terms of flavor, it is quite divine. The vermouth dulls all of the whiskey's edge leaving only the wonderful caramel flavoring. Typically, it's made with 3-4 shots of the principle, so one drink takes you a long way toward inebriation while looking classy the entire time. It's my new official drink.

CUDA, DirectX9, SSE3, and the Manhattan cocktail: all for the win!

PTX to SPU Translator

2008-11-08T12:04:00.004-05:00

Our semester project for Dynamic Compilation and Managed Runtimes is a code translator from NVIDIA's PTX virtual assembly language to IBM Cell SPU. Last night, we just demonstrated it by translating a kernel that computes the complex Givens rotation of a pair of values. The resulting Cell SPU assembly source was linked with a runtime platform we developed, and I executed it on the Cell processor of a PlayStation3 in PaSTEC, my favorite cluster.

We'll be writing a paper due in December that I'll post. Additionally, we'll continue it through next spring. Exciting.

Congratulations, Joe and Sara!

2008-11-03T02:10:00.005-05:00

Joe and Sara are now married in an exceedingly elegant wedding ceremony and reception.

First post.

HPEC 2008

2008-09-25T18:05:00.002-04:00

The conference is over, and my talk completed. It went well.

Still no Erdos Number. Need to think more.

GPU VSIPL

2008-09-20T18:33:00.003-04:00

I've released GPU VSIPL: http://gpu-vsipl.gtri.gatech.edu.

After HPEC 2008, I'll upload slides and the extended abstract we will present at the conference. Additionally, we plan to include our own test suite, extended coverage of VISPL Core Profile [matrices, linear algebra, and signal processing functions in 1- and 2-D], and a document on writing high-performance VSIPL applications with GPU VSIPL.

Exciting times.

Cluster Names

2008-09-12T11:46:00.003-04:00

In a meeting today, we got onto the topic of project names and then computer naming schemes. Any new cluster needs unique and memorable identifiers for each machine. Numbers are too sterile, so clever administrators strive to select a theme capable of generating thirty or more unique names that still convey membership of the group.

Ada suggested the "vogue" cluster with nodes named after designers. We then sat around tossing out ideas for the GPU cluster being proposed. Here were several of the more amusing ideas

* dead comedians (already in use by CoC)
* high-end automobiles
* Disney characters
* 3rd world dictators

Here are a few more to have on hand in case I ever need them:

* coffee terms (already in use by emma+andy)
* classic video games
* characters in Les Miserables (spelling might be prohibitive, and eponine can be expected to crash before your simulation ends...)
* characters in practically any classic novel (Catch22, Whuthering Heights..)

Audience participation: add to the list.

Wireframe

2008-09-06T02:01:00.003-04:00

I spent some time tinkering with the HTML canvas tag and got this working in not too much time: wireframe cube. Reminds me of high school.

That said, it's odd it took the web this long to include a 2D rendering target that's less sophisticated than GDI+.

Backx0r

2008-09-05T19:36:00.004-04:00

As many of you know, I bike from this apartment to school. It's a fun ride that has just enough hills to make each commute a workout. Consequence: I get roughly two 15-minute periods of vigorous exercise every weekday whether I want it or not. With showers in the CRC and a change of clothes on hand, it's fairly convenient.

Several things that aren't, though, are the gates at this apartment complex. They are triggered by induction loops, coils of wire embedded in the pavement that detect vehicles by measuring the change in inductance of the loop when an electrical conductor passes near by. The rear gate is especially challenging since repaving has hidden the coils, and I believe it is not particularly sensitive.

However, a bit of online research and a brief experiment have changed all that. I can open it with impunity now. No need for bizarre aluminum blocks taped to my feet.

--

School and research are going well. I'm involved with a new lab doing research into heterogeneous multicore systems. Additionally, I'm taking GPU VSIPL to dizzying new heights with support for matrices and various linear algebra solvers all getting more than 20x speedups from your trusty GPU. Additionally, in 20 short days, I'll be presenting an awards paper at HPEC 2008. This makes me nervous.

Coolest class ever: Dynamic Compilation. Our semester project will be CUDA=>IBM Cell/BE translator. Stop laughing.

--

I'm taking some time to organize all the code I've written in the past that may be effectively reused. Stand by for libkerr and kerrutils.

Hotlink

2008-08-05T17:24:00.001-04:00

Hotlink

Property Allocation

2008-07-22T19:13:00.003-04:00

So, the rigors of four independent people sharing one refrigerator and kitchen have finally caught up with this house. All de facto boundaries of personal "spaces" within our fridge have disappeared in the previous weeks. Finding my cream pushed behind someone's leftovers and spilled has empowered me to speak out.

I'm working on a solution for this and the other common spaces problem. People have laid claims to their storage areas, some of which are unreasonable. People leave dishes in the sink, also unreasonable. Messes get ignored and grow as a result, unreasonable.

What have the readers of this blog implemented when sharing living spaces with people you weren't sleeping with? Any pitfalls? How did you keep dishes out of the sink?

Oh, and the first person to cite Garret Hardin's specious "Tragedy of the Commons" gets a cock punching. We're not talking about sterilizing human beings to keep the Earth from being overpopulated. We're not even talking about consuming commonly held goods. I'm just designing a workable attitude toward property rights with respect to space and time to avoid inefficiency and to prevent people from resenting and hating the others. Perhaps it's a little late, but it's a good experiment.

Golub and Van Loan wear the clock so you can know the time

2008-07-21T11:57:00.002-04:00

Just thought I'd share the magic.


U = 
[
0.3594  -0.5577  0.5229  -0.4077  0.3128  0.1494 ;
0.4158  -0.4082  0.0545  0.4089  -0.5689  -0.4082 ;
0.4449  -0.1494  -0.4741  0.4072  0.2794  0.5577 ;
0.4449  0.1494  -0.4729  -0.4093  0.2786  -0.5577 ;
0.4158  0.4083  0.0558  -0.4076  -0.5697  0.4082 ;
0.3594  0.5577  0.5216  0.4088  0.3135  -0.1494 ;
];

B = 
[
-30.5509  -0.0006  -0.0000  -0.0000  -0.0000  0.0000 ;
-0.0000  -7.4641  0.0001  -0.0000  -0.0000  0.0000 ;
-0.0000  -0.0000  1.7912  0.0035  -0.0000  0.0000 ;
0.0000  -0.0000  0.0000  1.0000  0.0008  0.0000 ;
-0.0000  0.0000  0.0000  -0.0000  0.6579  0.0000 ;
-0.0000  0.0000  0.0000  -0.0000  0.0000  -0.5359 ;
];

V' = 
[
-0.3594  -0.4158  -0.4449  -0.4449  -0.4158  -0.3594 ;
0.5577  0.4082  0.1494  -0.1494  -0.4083  -0.5577 ;
0.5234  0.0540  -0.4746  -0.4724  0.0563  0.5211 ;
-0.4072  0.4092  0.4065  -0.4100  -0.4073  0.4093 ;
0.3126  -0.5687  0.2796  0.2784  -0.5699  0.3137 ;
-0.1494  0.4082  -0.5577  0.5577  -0.4082  0.1494 ;
];


U B V' = 
[
7.0000  6.0000  5.0000  4.0000  3.0000  2.0000 ;
6.0000  7.0000  6.0000  5.0000  4.0000  3.0000 ;
5.0000  6.0000  7.0000  6.0000  5.0000  4.0000 ;
4.0000  5.0000  6.0000  7.0000  6.0000  5.0000 ;
3.0000  4.0000  5.0000  6.0000  7.0000  6.0000 ;
2.0000  3.0000  4.0000  5.0000  6.0000  7.0000 ;
];
U and V' are unitary.

Working on making SVD converge as fast as possible and to minimize storage requirements. Right now, it holds a whole copy of the tridiagonal matrix B. Storing a lot of elements known to be zero is a bit wasteful, if you believe that information theory voodoo.

Yes, I realize the output B is not *quite* diagonal, but that's with four iterations of the diagonalization process (two Givens rotations per row) and no explicit zeroing of elements.

Problems to solve:
* minimize storage requirements for B.
* figure out whether it is possible to store U and V in factored Householder forms and accumulate the Givens rotations some other way (as opposed to forming U and V explicitly then rotating them to diagonalize as I do now).
* complex Householder reflections

Singular Venereal Disease

2008-07-17T17:14:00.003-04:00

Implemented singular value decomposition in a sense. It's not integrated into VSIPL but rather written as a VSIPL application. That'll change in the next few days.

As the reference implementation of VSIPL lacks this functionality, I'm told I can submit mine to VSIPL forum for distribution perhaps. That'd be a feather in my cap.

My HPEC presentation is taking shape. I have to cover a slide a minute as it is. Probably should condense it.

Emma's here. No more time for interwebs!

Weather Models

2008-07-14T12:09:00.005-04:00

Frequently, news articles covering global warming cite weather models that predict the Earth's climate will warm in the coming decades. The articles never mention whether said models have accurately predicted past meteorological trends given historical data. Sometimes articles explicitly claim the models have not. In all cases, the expectation for a model to explain the past is categorically ignored.

I'd like to propose what appears to be an equally valid weather model:


printf("Earth's climate will be hotter 100 years from now.\n");

This clearly achieves the modern standard of rigor for getting the model published. It was cheap to develop, and as a bonus, it accurately confirms what we already know about the future of our planet!

Surely you don't disagree with overwhelming popular opinion on the subject...

DC Exploration

2008-07-14T11:50:00.003-04:00

On Thursday, I attempted to visit the Whitehouse to conduct a self-guided tour. On the Metro, I began to sweat profusely and feel light-headed. By the time I approached McPhereson Square, bloodflow to my brain apparently decreased so much that my vision became impaired. I willed myself to remain upright, and when the doors opened I walked forward and groped around for a railing on the platform. I was completely blind by then. Bad news. I'm okay now though.

I need to transmit my medical care facility doctor's note to the person who arranged my tour so they don't think I'm a slacker who failed to show up because I don't value free things.

--

My parents and sister arrived later that evening. We ate Chinese down the street and planned our weekend. I had to work Friday, so they visited Smithsonian museums of their choice. We went to eat at Nathans in Georgetown again which ruled. Saturday, my dad and I drove out to Dulles to see the Steven Udvar-Hazy Center. This is perhaps the greatest single aviation exhibit I've ever seen, and I will gladly take any readers who visit and be just as thrilled to go a second time.

I have photos, but they require post-processing and downsampling. I'll try to get them up soon. I'm creating a flickr account. A separate post will chronicle what was seen with perhaps a way to prevent the strong-willed reader from avoiding spoilers.

Saturday night we ate at The Front Page in Ballston then drove down to the Iwo Jima Memorial after supper. This proved a wise choice as it was low key, parking was free, and the view was splendid. The Netherlands presented the US with a carillon tower in 1960 as a way of demonstrating gratitude for liberating them from the Nazis during 1944-45. A player comes Saturdays from 6pm-8pm and plays. With the gates unlocked, my sister ran up the stairs to obtain a 100ft higher view of Arlington National Cemetery, the Potomac, and the Mall across the river. Next week, I plan to drag Emma along closer to 6pm such that there is more light and take photographs.

Sunday, we returned to the National Gallery and I saw the only Da Vinci work in the United States: Ginevra de Benci. It is painted on both faces of a wooden panel and is mounted in a free-standing exhibit so you can see the reverse as well as the obverse. Ginepro is apparently the Latin term for juniper tree and was used to symbolize chastity. Now we see it as a precursor to gin.

With that, they left Sunday afternoon to split the trip in half. We ate well, and everyone seemed to have a splendid time.

Hype v Economics

2008-07-09T09:29:00.003-04:00

Scary

The element gallium is in very short supply and the world may well run out of it in just a few years. Indium is threatened too, says Armin Reller, a materials chemist at Germany’s University of Augsburg. He estimates that our planet’s stock of indium will last no more than another decade.

Not so scary [ArsTechnica]

One of the arguments trotted out in most posts focuses on the fact that the price of Indium has risen from around $95 in 2002 to over $1000 in 2006. What these posts fail to point out is that it was around $180 in 2000, so $95 appears to have been carefully chosen as a starting point. The end point is pretty suspicious as well, given that Indium's price was around $665 in May. Meanwhile, everyone fails to point out that the price of Gallium has remained flat in recent years.

To understand why the supply side has not increased production substantially with this increase in price, you have to look at how Indium is produced. Indium is extracted from raw materials mainly as a by-product of Zinc production. Because the price of Zinc has not increased recently, and may decrease in the near future, producers have little incentive to increase output.

Let's try to keep our heads, Internet.

July 4th Celebrations

2008-07-09T09:14:00.003-04:00

Part 1 of my July 4th plans included grilling. The day was sunny, clear, and cool. My roommate Carter and I proceeded to skewer a number of chunks of beef and chicken. We included slices of onion, peppers, and zucchini. With kabobs constructed, I attempted to get a fire roaring on our kettle grill. Rather than rest the fuel on a grate above air vents at the lowest point and rely on convection to produce a flow of oxygen, our grill is simply a large pan. Given the dimensions of the fire I attempted, the result was an anemic heat that demanded approximately 45-60 minutes to grill each kabob.

The result may have been worth it. Each cut was tender and juicy. All those involved concluded the grilling was a success.

Still, air holes will be nice.

--

I spent the rest of the day at home. Some of that time was spent coding. When it was time for fireworks, I remained indoors due to rain.

I spent Saturday indoors as well, though I went out for a midafternoon pedal down the bike paths of Arlington. A rather interesting rail follows I66 all the way to the Key bridge from George Mason Dr (my stoop, practically). Atlanta is missing out.

July 4

2008-07-03T14:44:00.004-04:00

[from Declaration of Independence]

We, therefore, the Representatives of the united States of America, in General Congress, Assembled, appealing to the Supreme Judge of the world for the rectitude of our intentions, do, in the Name, and by Authority of the good People of these Colonies, solemnly publish and declare, That these United Colonies are, and of Right ought to be Free and Independent States; that they are Absolved from all Allegiance to the British Crown, and that all political connection between them and the State of Great Britain, is and ought to be totally dissolved; and that as Free and Independent States, they have full Power to levy War, conclude Peace, contract Alliances, establish Commerce, and to do all other Acts and Things which Independent States may of right do. And for the support of this Declaration, with a firm reliance on the protection of divine Providence, we mutually pledge to each other our Lives, our Fortunes and our sacred Honor.

I shall observe this glorious occasion with kebobs and dry ice.

Nightlife

2008-06-30T11:33:00.006-04:00

Saturday night, my friend Nick Merryman graciously invited me to hang out with him, his fiancee, and their friends. This resulted in eating Mediterranean food, drinking tequila from several establishments, dancing in a nearly vacant underground disco, and smoking a hookah. I sneakily managed to hold a pull for the duration of a hookah cycle and exhaled just as the last person exhaled; this made for some surprises. Nick managed to handwhistle with smoke which was pretty amusing.

I need to figure out how to smuggle a portable whiteboard into bars. It would make discussions easier, and we wouldn't have to busily scribble illustrative plots onto napkins.

Nick reported he just sent a PCB out for fabrication which is pretty impressive. I've never done any *real* RF circuit design and consider that to be fairly hardcore. His board, among other things, has pinouts for a 320-pin FPGA.

--

I also need to arrange a rotisserie over a woodgrill in our backyard. My roommate seems to have cooking inclinations, so one weekend we should roast a large pile of meat.

--

On the GPU front, I got TRAC installed and am using it now. Someone at GTRI is using GPU VSIPL which is interesting.

DC v Heller

2008-06-26T11:04:00.007-04:00

DC v Heller

Held:

1. The Second Amendment protects an individual right to possess a firearm unconnected with service in a militia, and to use that arm for traditionally lawful purposes, such as self-defense within the home.
Pp. 2–53.

(a) The Amendment’s prefatory clause announces a purpose, but does not limit or expand the scope of the second part, the operative clause. The operative clause’s text and history demonstrate that it connotes an individual right to keep and bear arms. Pp. 2–22.

Well done.

Update (June 30)

I was going to put this in the comments but I figure it's better if more public.

What firearms can you register in DC now?

Semiautomatic handguns are not registered

No weapon can shoot more
than one shot by a single function of the trigger,
or semi-automatically shoot more than 12 shots
without manual reloading or be readily converted
or restored to do so.

I mention this just so people realize that Heller only labels broad and general bans as unconstitutional. Neither the courts nor legislatures have received much guidance by this decision to determine what other types of restrictions might also be unconstitutional. Clearly, it helps to establish the individual right to keep and bear arms and to clarify the self-defense purpose inherent in the 2nd Amendment. It is also helpful to state that "handguns" cannot be prohibited.

It has been observed that the ruling was close. This matters to an extent, but I don't think the narrow outcome is going to be the most significant source of legal activity to follow, at least in the next few years. District v Heller acknowledges the validity of some restrictions and does not go into detail about what those might be limited to. The laws are in a flux now, and both sides see plenty of opportunity to go forward.

With regard to "no semiautomatics," the 'readily converted or restored to' clause covers any weapon for which a >12 round magazine exists by my interpretation. You can buy a 45 round drum for 1911s on GunBroker for $60 or so. Everyone but Nic recognizes the world of difference between semiautomatics and manually fed weapons. This new kind of ban is going to be the new language and style of the debate.

Amusingly, it is not contradictory to support the Heller decision and yet defend the semiautomatic ban. Watch for that from both presidential candidates as they can easily appear to appease both sides. Terms like "common sense" will fly out of mouths as effortlessly as hot air.

Condition 1

2008-06-24T10:28:00.004-04:00

While simulating, I read a short piece on 'condition one' carrying of a 1911. I didn't exactly learn anything I didn't already agree with. In fact, at the end, the author suggested an experiment I've already performed to relieve my own anguish.

To get over whatever fears you may have of carrying a 1911 in condition one, cock and lock it with an empty chamber and carry as you normally do. When undressing and disarming, take note for yourself the position of the hammer and see that the weapon does not in fact fire itself. Repeat until you're finally brave enough to rack one in as it were before holstering.

You have been carrying, right?

Survival Instinct

2008-06-24T10:06:00.002-04:00

If [the levee outside St. Charles, MO] breaches, the river will swamp 100 homes in east Winfield, as well as 3,000 acres of farm fields, several businesses and a city ballpark. A muskrat that burrowed a hole in the soft ground released a geyser of water, and officials said it took nearly six hours Monday to choke off the leak.

Bastard muskrats. With the nearby water table literally above your head, you shouldn't make holes in the ground no matter your species.

Talking to the Police

2008-06-23T17:29:00.007-04:00

The reader(s) of this blog tend to tread the same grounds, but I thought I'd point out that I finally watched

Talking with the Police

It's pretty amusing and paints an idyllic portrait of law school. Prof. James Duane cites eight good reasons why talking to the police has no upside and plenty of downside.

8.) No matter how careful you are, even if you're completely innocent and do not implicate yourself in any way, the mere act of making a statement opens yourself up to problems if someone else bears false witness against you (perhaps by mistake and police later convince them their vague feeling is in fact a solid recollection). That's bad.

That said, in November 2006 I was more or less suspected of fraud. Someone (not me) used my name and whatever information that is available about me in the Real White Pages to make an online purchase with a stolen credit card. I agreed to help the police and by the end of the interview was permitted to use the investigating officer's computer for research to save everyone time.

I haven't heard about it since, but I wonder what direction he would have taken the investigation had I refused to say anything... as I understand it, search warrants aren't all that challenging to get if you ask the right court. That alone would have been a major pain in the ass to have my computer seized shortly before Deadweek.

Did performing the opposite of Prof. Duane's advice save me some trouble? His lecture suggests that the only reason things went as well for me as they did is because the officer didn't think I was responsible to begin with. "Have you ever had your mind changed by talking to a suspect, Officer Invited Guest? No?"

Comments or analysis?

South Carolina, Georgia Reciprocity

2008-06-23T12:24:00.002-04:00

The South Carolina House recently approved a bill establishing reciprocity for CCW permits issued by other states.

Alive

2008-06-23T11:23:00.003-04:00

I'm alive. My phone has been in quiet-and-ignored mode since Emma got here, and my laptop has been off. This was meant to be no slight to anyone. I simply wanted to prioritize my focus, and it was refreshing. I'll start returning calls and texts during my lunchbreak today.

We ate at El Ranchero, an Italian restaurant, Nathans (in Georgetown), P.F. Changs, and Andy's Bed and Breakfast. These are places in which dining alone would be somewhat awkward, but with Emma they couldn't have been better.

Emma seemed to get a kick out of Georgetown. It's a scenic area with plenty of interesting shops and well-dressed and exciting passers by. We stopped in at a few and thoroughly soaked up the atmosphere. Since we drove, we took an evening motor tour of the District passing by the Whitehouse, the Supreme Court, the Capitol, a brief and not necessarily planned excursion by the FDR Memorial, the Pentagon, the Air Force Memorial, and an inventive detour through west Arlington.

Then we drank whiskey that we finally managed to procure from the state-run Alcohol Beverage Control stores. It is located half a block from my workplace and closes early, so I prognosticate some lunchbreak liquor runs when I run low in the coming months; I can only hope I run into my colleagues while I'm there.

During our visit to the National Mall, we managed to cover the Hirshhorn Museum of Modern Art and then the National Gallery. We saw a variety of cool pieces at the Hirshhorn though were occasionally faced with the typical dilemma of modern art appreciation: are you looking at a work of art or a fire code regulation? We made sort of a whirlwind tour of the National Gallery seeing several rather old busts of important people (Washing-ton, for example) and a few positively grand paintings (Lake Lucerne being my favorite).

The Spy Museum and the Newseum are planned for the next trip.

Emma has photos, and she left me her pink point-and-shoot camera (the one that has incriminated so many of us on Facebook). I'll take higher resolution photos from now on.

It was a fun visit and I can't wait to fly her back here sometime soon. Now that she's back in ATL, your lives can go back to being interesting again.

Heh. You know you like abuse, gentle reader.