Compilers and More: Programming GPUs Today

By Michael Wolfe

October 8, 2008

In the not-too-distant past, ENIAC was programmed with switches and a plugboard. Stored program computers soon followed that allowed one to write a program, load it into the computer memory, and run it. Initially, those programs had to be written in or manually translated into binary machine code, but soon assembly languages and assemblers were developed to simplify the process.

Soon followed operating systems, multiprogramming, and the concept of an application binary interface (ABI). The ABI defines the interface between an application and the operating system, libraries and other components. One aspect of an ABI is to define a calling convention, including how arguments are passed to a function and where the return value can be retrieved. For instance, the x64 ABI defines that the first six integer or pointer arguments are passed in registers (%rdi, %rsi, %rdx, %rcx, %r8, %r9); the first eight floating point arguments (single or double precision) are passed in SSE registers (%xmm0 to %xmm7); and any remaining arguments are pushed on the stack (in right-to-left order). This allows up to 14 arguments to be passed in registers, which surely captures most function calls.

But not all; WRF, the Weather Research and Forecasting Model, is used for both atmospheric research and numerical weather prediction. A version of WRF is included in the SPEC CPU2006 suite. One routine (copying from the WRF source code) is “a mixed phase ice microphysics scheme” (WSM5), with 49 arguments; it calls a subroutine (WSM52D) to handle the two-dimensional physics with 47 arguments (19 integers, 18 floating point scalars, 10 floating point arrays). Imagine writing the routine call by hand in assembly language; it takes over 100 instructions just to marshall the arguments and put them in the right place.

Instead, the computing community created higher level programming languages. While the first compiler (for the A-0 system) was more like what we would today call a loader, programming languages and compilers have progressed to where we use many higher level languages (C, Java, Fortran, others too many to enumerate) with a great increase in productivity. Much programming is done without using a textual language at all; for instance, a spreadsheet is a form of a program, and various visual programming interfaces exist. Now, the routine call in WRF with 47 arguments takes one Fortran statement, much easier to write and maintain than the corresponding assembly code:

    CALL wsm52D(t, q(ims,kms,j), qci, qrs,                 &
            w(ims,kms,j), den(ims,kms,j),                  &
            p(ims,kms,j), delz(ims,kms,j), rain(ims,j),    &
            rainncv(ims,j),delt,g, cpd, cpv, rd, rv, t0c,  &
            ep1, ep2, qmin,                                &
            XLS, XLV0, XLF0, den0, denr,                   &
            cliq,cice,psat,                                &
            j,                                             &
            ids,ide, jds,jde, kds,kde,                     &
            ims,ime, jms,jme, kms,kme,                     &
            its,ite, jts,jte, kts,kte                      )

If programming in binary is akin to using fingers and teeth, and assembly language is like using sticks and stone knives, think of higher level languages as the power tools of programming.

Enter GPUs

The earliest GPUs were hardware graphics accelerators to handle line drawing, area fill, image transfer, and so on, offloading the CPU. The adoption of standardized libraries such as OpenGL and Direct3D drove the development of hardware 3D graphics accelerators, particularly with programmable shading capability. Since 2000, the programmability of the graphics accelerator chips has improved to the point where they can be used for nongraphics applications. This has been called GPGPU (General Purpose computation on GPUs). Early GPGPU programming used the existing graphics libraries, such as OpenGL, mapping between computing concepts (array, loop, execute) and graphics concepts (texture, kernel, draw). This is truly heroic programming. It’s like using a chain saw to carve blocks of ice: in the right hands, it can produce something beautiful, but one wrong mistake and all you have is ice cubes (or worse).

More recently, the programming research and development community have tried to come up with programming models that would map well onto GPUs and similar parallel computers, particularly stream programming, evidenced in several projects: StreamIt at MIT, Sh at Waterloo (which led to RapidMind), Brook at Stanford (which spun out briefly as PeakStream, and which AMD has adopted and extended as Brook+), and others.

The GPU programming model that has caught the most attention is NVIDIA’s CUDA. The language is an extension to C; the software includes compiler, libraries, and many examples. There is a large user community, including a few dozen universities that use it in course work. Moreover, the software is free, though it (obviously) only targets NVIDIA GPUs.

Another programming model, very similar to CUDA and being sponsored by Apple and others, is OpenCL. The programming models are so similar that I’ll only point out the differences here.

