GPU computing

Some news!

2010-07-10T20:16:00.001+02:00

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

DirectCompute Double precision Mandelbrot demo and more..

2010-07-05T04:41:00.002+02:00

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):

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

2010-07-04T21:06:00.003+02:00

*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..

ATI Stream SDK roadmap

2010-07-03T15:40:00.004+02:00

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%)

About AMD OpenCL 2.1!

2010-05-05T20:57:00.022+02:00

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?

Mandelbrot using OGL 4.0 features (double precision and precise keyword)

2010-04-07T04:57:00.016+02:00

http://dl.dropbox.com/u/1416327/mandeldouble.rar
above executable contains:
*uses gl_arb_gpu_shader5 in a float-float implementation with precise keyword for fixing agressive Nvidia compiler
*uses arg_gpu_shader_FP64 with doubles.. and fallbacks to doublepAMD on catalyst no ogl 4.0 drivers..
*normal mandelbrot implementation

on AMD 5850 with 1920x1080 res ati gl 4.0 drivers
I obtain:
*15fps using float-float approach..
*50fps using doubles with ati gl 4.0 drivers
*130fps using single precision
Note pre GL 4.0 drivers using doublepAMD attain 36fps on double precision now gl 4.0 drivers either doublepAMD or double attain 50fps..
You can deduce Gflop/s seeing glsl code.. it's very high..

I use #if 1 instead of #ifdef GL_arb_gpu_shader5 or shader_fp64 as then shaders work on Nvidia GL 3.3 drivers altough without doubles (instead double precision) and without precise keywork so float-float is still bad!
i.e. I force #pragma extension enable

Sorry for big exe is linked to Cg altough not usingly now it was used for correct disabling of optimization on Nvidia.. but it's not working now
program arguments are first pixel start horizontal offset for multimonitor setups second fullscreen or no then fragment and vertex shader and then zoom and x and y offset in mandelbrot..
It's used for showing a enough zoom for seeing diff between single and bigger precision either double precision or float-float.. last argument in use glsl or cg backend..
but as said cg is broken..

seems amd doesn't optimize so many as float-float without precise works ok!

AMD 5850 with ogl 4.0 drivers windows 7(with fps)
http://dl.dropbox.com/u/1416327/float-float.jpg
http://dl.dropbox.com/u/1416327/fp32.jpg
http://dl.dropbox.com/u/1416327/fp64.jpg
NVIDIA
bad float-float is similar to amd fp32 photo

fix for float-float-> use precise
I hope this goes well with Fermi OGL 4.0 drivers and also enable precise keywork for GL 3.0 hardware..
Cg has a trick for disabling optimizations so it's not needed..
search blog for more info..

vec2 dblsgl_add (vec2 x, vec2 y)
{
precise vec2 z;
float t1, t2, e;

t1 = x.y + y.y;
e = t1 - x.y;
t2 = ((y.y - e) + (x.y - (t1 - e))) + x.x + y.x;
z.y = e = t1 + t2;
z.x = t2 - (e - t1);
return z;
}

vec2 dblsgl_mul (vec2 x, vec2 y)
{
precise vec2 z;
float up, vp, u1, u2, v1, v2, mh, ml;

up = x.y * 4097.0;
u1 = (x.y - up) + up;
u2 = x.y - u1;
vp = y.y * 4097.0;
v1 = (y.y - vp) + vp;
v2 = y.y - v1;
//mh = __fmul_rn(x.y,y.y);
mh = x.y*y.y;
ml = (((u1 * v1 - mh) + u1 * v2) + u2 * v1) + u2 * v2;
//ml = (fmul_rn(x.y,y.x) + __fmul_rn(x.x,y.y)) + ml;

ml = (x.y*y.x + x.x*y.y) + ml;

mh=mh;
z.y = up = mh + ml;
z.x = (mh - up) + ml;
return z;
}

Some things I forgot..

2010-04-02T04:40:00.001+02:00

First is directcompute blog
http://www.yakiimo3d.com/
with nebularot code
also seems rigid body on gpu is starting
before physx sdk 3.0 and batman with a gtx480 will use it I have found on nvidia ftp:
BatmanAA_GTX480and470_PhysX_Patch.zip
dated 30 march 2010
and has rrb.dll that is
gpu accelerated rigid body dynamics v 1.0.0.1 dated 11 january 2010
depends on cudart 3.0 patch 9
exposes

AgPmDestroySourceConnection
AgPmEventEnabled
AgPmEventLoggingEnabled
AgPmSubmitEvent
PrbCreatePhysicsSDK
PrbFree
PrbGetPhysicsSDK
PrbMalloc
PrbMallocDEBUG
PrbReleasePhysicsSDK

similar to physxcore

AgPmDestroySourceConnection

AgPmEventEnabled

AgPmEventLoggingEnabled

AgPmSubmitEvent

NgCreateCoreSDK

NpCreatePhysicsSDK

NpGetFoundationSDK

NpGetPhysicsSDK

NpGetPhysicsSDKAllocator

NpGetUtilLib

NpReleasePhysicsSDK

NxCreateCoreSDK

NxGetValue

batman originial release had this?

Also note cudart 3.0 patch 9 is found in physx runtime 22 feb 2010

note we have cuda 3.0 rt dll
8 beta
9 physx
11 optix2b3
14 final

Megapost!

2010-04-02T02:13:00.000+02:00

Today fools{

*GTX 485 is 512 cores 3gbytes gddr5 and 850/1750 shaders..

*ati 5990 has 4 gpus in board..

*bulldozer benchmarks

}end fools..

ATI has released:

*5870 2gb 6 outputs

*GL 3.3/4.0 drivers (linux &win)

*GPU perfstudio 2.2

*AMD ADL SDK 3.0 (aka eyefinity sdk)

two stream documents:

*OpenCL Programming Guide

*GPU Computing: Past, Present and Future with ATI Stream Technology michael chu

lame to see backup slide cuda vs opencl..

*vaapi with h.264 decode on westmere cpus on git

well we have now h.264 gpu decode on linux via vaapi for intel nvidia and amd cards..

well amd with 5xxx not ok and intel g45 will wait until q3 2010..

also what about vc-1? ati and nvidia support is there even on 8800gt via latest vdpau..

intel will catch up?

and what about dual hd decode is working with every api/implemenation on latest gpu's all intel hd 2010 graphics amd 5xxx and gt240 and fermi have hardware suport for it..

what about h.264 mvc vaapi exposes it? i.e. api allows that and what about xvba,dxva and vdpau..

also now we have cuvid for mac even in x64 possible so cuvid will allow or allows mvc?

also now gnash vaapi support is integrated in trunk and compilable in mac and windows seems so we can

port vaapi to mac and win and even implement a cuvid vaapi wrapper?

this would allow mplayer and gnash to support gpu video decode on mac for nvidia cards for hd video and flash video..

Nvidia has released:

