What I'm expecting from GTC..

Well really I think I'm expecting to much altough in form of a lot of minor improvements in his software products (so I'm no expecting new architectures info (Maxwell) etc..), but anyway I have compiled a list of things so I can check later wheter NV is doing his work or not :-)
Of course it will be good even if all these pieces come in to place say over H1 2013..

*nvfx: new effects system open source, cross vendor support, etc.. was anounced at Siggraph and has an empty github site and also there is a talk at GTC so there is no better place and moment to upload to github.. This system also uses a more efficient OGL state management ext called NV_state_object better aligned to DX10-11 state managements via objects so seems also more like DSA management..
*Only consumer HW info may be GK114/6 archs info which may bring some new things as note even Titan has no DX11.1 profile support so hoping before Maxwell support says GK114 and such new 680 replacements must have it so there is some minor arch enhancement in graphics side..
Also can have one more thing.. see next point.. well say it briefly dynamic parallelism everywhere and from anywhere to anywhere.. (ANYWHERE={CUDA,OGL})
*new ogl exts: NV_state_object (DX11 like state objects) and some kind of dynamic parallelism for graphics APIs..
Regarding  this is interesting there as there is a patent on it and it's about exposing dynamic parallelism in graphics world which implies OGL in near future i.e. graphics shaders can create new draw calls and put on the dispatch manager queue..
Also for completeness what's holding NV from exposing launch graphics from compute kernels and dispatch compute kernels from graphics shaders.. Note seems NV_state_object is much needed in two cases (CUDA->OGL and OGL->OGL dispatch draw call cases) as some state env is needed in these cases CPU apis are not useful as it's GPU work and default OGL state may not be useful..