My last column discussed the GPU architectures and some of the problems facing programmers who want to compute on one. The first is to have an application with enough of the right type of parallelism to map onto the GPU. As the most parallelizable simple example, let’s see what it would take to port a matrix multiplication to the GPU.

In Fortran, a matrix multiplication looks like a triply-nested loop:

    do i = 1,n
      do j = 1,m
        do k = 1,p
          a(i,j) = a(i,j) + b(i,k)*c(k,j)

In C, we have to decide whether to store the arrays linearized in a long vector, or whether to use a vector of pointers, or whether we have the degenerate case with fixed size arrays. Let’s assume we use linearized arrays:

    for( int i = 0; i < n; ++i )
      for( int j = 0; j < m; ++j )
        for( int k = 0; k < p; ++k )
          a[i+n*j] += b[i+n*k] * c[k+p*j];

Matmul is a wonderful example to use when experimenting with loop optimizations, because it can be rewritten in so many ways. The three loops can be interchanged or reordered in six ways, strip-mined or tiled, parallelized and vectorized. To optimize for vector instructions, we want the i (stride-1) index innermost, to maximize the memory fetch/store bandwidth. For parallel multiprocessor or multicore execution, we want the j index outermost, so each processor or core is computing distinct columns of a. To optimize for cache memories, we want to tile all the loops, so the innermost loops compute a submatrix multiplication where the submatrices all fit in cache. An optimized, parallelized, vectorized matmul for a quad-core processor might look like:

    jts = j tile size;
    its = i tile size;
    kts = k tile size;
    parfor( int p = 0; p < 4; ++p ) /* parallel loop */
      for( int jt = p; jt < m; jt += 4*jts )
        for( int it = 0; it < n; it += its )
          for( int kt = 0; kt < p; kt += kts )
            for( int j = jt; j < min(m,jt+jts); ++j )
              for( int k = 0; k < min(p,kt+kts); ++k )
                for( int i = 0; i < min(n,it+its); ++i ) /* vector mode */
                  a[i+n*j] += b[i+n*k] * c[k+p*j];

So, even optimizing this for a modern parallel workstation or server takes significant work, knowledge of the memory hierarchy, and experimentation. In the past, programmers would have to do this all manually, though advanced compiler technology is now able to achieve this kind of optimization automatically.