*Nexus march beta (same as shown in GDC'10 so would allow d3d10 and d3d11 shader debug on Fermi..)

*Optix 2.0b3

*CUDA 3.0

*OGL 3.3 drivers

Still lacking

*Cg 3.0?

*OGL 4 drivers with ext_image_load_store and ext_image_atomic_counters support

*Linux Fermi drivers (win has 197.17)

*3d vision surround sdk

*3dtv hdmi 1.4 drivers

*256 drivers

*nv d3d11 sdk presumably has:

hair tess and water tess demos

*physx sdk 3.0 with rigid body on gpu and height field water as fermi launch demo?

*voltage tweakers software and max oc with it for gtx 480 (900mhz?) and 470(750/800mhz possible) and bencharmks

*optix 2 and nexus 1.0 final

*test voxilla demos and fp64 in cuda and opencl perf cud-z

is 1/4 of tesla products? can be hacked? see ptx code and cubin code..

gpu computing:

*cudart x64 for mac

*cuda-gdb for mac

*cuda-gdb support for ocl binaries

*promised nv official tools for diassembly and assembly of fermi binaries (new cubins old use decuda or also will support sm_1x binaries?) promised soon in sigg asia cuda perf optimization course..

*mac cuda-opengl efficient interop?

official perf

*tesselation 6-8x

*raytracing 3.5x

*sli near 2x on d3d11 games

*3d vision near 2x (see 3d vision blog)

ok but rops and texture power very low and seems tex units capped at half

as gf104 info surface has 64 tex units also..

nvidia agrees has gddr5 controller problems so no uses gddr5 5000mhz chips to 1250mhz..

470 seems use 4000mhz chips..

reviews notes:

*noticias3d has slides and perf vs 5870 with launch 8.66 drivers so can be good to test perf improvement overall as this would be the perf six months ago.. cat 10.2/10.3 have 10% perf improvement..

*ixbt uses rightmark geo shaders perf..

*anandtech has chen nqueen opencl perf. and folding@home new client but other site claims on 50% perf vs 2-4x improvement anand says

*review have new d3d11 bencharmk by sweden company

*sandra 2010 gpgpu benchmarks but double prec is bad..

*d3d11 games metro, heaven 2.0, dx 11 sdk tess demos, just cause2 benches..

*luxrays perf on beyond3d forums..

Apple has released 10.6.3 without amd cal libs (see pgi 10.2 with cal info saying aticalrt.dylib)

also seems to have almost ogl 3.0 for amd nvidia has some extension less and cpu driver lacks 3/4 extensions..

I have found fermi on ogl binary driver but not support really..

phoronix found ogl drivers has more than 50% perf degradation on 9400 (bad)

but should allow steam to run on mac well..

regarding opencl still no new headers for cuda sdk 3.0 issues and seems no big improvements as no mentioned on release

I have to test 10.6.3 with a cuvid x64 executable i have, optix 2.0b3 sdk, run fft opencl and ocean apple demos on both nvidia and ati gpus.. and run nvidia ocl ft3d sample which says has issues with apple opencl to see if fixed..

also ocl headers in ipad 3.2 sdk golden master?..

still seems no double support on opencl for nvidia and no image support for ati gpus on apple..

add that to no fix double prec on compute shaders..

double in ogl 4.0 (ext_gpu_shader_fp64):

Nvidia has released ogl 3.3 but 4.0 drivers will support fp64 on

gt275?

also double support is on 4850 cards on ati 4.0 drivers?

also will nvidia release wgl_nvx_dx_interop spec and ext_image_load_store extension on gl 4.0 drivers?

any extension more?

with that at least directcompute and ogl will allow 3d image writes.. opencl allows 2d image writes by default and cuda least good? with from pitch linear mem..

lacking is opencl 3d image writes extension and cuda surface functions removed from cuda 3.0beta.. I think they didn't work..

also a post is interesting in nvidia forums saying that now opencl using a writable texture seems to not

Iz3d 1.11 released has shutter support (i can't test in samsung 120hz because I have activation issues)

but I have found anaglyph which shows algorithm goes good d3d9,10 and 11 in directx sdk samples..

lame ati d3d11 mecha ladybug doesn't work ok..

mecha crashes and ladybug doesn't affect view..

nvidia compute shader ocean demo doesn't see good and 3d vision works 197.13 with that demo!

also some tesselation doesn't work

brief:

*32 bits ok 64 bits examples crash (its my system fault?)

*Youtube 1080p 3D HD works with internet explorer with flash 10.0 not 10.1 and with youtube in english mode!

*Windowed stereo mode works.

so nvidia has to add youtube 3d and windowed stereo mode support (for non quadro) in 256 magical drivers.. better if they add also nvidia 3dtv and hdmi 1.4 out for opengl qb for quadros..

Also diagnostic utility reports about ati aqbs surface format d3d which must be amd catalyst 10.3 3d support shows is not supported altough using catalyst 10.3 whql so seems I must have lcd setup to 120hz or finds a hdmi projector? anyway can't setup hz on catalyst cc now..

I would love to have cuda hook that allows to enable graphics interop trough host for tesla computing driver on windows and running kernel moduly only on linux to run nbody for example..

it's a shame ogl interop was through host if not run on same gpu on earlier versions not it returns error..

also for opencl which reports ogl interop..

both for d3d and ogl interop..

also would add a cubin to ptx on the fly for running nufft or fastest matmul cubin codes on fermi..

also test enabling cu-force-ptx-jti

Would be good to test d3d ocl interop with dxva 2.0 d3d9 tex? interop to build a open source badaboom..

I would love to see on a 8800gt or gt200 with vp2 (vc-1 vld not supported) where we have lower cpu usage if using cuvid, dxva or vdpau.. assuming all these handle it..

the same for dual stream hd and mvc when it gots out..

Currently I found lacking on AMD 5xxx:

*OCL image support

*OGL-OCL tex interop

*xvba 5xxx incorrect decoding

I would love to have a simple ogl qb driver with anaglyph output for testing porting gnash, mplayer etc.. to support 3d stereo rendering and youtube 3d on mac and linux..

Then port 3d vision to these oses..

note I have learnt from Unigine Heaven 2.0 that iz3d doesn't work from launcher but it has .bat files for launching the demo and with that iz3d works in d3d9, in d3d10 crashes as soon as activated and d3d11 depends but no sees good..

note seems windows demo compiled on 7 march has no support for amd old tesselator gl extension editing haven.cfg so doesn't work also doesn't work with amd ogl 4.0 drivers..

on linux you can use heaven 2.0 with ati tesselation as linux build is later..

I would like atioc utility on linux to overclock much than officially supported as msi afterburner does..

have to hook ati adl and see..

angle google code project is improving fast:

*now has ogl samples included with esut.h and support for loops in shaders etc..

*64bit requires

--- src/libEGL/Display.cpp (revision 49)

+++ src/libEGL/Display.cpp (working copy)

@@ -63,8 +63,8 @@

}

else

{

- EGLint minSwapInterval = 4;

- EGLint maxSwapInterval = 0;

+ int minSwapInterval = 4;

+ int maxSwapInterval = 0;

Index: src/libGLESv2/geometry/vertexconversion.h

===================================================================

--- src/libGLESv2/geometry/vertexconversion.h (revision 49)

+++ src/libGLESv2/geometry/vertexconversion.h (working copy)

@@ -122,7 +122,7 @@

static const std::size_t finalWidth = N+(N&1);

};

-template

+template

struct WidenToFour

samples require more changes also..

I have been trying to port

Crazy drivers:
amd:

cat 10.2 B_95228 3/2
cat 10.3b B_95437 5/2
cat 10.3 B_96537 3/3
10.3a B_97263 14/3
10.3 ogl4 B_97624 24/3
10.3b B_97763 25/3
10.4 shipping for ubuntu 10.4

nvidia

196.75 required nexus support
197 or higher ->ocl d3d interop
197.13 cuda 3.0 oficial ones and whql
197.15 ogl 3.3 driver
197.16 notebook verde driver with 3d vision external support
197.17 fermi launch press drivers
197.25 starcraft dx8 issues

geforce 256 in april with 3d vision surround

about ogl 3.3/4.0 drivers

ogl 3.3 samples released..

ogl 4.0

openglext and extensions viewer show ogl 3.3/4.0 extensions

google code gle,gloader load 4.0 extensions.. glew?

info released about http://developer.download.nvidia.com/opengl/specs/GL_EXT_gpu_memory_info.txt

Fermi post launch analysis:

lacks
http://forum.beyond3d.com/showpost.php?p=1414824&postcount=283
latest gpgpu releases:

*thrust 1.2
*jacket 1.3
*Folding@Home fermi with openmm? gpu3 client
*cudpp 1.1.1?

released:
*nvidia Design Garage
*supersonic sled

demos not public:
*Raging Rapids tech demo
*hair demo
*water tesselation demo
*d3d11 demo by sweden company

testing cufft I have found since 2.3 includes nufft cubin only improvements (nufft paper sc09)
nufft has test bench code for 256^3 fft trasnform.
cufft in sc09 has perf over 160gflops for 256x144x192
cufft 3.0 only superfast if power of two every dimension altough different..
if not 20-30glfops

have to test fft dx compute shader microsoft library..

amd 5850 in glext shows

*doesn't have:

GL_EXT_stencil_two_side?

GL_ARB_compatibility (3.1)-> seems present so it present if 3.1 queries?

GL_EXT_shader_image_load_store->present in dll!
accessorStore UAV_STORE
imageLoad imageStore

GL_ARB_shading_language_include->seems has include basic support!

has:

GL_EXT_vertex_attrib_64bit (no published spec)
GL_ARB_texture_compression_bptc->tiene ext
GL_EXT_shader_atomic_counters (no published spec)
imageAtomicAdd imageAtomicSub imageAtomicMin imageAtomicMax
GL_ARB_texture_swizzle->tiene ext_texture_swizzle
GL_ARB_texture_buffer_object_rgb32->tiene ext

propietary
add GL_AMDX_debug_output
amdx->GL_AMD_name_gen_delete
GL_AMD_conservative_depth

-------------------------------------------
Not implemented extensions in OpenGL 2.0:
GL_EXT_stencil_two_side

-------------------------------------------
Not implemented extensions in OpenGL 3.0:
GL_NV_depth_buffer_float->tiene arb_depth

-------------------------------------------
Not implemented extensions in OpenGL 3.1:
GL_ARB_compatibility

-------------------------------------------
Not implemented extensions in OpenGL 3.3:
GL_ARB_shading_language_include->no ned
GL_ARB_texture_swizzle->tiene ext_texture_swizzle

-------------------------------------------
Not implemented extensions in OpenGL 4.0:
GL_ARB_texture_buffer_object_rgb32->tiene ext

-------------------------------------------

*I can't speak but I have betas of both:

*OpenRL 1.0b2

Has Windows (x32,x64) libraries and Mac x32 only libraries

still lacking linux and mac x64 binaries..

remember optix has mac also but only x32..

no opencl bits found anywhere and support from now and only cpu release

but uses all my 8 cores..

would be nice to port optix to OpenRL samples and tutorial and viceversa..

or better make a OpenRL wrapper to Optix 2.0b3 with fermi support..

*Intel Compilter 12 (composer 2011)

cilk,#pragma vector size(4,8) etc..

vs2010 support

aes-ni for crc32 and better avx overall

and more..

ipp 7.0 beta

intel compiler 12 beta

tbb old version?

*I have code libecuda,libptx of PFC ptx emulator of UPC now..

trying for windows and update to ptx isa 2.0

*Still no gdebuggerCL

ptx 2.0 isa released:

includes ptx 1.5 info also (llvm ptx nvidia opencl compiler emits this code)

-> mainly adds separate tex and sampler setup also same __param stuff as functions arguments

-> also shows opencl has no name mangling for kernels ocl.. now testing if ptx with addc can be inserted on opencl i.e. conversor from cuda 3.0 ptx kernels to opencl ptx kernels would be good..

also a cubin to ptx is possible? would allow me to run fastest to date matmul on fermi as fermi doesn't run cubins..

there are some limitation? ask barra creator he has a tesla cubin simulator..

so I could theoretically go from cubin to ocl compatible ptx code..

ptx 2.0 shows for fermi

HAS (also implemented):

*d3d11 cs 5.0 integer instructions

*ldu

*unified address space ld loads

*surface functions (load and store)-> 3d image writes

has load with format or not and with format loads are not implemented and stores with format also not implemented excepting a b32 format

EXPOSES:

*recursion via..

*functions calls with stack (so recursion possible) without defining and abi

*calloc function

*variable args to functions

note this is not implemented in 2.0

lacking still are:

*jump to register/pointer or call to register/pointer (virtual functions?)

*host system calls malloc,printf..etc..

also cuda book shows:

*fermi predication based on

"A Comparison of Full and Partial Predicated Execution Support

for ILP Processors"

*fermi supports terminating kernels when you want (driver stability improvements?)

also for load balancing..

*cuda fermi implementations priorities.. virtual unified space can take years..

*virtual address space good with GMAC approach for unified unique address for CPU GPU mem now GPU address is unified

What's for CUDA 3.1 and OpenGL 3.3/4.1!

2010-03-21T20:37:00.000+01:00

Let's see CUDA 3.0 vs beta:

*adds full blas support
*opencl local atomics
*ocl i cuda d3d9-11 interop..
*updated guides since beta..
still no ptx 1.5,2.0 specs..
also nv-cl extensions published now: http://developer.download.nvidia.com/compute/cuda/3_0/toolkit/docs/opencl_extensions/cl_nv_compiler_options.txt

Interesting notes.

*Float16 (half) textures are supported in the runtime

*cublas complete i ieee754 complaint fermi

*SGEMM performance on Fermi-based GPU is 30% lower than expected.

It will be fixed in 3.1.

*The stability of the large-prime FFT transform (signals with a length

that is prime and >64k samples) is extremely variable, giving single-

precision accuracy in the range 0.005->0.025. In general, smaller signals

experience greater accuracy.

*This package will work MAC OSX running 32/64-bit.

* CUDA applications built in 32/64-bit (CUDA Driver API) is supported.

* CUDA applications built as 32-bit (CUDA Runtime API) is supported.

(10.5.x Leopard and 10.6 SnowLeopard)

Note: x86_64 is not currently working for Leopoard or SnowLeopard

*CUDA applications built with the CUDA driver API can run as either 32/64-bit applications.

* CUDA applications using CUDA Runtime APIs can only be built on 32-bit applications.

SDK Release 3.0 Final:

* Replaced 3dfd sample with FDTD3d (Finite Difference sample has been updated)

* Added support for Fermi Architecture (Compute 2.0 profile) to the SDK samples

* Updated Graphics/CUDA samples to use the new unified graphics interop

* Several samples with Device Emulation have been removed. Device Emulation is

deprecated for CUDA 3.0, and will be removed with CUDA 3.1.

* Added new samples:

concurrentKernels (Fermi Capability)

* Bug Fixes

have added simplempi also..
have to test with intel mpi 4.0

MAC notes:
cuda.dylib is 64bit and has 195API and 195 185 dylibs versioned as 195_96 or 185_55..
*has cuda-memcheck but no cuda-gdb
*cuda kext is fatbin with 64 bits and also cuda.dylib so cuda driver applications are compatible with 64 bits
and compilable..
note also can boot in 64 bit kernel due to kext..
cudart 32 bit
then we can in theory program a cudart wrapper over cuda driver and compile in 64 bits more
now than cudart is stateless and has interop with cuda driver mem alloc..

all needed is cublas and cufft to be 64 bits compile in that..

we have code for cudpp,thrust and cusp and in the meanwhile volkov matmul,fft and lapack codes

so all these can be compiled with 64 bits if we had a cudart 64 bit and see what's up..
well I have compiled cudadevicedrv and matmuldrv
(i'm the first in the world to have 64 bit cuda apple binaries? excepting at nvidia..?)
I have get rid of cutil though compiling to 64 bits would be no problem some notes:
nvcc on mac defaults to 32 bits vs gcc defaults on 64 bits on Snow leopard..
so for using 64bits you must use -m64 in nvcc..
but for cuda driver projects nvcc is of no use since you can use g++ for cuda driver api and compile cuda
files to ptx with nvcc -ptx

if you use nvcc with -m64 you get both cpu 64 bit code but also using -ptx you get ptx code

using 64 bit pointers for Fermi?

so you can use 32 bit pointers in Fermi is better use 32 bit pointers..
so matrixmuldrv use nvcc -ptx for 32bit pointers and use g++ (-m64) and you get
but cudamoduleloaddataex i get error
CUDA_ERROR_POINTER_IS_64BIT = 800, ///< Attempted to retrieve 64-bit pointer via 32-bit API function
loading ptx either if I use a nvcc -m64 or nvcc (all with -ptx) get this error..
so ptx with 32 or 64 bit pointers doesn't change that..
I have to compare files with 32 and 64 bit pointers to see differences also with sm_20..
also note for nvcc -m64 to work either if it not needed needs /usr/local/cuda/lib64 to exist..
so I have copied lib->lib64 or do a symlink..
so you can now run it..
I have to write tutorial of using cuda and nvcc and achieving macos fat binaries(i386 ad 64)
*I see nvcuvid library for mac in gpu computing sdk.. only 32 bits..
/C/common/lib
and /C/common/inc/cuvid

Anyway I have a libcuvid (vs libnvcuvid) for 64 bits /usr/local/cuda (where i have get from?)

*also a pref pane control panel with autoupdate and shows gpu driver version and cuda driver version..

note opencl samples on mac no work until 10.6.3..

good is opencl not definided behavior (implementation specific) for nvidia:

http://developer.download.nvidia.com/compute/cuda/3_0/toolkit/docs/NVIDIA_OpenCL_ImplementationNotes_3.0.txt

issues with mac..

opengl 4.1/3.3 perfect release:

*ext_direct_state_access

*ext_separate_shader_objects

*RW textures (3d also) ext_image_load_store

*binary shaders (gl es 2.0 api)

in theory you can use some ir from 3dlabs frontend compiler source..

or also translate to hlsl via som translator (amd hlsl2glsl?) and then use binary hlsl shader..

also a good translator..

http://code.google.com/p/angleproject/

has flex/bison glsl parser and also a glsl2hlsl translator (es 2.0)..

going from binary to dx il via:

fxc /dumpbin

but dx il to binary? also how from dx il->hlsl or glsl directly..

I also have found wine handles/parses more or less dxbc files..

/dlls/d3d10/effect.c

static HRESULT parse_shade

NV OGL extensions:

*fermi fuction pointers and recursion for glsl?

would be good addition to bindless extensions and shader buffer load..

CUDA 3.1:
*cuda-gdb OpenCL HW debugging support..
*pinned GPU mem interop with MPI Infiniband.. (spring10 in sc09)

*template for a DirectCompute project
Currently there is no template for a DirectCompute project, but NVIDIA will be
providing one soon.
*Fix perf of CUBLAS SGEMM by 30% faster on Fermi
*Fix CUFFT perf vs 3.0beta goes 180-190gflops to 150gflops
*provide official cudaasm/decuda or documentation about cubin/ELF format for SM_20 devices? also for sm_10?
*PTX 1.5, 2.0 docs?
*Updated opencl best practices for Fermi? cuda best.. guide is updated but for Fermi?

*Surface functions: RW textures with x,w addressing etc.. also 3d image writes.. headers and exported functions in beta but removed in final..

Also CUDA to CPU compiler or is gpuocelot mature enough and also mac and windows ports avaiable..
would be good a direct PTX2CPU code conversor and using gpuocelto lib as cudart and cuda api..

Mac

*add cuda-gdb (with ocl also) and OpenCL visual profiler

opencl mac no xutan 2 ejemplos

cuda opengl slow mac

ship

is going to work with fermi cuda.kext

*Related is first 195 series 197 whql driver for Quadros enabling OpenCL on these devices..

Adds support for CUDA 3.0 for improved performance in GPU Computing applications. See CUDA for more details.
This driver resolves fan speed issues reported with version 196.75 drivers.
Adds support for the Open Computing Language (OpenCL) 1.0 in Quadro FX Series x700 and newer as well as the FX4600 and FX5600.

*Nvidia mentions compute cluster driver but is 196.28 not updated since early feb.. anyway d3d interop
added finally is not nedeed here..
*
to pierre boudier you cansee ogl 4.0 drivers soon and also a image write and random access extension soon ala d3d11 rwtexture..
ubuntu 10.4 fglrx 8.72

fglrx-installer (2:8.721-0ubuntu1) lucid; urgency=low

* New upstream release:
- Restore compatibility with kernel 2.6.32 and xserver 1.7 (LP: #494699).
- Add Passive Stereo support on workstation (FireGL/FirePro) hardware.
- Add Eyefinity support (more than 2 monitors on Radeon HD 5xxx hardware).
Officially WS-only but should work on consumer boards as well.

GL_EXT_shader_subroutine GL_EXT_timer_query

Also what about 3d stereo on linux:
*3d vision for opengl qb on quadro with stereo connector is here..
*a 3dtv for linux so opengl qb can be output to hdmi 1.4 on linux? this can add working on low profile quadros as stereo connector is not needed (is not needed in 3d vision is Nvidia way of artificially limiting to super high end quadros well expect perhaps better synch..)
also if they add VDPAU h.264 MVC and you decrypt bluray3d with anydvd hd you will be able in theory to see it in linux gpu accelerated decoding and sending to tv's via hdmi 1.4..
let's see also how windows is handled as not dxva 2.0 support it mvc? also not cuvid so leet's see if they add it to cuvid also..
so seems all cyberlink will get some library by nvidia or what?
*ATI has hooks for d3d9,10? d3d11? in 10.3, also fglrx 8.72 add passive stereo for ogl qb (active stereo is here right?.. but for 120hz lcds also?)
let's see also how ati manages output to HDMI 1.4 tv's via either IZ3D partnership or what? in fact I expect iz3d only hooks d3d stereo and the amd will add some HDMI 1.4 stereo from this hooks so will be good a sdk or documentation of this hooks..
Also Nvidia will be good publishing stereo sdk (promised in gdc2010) and hope also this hooks (d3d9-11) will work with 3dtv and output to hdmi 1.4 tvs.. In fact yes as Avatar and 3d stereo vision use this hooks presumably..
mac is out in this scope..

also nvidia can be late with fermi but not with software supporting it..
now d3d11 is with cs5.0 here and also we have now d3d11 interop for cuda in 3.0 and d3d11 interop with opencl extension and also optix d3d11 interop..
We have d3d11 interop with:
*CUDA 3.0
*OpenCL
*Optix
HW debugging:
Nsight.
All need to be released is nsight which will also bring d3d11 support (hw debug and profile) wii be good to hw debug cuda, d3d11 cs, cuda with d3d11 interop, and trace opencl and opengl (4.0? will be traced?)..

also cg 3.0 will have support for d3d11? and also sm5.0 opengl 4.0 support? i.e. tesselation shaders with glsl output?
note cgc 3.0 is shipping on tegra sdk and also as part of nvidia drivers 195 opengl compiler..
I have seen cgfx working with optix and cuda in a blog so hope they ship example soon..
http://lorachnroll.blogspot.com/2010/03/mixing-nvidia-technologies-thanks-to.html

GPU: GF100 @ 700MHz
- CUDA cores: 480 @ 1401MHz
- Memory: 1536MB GDDR5 @ 1848MHz 384-bit
- TDP: 250W
GeForce GTX 470:
- GPU: GF100 @ 607MHz
- CUDA cores: 448 @ 1215MHz
- Memory: 1280MB GDDR5 @ 1674MHz 320-bit
- TDP: 225W
- Price: $349US

- 3D APIs: OpenGL 4.0 and Direct3D 11
- GPU Computing: OpenCL, CUDA and DirectCompute
- 3-way SLI support

GeForce GTX 480 : 480 SP, 700/1401/1848MHz core/shader/mem, 384-bit, 1536MB, 250W TDP, US$499

GeForce GTX 470 : 448 SP, 607/1215/1674MHz core/shader/mem, 320-bit, 1280MB, 225W TDP, US$349

Note also we have like GLSL and OCL vec4 and other C++ libraries:
*GLM has GLSL strict compliance..
even with GMX experimental extensions we have SIMD implementations..
*DX SDK feb 2010 has XNAMATH 2.02 SIMD math library
also read:
http://www.gamasutra.com/view/feature/4248/designing_fast_crossplatform_simd_.php

HDR good maps:

http://www.hdrlabs.com/sibl/archive.html

Nvidia employess blogs:

http://timothylottes.blogspot.com/
http://jamesdolan.blogspot.com/
http://industrialarithmetic.blogspot.com/
http://castano.ludicon.com/blog/

http://twitter.com/castano

http://twitter.com/tmurray_cmpxchg

showing max cuda mem:

http://forums.nvidia.com/index.php?showtopic=102682 cuda maxmem

caustics patents:

US patent applications: 20090096788, 20090096789, and especially 20090128562,

The LLVM 2.7 binaries are available for testing:

http://llvm.org/pre-releases/2.7/pre-release1/

http://amnoid.de/tmp/clangtut/tut.html

http://lists.cs.uiuc.edu/pipermail/cfe-dev/2009-May/005167.html

http://synopsis.fresco.org/

Performance inconsistencies when testing various bit-counting methods

ubuntu cheat cube:119834-cheat-cube-ub

ie9 VML to SVG Migration Guide

windows phone 7:

*xna ctp 4.0 avaiable works with pc but only reach profile not hidef..

*unlocked image with all apps instructions on a blog..

*petzold samples and book excerpt avaiable..

*also sqlite port ->csharp-sqlite.wp

Windows 7 XP Mode now has support for CPUs without virtualization VT-D support..

Windows 7 SP1 virtualization news:

With Microsoft RemoteFX, users will be able to work remotely in a Windows Aero desktop environment, watch full-motion video, enjoy Silverlight animations, and run 3D applications," Microsoft's Max Herrmann writes, "All with the fidelity of a local-like performance when connecting over the LAN."

cuda will work with it? i.e. no need for compute cluster driver and also ogl,dx and interop support..

Q: Will RemoteFx support also OpenGL hardware acceleration which is the 3D high level API used by professional applications like CAD systems or medical applications ?

A: RemoteFX will support certain OpenGL applications. However, as the development of RemoteFX is still ongoing, it is too early to provide any specifics at this point.
Q: Are you plan to introduce RemoteFX also for Windows 7 because their are many scenarios where the remote system is not a server but a high end workstation ?
A: RemoteFX has been designed as a Windows Server capability to support the growing demand for multi-user, media-rich centralized desktop environments. Windows 7 will be supported as a virtual guest OS under Hyper-V.

Dynamic Memory is an improvement to Hyper-V which allows users to pool all available physical host memory together, and dynamically allocate it to virtual machines. In other words, if the workload changes, VMs can get access to extra memory without having to shut them down.

XNA forums:

Updated list of D3D12 suggestions

Unable to perform a recursive call with DirectCompute?

How to AttachBuffersAndPrecompute to ID3DX11FFT

RWStructuredBuffer counter

The IncrementCounter is faster than IterlockedAdd(Buffer[0], 1) in 4 times.

Gamefest 2010 presentations?

D3D11 / D2D Interoperativity

329M pairs/sec radix sort performance, 408M keys/sec - crushes CUDPP numbers

AppendStructuredBuffer driver bug?

How to debug DirectX 11 Compute Shaders?

Creating a Shared Surface with DXGI

atomic
I have some questions about RWStructuredBuffer:
1. How to copy hidden counter to system memory? CopyStructureCount
2. How to reset the counter to zero? last argument of OMSetRenderTargetsAndUnorderedAccessViews
3. Why the performance of this counter is much more than the performance of InterlockedAdd at the element buffer? (HD 5670)
The IncrementCounter is faster than IterlockedAdd(Buffer[0], 1) in 4 times.
How to AttachBuffersAndPrecompute to ID3DX11FFT?

http://gephi.org/

http://forums.xna.com/forums/t/49607.aspx
Thank you. I forgot about debug version of the D3DCSX. Debug message proved to be helpful. For the record: 1. The number of buffers attached must be exactly the same as in D3DX11_FFT_BUFFER_INFO. 2. The views MUST be created with the D3D11_BUFFER_UAV_FLAG_RAW flag (although it wasn't mentioned in documentation).

The Chrome dev channel release has support for an Open GL ES 2.0 interface

for Native Client. This is something we said we would do sometime last year.

When we consider it stable, documented etc. we will do more of an

announcement.

Google are announcing that NaCl now also supports x86-64 and ARM.

http://www.osnews.com/story/23021/Native_Client_Portability_Almost_Native_Graphics_Layer_Engine

NaCl_SFI:Adapting Software Fault Isolation to Contemporary CPU

Architectures

pnacl: Portable Native Client Executables

from GDC:

this are also graphics API translations:

Cider & Cedega: Direct3D on OpenGL

GameTree.tv: Direct3D on OpenGL ES

SwiftShader: DX Software Rendering (also WARP)

ANGLE Project: WebGL (OGL ES 2.0) on Direct3D

now we need GPGPU apis so:

cuda on opencl?

cuda on cal?

directcompute on opencl?

opencl on directcompute?

posted on opengl and cuda forums:

Questions to nvidia:

*Is Nvidia going to expose ext_gpu_shader_fp64 on GT2xx hardware with double precision or is for d3d11 hardware?

For example gtx275

AMD seems to support double precision on GLSL via doublepAMD even on 4850 cards..

Also is Nvidia with initial GL 4.0 drivers going to finally expose documentation for wgl_nv_dx_interop and have the shown at gtc texture writting and random access support?

via ext_image_load_store?

Please post PTX 1.5 and 2.0 documents..

Also I'm summing here things promised soon by Nvidia so let's see how much it takes before we get:

*cuda-gdb support for hardware debugging of OpenCL kernels

*cuda-gdb GPU debugger for Mac (with OpenCL support also)

Mac related:

Is mac 64 supported?

This package will work MAC OSX running 32/64-bit.

CUDA applications built in 32/64-bit (CUDA Driver API) is supported.

CUDA applications built as 32-bit (CUDA Runtime API) is supported.

(10.5.x Leopard and 10.6 SnowLeopard)

Note: x86_64 is not currently working for Leopoard or SnowLeopard

UDA applications built with the CUDA driver API can run as either 32/64-bit applications.

CUDA applications using CUDA Runtime APIs can only be built on 32-bit applications.

My mac notes:

nvcc matrixMul_kernel.cu matrixMulDrv.cpp -I../../common/inc/ ../../lib/libcutil_i386.a matrixMul_gold.cpp -Xlinker /usr/local/cuda/lib/libcuda.dylib

nvcc matrixMul_kernel.cu -c -m64

g++ matrixMul_gold.cpp matrixMulDrv.cpp -I../../common/inc/ -I$CUDA_INC_PATH -L$CUDA_LIB_PATH /usr/local/cuda/lib/libcuda.dylib ../../lib/libcutil_i386.a

para nvcc -m64 crea lib64 con copia de lib

nvcc -m64 deviceQueryDrv.cpp -I../../common/inc/ -I../../../shared/inc -Xlinker /usr/local/cuda/lib/libcuda.dylib

quita cut

nvcc defaults 32 bits

gcc defaults 64

g++

g++ deviceQueryDrv.cpp -I../../common/inc/ -I../../../shared/inc /usr/local/cuda/lib/libcuda.dylib -I$CUDA_INC_PATH

//#include

#define CU_SAFE_CALL_NO_SYNC(a) a

//CUT_EXIT(argc, argv);

export CUDA_BIN_PATH=/usr/local/cuda/bin

export CUDA_LIB_PATH=/usr/local/cuda/lib

export CUDA_INC_PATH=/usr/local/cuda/include

export PATH=$PATH:/usr/local/cuda/bin

raw data..

2010-03-18T19:32:00.000+01:00

games:
*metro 2033 and just cause 2 demo avaiable! (fermi launch titles?)
*assasins creed2 and bad company 2 this month also..
*Command & Conquer 4: Tiberian Twilight
*3d vision cd 1.23 has direct3d11 support! (so list support for d3d11 fermi supersleddemo)

iexplore 9 preview with direct2d directwrite support

*3D texture based separable convolution, extension of SDK example
code:
http://forums.nvidia.com/index.php?showtopic=163382

*bin format for fermi is similar ptx: post luebke on gpgpu-sim mailing list
one guy from pathscale says he has all info on this and other low level info presumably PTX 1.5,2.0 specs (bin format spec?) and also info for open source cuda driver for BSD etc..?
*gpgpu-3 papers avaiable!
http://www.ece.neu.edu/groups/nucar/GPGPU/GPGPU-FinalProgram.pdf
*CULA 1.2 avaiable with some eigenvectors/values stuff..

*"GPU Sample Sort" paper for the upcoming IPDPS 2010 conference?

It is possible to achieve much higher sorting rates for NV devices than with the Satish/CUDPP methods. You might be interested in our radix CUDA sorting results here at UVA. We demonstrate 480M pairs/sec, and 550M keys/sec on our GTX285 (with other devices evaluated as well). Interestingly enough, our keys-only results on the NV GT200 architecture are superior to the cycle-accurate sorting results from the (defunct) 32-core Larrabee.

Where is source?

http://www.cs.virginia.edu/~dgm4d/papers/RadixSortTR.pdf

Other sorting new papers:
*Revisiting Sorting for GPGPU Stream Architectures
"GPU Sample Sort" paper for the upcoming IPDPS 2010
N. Leischner, V. Osipov, and P. Sanders. GPU sample sort. In Proc. Int'l Parallel and Distributed Processing Symposium (IPDPS), to appear, 2010 (currently available at http://arxiv1.library.cornell.edu/abs/0909.5649).

*CUFFT does support streams... and seems has 3d ffts perf improvements of sc08 paper included so
apple fft code seems now work on Nvidia OpenCL but offer 2x-3x perf disadvantage vs cufft..

2d to 3d video conversion:
we have reald and other directshow plugin..
now:
arsoft sim 3d plus hd coming q2..
and powerdvd 10..
*TrueTheater™ Stabilizer
*TrueTheater™ 3D
*TrueTheater Noise Reduction

PowerDVD 10 Mark II: Consumers who purchase PowerDVD 10 Ultra 3D will receive a FREE UPGRADE that enables support of the Blu-ray 3D format and 2D to 3D conversion of video files. Available this summer.
Blu-ray 3D playback requires FREE "Mark II" upgrade which will be available soon.

lot of betas coming:
qt 4.7
intel compiler 12
vmware workstation 7.1
other march:
openrl
heaven 2.0

http://www.cs.utk.edu/~dongarra/WEB-PAGES/cscads-libtune-09/
1st CUDA Developers' Conference
http://www.smithinst.ac.uk/Events/CUDA2009
see
"Looking after the 7 dwarfs: numerical libraries / frameworks for GPUs" Mike Giles
also
"The Art of Performance Tuning for CUDA and Manycore Architectures"
David Tarjan (NVIDIA)
Kevin Skadron (U. Virginia)
Paulius Micikevicius (NVIDIA)

cudpp 1.1.1 svn has fermi support
cusp has amg geometric multigrid..
http://forums.nvidia.com/index.php?showtopic=163382&st=0&#entry1022104

See DirectX 9.0 on OpenGL ES 2.0 ->http://www.gametree.tv/ linux sdk

Coming in Spring 2010, the GameTree.tv Publishing SDK for Intel CE hardware will include the tools you need to optimize and debug your game for the GameTree.tv Gaming Platform, plus the ability to order Intel CE hardware.Developer Tools & Documentation     available     available
OpenGL ES 1.1 and 2.0
- Windows Game Development and Emulation
- Linux Desktop Runtime SDK     available     available
Direct3D® support
- Fixed-Function
- Shader Model 1.0 and 2.0 API
- Linux Desktop Emulation SDK     available     available
Debugging With Visual Studio     Coming March 2010     available
GameTree.tv Developer Forums     Coming Soon     available
Publish Games For Commercial Sale
Detailed Hardware Setup Documentation
Hardware Order Process
Developer Relations Support

fglrx 8.72.5 has ubuntu 10.4 support and opengl 3.2.97xx (opengl 3.3/40 partial support?)

Nvidia theater GDC notes:
dmm2
dmm2 free 1500 objects (star unleashed not uses more) max, has interop with physx and bullet adds
also directcompute and opencl simulation

shipping september october beta
still not ready plastic simulation and fracture mode.. calculates stress on volume so physical based break..
uses fp32 for gpu support and sse..

3d vision on unreal engine 3 shipping in april..
3d vision sdk soon code samples etc developer tricks for surround
surround recommends gfx400 in sli i "release 256 driver"

khrnos gdc sessions published has
info physics amd opencl sph and soft bodies no rigid bodies this is bullet work..
also fem simulation is dmm2 work..
no more interesting talk slides?: fft profiling for OpenCL by Nvidia employee

physxlab with destruction (precalculated) is beta now with unreal engine 3 integration

new unigine 2.0 this month on 26 has Linux support? and Windows OpenGL tesselation support with Fermi /5xxx cards?
nsight 480gtx 8marzo release
nexus 1.0 opengl and opencl analyzer not hardware debugger but like gdebugger gl+cl

fermi games just cause 2 (d3d10 only) metro 2033 (d3d11 optional)
http://nvidia.fullviewmedia.com/gdc2010/agenda.html
opengl 4.0 extensions viewer and glew in trunk support!
assasins cred2, badcompany 2
ati open 3d
nvidia 3dtv

cuda and visual studio:

QUOTE

- create empty cuda projects trough "project.."

You can just create an ordinary console project and then add .cu files to this project (see next point).

QUOTE

- add new .cu files through "add new item" (renaming c++ or txt in .cu files causes build errors)

If you add the CUDA build rules (Cuda.rules, distributed with the SDK) then VS will automatically detect the .cu files and pass them to nvcc to compile these to standard .obj files, the standard linker (link.exe) will then link these with the rest of your application's .obj files.

QUOTE

- doesn't highlight code in .cu files

See the instructions in (SDK_INSTALL_DIR)\C\doc\syntax_highlighting\visual_studio_8

QUOTE

- must copy a thousand times cutil64.dll around till it releases the program ...

Cutil is used to minimise code replication between the SDK samples. I'd advise understanding what you actually need and implementing it yourself. For example, most people only want the cuda safe call macros and you would be better off handling the error in a manner suitable for your app rather than just calling exit().

QUOTE

- must add a "thousand" new libraries not to cause build errors

By "thousand" do you mean one (cudart.lib)?! Ok, so you're using cutil so you need cutil64.lib too. But by definition using any library (and the CUDA API is provided through a library) you have to link with libraries.

QUOTE

- and even then its not sure if it runs

Can't help with that one (without more info).

I would advise the following.

Preparation:

Set up syntax highlighting
Set up Intellisense

Development:

Create a new, empty, console project (or you can use an existing project if you have one
Add your .c, .cpp and .cu files
Add the Cuda.rules
Modify C/C++ code generation to use /MT in release, /MTd in debug
Do the same for the Cuda code generation
Add cudart.lib to all configurations (i.e. release and debug)
Build, run, debug etc.

Proceedings of 24th IEEE International Parallel and Distributed Processing Symposium

gpu papers:

Session 2: Scientific Computing with GPUs Improving Numerical Reproducibility and Stability in Large-Scale Numerical Simulations on GPUs
Implementing the Himeno Benchmark with CUDA on GPU Clusters
Direct Self-Consistent Field Computations on GPU ClustersParallelization of Tau-Leap Coarse-Grained Monte Carlo Simulations on GPUs

A High-Performance Fault-Tolerant Software Framework for Memory on Commodity GPUs

Sort
High Performance Comparison-Based Sorting Algorithm on Many-Core GPUs
GPU Sample Sort
Highly Scalable Parallel Sorting

Session 9: Software Support for Using GPUs 26
Object-Oriented Stream Programming using Aspects
Optimal Loop Unrolling For GPGPU Programs
Speculative Execution on Multi-GPU Systems
Dynamic Load Balancing on Single- and Multi-GPU Systems

Fisheye Lens Distortion Correction on Multicore and Hardware Accelerator Platforms .. . . 37
Large-Scale Multi-Dimensional Document Clustering on GPU Clusters

Dynamically Tuned Push-Relabel Algorithm for the Maximum Flow Problem on CPU-GPU-Hybrid Platforms .

Optimization of Linked List Prefix Computations on Multithreaded GPUs Using CUDA

Inter-Block GPU Communication via Fast Barrier Synchronization

What's left in OpenGL 4.0? and more raw info..

2010-03-18T19:01:00.000+01:00

Somedays ago OGL 3.3 and 4.0 specs were published and a set of equivalent ARB extensions were put on registry where GLSL 3.3 and 4.0 were released.. now ogl 4.0 compatibility spec is +600 pages long core is 420 pages..

Other things:

*OGL 4.0 quick reference card

http://www.khronos.org/files/opengl4-quick-reference-card.pdf

*new glext.h and gl3.h updated
*glloader, glew on svn, and opengl extensions viewer for 3.3/4.0 already support it..
wait for sdl, smfl ..
*Waiting Fermi drivers on launch day..

remember all ARB extension no vendor or EXT..

No arb extensions included in 3.3/4.0 spec are:

GL_ARB_shading_language_include
GL_ARB_texture_compression_bptc

so HDR D3d11 texture format not required for ogl 4.0..

also lost is #include in shaders..

5xxx series include ogl4.0 emulating double on cpu? better with double-float emulation..

Last Nvidia found:

GL_EXT_shader_image_load_store

GL_EXT_vertex_attrib_64bit

and amd:

GL_EXT_shader_atomic_counters

are not found..

AMD 10.3 includes also first extension blend_func_extended..

GL_EXT_vertex_attrib_64bit adds vertex attribs:

so now fp64 is only for uniforms and passing not vertex attribs

remember no double rendertargets tex formats simlar to d3d11..

GL_EXT_shader_image_load_store allow write to random access to texes RWtexture3d
amd has amdx_random_access_target

ARB_blend_func_extended is called dual source blending in DX10, but got dropped in DX11..

We have tesselation shaders, dynamic shader linkage and compute interop with OCL..

still lacking vs d3d11 is:
*multi-threaded rendering:
remember only creation of resources in current drivers.. no parralel command list creation
is driver or hardware issue?

*random access load/store/atomic to texes->GL_EXT_shader_image_load_store amdx_random_access_target+GL_EXT_shader_atomic_counters RWtexture3d
*lacking atomic access to texs and mem barriers in fragment shaders: DeviceMemoryBarrier in d3d11
*GL_AMD_conservative_depth adds:

Conservative oDepth - This algorithm allows a pixel shader to compare the per-pixel depth value of the pixel shader with that in the rasterizer. The result enables early depth culling operations while maintaining the ability to output oDepth from a pixel shader.

So people on OGL forums are criticizing lack of:
*multi-threaded rendering
*shader binaries for avoid compilation preferibly crossvendor and plaform as DX IL DXBC (which is almost 100% compatible with ATI's IL)

*direct state access
* Epic fail for GL_ARB_sampler_objects as no glsl support..

I lack:

*ext_separate_shader_objects

The ability to separate program objects is only going to become increasingly more relevant.

*nv_texture_barrier

crossprocess texture sharing?

Support for programmable offsets in gather is there see 2x speedup in Fermi whitepaper and tesselation

fermi test would be good

fermi:

196.78 drivers support fermi..

full support for OGL 4.0 in fermi launch..

stocasthic transpareny i3d 2010 has fermi perf on this algorithm via ogl sample_shading 10.1 extension

GLwgl_dx_interop

GL_NVX_gpu_memory_info

GL_NV_gpu_program4_1

published then?

try openrl with opencl on fermi..

opencl drivers at fermi launch will have:

1.cuda 3.0 final

Fermi Direct3D 11 interoperability

Fermi HW Profiler support in OpenCL Visual Profiler

Complete BLAS lib, now with complex routines

cuda-gdb support for JIT compiled kernels

add

C++ Class Inheritance

C++ Template Inheritance

Unified interoperability API for Direct3D and OpenGL

OpenGL texture interoperability

2 with new opencl driver support:

*pragma unroll

*local atomics

*icd final

*d3d9/10/11 support

fxc interface has interface support but functions inside it how are called

see "CUDA_Developer_Guide_for_Optimus_Platforms"

http://www.stumblingahead.com/blog/?p=66 talking about tesselation soon..

2010 conferences GPU papers:

*PPOP

*GDC 2010

*I3D 2010

*GPGPU-3

*ASPLOS

MacroSS: Macro-SIMDization of Streaming Applications,

COMPASS: A Programmable Data Prefetcher Using Idle GPU Shaders,

"Investigating the Impact of Code Generation on Performance Characteristics of Integer Programs."

EUROGRAPHICS 2010

SIGGRAPH 2010

Interesting new/coming books:

*Game Programming Gems 8

*gpu computing gems 2010?

*Game Engine Gems 1, Volume One

*Programming Massively Parallel Processors: A Hands-

*GPU Pro: Advanced Rendering Techniques

*Multigrid Methods on GPUs

*Game Coding Complete, Third Edition

*Video Game Optimization

*Game Engine Architecture

*Real-Time Cameras

Programming Game AI by Example

Comments:

GL_ARB_shading_language_include-> glsl acepta #include i compilarshaderincludepaths fija <> paths de busqueda

GL_ARB_texture_compression_bptc

textures d3d 11 -> compressor incluido mejor offline

GL_ARB_blend_func_extended

permite usar dos salidas de fragment shader como color in i blend factors

mira ejemplo ventana color reflectiva en un paso usando con rops

GL_ARB_explicit_attrib_location->

fija en glsl explicito como las variables entre shaders se pasan e

GL_ARB_occlusion_query2

permite una boleana para si algo pasa o no

GL_ARB_sampler_objects

BindSampler( uint unit, uint sampler );

When a sampler object is bound to a texture unit, its state supersedes that

of the texture object bound to that texture unit. If the sampler name zero

is bound to a texture unit, the currently bound texture's sampler state

becomes active. A single sampler object may be bound to multiple texture

units simultaneously.

no cambia glsl a hlsl con tex.sampler

GL_ARB_shader_bit_encoding

con esto puedo usar fast float to int de spap paper kun zhou que coge bits

de float i haciendo cosas consige abs, float2int de valor ,etc..

To obtain signed or unsigned integer values holding the encoding of a

floating-point value, use:

genIType floatBitsToInt(genType value);

genUType floatBitsToUint(genType value);

Conversions are done on a component-by-component basis.

GL_ARB_texture_rgb10_a2ui

GL_ARB_texture_swizzle

GL_ARB_timer_query

GL_ARB_vertex_type_2_10_10_10_rev

GL_ARB_draw_indirect

compute interop

void DrawArraysIndirect(enum mode, const void *indirect);

nuevo buffer object

DRAW_INDIRECT_BUFFER

que hay bindeao

se usa como datos del num elementos etc..

que no

pues el puntero indirect se usa?..

GL_ARB_gpu_shader5

GL_ARB_gpu_shader_fp64

Should double-precision fragment shader outputs be supported?

RESOLVED: Not in this extension. Note that we don't have

double-precision framebuffer formats to accept such values.

GL_ARB_shader_subroutine

GL_ARB_tessellation_shader

GL_ARB_texture_buffer_object_rgb32

GL_ARB_transform_feedback2

1.transform feedback objects

2.pause and resume transform feedback

3.ability to draw primitives captured in transform feedback mode without querying the captured

primitive count

DrawTransformFeedback()

GL_ARB_transform_feedback3

unreal 3 news:

*palm webos and iphone support (on mac?)

*3d vision support

http://www.chw.net/2010/02/29-incomodas-preguntas-para-nvidia-sobre-gf100/

AMD Open Physics Initiative Expands Ecosystem with Free DMM for Game Production and Updated version of Bullet Physics

Apple adopts DirectX 11 GPUs, buys AMD Radeon HD 5750

apple news:

*99 dev program

*valve games to mac next month and monkey island 2 se..

*6core macpro next week (12 core?)Mac Pro 'hexacore' Xeon Core i7-980x coming Tuesday

reviews on anandtech 980 gulftown with aes today..

*amd 5750 imac in june? adds opengl 4.0 and ocl full support for mac..

so 10.6.4 will support amd 5xxx

*iphone 4.0 multitasking support

*10.6.3 this month?

CUDA:cuda-gdb gpu support and visual profilers,64 bit and efficient gl interop soon?

http://pasco2010.imag.fr/images/poster_pasco2010.pdf

http://unlimiteddetailtechnology.com/

roxio cienplayer 3d

CLyther = Python + OpenCL

amd open physics (free dmm 2.0 with ocl) and open stereo(qbf stereo for radeon?)

also eyefinity sdk coming soon..

ticker tape avaiable

pgi insider feb 2010 volume

http://www.pgroup.com/lit/articles/insider/v2n1a3.htm

says new fermi support and data region things..

XNA 4.0 winpho 7 tegra2 soon..

Yellow Dog Enterprise Linux for CUDA

http://ydl.net/cuda/iso/YDELforCUDA-6.2-20100302-DVD.iso download free for students

Jenkins Software Announces Data Mining Tool for Game Developers

As a further enhancement, AMD has developed new parallel GPU accelerated implementations of Bullet Physics’ Smoothed Particle Hydrodynamics (SPH) Fluids and Soft Bodies/Cloth. The new code written in OpenCL and Direct Compute will be contributed as open source.

OpenGL usage from an ISV perspective

intel gpa 3.0

nity Announces 3.0 Platform, Support For PS3, iPad, And Android

Valve Confirms Mac Versions Of Steam, Valve Game

http://www.raknet.net/echochamber

Erwin Coumans - SONY - Porting existing code to OpenCL

Ben Gaster AMD and Avi Shapira - Graphic Remedy - Debugging fluid dynamics on OpenCL

Greg Smith - NVIDIA - FFT and OpenCL Profiling

http://www.arm.com/community/software-enablement/google/solution-center-android.php

http://realworldtech.com/forums/index.cfm?action=detail&id=108017&threadid=108017&roomid=2

I can only say that at CAL level (and obviously OpenCL built upon CAL) there are numerous problems with multiple GPUs.

Definitely you're need one thread and one context per each GPU to make it working. But it itsn't enough because almost every CAL function isn't thread safe, thus calling calResMap() (which is the only to get access to local GPU memory) in one thread blocks all other threads/contexts.

And (as I've already wrote at these forums), OpenCL using calCtxWaitForEvent() function instead of CPU burning loop

while (calCtxIsEventDone(calCtx, e) == CAL_RESULT_PENDING);

to wait for GPU kernel completion.

But this calCtxWaitForEvent() also blocks every context currently running. This especially noticeable when there are different devices at system (like 5770+4770). So basically it's simply impossible to asynchronously work with multiple GPUs within single process.

All above things applies to windows version of CAL, never tried linux one.

Yup, and I use 1 thread per GPU too. So 1 thread, 1 context, 1 queue for each GPU. I tried other configurations but they weren't working (i.e. not running in parallel).

Why on HD4870 with 512 MB onboard RAM only 128 available to OpenCL ???

http://forums.amd.com/devforum/messageview.cfm?catid=390&threadid=128846&enterthread=y

MacroSS: Macro-SIMDization of Streaming Applications,

COMPASS: A Programmable Data Prefetcher Using Idle GPU Shaders,

"Investigating the Impact of Code Generation on Performance Characteristics of Integer Programs."

http://ctk-dev.sourceforge.net/

gmac

http://ctk-dev.sourceforge.net

http://code.google.com/p/fluidic/

http://otoy.com/

http://www.gameenginegems.com/

We're excited to announce a new addition to the Palm® webOS™ development platform: the webOS Plug-in Development Kit (PDK) lets developers extend their webOS applications by writing plug-ins in C or C++. The webOS PDK makes it easy for developers to leverage existing code and exposes new capabilities — including high-performance 3D graphics.

http://code.google.com/p/gyp/source/checkout

GPU computing toys!

2010-03-07T18:12:00.016+01:00

Hi I would like to release some lame but hopefully useful tools:
https://dl.dropbox.com/u/1416327/cld3d.rar

First OCL D3D interop headers and spec for Nvidia and AMD and a tool for checking current status:
the headers are in h
and are for d3d9,10,11 for NV and d3d9,10 for AMD..
#include for every d3d version and call initcld3d() in your code and voila you have the
d3d stuff..
if you #define INCAMD you have even amd functions included and can avoid amd headers..

with these I have complied four exes named cl_xx_interop which check d3d 9,9Ex,10 and 11..
they check extension reporting, try to create a shared context in some ways and then associate a d3d object and textures to ocl and aquire and release it prior to use..

Also cl_d3d10_interop build shows image formats avaiable to OpenCL images see next post..

Testing OCL-D3D11 interop
Checking D3D interop extensions support for platform: NVIDIA Corporation
nv D3D 9 interop extension: Found.
nv D3D 10 interop extension: Found.
nv D3D 11 interop extension: Found.

Using device: GeForce GTX 275
Enabling texture interop checks: image support is supported.
clGetDeviceIDsFromD3D11NV pointer: Found
and it works! (returns d3d associated ocl device)
clCreateFromD3D11BufferNV pointer: Found
clCreateFromD3D11Texture2DNV pointer: Found
clCreateFromD3D11Texture3DNV pointer: Found
clEnqueueAcquireD3D11ObjectsNV pointer: Found
clEnqueueReleaseD3D11ObjectsNV pointer: Found
Testing context creation with
no dev (clCreateContextFromType): OK.
dev info (getdeviceids): OK.
dev info (clGetDeviceIDsFromD3DNV CL_PREFERRED_DEVICES_FOR_D3D9_NV): OK.
Testing clCreateFromD3D11BufferNV: OK.
Testing aquire release stuff: Ok.. releasing it: Ok.
Testing clCreateFromD3D11Texture2DNV: OK.
Testing aquire release stuff: Ok.. releasing it: Ok.
Testing clCreateFromD3D11Texture3DNV: OK.
Testing aquire release stuff: Ok.. releasing it: Ok.

Also I contains a optd3d which displays the four optional d3d11 features (cap bits):

In my gtx 200 displays:

multithreaded comand lists: 0
multithreaded Concurrent Creates: 1
Double precision: 0
Compute Shader: 1

in ATI 5850 displays:

multithreaded comand lists: 0

multithreaded Concurrent Creates: 1

Double precision: 1

Compute Shader: 1

Anyway double prec is not working with loops..

This shows multithreaded command lists are still not supported by ATI (are this supposed to be a implementation issue or a hardware limitation..)

Equal to Nvidia and upcoming Fermi..

I include a CLinfo not mine but for checking CL info..

report.bat create a report.txt with the info of all this executables..

I also include 2dbench for cheking GDI in Windows 7 perf issues.. AMD will fix in Catalyst 10.4..

There is a high efficient matmul for CUDA and AMD cards and peakflops for AMD cards..

%
% compute C = A*B, A:mxk, B:kxn, C:mxn
%
% cubin file = ../method1/decuda_ldsb32_cudasm.cubin
% kernel function = method1_variant_sgemmNN
% use device: GeForce GTX 275
% m=n=k gpu_time (ms) flops (Gflops/s)
   32 0.044 1.391
   128 0.120 32.451
   224 0.194 107.870
   320 0.302 201.802
   416 0.445 301.033
   512 0.619 403.979
   608 1.277 327.914
   704 1.582 410.719
   800 2.618 364.210
   896 3.135 427.439
   992 4.401 413.123
   1088 6.014 398.868
   1184 6.981 442.860
   1280 8.751 446.365
   1376 10.911 444.746
   1472 13.403 443.262
   1568 16.377 438.470
   1664 18.901 454.051
   1760 22.437 452.594
   1856 25.820 461.218
   1952 31.233 443.566
   2048 33.317 480.229
   2144 39.834 460.841
   2240 44.989 465.337
   2336 51.643 459.765
   2432 56.514 474.095
   2528 64.183 468.859
   2624 72.540 463.923
   2720 79.686 470.387
   2816 85.826 484.626
   2912 96.003 479.094
   3008 108.801 465.942
   3104 121.579 458.181
   3200 126.446 482.699
   3296 138.522 481.473
   3392 153.544 473.440
   3488 168.797 468.268
   3584 177.873 482.085
   3680 193.298 480.227
   3776 212.160 472.675
   3872 229.596 470.947
   3968 246.403 472.280
   4064 260.086 480.699

clock 1620

% m=n=k gpu_time (ms) flops (Gflops/s)

32 0.040 1.516

128 0.108 36.044

224 0.173 120.900

320 0.265 229.925

416 0.393 341.338

512 0.535 467.090

608 1.107 378.021

704 1.371 474.163

800 2.270 420.030

896 2.751 486.983

992 3.804 477.992

1088 5.205 460.925

1184 6.003 514.983

1280 7.609 513.393

1376 9.396 516.463

1472 11.555 514.134

1568 14.145 507.666

1664 16.427 522.442

1760 19.387 523.784

1856 22.182 536.854

1952 26.860 515.777

2048 28.642 558.623

2144 34.530 531.627

2240 39.585 528.868

2336 44.440 534.292

2432 49.141 545.226

2528 55.274 544.429

2624 63.241 532.134

2720 68.451 547.592

2816 74.160 560.865

2912 82.945 554.516

3008 94.150 538.449

3104 104.581 532.653

3200 108.907 560.436

3296 119.277 559.158

3392 131.982 550.785

3488 146.003 541.376

3584 154.088 556.502

3680 166.307 558.166

3776 184.523 543.469

3872 198.692 544.196

3968 214.158 543.390

4064 223.720 558.838

it's a cubin so will not work in fermi
5850 stock

flopspeak.exe
Device 0
target 8
localRAM 1024 MB
uncachedRemoteRAM 2047 MB
cachedRemoteRAM 2047 MB
engineClock 725 MHz
memoryClock 1000 MHz
wavefrontSize 64
numberOfSIMD 18
doublePrecision 1
localDataShare 1
globalDataShare 1
globalGPR 1
computeShader 1
memExport 1
pitch_alignment 256
surface_alignment 4096
Device 0: execution time 7913.45 ms, achieved 2041.80 gflops
oc 950mhz

flopspeak.exe

engineClock 950 MHz
memoryClock 1000 MHz

Device 0: execution time 6039.35 ms, achieved 2675.40 gflops

matmul.exe 2048 2048 100

Device 0: execution time 1415.08 ms, achieved 1214.06 gflops
oc 950mhz
Device 0: execution time 1114.06 ms, achieved 1542.09 gflops

UPDATE 1:
Nvidia and ATI working together!
opencl.dll from ati sdk 2.01

Found 2 platform(s).
platform[01104BA0]: profile: FULL_PROFILE
platform[01104BA0]: version: OpenCL 1.0 CUDA 3.0.1
platform[01104BA0]: name: NVIDIA CUDA
platform[01104BA0]: vendor: NVIDIA Corporation
platform[01104BA0]: extensions: cl_khr_byte_addressable_store cl_khr_icd cl_khr_
gl_sharing cl_nv_d3d9_sharing cl_nv_d3d10_sharing cl_nv_d3d11_sharing cl_nv_comp
iler_options cl_nv_device_attribute_query cl_nv_pragma_unroll
platform[01104BA0]: Found 1 device(s).
   device[01104C08]: NAME: GeForce GTX 275
   device[01104C08]: VENDOR: NVIDIA Corporation
   device[01104C08]: PROFILE: FULL_PROFILE
   device[01104C08]: VERSION: OpenCL 1.0 CUDA
   device[01104C08]: EXTENSIONS: cl_khr_byte_addressable_store cl_khr_icd c
l_khr_gl_sharing cl_nv_d3d9_sharing cl_nv_d3d10_sharing cl_nv_d3d11_sharing cl_n
v_compiler_options cl_nv_device_attribute_query cl_nv_pragma_unroll cl_khr_glob
al_int32_base_atomics cl_khr_global_int32_extended_atomics cl_khr_local_int32_ba
se_atomics cl_khr_local_int32_extended_atomics cl_khr_fp64
   device[01104C08]: DRIVER_VERSION: 196.75

   device[01104C08]: Type: GPU
   device[01104C08]: EXECUTION_CAPABILITIES: Kernel
   device[01104C08]: GLOBAL_MEM_CACHE_TYPE: None (0)
   device[01104C08]: CL_DEVICE_LOCAL_MEM_TYPE: Local (1)
   device[01104C08]: SINGLE_FP_CONFIG: 0x3e
   device[01104C08]: QUEUE_PROPERTIES: 0x3

   device[01104C08]: VENDOR_ID: 4318
   device[01104C08]: MAX_COMPUTE_UNITS: 30
   device[01104C08]: MAX_WORK_ITEM_DIMENSIONS: 3
   device[01104C08]: MAX_WORK_GROUP_SIZE: 512
   device[01104C08]: PREFERRED_VECTOR_WIDTH_CHAR: 1
   device[01104C08]: PREFERRED_VECTOR_WIDTH_SHORT: 1
   device[01104C08]: PREFERRED_VECTOR_WIDTH_INT: 1
   device[01104C08]: PREFERRED_VECTOR_WIDTH_LONG: 1
   device[01104C08]: PREFERRED_VECTOR_WIDTH_FLOAT: 1
   device[01104C08]: PREFERRED_VECTOR_WIDTH_DOUBLE: 1
   device[01104C08]: MAX_CLOCK_FREQUENCY: 1404
   device[01104C08]: ADDRESS_BITS: 32
   device[01104C08]: MAX_MEM_ALLOC_SIZE: 229998592
   device[01104C08]: IMAGE_SUPPORT: 1
   device[01104C08]: MAX_READ_IMAGE_ARGS: 128
   device[01104C08]: MAX_WRITE_IMAGE_ARGS: 8
   device[01104C08]: IMAGE2D_MAX_WIDTH: 8192
   device[01104C08]: IMAGE2D_MAX_HEIGHT: 8192
   device[01104C08]: IMAGE3D_MAX_WIDTH: 2048
   device[01104C08]: IMAGE3D_MAX_HEIGHT: 2048
   device[01104C08]: IMAGE3D_MAX_DEPTH: 2048
   device[01104C08]: MAX_SAMPLERS: 16
   device[01104C08]: MAX_PARAMETER_SIZE: 4352
   device[01104C08]: MEM_BASE_ADDR_ALIGN: 256
   device[01104C08]: MIN_DATA_TYPE_ALIGN_SIZE: 16
   device[01104C08]: GLOBAL_MEM_CACHELINE_SIZE: 0
   device[01104C08]: GLOBAL_MEM_CACHE_SIZE: 0
   device[01104C08]: GLOBAL_MEM_SIZE: 919994368
   device[01104C08]: MAX_CONSTANT_BUFFER_SIZE: 65536
   device[01104C08]: MAX_CONSTANT_ARGS: 9
   device[01104C08]: LOCAL_MEM_SIZE: 16384
   device[01104C08]: ERROR_CORRECTION_SUPPORT: 0
   device[01104C08]: PROFILING_TIMER_RESOLUTION: 1000
   device[01104C08]: ENDIAN_LITTLE: 1
   device[01104C08]: AVAILABLE: 1
   device[01104C08]: COMPILER_AVAILABLE: 1
platform[0313A434]: profile: FULL_PROFILE
platform[0313A434]: version: OpenCL 1.0 ATI-Stream-v2.0.1
platform[0313A434]: name: ATI Stream
platform[0313A434]: vendor: Advanced Micro Devices, Inc.
platform[0313A434]: extensions: cl_khr_icd
platform[0313A434]: Found 2 device(s).
   device[0338CA70]: NAME: Intel(R) Core(TM) i7 CPU 920 @ 2.67GHz
   device[0338CA70]: VENDOR: GenuineIntel
   device[0338CA70]: PROFILE: FULL_PROFILE
   device[0338CA70]: VERSION: OpenCL 1.0 ATI-Stream-v2.0.1
   device[0338CA70]: EXTENSIONS: cl_khr_icd cl_khr_global_int32_base_atomic
s cl_khr_global_int32_extended_atomics cl_khr_local_int32_base_atomics cl_khr_lo
cal_int32_extended_atomics cl_khr_byte_addressable_store
   device[0338CA70]: DRIVER_VERSION: 1.0

   device[0338CA70]: Type: CPU
   device[0338CA70]: EXECUTION_CAPABILITIES: Kernel
   device[0338CA70]: GLOBAL_MEM_CACHE_TYPE: Read-Write (2)
   device[0338CA70]: CL_DEVICE_LOCAL_MEM_TYPE: Global (2)
   device[0338CA70]: SINGLE_FP_CONFIG: 0x7
   device[0338CA70]: QUEUE_PROPERTIES: 0x2

   device[0338CA70]: VENDOR_ID: 4098
   device[0338CA70]: MAX_COMPUTE_UNITS: 8
   device[0338CA70]: MAX_WORK_ITEM_DIMENSIONS: 3
   device[0338CA70]: MAX_WORK_GROUP_SIZE: 1024
   device[0338CA70]: PREFERRED_VECTOR_WIDTH_CHAR: 16
   device[0338CA70]: PREFERRED_VECTOR_WIDTH_SHORT: 8
   device[0338CA70]: PREFERRED_VECTOR_WIDTH_INT: 4
   device[0338CA70]: PREFERRED_VECTOR_WIDTH_LONG: 2
   device[0338CA70]: PREFERRED_VECTOR_WIDTH_FLOAT: 4
   device[0338CA70]: PREFERRED_VECTOR_WIDTH_DOUBLE: 0
   device[0338CA70]: MAX_CLOCK_FREQUENCY: 2698
   device[0338CA70]: ADDRESS_BITS: 32
   device[0338CA70]: MAX_MEM_ALLOC_SIZE: 536870912
   device[0338CA70]: IMAGE_SUPPORT: 0
   device[0338CA70]: MAX_READ_IMAGE_ARGS: 0
   device[0338CA70]: MAX_WRITE_IMAGE_ARGS: 0
   device[0338CA70]: IMAGE2D_MAX_WIDTH: 0
   device[0338CA70]: IMAGE2D_MAX_HEIGHT: 0
   device[0338CA70]: IMAGE3D_MAX_WIDTH: 0
   device[0338CA70]: IMAGE3D_MAX_HEIGHT: 0
   device[0338CA70]: IMAGE3D_MAX_DEPTH: 0
   device[0338CA70]: MAX_SAMPLERS: 0
   device[0338CA70]: MAX_PARAMETER_SIZE: 4096
   device[0338CA70]: MEM_BASE_ADDR_ALIGN: 32768
   device[0338CA70]: MIN_DATA_TYPE_ALIGN_SIZE: 128
   device[0338CA70]: GLOBAL_MEM_CACHELINE_SIZE: 64
   device[0338CA70]: GLOBAL_MEM_CACHE_SIZE: 65536
   device[0338CA70]: GLOBAL_MEM_SIZE: 1073741824
   device[0338CA70]: MAX_CONSTANT_BUFFER_SIZE: 65536
   device[0338CA70]: MAX_CONSTANT_ARGS: 8
   device[0338CA70]: LOCAL_MEM_SIZE: 32768
   device[0338CA70]: ERROR_CORRECTION_SUPPORT: 0
   device[0338CA70]: PROFILING_TIMER_RESOLUTION: 1
   device[0338CA70]: ENDIAN_LITTLE: 1
   device[0338CA70]: AVAILABLE: 1
   device[0338CA70]: COMPILER_AVAILABLE: 1
   device[04A30050]: NAME: Cypress
   device[04A30050]: VENDOR: Advanced Micro Devices, Inc.
   device[04A30050]: PROFILE: FULL_PROFILE
   device[04A30050]: VERSION: OpenCL 1.0 ATI-Stream-v2.0.1
   device[04A30050]: EXTENSIONS: cl_khr_global_int32_base_atomics cl_khr_gl
obal_int32_extended_atomics cl_khr_local_int32_base_atomics cl_khr_local_int32_e
xtended_atomics
   device[04A30050]: DRIVER_VERSION: CAL 1.4.556

   device[04A30050]: Type: GPU
   device[04A30050]: EXECUTION_CAPABILITIES: Kernel
   device[04A30050]: GLOBAL_MEM_CACHE_TYPE: None (0)
   device[04A30050]: CL_DEVICE_LOCAL_MEM_TYPE: Local (1)
   device[04A30050]: SINGLE_FP_CONFIG: 0x6
   device[04A30050]: QUEUE_PROPERTIES: 0x2

   device[04A30050]: VENDOR_ID: 4098
   device[04A30050]: MAX_COMPUTE_UNITS: 18
   device[04A30050]: MAX_WORK_ITEM_DIMENSIONS: 3
   device[04A30050]: MAX_WORK_GROUP_SIZE: 256
   device[04A30050]: PREFERRED_VECTOR_WIDTH_CHAR: 16
   device[04A30050]: PREFERRED_VECTOR_WIDTH_SHORT: 8
   device[04A30050]: PREFERRED_VECTOR_WIDTH_INT: 4
   device[04A30050]: PREFERRED_VECTOR_WIDTH_LONG: 2
   device[04A30050]: PREFERRED_VECTOR_WIDTH_FLOAT: 4
   device[04A30050]: PREFERRED_VECTOR_WIDTH_DOUBLE: 0
   device[04A30050]: MAX_CLOCK_FREQUENCY: 725
   device[04A30050]: ADDRESS_BITS: 32
   device[04A30050]: MAX_MEM_ALLOC_SIZE: 268435456
   device[04A30050]: IMAGE_SUPPORT: 0
   device[04A30050]: MAX_READ_IMAGE_ARGS: 0
   device[04A30050]: MAX_WRITE_IMAGE_ARGS: 0
   device[04A30050]: IMAGE2D_MAX_WIDTH: 0
   device[04A30050]: IMAGE2D_MAX_HEIGHT: 0
   device[04A30050]: IMAGE3D_MAX_WIDTH: 0
   device[04A30050]: IMAGE3D_MAX_HEIGHT: 0
   device[04A30050]: IMAGE3D_MAX_DEPTH: 0
   device[04A30050]: MAX_SAMPLERS: 0
   device[04A30050]: MAX_PARAMETER_SIZE: 1024
   device[04A30050]: MEM_BASE_ADDR_ALIGN: 4096
   device[04A30050]: MIN_DATA_TYPE_ALIGN_SIZE: 128
   device[04A30050]: GLOBAL_MEM_CACHELINE_SIZE: 0
   device[04A30050]: GLOBAL_MEM_CACHE_SIZE: 0
   device[04A30050]: GLOBAL_MEM_SIZE: 268435456
   device[04A30050]: MAX_CONSTANT_BUFFER_SIZE: 65536
   device[04A30050]: MAX_CONSTANT_ARGS: 8
   device[04A30050]: LOCAL_MEM_SIZE: 32768
   device[04A30050]: ERROR_CORRECTION_SUPPORT: 0
   device[04A30050]: PROFILING_TIMER_RESOLUTION: 1
   device[04A30050]: ENDIAN_LITTLE: 1
   device[04A30050]: AVAILABLE: 1
   device[04A30050]: COMPILER_AVAILABLE: 1
UPDATE 2:
DX formats included in optd3d

GPGPU Image support!

2010-03-07T17:44:00.003+01:00

1. D3D
In doc there is a table "Hardware Support for Direct3D 11 Formats"

Format(DXGI_FORMAT_*)	# Bits	Format Target
Format(DXGI_FORMAT_*)	# Bits	1	2	3	4	5	6	7	8	9	10	11	12	13	14	15	16	17	18	19	20	21	22	23	24	25	26	27	28	29	30	31	32	33	34	35	36	37	38
UNKNOWN	0	X																				X										X
R32G32B32A32_TYPELESS	128					X	X	X	X							X																X							X
R32G32B32A32_FLOAT	128	X	X		X	X	X	X	X	X	X			X		X	X	X	X				X	X								X	X	o	o	X	X		X
R32G32B32A32_UINT	128	X	X		X	X	X	X	X	X						X		X					X	X								X	X	o	o		X		X
R32G32B32A32_SINT	128	X	X		X	X	X	X	X	X						X		X					X	X								X	X	o	o		X		X
R32G32B32_TYPELESS	96					X	X	X	X							X																X							X
R32G32B32_FLOAT	96	X	X		X	X	X	X	X	X	o			o		X	o	o	o¹													X	X	X	o	X	X		X
R32G32B32_UINT	96	X	X		X	X	X	X	X	X						X		o														X	X	X	o		X		X
R32G32B32_SINT	96	X	X		X	X	X	X	X	X						X		o														X	X	X	o		X		X
R16G16B16A16_TYPELESS	64					X	X	X	X							X																X							X
R16G16B16A16_FLOAT	64	X	X			X	X	X	X	X	X			X		X	X	X	X				X	X								X	X	X	o	X	X	X	X
R16G16B16A16_UNORM	64	X	X			X	X	X	X	X	X			X		X	X	X	X				X	X								X	X	X	o	X	X		X
R16G16B16A16_UINT	64	X	X			X	X	X	X	X						X		X					X	X								X	X	X	o		X		X
R16G16B16A16_SNORM	64	X	X			X	X	X	X	X	X			X		X	X	X	X				X	X								X	X	X	o	X	X		X
R16G16B16A16_SINT	64	X	X			X	X	X	X	X						X		X					X	X								X	X	X	o		X		X
R32G32_TYPELESS	64					X	X	X	X							X																X							X
R32G32_FLOAT	64	X	X		X	X	X	X	X	X	X			X		X	X	X	X				X	X								X	X	X	o	X	X		X
R32G32_UINT	64	X	X		X	X	X	X	X	X						X		X					X	X								X	X	X	o		X		X
R32G32_SINT	64	X	X		X	X	X	X	X	X						X		X					X	X								X	X	X	o		X		X
R32G8X24_TYPELESS	64					X	X		X							X																X							X
D32_FLOAT_S8X24_UINT	64					X	X		X							X				X												X	X	X	o				X
R32_FLOAT_X8X24_TYPELESS	64					X	X		X	X	X	X		X	X	X																X					X		X
X32_TYPELESS_G8X24_UINT	64					X	X		X	X						X																X					X		X
R10G10B10A2_TYPELESS	32					X	X	X	X							X																X							X
R10G10B10A2_UNORM	32	X	X			X	X	X	X	X	X			X		X	X	X	X				X	X								X	X	X	o	X	X	X	X
R10G10B10A2_UINT	32	X	X			X	X	X	X	X						X		X					X	X								X	X	X	o		X		X
R10G10B10_XR_BIAS_A2_UNORM	32						X																									X						X	X
R11G11B10_FLOAT	32	X	X			X	X	X	X	X	X			X		X	X	X	X				X	X								X	X	X	o	X	X
R8G8B8A8_TYPELESS	32					X	X	X	X							X																X							X
R8G8B8A8_UNORM	32	X	X			X	X	X	X	X	X			X		X	X	X	X				X	X								X	X	X	o	X	X	X	X
R8G8B8A8_UNORM_SRGB	32					X	X	X	X	X	X			X		X	X	X	X													X	X	X	o	X	X	X	X
R8G8B8A8_UINT	32	X	X			X	X	X	X	X						X		X					X	X								X	X	X	o		X		X
R8G8B8A8_SNORM	32	X	X			X	X	X	X	X	X			X		X	X	X	X				X	X								X	X	X	o	X	X		X
R8G8B8A8_SINT	32	X	X			X	X	X	X	X						X		X					X	X								X	X	X	o		X		X
R16G16_TYPELESS	32					X	X	X	X							X																X							X
R16G16_FLOAT	32	X	X			X	X	X	X	X	X			X		X	X	X	X				X	X								X	X	X	o	X	X		X
R16G16_UNORM	32	X	X			X	X	X	X	X	X			X		X	X	X	X				X	X								X	X	X	o	X	X		X
R16G16_UINT	32	X	X			X	X	X	X	X						X		X					X	X								X	X	X	o		X		X
R16G16_SNORM	32	X	X			X	X	X	X	X	X			X		X	X	X	X				X	X								X	X	X	o	X	X		X
R16G16_SINT	32	X	X			X	X	X	X	X						X		X					X	X								X	X	X	o		X		X
R32_TYPELESS	32					X	X	X	X							X					X											X							X
D32_FLOAT	32					X	X		X							X				X												X	X	X	o				X
R32_FLOAT	32	X	X		X	X	X	X	X	X	X	X		X	X	X	X	X	X				X	X	X				X			X	X	X	o	X	X		X
R32_UINT	32	X	X	X	X	X	X	X	X	X						X		X					X	X	X	X	X	X	X	X	X	X	X	X	o		X		X
R32_SINT	32	X	X		X	X	X	X	X	X						X		X					X	X	X	X	X	X	X	X	X	X	X	X	o		X		X
R24G8_TYPELESS	32					X	X		X							X																X							X
D24_UNORM_S8_UINT	32					X	X		X							X				X												X	X	X	o				X
R24_UNORM_X8_TYPELESS	32					X	X		X	X	X	X		X	X	X																X					X		X
X24_TYPELESS_G8_UINT	32					X	X		X	X						X																X					X		X
R8G8_TYPELESS	16					X	X	X	X							X																X							X
R8G8_UNORM	16	X	X			X	X	X	X	X	X			X		X	X	X	X				X	X								X	X	X	o	X	X		X
R8G8_UINT	16	X	X			X	X	X	X	X						X		X					X	X								X	X	X	o		X		X
R8G8_SNORM	16	X	X			X	X	X	X	X	X			X		X	X	X	X				X	X								X	X	X	o	X	X		X
R8G8_SINT	16	X	X			X	X	X	X	X						X		X					X	X								X	X	X	o		X		X
R16_TYPELESS	16					X	X	X	X							X																X							X
R16_FLOAT	16	X	X			X	X	X	X	X	X			X		X	X	X	X				X	X								X	X	X	o	X	X		X
D16_UNORM	16					X	X		X							X				X												X	X	X	o				X
R16_UNORM	16	X	X			X	X	X	X	X	X	X		X	X	X	X	X	X				X	X								X	X	X	o	X	X		X
R16_UINT	16	X	X	X		X	X	X	X	X						X		X					X	X								X	X	X	o		X		X
R16_SNORM	16	X	X			X	X	X	X	X	X			X		X	X	X	X				X	X								X	X	X	o	X	X		X
R16_SINT	16	X	X			X	X	X	X	X						X		X					X	X								X	X	X	o		X		X
R8_TYPELESS	8					X	X	X	X							X																X							X
R8_UNORM	8	X	X			X	X	X	X	X	X			X		X	X	X	X				X	X								X	X	X	o	X	X		X
R8_UINT	8	X	X			X	X	X	X	X						X		X					X	X								X	X	X	o		X		X
R8_SNORM	8	X	X			X	X	X	X	X	X			X		X	X	X	X				X	X								X	X	X	o	X	X		X
R8_SINT	8	X	X			X	X	X	X	X						X		X					X	X								X	X	X	o		X		X
A8_UNORM	8					X	X	X	X	X	X			X		X	X	X	X				X	X								X	X	X	o	X	X
R9G9B9E5_SHAREDEXP	32					X	X	X	X	X	X			X		X																X
R8G8_B8G8_UNORM	16					X	X	X	X	X	X			X		X																X
G8R8_G8B8_UNORM	16					X	X	X	X	X	X			X		X																X
BC1_TYPELESS	4						X	X	X							X																X							X
BC1_UNORM	4						X	X	X	X	X			X		X																X							X
BC1_UNORM_SRGB	4						X	X	X	X	X			X		X																X							X
BC2_TYPELESS	8						X	X	X							X																X							X
BC2_UNORM	8						X	X	X	X	X			X		X																X							X
BC2_UNORM_SRGB	8						X	X	X	X	X			X		X																X							X
BC3_TYPELESS	8						X	X	X							X																X							X
BC3_UNORM	8						X	X	X	X	X			X		X																X							X
BC3_UNORM_SRGB	8						X	X	X	X	X			X		X																X							X
BC4_TYPELESS	4						X	X	X							X																X							X
BC4_UNORM	4						X	X	X	X	X			X		X																X							X
BC4_SNORM	4						X	X	X	X	X			X		X																X							X
BC5_TYPELESS	8						X	X	X							X																X							X
BC5_UNORM	8						X	X	X	X	X			X		X																X							X
BC5_SNORM	8						X	X	X	X	X			X		X																X							X
B8G8R8A8_TYPELESS	32					X	X	X	X							X																X							X
B8G8R8A8_UNORM	32	X	X			X	X	X	X	X	X			X		X	X	X	X													X	X	X	o	X	X	X	X
B8G8R8A8_UNORM_SRGB	32					X	X	X	X	X	X			X		X	X	X	X													X	X	X	o	X	X	X	X
B8G8R8X8_TYPELESS	32					X	X	X	X							X																X							X
B8G8R8X8_UNORM	32	X	X			X	X	X	X	X	X			X		X	X	X	X													X	X	X	o	X	X		X
B8G8R8X8_UNORM_SRGB	32					X	X	X	X	X	X			X		X	X	X	X													X	X	X	o	X	X		X
BC6H_TYPELESS	8						X	X	X							X																X							X
BC6H_UF16	8						X	X	X	X	X			X		X																X							X
BC6H_SF16	8						X	X	X	X	X			X		X																X							X
BC7_TYPELESS	8						X	X	X							X																X							X
BC7_UNORM	8						X	X	X	X	X			X		X																X							X
BC7_UNORM_SRGB	8						X	X	X	X	X			X		X																X							X

Buffer
Input Assembler Vertex Buffer
Input Assembler Index Buffer
Stream Output Buffer
Texture1D
Texture2D
Texture3D
TextureCube
Shader ld
Shader sample (any filter)
Shader sample_c (comparison filter)
Shader sample (mono 1-bit filter)
Shader gather4
Shader gather4_c
Mipmap
Mipmap Auto-Generation
RenderTarget
Blendable RenderTarget
Depth/Stencil Target
Raw UAV and SRV
Structured UAV and SRV
Typed UAV
UAV Typed Store
UAV Typed Load
UAV Atomic Add
UAV Atomic Bitwise Ops
UAV Atomic Cmp Store or Cmp Exch
UAV Atomic Exchange
UAV Atomic Signed Min or Max
UAV Atomic Unsigned Min or Max
CPU Lockable
4x Multisample RenderTarget
8x Multisample RenderTarget
Other Multisample Count RT
Multisample Resolve
Multisample Load
Display Scan-Out
Cast Within Bit Layout

A API for getting supported formats is ID3D11Device::CheckFormatSupport..

Would be good to write a program for checking the formats supported by AMD and Nvidia..

In CUDA 3.0 you have:

1, 2 or 4 components:
*Signed or unsigned 8-, 16- or 32-bit integers (18)
*16-bit floats (currently only supported through the driver (6)
API), or 32-bit floats
24 tex formats

For CUDA-GL interop (from forums):

works for FP textures:
XXXX = R,RG,RGB or RGBA
YY = 16 or 32

i.e. 8 FP formats

works for integer texes:
XXXX = R,G,RGB or RGBA
YY = 8,16 or 32
ZZ = I or UI

i.e. 24 FP formats

depth renderbuffers doesn't work I don't know if color renderbuffers work I assume yes at least for CUDA 3.0 final..

use:
glGenTextures(1,&tex);
glBindTexture(GL_TEXTURE_2D , tex);
glTexImage2D(GL_TEXTURE_2D , 0 , GL_XXXXYYF , width , height , 0 , GL_RGBA , GL_FLOAT , 0);
glTexParameteri(GL_TEXTURE_2D , GL_TEXTURE_MIN_FILTER , GL_NEAREST);
cudaGraphicsGLRegisterImage (&resource , tex , GL_TEXTURE_2D , cudaGraphicsMapFlagsNone);
for integer texes change to that:
glTexImage2D(GL_TEXTURE_2D , 0 , GL_XXXXYYZZ , width , height , 0 , GL_RGBA_INTEGER , GL_UNSIGNED_BYTE , 0);
notes:
Notice that it is important to set the minification filter to GL_NEAREST.
In conclusion, it looks like the cudaGraphicsGL interface is working for most formats, excluding normalized internal formats such as the commonly used GL_RGBA8 format.

cuda GL allows to use RGB texes altough CUDA seems not from DOC!

OCL DX interop for Nvidia:

------------------------------------------------------------------
   DXGI Format cl_channel_order cl_channel_type
   ------------------------------ ---------------- ---------------
   DXGI_FORMAT_R32G32B32A32_FLOAT CL_RGBA CL_FLOAT
   DXGI_FORMAT_R32G32B32A32_UINT CL_RGBA CL_UNSIGNED_INT32
   DXGI_FORMAT_R32G32B32A32_SINT CL_RGBA CL_SIGNED_INT32
   DXGI_FORMAT_R16G16B16A16_FLOAT CL_RGBA CL_HALF_FLOAT
   DXGI_FORMAT_R16G16B16A16_UNORM CL_RGBA CL_UNORM_INT16
   DXGI_FORMAT_R16G16B16A16_UINT CL_RGBA CL_UNSIGNED_INT16
   DXGI_FORMAT_R16G16B16A16_SNORM CL_RGBA CL_SNORM_INT16
   DXGI_FORMAT_R16G16B16A16_SINT CL_RGBA CL_SIGNED_INT16
   DXGI_FORMAT_R8G8B8A8_UNORM CL_RGBA CL_UNORM_INT8
   DXGI_FORMAT_R8G8B8A8_UINT CL_RGBA CL_UNSIGNED_INT8
   DXGI_FORMAT_R8G8B8A8_SNORM CL_RGBA CL_SNORM_INT8
   DXGI_FORMAT_R8G8B8A8_SINT CL_RGBA CL_SIGNED_INT8
   DXGI_FORMAT_R32G32_FLOAT CL_RG CL_FLOAT
   DXGI_FORMAT_R32G32_UINT CL_RG CL_UNSIGNED_INT32
   DXGI_FORMAT_R32G32_SINT CL_RG CL_SIGNED_INT32
   DXGI_FORMAT_R16G16_FLOAT CL_RG CL_HALF_FLOAT
   DXGI_FORMAT_R16G16_UNORM CL_RG CL_UNORM_INT16
   DXGI_FORMAT_R16G16_UINT CL_RG CL_UNSIGNED_INT16
   DXGI_FORMAT_R16G16_SNORM CL_RG CL_SNORM_INT16
   DXGI_FORMAT_R16G16_SINT CL_RG CL_SIGNED_INT16
   DXGI_FORMAT_R8G8_UNORM CL_RG CL_UNORM_INT8
   DXGI_FORMAT_R8G8_UINT CL_RG CL_UNSIGNED_INT8
   DXGI_FORMAT_R8G8_SNORM CL_RG CL_SNORM_INT8
   DXGI_FORMAT_R8G8_SINT CL_RG CL_SIGNED_INT8
   DXGI_FORMAT_R32_FLOAT CL_R CL_FLOAT
   DXGI_FORMAT_R32_UINT CL_R CL_UNSIGNED_INT32
   DXGI_FORMAT_R32_SINT CL_R CL_SIGNED_INT32
   DXGI_FORMAT_R16_FLOAT CL_R CL_HALF_FLOAT
   DXGI_FORMAT_R16_UNORM CL_R CL_UNORM_INT16
   DXGI_FORMAT_R16_UINT CL_R CL_UNSIGNED_INT16
   DXGI_FORMAT_R16_SNORM CL_R CL_SNORM_INT16
   DXGI_FORMAT_R16_SINT CL_R CL_SIGNED_INT16
   DXGI_FORMAT_R8_UNORM CL_R CL_UNORM_INT8
   DXGI_FORMAT_R8_UINT CL_R CL_UNSIGNED_INT8
   DXGI_FORMAT_R8_SNORM CL_R CL_SNORM_INT8
   DXGI_FORMAT_R8_SINT CL_R CL_SIGNED_INT8

OCL supported textures see my program:
clGetSupportedImageFormats
use

void getimageinfo(cl_context context,cl_mem_flags m,cl_mem_object_type te)

{

size_t num_entries; cl_image_format *image_formats;

cl_int status=clGetSupportedImageFormats (context,m,te,0,NULL,&num_entries);

if(status==CL_SUCCESS&&num_entries>0)

{

image_formats=(cl_image_format*)malloc(num_entries*sizeof(cl_image_format));

status=clGetSupportedImageFormats (context,m,te,num_entries,image_formats,NULL);

if(status==CL_SUCCESS)

{

int o,t;

int i,j;

cl_int orders[]={CL_R, CL_A,CL_INTENSITY, CL_LUMINANCE,CL_RG, CL_RA,CL_RGB,CL_RGBA,CL_ARGB, CL_BGRA};

char *or[]={"CL_R", "CL_A","CL_INTENSITY", "CL_LUMINANCE","CL_RG", "CL_RA","CL_RGB","CL_RGBA","CL_ARGB", "CL_BGRA"};

cl_int types[]={

CL_SNORM_INT8 , CL_SNORM_INT16, CL_UNORM_INT8, CL_UNORM_INT16, CL_UNORM_SHORT_565, CL_UNORM_SHORT_555, CL_UNORM_INT_101010,CL_SIGNED_INT8,

CL_SIGNED_INT16, CL_SIGNED_INT32, CL_UNSIGNED_INT8, CL_UNSIGNED_INT16, CL_UNSIGNED_INT32, CL_HALF_FLOAT, CL_FLOAT};

char * tt[]={"CL_SNORM_INT8" ,"CL_SNORM_INT16","CL_UNORM_INT8","CL_UNORM_INT16","CL_UNORM_SHORT_565","CL_UNORM_SHORT_555","CL_UNORM_INT_101010",

"CL_SIGNED_INT8","CL_SIGNED_INT16","CL_SIGNED_INT32","CL_UNSIGNED_INT8","CL_UNSIGNED_INT16","CL_UNSIGNED_INT32","CL_HALF_FLOAT","CL_FLOAT"};

for(i=0; i

{

for(j=0; j

{

if(image_formats[i].image_channel_order==orders[j])

o=j;

}

for(j=0; j

{

if(image_formats[i].image_channel_data_type==types[j])

t=j;

}

printf("Format %d: %s, %s\n",i,or[o],tt[t]);

}

free(image_formats);

}

AMD and Nvidia return same for all args cl_mem_flags flags read or write only and cl_mem_object_type image_type set to 2d or 3d.. perhaps 3d write could report 0?

Nvidia:

Format 0: CL_R, CL_FLOAT
Format 1: CL_R, CL_HALF_FLOAT
Format 2: CL_R, CL_UNORM_INT8
Format 3: CL_R, CL_UNORM_INT16
Format 4: CL_R, CL_SNORM_INT16
Format 5: CL_R, CL_SIGNED_INT8
Format 6: CL_R, CL_SIGNED_INT16
Format 7: CL_R, CL_SIGNED_INT32
Format 8: CL_R, CL_UNSIGNED_INT8
Format 9: CL_R, CL_UNSIGNED_INT16
Format 10: CL_R, CL_UNSIGNED_INT32
Format 11: CL_A, CL_FLOAT
Format 12: CL_A, CL_HALF_FLOAT
Format 13: CL_A, CL_UNORM_INT8
Format 14: CL_A, CL_UNORM_INT16
Format 15: CL_A, CL_SNORM_INT16
Format 16: CL_A, CL_SIGNED_INT8
Format 17: CL_A, CL_SIGNED_INT16
Format 18: CL_A, CL_SIGNED_INT32
Format 19: CL_A, CL_UNSIGNED_INT8
Format 20: CL_A, CL_UNSIGNED_INT16
Format 21: CL_A, CL_UNSIGNED_INT32
Format 22: CL_RG, CL_FLOAT
Format 23: CL_RG, CL_HALF_FLOAT
Format 24: CL_RG, CL_UNORM_INT8
Format 25: CL_RG, CL_UNORM_INT16
Format 26: CL_RG, CL_SNORM_INT16
Format 27: CL_RG, CL_SIGNED_INT8
Format 28: CL_RG, CL_SIGNED_INT16
Format 29: CL_RG, CL_SIGNED_INT32
Format 30: CL_RG, CL_UNSIGNED_INT8
Format 31: CL_RG, CL_UNSIGNED_INT16
Format 32: CL_RG, CL_UNSIGNED_INT32
Format 33: CL_RA, CL_FLOAT
Format 34: CL_RA, CL_HALF_FLOAT
Format 35: CL_RA, CL_UNORM_INT8
Format 36: CL_RA, CL_UNORM_INT16
Format 37: CL_RA, CL_SNORM_INT16
Format 38: CL_RA, CL_SIGNED_INT8
Format 39: CL_RA, CL_SIGNED_INT16
Format 40: CL_RA, CL_SIGNED_INT32
Format 41: CL_RA, CL_UNSIGNED_INT8
Format 42: CL_RA, CL_UNSIGNED_INT16
Format 43: CL_RA, CL_UNSIGNED_INT32
Format 44: CL_RGBA, CL_FLOAT
Format 45: CL_RGBA, CL_HALF_FLOAT
Format 46: CL_RGBA, CL_UNORM_INT8
Format 47: CL_RGBA, CL_UNORM_INT16
Format 48: CL_RGBA, CL_SNORM_INT16
Format 49: CL_RGBA, CL_SIGNED_INT8
Format 50: CL_RGBA, CL_SIGNED_INT16
Format 51: CL_RGBA, CL_SIGNED_INT32
Format 52: CL_RGBA, CL_UNSIGNED_INT8
Format 53: CL_RGBA, CL_UNSIGNED_INT16
Format 54: CL_RGBA, CL_UNSIGNED_INT32
Format 55: CL_BGRA, CL_UNORM_INT8
Format 56: CL_BGRA, CL_SIGNED_INT8
Format 57: CL_BGRA, CL_UNSIGNED_INT8
Format 58: CL_ARGB, CL_UNORM_INT8
Format 59: CL_ARGB, CL_SIGNED_INT8
Format 60: CL_ARGB, CL_UNSIGNED_INT8
Format 61: CL_INTENSITY, CL_FLOAT
Format 62: CL_INTENSITY, CL_HALF_FLOAT
Format 63: CL_INTENSITY, CL_UNORM_INT8
Format 64: CL_INTENSITY, CL_UNORM_INT16
Format 65: CL_INTENSITY, CL_SNORM_INT16
Format 66: CL_LUMINANCE, CL_FLOAT
Format 67: CL_LUMINANCE, CL_HALF_FLOAT
Format 68: CL_LUMINANCE, CL_UNORM_INT8
Format 69: CL_LUMINANCE, CL_UNORM_INT16
Format 70: CL_LUMINANCE, CL_SNORM_INT16

AMD:

Format 0: CL_RGBA, CL_UNORM_INT8
Format 1: CL_RGBA, CL_UNORM_INT16
Format 2: CL_RGBA, CL_SIGNED_INT8
Format 3: CL_RGBA, CL_SIGNED_INT16
Format 4: CL_RGBA, CL_SIGNED_INT32
Format 5: CL_RGBA, CL_UNSIGNED_INT8
Format 6: CL_RGBA, CL_UNSIGNED_INT16
Format 7: CL_RGBA, CL_UNSIGNED_INT32
Format 8: CL_RGBA, CL_HALF_FLOAT
Format 9: CL_RGBA, CL_FLOAT
Format 10: CL_BGRA, CL_UNORM_INT8

OCL-GL interop I don't know:
for Nvidia is either the 70 above or the CUDA-GL supported formats or the GL equivalent of CUDA interop.. I suspect the CL_RGB ones supported..
for AMD either the CL image ones or CAL DX interop ones
I suspect RGB formats
AMD CAL:

CAL has textures exposed and CAL DX interop would be good to explore..

CUDA 3.0 and Nexus in VS 2010, CUDA on FreeBSD 8.0 and much more!

2010-03-05T17:54:00.001+01:00

Interesting Nvidia threads:
1.Nexus: Unofficial Nexus / Visual Studio 2010 integration
http://forums.nvidia.com/index.php?showtopic=161096
->enables also cuda 3.x compiling with vs 2010!
this is awesome brings vs2010rc+cuda 3.0+nexus and also project templates for cuda and nexus apps!

It patches the vsvars32.bat file to read "Setting environment for using Microsoft Visual Studio 2008 x86 tools" instead of "Setting environment for using Microsoft Visual Studio 2010 x86 tools" to get around nvcc's Visual C++ version detection; otherwise it fails with this message: "nvcc fatal : nvcc cannot find a supported cl version. Only MSVC 8.0 and MSVC 9.0 are supported". It also creates the vcvarsamd64.bat file to make 64-bit builds work, or otherwise nvcc files with "nvcc fatal : Visual Studio configuration file '(null)' could not be found" (see this thread).

Nexus news: no DX9 and OGL in initial release:

DX10 is currently supported and DX11 will be available in the Beta 2 release which is scheduled for early March.DX10 and DX11 are the graphics APIs of choice for 1.0. Interesting that you feel OpenGL is favored, as full support for OpenGL won't be in the 1.0 release. The Beta 1 was focused on Compute - The Beta 2 will be released just before GDC and will bring full DX10, DX11 debugging and profiling into Visual Studio. This is a dream for game and graphics developers - perfhud on steroids. OpenGL support will come out *sometime* in late 2010. Pro version: (paid version) In addition to premium support, platform analysis (cpu+gpu correlated timeline) and advanced debugging capabilities will be available only in the pro version.

2.Feature Request: Support simultaneous native and CUDA debugging

I have noticed that it is not possible to debug CUDA and native code simultaneously on the same Visual Studio instance. I tried starting debugging through the Start CUDA Debugging option and then attaching the native debugger to the running process (inserting a 10 second sleep at the start of the program helped make this easier), but as soon as a breakpoint on a CUDA kernel is hit Visual Studio freezes.

I've been able to debug native and CUDA code on the same process simultaneously by having two Visual Studio instances open, and it works very well. I think it would be very valuable to be able to step through native code and device code on the same session, much like mixed debugging works with .NET.

2. Eclipse Plugin for CUDA and QT development

http://forums.nvidia.com/index.php?showtopic=160564

we developed a plugin for Eclipse, which comortably allows CUDA and QT development. It provides three toolchains, which can be used to compile CUDA and/or QT sources.

Features include:

- Error Parsing
- Dependency Calculation
- Automatic invocation of all tools
- ...

http://www.ai3.uni-bayreuth.de/software/eclipsecudaqt/index.php

Fastest CUDA reduction code to date! (following news last week of fastest matmul for GT200 and for AMD in C like language)
http://forums.nvidia.com/index.php?showtopic=160196

My simple but speedy reduction code (runs 106.4GB/s on GTX 295), 106.4/111.9=95.1% to the peak bandwidth good reduction code 5ms for 150m integers.
Testing with different input size I can see that your code is significantly slower if size is less than 16M, about the same speed with 32M and faster with more than 32M on the GTX 260.
Seems my code can beat SDK recution on every input size provided that the parameters M and K are properly choosed. Here is a detailed result for different M and K on different input sizes, and the performance for the SDK reduction with the same sizes are also listed.
gtx 295 1 core.
[size=~512K]
My code (M=240, N=64, K=34): 60.3GB/s (23.3% faster)
SDK reduction (size=1<<19): 48.9GB/s
[size=~1M]
My code (M=240, N=64, K=69): 76.0GB/s (15.8% faster)
SDK reduction (size=1<<20): 65.6GB/s
[size=~2M]
My Code (M=240, N=64, K=137): 86.6GB/s (9.3% faster)
SDK reduction (size=1<<21): 79.2GB/s
[size=~4M]
My Code (M=240, N=64, K=273): 94.2GB/s (5.8% faster)
SDK reduction (size=1<<22): 89.0GB/s
[size=~8M]
My Code (M=240, N=64, K=546): 99.5GB/s (5.0% faster)
SDK reduction (size=1<<23): 94.8GB/s
[size=~16M]
My Code (M=240, N=64, K=1092): 103.1GB/s (5.5% faster)
SDK reduction (size=1<<24): 97.7GB/s
[size=~32M]
My Code (M=240, N=64, K=2184): 104.9GB/s (6.3% faster)
SDK reduction (size=1<<25): 98.7GB/s
[size=~64M]
My Code (M=480, N=64, K=2184): 105.8GB/s (7.3% faster)
SDK reduction (size=1<<26): 98.6GB/s
[size=~128M]
My Code (M=720, N=64, K=2912): 106.4GB/s (9.1% faster)
SDK reduction (size=1<<27): 97.5GB/s

*cuda_wrapper
The CUDA wrapper library provides means for an efficient resource sharing and resource protection on multi-user GPU clusters.It implements the following functionality:1) Virtualization of the physical GPU devices2) Ensuring NUMA affinity for GPUs
http://sourceforge.net/projects/cudawrapper/

It's suposed to show that allocating resources and freeing and allocating new ones show mem intact as left by last object so no privacy in this sense!

*Seems depth buffers/renderbuffers are not supported by GL interop in CUDA 3.1

so no tex where format is GL_DEPTH_COMPONENT32 in
cudaGraphicsGLRegisterImage (&resource, tex , GL_TEXTURE_2D , cudaGraphicsMapFlagsNone);
Also remember this is a post showing current color formats seems to be R,RG,RGB,RGBA in float,float16 and uint8 formats more or less similar to published OpenCL DX interop formats.. Good to write a tool that writes current formats on OpenCL as there is a function for it for seeing if DX interop disables some formats on Nnvida hard at least.

*cuda on freebsd 8.0!
Inter-kernel communication is not supported under pain of me glaring at you really hard.
The recipe is:
FreeBSD 8.0 + NVidia driver 195.22 + CUDA 3.0
Also linprocfs and linsysfs should be mounted

uname -a
FreeBSD av429635.oops 8.0-RELEASE FreeBSD 8.0-RELEASE #0: Sun Feb 7 17:30:12 MSK 2010 root@av429635.oops:/usr/src/sys/i386/compile/ALECN80 i386

mount /compat/linux/proc/
mount /compat/linux/sys/

Well, not exactly "CUDA work on FreeBSD" - it's Linux program that use Linux libs under Linuxlator on FreeBSD that works.
Also I didn't try to compile Cuda programs yet, I've just launched programs pre-compiled on Linux (Debian)

also seems 190 drivers showed info with cudadeviceinfodrv but not created context.

*ipad: cpu is cortex a8 1ghz 1 core (same as 3gs) but stripped
gpu is power sgx variant.. but slow for pixel resolution (perhaps is 535 or 530 worse or 540 I doubt so)
so tegra2 is a lot better in cpu and gpus side seems
perhaps flash does`'t work by custom gpu altough using PowerVr IP.. as OMAP3 or 4 has been shown with Flash 10.1 video acceleration in MWC..

*"Optimus Works Perfectly With Intel Wireless Display (WiDi)"

A perfect notebook must have it!
Still in WiDi mode you lose 3d 120hz via HDMI and also it hasn't HDCP so no Bluray..
I hope next Widi has HDCP and also HDMI 1.4 so 3d also work but will require double bandwith and seems to stress current wifi..
My question is with Optimus where Nvidia sends to Intel IGP if it have teoretically 3d screen built in with 120hz will work Nvidia 3d Vision and what about if it has DVI dual link output and I connect to 3d 120hz display? I suspect the answer is the same at least the technical hurdles seem to be.. and I think correctly is hard as is a PCI Express transfer and seems 1Gbytes/s is currently used for 60hz? so at least this would put more streess but entirelly doable if Intel IGP recognizes special 120hz modes of LCD and acts acordingly..
Also all requires Windows 7 (Optimus requires as it has two graphic drivers different IHVs at the same time and Widi seems to require 7 x64)
Also will work with macbookpro optimus laptops widi? It would require support from Intel as is using some MyWifi tech so must see.. Perhaps Apple waits for LightStage optical video outputs no wireless tech..
A dream notebook in graphics must have a d3d11 with 3d (so Fermi) also with standard 3d outputs so HDMI 1.4 outputs and 3d 120hz builtin screen and Optimus and possibly Widi better at least with HDCP support.. let's wait how long it takes to arrive to that I hope least than a year..
*Optimus has nvgpustateviewer tool that shows if Nvidia GPU is activated or not. Where to download?

intel widi no hdcp so no bluray viewing of course not 3d but optimus compatible now
similar a PERFECT 3D PROJECTOR
*720p at least
*Broad 3d support: 3d via hdmi 1.4, dlplink, 3d vision compatible
*hdcp support
so it can output 3d vision, PS3 3d games (HDMI 1.4) and Bluray 3d(HDMI 1.4+hdcp)
Now Acer and Viewsonic support all but HDMI 1.4.. so no ps3+bd3d support..

current projector is hdcp so bluray and 3d 120hz via duallink dvi or hdmi so no hdmi 1.4 3d spec support for projecting ps3 games bluray players output,etc..

iz3d 1.11 coming soon using catalyst 10.3 3d hooks for better multimon support (3d vision surround?) and possibly crossfire and also bringing d3d10 support for games

HYDRA in AMD chipset shown with GTX275+5870 are using improved Hydra 1.5 driver with better Mix mode.. it would be interesting to see how perf and compatibilty improves over time (i.e. see the hardware potential once all software issues remain solved/tuned..)

Regarding Widi:
"The software drivers that work with Intel® Wireless Display only apply to Microsoft Windows 7 64-bit*.
Intel® PROSet/Wireless WiFi Connection Utility for Windows 7 64-Bit for Intel Wireless Display
Requires special Proset driver:
Wireless Driver:
Drivers and management software for Microsoft Windows 7 64-bit OS*.
NOTES:
http://www.intel.com/support/wireless/wtech/iwd/sb/CS-031109.htm

-The ZIP file is provided with Intel® My WiFi Technology enabled.
-Intel® My WiFi Technology has the following requirements:
-Intel® Centrino® Ultimate-N 6300, Intel® Centrino® Advanced-N 6200, Intel® Centrino® Advanced-N + WiMAX 6250, Intel® WiFi Link 1000, Intel® WiFi Link 5300, or Intel® WiFi Link 5100
-Minimum of Intel® PROSet/Wireless WiFi Connection Utility 13.0.0.0 on Microsoft Windows 7*
NOTE: Intel® Wireless Display requires one of the following products:
-Intel® Centrino® Ultimate-N 6300
-Intel® Centrino® Advanced-N 6200
-Intel® Centrino® Advanced-N+WiMAX 6250
NOTE: Features removed from this version:
Wake on Wireless LAN is not present in this version of the application.
the Intel® My WiFi Technology application is not supported for Windows Vista. This feature is available on Windows 7 only.
For the latest driver for the Intel® PROSet/Wireless WiFi Connection Utility (for Intel® Centrino® Advanced-N 6200). Intel recommends that you use the latest drivers for best performance.

intel media sdk 1.5rc

See http://software.intel.com/en-us/articles/intel-media-software-development-kit-intel-media-sdk/

Its going to support Intel Media SDK H.264 MVC codec of 3D Bluray either via GPU video processors or if they not support via optimized multithreaded SSE enabled code..
Also similar to CPU h.264 encoding support is going to be a 3D MVC encoder?

*Shader Model 5 (see Shader Model 5) vs OpenCL kernels:

Common (more or less):
Doubles with denorms
Reduced-precision reciprocal
Shader conversion instructions - fp16 to fp32 and vice versa
Structured buffer, which is a new type of buffer containing structured elements.
of which some things not present in OpenCL kernels
Resinfo on buffers
Count bits set instruction
Find first bit set instruction
Carry/Overflow handling
Bit reversal instructions for FFTs
Conditional Swap intrinsic
Also Dispatch indirect
remember it's about reading from GPU buffer the grid size to launch still requires CPU to launch the kernel..
but I doesn't require reading about 3 integers of grid which being so much size the PCI transaction still would be 1k? and add a lot of latency and add a CPU GPU synch point.. remember still no block size
at runtime kernel must be compiled for a fixed block size.
it' a evolution(?) of Draw Indirect - Direct3D 10 implements DrawAuto, which takes content (generated by the GPU) and renders it (on the GPU). Direct3D 11 generalizes DrawAuto so that it can be called by a Compute Shader using DrawInstanced and DrawIndexedInstanced.

* gDEBugger CL is a new and exciting product; it brings all of gDEBugger's Debugging and Profiling capabilities to the OpenCL developer's world. gDEBugger CL, now in beta testing, supports all OpenCL implementations on Windows, Mac OS X and Linux. The upcoming gDEBugger iPhone version includes on-device debugging and profiling abilities, running in real-time and letting developers optimize their game on the actual iPhone device. gDEBugger iPhone displays invaluable inside information such as iPhone's GPU, CPU, graphic driver and operating system performance counters.

New in Nvidia 196.75 drivers!

2010-03-03T03:09:00.001+01:00

opencl now has d3d interop:
cl_nv_d3d9_sharing cl_nv_d3d10_sharing cl_nv_d3d11_sharing
cl gl get info khr now?

new ogl extensions (dx10.1 and 11)

nv ones:
GL_NV_vertex_attrib_64bit
GL_NVX_gpu_memory_info
GL_NV_gpu_shader5
GL_NV_gpu_program5
GL_NV_gpu_program_fp64
GL_NV_gpu_program4_1

present in cat 10.3 also
GL_EXT_gpu_shader_fp64
GL_EXT_texture_buffer_object_rgb32
GL_EXT_tessellation_shader
GL_EXT_shader_subroutine
GL_EXT_gpu_shader5
GL_EXT_texture_compression_bptc

nv present:
GL_EXT_draw_indirect
GL_EXT_shader_image_load_store
GL_EXT_vertex_attrib_64bit
GL_EXT_transform_feedback3
GL_EXT_transform_feedback2

amd present:
GL_EXT_shader_atomic_counters
GL_AMD_conservative_depth

good all 195 drivers nv_
texture_buffer_object_rgb32
tessellation_shader
shader_subroutine

migrated to ext now all functionality is ext_ stuff!
nv lacks ext_shader_atomic_counters
and amd well the most imp:
GL_EXT_draw_indirect
GL_EXT_shader_image_load_store
and perhaps
GL_EXT_vertex_attrib_64bit

GPU computing in a browser, and other news..

2010-03-03T03:08:00.002+01:00

Seems just after WebGL brings GPU graphics APIs to browsers (OpenGL ES) people are
asking about WebCL bringing similar to OpenCL to browser..

1.One way is via webgl using the old GPGPU using Graphics API tricks..

see here for a matmul running on GPU->learningwebgl.com

2.In the meantime seems DirectX11 plugin coming to browsers
"DirectX 11 3D Games coming to Browsers with Vision Engine 8"
this should allow compute shaders on a browser! altough not mentioned only tesselation is mentioned..

3. Search Jetpack CUDA plugin..
http://mozillalabs.com/jetpack/2010/01/25/elevating-javascript-performance-through-gpu-power/

4.nacl gpu computing opencl. Should be possible with Google NaCl for Chrome via OpenCL,CUDA ,etc.. ?

WebGL support for QtWebKit:

Related is Webkit in QT trunk has WebGL support working on N900.
Also a new backend for audio/video elements using the Qt Multimedia framework and initial work on
WebGL support.

NOTES:

* Works only when accelerated composition is not enabled in compilation
* Added --webgl command line switch to QGVLauncher, added toggle button to
QtLauncher
Why GraphicsLayer (accelerated composition layer) doesn't handle WebGL? Missing
methods:

* setContentsToGraphicsContext3D
* setGraphicsContext3DNeedsDisplay

I thinks adding support for content caching is subtask for this. WebGL support
can be tested by compiling with WTF_USE_ACCELERATED_COMPOSITING=0.

WebGL and Khronos stuff:
https://cvs.khronos.org/svn/repos/registry/trunk/public/webgl/sdk/tests/webgl-conformance-tests.html
http://www.khronos.org/webgl/wiki/Debugging
https://cvs.khronos.org/svn/repos/registry/trunk/public/index.php

opera 10.5 released: has faster javascript to date until minefield picks new *monkey
now concentrate on webgl..

Related is I forgot to blog about another amazing development, the DirectWrite & Direct2D landing in the Minefield nightly builds (though pref'd off.)

Enabling:
1.Enter 'about:config'
2.Click through the warning, if necessary
3.Enter gfx.font in the 'Filter' box
4.Double-click on 'gfx.font_rendering.directwrite.enabled' to set it to true
5.Below this, right click and select New > Integer to add a pref setting
6.Enter 'mozilla.widget.render-mode' for the preference name, 6 for the value
7.Restart
(To disable, set gfx.font_rendering.directwrite.enabled to false, delete mozilla.widget.render-mode, then restart.)

Old news:
"OpenSceneGraph-2.9.6 released, introduces OpenGL ES + OpenGL 3.x support!"

In March expect:
Launch of Fermi I would want PhysX 3.0,CG 3.0,CUDA 3.0final , OGL fermi d3d11 extensions and 200 drivers with CUDA 3.1 beta and OCL new d3d extensions and 3d image writes and 3d vision windowed, youtube 3d browser support and vdpau glx interop enhancements.
Unigine heaven linux demo
iz3d 1.11 direct3d 10 drivers
Events: Cebit first week, GDC second week, 15march 3rd gpgu workshoop
Snow Leopard 10.6.3
OpenRL public release
Ubuntu 10.4 beta with fglrx 8.72 beta which seems allows 3rdmultiple vendors ATI and Nvidia working together (PCI arbitration) similar to in Windows and Mac..
Joint that with an improved switcheroo patch for kernel 2.6.34 and that could possibly bring similar to Optimus on Linux.. reportedly Apple is working on Optimus tech for MacOS in next Macbook's possibly this spring.. that implies some GT2xx chip in Apple laptops.. I would hope some laptops with amd 5xxx mobility as Fermi seems long ago (well perhaps an anuncement this summer) which would bring D3D11 features but macoSX doesn't exploit's it currently barely OGL 3.x.. I hope OGL support pickup's fast on Macos and 10.6.4/5 has OGL 3.2 with ARB extensions bring parity to d3d 10.1.That would add 5xxx drivers for hackinttoshes but that can possibly come with Dual Xeon 6 cores Westmere Mac Pro coming perhaps this month also..

Also raw info:

Try less frequently than every Draw, e.g. after every important Release(), calling ClearState(), then Flush().
I assume D3D runtime check reference counters while Flush() or Present() is invoked.
yield
I am just dipping into Direct3D and i have created a app which is able to display UYVY video using D3DXLoadSurfaceFromMemory(), StrechRect() and Present() API's.

I'd like to capture and scale screenshots directly in video memory without having to rely on the CPU too much.

I've looked at GetFrontBufferData, but this copies data into system memory.

Any suggestions on how to go about doing this?

FYI I'm brand new to the Direct3D API, so nothing is to obvious to mention.

Thanks!

I have found 3dlabs GLSL frontend compiler and other utilities is on orange book web page..

cudaGetDeviceCount doesn't create a context. The first cudaMalloc will create a context. If you want to force context creation before a cudaMalloc, use cudaFree(0).
mira dc5.0 code para mul24 i IL en gpu shader analyzer para ver que pasa tambien
mads con float, integer i double

cg 2.2 gs_simple uses glsl geo shader and works with ati! so seems this would work CG on MacOSX on Nvidia and ATI also (remember Nvidia doesn't worked on Mac)..
this seems to show example to use that GLSL geo shader support on other programs (only this works).. also this would be final cg release and now waiting cg 3.0 with domain and hull shaders and fermi instructions..

New findings and questions..

2010-03-01T20:52:00.003+01:00

Regarding DX IL:
Well I can only generate with fxc, right?.. also seems I can't feed DX IL to DX via fxc or D3DCompile or CreateComputeshader? seems no.. then what's is for excluding IHVs for doing drivers for it as base.. so no IL modification and compiling from that?ATI SKA also gets it but doesn't generate from it..
Also is DX IL spec public or anywhere avaiable?

Regarding OGL-DX interop trough OCL:
having new DX extensions for OCL Nvidia published only and AMD shipping is possible to
use for OGL-DX interop? (using createcontex with cl_context_properties having both ogl context and d3d context stuff)
It will work someday? one vendor at least? ogl extension says can be possible..
also what about wgl_dx_interop is going to be supported on Vista/7 and d3d9,10,11..
going to be introduced (at least spec txt) in fermi gl extensions this month?

Regardinng OCL binaries
Found AMD OpenCL 2.01 supports binaries (both CPU and GPU targets) getting and building from that altough AMD release notes list that as a lacking feature..
perhaps since 2.0..
target CPU binary should be cross CPU i.e. work with all CPUs (AMD,INtel) across generations.. even Atoms..
there is a flag for only SSE2 requirement obviating current sse3 it will generate only sse2 code and run even on p4?..
GPU support is good but worse than Nvidia first binary chars are CLBC (cl byte code? similar to DXBC) and has assembly device code so I use 5xxx will not work on 4xxx would be better AMD IL so would work on all GPUs supported..
well at least seems that OCL generates AMD IL v2 in my 5xxx and I don't know if this works on 4xxx..
Also seems ELF binary and also has other info than code so you can't modify code as some headers will show code size etc..
How OCL GPU binaries compare to ELF CAL binaries with Calclassemble?..
Are the formats going to be published simiar to CAL ELF binaries.. well at least they were some time ago but I don't know if they are up to date or possible now that seems device assmebly is not possible or at least not supported officialy on 5xxx..
Also remember Nvidia gets PTX so should work current OCL binaries with Fermi acording to Fermi compatiblity guide..
also straight ptx allows modificating code.. possible but spec 1.5 still not published (this month?)
Anyway I didn't mention last time but with decuda git now having most GT 200 arch instructions (SM 1.3) you teoretically could write a CUDA wrapper that intercept cubin and using decuda get PTX which you feed to CUDA stack.. don't know why Nvidia doesn't do that.. well they must have reason regarding precision,
mul24 is not native instruction,etc..

I have ported/fixed also swan to windows and added better opencl translation from cuda kernels..
Trying to get CAL++ fiexs for windows also..

Todays news:
*cebit: Geforce 480 boxes show 1.5gb ram 8pin+6pin connector..
ATI competition will be a 950mhz 5000mhz 5870 and 5970 with 4gb at 850mhz
also seems a Computex Dual Fermi possible by Asus..
*http://www.geosenseforwindows.com/ supplies a sensor driver for Windows for using location apis
gives a demo google maps enabled.. works with weather gadget..
Then I hope QT Location API in mobilty pack has win7 location api support..
*cebit: gigabyte shows laptop with docking station having nvidia gtx2xx for laptops and netbook with multitouch and tablet convertible
*Hardware accelerated graphics and text in Firefox directwrite and 2d in nightly firefox for windows 7
*glu3 soon.
Old news:
*Flash 10.3 beta 3 supports GPU decoding for fluid HD youtube on netbooks with GMA500 (720p) and Broadcom CrystalHD (1080p) with new gma500 and CrystalHD new drivers..
as it's based on DXVA seem now they have proper DXVA on drivers.. it's 1 or dxva 2? i suppose 1 as it works on XP also but can be on vista uses dxva 2.0?..
*C3DL 2.0 now WebGL and beyond
*OpenScreenGraph 1.96 supports OGL ES 1.x and 2.0 and GL 3.x and Iphone coming soon..

OCL tip:
Images on today's hardware have caches, so you get most of the benefits of local memory without the difficulty. The caches are small (~32kB L1, ~768kB L2) so you need a lot of locality to make it work.
Writing to images is very slow. Avoid it if you can.

Reading Fermi CUDA stuff!

2010-02-26T20:30:00.005+01:00

Fermi comp guide:
CUBINs are only compatible forward up to major revision so 1.x cubins only work on Tesla arch not Fermi..
Nvidia 195 drivers and up support forcing JIT compilation of kernels in PTX for execution setting CUDA_FORCE_PTX_JIT, which is a way of checking for Fermi support of CUDA programs. I.e. if executable doesn't contain PTX codes will fail..
For optimizing for Fermi arch (since CUDA 3.0) is better to add explicitly code=compute_20 to code=compute_10 so it generates better code (?) also add sm_10 sm_20 for cubins.
Seems that cubin are cached by the driver so generated once and survive reboots, crashes,etc.. (where are they stored?)..
For CUDA driver API use nvcc -ptx and load using cumoduleloaddataex..
Since CUDA 2.1 compiling with arch=sm_xx (default standard (?)) evaluates to code=sm_xx (cubin) and code=compute_xx (ptx) so PTX code is inserted..

Fermi tuning guide& Programming guide 3.0:
*New graphics interop API tex interop and DX11 supports: (pags 37 cudart 63 driver api)
*interop cudart driver api(pag 72):
->allocate mem with whatever API
->if initiated context with driverapi first CUDA runtime call doesn't create context->cublas cufft work from CUDA driver api
-don't work with emulation and cuCtx{push,pop} functions..
*Use concurrent kernels:
check cudaGetDeviceProperties() concurrentKernels
and use multiple streams..

(up to 4 ONLY and from only one context? I supposed up to 16 as is the number of SMs so seems one kernel per Graphics processing cluter and not per SM) also from the same context invalidates running multiple parallel CUDA executables for extracting more perf (so no similar to use CPU cores running multiple single threaded apps. This is a shame as hardware has fast context switching but with bad coded CUDA program in parallel only aleviates overhead in switching it but no in parallel..
I suppose it's software implementation issues and fixed in software in CUDA 3.x or if not would be fixed for Fermi 2 so we can run as kernels as SMs and any number on contexts in parallel altough possibly every SM can run only one context..
*Arithmetic Instruction perf table (pag 90):
remember tesla 8 cores per sm's and good ops execute one warp in 4 clocks so 8 inst/s per SM.
Fermi is 32 cores per SM. So 16 Sms.
Note 32 bit integer is on Fermi as good as floating point so imad=mad in perf.. must see.
*All global mem and shared mem access is done per warp not half-warp as before so check all goes well.
Shared mem is expanded to 32 banks.
Now with cache global mem coalescing seems less a requirement and also shared mem is much better as only bank conflict are "when 2 o more threads request data in different words to same bank" i.e it has multiple words broadcast, etc.. So 8-bit access,16 bit , 32(always fast) 64bit (doubles) and even 96bit (was always) sequential access is good now as 32bit and also 32bits with 8bit offset for example.. well I don't know if a 8 bit offset for 32 bit words is bad or not as that would require breaking every word in two banks and don't know if that is served jointly or not but I pressume runs without bank conflicts fast!
*Similar to DC which seems to require knowing at compile time(DispatchIndirect is only for grid size?):

and OpenCL having __attribute__((reqd_work_group_size(X, Y, Z)))

The workgroup size that must be used as the local_work_size argument to clEnqueueNDRangeKernel. This allows the compiler to optimize the generated code appropriately for this kernel.

CUDA introduces __launchbounds__ to be appended to kernels for specifying min blocks per SM desired ocuppancy max workgroup size so it can optimize register (spilling) usage..

*By default (i.e. compiling source programs without change) L1 cache size will be 16kb so shared mem would be increased 3x per SM.. The function for setting is
*We know global mem was cached by hardware cache (L2) and know that there is a L1 cache at least 16kb in size.. I presumed this was used for caching global mem but turns out that L1 caching of global mem can be disabled (compile time would be better at runtime).. So what is used L1 for? local mem for register spilling for example that can't be disabled..
*a read only place from global mem (like const variables in C++) used along all threads in kernel is cached using constant cache (doesn't require __constant address space)..
*Don't use 24 int integers are slow on CUDA check at compile time with CUDA_ARCH only device code
but guide says..
*FP ops are higher precision so results can differ from Tesla
*As Fermi supports 64bit address space if passing -m64 to nvcc compiles to 64 host code and device code which is slower than 32 bits.. So if you don't need 64 address space but compile to 64 bits host (i.e. the GPUs your program would run are less than 4gb or program needs already less than so compile) compile separately kernel code from host code..
*CUDA C++:
function overloading: f(int a) f(double a)
default parameters: f(a,b=0);
namespaces: namespace nv{ int a;} namespace ati{ int a;} nv::a=2; ati::a; using namespace nv; a=3;(nv )
operator overloading uchar4 operator+() uchar4 a,b,c; c=a+b;
implicit, explicit and specialized templates: f(x) or int x; f(x) and f(x){return(2);} f ret(3);
Fermi stuff:
classes and functors.

Seems support for virtual functions is missing yet and function pointers.. but coming..
Recursion and mem allocation inside kernel still lacking and coming much later (?)..
Remember all that supported in hardware..

Search fermi new insts in b.5, b.6, b.11 (103,104)
I don't know if b.12 is new __prof_trigger which exposes 8 counters which are incremented per warp each time and can be queried by profiler.. would be good if you can read with another inst in kernel? must think..
b.14 has launchbound(pag 112) doc.
Appendix G has the architecture feature chart (G1)
LACKING documentation from the guide:
*Launching of 3D grids! (well in 102 b.4 you find griddim is dim3 type but in pag 8 2.2 you see blocks are 1d or 2d thing and well in b.13 in 111 is said grid is dim3 but .z=1) (DC 5.0 has it OpenCL model (the API) supports that also)
*Surface functions (I hope are no left for CUDA 3.1 or later as Fermi supports it and even Tesla as is used for RWTexture in DC and image writes in OpenCL driver)
*Info that Fermi allows D2H H2D simultaneous transfers via async functions (check concurrent bancwith 1.1)

Also somethings I was unaware of:
use of __restrict__ in cuda pointers and some SLI info about cuda, SLI and D3D graphics interop..

Questions about OpenCL AMD d3d9 interop!

2010-02-25T16:47:00.000+01:00

Which is the correct way? is API stable?
See it http://forums.amd.com/devforum/messageview.cfm?catid=390&threadid=128467&enterthread=y

Hi I have coded some example trying to see how d3d interop works..
I see up to three APIs one is
KB91 - Additional Header File Required For Preview Feature: ATI Stream SDK v2.01 Support for OpenCL™ / Microsoft® DirectX® 9 & 10 Interoperability
wich shows in cl_amd.hpp clEnqueueReleaseExternalObjects similar to clEnqueueReleaseGLObjects for gl interop..
this seems to be the correct way according to KB but I can't find clEnqueueReleaseExternalObjects with clGetExtensionFunctionAddress so I also
in cl_d3d9.h I see
clEnqueueAcquireD3D9ObjectsKHR which is another way..
this is found by clGetExtensionFunctionAddress
then for buffer interop similarly we found two functions:
clCreateFromD3D9BufferKHR
and then below is
"//
// Legacy AMD CL-D3D9 interop extension
//"
with
clCreateFromD3D9Buffer
function.
With clGetExtensionFunctionAddress I found clCreateFromD3D9BufferKHR which is the correct I think.
also I'm asking for correct texture interop(yeah I know is image support required but I'm talking having correct/stable source code here not testing)
I must use this clCreateFromD3D9TextureKHR?
For testing I create context with
cl_context_properties cps[6] =
    {
        CL_CONTEXT_PLATFORM,
        (cl_context_properties)platform,
        CL_CONTEXT_D3D9_DEVICE,
        (cl_context_properties)g_pd3dDevice,
        0
    };
with a D3d device created and then use
CL_API_ENTRY cl_mem (CL_API_CALL
*myclCreateFromD3D9TextureKHR)(
    cl_context          /* context */,
    cl_mem_flags        /* flags */,
    IDirect3DTexture9 * /* texture */,
    HANDLE              /* shared_handle */,
    UINT                /* miplevel */,
    cl_int *            /* errcode_ret */);
myclCreateFromD3D9TextureKHR=(P2)clGetExtensionFunctionAddress("clCreateFromD3D9TextureKHR")
as clCreateFromD3D9TextureKHR is defined already in header.
then calling
myclCreateFromD3D9TextureKHR(context,CL_MEM_READ_WRITE,g_inputTex,g_handle,0,&status);
where g_inputTex is created and g_handle is the last parameter returned by CreateTexture. Is this correct?
In this way I get runtime error after enabling images with "set GPU_IMAGES_SUPPORT=1"
setting the Handle shader parameter to NULL
myclCreateFromD3D9TextureKHR(context,CL_MEM_READ_WRITE,g_inputTex,NULL,0,&status);
returns an CL_INVALID_D3D_OBJECT
What's the correct API? Is that stable?
Is there any sample showing interop?
is d3d11 interop coming? as nvidia has one d3d11 extension published in khronos registry?..

News 25/2!

2010-02-25T16:46:00.000+01:00

*gpu-z 0.3.9 fixes opencl ati reporting!
*The Wind Top desktop has 24inch 3D 120hz FullHD multitouch monitor! Seems the first!
Jointly with Dell u2711 you have all the things I want from monitors in just two monitors
(well dell only adds 10 bit color, 27inch and 2560x1440 res)!

The Wind Top desktop hasHD (1080p resolution) displays range from 19- to 24-inches. At the top of the list are the Wind Tops AE2420 and AE2280, 22- and 24-inch multi-touch displays respectively, equipped with processors up to Intel Core i7. The 24-inch model features a 120Hz LED display that pairs with 3D shutter glasses. (That 3D trend isn’t dying off so fast.)

Nexus and C#

Yes, you can use Parallel Nsight/Nexus to debug CUDA C kernels written in C# or other CPU languages, but Nsight doesn't directly support the C# project type yet.
So to use CUDA.NET with Nsight, you'll need to create a dummy C++ project whose 'command' in your Nexus User Properties to your C# executable.
Then do Nexus Menu Start CUDA Debugging in Visual Studio, and you should be off and running. AFAIK, you'll still need to program the actual GPU code in CUDA C.

Pages with GPU computing stuff!
see the new? http://developer.nvidia.com/object/gpucomputing.html
you have 3 guides with Fermi stuff!

In the programming guide didn't mention that GF100 is capable of simultaneous transfers of cuMemcpyDtoHAsync and cuMemcpyHtoDAsync. I've added this to my good ol' concurrent bandwidth test and will be updating that in the near future.

search concurrent bandwidth test 1.1 for Fermi!

Missing is CUDA Developer Guide for Optimus Platforms.

__global__ function parameters are passed to the device:
* via shared memory and are limited to 256 bytes on devices of compute
capability 1.x,
* via constant memory and are limited to 4 KB on devices of compute capability
2.0.

others:
http://www.directx11tutorials.com/
[JumpToDX11-11] DirectCompute
http://vsts2010.net/220
http://www.opengpu.org/bbs/archiver/

Ivan Golubev is the blog to follow for Crypto and integer ops on GPUs!
http://www.golubev.com/blog/
He says he has added bitalign AMD IL v2 for MD5 and SHA1 cracking on 5xxx GPUs has a post estimating perf of even Fermi GPUs..
search ighashgpu 0.70 it has this support test md5 and sha1 perf:
ighashgpu.exe /h:96b13dbbc9f3bc569ddad9745f64b9cdb43ea9ae /t:sha1 /c:sd /max:7
ighashgpu.exe /h:cbe1d6d5800ec1e03a5f2a64882a0d41 /t:md5 /c:sd /max:7
In post around end January you can find also SSE code used in her program..
VS CUDA:

You should be able to implement bit rotations using the bit-align instruction introduced with Direct3D 11 and supported on both Fermi and Cypress (computes ((a:b) >> c) & 0xffffffff, where a:b is the concatenation of two 32-bit operands).
This adds nothing to the "NVIDIA vs. AMD" debate, but should provide a nice further improvement compared to the previous generation.

Maybe some other tricks are possible...
For instance both G80 and Fermi support free binary negation of operands to logic instructions (allowing NOR, NAND, NXOR, ANDN...), and Fermi supports a left shift followed by an addition as a single instruction.

Edit: also, there is always the MAD24 instruction for computations such as 5*i+1 (much faster than adds).

Benchmar reveiws has NVIDIA nTeresting: 22 February 2010!

Limitations in OpenCL
1. Can i include C inline assembly code in my openCL code?
2. Does OpenCL support addtion and subtraction with carry?
in AMD also current limitations:
Lacking Pinned mem!
uses one UAV for all allocations so max 256Mbytes usage!

Nvidia has not this two limitations no through DirectCompute!
Regarding the two OCL limitations modify CAL++ author includes in TODO list and second is assembly instruction on 5xxx so when in AMD IL author can add!
Also Nvidia trough CUDA there is a ADDC enabled compiler referenced in previous posts and also
inline assembly is unofficialy supported in CUDA!
In Nvidia OCL you can modify code PTX on the fly and add addc and feed them!

How to wait for kernel finalization without CPU usage (from Golubev blog):
CUDA create context with CU_CTX_BLOCKING_SYNC
CAL Specifically there is an undocumented feature calCtxWaitForEvent
True ATI again planted a dog - GPU kernel compiled Catalyst 9.12 are 10% slower on RV8 × 0. and somewhere in the 2-3 times slower on RV7X0. It happened due to the fact that now the ATI CAL compiler aggressively unroll !absolutely everything, so that the kernel will become the size of a few hundred KB, did not interfere in the cache ... and everything is covered

OpenCL for FreeBASIC: http://shiny3d.de/libs/fbOpenCL.zip
Remember there is also for FreePascal and Delphi!

5 Questions -- Implementing a bunch of OpenCL tools

Texture sharing
I thing you must use in OpenCL d3d interop..

http://msdn.microsoft.com/en-us/library/ee418929%28VS.85%29.aspx

ID3D10Device::OpenSharedResource
To share a resource between a Direct3D 9 device and a Direct3D 10 device the texture must have been created using the pSharedHandle argument of CreateTexture. The shared Direct3D 9 handle is then passed to OpenSharedResource in the hResource argument.

The following code illustrates the method calls involved.

sharedHandle = NULL; // must be set to NULL to create, can use a valid handle here to open in D3D9
pDevice9->CreateTexture(..., pTex2D_9, &sharedHandle);
...
pDevice10->OpenSharedResource(sharedHandle, __uuidof(ID3D10Resource), (void**)(&tempResource10));
tempResource10->QueryInterface(__uuidof(ID3D10Texture2D), (void**)(&pTex2D_10));
tempResource10->Release();
// now use pTex2D_10 with pDevice10


Textures being shared from D3D9 to D3D10 have the following restrictions.

    * Textures must be 2D
    * Only 1 mip level is allowed
    * Texture must have default usage
    * Texture must be write only
    * MSAA textures are not allowed
    * Bind flags must have SHADER_RESOURCE and RENDER_TARGET set
    * Only R10G10B10A2_UNORM, R16G16B16A16_FLOAT and R8G8B8A8_UNORM formats are allowed

Interesting post: http://software.intel.com/en-us/articles/copying-accelerated-video-decode-frame-buffers/ vlc 1.1 is using that approach I think and also MPC Home cinema it seems!
vlc 1.1 is doing that!

Final round of Tesla Compute Cluster driver testing:
*CUDA H264 GPU video encoding work through MediaCoder
*vreveal works (clean video, sharpness)
issues:
stabilization: Gray uniform colors
contrast: i get pink color
*Badaboom fails with:
.GPU 0: ATI Radeon HD 5800 Series
FATAL:There is no GPU device supporting CUDA.
(Altough there supports TCC CUDA)

Currently the global memory available is the value returned by CL_DEVICE_GLOBAL_MEM_SIZE in device query. Full physical memory is expected to be available in one of the upcoming releases.
Global buffer is 128bit aligned addresses, UAV's are byte aligned and on 5XXX series of cards you can have up to 9 UAV's per kernel. Also through UAV's you can do byte addressable writes with the UAV arena and also atomic operations. None of these can be done on the global buffer path.

Global buffer is 128bit aligned addresses, UAV's are byte aligned and on 5XXX series of cards you can have up to 9 UAV's per kernel. Also through UAV's you can do byte addressable writes with the UAV arena and also atomic operations. None of these can be done on the global buffer path.
it is easier to burst using global memory as it is an implicit 128 bit write versus an implicit 32bit write on UAV.

3 new tools!

2010-02-25T16:39:00.002+01:00

3 New GPU tools!

Swan: A simple tool for porting CUDA kernels to OpenCL

What is it?
Swan is a small tool that aids the reversible conversion of existing CUDA codebases to OpenCL. It does several useful things:
Translates CUDA kernel source-code to OpenCL.
Provides a common API that abstracts both CUDA and OpenCL runtimes.
Preserves the convenience of the CUDA <<< grid, block >>> kernel launch syntax by generating C source-code for kernel entry-point functions.

Why might you want it?
Possible uses include:
Evaluating OpenCL performance of an existing CUDA code.
Maintaining a dual-target OpenCL and CUDA code.
Reducing dependence on NVCC when compiling host code.
Support multiple CUDA compute capabilities in a single binary

Limitations
It's not a drop-in replacement for nvcc. Host code needs to have all kernel invocations and CUDA API calls re-written.
Swan does not support a few things. In particular:
CUDA C++ templating in kernel code.
OpenCL Images/Samplers (analogous to Textures).
Multiple device management in a single process.
Compiling kernels for the CPU.
CUDA device-emulation mode.

Furthermore, it's a work in progress. It works for our code but no promises it will for yours

Cloo 0.6.2

A new version of Cloo is out.
It introduces a tracking mechanism for kernel arguments (sampler or memory objects) which prevents them from being claimed by the GC in case the user application doesn't refer to them in later code. This behaviour has been backported to the existing Set*Argument methods since it is safer. You can override auto-tracking using the newly added overloads.
A critical bug affecting image read operations together with some other minor glitches were fixed.
As for breaking changes rename any ComputeImage.PixelSize to ElementSize and you're good to go.

Clootils have been improved, too. Now, you can take advantage of some bells and whistles which control the program building behavior.

CAL++ v. 0.8 release
anouncement

C++ to IL generator/compiler with C++ bindings for CAL
http://sourceforge.net/projects/calpp/

The CAL++ library has been just released. Project homepage is located here http://sourceforge.net/projects/calpp/ .

The project consist of two main components. One is C++ binding for CAL ( it's really much easier to develop new CAL applications using bindings ) and second is C++ to IL generator/compiler.

The C++ generator/compiler has syntax very similar to OpenCL ( with few necessary exceptions ). Also it supports all devices which can run CAL kernels ( finally OpenCL like language for 3xxx ).

It has some advantages over OpenCL compiler. To name few

- it's much closer to CAL - it allows to write code which is almost as good ( or as good ) as handwritten IL. Look at the matrix multiplication example - it has almost the same ISA as prunedtree original code ( it differs only where I've added some changes ).

- Advantage of using C++. I really wouldn't like to use double-double ( or quad float ) technique without C++.

- Powerfull control over loop unroling and code selection ( at IL compilation time ). The C++ language acts like preprocesor.

- It has LDS support for 4xxx, doubles, etc. And if something is missing it can be added really easy.

But as always there are some pitfalls to this approach

- it isn't OpenCL . Having standard is always usefull.

- Only partial support for structs ( it can be much improved but never as good as OpenCL ).

- CAL++ is much closer to IL and some more knowledge about IL is required to achive full potential ( hmmm I think this is also the case with OpenCL ).

- optimization is only performed by CAL IL compiler ( which isn't that good ).

With the library there are some examples included. I think the fastest matrix multiplication might be a small gem here .

I hope that CAL++ will be usefull to someone .

Doesn't compile under Windows MSVC 2008!
Use 0.8a for GCC 4.4!
QA:

1. Have you tested on Windows?

No. But with the exception to C++ compiler problems it should work ( there is nothing platform specific in the code ).

2. Also have you added 24 bit integer instructions? they are useful for getting thread id fast for example..

CAL++ is converting code to IL. So 24 operations need to be available in CAL IL. And unfortunatelly it isn't the case.

I'm thinking how hard is to add also GDS?

Using anything that isn't available in IL is really hard ( or close to impossible ).

When CAL supported ISA assembler compilation ( 3xxx family ) you could generate ISA ASM. I would call it really, really hard as you need to be aware of many architecture limits ( and those informations simply aren't available ).

But for 4xxx, 5xxx family to use ISA requires to write your own driver stack ( as CAL doesn't support asm any more ) - I think it's simply impossible at the moment.

" It cannot be compiled at the time as it depends on some CAL Vector/Matrix classes which aren't available for public use." are this AMD NDA code or is your own code?

It's my own code, but it's far from being ready. For vectorquantization example is can be easily replaced by Image2D with simple functions to fill data.

are you using any magic in it? or I can code some wrappers?..

The Matrix/Vector code is using a little bit of magic . Any vector/matrix expression ( like vec_a = 3*vec_b + vec_c + log(vec_d) ) is converted to proper kernel ( trick with using templates for delayed execution ) and executed on gpu. It saves a lot of time with writing custom kernels .

From TODO:

1. Add UAVs support,logical operations and more double math functions and as_typen conversion

2. Add il_asm function ( usage example: il_asm("mov %1,%2", v1, v2); would generate "mov r1,r2" )

3. Add documentation and more examples

4. Easier to use local cal arrays, and more user friendly code for IL creation functions

Ideas for porting algos to GPU:AVX SSE and MMX ports!

2010-02-25T16:09:00.000+01:00

Hi this can be seen as crazy but some research of year 96 can be useful in thinking what Intel thought
were heavy useful algos that could offer improved perf using SSE,MMX,AVX!

For AVX there is an AVX site containing a lot of posts:
some new are from January offering general CRC perf spee using pcmuldq on Westemere!
also some AVX report numbers using Sandy Bridge silicon!
For SSE see:
http://www.datasheetarchive.com/datasheet-pdf/1070.html
especially intel reports 802-833 here you can see
"Increasing the Accuracy of the Results from the Reciprocal and Reciprocal Square Root"
Instructions using the Newton-Raphson Method..
which in fact is redeferenced in gpu gems3 nbody

MMX manuals here:
http://www.tommesani.com/IntelAppNotes.html
http://software.intel.com/en-us/articles/mmxt-technology-manuals-and-application-notes/

About ATI and Nvidia drivers (OCL included)!

2010-02-25T16:08:00.001+01:00

Hi
I have been investigating AMD and Nvidia drivers..

for 10.3 there are 3d hooks support for 120hz monitors but is d3d9 d3d10 or d3d11 enabled? or both? what's the API? and also what about only fullscreen support as 3d vision d3d support or windowed support? 3d vision windowed is soon..

also I have checked 10.3 and has bugs in gdebugger 5.5 (8.68 no perf counters found),gdi (still slow on aero w7),heaven ogl(issues half screen as early drivers but I think hotfix 9.12 worked fine..)

nvapi should be coming with 200 series with gpu usage apis?

first for AMD you can use some components without installing the complete driver for example
opengl drivers and ati cal drivers I think..
note this is no go in Nvidia where every component has to be from the same version..
well I don't remember opencl if it's equal since tesla computing driver not includes nvcompiler.dll
but I have found you can enable OpenCL with tesla computing driver just use opencl.dll from amd sdk 2.01 the problem it didn't work is because I have opencl.dll from older nvidia driver and installing tcc driver over it didn't remove it and amd installer doesn't overwrite (but it should!)..
to fix it del opencl.dll and reinstall or:
goto C:\ATI\Support\streamsdk_2-0-1_win764\Packages\Apps and in dev dir drop msi to
http://dl.dropbox.com/u/1416327/extractmsi.bat and search in temp3 opencl.dll..
Now I have learned some things: OpenCL works with tcc, is clever enoguh to disable cl_khr_sharing (or it's some weird issue that seems superintelligent?) and also that introduces d3d9 interop, reports icd extension and exntesions report some unroll extensions I didn't konw if where on 195 first drivers.. also
note d3d9 interop in tcc won't work so also if were intelligent would be disable on tcc..
finally I have found aes amd sample now works and some other demos that didn't work first mandelbrot I think.. I don't know if it's due to 2.01 source improvements or Nvidia improvements or both!.
Note OpenCL Nvidia is super good now with some weird issues (amd aes sample, volume 3d demo fast on w 7.. still to check is functions with no parameters as HelloCL and Apple FFT d(and ocena) emo working?) fixed and also to be ICD complaint (now amd+nvida works with 2.01 opencl.dll)and d3d9 interop (i have not checked if working) I think almost all of these are 196 ocl improvements.. no it's left for ocl is d3d10 and half support and for Fermi 3d_image_writes and d3d11 interop and better perf?
Also ocl d3d interop enables gl-dx interop?
also I have tested cuvid if working and I can't see OGL CUVID example working and I have selected preferCUDA (no default preferVP (video processor) as that return errors on init same as preferDXVA(but this doesn't worked with normal drivers)) with that I have to cuGLctxcreate to cuctxcreate and remove all cuGL functions.. should work.. but returns error in hadledecodepicture with context_invalid.. theortically should work if only CUDa cores are used..
What about win7 mft and CUVENC I think this should work as is cuda kernels..

overclock tcc doesn't work and clock reporting but with Nvidia CUDA you can get stream processor clocks at least.. anyway fan speed works..
Lastly seeing from icd spec I think all that remains for a non Khronos member to implement it are an ordered list of ocl functions in dispatch struct..

Also reported version in ccc of 2d, d3d and ogl drivers are this files:
2d->atikmdag.sys 2d version 8.01.01.1010
3d->atiumdag.dll atiumd64 d3d 735
ogl->atiogl.. 9606

All files found in 10.3 beta (new is for post 9.12 files i.e. not found in 9.12 due to crossfire restructuring almost for sure all inclde):
ati2edxx.dl_-> ati external device utility syswow64
ati2erec.dl_

AtiEDUGetThermalApiVersion
AtiEDUEnumApiSupportedDevices
AtiEDUEnumSupportedExternalDevices
AtiEDUGetExtDeviceInfo
AtiEDUOpenAdapterHandle
AtiEDUCloseAdapterHandle
AtiEDUInitializeThermal
AtiEDUSetThermalRemoteTemperatureOffset
AtiEDUSetThermalRemoteTemperatureHighSetPoint
AtiEDUSetThermalRemoteTemperatureLowSetPoint
AtiEDUSetThermalRemoteTemperatureCriticalSetPoint
AtiEDUGetThermalRemoteTemperatureOffset
AtiEDUGetThermalRemoteTemperatureHighSetPoint
AtiEDUGetThermalRemoteTemperatureLowSetPoint
AtiEDUGetThermalRemoteTemperatureCriticalSetPoint
AtiEDUGetThermalRemoteTemperature
AtiEDUThermalEnableInterrupt
AtiEDUThermalDisableInterrupt
AtiEDUGetAdapterTemperatureOffset
AtiEDUGetThermalRemoteTemperatureFP

ati adl sdk
atiadlxx.dl_ system32
atiadlxy.dl_ syswow64
ADL_Workstation_Stereo_Get
ADL_Workstation_Stereo_Set

ADL_Workstation_AdapterNumOfGLSyncConnectors_Get
ADL_Workstation_Caps
ADL_Workstation_DisplayGLSyncMode_Get
ADL_Workstation_DisplayGLSyncMode_Set
ADL_Workstation_DisplayGenlockCapable_Get
ADL_Workstation_GLSyncCounters_Get
ADL_Workstation_GLSyncGenlockConfiguration_Get
ADL_Workstation_GLSyncGenlockConfiguration_Set
ADL_Workstation_GLSyncModuleDetect_Get
ADL_Workstation_GLSyncModuleInfo_Get
ADL_Workstation_GLSyncPortState_Get
ADL_Workstation_GLSyncPortState_Set
ADL_Workstation_LoadBalancing_Caps
ADL_Workstation_LoadBalancing_Get
ADL_Workstation_LoadBalancing_Set
ADL_Workstation_Stereo_Get
ADL_Workstation_Stereo_Set

aplication profiles
atiapfxx.blb system32
atiapfxx.ex_ system32

no se (old)
atibtmon.ex_ system32 ati brigthnes monitor

ati cal
aticalcl.dl_
aticalcl64.dl_
aticaldd.dl_ OK
aticaldd64.dl_
aticalrt.dl_
aticalrt64.dl_

crossfire (new)
aticfx32.dl_ ati radeon d3d11 driver syswow64
aticfx64.dl_ ati radeon d3d11 driver system32

old
atidemgx.dll graphics demsystem32 (catalyst control center)
atidxx32.dl_ d3d 11 driver syswow64
atidxx64.dl_ d3d 11 driver system32
atieclxx.ex_ AMD external events client module (ccc) system32
atiedu64.dl_ ati external device utility system32 ati2edxx.dl
atiesrxx.ex_ AMD external events client module (ccc) system32

no se (new)
atig6pxx.dl_ powerxpress vista ogl (thunk) driver system32
atig6txx.dl_ powerxpress vista ogl driver syswow64
atigktxx.dl_ powerxpress vista ogl (thunk) driver syswow64
atiglpxx.dl_ powerxpress vista ogl driver system32

old
atiicdxx.da_
atikmdag.sy_
atikmpag.sy_
atimpc32.dl_ radeon pcom universal driver syswow64
atimpc64.dl_ radeon pcom universal driver sys32
atimuixx.dl_ multilanguage dppe dll
atio6axx.dl_ ati opengl driver system32
atiodcli.ex_ no se
atiode.ex_ no se

ogl driver syswow64
atiogl.xml
atioglxx.dl_

new
atipblag.dat contiains list (3DMark06*.exe 3DMark2001.exe 3DMark2001SE.exe 3DMark03.exe 3DMark05.exe ..)

atipdl64.dl_ --
atipdlxx.dl_ ati desktop cwddedi syswow64 old adl lib
atitmm64.dl_ tmm clone control module

new
atitmp64.dl_
atiu9p64.dl_ -
atiu9pag.dl_ powerxpress vista user mode driver (d3d9?) syswow64

old
atiumd64.dl_ readeon directx universl driver system32
atiumd6a.ca_ dat64
atiumd6a.dl_ video acceleratrion universal driver
atiumdag.dl_ readeon directx universl driver syswow64
atiumdva.ca_ dat32
atiumdva.dl_ video acceleratrion universal driver syswow64

new
atiuxp64.dl_ -
atiuxpag.dl_ powerxpress vista user mode driver (d3d10?) syswow64

I see amdpcom32 radeon pcom universal driver syswow64

ati2erec.dl_ atitmp64.dl_
atikmdag.sys ati radeon kernel model driver
atipmdag.sys ati radeon kernel model driver
atikmpag.sys mini port driver

Also found is some functions of GL_EXT_direct_state_access extensions are found (is this useful?)
this is a good extensions so is good to know..
seems GL_ARB_compatibilty is not found..

GL_EXT_direct_state_access: MISSING
---------------------------
glBindMultiTextureEXT: OK
glCheckNamedFramebufferStatusEXT: OK
glClientAttribDefaultEXT: OK
glCompressedMultiTexImage1DEXT: OK
glCompressedMultiTexImage2DEXT: OK
glCompressedMultiTexImage3DEXT: OK
glCompressedMultiTexSubImage1DEXT: OK
glCompressedMultiTexSubImage2DEXT: OK
glCompressedMultiTexSubImage3DEXT: OK
glCompressedTextureImage1DEXT: OK
glCompressedTextureImage2DEXT: OK
glCompressedTextureImage3DEXT: OK
glCompressedTextureSubImage1DEXT: OK
glCompressedTextureSubImage2DEXT: OK
glCompressedTextureSubImage3DEXT: OK
glCopyMultiTexImage1DEXT: OK
glCopyMultiTexImage2DEXT: OK
glCopyMultiTexSubImage1DEXT: OK
glCopyMultiTexSubImage2DEXT: OK
glCopyMultiTexSubImage3DEXT: OK
glCopyTextureImage1DEXT: OK
glCopyTextureImage2DEXT: OK
glCopyTextureSubImage1DEXT: OK
glCopyTextureSubImage2DEXT: OK
glCopyTextureSubImage3DEXT: OK
glDisableClientStateIndexedEXT: OK
glDisableClientStateiEXT: MISSING
glDisableVertexArrayAttribEXT: MISSING
glDisableVertexArrayEXT: MISSING
glEnableClientStateIndexedEXT: OK
glEnableClientStateiEXT: MISSING
glEnableVertexArrayAttribEXT: MISSING
glEnableVertexArrayEXT: MISSING
glFlushMappedNamedBufferRangeEXT: MISSING
glFramebufferDrawBufferEXT: OK
glFramebufferDrawBuffersEXT: OK
glFramebufferReadBufferEXT: OK
glGenerateMultiTexMipmapEXT: OK
glGenerateTextureMipmapEXT: OK
glGetCompressedMultiTexImageEXT: OK
glGetCompressedTextureImageEXT: OK
glGetDoubleIndexedvEXT: OK
glGetDoublei_vEXT: MISSING
glGetFloatIndexedvEXT: OK
glGetFloati_vEXT: MISSING
glGetFramebufferParameterivEXT: OK
glGetMultiTexEnvfvEXT: OK
glGetMultiTexEnvivEXT: OK
glGetMultiTexGendvEXT: OK
glGetMultiTexGenfvEXT: OK
glGetMultiTexGenivEXT: OK
glGetMultiTexImageEXT: OK
glGetMultiTexLevelParameterfvEXT: OK
glGetMultiTexLevelParameterivEXT: OK
glGetMultiTexParameterIivEXT: OK
glGetMultiTexParameterIuivEXT: OK
glGetMultiTexParameterfvEXT: OK
glGetMultiTexParameterivEXT: OK
glGetNamedBufferParameterivEXT: OK
glGetNamedBufferPointervEXT: OK
glGetNamedBufferSubDataEXT: OK
glGetNamedFramebufferAttachmentParameterivEXT: OK
glGetNamedProgramLocalParameterIivEXT: MISSING
glGetNamedProgramLocalParameterIuivEXT: MISSING
glGetNamedProgramLocalParameterdvEXT: OK
glGetNamedProgramLocalParameterfvEXT: OK
glGetNamedProgramStringEXT: OK
glGetNamedProgramivEXT: OK
glGetNamedRenderbufferParameterivEXT: OK
glGetPointerIndexedvEXT: OK
glGetPointeri_vEXT: MISSING
glGetTextureImageEXT: OK
glGetTextureLevelParameterfvEXT: OK
glGetTextureLevelParameterivEXT: OK
glGetTextureParameterIivEXT: OK
glGetTextureParameterIuivEXT: OK
glGetTextureParameterfvEXT: OK
glGetTextureParameterivEXT: OK
glGetVertexArrayIntegeri_vEXT: MISSING
glGetVertexArrayIntegervEXT: MISSING
glGetVertexArrayPointeri_vEXT: MISSING
glGetVertexArrayPointervEXT: MISSING
glMapNamedBufferEXT: OK
glMapNamedBufferRangeEXT: MISSING
glMatrixFrustumEXT: OK
glMatrixLoadIdentityEXT: OK
glMatrixLoadTransposedEXT: OK
glMatrixLoadTransposefEXT: OK
glMatrixLoaddEXT: OK
glMatrixLoadfEXT: OK
glMatrixMultTransposedEXT: OK
glMatrixMultTransposefEXT: OK
glMatrixMultdEXT: OK
glMatrixMultfEXT: OK
glMatrixOrthoEXT: OK
glMatrixPopEXT: OK
glMatrixPushEXT: OK
glMatrixRotatedEXT: OK
glMatrixRotatefEXT: OK
glMatrixScaledEXT: OK
glMatrixScalefEXT: OK
glMatrixTranslatedEXT: OK
glMatrixTranslatefEXT: OK
glMultiTexBufferEXT: OK
glMultiTexCoordPointerEXT: OK
glMultiTexEnvfEXT: OK
glMultiTexEnvfvEXT: OK
glMultiTexEnviEXT: OK
glMultiTexEnvivEXT: OK
glMultiTexGendEXT: OK
glMultiTexGendvEXT: OK
glMultiTexGenfEXT: OK
glMultiTexGenfvEXT: OK
glMultiTexGeniEXT: OK
glMultiTexGenivEXT: OK
glMultiTexImage1DEXT: OK
glMultiTexImage2DEXT: OK
glMultiTexImage3DEXT: OK
glMultiTexParameterIivEXT: OK
glMultiTexParameterIuivEXT: OK
glMultiTexParameterfEXT: OK
glMultiTexParameterfvEXT: OK
glMultiTexParameteriEXT: OK
glMultiTexParameterivEXT: OK
glMultiTexRenderbufferEXT: OK
glMultiTexSubImage1DEXT: OK
glMultiTexSubImage2DEXT: OK
glMultiTexSubImage3DEXT: OK
glNamedBufferDataEXT: OK
glNamedBufferSubDataEXT: OK
glNamedCopyBufferSubDataEXT: MISSING
glNamedFramebufferRenderbufferEXT: OK
glNamedFramebufferTexture1DEXT: OK
glNamedFramebufferTexture2DEXT: OK
glNamedFramebufferTexture3DEXT: OK
glNamedFramebufferTextureEXT: OK
glNamedFramebufferTextureFaceEXT: OK
glNamedFramebufferTextureLayerEXT: OK
glNamedProgramLocalParameter4dEXT: OK
glNamedProgramLocalParameter4dvEXT: OK
glNamedProgramLocalParameter4fEXT: OK
glNamedProgramLocalParameter4fvEXT: OK
glNamedProgramLocalParameterI4iEXT: MISSING
glNamedProgramLocalParameterI4ivEXT: MISSING
glNamedProgramLocalParameterI4uiEXT: MISSING
glNamedProgramLocalParameterI4uivEXT: MISSING
glNamedProgramLocalParameters4fvEXT: OK
glNamedProgramLocalParametersI4ivEXT: MISSING
glNamedProgramLocalParametersI4uivEXT: MISSING
glNamedProgramStringEXT: OK
glNamedRenderbufferStorageEXT: OK
glNamedRenderbufferStorageMultisampleCoverageEXT: MISSING
glNamedRenderbufferStorageMultisampleEXT: OK
glProgramUniform1fEXT: OK
glProgramUniform1fvEXT: OK
glProgramUniform1iEXT: OK
glProgramUniform1ivEXT: OK
glProgramUniform1uiEXT: OK
glProgramUniform1uivEXT: OK
glProgramUniform2fEXT: OK
glProgramUniform2fvEXT: OK
glProgramUniform2iEXT: OK
glProgramUniform2ivEXT: OK
glProgramUniform2uiEXT: OK
glProgramUniform2uivEXT: OK
glProgramUniform3fEXT: OK
glProgramUniform3fvEXT: OK
glProgramUniform3iEXT: OK
glProgramUniform3ivEXT: OK
glProgramUniform3uiEXT: OK
glProgramUniform3uivEXT: OK
glProgramUniform4fEXT: OK
glProgramUniform4fvEXT: OK
glProgramUniform4iEXT: OK
glProgramUniform4ivEXT: OK
glProgramUniform4uiEXT: OK
glProgramUniform4uivEXT: OK
glProgramUniformMatrix2fvEXT: OK
glProgramUniformMatrix2x3fvEXT: OK
glProgramUniformMatrix2x4fvEXT: OK
glProgramUniformMatrix3fvEXT: OK
glProgramUniformMatrix3x2fvEXT: OK
glProgramUniformMatrix3x4fvEXT: OK
glProgramUniformMatrix4fvEXT: OK
glProgramUniformMatrix4x2fvEXT: OK
glProgramUniformMatrix4x3fvEXT: OK
glPushClientAttribDefaultEXT: OK
glTextureBufferEXT: OK
glTextureImage1DEXT: OK
glTextureImage2DEXT: OK
glTextureImage3DEXT: OK
glTextureParameterIivEXT: OK
glTextureParameterIuivEXT: OK
glTextureParameterfEXT: OK
glTextureParameterfvEXT: OK
glTextureParameteriEXT: OK
glTextureParameterivEXT: OK
glTextureRenderbufferEXT: OK
glTextureSubImage1DEXT: OK
glTextureSubImage2DEXT: OK
glTextureSubImage3DEXT: OK
glUnmapNamedBufferEXT: OK
glVertexArrayColorOffsetEXT: MISSING
glVertexArrayEdgeFlagOffsetEXT: MISSING
glVertexArrayFogCoordOffsetEXT: MISSING
glVertexArrayIndexOffsetEXT: MISSING
glVertexArrayMultiTexCoordOffsetEXT: MISSING
glVertexArrayNormalOffsetEXT: MISSING
glVertexArraySecondaryColorOffsetEXT: MISSING
glVertexArrayTexCoordOffsetEXT: MISSING
glVertexArrayVertexAttribIOffsetEXT: MISSING
glVertexArrayVertexAttribOffsetEXT: MISSING
glVertexArrayVertexOffsetEXT: MISSING
I don't know if posted but error printing for ocl errors is in oclutils.cpp nvida sdk
// Helper function to get error string
// *********************************************************************
const char* oclErrorString(cl_int error)
{
   static char errorString[][64] = {
   "CL_SUCCESS",
   "CL_DEVICE_NOT_FOUND",
   "CL_DEVICE_NOT_AVAILABLE",
   "CL_COMPILER_NOT_AVAILABLE",
   "CL_MEM_OBJECT_ALLOCATION_FAILURE",
   "CL_OUT_OF_RESOURCES",
   "CL_OUT_OF_HOST_MEMORY",
   "CL_PROFILING_INFO_NOT_AVAILABLE",
   "CL_MEM_COPY_OVERLAP",
   "CL_IMAGE_FORMAT_MISMATCH",
   "CL_IMAGE_FORMAT_NOT_SUPPORTED",
   "CL_BUILD_PROGRAM_FAILURE",
   "CL_MAP_FAILURE",
   "",
   "",
   "",
   "",
   "",
   "",
   "",
   "",
   "",
   "",
   "",
   "",
   "",
   "",
   "",
   "",
   "",
   "CL_INVALID_VALUE",
   "CL_INVALID_DEVICE_TYPE",
   "CL_INVALID_PLATFORM",
   "CL_INVALID_DEVICE",
   "CL_INVALID_CONTEXT",
   "CL_INVALID_QUEUE_PROPERTIES",
   "CL_INVALID_COMMAND_QUEUE",
   "CL_INVALID_HOST_PTR",
   "CL_INVALID_MEM_OBJECT",
   "CL_INVALID_IMAGE_FORMAT_DESCRIPTOR",
   "CL_INVALID_IMAGE_SIZE",
   "CL_INVALID_SAMPLER",
   "CL_INVALID_BINARY",
   "CL_INVALID_BUILD_OPTIONS",
   "CL_INVALID_PROGRAM",
   "CL_INVALID_PROGRAM_EXECUTABLE",
   "CL_INVALID_KERNEL_NAME",
   "CL_INVALID_KERNEL_DEFINITION",
   "CL_INVALID_KERNEL",
   "CL_INVALID_ARG_INDEX",
   "CL_INVALID_ARG_VALUE",
   "CL_INVALID_ARG_SIZE",
   "CL_INVALID_KERNEL_ARGS",
   "CL_INVALID_WORK_DIMENSION",
   "CL_INVALID_WORK_GROUP_SIZE",
   "CL_INVALID_WORK_ITEM_SIZE",
   "CL_INVALID_GLOBAL_OFFSET",
   "CL_INVALID_EVENT_WAIT_LIST",
   "CL_INVALID_EVENT",
   "CL_INVALID_OPERATION",
   "CL_INVALID_GL_OBJECT",
   "CL_INVALID_BUFFER_SIZE",
   "CL_INVALID_MIP_LEVEL",
   "CL_INVALID_GLOBAL_WORK_SIZE",
   };
   return errorString[-error];
}
also ogl qbf stereo is not enabled seeing glwinfo
seeing ogl driver depends on adl but not workstatinon_setstereo get stereo o caps functions used..
also

set OGL_FORCE_ASIC_ID=37956
set OGL_FORCE_ASIC_ID=68BE
set OGL_FORCE_ASIC_ID=0x68BE
set OGL_FORCE_ASIC_ID=26814

not seen
set OGL_FORCE_ASIC_ID=9444
"ATI FirePro V8750 (FireGL)" = ati2mtag_R7XGL, PCI\VEN_1002&DEV_9444
mine is
"ATI Radeon HD 5800 Series " = ati2mtag_Evergreen, PCI\VEN_1002&DEV_6899
set OGL_ENABLE_FORCE_ASIC_ID=1

tested
glewinfo for gpu name and
visualinfo.exe

installing 8.68.3 firepro driver has 30bit support and stereo:
seen

HKR,, DisableOGL10BitPixelFormats, %REG_DWORD%, 0
HKR,, Gxo30BppPanels, %REG_BINARY%, 15,C3,76,17,15,C3,78,17

in installation diff versus cataluyst

HKR,, DALNonStandardModesBCD1, %REG_BINARY%,12,80,07,68,00,00,00,00,12,80,09,60,00,00,00,00,16,00,12,00,00,00,00,70,17,92,13,44,00,00,00,00,18,00,14,40,00,00,00,00,18,56,13,92,00,00,00,00

HKR,, DALRULE_AllowNativeModeAsDefaultModes, %REG_DWORD%, 1

GCORULE_ExtTMDSReduceBlankTiming, %REG_DWORD%, 1

"ATI FireGL V3600" = ati2mtag_RV630GL, PCI\VEN_1002&DEV_958D

"ATI FireGL V5600" = ati2mtag_RV630GL, PCI\VEN_1002&DEV_958C

"ATI FireGL V7600" = ati2mtag_R600GL, PCI\VEN_1002&DEV_940F

"ATI FireGL V7700" = ati2mtag_RV630GL, PCI\VEN_1002&DEV_9511

"ATI FireGL V8600" = ati2mtag_R600GL, PCI\VEN_1002&DEV_940B

"ATI FireGL V8650" = ati2mtag_R600GL, PCI\VEN_1002&DEV_940A

"ATI FirePro 2260" = ati2mtag_RV610, PCI\VEN_1002&DEV_95CF

"ATI FirePro 2260 " = ati2mtag_RV610, PCI\VEN_1002&DEV_95CE

"ATI FirePro 2450" = ati2mtag_RV610, PCI\VEN_1002&DEV_95CD

"ATI FirePro V3700 (FireGL)" = ati2mtag_RV620GL, PCI\VEN_1002&DEV_95CC

"ATI FirePro V3750 (FireGL)" = ati2mtag_R7XGL, PCI\VEN_1002&DEV_949F

"ATI FirePro V5700 (FireGL)" = ati2mtag_R7XGL, PCI\VEN_1002&DEV_949E

"ATI FirePro V7750 (FireGL)" = ati2mtag_R7XGL, PCI\VEN_1002&DEV_949C

"ATI FirePro V8700 (FireGL)" = ati2mtag_R7XGL, PCI\VEN_1002&DEV_9456

"ATI FirePro V8750 (FireGL)" = ati2mtag_R7XGL, PCI\VEN_1002&DEV_9444

[ati2mtag_R6xGL_SoftwareDeviceSettings]

HKR,, OGL_Specific_NA, %REG_SZ%, 1

HKR,, CatalystAI_NA, %REG_SZ%, 1

HKR,, APISpecific_NA, %REG_SZ%, 1

HKR,, TemporalAAMultiplier_NA, %REG_SZ%, 1

HKR,, Main3D_NA, %REG_SZ%, 1

HKR,, VPURecover_NA, %REG_SZ%, 1

HKR,, SmartGart_NA, %REG_SZ%, 1

Shaders: measuring perf, source translation and parsing different languages!

2010-02-25T16:07:00.005+01:00

Hi,
I hope to be pretty exhaustive of options for parsing and translating between graphics and compute shaders ( some open source)

For DX shaders:
*GPU Shader analyzer (AMD ONLY)(get DX IL and get AMD IL and GPU assembly from DirectCompute shaders or graphic shaders): now is dx11 compatible and 5xxx series and has compute domain and hull shaders..
*fxc and D3DCompile API and lib: get DX IL and bytecode from DX shaders (multivendor)
note I don't know how to build from DX IL as D3DCompile doesn't accept and also not CreateComputeShader.. finally GPU SA doesn't want to eat too..
it's a fxc option or API for going from DX IL to DX BC so I can optimize DX IL and then compile to DX ByteCode and feed that to a compute shader?
at least if you have source you can see AMD IL, DX IL and R800 assembly and

teoretically you can get AMD IL from Compute shaders using GPU SA and feed into OpenCL when it support getting binaries and building from it (or intercepting now llc or something like that).. so in OpenCL you can modify generated assembly soon (on Nvidia now..)
so you can compare quailty of generated code..
Also now you can feed DX shaders to GPU SA and a equivalent OCL shader through SKA and compare AMD IL, assembly and even all the info of ALU/tex kernels/s etc..
I have tested simple vectoradd and quality is the same (kernels/s) altough AMD IL from OCL seems much longer..
Parsing HLSL: you have Nvidia CG compiler source so CG=HLSL in 99% so you have parser and front end compiler code.. (I think it has some flex bison things)

There was a AMD HLSL which was extension to HLSL having scatter doubles etc..
http://coachk.cs.ucf.edu/courses/CDA6938/s08/UCF-2008-02-01b.pdf
http://coachk.cs.ucf.edu/courses/CDA6938/s08/AMD_IL.pdf

Now included in compute shader 5.0 and pixel shader 5.0 in DX11 all functionality also included in upcoming GLSL ext_gpu_shader5 I presume (I can't find AMD HLSL compiler anywhere so I think efforts migrated to Brook+ efforts and AMD IL):

Note for Nvidia there is a tool similar to GPU SA but I think it's payed (ShaderPerf, perfkit can't I think..)
Have to see if Nexus will have PTX code from shaders or anything like that..
Also DX11 support is missing naturally in all tools (Perfkit, shaderperf,etc..)

HLSL<->GLSL source to source translation:
hlsl2glsl-v0.9 (OpenGL ES also source code)
babelshader

Only pixel vertex shaders..

GLSL
GPU Shader analyzer (get AMD IL from DirectCompute shaders or graphic shaders)
you can use HLSL->GLSL and using GPU SKA compare quality of generated GLSL vs HLSL AMD IL or assembly code..

Parsing GLSL:
you have flex and bison almost from spec (tokens and grammar)..

3d labs glsl validate and front end compiler open source..
(i can't find)

hlsl2glsl-v0.9..

Brook+ is open source
has AMD IL code gen source and brook parsing (ctool based)

Another thing is measuring perf of not shaders of whole thing with OpenGL gdebugger, mac OpenGL perf libs, AMD GPU 2.1, Nvidia Nexus and GL perf API and libs(perfkit sdk),
For having similar to OpenCL CAL lib see CAL++.
For porting CUDA to OpenCL there are to guides from Nvidia and AMD and:
Experiences porting from CUDA to OpenCL
Presentation at the Daresbury Machine Evaluation Workshop, 2009
also a tool:
Swan: A simple tool for porting CUDA kernels to OpenCL
A good OpenCL to DirectCompute driver wuold be good!

Enabling OpenCL Image support on AMD GPUs!

2010-02-19T21:20:00.003+01:00

Well I have been holding this trick on my head for over a month now..
More info on my blog coming soon: oscarbg.blogspot.com
Really you can enable image support set:
set GPU_IMAGES_SUPPORT=1
or export GPU_IMAGES_SUPPORT=1 in linux
tested on 5870 and amd stream 2.0 and hotfix 9.12 only works for
2d images..
Similarly you can enable byte_addresable_support (but seems is not using RAW UAVs) and some Nvidia samples work (histogram64) with GPU_BYTE_ADDRESSABLE_STORE
Also doules extension reporting
GPU_DOUBLE_PRECISION
and gl
CL_KHR_GL_SHARING

Running QT everywhere!

2010-02-19T21:13:00.000+01:00

TODO: post links for every thing..
I have just found a lot of platform running QT!
Last QT 4.6.2 ships with win32 bin, mac(32,64) and linux(32,64)!
You can build for win64 but is long and qt 4.6 win64 binaries ara avaiable on google code since today!
If you use VS install latest qt vs ide 1.1.4
you can build also with qt creator 1.3.1
For mobiles:
you have symbian and maemo and now meebo (moblin+maemo)
Also I have found tegra2 board working in qt blog post! (android? windows ce? linux?)
Also you have a google nacl port (for Chrome browser or IE via frame) in qt labs blog!
A port to kindle amazon is online also!
And in MWC has been shown working with remaining mobile GPUs:
*omap4(sgx 540)
*st u8500 (mali gpu)
which jointly with
tegra2 (nvidia gpu)
show is everywhere..
Well for Android you have a QT port also:
you need custom NDK with STL port included if you want..

Why I'll choose Qt GUI and not Android one?
1. The speed, Qt is more powerful and it's much more faster.
2. The features, just look at http://doc.trolltech.com/4.6/qtgui.html.
3. Declarative UI.
4. The API is very robust and stable.
5. IMHO Qt is written in a superior language. I don't like java :P. I
think if you'll ask java about me it will give you the same answer :P.
(Ok here I'm jocking).
6. etc.

only left is ipod ipad but in progress:
http://www.qt-iphone.com/Roadmap.html
currently QtCore mostly done, QtGUI hard as cocoa touch!=coca
also would be good to have all QT multitouch support and Mobility APIs just anounced as Location +Sensors+Camera API..
but this is easier said than done
then I can programm for QT for everything..

Parallel algorithms avaiable on CUDA,OCL,DC,CAL: status update

2010-02-19T17:37:00.002+01:00

lin alg status update:
Matmul:
CUDA: CUBLAS (no code) Volkov (code) and yesterday post (assembly fastest to date 480 gflops)
CAL: beyond3d cal 1tflop matmul post
OCL: hazeman post above uses port of cal code to propietary but similar to CL code..
DC: bernaclejunior testing with doubles doesn't worked (XNA forums)
Matvec:
CUDA: CUBLAS (closed) and some papers use custom code (magma, paper mid 2008) (as 20-50% faster)
OCL: Bealto post above (high efficient on AMD and ATI) should be easy to port DC
Sparse matvec:
CUDA: CNC,CUSP,etc..
OCL,DC: BernacleJunior post on AMD and XNA forums (working on it)..

FFT:
CUDA: CUFFT 2 papers at SC08 having higher perf 3d ftts and 2d paper ->d3dCx
DC: has lib
OCL:
Apple code is 2x-3x slower than CUFFT seems (on Nvidia Linux )(also 10.6.2 is slow go see 10.6.3..)
on AMD doesn't work for size >512^2 in 2.0 or 2.01 fixed internally seems..
AMD 2.01 sample is hard coded 1024 perf?

Sort:
CUDA: CUDPP, CUDA sample (code)
OCL,DC: BernacleJunior post on AMD and XNA forums.He claims near 400Mkeys/s on vs state of the art Nvidia sorting less 200mkeys on GTX285.
Also reportedly Lee Hows has fast code working!

also CUDPP has triangular solvers and soon graph algos and hashes..

More news!

2010-02-19T17:08:00.015+01:00

I have left some news and some news:
UPDATE:

1. AMD SKA allows getting AMD IL without having AMD GPU and also see tex:alu ratio, and other info for all AMD GPUs at the same time
2. AMD SDK ships utils source so now Nvidia and AMD OCL SDKs can be compiled in VS2010!
3. gdebugger 5.5 doesn't detect amd perf counters with 10.2 I think with 9.12 hotfix worked
not working with 10.3 beta

*See next post http://oscarbg.blogspot.com/2010/02/parallel-algorithms-avaiable-on.html
*Fermi X2 on track, possible launch date is May!
*In 2-3 weeks we have 5830 (high perf low budget card) and 2GB 6 miniDP 5870 card on 11 march!
*catalyst 10.3 beta leak avaiable go search for it! (8.71.3 CAL 556)
*gpu computing gems call
*matmul by hazeman:
it's a assembly->c port similar to 1tflops mamtmul cal example it's bad it uses her own C->IR compiler but easy port OCL? and what about perf?
*bernaclejunior is doing good job regarding sort and sparse matvec on OCL,DC..
He claims near 400Mkeys/s on vs state of the art Nvidia sorting less 200mkeys on GTX285.
Also reportedly Lee Hows has fast code working!
Some intermediate code posted on XNA and AMD forums but still not the best..
*Matvec mul high perf OCL code from Bealto (AMD and Nvidia tested).
*I have tested cubin optimized matmul code and I get 480gflop/s not bad from 380gflop/s
and also I have seen tesla computing driver no supports overclokcing in evga precision..
also gpu-z and evga not read core speed and mem speed and also not gpu usage and mem info anyway
temperature and fan speed is ok..
It's very long so (tested on vc2010rc1)
change in autoprofile:

profile_sgemm_square("../method1/decuda_ldsb32_cudasm.cubin", "method1_variant_sgemmNN", &method1_DrvWrapper, cat(OUTPUT_DIR,"method1/variant_threads320.txt") );
profile_general_sgemm_square("../method6/decuda_ldsb32_cudasm.cubin", "method6_variant_sgemmNN", &method6_DrvWrapper, cat(OUTPUT_DIR,"method6/variant_threads320.txt") );
profile_general_sgemm_square("../method7/decuda_ldsb32_cudasm.cubin", "method7_variant_sgemmNN",
&method7_DrvWrapper, cat(OUTPUT_DIR,"method7/variant_threads320.txt") );
profile_sgemm_square("../method8/decuda_ldsb32_cudasm.cubin", "method8_variant_sgemmNN",
&method8_DrvWrapper, cat(OUTPUT_DIR,"method8/variant_threads256.txt") );

variants are the fastest and 1 is the best (480gflops/s). also set:

for( n1 = 32 ; n1 <= 4096 ; n1+=96) for( n1 = 5 ; n1 <= 4096 ; n1++)

result is 100x test speed in

->profile_general_sgemm_square(profile_general_sgemm_suqare.cpp,profile_sgemm_suqare.cpp)
->profile_CUBLAS_overN

* I also have tested voxel sparse demo and fixed for tcc but building 1gb samples crashes on
ball example no mem with x32 release exe but x64 crashes anyway have to fix..
Found also sibenik and Fairy scenes but I don't know how to build sibenik-d example displacament
mapped using bump map texs(?)
I have to test it..

Antialiasing in Deferred shading GL code

new GL multivendor SM5.0 info found in 10.3:
*GL_EXT_tessellation_shader

gl_TessCoord gl_TessLevelOuter gl_TessLevelInner

*GL_EXT_shader_subroutine
*GL_EXT_gpu_shader5

memoryBarrier bitCount findLSB findMSB bitfieldReverse bitfieldInsert bitfieldExtract floatBitsToInt floatBitsToUint intBitsToFloat uintBitsToFloat

*GL_EXT_gpu_shader_fp64
new in 10.3:
*GL_EXT_shader_atomic_counters
GL_MAX_ATOMIC_COUNTERS_EXT

glResetAtomicCounter
check fail: index must be a constant in atomic counter functions
gl_MaxAtomicCountersEXT
atomicCounterIncrementEXT atomicCounterDecrementEXT atomicCounterEXT
imageAtomicAdd imageAtomicSub imageAtomicMin imageAtomicMax imageAtomicIncWrap imageAtomicDecWrap imageAtomicAnd imageAtomicOr imageAtomicXor imageAtomicExchange imageAtomicCompSwap

GL_EXT_texture_compression_bptc (replaces amd extensions)
GL_AMD_conservative_depth

OpenCL:
1.pinned mem enabled on nvidia via:

I use

host_mem = clCreateBuffer(context,
CL_MEM_ALLOC_HOST_PTR | CL_MEM_READ_WRITE,
size,NULL,&ocl_err);
*ptr = (void*)clEnqueueMapBuffer(cmd_queue,host_mem,
CL_TRUE,CL_MAP_READ|CL_MAP_WRITE,
0,size,0,NULL,&evt,&ocl_err);

to create page locked memory using the NVIDIA driver, where it works fine. However on my AMD card this makes no difference to malloced memory.

AMD guys confirm still not working.
2. cvs with icd code is for Khronos members spec updated
3. new cl headers at khronos list funcaddress used by ICD.. has cl_ext.h and cl_gl_ext.h
http://www.khronos.org/registry/cl/ headers
4. OpenCL 2.01 on Ubuntu 9.10:

You indeed have to boot with "nopat" or use Catalyst 10.2, when it becomes available. CAL version >= 1.4.553 to get this working without "nopat" option.

XvBA and other linux video decoding updates:
*vaapi guy working on Crystal HD support?
For Crystal HD demos:
Crystal HD SDK from GIT, as of 2010/02/15.

*at least supported in basic samples
*xvba now working for vlc 1.1git and gnash via updates for xvba-vaapi and gnash

Status of Xbva:
Works with MPLAYER (ass subtitles included),VLC and GNASH!
issues:
1.First broken decode in 5xxx..
2.Deinterlacing is broken in XvBA. It's the second most critical bug that has to be fixed by the
end of April.
(Only bob deinterlacing at this time. More elaborated deinterlacers are not, and won't be, exposed to the public builds of xvba-video.)

Changelog:

Version 0.6.5 - 08.Feb.2010
* Add brightness/contrast/hue/saturation display attributes
* Fix vaPutSurface() window resize. e.g. when switching to full-screen mode
* Allow vaPutSurface() to render to multiple drawables from a single surface

Notes:
- My ProcAmp adjustments are probably not fully correct. e.g. hue doesn't preserve luminance yet. Besides, this uses an extra FBO.
- The last change workarounds a bug in the driver and now makes it possible to use VA-API acceleration with Gnash with the the AGG renderer. However, this exhausts another performance problem (flickering in windowed mode) of the driver. You can workaround that with XVBA_VIDEO_PUTSURFACE_FAST set to "yes" or "1". The semantics are not fully equivalent and can cause problems, hence it's disabled by default though it's designed to work with Gnash and MPlayer.
There is already native VA-API support for G45. At this time, it only does MPEG-2 VLD, i.e. full video decode. Intel is working on H.264 support and this should be available by Q2. I don't think there is any H.264 video decoding at Gallium3D level yet, so VDPAU / VA-API support would be useless at this time.

Version 0.6.6 - 11.Feb.2010
* Fix XvBA objects destruction for fglrx >= 8.70.3
* Fix vaPutImage() to a surface used for decoding
* Fix vaGetImage()/vaPutSurface() with surface dimensions not a multiple of 16
* Fix rendering of VA subpictures that were previously deassociated

The third change is actually two different workarounds for a single and major flaw in XvBA. I have not fully regression tested but this looks OK for MPlayer, Gnash and VLC. This should fix Kano problems.
The fourth change is a fix for MPlayer/VA-API with ASS support, and that I will probably upload tomorrow. I have to check against the latest Intel drivers first. NVIDIA is already fine.

With this mplayer-vaapi snapshot and xvba 0.6.6 ASS works!

Version 0.6.7 - 18.Feb.2010
* Use fail-safe values for H.264 videos encoded over HP@L4.1
* Fix hue rotation to preserve luminance
* Fix internal contrast range to [ 0.0f .. 10.0f ]
* Fix rendering of multiple subpictures per surface
* Fix vaCopySurfaceGLX() for surfaces with dimensions not a multiple of 16

- The first change ensures that we don't crash or do weird things if we throw unsupported H.264 contents to the decoder. Wel, it
tries to get things on a safer side, without really fixing it.

- The ProcAmp changes are probably still not correct but this looks better for contrast and hue rotation.

- The fourth change fixes rendering of multiple subpictures per surface. In particular, you can now have OSD + EOSD + ProcAmp
adjustment bars (3 subpictures) in MPlayer without crashing the application.

- The last change is a workaround for a serious XvBA flaw, now implemented in vaCopySurfaceGLX(). e.g. for mplayer -vo vaapi:gl -va
vaapi. As a side effect, this would also workaround another limitation in the future iteration (0.6.8) whereby only GL_BGRA textures
are supported at this time.

Mplayer vaapi

Version 2010.02.12
* Fix YV12 rendering for SW codecs
* Add EOSD support (ASS subtitles)
* Add compatibility with original VA-API 0.29
* Add support for -geometry +xxx+yyy (Adam Strzelecki)

For EOSD & AMD, you need xvba-video >= 0.6.6.

Learned from voxel rendering demo code: CUDA 3.0 how to change cache size (for Fermi) function found!

2010-02-18T21:21:00.001+01:00

its in voxel code:
\efficient-sparse-voxel-octrees\src\framework\base\dllimport.inl
cuFuncSetCacheConfig
cuFuncSetCacheConfig, (CUfunction hfunc, CUfunc_cache config), (hfunc, config))
also other functions i didn't know in:
cuGraphicsSubResourceGetMappedArray
cuGetExportTable

Also they don't use GLEW and initialize..
other tricks:
CPU trick:

// Force the main thread to run on a single core.
SetThreadAffinityMask(GetCurrentThread(), 1);

GPU trick:

flags |= CU_CTX_SCHED_SPIN; // use sync() if you want to yield
#if (CUDA_VERSION >= 2030)
flags |= CU_CTX_LMEM_RESIZE_TO_MAX; // reduce launch overhead with large localmem
#endif
what about CU_CTX_LMEM_RESIZE_TO_MAX?

Also Voxel raycasting demo has good code supports Stereo OpenGL rendering and GUI controls!! for Quadros!
and good code multisampling..

also you can see functions added since 2.1:

#if (CUDA_VERSION >= 2020)
FW_DLL_IMPORT_RETV( CUresult, CUDAAPI, cuDriverGetVersion, (int *driverVersion), (driverVersion))
FW_DLL_IMPORT_RETV( CUresult, CUDAAPI, cuMemHostAlloc, (void **pp, size_t bytesize, unsigned int Flags), (pp, bytesize, Flags))
FW_DLL_IMPORT_RETV( CUresult, CUDAAPI, cuMemHostGetDevicePointer, (CUdeviceptr *pdptr, void *p, unsigned int Flags), (pdptr, p, Flags))
FW_DLL_IMPORT_RETV( CUresult, CUDAAPI, cuFuncGetAttribute, (int *pi, CUfunction_attribute attrib, CUfunction hfunc), (pi, attrib, hfunc))
FW_DLL_IMPORT_RETV( CUresult, CUDAAPI, cuTexRefSetAddress2D, (CUtexref hTexRef, const CUDA_ARRAY_DESCRIPTOR *desc, CUdeviceptr dptr, unsigned int Pitch), (hTexRef, desc, dptr, Pitch))
FW_DLL_IMPORT_RETV( CUresult, CUDAAPI, cuWGLGetDevice, (CUdevice *pDevice, HGPUNV hGpu), (pDevice, hGpu))
#endif

#if (CUDA_VERSION >= 2030)
FW_DLL_IMPORT_RETV( CUresult, CUDAAPI, cuMemHostGetFlags, (unsigned int *pFlags, void *p), (pFlags, p))
FW_DLL_IMPORT_RETV( CUresult, CUDAAPI, cuGLSetBufferObjectMapFlags, (GLuint buffer, unsigned int Flags), (buffer, Flags))
FW_DLL_IMPORT_RETV( CUresult, CUDAAPI, cuGLMapBufferObjectAsync, (CUdeviceptr *dptr, unsigned int *size, GLuint buffer, CUstream hStream), (dptr, size, buffer, hStream))
FW_DLL_IMPORT_RETV( CUresult, CUDAAPI, cuGLUnmapBufferObjectAsync, (GLuint buffer, CUstream hStream), (buffer, hStream))
#endif

#if (CUDA_VERSION >= 3000)
FW_DLL_IMPORT_RETV( CUresult, CUDAAPI, cuMemcpyDtoDAsync, (CUdeviceptr dstDevice, CUdeviceptr srcDevice, unsigned int ByteCount, CUstream hStream), (dstDevice, srcDevice, ByteCount, hStream))
FW_DLL_IMPORT_RETV( CUresult, CUDAAPI, cuFuncSetCacheConfig, (CUfunction hfunc, CUfunc_cache config), (hfunc, config))
FW_DLL_IMPORT_RETV( CUresult, CUDAAPI, cuGraphicsUnregisterResource, (CUgraphicsResource resource), (resource))
FW_DLL_IMPORT_RETV( CUresult, CUDAAPI, cuGraphicsSubResourceGetMappedArray, (CUarray *pArray, CUgraphicsResource resource, unsigned int arrayIndex, unsigned int mipLevel), (pArray, resource, arrayIndex, mipLevel))
FW_DLL_IMPORT_RETV( CUresult, CUDAAPI, cuGraphicsResourceGetMappedPointer, (CUdeviceptr *pDevPtr, unsigned int *pSize, CUgraphicsResource resource), (pDevPtr, pSize, resource))
FW_DLL_IMPORT_RETV( CUresult, CUDAAPI, cuGraphicsResourceSetMapFlags, (CUgraphicsResource resource, unsigned int flags), (resource, flags))
FW_DLL_IMPORT_RETV( CUresult, CUDAAPI, cuGraphicsMapResources, (CUgraphicsResource *resources, CUstream hStream), (resources, hStream))
FW_DLL_IMPORT_RETV( CUresult, CUDAAPI, cuGraphicsUnmapResources, (unsigned int count, CUgraphicsResource *resources, CUstream hStream), (count, resources, hStream))
FW_DLL_IMPORT_RETV( CUresult, CUDAAPI, cuGetExportTable, (const void **ppExportTable, const CUuuid *pExportTableId), (ppExportTable, pExportTableId))
FW_DLL_IMPORT_RETV( CUresult, CUDAAPI, cuGraphicsGLRegisterBuffer, (CUgraphicsResource *pCudaResource, GLuint buffer, unsigned int Flags), (pCudaResource, buffer, Flags))
FW_DLL_IMPORT_RETV( CUresult, CUDAAPI, cuGraphicsGLRegisterImage, (CUgraphicsResource *pCudaResource, GLuint image, GLenum target, unsigned int Flags), (pCudaResource, image, target, Flags))
#endif

currently fails with CUDA Compute Cluster driver:
in CudaModule::staticInit(void)
change that:
checkError("cuGLCtxCreate", cuGLCtxCreate(&s_context, flags, s_device));
by
if(tcc)
{
checkError("cuCtxCreate", cuCtxCreate(&s_context, flags, s_device));
//res = cuGLInit();
}
else
checkError("cuGLCtxCreate", cuGLCtxCreate(&s_context, flags, s_device));
cuglinit perhaps needed but depecrated anyway
changed in cuInit(0); or after cuctxcreate?

also if tcc was more smart would work and fallback
to host interop as CUDA already does so I think directly
all CUDA GL functions return error in tcc..

anyway thanks good code change:
Buffer::Hint_CudaGLin CudaRenderer::CudaRenderer(void) to Buffer::Hint_None
so
: m_frameBuffer (NULL, 0, Buffer::Hint_None),//Buffer::Hint_CudaGL),

A month of news!

2010-02-18T19:49:00.003+01:00

So here it goes all random news I consider interesting in this past month:
* AMD CAL libs coming to MAC? In PGI 10.2 pgaccelinfo includes -ati -amd to report ati accelerators info.. This is in Mac release too.. and says libamdcalcl.dylib not found.. so seems
is not working?
This will close the hole of having standard OpenCL in 3 OSes and also CUDA and CAL on three Oses also..
Remember related news is PGI interested in using Noveau stack as base of enabling GPU computing stack thorugh it for OpenSolaris and FreeBSD after Nvidia spoke about Solaris and demonstrated(?) in GTC08.. but now is dead..
For Windows pgaccelinfo working copy aticalcl.dll to libamdcalcl.dll in dir and also calrt and it works.. So seems really PGI has AMD CAL for MAC.. as has linked the dylib no?
I hope they don't spend too energies working on it since OpenCL is better target for PGI accelerator model..
perhaps is good mail streamdeveloper amd dot com asking it..
* Par4All allows autopar for CUDA, etc..
* After AMD assembly matmul kernel achieved 1tflop on 48xx hardware and 58xxx should be 2tflops
now we have assembly optimized matmul for Nividia having 10-20% better perf.
search "Hand-Tuned SGEMM on GT200 GPU, 10% ~ 20% improvement of SGEMM"
allows 512 gflops gtx 285-> 1tflop matmul for Fermi? (like larrabe? mira acm video)
it has code and report..
Also has trick that using asm("") in cu kernel including PTX works via nvcc due to Open64 features..
*Nvidia has released updated videos on Youtube demos of "fluid demo" for fermi launch, Parallel Nsight (nexus) and one about Sled demo talking a enginner about it..
* I still don't know if some opencl.dll from Khronos works for Nvidia and AMD cards simultaneously..
some one says 2.01 opencl.dll works for two simultaneous..
I don't know but seems AMD works with Nvidia opencl.dll if I have Tesla Computing driver
Related khronos icd released tough things to remember are now you can program compatible OpenCL ICD with the doc and also that through ICD some functions which can not be resolved to concrete platform as unload compiler are "no operation"..
Another thing is that 2.01 dll seems has d3d10,d3d9 interop functions(?) or this are getted via
ICD that supports functions not exported through it, I must see..
also Nvidia has d3d11 interop what about AMD?..
Also spec has some cvs links from Khronos for getting some code (ICD loader code?) so someone can mail khronos jon leech for ex. for khronos cvs icd password..
*cudpp has now triangular solvers from 2010 paper..
still waiting for adding sa2009 paper hash functions..
Also a survey has been released saying in which to devote more energies: double supp, graph functions etc..
*bad article by demerijan about Fermi
http://www.semiaccurate.com/2010/02/17/nvidias-fermigtx480-broken-and-unfixable/
but Nvidia seems confident and set clocks for Fermi this week and seems also mid range and other cards taped out some time ago..
*cusp progressing towards dense math(?) has matmul dense and lu solve seems.
* Still clGetGLContextInfoKHR not usable altough present in header 2.01 (was it before in 2.0?)
also some string in Khronos ICD dll but no in lib and dll's really..
*Linux news:
Catalyst 10.2 has direct2d based acceleration search phoronix
also now Noveau has Galluim 3d support in Fedora 3 (working OpenGL ES 2.0 and OpenVG state trackers?)
Heaven benchmark for Linux coming in March for GDC? new version for sure (support for Fermi seems also as now Catalyst 10.2 shows all big sm5.0 features going trough EXT as double support (ext_fp64), shader model 5.0 (ext_shader5), tesselation stuff (ext_tessaltion_shader)
still no standard ext's for HDR new tex compression shipping but no doc and also similar for radnom accest target..
*OpenGL and OpenVG demos:
Some nice code and tutorials found on web:
->OpenGL geometry shader one pass texture cubemap render (3 ways)
->OpenGL GEO culling ->from 2.1billion to 2 million works ATI and Nvidia is 3.2 code..
->Complex OpenVG demo from SA 2009 Khronos presentation (animation)
->OpenGL uniforms vs texture objects.
->Hardware Tessellation on Radeon in OpenGL (geeks3d):
says there are two tesselators in 5xxx extensions
->Mali SDK UI 2.3, Tegra Khronos SDK..
->Code from Stanford Iphone GL ES course
http://www.khronos.org/news/multimedia/optimizing-opengl-for-iphone-stanford-university
-> OpenGL 3.2 samples:
http://nopper.tv/opengl_3_2.html
g-truc ->OpenGL 3 Samples Pack 1.2.1 released
*Also seems WebGL released spec at GDC09 as some talks from Khronos.. also Firefox 3.7 will have it and roadmaps plan for mid year now at alpha 2.
*I have found on ACM video rattner sc09 shows Larrabee demo matmul and sparse math..
More videos are from AMD OpenCL PHD boy..
*Would be nice if optix gets upgraded for:
->Breadth first abd packet ray compression via sort paper EG2010.. improves kernels Timo and Aila used in Optix?
improves raytracing 2x-4x shadows kernels
->Include Sparse Voxel Raycasting I3D 2010 paper
->OpenRL compatibilty.. see diferences are small..
Regarding id3 2010 for me in only remains to be seen stocastic transparency bi Enderton..
Also OpenRL is going to Khronos similar to OpenCL by Apple was.. must check similarities to OpenRT (previous standard )
* Seems AMD drivers for Windows 7 in GDI mode has a bug:
In the same artice some info on GDI accelerated on XP and 7 but not vista..
Also in 7 is in Aero only..
gdi bug 5xxx series:
http://www.tomshardware.com/reviews/2d-windows-gdi,2547-15.html
AMD has supplied hotfix and seems 10.2 WHQL doesn't contain it so perhaps 10.3? or 10.4
good theory about gdi on Windows.. disabled in Vista..
download 2dbench de tomshardware for checking perd..
* opennl 3.0 released having CUDA numerical libraries (CNC and CUSP similar?)
* Sparse voxels octree I3D 2010 paper avaiable and extended NV tech report '10 #1 with more photos and gtx285 perf.. also video and code avaiable in google code cuda voxel raycasting project..
see realtimerendering blog post..
* tegra2 full sdk
has now Android 2.1 images and Khronos full SDK (tegra khronos sdk)..
also seems video compression via OpenMAX in Linux and Android already?..
* Current Catalyst are 10.2 (8.70.2) whql and avaiable 8.70.3( only changes OPenGL version no cal no d3d)
beta given to press 10.3 is ati 8.71.3..
Now about it
3d hooks info is needed and good if enable opengl qb stereo on radeon..
better a sdk as with sample of d3d driver hooks similar to 3d vision is used in Avatar..
*There is a gpu-z enabled opencl ati I don't know if checks correctly or only enables ok..
*Now there are GDC 2010 info from Nvidia in developer.nvidia.com and from intel gdc 2010
From Intel expect:
->GPA 3.0
You'll see in-depth, real-time demos of GPA 3.0, including the much anticipated advanced
thread/task timeline that helps optimize task-based threading. New features such as automated
summarization of your game engine’s performance on multi-core CPUs, the DirectX API, and the
GPU will have you breath-ing a sigh of relief. Platform performance analysis has finally
arrived.
->Intel C++ Compiler version 12 info
This session in-cludes a review of the new automatic vectorization features in the upcoming
Intel C++ Compiler version 12.
->Tickertape
Shows a highly-threaded particle system with orientable quads — like paper in a parade. Particles are affected not only by gravity, but also by air resistance and wind.
*Book Programing ... by Kirk released is CUDA book..
Materials are here:
http://www.elsevierdirect.com/companion.jsp?ISBN=9780123814722
There is also a 3 chapter sample..
*In Khronos I have found a OpenCL NVIDIA build of 2010-02-03
Released soon?
Also a ARM Cortex A9 one:
Samsung Electronics 2010-02-03 OpenCL_1_0
Embedded Linux System with SAMSUNG OpenCL Library with OpenCL running on a ARM Cortex-A9
MPCore CPU.
* Realistic Demo Crymod: Widet2_Benchmark_alpha.7z
*From Caustics:
"due to be released in March"

OpenRL™ SDK Public BETA Registration

Caustic Graphics is about to achieve our next major milestone in bringing cinema quality graphics to every display. We are introducing our OpenRL SDK V1.0 restricted BETA release this week, which is the first implementation of our Open Ray Tracing Language (OpenRL) specification. The OpenRL SDK also includes our new OpenRL shading language (RLSL), which is based on GLSL and provides run-time compiled programmable shaders for ray tracing.

Similar to OpenGL for rasterization, the OpenRL specification is a framework for writing ray tracing applications that execute across heterogeneous compute platforms. Today there is no open standard, cross-platform API for ray tracing. Consequently developers must program their ray tracing applications "to the metal" or accept “vendor lock-in” by using a proprietary closed standard that is limited to a specific subset of hardware.

Later this year, we will be proposing the OpenRL specification as an open standard to the non-profit technology consortium, the Khronos Group. Moreover, we will actively solicit and support the introduction of third-party implementations of OpenRL. In the meantime, we are pleased to introduce the first implementation of the OpenRL specification, which we are calling the OpenRL SDK.

Some quick facts and features slated for the OpenRL SDK:
OS support for Windows, Mac OS X, and Linux;
Uses all OpenCL-based GPUs (e.g., AMD, nVidia, S3) and x86 CPUs (AMD, Intel) simultaneously;
Adding more compute delivers an immediate and nearly linear performance boost;
Plugging in one or more CausticOne or CausticTwo cards delivers the ultimate in ray tracing acceleration.
Target markets include but are not limited to, Film, Video, Games, Transportation, Education, Consumer Products, Architecture, Engineering, and Construction.

We would like to invite you to participate in our OpenRL BETA public program, slated for release this quarter. The OpenRL SDK Public BETA program will include free access to our developer forum where you can post your questions and answers to the OpenRL SDK, RLSL, CausticOne and CausticTwo.

Fill out the form below. Upon release we will send you an email with instructions to download the OpenRL SDK.

P.S. - For those of you who signed up for the CausticRT Emulator, well don't fret. The OpenRL SDK name supersedes CausticRT and CausticGL, whose names will be retired upon release of the production version of the OpenRL SDK.

So OpenCL based and submitting to Khronos..
S3 support intigues me as no driver supports it?

*gdebugger 5.5 with new AMD support for (Catalyst 9.12 and up) performance counters
Also gdebugger cl in beta soon..
*ati OpenCl released 2.01
at least fixes pcchen 8 - knights demo ..
Still no bugs for Apple FFT code fixed but reportedly fixed internally by AMD..
Still not now if OpenCL OpenMM is fixed and about early pyrit builds that now have contermeasures..

*10.6.3 check opengl 3.2 nvidia doubles and cl ati image and ati cal

RAW:
catalyst 10.2 i 10.3 news (8.71.3) 3d qb for d3d (can enable qb 3d ogl via ocl dx ogl interop?)
58xx xbvau not work but patch similar to 4xxx card bug earlier will fix it
fglrx 10.4 ubuntu driver fixed by then..
pgi 10.2 pgaccelinfo has cal info and libamdcal.dylib not found (amd has cal for mac?)
gdc eyefinity sdk?

Catalyst 10.2 has 181 GL extensions!

3 new, 1 EXT, 2 ARB:
GL_ARB_blend_func_extended - more enhancements to blending? whats left in DX10/11 that OGL doesn't have?

GL_ARB_fragment_coord_conventions - DX9 compatibility (wasn't this in OpenGl 3.2?!? still missing transform_feedback2) no estaba en 9.12 hotfix

GL_EXT_texture_buffer_object_rgb32 - this one is interesting as GL_ARB_texture_buffer_object already lists all the RGBA32 F, I, and UI.
ojo vi en fermi 195 drivers

Also I note that 2 amd extensions have been documented:
http://www.opengl.org/registry/specs/AMD/seamless_cubemap_per_texture.txt - when did this get added?
http://www.opengl.org/registry/specs/AMD/shader_stencil_export.txt - from 10.1

Wonder how far away we are from GL 3.3. Still haven't seen DX11 stuff yet, but they must be working on it!

Can't see any sign of the rumored (or under NDA) per-game application profile support yet in CCC. Supposed to be in 10.2...

tesla computing driver released 19.628 64 bits windows 2800 r2: opencl support?, nexus with ati?compute exclusive timeout
*Still no compiler no doubles feb 2010 directx sdk
*fermi 4x slowdown doubles

About Tesla computing driver!

2010-02-18T15:39:00.005+01:00

Hi boys,
I'm becoming increasingly lazy in publishing stories.. sorry for that..
A good megacompilation is coming this week..
Anyway today is old news for installing Tesla Computing driver (196.28) see info (slide 35)

Tesla Compute Cluster (TCC)Driver
Enables Windows HPC on Tesla
Enables Tesla without a NVIDIA graphics card with Windows 7, Server 2008 R2, Windows Vista, Server 2008
Only Tesla 8-series, 10-series and 20-series supported
Only works with CUDA
Does not support OpenGL and DirectX
Available in beta now, release in Jan 2010
Enables the following features under Windows with CUDARDP (Remote Desktop)
Launch CUDA applications via Windows Services
No Windows Timeout issues
No penalty on launch overhead
KVM-over-IP enabled (CPU Server on-board graphics chipset enabled)

Teoretically is for Tesla and Windows 2008R2 64 bit only but I have succefussly installed on Windows 7 on GTX 275..
so installing compute driver is possible on Geforce
just locate NVWD.inf and add
under [NVIDIA_SetA_Devices.NTamd64.6.0]
%NVIDIA_DEV.05E6.01% = Section001, PCI\VEN_10DE&DEV_05E6
and under [NVIDIA_SetA_Devices.NTamd64.6.1]
%NVIDIA_DEV.05E6.01% = Section002, PCI\VEN_10DE&DEV_05E6
this is for GTX 275 for others locate in inf for your card..

Only supports CUDA at the moment (CUDA C(++)).. of course I have checked and PGI Fortran detects it so also CUDA Fortran..
OpenCL doesn't work currently but I think they will add support for it in 200 series..
DirectCompute I have not many hopes.. as device detection is for graphic devices and doesn't expose DirectX..

Have to test Badaboom for seeing if CUVID works and also what about CUDA video encoding via kernels is working (or CUVENC..)
I have to test if OptiX works (I hope but not graphic demos so save render to a file and check)
Also what about PhysX.. Theoretically should work as no interop with graphics is currently enabled at least this enables ATi rendering+PhysX work..
Have to test..
At least has no graphics API dependencies

Also what about other CUDA strong programs as Vreveal..

Last would be good if supported Nexus as I have ATI+Nvidia but Nvidia normal driver only enables CUDA if you extended desktop to *at least ONE Nvidia GPU*..
i.e. if you have 2 nvidias you can use Nexus but ATI+nvidia no work as extending desktop crashes GPU debugger and no extending it doesn't enable Nvidia driver (Windows 7 limitation?)..

Official response:
The Nexus Beta currently only officially supports the 195.62 driver from nvidia.com. Regarding support of the TCC driver, it is not currently supported, but Nexus debugging support using TCC is something we are considering for a future build.

Posted in link above:

Hi,
some questions:
I have tested on a GTX 275 (adding device id to the driver inf) with AMD 5850 as display good work..
(hmm I hope by saying that please don't block this possibilty similar to how are you going to block double prec potential by slowing down 4x on geforce fermi cards..)
By "Only CUDA is supported in this release" you mean OpenCL is supported right now?
I have tested OpenCL ocldeviceQuery and fails to search platform ID and is clear the driver doesn't include opencl.dll nor nvcompiler.dll..
Using that dlls from 196.34 don't work altough nvcuda.dll seems to have Khronos ICD entry points..
are you going to support OpenCL soon on Tesla driver?
I have not tested but are CUDA programs using textures mean to work? I assume yes altough a graphics feature..
Also what about for DirectCompute? i.e for DirectCompute apps not using graphics are you going to support it?
And finally Nexus, I have access to Nexus beta which supports by using two Nvidia cards in one PC debugging in one computer.. is this supported by using as display device an ATI card now one Nvidia card has not to have extended desktop
I hope you add support for Nexus for Tesla Computing driver in case it isn't supported right now..