Any news regarding CUDA 5.5 support? ;)

Reasons for lacklustre performance with Cuda 5.5 direct builds were investigated and identified, prompting major redesign intended for deeper in x42 to be pulled forward, in front of other plans.

I'm glad that You (you people) identified them.

Would any of those include one of these:

0) What CUDA (NV) functios are concidered 'special'? There are a limited number of cores for special instructions/operations.

1) autonomous thread invocation from within a thread (sparse/dense/needed) and processing other threads during memmoves (that can be initiated by GPU too)

2) __sincosf/sqrtf/... calls in calc chirp data or any other place: is there a repeating pattern for the argument of the call (caching)? Do any/some of the threads/blocks invoked start at the beginning or in the middle of a repeating pattern (time goes on at specific intervals)?

3) Transpose - could that be avoided if the next call (and programmer) would know that the input is in a non transposed form. Could something be done twice so that every other call could be avoided and it would only be inside a function when some data is in an unorthodox form.

4) active Warps/blocks/threads/registers/TITAN/780/+

5) (CRAY_)CUDA_PROXY on linux and running multiple tasks

6) sometimes a small number of register spills does not hurt performance

7) More active striding (thread or blockId .x*prime)&31 of writes + __shared__ buffering

... I'm sorry. I know You know.
pre post edit: (this was written in a foreign language after a few beers).

---

To overcome Heisenbergs:
"You can't always get what you want / but if you try sometimes you just might find / you get what you need." -- Rolling Stones

I'm glad that You (you people) identified them.

Would any of those include one of these:

0) What CUDA (NV) functios are concidered 'special'? There are a limited number of cores for special instructions/operations.

Not really compute bound anywhere at the moment on slow to fast GPUs, but a definite goal. special instructions for gk110 come in that category.

1) autonomous thread invocation from within a thread (sparse/dense/needed) and processing other threads during memmoves (that can be initiated by GPU too)

I explored if this dynamic parallelism feature might be useful down the line. It's of possible benefit for further reducing host communications down the line, though other bottlenecks prevent advantage there for the moment.

2) __sincosf/sqrtf/... calls in calc chirp data or any other place: is there a repeating pattern for the argument of the call (caching)? Do any/some of the threads/blocks invoked start at the beginning or in the middle of a repeating pattern (time goes on at specific intervals)?

While the nVidia original implementation used single elements per thread, my optimised versions I made scale with multiple elements per thread and so paid attention to strides etc. fortunately that part ws simple & scaled without mods to new cards from Fermi onwards. Fixing some precision in nvidia's implementation required some research& design, to emulate double precision math better &faster for this area, on devices that don't have it.

3) Transpose - could that be avoided if the next call (and programmer) would know that the input is in a non transposed form. Could something be done twice so that every other call could be avoided and it would only be inside a function when some data is in an unorthodox form.

The nVidia implementation avoids transpose, with a comment there saying they didn't need it. It turns out this was a bad design choice, that persists up to now. nVidia GPUs will experience bank or segment conflicts at various cache or memory levels if power of 2 strides are accessed. That's the dominant cause of the VLAR problem (vertical strided random accesses in the pulsefinding)

4) active Warps/blocks/threads/registers/TITAN/780/+

The original code used fixed sizes arranged for G80 architecture. I've made most of the important ones automatically scale. This occupancy/utilisation though is dominated by the other transfer & performance issues, induced by system driver latencies. Latency hiding techniques are being explored with the aim of becoming compute bound again.

5) (CRAY_)CUDA_PROXY on linux and running multiple tasks

I've been looking at various possibilities, given that I want heterogeneous clustering for x42 'superdreadnaught'. Several MPI based implementations look like good starting points, but I require fault tolerance. Ideally I'd like pipeline or a number of other process or lightweight thread mapping capability, to include outriggers in complex topologies. Future massive GBT AP's possible, might have huge autocorrelations that could benefit from say a hypercube arrangement of virtual nodes at process, thread, host or some mix on different devices across a network. A similar consideration is being give to different granularities, such as if tasks stayed small& you want multiple per node.

6) sometimes a small number of register spills does not hurt performance

Should not do no, since most of the processing is either global memory or bus bound at the moment, though the 50% compute threshold goal is looming near already, even prior to extensive latency hiding. When it passes in to compute bound territory, then hot kernels will need looking at with those instruction level choices in mind.

7) More active striding (thread or blockId .x*prime)&31 of writes + __shared__ buffering

I due use some loads to shared mem buffers, though turns out, after a lot of experimentation & research, the Cuda best practices guide gave(gives) incomplete pictures. Shared memory is pretty slow compared to registers, and has some bank conflict issues that are tricky to deal with. volkov detailed much of this in a presentation to nVidia, then his work went toward CUFFT 2.3+ and CUBLAS. I tend to go for low occupancy max bandwidth now for the lightweight kernels, though combining cascade kernels together later I will probably revisit this once compute density is high enough. In fact GK110 has some new instructions that replace the need for shared memory use in special cases that do apply to us in a number of key areas (reductions mostly). I've yet to use those, and they should replace some pretty large chunks of my hand optimised kernels with hardware instructions.