But we want to compute the matmul on the GPU, using CUDA. Let’s list the steps we must take in our program to get there.

  • Initialize the GPU; since we only have to do this once for the whole application, I’ll ignore this step.
  • Allocate memory on the GPU. We’ve already allocated the memory (explicitly or implicitly) on the CPU for the arrays, but the GPU executes from its own separate memory. So, the first thing we must do is allocate memory for new copies of the data on the GPU. In concept, it’s just like executing a malloc on the GPU, but things are never quite so simple. We can start by simply allocating linear GPU memory:

      cudaMalloc( &dev_a, n*m*sizeof(float) );
      cudaMalloc( &dev_b, n*p*sizeof(float) );
      cudaMalloc( &dev_c, p*m*sizeof(float) );

    However, we may decide (or find) that the matrix columns aren’t aligned on 64-byte boundaries (we’re using column-major storage in our example). Since aligned memory accesses are faster than unaligned, we can force alignment by using a different allocation routine:

      cudaMallocPitch( &dev_a, &pitch_a, n*sizeof(float), m );
      cudaMallocPitch( &dev_b, &pitch_b, n*sizeof(float), p );
      cudaMallocPitch( &dev_c, &pitch_c, p*sizeof(float), m );

    This returns the allocated (aligned) size for the first dimension (the pitch), given the requested sizes of the two dimensions. There is a third option, allocating CUDA Arrays and mapping them into textures, which I’ll not discuss right now.

  • Move data to the GPU. The b and c matrices on the host must be copied from host memory to GPU memory. Even though our example loops don’t actually initialize the value of a to zero, we’ll assume that’s what we wanted, so we only have to move b and c. The actual data movement will be done with a hardware DMA transfer. Hardware DMA doesn’t know about virtual memory and is optimized to move large contiguous chunks of memory across the PCI bus. We can ignore that issue and just move the data with specialized memcpy call:

      cudaMemcpy2D( dev_b, pitch_b, b, n*sizeof(float), n*sizeof(float), p,
          cudaMemcpyHostToDevice );
      cudaMemcpy2D( dev_c, pitch_c, c, p*sizeof(float), p*sizeof(float), m,
          cudaMemcpyHostToDevice );

    The arguments give the destination pointer and pitch, the source pointer and pitch, the two dimension sizes, and copy direction. If we want to optimize the data transfer, we can allocate the host arrays in page-locked (pinned) memory. This makes sure the arrays don’t get paged out by the virtual memory manager. The disadvantage is that pinning large amounts of memory reduces the amount of memory available for paging, potentially reducing performance for other applications running at the same time. CUDA provides handy routines to allocate and free pinned host memory. OpenCL seems to provide the ability to allocate and copy data in a single function call.

  • Select the kernel domain. As I mentioned last time, the GPU actually executes a (usually small) scalar kernel program on each point of a multidimensional domain. The selected domain affects both the host program (a little) and the kernel program (a lot). Moreover, the domain determines how much of what kind of parallelism is being used. I’m going to expand on this point more in my next column, but for now let’s assume we’ve chosen to execute the i and j loops in parallel. This gives us a kernel domain of nxm, where the body of the kernel is the k loop.
  • Write the GPU kernel. Again, I’ll expand on this next time around, but the kernel might look like:

    __global__ void mmkernel( float* a, float* b, float* c,
      int pitch_a, int pitch_b, int pitch_c,
      int n, int m, int p )
      int i = blockIdx.x*32+threadIdx.x;
      int j = blockIdx.y;
      float sum = 0.0;
      for( int k = 0; k < p; ++k )
      sum += b[i+pitch_b*k] * c[k+pitch_c*j];
      a[i+pitch_a*j] = sum;

    Recalling the last column, the kernel will run 32 copies in SIMT mode, with n/32xm thread blocks executing in parallel on the various GPU multiprocessors. Note: this is a particularly unoptimized matrix multiplication kernel, but it should work.

  • Run the kernel. Here is one point where the CUDA model really shines. Running the kernel actually takes several steps, which are nicely hidden by the CUDA compiler. The steps include:
    • Load the kernel to the GPU. When we run a program on your CPU, we depend on the operating system to load our program and prepare it for execution. If we use explicit shared objects or dynamic load libraries, our program will search for the object and load it at run time. GPU kernels use a similar model; the GPU kernel must be downloaded by the application from the host to the GPU. One significant benefit is the GPU kernel is typically stored in a portable format; when it is downloaded, it is translated and optimized for the specific GPU installed in the system. This lets us run the same program on systems with different GPUs, without having to recompile or reoptimize the application. Since the GPU manufacturers come out with new models every 9-12 months, this works to our advantage.
    • Define the execution domain; we’ve already decided on the domain, now we have to put it in the program.
    • Pass the arguments to the kernel from the host.
    • Launch the kernel. The execution can proceed asynchronously with the host, and the host can test whether the kernel has completed across its whole domain.

    In CUDA, this is done with a few lines:

       dim3 threads( 32 );
       dim3 grid( n/32, m );
       mmkernel<<< grid, threads >>>( dev_a, dev_b, dev_c,
                   pitch_a, pitch_b, pitch_c, n, m, p );

    The NVCC compiler translates this into the steps outlined above.

    OpenCL is not so convenient. Since it is library-based, it can’t depend on a compiler to simplify the steps. Instead, we will have to do each step explicitly, something approaching:

      /* program is a prebuilt kernel program */
      kernel = clCreateKernel( program, “mmkernel” );
      grid[0] = n;
      grid[1] = m;
      threads[0] = 32;
      threads[1] = 1;
      /* context is the GPU compute context */
      range = clCreateNDRangeContainer( context, 0, 2, grid, threads );
      clSetKernelArg( kernel, 0, dev_a, sizeof(dev_a), NULL );
      clSetKernelArg( kernel, 1, dev_b, sizeof(dev_b), NULL );
      clSetKernelArg( kernel, 2, dev_c, sizeof(dev_c), NULL );
      clSetKernelArg( kernel, 3, pitch_a, sizeof(pitch_a), NULL );
      clSetKernelArg( kernel, 4, pitch_b, sizeof(pitch_b), NULL );
      clSetKernelArg( kernel, 5, pitch_c, sizeof(pitch_c), NULL );
      clSetKernelArg( kernel, 6, n, sizeof(n), NULL );
      clSetKernelArg( kernel, 7, m, sizeof(m), NULL );
      clSetKernelArg( kernel, 8, p, sizeof(p), NULL );
      /* queue is a GPU work queue */
      clExecuteKernel( queue, kernel, NULL, range, NULL, 0, NULL );

  • Wait for the kernel to finish. If we have a more complex computation, we might queue up several kernels; they will execute one after the other as they finish. In that case, we only have to wait until the last kernel is done.
  • Move results back from the GPU to the host. This, after all, is why we are doing the computation, to get the results. This is simply the inverse of loading data onto the GPU:

       cudaMemcpy2D( a, n*sizeof(float), dev_a, pitch_a, n*sizeof(float), m,
          cudaMemcpyDeviceToHost );

  • Free the device memory. Again, the inverse of allocation:

   cudaFree( dev_a );
   cudaFree( dev_b );
   cudaFree( dev_c );

