Engineering with a Global Dataset

Engineering with a Global Dataset

Strava Labs

Live activity uploads.
Engineering with a Global Dataset
The Strava dataset is growing by millions of activities a day and comprises trillions of GPS data points.
Strava Labs showcases interesting side projects enabled by this dataset. These projects are independent of Strava and may be added, removed or break at any time.
Interactive and experimental software the team is building.
Read about our technology, culture and latest projects.
A comprehensive portal for information about the Strava API.

Featured Projects and Articles
Global Heatmap
Updated for 2017! Global Heatmap
Over 1 billion activities, 13 trillion data points create the ultimate map of athlete playgrounds.
The Clusterer
All 400 million Strava activities from 2017, clustered! Explore worldwide clusters by distance and activity type.
Project KODOS
Project KODOS
Detailed statistics on your personal data.
2014 vs. 2015 Heatmap
Strava, 2014 vs. 2015
Roads change, people change. Compare the 2014 and 2015 heatmaps and see the change.
Playback your activity along with those riding near you. Find who you passed or ‘flew by’ while riding.
Auto-drawing map geometry from the Strava global dataset.
Achievement Map
Achievement Map
Your cycling KOM/QOMs and running CRs highlighted on an interactive map.
see all projects »
Strava Labs is where Strava engineering occasionally releases side projects independent of our product and share our processes and stories. Projects may be added, removed or break at any time, so use them at your own risk.

© 2016 Strava, Inc. All rights reserved. Terms Privacy Policy @StravaEng

Achieved Occupancy

Achieved Occupancy

Open topic with navigation

You are here: Developer Tools > Desktop Developer Tools > NVIDIA Nsight Visual Studio Edition > Achieved Occupancy
Achieved Occupancy
(Undefined variable: MyVariables.NsightVSEMainHeader)
Send Feedback
For all CUDA kernel launches recorded in both Profile and Trace modes, the Occupancy experiment detail pane shows “Theoretical Occupancy”, the upper limit for occupancy imposed by the kernel launch configuration and the capabilities of the CUDA device. The Achieved Occupancy Profile mode experiment measures occupancy during execution of the kernel, and adds the achieved values to the Occupancy experiment detail pane alongside the theoretical values. Additional graphs show achieved occupancy per SM, and illustrate how occupancy can be controlled by varying compiler and launch parameters.

Definition of Occupancy
The CUDA C Programming Guide explains how a CUDA device’s hardware implementation groups adjacent threads within a block into warps. A warp is considered active from the time its threads begin executing to the time when all threads in the warp have exited from the kernel. There is a maximum number of warps which can be concurrently active on a Streaming Multiprocessor (SM), as listed in the Programming Guide’s table of compute capabilities. Occupancy is defined as the ratio of active warps on an SM to the maximum number of active warps supported by the SM. Occupancy varies over time as warps begin and end, and can be different for each SM.

Low occupancy results in poor instruction issue efficiency, because there are not enough eligible warps to hide latency between dependent instructions. When occupancy is at a sufficient level to hide latency, increasing it further may degrade performance due to the reduction in resources per thread. An early step of kernel performance analysis should be to check occupancy and observe the effects on kernel execution time when running at different occupancy levels.

Theoretical Occupancy
There is an upper limit for active warps, and thus also for occupancy, derivable from the launch configuration, compile options for the kernel, and device capabilities. Each block of a kernel launch gets distributed to one of the SMs for execution. A block is considered active from the time its warps begin executing to the time when all warps in the block have exited from the kernel. The number of blocks which can execute concurrently on an SM is limited by the factors listed below. The upper limit for active warps is the product of the upper limit for active blocks and the number of warps per block. Thus, the upper limit for active warps can be raised by increasing the number of warps per block (defined by block dimensions), or by changing the factors limiting how many blocks can fit on an SM to allow more active blocks. The limiting factors are:

Warps per SM
The SM has a maximum number of warps that can be active at once. Since occupancy is the ratio of active warps to maximum supported active warps, occupancy is 100% if the number of active warps equals the maximum. If this factor is limiting active blocks, occupancy cannot be increased. For example, on a GPU that supports 64 active warps per SM, 8 active blocks with 256 threads per block (8 warps per block) results in 64 active warps, and 100% theoretical occupancy. Similarly, 16 active blocks with 128 threads per block (4 warps per block) would also result in 64 active warps, and 100% theoretical occupancy.

Blocks per SM
The SM has a maximum number of blocks that can be active at once. If occupancy is below 100% and this factor is limiting active blocks, it means each block does not contain enough warps to reach 100% occupancy when the device’s active block limit is reached. Occupancy can be increased by increasing block size. For example, on a GPU that supports 16 active blocks and 64 active warps per SM, blocks with 32 threads (1 warp per block) result in at most 16 active warps (25% theoretical occupancy), because only 16 blocks can be active, and each block has only one warp. On this GPU, increasing block size to 4 warps per block makes it possible to achieve 100% theoretical occupancy.

Registers per SM
The SM has a set of registers shared by all active threads. If this factor is limiting active blocks, it means the number of registers per thread allocated by the compiler can be reduced to increase occupancy (see __launch_bounds__). Kernel execution time and average eligible warps should be monitored carefully when adjusting registers per thread to control occupancy. The performance gain from improved latency hiding due to increased occupancy may be outweighed by the performance loss of having fewer registers per thread, and spilling to local memory more often. The best-performing balance of occupancy and registers per thread can be found experimentally by tracing the kernel compiled with different numbers of registers per thread, controlled via __launch_bounds__.