Pulsefinding (transposing for horizontal access) & host transfers (driver latency induced, particularly WDDM) are the big targets for now. These combined account for around 40% of elapsed time at mid to high angle ranges, and higher than 70-80% at VLAR, mostly pulsefinding issues there.

Jason

---

"Living by the wisdom of computer science doesn't sound so bad after all. And unlike most advice, it's backed up by proofs." -- Algorithms to live by: The computer science of human decisions.

Thank You Jason,

We're all expecting now.
You'll deliver a full matured ba.. product.

I hope it'll be an ET. (Extreme Throughput)

--
Nothing to say here.

---

To overcome Heisenbergs:
"You can't always get what you want / but if you try sometimes you just might find / you get what you need." -- Rolling Stones

A short reminder, the latest Nvidia driver version 331.58 shows even CUDA 6.0!

22/10/2013 22:51:04 |  | CUDA: NVIDIA GPU 0: GeForce GT 640 (driver version 331.58, CUDA version 6.0, compute capability 3.0, 2048MB, 2017MB available, 692 GFLOPS peak)
22/10/2013 22:51:04 |  | CUDA: NVIDIA GPU 1: GeForce GT 430 (driver version 331.58, CUDA version 6.0, compute capability 2.1, 512MB, 498MB available, 288 GFLOPS peak)
22/10/2013 22:51:04 |  | OpenCL: NVIDIA GPU 0: GeForce GT 640 (driver version 331.58, device version OpenCL 1.1 CUDA, 2048MB, 2017MB available, 692 GFLOPS peak)
22/10/2013 22:51:04 |  | OpenCL: NVIDIA GPU 1: GeForce GT 430 (driver version 331.58, device version OpenCL 1.1 CUDA, 512MB, 498MB available, 288 GFLOPS peak)

Maybe this is a little stimulus? ;)

---

Aloha, Uli

LoL, no signs of a Cuda 6.0 tools yet ;)

---

"Living by the wisdom of computer science doesn't sound so bad after all. And unlike most advice, it's backed up by proofs." -- Algorithms to live by: The computer science of human decisions.

I mentioned it to Jason a week ago, at present he has no idea what the changes are, let alone having the Cuda 6 SDK.

Edit: beaten by the man himself.

Claggy

A short reminder, the latest Nvidia driver version 331.58 shows even CUDA 6.0!

22/10/2013 22:51:04 |  | CUDA: NVIDIA GPU 0: GeForce GT 640 (driver version 331.58, CUDA version 6.0, compute capability 3.0, 2048MB, 2017MB available, 692 GFLOPS peak)
22/10/2013 22:51:04 |  | CUDA: NVIDIA GPU 1: GeForce GT 430 (driver version 331.58, CUDA version 6.0, compute capability 2.1, 512MB, 498MB available, 288 GFLOPS peak)
22/10/2013 22:51:04 |  | OpenCL: NVIDIA GPU 0: GeForce GT 640 (driver version 331.58, device version OpenCL 1.1 CUDA, 2048MB, 2017MB available, 692 GFLOPS peak)
22/10/2013 22:51:04 |  | OpenCL: NVIDIA GPU 1: GeForce GT 430 (driver version 331.58, device version OpenCL 1.1 CUDA, 512MB, 498MB available, 288 GFLOPS peak)

Maybe this is a little stimulus? ;)

It showed up in the 331.40 beta driver.

---

http://www.anandtech.com/show/7515/nvidia-announces-cuda-6-unified-memory-for-cuda

Nice of them to leak press info before authorised closed beta even get the download.

---

"Living by the wisdom of computer science doesn't sound so bad after all. And unlike most advice, it's backed up by proofs." -- Algorithms to live by: The computer science of human decisions.

Now to be clear here, CUDA 6’s unified memory system doesn’t resolve the technical limitations that require memory copies – specifically, the limited bandwidth and latency of PCIe – rather it’s a change in who’s doing the memory management. Data still needs to be copied to the GPU to be operated upon, but whereas CUDA 5 required explicit memory operations (higher level toolkits built on top of CUDA withstanding) CUDA 6 offers the ability to have CUDA do it instead, freeing the programmer from the task.

Very doubtful advantage IMHO. If one wants higher level of abstraction one could use some of C++ libraries over CUDA that do such abstraction of actual memory resourses. AMP, for example.

"Zero copy" operation was the real advantage cause it eleminated some of data transfers. This one doesn't eliminate data transfers, it just hides them from programmer (and adds possibility of inefficiency, driver dependent).

---

SETI apps news
We're not gonna fight them. We're gonna transcend them.