CUDA (and, from reports, OpenCL) lets us program the GPU using a familiar language, C. However, a great deal of the programming is done using library calls, and the computation itself, the matrix multiplication, is divided into two parts: the kernel, on the GPU, and its invocation on the host. What was a multi-dimensional loop, which modern compilers are pretty good at optimizing, has turned into dozens of lines of code to manage memory, move data, and deal with the architectural specifics of the GPU.

I’m sure some (or most, or perhaps all) of the readers will at this point say “You shouldn’t be programming matrix multiplication anyway; just call a library routine.” That’s true; in fact, NVIDIA provides versions of the BLAS routines, including SGEMM, which give very good performance. But, if it takes this much effort to get a matrix multiplication moved to the GPU, imagine how much effort it takes to move a real computation (say, 180,000 lines of WRF, or even just the 10,000 lines of microphysics). Here I had only three arrays with regular access patterns and full freedom to parallelize the loops. Can you begin to see the difficulties?

And I haven’t even begun the real programming process, such as handling error returns from the runtime calls. Given the CUDA or OpenCL code, how portable will it be, how will I maintain it to keep it at the peak of efficiency?

As I mentioned above, the CUDA NVCC compiler simplifies some of the coding details. Compilers are good at bookkeeping, organizing details about memory addressing, alignment, and so on. OpenCL seems to be a step backwards, from a compiler-oriented solution aimed at raising the level of programming to a library-oriented solution aimed at giving low-level control to the programmer.

Low-level control isn’t bad; that would be like saying assembly language is bad. It isn’t bad, but it’s only necessary for a very small bit of the programming we do today. If high level languages are the power tools of programming, we seem to be taking a step back to hand drills and saws. Many woodworkers prefer hand tools, and they can make beautiful furniture, but the cost is high and productivity is low. We need hand tools, but we’ll be much more productive with better power tools.

In my next column, I’ll look at the matrix multiplication kernel, exploring various ways to optimize for parallelism and memory bandwidth, and presenting performance. Whether you think programming the host side of a GPU program is hard, or just work, you’ll be entertained, enlightened or frightened when you see what goes into the GPU side. I’m on record as saying that parallel programming isn’t easy and never will be, but we can and should develop the tools and training to turn this Acceleration Nightmare into something about as scary as a Halloween Haunted House.

Subscribe to HPCwire's Weekly Update!

Be the most informed person in the room! Stay ahead of the tech trends with industy updates delivered to you every week!

Supercomputers Streamline Prediction of Dangerous Arrhythmia

June 2, 2020

Heart arrhythmia can prove deadly, contributing to the hundreds of thousands of deaths from cardiac arrest in the U.S. every year. Unfortunately, many of those arrhythmia are induced as side effects from various medicati Read more…

By Staff report

Indiana University to Deploy Jetstream 2 Cloud with AMD, Nvidia Technology

June 2, 2020

Indiana University has been awarded a $10 million NSF grant to build ‘Jetstream 2,’ a cloud computing system that will provide 8 aggregate petaflops of computing capability in support of data analysis and AI workload Read more…

By Tiffany Trader

10nm, 7nm, 5nm…. Should the Chip Nanometer Metric Be Replaced?

June 1, 2020

The biggest cool factor in server chips is the nanometer. AMD beating Intel to a CPU built on a 7nm process node* – with 5nm and 3nm on the way – has been instrumental to AMD’s datacenter market resurgence. Nanomet Read more…