Shared Memory per SM
The SM has a fixed amount of shared memory shared by all active threads. If this factor is limiting active blocks, it means the shared memory needed per thread can be reduced to increase occupancy. Shared memory per thread is the sum of “static shared memory,” the total size needed for all __shared__ variables, and “dynamic shared memory,” the amount of shared memory specified as a parameter to the kernel launch. For some CUDA devices, the amount of shared memory per SM is configurable, trading between shared memory size and L1 cache size. If such a GPU is configured to use more L1 cache and shared memory is the limiting factor for occupancy, then occupancy can also be increased by choosing to use less L1 cache and more shared memory.

Achieved Occupancy
Theoretical occupancy shows the upper bound active warps on an SM, but the true number of active warps varies over the duration of the kernel, as warps begin and end. As explained in Issue Efficiency, an SM contain one or more warp schedulers. Each warp scheduler attempts to issue instructions from a warp on each clock cycle. To sufficiently hide latencies between dependent instructions, each scheduler must have at least one warp eligible to issue an instruction every clock cycle. Maintaining as many active warps as possible (a high occupancy) throughout the execution of the kernel helps to avoid situations where all warps are stalled and no instructions are issued. Achieved occupancy is measured on each warp scheduler using hardware performance counters to count the number of active warps on that scheduler every clock cycle. These counts are then summed across all warp schedulers on each SM and divided by the clock cycles the SM is active to find the average active warps per SM. Dividing by the SM’s maximum supported number of active warps gives the achieved occupancy per SM averaged over the duration of the kernel, which is shown in the Achieved Occupancy Chart. Averaging across all SMs gives the overall achieved occupancy, which is shown alongside theoretical occupancy in the experiment details pane.

Causes of Low Achieved Occupancy
Achieved occupancy cannot exceed theoretical occupancy, so the first step toward increasing occupancy should be to increase theoretical occupancy by adjusting the limiting factors. The next step is to check if the achieved value is close to the theoretical value. The achieved value will be lower than the theoretical value when the theoretical number of active warps is not maintained for the full time the SM is active. This occurs in the following situations:

Unbalanced workload within blocks
If warps within a block do not all execute for the same amount of time, the workload is said to be unbalanced. This means there are fewer active warps at the end of the kernel, which is a problem known as “tail effect”. The best solution is to try having a more balanced workload among the warps in each block.
Unbalanced workload across blocks
If blocks within a grid do not all execute for the same amount of time, this workload is also unbalanced, but the efficiency of the device can be improved without having to change to a more balanced workload. Launching more blocks will allow new blocks to begin as others finish, meaning the tail effect does not occur inside every block, but only at the end of the kernel. If there are not more blocks to launch, running concurrent kernels with similar block properties can achieve the same effect.
Too few blocks launched
The upper limit for active blocks per SM is determined by the theoretical occupancy, but that calculation does not account for a launch with fewer than that number of blocks per SM. The number of SMs on the device times the maximum active blocks per SM is called a “full wave”, and launching less than a full wave results in low achieved occupancy. For example, on a device with 15 SMs, and a configuration expecting 100% theoretical occupancy with 4 blocks per SM, a full wave would be 60 blocks. Launching only 45 blocks (assuming a balanced workload) will result in approximately 75% achieved occupancy.
Partial last wave
The SM has a maximum number of warps that can be active at once. Since occupancy is the ratio of active warps to maximum supported active warps, occupancy is 100% if the number of active warps equals the maximum. If this factor is limiting active blocks, occupancy cannot be increased. For example, on a GPU that supports 64 active warps per SM, 8 active blocks with 256 threads per block (8 warps per block) results in 64 active warps, and 100% theoretical occupancy. Similarly, 16 active blocks with 128 threads per block (4 warps per block) would also result in 64 active warps, and 100% theoretical occupancy.
Varying Block Size

Shows how varying the block size while holding other parameters constant would affect the theoretical occupancy. The circled point shows the current number of threads per block and the current upper limit of active warps. Note that the number of active warps is not the number of warps per block (that is threads per block divided by warp size, rounded up). If the chart’s line goes higher than the circle, changing the block size could increase occupancy without changing the other factors.

Varying Register Count

Shows how varying the register count while holding other parameters constant would affect the theoretical occupancy. The circled point shows the current number of registers per thread and the current upper limit of active warps. If the chart’s line goes higher than the circle, changing the number of registers per thread could increase occupancy without changing the other factors.

Varying Shared Memory Usage

Shows how varying the shared memory usage while holding other parameters constant would affect the theoretical occupancy. The circled point shows the current amount of shared memory per block and the current upper limit of active warps. If the chart’s line goes higher than the circle, changing the amount of shared memory per block could increase occupancy without changing the other factors.

Achieved Occupancy Per SM

The achieved occupancy for each SM. The values reported are the average across all warp schedulers for the duration of the kernel execution. The line across all bars is the average, which is the number reported as Achieved Occupancy in the other tables.

Low occupancy is not a problem in itself, but it usually results in having too few eligible warps. If the percentage of cycles with no eligible warp in the Warp Issue Efficiency chart is high …
… try to increase the number of active warps if possible. In many cases increasing the number of active warps will result in an larger pool of eligible warps. If …
… the theoretical occupancy is low, try to optimize the execution configuration of the kernel launch, using the Occupancy table to identify which factor(s) are limiting occupancy. If you are register limited do not rule out experimenting with launch bounds to increase occupancy, even if this results in some register spilling.
… the achieved occupancy is well below the theoretical occupancy, check the Instruction Statistics experiment for highly unbalanced workloads or tail effects. Potential strategies may include splitting the kernel grid in a more fine granular way, distribute work across the blocks in a more balanced way, avoiding gathering the final result on a single block, warp, or thread.
… the Pipe Utilization experiment shows a particular pipeline is already fully utilized, increasing active warps is unlikely to results in more eligible warps, because all additional active warps will stall trying to access the oversubscribed pipeline. In this case, try to reduce the load on this pipeline or investigate if the expected peak performance for the target hardware is already reached.