Also please upload documentation on NV_GPU_shader5_memory_extended  shipping in 313 drivers altough I suspect is for exposing cache modifiers to load store operations supported on CUDA already like load non cached,load cached, etc..
Note my previous post asking for NV to expose all compute functionality (ISA richness in this case) to OGL compute shaders via at least now lacking PTX ISA instructions and also via some asm() function (which is reserved already in GLSL and usable in OCL kernels in NV and even AMD!(this is new for me I found last month and you can use AMDIL altough I haven't been able in exposing clock cycle counter to work yet))
*Grid SDK: well I'm interested in frame capture APIs not cloud stuff.. related I see OGL support for NVENC is being implemented so some update for NVENC will be good..
*OpenGL SDK: well one seems overdue (exposing advanced usage cases of OGL 4.x features) and a tess sample was released soon this year.. One deferred+ sample would be good..
*Cg 3.2: I want  glsl 4.3 support integrated into Cg for some things I'm working and Cg 3.1 is almost one year old.. also I think if support for cg compute shader is or not implemented ( as said in Cg language/runtime) will say much of wheter Cg is dead or not.. Also what about bindless texes in Cg?
*cuda 5.1 I suggested to NV team in late October equaling CUDA to OGL compute shader so support for compressed texes, depth textures, msaa textures (even depth ones..).. Note some of this are in OCL 1.2 exts release in SA 2012.. And also expose similar functionality to all remaining OCL 1.2 new exts  in case support avaiable in HW or easy to do by runtime like terminate kernel, out of bounds stuff, memory initialization etc..
One thing that I forgot at the time:
Expose atomic counters (now are shipping on OGL compute world) on CUDA and OCL (like AMD does on OCL) this are equivalent to atomadd(ptr,1) but an order of magnitude faster than global atomics at least on Fermi (not know in Kepler) and they are the foundation of "hardware accelerated queues" not? I remember how when NV readied OGL 4.2 beta drivers atomic counters were slow and then after some month or so they get tremendous speedup and they deserved special instruction exposed in NV OGL assembly language..
*cuda compiler sdk seems is going final and I think will bring up to date to CUDA 5.1 or 6.0 whatever they may end naming new CUDA release. (hope also gets up to date LLVM/Clang integration so 3.2 and/or 3.3)
*cuda.lang: Well I want to play with these for a long time.. motivation well bring more an offline compilation model to CUDA like OpenCL and basically avoid needing in Windows VS installed for realtime compilation of CUDA kernels: could be useful to dynamic compilation of Optix shaders (like OpenRL) and also for research software of nvidia like CUDAraster, VoxelPipe, etc..
*Shipping all Physx stuff from last GTC and GDC into production:
->apex 1.3 (bring realtime fracture support done entirely on gpu to existing RGB support)
->physx 3.3 (rigid bodies on gpu and perhaps even fracture like APEX)
Hope at least by GDC which is later we will get all of these in beta form..
One anoying thing for me at least is that Physx GPU interop with graphics APIs isn't avaiable (altough yes in APEX).. which anoys me is that APEX is Physx under the hood so please also expose GPU buffers of result simulation of GPU modules like cloth, fluid, and soon rigid bodies..
*optix 3.1 preview-> bring some gk110 perf improvements.. seems current Optix doesn't exercise all potential judging from perf numbers on Nv forums (barely better than GTX 680?)..
*cuda roadmap nda discussion: Well it was anyway a surprise to see NV invited me to a NDA discussion of future roadmap (hope saying it isn't NDA :-)) at GTC.. I can't attend but I hope they will be talking about how to expose unified CPU/GPU in CUDA and potentially new ISA sm_40?
*volume render solution
*ocl 1.2 in drivers: well with OCL 2.0 spec coming perhaps at Siggraph it's time to implement OCL 1.2 in NV drivers? In time with new CUDA support?
*nsgiht 3.0 final and 3.1 preview: After GLSL native debug I want (really more than I need right now but anyways soon will need..) (VS2012 support, OGL 4.3 support with compute shaders, and my biggest desire is for a unified host and device debugging experience like that ships in Nsight Eclipse edition)
For Eclipse edition I hope they add single GPU debugging with software preemption much like her older brother and also OGL debugging with that basically GPU debugging is perfect for me on Windows and Linux and all that remains is GPU true software preemption..

Read More

Exposing advanced PTX ISA in OGL compute..

Well that is an extended answer to Timothy post in http://timothylottes.blogspot.com.es/2012/12/storing-objects-on-gpu.html ..

Hi Tim,
well there is a lot of compute stuff not present currently on OGL compute shaders.. you have really opened Pandora's box here..
well seriously..
first please apologize for extension of comment but hope that's motivates Nvidia OGL team to implement "simple" way of exposing "all" current ISA functionality compared to PTX level...
also before you get tired and bored of reading I would like to ask how you to confirm you measure graphics->compute->graphics context switch is simply something as Draw(); DispatchCompute();
Draw(); with empty computekernel as you say.. also should something as DrawIndirect(); DispatchIndirect(); DrawIndirect(); improve/lower timings of contex switching time?..
As you say some native instructions aren't present and also some special registers (in PTX parlance chapter 9 at ptx manual) but also some CUDA runtime level functionality..
I'm going to expose now is use cases of low level PTX isa instructions and special registers are used presently in very low-level optimized GPGPU codes most originated from Nvidia research guys..
After all that examples my conclusion is IMHO, as Nvidia currently supports assembly level shaders (NV_gpu_shader5 NV_gpu_compute_shader5) and even CUDA doesn't take
effort to expose all things from PTX world to "high level" CUDA C kernels NV should follow CUDA approach and should do two things:
1. expose all PTX insts+special registers currently not exposed to NV assembly shaders in it.. may be GL_NV_compute_program5_extended or something like these..
the idea here is it's that overkill to do something similar to NV_shader_atomic_float extension for exposing atomic floats (altough this extension does more and defines
GLSL functions in addition to ISA ADD.F32 variant in shader assembly:
float atomicAdd(float *address, float data);
but this limitation can also be avoided by adding only another extension..
2. expose as motivated below a asm() function in OGL shaders (like NV_gpu_inline_asm) that work exactly like currently supported CUDA kernels (and even OpenCL! since 4.x(1?) drivers)..
   note that's the way that OpenCL is able to exploit as much as CUDA new GPU cores ISA enhancements..
   For OpenCL see inlinePTX kernel code ->asm("mov.u32 %0, %%laneid;" : "=r"(laneid));
with that as said NV_shader_atomic_float extension should be not needed using:
#pragma extension NV_gpu_inline_asm
float atomicAdd(float *address, float data)
{
 { float v; asm("ATOM.ADD.F32 %0, %1, %2;" : "=r"(v) : "r"(adress), "r"(data)); return v; }
}
 some low level optimized CUDA codes that come to my mind are from NV reesearch as cudaraster, raytracing,and also  http://code.google.com/p/cudadma/ these are examples that show
 that possiblity of making some optimized ports of these codes is tied to accessing being able to access these low level ins as you found on MGPUsort..

 *Cuda dma:
 well this code uses generalized named barriers which allow also for not locking between diferent barriers (SM20 and up) and think counting threads passed and waiting at
 barrier(syncthreads_count()).. see bar in PTX manual, bar.arrive and example PTX code..
 also project home now shows GK110 instruction usage->
 " The new instances will make use of the ldg intrinsic for supporting many more outstanding loads in flight by issuing them through the texture cache"
 *cudaraster/raytracing/voxelpipe NVR projects (please note these are specially interesting as they are related to forward looking graphics/compute integration and implementation in OGL compute should allow tighter integration
and possibly better improvements to graphics assuming someday hardware queues between stages get exposed in OpenGL):
 we find code like these:
 __declspec(__device__) __inline U32   getLaneMaskLt           (void)                  { U32 r; asm("mov.u32 %0, %lanemask_lt;" : "=r"(r)); return r; }
  ->latest version optimzed for kepler also uses new vector simd instructions (see hpg2012 poster)
  they seem to define also prmt:
__declspec(__device__) __inline U32   prmt     (U32 a, U32 b, U32 c)   { U32 v; asm("prmt.b32 %0, %1, %2, %3;" : "=r"(v) : "r"(a), "r"(b), "r"(c)); return v; }
__declspec(__device__) __inline U32   add_cc                  (U32 a, U32 b)          { U32 v; asm("add.cc.u32 %0, %1, %2;" : "=r"(v) : "r"(a), "r"(b)); return v; }
add with carry for bignum computations one thing I'm sure is that addc isn't exposed in variant where carry is stored in special carry flag bit (CC.CF)
but not sure if graphics shaders expose currently (think yes but carry must be saved as integer and then added so this sequence for 128 bit addition isn't as efficient)
add.cc.u32 x1,y1,z1; // extended-precision addition of
addc.cc.u32 x2,y2,z2; // two 128-bit values
addc.cc.u32 x3,y3,z3;
 addc.u32 x4,y4,z4;

Another interesting instruction used in some papers is bypassing or forcing caching loads per instruction etc(8.7.6.1 Cache Operators).. found in cudaraster also:
__declspec(__device__) __inline uint4 cachedLoad              (const uint4* p)        { uint4 v; asm("ld.global.ca.v4.u32 {%0, %1, %2, %3}, [%4];" : "=r"(v.x), "=r"(v.y), "=r"(v.z), "=r"(v.w) : "l"(p)); return v; }
__declspec(__device__) __inline void  cachedStore             (U32* p, U32 v)         { asm("st.global.wb.u32 [%0], %1;" :: "l"(p), "r"(v)); }
__declspec(__device__) __inline U32   uncachedLoad            (const U32* p)          { U32 v; asm("ld.global.cg.u32 %0, [%1];" : "=r"(v) : "l"(p)); return v; }
__declspec(__device__) __inline void  uncachedStore           (uint4* p, uint4 v)     { asm("st.global.cg.v4.u32 [%0], {%1, %2, %3, %4};" :: "l"(p), "r"(v.x), "r"(v.y), "r"(v.z), "r"(v.w)); }
Finally some codes use %clock and %clock64 variant register for true kernel cycles counting..
We only need to see PTX manual as said for me interesting lacking instructions are (sorry if something is already present):

mul24 (specially useful for thread index,sad(for video operations),bfind,popc, popcll (like clock and others we have here 64bit instruction which allows 2x faster)
The last one was shown also of good use in computer vision (object recognition) http://nvidia.fullviewmedia.com/gtc2010/0923-c-2209.html) sadly no slides only video but
new kepler shufl instructions which NV said improved perf of galaxy code shown in GTC..
and special registers:
%lanemask_eq, %lanemask_le, %lanemask_lt, %lanemask_ge, %lanemask_gt
%clock, %clock64
I'm not even asking about more general features in kernels features such as dynamic malloc, printf support, and true function call in kernels etc..
Well really function call and recursion stuff is a wonderful example of how really simple is to
add support for things by adding only assembly language extensions and using asm blocks without changing GLSL parser  :
asm(
".func foo { call foo; .. ret; } … call foo;")
and not even advanced CUDA runtime compute features such as:
*tuning L1 cache size vs shared mem size (16/48 32/32 48/16 currenlty),
*how to expose concurrent kernels (really usefull in sm_35) to OpenGL without stream model
*host pinned mem (use cpu mem from kernel on the fly without transfers) ->should enable out of core compute algorithms..
 (that's also interesting as AMD has one extension for it GL_AMD_pinned_memory which works even on discrete GPUs and searching while ago
 from nv found NV_sysmem_buffer but hacking to code to use it seems non functional)
*P2P functionality of accesing from kernels mem from other GPUs etc..
(note this is another interesting thing as AMD currently exposes that in OGL world GL_AMD_bus_addressable_memory but seems FirePro only so that would
be possible once GL 4.3 drivers ship also they enable that in OCL world this month with CL_AMD_bus_....) they call it DirectGMA..
Also lacking vs CUD/OpenCL but that should be addressed by OpenGL ARB as isn't NV specific
is allow to programatically exposing work group size like by requiring local_size_x variables not to be constant like (new extension like GL_NV_uniform_local_size)

uniform int local_size_x_from_app;
layout (local_size_x = local_size_x_from_app) in;

or better new function
DispatchComputeSetLocalSize(globx,globy,globz,locx ,locy,locz)
equivalent of CUDA support

and finally also some needed for serious compute similar to usage HPC programs is allowing unroll marks to compiler to unroll selected loops
#pragma unroll
--END of answer--

thah shows in some aspects OGL compute support exposed in 4.3 is really very simple somewhat like CUDA 1.0 and even worse in some aspects specially in regard to exposing ISA of current NV compute cores..
For example some CUDA codes use %clock to measure at cycle count of their kernels.. that's supposedly was shipping in CUDA 1.0 altough using PTX code.. today we can use asm() blocks to  define a clock function in CUDA C code and use that similar to CPU world using asm() blocks around TSC counter..


 

Read More

Some news!

News:
*Gpu computing gems 1 or GPU gems 4 source code already avaiable in gpucomputing.net:
Book for November..
Right now:

Title
A Programmable Graphics Pipeline in CUDA for Order Independent Transparency		1 new	07-10-2010
High Performance Iterated Function Systems		0 new	07-02-2010
CUDA Implementation of the Tree-based Barnes Hut n-Body Algorithm		0 new	07-01-2010
Connected Component Labeling in CUDA - demo+code		0 new	06-30-2010
A Practical Guide toMassively ParallelMonte Carlo Simulations: The Ising Model		0 new	06-30-2010
Parallel LDPC Decoding using CUDA		0 new	06-30-2010
Path Regeneration for Random Walks		0 new	06-30-2010
GPU Gems 4: Deformable Volumetric Registration using B-splines Source Code		0 new	06-30-2010
Monte Carlo Photon Transport on the GPU		0 new	06-30-2010
Lattice-Boltzmann Lighting Models - Source Code		0 new	06-30-2010
RNA folding GPU		0 new	06-30-2010
Haar Classifiers for Object Detection with CUDA: Pixel-parallel processing kernel		0 new	06-29-2010
Multiclass Support Vector Machine		0 new	06-29-2010
Parallelization of the x264 encoder using OpenCL		0 new	06-21-2010
Cone-Beam CT image reconstruction using the Katsevich Algorithm		0 new	06-21-2010
Line forward projection on CUDA		0 new	06-11-2010

seems MareNostrum getting a rack of Fermis perhaps with IBM Power7

see now Nvidia would have to publish a PowerPC arch CUDA driver?

Or using PathScale with full open source based computing stack..
avaiable here branch from noveau:

http://github.com/pathscale/pscnv/commits/master

Seems Nvidia TCC supporting driver Fermi in IBM web site version 197.81

Catalyst 10.8 beta seems avaiable 10.7 coming 21/7..


Physx 3.0 coming with CPU improvements:
*auto threading
*sse enabled by default
Mafia has new runtimes NVIDIA PhysX driver: 10.04.02_9.10.0522.
Mueller has post paper of Fermi launch demo using water heigh fields plus particles..
Two other papers interesting from Nvidia research are:

HLBVH: Hierarchical LBVH Construction for Real-Time Ray Tracing
PantaRay: Fast Ray-traced Occlusion Caching of Massive Scenes

Hwu based course from Stanford:
http://code.google.com/p/stanford-cs193g-sp2010/wiki/ClassSchedule

Two interesting conferences program avaiable:

PACT
has intel gpu paper demystifying ..
also Revisiting Sorting for GPGPU Stream Architectures
which achieves near 500mkeys/s on gt200..



there is a workshop on gpus
http://informatik.technikum-wien.at/gpusca/
and web doesn't work.

The Nineteenth International Conference on
Parallel Architectures and Compilation Techniques (PACT)
Vienna, Austria, September 11-15, 2010

Interesting papers:
Scalable Thread Scheduling and Global Power Management for Heterogeneous Many-Core Architectures
Dynamically Managed Multithreaded Reconfigurable Architectures for Chip Multiprocessors
WAYPOINT: Scaling Coherence to Thousand-core Architectures
Scalable Hardware Support for Conditional Parallelization
Less is More: Trading off Work-Efficiency for Scalability in Irregular Programs
Revisiting Sorting for GPGPU Stream Architectures
D. Merrill, A. Grimshaw
An Integer Programming Framework for Optimizing Shared Memory Use on GPUs
W. Ma, G. Agrawal
DMATiler: Revisiting Loop Tiling for Direct Memory Access
A Software-SVM-based Transactional Memory for Multicore Accelerator Architectures with Local Memory
Automatic Vector Instruction Selection for Dynamic Compilation
An OpenCL Framework for Heterogeneous Multicores with Local Memory

SC10

I would like to review this papers:
Scalable Tile Communication-Avoiding QR Factorization on Multicore Cluster Systems
Parallel Fast Gauss Transform
Overlapping Methods of All-to-All Communication and FFT Algorithms for Torus-Connected Massively Parallel Supercomputers
The Multi-Scale Heart Simulation on Massively Parallel Computers
Using 3.5-D Blocking Optimization for Stencil Computations on Modern CPUs and GPUs
An 80-Fold Speedup, 15.0 TFlops, Full GPU Acceleration of Non-Hydrostatic Weather Model ASUCA Production Code
Exploiting 162-Nanosecond End-to-End Communication Latency on Anton
Strider: Runtime Support for Optimizing Strided Data Accesses on Multi-Cores with Explicitly Managed Memories
Multithreaded Asynchronous Graph Traversal for In-Memory and Semi-External Memory
OpenMPC: Extended OpenMP Programming and Tuning for GPUs
Scalable Graph Exploration on Multicore Processors
The 48-core SCC processor: the programmer’s view
Exploring a Novel Gathering Method for Finite Element Codes on the Cell/B.E. Architecture
Reducing Multicore Bandwidth Requirements for Combinatorial Multigrid
Diagnosis, Tuning and Redesign for Multicore Performance: A Case Study of the Fast Multipole Method
Scaling Hierarchical N-Body Simulations on GPU Clusters
Size Matters: Space/Time Tradeoffs to Improve GPGPU Applications Performance
The Sharing Tracker: Using Ideas from Cache Coherence Hardware to Reduce Off-Chip Memory Traffic with Non-Coherent Caches

Read More

DirectCompute Double precision Mandelbrot demo and more..

In addition to first demo using double precision on GL 4.0 here now on DirectCompute:
THIS DEMO NEEDS DX JUNE 2010 RUNTIMES
so update if needed

this test on AMD shows a ATI DirectCompute DPFP bug.. it shows incorrect rendering..

Also note I learned DirectCompute doesn't admit division with doubles so I have to change /2 with *0.5.
Nvidia Fermi works OK!
DirectCompute Double precision Mandelbrot (includes source based almost 100% on Voxilla demo):

use test.bat app starts at big zoom so it shows DP in action.. if you exit with esc then shows same rendering at SPFP.. note with mouse you can zoom in out.. 

bat calls mandel.exe 0 for SP or mandel.exe 1 for DP..

Also note I expected better perf for AMD than Nvidia but two work very slow i.e. Nvidia runs at full speed (i.e. capped 8x vs Teslas) but AMD has perf issues as it should run at least 3-4x vs Nvidia Fermi..
Also has vector mode running somewhat faster than scalar shader (but not much could run up to 4x faster if compiler didn't extract perf of scalar code but runs not much faster compared to SP where vector code outperforms scalar code by a higher amount).. fermi perf is unaffected by using vector code..

Correct behavior:

Double precision (on GTX 470)

See full Window

Single precision

On AMD 5850 DP renders as (i will post image soon):

Related also I patched Nvidia Physx Demo to work on AMD changing GLSL code using Cg non standard functions.. it exhibits some OpenGL bugs.

Instructions:

Download Nvidia Physx Demo here ((select FLUIDS: TECHNOLOGY DEMO)

and use this exectuable for running on AMD cards (extract on demo dir).
It shows artifacts on AMD card not on rendering but on desktop outside of program window..
On AMD 5850 DP bad renders as (i will post image soon):

Read More

A lot of things you probably don't know.. and a worth it..

*TCC support for GF100 products will be out next week also this drivers will add support for simultaneously running this drivers with normal graphics drivers (that support OGL,DX,DXVA,etc..) I suspect graphics and TCC driver will have to have same version as both write dll's in windows system..
I hope still inf trick works so I can enable on Geforce Fermi and also that this works with Nsight also.. anyway is not severe as 25x drivers seems to add support for CUDA cards (Geforces even) without extending desktop on it so kernels exec time needn't be time limited for TDR.. before it required to use two Nvidia cards and one can be not desktop extended but if you used say a ATI card and a Nvidia card without desktop extended on Nvidia so to use Nsight for example (which requires no desktop extended) it will fail since CUDA will not find a CUDA card..
*There is support for Fermi on MacOs right now on Nvidia 19.5.8f03 drivers released month before but wuthout reposting so have NVDAGF100HAL.kext..
Anyway it only works OGL support as both CUDA and OCL don't use it..
I have to use NVloader injector which anyway doesn't work with Fermi on 64 bit kernel mode.. note gf 275 works in 64 bit with this injector also..
note i wanted to fix and all I found was a cuGetExportTable and something like MacCompatibiltyTID used by a checkcompatibility executable perhaps fixing it will work..
One in Nvidia forums assumed OCL broken fixed creating a OGL context beforce searching for OCL devices (oclgetdevice) but this trick didn't work..
*Storing ELF binaries instead of CUBIN deletes use of decuda hopefully one very interesting solution is..

*Seeing MAGMA webinar seems big release for SC2010 with some big features check magma presentation for what to expect..
*Physx 3.0 nearing to launch as Physx Visual Debugger includes support for it in release note says..

Note this brings concurrent kernels support for Fermi for improved perf on physics simulations.. hopefully also includes wrinkle meshes feature studied by Mueller.

Note also GPU AI notes once Function pointers supported on CUDA will use it so expect a new release sometime optimized even more for Fermi too..

Probably anuonced at Siggraph.. even launching later..

Hope too see also APEX shipping for other than Big AAA games i.e. downloadable for everyone..

Lastly I expect Optix 2.0 and Cg 3.0 final  for Siggraph and let's see also in time OpenRL with OpenCL support for GPUs would be interesting for ATI.. Note also Luxrender GPU 1.6 brings Stocasthic Photon Mapping and uses OCL on ATI GPUs also..

*Nsight also is moving fast from beta in early June now is RC state.. launching at siggraph?

*ATI Doubles on DirectCompute are broken.. altough feature flag is supported..

now we can test it with June DX compiler before it was broken for doubles inside control flow (loops, if,etc..)
Mainly compiling works but rendering shows issues vs Fermi which supports nicely..
Download my code.. (coming soon..)
*ATI GLSL driver is somewhat broken at least seems to geometry shaders as I fixed Nvidia Physx fluid demo to use non Cg code on GLSL code and some other fix related to point rendering and now seems to work but not without instabilities present as noise in screen even outside the window it fills..
Download ant test.. (coming soon..)
Also GLSL driver don't implement fetching integer textures with integer coordinates (texel2Dfetch( itex))
*CUDA 3.1 ships with three interesting examples: one is oclTridiagonal a fast tridiagonal solver.. interesting for a DoF cinematic renderer as in Metro using OCL/OGL..
other one is oclCopyComputeOverlap shows two things one is that concurrent kernel and exec is possible in OCL.. via command queues also shows there is an issue in 25x drivers that prevent full scaling I think good is 30% faster code and I obtain 20% on 25x drivers.. on 197 drivers I obtain 30%..
note that on both ATI and Apple platforms even with Nvidia GPUs exhibit no scaling and even negative scaling (-15%)
Good is that is fixed issue in 258.19 OCL 1.1 preview drivers with report CUDA 3.2 so I obtain back 30% overlap.. Note that other 258 drivers don't work (as they report older CUDA code 3.1 and OCL 1.0)..
One more interesting thing is that supposedly even dual dma engine is suposed to work on ocl so overlap would be 50%.. seems restricted to Tesla but Nvidia has been less detailed than double capping on Geforce..
Luckily I have a trick for you 197.44 driver seem to support Dual DMA engine on Geforce Fermi too!
This is OGL 4.0 driver so all you lost to current 256 drivers is CUDA 3.1 features only.. Linux also use OGL 4.0 driver on developer.nvidia.com and you have it...
Note also 197.75 etc don't work only work with this..
*So seems DUAL DMA engine is broken/disabled on Geforce Fermi without any reason other than economical..
*CUDA simpleStream seems to show broken streams on Fermi but it's due to not sending enough work.. a simple fix..
*Matmul by Lschien is one of the fastest ones for CUDA but it fails currently on fermi due to using cubins with obtained modifing tesla asm via decuda cudaasm.. thanks god seems related to volatile keyword don't working correctly pre cuda 3.0.. author suggest a fix assuming this works that uses cuda variant 6.. I have tested and it works so it's fixed I obtain near 850Gflops on Fermi 470 at 1650Mhz..
*Lot of soft updated to CUDA 3.x even 3.1 right now: NPP 3.1,CULA 2.0, JACKET 1.4,OpenMM 2.0 on Zephyr SVN, Gromcas 4.5 beta,GMAC, etc..

More news:

Also Nvidia has released a lot of drivers on 256 brach lets see rough differences/progression:
197.44 first OGL 4.0 driver and also unique supporting Dual DMA engine on Fermi on on Tesla/Quadro boards.. also has no issues in single dma..
256 add cuda 3.1 currently all has issues in concurrent kernel and exec on Fermi at least on OCL
257.15 bluray3d
257.19 nsight june beta drive
257.21 whql (supports nsight)
257.29 ion support accelerated dxva flash with pciex 1x devices
258.18 ocl 1.1 beta (says cuda 3.2!) fixes oclCopyCompute issues (but single DMA on Fermi)
258.48 first supporting Quadro Fermis..

258.69 shipping with 3d vision surround (Nvidia ntersect says youtube 3d support coming soon.. also I hope they add windows DX 3d vision support soon..)

Some other striking news :-) are:
*OpenCurrent 1.1 ships with CUDA 3.0 and multigpu code..
well I have been testing with CUDA 3.1 because I have Ubuntu 9.10 and with CUDA 3.1 GCC 4.4 works ok (so Ubuntu 10.4 is right also..) and has some issue related to now supporting true functions I think I must add some static to a function as cuda 3.1 release notes porting guide says.. with CUDA 3.0 GCC 4.4 doesn't work so I have to check with a Ubuntu 9.04 if I don't fix..
*OpenMP to CUDA compiler is avaiable in Cetus 1.2.
*PGI 10.6 is avaiable integer support in kernels and VS 2010 support at least.

I have tested GATLAS and is good at least 260 gflops on a gtx 275.. and I tested on MAC so at least works in Lin and Mac without much work and says author with 5870 and stream 2.1 achieves some image kernels 1,3 tflops so similar to cal++ matmul in OpenCL! have to test or modify code(?) for double testing..

Some tricks and work to do:

RAW DATA:
I know its lame but at least you can emulate 3d image writes on cuda with surfaces using ptx 3d tricks (post later).
I have to put a sample of CUVID on MAC.
SimpleStreams in cuda seems fermi bad in forums says increase work to 500.
matmul chien says put volatile and check (works!)
bsgp fermi support checking mail with author..
sparse matrix ati code test on fermi..

See fermi benchmarks:
nvidia benchmarks in blog
openvidia benchmarks..
cula blog
jacket blog
same papers of hpg2010 presentations billeter scattering and aov mcguire..
seems also code of rasterization and color stocastic shadow map coming soon..

Read More

ATI Stream SDK roadmap

I have found a roadmap of ATI Stream SDK till end of year:

DISCLAIMER: It's on Internet and found with some luck.. no breaking of NDA

Let's talk about it..
currently AMD OpenCL lacks:
*opengl interop issues:images interop issues (for example copy buffer to image where image is opengl tex acquired doesn't work)
*expose multiple component images (other than rgba)
*DX interop
*expose all graphics mem (currently 128-256mb)
*Catalyst integration

Stream SDK 2.2 Adds:
*OCL 1.1 (3 component vectors is part and image support ocl 1.1 is multiple component images (r,rg,rgb))
*DX10 interop (seems only that no dx9 or dx11 as Nvidia has)
*mem fences don't generate unneeded barrier isa instructions
*append buffers (what about also about GDS extension)
*seems atomics ocl 1.1 is nothing new? and offline compilation goes final from preview and dpfp adds fma as others are supported now(?)
dpfp fma should allow peak test kernels in benchmarks showing high numbers.. near 400-500gflop/s..

A lot more interesting is 2.3:
*In process compilation of OpenCL kernels means no shipping LLVM compilers (llc,etc..) and hopefully means will be integreated in atiocl.dll so it can ship OpenCL builtin in Catalyst 10.12..
*Library models
*C++ template support in kernels (I hope this means you can specify at least kernels args depeding on template argument for supporting double and float kernels with one code for example similar to CUDA support)
*Adds trig DPFP routines (but still no complete DPFP support seems so horrible as Nvidia shiping since October 2009 and AMD said support coming gradually since end 2009)
The more interesting is last three:
*FFT library: why not also a blas lib, I suspect is ocl based as directcompute has its fft lib
also is going to be part of acml? currently matmul in acml gpu is cal based..
At least I hope to be only binary library and also for Win and Lin so for Mac I hope somehow we can extract  OpenCL kernels or create a wrapper around it and use Wine or something like this to test perf on MAC on AMD boards is correct..
*OpenPhysics: well at least some to play, I expect cloth, soft body and SPH particles support in OpenCL and/or DirectCompute.. well in bullet site there is a preliminary executable with cloth demo and AMD worker talking about state of soft body support (http://code.google.com/p/bullet/issues/detail?id=390#c3) seems since last week also we have directcompute and opencl code for both cloth and soft body in trunk..
Also by September we will have DMM 2.0 as said in GDC that has some OpenCL love for this rigid body+fracture simulatior..
*OpenDecode UVD: Well a cuvid/vdpau library for AMD boards.. Nvidia has put lot of love to GPU video decoding and interop with CUDA/OpenGL with CUVID for Win and Mac and VDPAU for Linux..
VDPAU has since 256 drivers efficient OpenGL and CUDA interop.. CUVID has by def efficient CUDA interop and fast OpenGL/DX interop in Windows.. CUVID for MAC only seems good for feeding data to CUDA as OpenGL interop in MAC is slow right now (and has been so, since ever)..
I expect this brings fast interop to OpenCL on Win and Lin and that adds to DXVA DX interop on Win and AMD xvBA on Linux which VAAPI wrapper seems to provide fast OGL interop..
So Mac seems left but I hope recent video acceleration API on 10.6.3 supports AMD 5xxx cards when released and also that VC1 support is added in addition to h264.. I think this provides fast path to OpenGL textures so as OpenCL/OpenGL interop is fast on Apple provides also OpenCL interop on that platform..
Another thing is if Dual Stream acceleration will be exposed and supported.. on Nvidia I think both DXVA,CUVID and VDPAU expose with a GTX 470 at least..
Also related is Catalyst 10.7 having improved support for VLC 1.1.1 DXVA decoding for AMD cards which I presume relates to fast path GPU/CPU sending of frames works..
Remember also last month Nvidia released a ION driver (257.29) improving perf with DXVA on ION with PCIex x1 as Flash requires (GPU->CPU->GPU roundtrip)..

What's left after OCL 1.1 and stream sdk 2.3:
Well I expect Global Data Share and shared registers extensions,3d image writes, true complete DPFP support (cl_khr_fp64), complete BLAS and FFT lib (as CUBLAS and CUFFT in CUDA),  pinned mem working, host mem accessible from GPU extension, gather4 instructions for image support in OpenCL, and working concurrent kernel and mem transfers (i.e. concurrency in oclCopyCompute CUDA 3.1 example >=20%)

Read More

About AMD OpenCL 2.1!

AMD is progressing good and now we have an OpenCL stack with a lot features/optional extensions published and even AMD propietary ones:
regarding supported extensions:
*Image support: well only on 5xxx GPU (i don't know but I expect for CPUs also support as Apple CPU implementation? 4xxx don't expect but should be possible (CAL supports image/textures on 4xxx)))
right now only RGBA formats: but only supports 10/11 formats which are the obligatory ones (Nvidia has 7x).. well all rgba 4 channels so some Nvidia examples won't work..
well in 2.01 you can use export or set GPU_IMAGES_SUPPORT and get it on 5xxx..
no support on CPU also..
2.1 really has 3d tex support (didn't work in 2.01 hack)..
You can test Nvidia ocl samples oclVolumeRender and oclsimpletexture3d if you change samples to load on a 4 channel tex:
basically change in initCLvolume or oclsimpletexture3d h_volume to use 4 channel in initCLvolume:
volume_format.image_channel_order = CL_RGBA;
volume_format.image_channel_data_type = CL_UNORM_INT8;
       
        uchar * h_volume2=(uchar *)malloc(volumeSize[0] * volumeSize[1]*4*volumeSize[2]);
        for(int i=0; i<(volumeSize[0] * volumeSize[1]*volumeSize[2]); i++)
        h_volume2[4*i]=h_volume[i];
        d_volumeArray = clCreateImage3D(cxGPUContext, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, &volume_format,
                                        volumeSize[0],volumeSize[1], volumeSize[2],
                                        volumeSize[0]*4,volumeSize[0] * volumeSize[1]*4,
                                        h_volume2, &ciErrNum);
   
Also a bug mentioned in developer notes is linear filtering can't work if setted constant via
constant sampler_t volumeSampler = CLK_NORMALIZED_COORDS_TRUE | CLK_ADDRESS_CLAMP | CLK_FILTER_LINEAR;
(also note CUDA 3.0 final has a bugs regarding linear filtering on 3d tex samples and Nvidia and AMD OpenCL  samples aren't working on other IHV OCL because some need constant or __const samplers and others not work with that I don't remember)
so I have to comment this sample in volumesample (simpletex3d does the right ting) in cl shader and setting via adding a parameter
__kernel void
d_render(__global uint *d_output,
         uint imageW, uint imageH,
         float density, float brightness,
         float transferOffset, float transferScale,
         __constant float* invViewMatrix
 #ifdef IMAGE_SUPPORT
          ,__read_only image3d_t volume,
          __read_only image2d_t transferFunc,
          sampler_t volumeSampler

 #endif
         )
then you can add form simpletex
case 'f':
            linearFiltering = !linearFiltering;
            ciErrNum = clSetKernelArg(ckKernel, 10, sizeof(cl_sampler), linearFiltering ? &volumeSamplerLinear : &volumeSamplerNearest);
            shrLog("\nLinear Filtering Toggled %s...\n", linearFiltering ? "ON" : "OFF");
            oclCheckErrorEX(ciErrNum, CL_SUCCESS, pCleanup);
            break;
to keyboard gl..
Also checked simultaneous image and opengl interop and it worked..
http://dl.dropbox.com/u/1416327/clinterop2.c
define USEGL or not to check image support or simultaneous image and opengl interop.. (clcreateimageformgltexture..)
note in gl interop there is some image updown and some greener image but Nvidia OCL gets similar output so I have to revise code but for now is working..
Lastly what's lacking is 3d image write support but using a sample (using amd new simple image has some more or less disabled 3d texture write test so changing a few simple lines you can test)..
currently I see even cl shader compiler has imagewrite 3d signature so when changing code if you pass a 3d image object and using int2 for coords says it needs a int4 arg.. changing the code the error you have is "I can't find builtin function #xyz" so seems all is well in place.. including  the #pragma enable image 3d writes fails saying extension not know but anyway seems like perhaps next version has this support and implementation more advanced than Nvidia?
GL Interop:
Well AMD example has VBO example and works.. using oclPostprocessGL as PBO example also works..
Even  changing code in these two demos for creating VBO and PBO GL objects before CL context creation works and that shouldn't work as is said to be a limitation..
So seems current limitation is GL context before CL context which is per spec as createcontext needs gl context..

Also as said before we have image support GL interop working..
Byteaddresable well works but at IL level seems is some and and or masks so hardware has no native byte addressing also by the fact that IL shows UAV which is a dx concept that needs 32bit aligned accesses so I think not native also even UAV DX byte buffer allows byte addressing but as said at 32bit aligned.. general UAV a like int vectors so a[1] is as a byte pointer a[4]..
I have to see how can AMD fight against race conditions if not native when multiple threads write bytes in same word as if doing RMW must use atomics?.. and overhead

What troubles me the most is that Apple demos as GL interop fails but with GL interop is using image support
and also some copyimagetobuffer or buffertoimage so I have to see if is GL interop problem, image support problem or copy problem.. then I will release it..

regarding samples it has new boxgl sample not mentioned:

AMD is publishing a lot of extensions (some very simple):
*amd_printf: I have checked and now works now with Visual Studio (2.01 with Linux?)
if you don't enable explcitilly compiler fails.. previously no way to disable it..
*amd_fp64: GPU no changes as 2.01 so +-/* only and in CPU whealth of features but no conformance and strictness so no dmad i thing so how can GPCbenchmark get so high gflops in doubles without mad I don't know.. also I don't know if mads are generated for integers now as it seems to use it but last time i checked 2.00 in january didn't use that.. also what about mad24
*amd_media_ops: could obviate pyrit cal++ implementation that on trunc or svn has 2x-3x improvement over OpenCL due to to bitalign use now you can use on opencl now.. would be interesting to see if cpyrit gets support now that on trunc also code used rotate opencl native instruction for better possibly perf with ISAs having it.. also it has SAD support that was anounced by AMD to add to OpenCL on 5xxx launch
in binary there are hints of:

amd_vector3 I assume defines float3 or no.. i think nvidia hasn't it even unofficially so good to have..
also some apple demos #define float3 so good to be able to disable to it similar to printf as this code should now work on AMD without any modification..
amd_atomic_counters similar to unpublished glsl atomic_counters?
also ext_device_fission is currently lacking extension ocumentation.. and only cpu but seems to expose concurrent kernels on Fermi GPUs so hope Nvidia supports it.. anyway it's a shame using two or more commanq queues aren't able to extract perf in Nvidia as Nvidia supports it in CUDA via streams which is a similar concept.. I have to post the code I coded to check it..

Also now I have found trick to enable fully working  GLSL sprites used in Nvidia OCL samples and Particles demos simply by changing in fragment shaders tex_coord[0] glPointCoord..
(thanks pboudier AMD forums) before you can show as point redering particles use 'p' key or with menu optiuon..

Regarding samples interop many bugs are fixed but only remain the ones due to architectural differencees:
mainly warp related, shared mem size, workgroup size and other out of resources limitations (register stack?) etc..
Particles and Sort on AMD examples need a fix I posted some time ago..

Biggest complains/suggestions and bugs/limitations are:
*Byte addresable HW native? thread race conditions issues with different byte and same word by multiple threads or perf issues due to atomics usage?
*More image formats support (at least R and RG with half float, float and int8/16)
*3d image writes
*d3d9 and 10 interop: disabled in 2.1 (worked in 2.01?) supposedly coming in next version in Q3 anyway a new khr d3d10 extension is published on Khronos which is similar to nvidia but differs in supported a shared handle parameter and a flag in device info saying if it will get improved perf interop with a shared handle..
Would be good KHR d3d9 and D3d11 extensions as Nvidia and AMD supporting it..
for example DXVA->opencl via this extensions should enable MultiIHV via badabooms in the decoding part and perhaps full using MFT GPU encoders..
*Doubles still lacky on GPU (+-/*) and not conformant on GPU
*No device fission on GPU as AMD shared it's stream processors have support for it in HW at least the 80 shaders blocks so 20 conc kernels in 5xxx theoretically possible.. but I think is a CAL API moslty limiation or AMD IL so can take a while to fix?

So biggest Nvidia remaininglimitations now are:
*3d image writes
*Conc kernels on Fermi->No device fission on GPU  or using multiple command streams..
Also Dual DMA is usable?

Read More