By Doug Black

Supercomputer-Powered Protein Simulations Approach Lab Accuracy

June 1, 2020

Protein simulations have dominated the supercomputing conversation of late as supercomputers around the world race to simulate the viral proteins of COVID-19 as accurately as possible and simulate potential bindings in t Read more…

By Oliver Peckham

HPC Career Notes: June 2020 Edition

June 1, 2020

In this monthly feature, we'll keep you up-to-date on the latest career developments for individuals in the high-performance computing community. Whether it's a promotion, new company hire, or even an accolade, we've got Read more…

By Mariana Iriarte

AWS Solution Channel

Computational Fluid Dynamics on AWS

Over the past 30 years Computational Fluid Dynamics (CFD) has grown to become a key part of many engineering design processes. From aircraft design to modelling the blood flow in our bodies, the ability to understand the behaviour of fluids has enabled countless innovations and improved the time to market for many products. Read more…

Supercomputer Modeling Shows How COVID-19 Spreads Through Populations

May 30, 2020

As many states begin to loosen the lockdowns and stay-at-home orders that have forced most Americans inside for the past two months, researchers are poring over the data, looking for signs of the dreaded second peak of t Read more…

By Oliver Peckham

Indiana University to Deploy Jetstream 2 Cloud with AMD, Nvidia Technology

June 2, 2020

Indiana University has been awarded a $10 million NSF grant to build ‘Jetstream 2,’ a cloud computing system that will provide 8 aggregate petaflops of comp Read more…

By Tiffany Trader

10nm, 7nm, 5nm…. Should the Chip Nanometer Metric Be Replaced?

June 1, 2020

The biggest cool factor in server chips is the nanometer. AMD beating Intel to a CPU built on a 7nm process node* – with 5nm and 3nm on the way – has been i Read more…

By Doug Black

COVID-19 HPC Consortium Expands to Europe, Reports on Research Projects

May 28, 2020

The COVID-19 HPC Consortium, a public-private effort delivering free access to HPC processing for scientists pursuing coronavirus research – some utilizing AI Read more…

By Doug Black

$100B Plan Submitted for Massive Remake and Expansion of NSF

May 27, 2020

Legislation to reshape, expand - and rename - the National Science Foundation has been submitted in both the U.S. House and Senate. The proposal, which seems to Read more…

By John Russell

IBM Boosts Deep Learning Accuracy on Memristive Chips

May 27, 2020

IBM researchers have taken another step towards making in-memory computing based on phase change (PCM) memory devices a reality. Papers in Nature and Frontiers Read more…

By John Russell

Hats Over Hearts: Remembering Rich Brueckner

May 26, 2020

HPCwire and all of the Tabor Communications family are saddened by last week’s passing of Rich Brueckner. He was the ever-optimistic man in the Red Hat presiding over the InsideHPC media portfolio for the past decade and a constant presence at HPC’s most important events. Read more…

Nvidia Q1 Earnings Top Expectations, Datacenter Revenue Breaks $1B

May 22, 2020

Nvidia’s seemingly endless roll continued in the first quarter with the company announcing blockbuster earnings that exceeded Wall Street expectations. Nvidia Read more…

By Doug Black

Microsoft’s Massive AI Supercomputer on Azure: 285k CPU Cores, 10k GPUs

May 20, 2020

Microsoft has unveiled a supercomputing monster – among the world’s five most powerful, according to the company – aimed at what is known in scientific an Read more…

By Doug Black

Supercomputer Modeling Tests How COVID-19 Spreads in Grocery Stores

April 8, 2020

In the COVID-19 era, many people are treating simple activities like getting gas or groceries with caution as they try to heed social distancing mandates and protect their own health. Still, significant uncertainty surrounds the relative risk of different activities, and conflicting information is prevalent. A team of Finnish researchers set out to address some of these uncertainties by... Read more…

By Oliver Peckham

[email protected] Turns Its Massive Crowdsourced Computer Network Against COVID-19

March 16, 2020

For gamers, fighting against a global crisis is usually pure fantasy – but now, it’s looking more like a reality. As supercomputers around the world spin up Read more…

By Oliver Peckham

[email protected] Rallies a Legion of Computers Against the Coronavirus

March 24, 2020