previous topic next topic
NVIDIA GameWorks Documentation Rev. 1.0.150630 ©2015. NVIDIA Corporation. All Rights Reserved.

grep your way to freedom

grep your way to freedom

/* HACK */
grep your way to freedom
JAN 2018
I was recently eating lunch with someone who asked a question that I became somewhat obsessed with: what happens when you grep a file and append the contents to the same file? It’s one of these nasty cases of modifying a file that you’re reading from.

Huge thank you to Mehmet Emre & Antal Spector-Zabusky for pair programming down this rabbit hole!

Problem Spec
To be precise, what happens when you run the following code snippet?

# Make a junk file that just has the char a in it
echo “a” > test.txt

# Grep for a & append back to the same file
grep “a” test.txt >> test.txt
Okay before you run it, what are some reasonable things that could happen? And which of those do you think actually happens?

This nice photo (of the magical Zion Narrows!) is here to avoid giving it away in case you want to think about it… scroll on for some possible answers.

Possible Behavior
Okay, here are some reasonable things that could happen:

grep could finish scanning the file, then append. If this is the case, the grep would terminate, and if we ran cat test.txt we’d see that the test file has two a’s.

grep could read from the file, find an ‘a’, and then append that ‘a’ to the file. It could then check if there was any more file to read, and see the ‘a’ it just wrote. If this is the case, grep would get stuck in an infinite loop, and if we killed it we would see that the test file has many a’s in it.

This could be an error. If this is the case, grep could do nothing and return an error code.

Some exciting race condition that causes either a non-deterministic mix of 1 & 2, or maybe a deadlock.

Okay so what actually happens?
… it depends!

First thing it depends upon is your OS. If you run the GNU grep (ie, on a Linux box) you get this nice error message:

grep: input file ‘test.txt’ is also the output
I was actually not expecting any implementation to do error handling like that. That’s nice!

If you run the BSD grep (ie, on a Mac) you get:

$ echo “a” > test.txt
$ grep “a” test.txt >> test.txt
$ cat test.txt
So it’s the case where it just terminated. Okay, well that’s boring, so let’s rinse and repeat by running the same command, and then printing out how many lines the file has:

And about 15 up-arrow+enters later:

We get stuck in an infinite loop!

How fun! We’ll look into why this is happening in a second, but first,

Follow up question
What happens if you try the same kind of command but with cat instead of grep?

# Make another junk file that just has the char a in it
echo “a” > cattest.txt

# Cat & append back to the same file
cat cattest.txt >> cattest.txt
Follow up answer
We get stuck in an infinite loop right off the bat. Inconsistent behavior! Excellent.

Okay, let’s look into what’s actually happening. As you might guess, cat is simpler than grep, so let’s figure out what’s going on there first.

Since I’m running a mac, I’ll be referring to Apple’s recently opensourced code for cat which unfortunately doesn’t have line numbers- I’ll copy the code we care about below but if you want to follow along in the file, you can control-F for write.

while ((nr = read(rfd, buf, bsize)) > 0)
for (off = 0; nr; nr -= nw, off += nw)
if ((nw = write(wfd, buf + off, (size_t)nr)) > test.txt
If we run grep with the line-buffered flag then we get the same behavior as cat – it goes into an infinite loop even for very small files. So clearly the answer has something to do with buffering.

Line buffering
So if we check the grep source for something about line buffering, we’ll find that the –line-buffered flag calls something called setlinebuf

if (lbflag)
which sets stdout to have this behavior:

line buffered characters are saved up until a newline is output or input is read from any stream attached to a terminal device (typically stdin).

Cool. So if it spins with the –line-buffered flag, then probably what’s causing that is that the write buffer is being flushed rather than the read buffer. To verify this, let’s snoop on the read & write calls. On mac I’ll call

sudo rwsnoop -n grep
to start watching for read/write calls; this will give me info about any read/write calls that grep performs. On Linux you can use strace. (Sidenote: you’ll need to let dtrace through SIP to get this to work)

Okay, so after running that in one shell, I’ll pop open another shell and repeat this whole grep experiment from the beginning.

Sweet so we can see that we’ve got two read calls from grep, the first of which reads 2 bytes, the second of which reads 0 bytes. Then we’ve got a write of 2 bytes. That seems pretty consistent with the behavior we’re seeing. Running a few more rounds:

Basically the same story; we can see that now we’re reading 4 bytes, and 8 bytes- again exactly what we’d expect. Skipping forward to where we go into the infinite loop:

The last round before we get stuck we can see that we’re doing two reads of size 32KB before getting a read of size 0 and doing the write. This means the looping we’re seeing is definitely not caused by buffering on the read buffer!

When we do start looping, we do three reads of size 32KB and then a 64KB write; then repeat! Okay, so now all I want to know is where are these buffer sizes coming from?

grep source
First up: what’s grep doing? Well if we look at the main it just processes a ton of flags, and then (very near the bottom) calls procfile.

procfile is in util.c along with just about everything else we’re going to care about. procfile does two things we care about, first it calls grep_open (which will read the file) and then procline (which will write the results).

* Opens a file and processes it. Each file is processed line-by-line
* passing the lines to procline().
procfile(const char *fn)
{ …
f = grep_open(fn);

/* Process the file line-by-line */
if ((t = procline(&ln, f->binary)) == 0 && Bflag > 0) {

Where grep reads
grep_open is in file.c. grep_open calls grep_refill:

* Opens a file for processing.
struct file *
grep_open(const char *path)

/* Fill read buffer, also catches errors early */
if (bufrem == 0 && grep_refill(f) != 0)

grep_refill does a read into a buffer:

grep_refill(struct file *f)
{ …
nr = read(f->fd, buffer, MAXBUFSIZ);
MAXBUFSIZ is defined at the top of this file:

#define MAXBUFSIZ (32 * 1024)
and is the exact size of the reads we were seeing. Hooray!

To recap, here’s the control flow:

main -> procfile -> grep_open -> grep_refill -> read(…, MAXBUFSIZ)

Where grep writes
From before, recall that procfile, called procline. procline eventually calls printline:

* Processes a line comparing it with the specified patterns. Each pattern
* is looped to be compared along with the full string, saving each and every
* match, which is necessary to colorize the output and to count the
* matches. The matching lines are passed to printline() to display the
* appropriate output.
static int
procline(struct str *l, int nottext)

printline(l, ‘:’, matches, m);
which, if you scroll all the way down, on the penulitimate line, calls fwrite:

* Prints a matching line according to the command line options.
printline(struct str *line, int sep, regmatch_t *matches, int m)

fwrite(line->dat, line->len, 1, stdout);
So where cat was doing an honest-to-god write, grep goes through fwrite. So whatever buffer size we’re using is somehow tied up in the fwrite source.

To summarize, here’s the control flow:

main -> procfile -> procline -> printline -> fwrite

fwrite buffer size
Figuring out where fwrite sets it’s buffer size was more challenging- props here to Mehmet for his spelunking skills. Here I’ll show you the relevant control flow up front:

fwrite –> __sfvwrite –> prep_write –> __swsetup –> smakebuf –> __swhatbuf

I won’t walk you through all those layers of control flow but you can definitely check it out yourself if you want; starting at fwrite and ending at __swhatbuff.

Interesting Side-note! The implementation of makebuf.c in the latest apple OS (High Sierra) is rather different than in the implementation in an older version. Perhaps because Apple redid their file system for High Sierra? I know this because I went through an existential crisis (is everything I know a lie?) trying to figure out why the behavior I was seeing was inconsistent with the source code before I figured out it was the wrong version!

Okay, anyways, let’s take a look at the __swhatbuff that’s in High Sierra, because that’s what I’m running:

# define MAXBUFSIZE (1 << 12)
# define MAXBUFSIZE (1 <_file _file, &st) < 0) {
*couldbetty = 0;
*bufsize = BUFSIZ;
return (__SNPT);

/* could be a tty iff it is a character device */
*couldbetty = (st.st_mode & S_IFMT) == S_IFCHR;
if (st.st_blksize _blksize = *bufsize = st.st_blksize > MAXBUFSIZE ? MAXBUFSIZE : st.st_blksize;
return ((st.st_mode & S_IFMT) == S_IFREG && fp->_seek == __sseek ?
__SOPT : __SNPT);

What I think is going on here is that we’re setting *bufsize which is being “returned” to the caller through the pointer. The first two “ifs” don’t apply to us- we’re not dealing with an empty file, and we’re not a character device. So in the final block of code, we set *bufsize to 2 different values depending on whether st.st_blksize > MAXBUFSIZE, either to


or to


And actually, the top of the file says MAXBUFSIZE is 1 << 16 (which is 64KB) so we’re golden. For fun we can check that the st_blksize is larger than 64KB to be sure that’s the branch we’re taking. A quick c script that looks a lot like __swhatbuf will do the trick:


int main()
struct stat st;
FILE *fp;
fp = fopen(“test.txt”, “r”);
fstat(fp->_file, &st);
printf(“st_blksize was: %i\n”, st.st_blksize);
return 0;
Compiling and running it:

$ gcc deleteme.c & ./a.out
st_blksize was: 4194304
Sweet, so st_blksize is 4MB which is definitely bigger than 64KB. So that’s why that’s the transition point between terminating & going into an infinite loop! 🎉

1) Weird things happen when you read from files you write from.

2) The behaviors we saw had to do with buffering of the write buffer.

3) It’s rad that you can read open-source code to figure out exactly what’s going on. Grep your way to freedom!

I was looking at both the freebsd & apple’s open source implementations of fwrite and they were almost identical… except that freebsd’s version has both an integer overflow check & a lock. It’s a little alarming that apple’s version is missing those…it’d be interesting to see if you can get some weird behavior to trigger.

Thanks also to Dan Luu & Vaibhav Sagar for their feedback on a draft, and to Tobin Yehle for both proofreading a draft and figuring out I was looking at the wrong version of the makebuf.c file!
Github & Bitbucket

Charge Angels

Charge Angels

The EV Ecosysem

The First French Smart Grid Network

The Application Architecture

Our Solution

Our Partners

The Team
SAP Labs Mougins Runs See it live At SAP Labs Mougins* *not compatible with Internet Explorer

Empower charging capacity
Chargers are used more efficiently at their maximum capacity

Sharing Energy accross the Grid
Central Monitoring Software for all Energy Sources

Promote renewable energy
Solar, Wind Energy must be stored on premise and firstly proposed to the user

Digitalisation of Energy: A Vision Becoming Today A Reality

Our Application is interoperable with different types of chargers and brands (OCPP 1.6 compliant)

ChargeAngels Status
Display the realtime status of the evse landscape (available, occupied, charging, failure), handle user access

ChargeAngels Energy
Real time Display of Power Consumption and Total Charging Energy. Charging Profiles, Power limitation and Load balancing with the Grid

ChargeAngels Booking
Centralized user booking, BookNow@ Button by end of charge prediction, Charging car profiles, Users preferences

ChargeAngels User Experience
Realtime notifications for users to inform status of charge (Big data predictive Algorithms). Operations notifications for carfleet Administrators

News: Alpes-Maritimes Climate-Energy Trophée 2017
Your company wants to be part of the Electromobility Ecosystem ? Join The Smart Grid Club NCA
Schneider EVLink 22Kw EVSE
OCPP Standard
ADEME Sophia – FabMob

© – 2017 – All rights reserved (INPI Branded) – SIRET: 821 561 230 00011Design: HTML5 UP

It’s Time to Stop Buying Phones from OnePlus

It’s Time to Stop Buying Phones from OnePlus

OnePlus hasn’t been making Android phones for a very long time, but over its four years in existence it has made numerous screw ups. This all finally led up to last week’s massive credit card breach. It’s time we stop trusting a company with such clearly low standards.

Who Is OnePlus?
Founded just a hair over four years ago—at the end of 2013—OnePlus is a subsidiary of Chinese electronics manufacturer Oppo. The initial idea for the company was admirable: to make smartphones with high end components and features without the high end price tag. That’s something we can all get behind, and OnePlus was received with much fanfare in the Android community.

And you know what? For the most part, OnePlus makes good smartphones. They’re packed with excellent hardware and review exceptionally well. But the problem isn’t with the phones—it’s the company itself.

April 2014: Smash the Past
The first OnePlus phone—the OnePlus One—was released in April of 2014 by invite-only, just four months after the company’s formation. That’s a good turnaround time, but this also marks the beginning of a long list of genuinely questionable decisions by the company.

In an attempt to get people excited for their first smartphone, OnePlus launched a very poorly thought-out contest called “Smash the Past,” where they wanted users to smash their current smartphones. On video. Nothing about that sounds like a good idea out of the gate, but it gets worse.

Here’s how the promotion was supposed to work: users would apply to be a participant in the program, telling the company how they’d smash their phones. If chosen, they would then smash their phone in the described manner, on camera. Then they could buy a OnePlus One for one dollar.

The list of phones was also limited—they didn’t want you smashing some junker in order to get a OnePlus One for a buck. Nope, it had to be a high-end smartphone for the time, like an iPhone 5, Samsung Galaxy Note 3, Nexus 5, or Moto X. Sound crazy yet?

Here’s what actually happened: people misunderstood, because of course they did. People smashed their phones on camera not only before they were chosen to do so, but also before the contest even started. Now, this could be attributed to people just not reading or understanding, but the whole thing could’ve been avoided had OnePlus not launched such a dumbass campaign in the first place.

Why Do Lithium-Ion Batteries Explode?
Furthermore, even if you give them the benefit of the doubt on that, telling people to smash their phones is a terrible idea in the first place. Phones are full of chemicals that are safe when housed within their intended enclosure, but prone to explode when put in certain conditions…like being smashed to hell. This wasn’t just dumb, it was flat-out dangerous.

It’s been almost four years since that campaign, and I still can’t believe it was real.

Due to a high level of backlash, the company did try to make it better by later allowing winners of the contest to donate their old handset instead of smashing it, which is a lot better. Why didn’t they just do that in the first place?

You’d think that the company would’ve learned its lesson after that. Nah.

April/May 2014: The Invite System
The OnePlus One was originally only available through an invite-only system. This was put in place to combat the limited quantities available at release, but it just made things needlessly complex.

In order to get an invite, customers had to jump through a series of hoops with things like contests. All just to buy the damn phone in the first place. Once they bought a phone, those customers were given a limited number of invites to hand out to friends and family. If you were one of the lucky ones to get an invite, you had 24 hours to use it. After that, it was gone. God forbid you were out of the house that day, or on a family vacation.

I mean, I get stock shortages and whatnot, but using an invite system to “allow” people to give you money is arrogant at best—especially for a company with zero track record—and totally moronic at worst. I’m not suggesting that the current system of “first come, first served” is the best way to handle limited handset quantities either, but making customers work for the chance to buy a phone was confusing and convoluted. When someone has to make a YouTube video explaining how to buy a product, you’re doing something wrong.

May 2014: Yellow Screens and Warranty Issues
Shortly after the OnePlus One was released, some users complained of screen yellowing. The company claimed that a yellow screen was “not a quality issue and not covered under warranty.” Given the recent Pixel 2 XL screen issues and similar response from Google, I won’t point this out as an uncommon reaction. That said, Google did ease users’ minds by increasing the Pixel 2 warranty to an unprecedented two years, so at least there was a response.

I think the users who had this issue had every right to be upset, and the company should’ve done more to address it. The bottom line here is this: if there’s a clear issue with a display as soon as the phone comes out of the box, it needs be investigated at the very least.

August 2014: The Ladies First Campaign
In what might be the shortest contest in OnePlus history, an ill-conceived campaign just for women called “Ladies First” was launched in August 2014.

The idea was this: women (and women only) would draw the OnePlus logo on their body or a piece of paper, then take a picture of themselves and upload it to the OnePlus forum. The pictures would then get upvoted by forum members—mostly men—and the top 50 got a free OnePlus t-shirt and an invite to buy the phone.

Holy shit.

If you thought the “Smash the Past” campaign was stupid, it only took them four months to make an even worse one. Objectifying and sexist, Ladies First received a massive backlash as soon as it was announced—so much that the company cancelled it within hours. The company then claimed it was a “very misguided effort by a few isolated employees.” Right.

November 2014 – April 2015: Cyanogen Woes
Out the box, the OnePlus One ran Cyanogen OS—a custom fork of Android that was once the king of the custom ROM world. At the time, the founders of CyanogenMod had taken the ROM’s fame and tried to monetize it by building a company around it: Cyanogen, Inc. This company was responsible for building the operating system that would run on the One.

This is one of the things that made the OnePlus One so appealing to Android die-hards. This also later turned into a nightmare for OnePlus.

While no fault of OnePlus itself, Cyanogen, Inc. told the company in November of 2014 that it had entered into an exclusive deal with Micromax to produce the operating system for an upcoming line of phones for release in India. This came at a bad time for OnePlus, as it caused a temporary ban of import and sales of the the OnePlus One in India. The ban was lifted a week later as OnePlus decided to release its own custom version of Android called Oxygen OS.

This was the beginning of the end for OnePlus and Cyanogen, but the relationship took six months or so to really come to an end. The two companies reportedly butted heads a lot, with Cyanogen claiming in the end that OnePlus used the Cyanogen name to garner popularity—which may or may not be true, but I’m sure it didn’t hurt OnePlus’ recognition early on—saying it “built their brand on the back of Cyanogen.” Regardless of how you feel about it, that’s a bad look for any company—including both involved here.

The breakup between the two companies was very public and very ugly. It ultimately ended with the companies ending their relationship, and OnePlus using its OxygenOS moving forward.

That’s pretty rough for users who bought the phone in part because it was running Cyanogen OS, only for the partnership to end a year later. In comparison to Cyanogen, OxygenOS was less customizable, but still retained a lot of the stock Android look and feel users had grown to love about the handset’s former OS.

This was the end to a very rocky first year for a new company. Most other small companies wouldn’t have survived all the backlash and adversity thrown its way in such a short period of time, but OnePlus somehow prevailed.

August 2015: OnePlus 2 and More Invite System Garbage
Like most companies with a reasonably successful smartphone, the company followed up the first One with…the cleverly-named OnePlus 2.

Despite the company carrying the tagline “Never Settle,” the OnePlus 2 was released without NFC (near field communication)—a feature that was considered a staple for flagship phones of the day—and without wireless charging. This caused backlash in the Android community, though OnePlus claimed not enough OnePlus owners used NFC to justify its inclusion.

Also like the launch of the One, the 2 was released with an invite-only buying system. While they didn’t have any stupid contests like with the One, the 2 had its own share of issues where the invite and sale system is concerned—mostly that users couldn’t really buy the phone.

Initially, OnePlus promised a “new and improved” invite system, including 30-50 times more invites than were available with the OnePlus One. The thing is, it didn’t play out like that. North American orders were delayed by 2-3 weeks, and they also found issues with the materials used on the USB cables not being up to snuff. As a result, they had to slow the rollout of invites to “closely monitor and act on user feedback.”

So after once again failing to deliver, OnePlus CEO Carl Pei posted an apology in the OnePlus forums about how the company “messed up” the launch, noting that it took them a month after their target date to start shipping phones in “meaningful quantities.”

This company is basically a series of poor decisions and subsequent apologies…with some smartphones mixed in.

November 2015: OnePlus Sells Garbage USB-C Cables

Watch Out: How to Buy a USB Type-C Cable That Won’t Damage Your Devices
OnePlus ventured out from just smartphones, offering accessories like USB-C cables and USB-C to Micro-USB adapters—both of which were confirmed to be out of compliance with USB-C standards by trusted Google engineer and USB-C expert Bensen Leung. In short, using the cable or adapter had the potential to fry the power source thanks to the dodgy resistors in their construction.

Once again, OnePlus found itself apologizing for the mistake and offering refunds—but only to customers who purchased the USB-C cable, not the USB-C to Micro-USB adapter (which was equally as bad as the cable). It also noted that the cable and adapter were both safe to use with the OnePlus 2…just not other phones. Talk about a recipe for disaster.

June 2016: IMEI Data Being Sent to OnePlus Servers Over an Unencrypted Connection
On pretty much any Android phone, when you check for operating system updates the phone contacts the manufacturers servers to see if there’s new software available. Pretty common.

On the OnePlus 3, however, the phone was also sending the IMEI—that’s a numeric value that uniquely identifies that exact phone—over an unencrypted connection. That means a value that can connect your phone to your person was being sent over an open connection to OnePlus’ servers.

To make this even more interesting, it was also discovered that a proper IMEI wasn’t even necessary in order for the device to receive an update package. To test this, one OnePlus forum user sent a test request to OnePlus’ update server with an unusable IMEI, and an update package was returned.

It’s worth mentioning that this isn’t a significant issue on its own—just another questionable decision.

January 2017: OnePlus Gets Caught Cheating on Benchmarks
Benchmark scores used to be a hot topic on Android, so the better number a phone could produce, the better that phone seemed to end users.

With that in mind, benchmark scores on the OnePlus 3t were manipulated to be higher than the actual performance would indicate. OnePlus was apparently targeting specific applications by name and pushing the CPU into a specific scaling mode to push the stores higher than they would normally have been.

It’s also worth mentioning that other manufacturers were found guilty of the same thing with that research, and manufacturers like Samsung, HTC, Sony, and LG were all found guilty of doing the same thing back in 2013. So it wasn’t a unique offence, but something that hadn’t really been an issue for several years.

June 2017: OnePlus Gets Caught Cheating on Benchmarks…Again
After being “caught” for cheating on benchmarks with the OnePlus 3t, one could assume it wouldn’t be an issue again. But it was, because OnePlus got busted for maximizing benchmark scores again with the OnePlus 5.

This time the scores were accused of being boosted by as much as 5%. There’s an incredibly detailed analysis and write-up on the subject over at XDA, so I’d recommend taking a look at that if you’re interested in the gory details.

June 2017: OnePlus 5’s Display Confirmed to be Installed Upside-Down
Users of the OnePlus 5 noticed some odd “jellying” when scrolling on the phone, but it was unclear why this was happening. It was soon discovered why—the screen was mounted upside-down. Deliberately.

Because it was upside down, the screen refreshed from bottom to top (instead of the top to bottom), causing some interesting issues when scrolling. It didn’t appear to affect all units, but it was pretty obvious on the ones it did.

You may be wondering why the display was intentionally placed upside-down, and for that I turn to XDA for some well-researched speculation:

If you’ll take a look back at any of the thorough teardowns of the smartphone, you may notice that the display controller IC is located at the bottom. In order to compensate for the module’s placement, OnePlus flipped the display panel so the display cable would easily reach the motherboard and none of these components would interfere with other elements at the top of the device. But why would they need to do this all in the first place?

Take a look at what is placed at the top of the smartphone – the dual camera and some antennas. Like with any decision involving where to place components in a smartphone, it likely came down to space considerations. With limited space, the company had to decide where to place each component so everything would fit. Since the dual lens camera, which is new to the OnePlus line-up, takes up more space than a single lens camera it is possible the company moved the motherboard – and hence flipped the display panel – in order to accommodate the new camera module.

And there you have it.

July 2017: The OnePlus 2 Reaches End of Life Early
In June of 2014—after telling users that Nougat would be available for the OnePlus 2—OnePlus confirmed that the 2 would not get the Nougat update and had in fact reached its end of life at Marshmallow. It’s sadly common for Android phones to not get updates, but it was particularly crappy of OnePlus to promise one thing, and then renege.

July 2017: OnePlus 5 Devices Rebooted During 911 Calls
In 2017, a OnePlus 5 owner saw a building on fire, tried to call 911, and the phone rebooted. Twice.

How to Properly Test 911 Services on Your Cell Phone
It turns out the OnePlus 5 had a memory flaw that would cause it to reboot when emergency services were contacted, which is a massive issue (if that isn’t obvious). If there’s any time a cell phone needs to work, it’s during an emergency. Even phones without a SIM card are supposed to be able to make emergency calls.

Fortunately, the company rolled out a fix pretty quickly. But such an issue should not have existed in the first place.

October 2017: OnePlus Collects Private Data Without Approval
In October of 2017, it was revealed that OxygenOS was collecting data about device usage—a fairly common thing among smartphone manufacturers. But within this data was the device’s serial number, meaning an user’s identity could the be connected with said data.

The company claimed it was sending two separate sets of data—one for device usage and analytics data, the other with device information (serial number) for “better after-sales support.” It was also noted that all data was transmitted over HTTPS for security.

The thing is, this wasn’t really the question. The real issue here is that OnePlus was doing all this without user agreement—taking users’ data and shipping it back to the mothership without consent.

Just a few days after this data collection was revealed, OnePlus responded to the backlash by limiting the amount of data collected moving forward.

By the end of October, all OnePlus phones running OxygenOS will have a prompt in the setup wizard that asks users if they want to join our user experience program. The setup wizard will clearly indicate that the program collects usage analytics. In addition, we will include a terms of service agreement that further explains our analytics collection. We would also like to share we will no longer be collecting telephone numbers, MAC Addresses and WiFi information.

Like so many things in the past, this is a response to an action that shouldn’t have been an issue to begin with.

November 2017: Another, More Serious Security Issue
Just a month after the discovery of OnePlus collecting user data without approval, another vulnerability was found that allowed many OnePlus phones to be rooted without unlocking the bootloader, through a backdoor called EngineerMode.

The Case Against Root: Why Android Devices Don’t Come Rooted
OnePlus claimed the vulnerability wasn’t that big of a deal as it only worked with ADB, which requires USB debugging to be enabled in Developer Options (this is also disabled by default on Android devices). Security researchers at NowSecure looked deeper into the issue and provided a more in-depth explanation of its capabilities here. XDA also has a good write up of EngineerMode’s function and how this exploit worked here.

Essentially, an attacker would need physical access to the device to easily achieve root access and execute malicious code or commands, making this one of the less terrible vulnerabilities that we’ve seen.

At first, it was believed EngineerMode was a Qualcomm app, but after investigation, Qualcomm claimed it wasn’t theirs. Curious.

OnePlus quickly patched the vulnerability by removing EngineerMode.

January 2018: A Massive Credit Card Breach
In January 2018 OnePlus officially announced a massive breach in which 40,000 customers’ credit card info was stolen. The actual breach happened between November of 2017 and January 2018, at which point OnePlus finally figured out what was going on and halted credit card transactions.

OnePlus will be providing affected customers with a year of free credit monitoring, which is a paltry restitution. The damage here isn’t so easily remedied, and each user is going to have to deal with the repercussions of a stolen credit card.

Why Do We Keep Giving This Company Money?
That’s a long, long list of issues. They started off as poor decisions from a a young company—asking customers to destroy their phones or ladies to post selfies as part of a contest are both foolhardy at best, but not all that damning.

But then the issue continued to get worse. Selling USB-C cables that could literally damage (or destroy) user hardware and collecting user data without permission are both bad. Phones rebooting during 911 calls and backdoors that allowed easy root access by attackers are worse.

Having an issue where credit card information is stolen that stays open for over two months before it gets noticed? That’s pretty awful.

Look, I get why Android fans like OnePlus. They do release good hardware at great prices—and save for the past issues with data collection and whatnot, their software seems to be enjoyed by most who use it too.

Furthermore, I understand that any of these issues on their own isn’t the end of the world—in fact, some of the have happened to other companies we love and trust.

But taken all together, this is a pretty long list of problems, and in its four(ish) years of existence, OnePlus has shown time and time again that it doesn’t know what it’s doing and that it can’t be trusted. This company consistently shows its lack of responsibility towards customers—both potential and current. And yet people continue to fawn over them.

If you’re looking for good phones at a good price, there are other options out there. The Motorola Moto X4 is an excellent phone for just $400. The Essential PH-1 started off rocky, but a series of software updates and an attractive price drop makes it an excellent option that seems to keep getting better. It’s very comparable to the OnePlus 5t for the same price.

Maybe OnePlus will redeem themselves, but at this point, we can’t recommend buying from them until they clean up their act—and prove they can keep it that way long-term. For now, it’s time to stop trusting OnePlus with your personal information, your data, and your money. It’s time to stop settling.

Kwiggle: Das kompakteste Faltrad der Welt

Kwiggle: Das kompakteste Faltrad der Welt

Newsletter newsletter-subscribe-icon


Das kompakteste Faltrad der Welt
Stellen Sie sich ein Faltrad vor, das ca. 8,5 kg wiegt, in das Handgepäck eines Flugzeuges passt, in 10 Sekunden zusammengefaltet bzw. auseinandergeklappt werden kann und mit dem Sie ganz entspannt 20-25 km/h fahren können.


Kwiggle® ist das weltweit einzige Fahrrad, das diese einmalige Kombination aus Fahrtauglichkeit, Kompaktheit und Gewicht ermöglicht.

Kwiggle® ist ein idealer Begleiter auf Reisen, ideal geeignet für Berufspendler und so kompakt und leicht, dass es in jeder Wohnung und auch im Tragerucksack Platz findet.

Erst patentierte Antriebstechnologien, ein patentiertes Faltsystem und eine patentierte Sitzeinrichtung ermöglichen diese einzigartige Kompaktheit.

Kwiggle®: Die wohltuende Form des aufrechten Fahrradfahrens
Mit Kwiggle® fahren Sie erstmals aufrecht, fast stehend und doch angenehm abgestützt.
Kwiggle® fördert die Beweglichkeit des Rückens und vermeidet Schulter- und Nackenverspannungen. Fahrradfahren wird zu einer angenehmen Mischung aus Biken und Walken und kann sowohl gemütlich als auch fitnessfördernd sein.
Kwiggle® eignet sich für Outdoor-Fitness ebenso wie zur Rückentherapie.
Kwiggle®: Ein Fahrrad für alle
Egal, wie groß Sie sind. Dank der einstellbaren Sitzeinrichtung kann jeder auf demselben Kwiggle® fahren. Der besondere Pfiff Kwiggle® faltet sich trotzdem immer auf das gleiche Faltmaß zusammen.

Kwiggle®Bike GmbH • Altenbekener Damm 61 • 30173 Hannover • Tel. 0176 – 64771438 •

Kwiggle-Bike© 2016 Kwiggle Bike GmbH

Manta5 Hydrofoiler

Manta5 Hydrofoiler






Product Specifications
It took the best aeronautical, hydrodynamic, mechanical, electrical engineers to overcome the complexities of hydrofoils and propeller positioning to create more speed, better balance, and the ability to submerge launch- a world first.

Carbon Fibre Hydrofoils & Propellor

The knowledge gained from designing numerous combinations of foils refined and pioneered the new, smooth, easy to use hydrofoil combination. It strikes the perfect balance between stability and manoeuvrability in all types of water-terrain.

Hoerner Tip design to minimise drag and extend effective wingspan.
Easily assembled and disassembled with two Allen-key bolts into the bayonet mount and front tiller.
Rear wingspan measures 2m. Front wing, 1.2m
Nicknamed the “general purpose foils”, and can be used on a the widest range of ‘water-terrains’ E.g. Lakes, rivers, and the ocean.
Carbon fibre reinforced nylon propeller creates a perfect balance producing thrust at low speed and reducing drag at top speed.
Superior Lithium Battery Technology and 400W E-Bike Motor
We’ve worked with the best engineers all over the world to create our IP67 rated battery that’s designed to safely withstand the rugged marine environment. At the same time, we have maintained a user-friendly approach for removable recharging on any standard power outlet.

The 400W E-Bike motor has been waterproofed to an IP67 rating using world-first techniques. In addition, the motor has undergone rigorous destructive testing to guarantee reliability. Revered for its power to weight ratio this lightweight motor is also one of the quietest.

Cutting-edge waterproofing techniques ensure battery and motor longevity in both fresh and saltwater.
From one charge the battery offers an hour of pedal assistance for an 85kg rider in average conditions.
Battery housing indicators display riding modes, battery status and flight level – allowing riders to optimise their planing speed and ride time.
Detachable battery pack for easy recharge into any standard wall socket.
Brushless motor provides pedal assist
BMS protection against overcharging, overcurrent, excess voltage, deep discharge, short-circuit and overheating.

Bike Frame, Buoyancy Modules, Gearbox and Drivetrain
A modular design means riders have a cost-effective upgrade path for future sporting and leisure components. Should riders wish to customise their hydrofoilers they can select new foil combinations, buoyancy modules, motors, new propellers and battery packs as they enter the market.

The TIG welded 6061-T6 aluminium frame is an aircraft grade alloy revered for its strength and weight advantages.
Chain & shaft-drive hybrid with industry standard components for easy replacement.
Buoyancy modules designed to streamline the hydrofoiler both under and above water.
Proprietary gearboxes have been designed especially to keep the streamlined body profiles narrow to reduce drag.
The drivetrain is designed with user-serviceability in mind. The crankset sub-assembly, the gearbox sub-assembly, and the propeller/shaft sub-assembly can all be removed independently of each other allowing for easy maintenance.
Be first in line for pre-sales
To ensure we provide the best customer service Manta5 will only produce a limited number of Hydrofoilers for public release.
To be the first in line please register your interest today.


Our Journey
Privacy Policy
Sign up for exclusive event invites, demos and official pre-sale announcements.

email address