Last week, we highlighted [email protected], a massive, crowdsourced computer network that has turned its resources against the coronavirus pandemic sweeping the globe – but [email protected] isn’t the only game in town. The internet is buzzing with crowdsourced computing... Read more…

By Oliver Peckham

Global Supercomputing Is Mobilizing Against COVID-19

March 12, 2020

Tech has been taking some heavy losses from the coronavirus pandemic. Global supply chains have been disrupted, virtually every major tech conference taking place over the next few months has been canceled... Read more…

By Oliver Peckham

Supercomputer Simulations Reveal the Fate of the Neanderthals

May 25, 2020

For hundreds of thousands of years, neanderthals roamed the planet, eventually (almost 50,000 years ago) giving way to homo sapiens, which quickly became the do Read more…

By Oliver Peckham

DoE Expands on Role of COVID-19 Supercomputing Consortium

March 25, 2020

After announcing the launch of the COVID-19 High Performance Computing Consortium on Sunday, the Department of Energy yesterday provided more details on its sco Read more…

By John Russell

Steve Scott Lays Out HPE-Cray Blended Product Roadmap

March 11, 2020

Last week, the day before the El Capitan processor disclosures were made at HPE's new headquarters in San Jose, Steve Scott (CTO for HPC & AI at HPE, and former Cray CTO) was on-hand at the Rice Oil & Gas HPC conference in Houston. He was there to discuss the HPE-Cray transition and blended roadmap, as well as his favorite topic, Cray's eighth-gen networking technology, Slingshot. Read more…

By Tiffany Trader

Honeywell’s Big Bet on Trapped Ion Quantum Computing

April 7, 2020

Honeywell doesn’t spring to mind when thinking of quantum computing pioneers, but a decade ago the high-tech conglomerate better known for its control systems waded deliberately into the then calmer quantum computing (QC) waters. Fast forward to March when Honeywell announced plans to introduce an ion trap-based quantum computer whose ‘performance’ would... Read more…

By John Russell

Leading Solution Providers

SC 2019 Virtual Booth Video Tour



Tech Conferences Are Being Canceled Due to Coronavirus

March 3, 2020

Several conferences scheduled to take place in the coming weeks, including Nvidia’s GPU Technology Conference (GTC) and the Strata Data + AI conference, have Read more…

By Alex Woodie

Exascale Watch: El Capitan Will Use AMD CPUs & GPUs to Reach 2 Exaflops

March 4, 2020

HPE and its collaborators reported today that El Capitan, the forthcoming exascale supercomputer to be sited at Lawrence Livermore National Laboratory and serve Read more…

By John Russell

‘Billion Molecules Against COVID-19’ Challenge to Launch with Massive Supercomputing Support

April 22, 2020

Around the world, supercomputing centers have spun up and opened their doors for COVID-19 research in what may be the most unified supercomputing effort in hist Read more…

By Oliver Peckham

Cray to Provide NOAA with Two AMD-Powered Supercomputers

February 24, 2020

The United States’ National Oceanic and Atmospheric Administration (NOAA) last week announced plans for a major refresh of its operational weather forecasting supercomputers, part of a 10-year, $505.2 million program, which will secure two HPE-Cray systems for NOAA’s National Weather Service to be fielded later this year and put into production in early 2022. Read more…

By Tiffany Trader

15 Slides on Programming Aurora and Exascale Systems

May 7, 2020

Sometime in 2021, Aurora, the first planned U.S. exascale system, is scheduled to be fired up at Argonne National Laboratory. Cray (now HPE) and Intel are the k Read more…

By John Russell

Summit Supercomputer is Already Making its Mark on Science

September 20, 2018

Summit, now the fastest supercomputer in the world, is quickly making its mark in science – five of the six finalists just announced for the prestigious 2018 Read more…

By John Russell

Fujitsu A64FX Supercomputer to Be Deployed at Nagoya University This Summer

February 3, 2020

Japanese tech giant Fujitsu announced today that it will supply Nagoya University Information Technology Center with the first commercial supercomputer powered Read more…

By Tiffany Trader

Australian Researchers Break All-Time Internet Speed Record

May 26, 2020

If you’ve been stuck at home for the last few months, you’ve probably become more attuned to the quality (or lack thereof) of your internet connection. Even Read more…

By Oliver Peckham

  • arrow
  • Click Here for More Headlines
  • arrow
Do NOT follow this link or you will be banned from the site!
Share This