[codedivine]

Has anyone experimented with C++ AMP yet?

I downloaded and installed Visual Studio 11 beta, and tried some simple examples.
Code is certainly cleaner, simpler than say OpenCL because a lot of the cruft is now implicit though support for some advanced scheduling etc. seems less clear.
Performance is decent, requires a lot of manual optimization of kernels but that is no different than OpenCL.

While the API is not really introducing any new compute capabilities over say OpenCL, and is limited to VS for the moment, but I do like that MS has put some thought into design and deployment.

1. Code is mostly statically compiled, guaranteed (almost) to work on any DX11 card, with the app vendor not having to worry about whether or not OpenCL driver is working etc.

2. TDRs (timeout detection and recovery) are a big problem for some OpenCL apps, with no good cross-vendor guarantees really. I have had my OpenCL apps crash because the kernel happened to run longer than say 2 seconds, invoking a TDR, and app not being given a chance to respond. Given the huge variety of hardware out there, if I am distributing my app binary, making sure that my kernel does not invoke TDR on any machine at all, is very hard. The only way is to ask end-users to disable TDR by editing registry.

In C++ AMP, if a TDR happens, an exception is thrown which your app can catch and respond to. On Windows 8, you will also get the ability to programmatically declare that you want a card all-to-yourself disabling TDRs as long as that is not connected to a display for example. Finer grained context switching is also coming in the future with WDDM 1.2

3. You do get the benefits of integration into VS, which IMO is a great IDE.

However, the big downside is of course that it is new, with not much of a library ecosystem and completely dependent on VS and Windows for now.

[AlexV]

Originally Posted by codedivine

Has anyone experimented with C++ AMP yet?

I've done a bit of work with it. I love it quite a bit, to be honest. It just fits in far more naturally than everything else I've tried.

What do you have in mind WRT the more advanced scheduling aspect? I don't think CL's:
enqueMadnessandDamnationOutOfOrderOutOfSightOutOfC haracters(CL_GL_KL_DOOT_DOOT_EXT_KHR_ATI_AMD_NVIDI A_APPLE_SGI_S3_SIS_PRINT_LULZ, *ptr_to_struct, *ptr_to_array_of_structs_of_ptrs_to_enums, ...)
queuing mechanism is notably more advanced, it's just annoyingly verbose (this seems to apply in general).

You have control over how the submission queues attached to one accelerator_view or another are created (either immediate or deferred), and you have a few async ops that return convenient std::future<void>s(IIRC). There are also means to create markers in the API, for a bit of added oomph. Of course, the deferred queuing aspect is reliant on having deferred contexts working in D3D drivers, and that's not all that certain these days (I think NV are the only ones who do it).

[Jawed]

What does AMP do that C++ bindings for OpenCL or AMD's static C++ for OpenCL not do?

As long as AMP is Windows only, it's not going anywhere is it?

[pcchen]

Originally Posted by Jawed

What does AMP do that C++ bindings for OpenCL or AMD's static C++ for OpenCL not do?

OpenCL is significantly more verbose than AMP. Of course, almost everything AMP can do, you can do it with OpenCL. However, it's just much easier with AMP

As long as AMP is Windows only, it's not going anywhere is it?

Well, Microsoft does not restrict anyone from making their own AMP implementations, so this is hardly a problem if there's demand for AMP.

[AlexV]

Originally Posted by Jawed

What does AMP do that C++ bindings for OpenCL or AMD's static C++ for OpenCL not do?

As long as AMP is Windows only, it's not going anywhere is it?

The answer to the first point is quite simple: be productive and portable. If I want to use something non-portable (as AMD's static C++ stuff is), I might as well use CUDA, because that's actually pretty swanking cool - I don't want to use something non-portable. That's without even considering how that particular proposal misses a few points about what makes C++ useful (note: If I have to do the name mangling manually, and I still do dumb passing around of kernel args, it's not the right call). As for the C++ bindings, you can only wrap so much, the underlying API is the problem. They're also not the most well documented or most solid things in the world (I need a macro def for them to work with 1.2? some stuff dubiously breaks? not nice! It's a good effort, but hamstrung by the underlying API and the lack of resources).

As long as AMP is Windows only it's only going to be the choice for consumer apps, because it's much better to bet on MS actually putting something useful out and having some sort of direction, as opposed to hoping that the perpetual tug of war that is Khronos will eventually figure it out. That's a pretty big place to go to, IMHO, as it's the only other avaliable slot. CUDA has the "pro" people - the ones who actually have money in this, not the ones who just need to get another paper or thesis out to finish their terms, but the ones who are ready to give an IHV a large wad of cash. It's not by accident, or by lack of people trying to use OCL.

[Dade]

Originally Posted by AlexV

I don't want to use something non-portable.

I'm too .. and this is exactly the reason I would pick OpenCL over C++ AMP: it is available only for Windows, how can you consider portable something that is available for a single platform ?

I agree on the the "productive" argument: developing something more complex than a few lines long OpenCL kernel is very time consuming. OpenCL definitively needs to introduce a lot more features in this area.

Originally Posted by AlexV

As long as AMP is Windows only it's only going to be the choice for consumer apps, because it's much better to bet on MS actually putting something useful out and having some sort of direction

I have some doubt, GPU computing is driven by hardware and MS has no control at all. AMD/Intel/Apple/IBM should have more interest and the tools to push and innovate OpenCL.

[entity279]

Originally Posted by Dade

I'm too .. and this is exactly the reason I would pick OpenCL over C++ AMP: it is available only for Windows, how can you consider portable something that is available for a single platform ?

The context in which portability was discussed was AMP vs AMD's C++ OCL bindings. Not AMP vs OCL.

[Drazick]

Could you show some examples of code?
For instance, apply Gaussian Blur on an image or something.

Thanks.

[pcchen]

There are several samples provided by Microsoft here:

http://blogs.msdn.com/b/nativeconcur...-download.aspx

[Ethatron]

I also have some short fragment of OpenMP vs. CR (Concurrency runtime) vs. AMP:

      /* forward/backward-z */
      ULONG v = 0x00808000 | (TCOMPRESS_NINDEP(format) ? 0x00 : 0xFF);
      ULONG c =              (TCOMPRESS_SIDES (format) ? 0x00 : 0xFF) << 24;

#if    defined(USE_AMP)
      Concurrency::extent<1> num(texo.Width * texo.Height);
      Concurrency::array_view<ULONG, 1> sArr(num, sTex);

      Concurrency::parallel_for_each(num, [=](index<1> elm) restrict(amp) {
    ULONG t = sArr[elm] | c;

    /* black matte to forward-z */
    /**/ if (( t & 0x00FFFFFF) == 0)
      t = (t | v);
    /* white matte to forward-z (can't be done here, as white is valid for partial derivatives)
    else if ((~t & 0x00FFFFFF) == 0)
      t = (t & v); */

    sArr[elm] = t;
      });
#elif    defined(USE_CCR)
      parallel_for(0, (int)(texo.Width * texo.Height), [=](int elm) {
    ULONG t = sTex[elm] | c;

    /* black matte to forward-z */
    /**/ if (( t & 0x00FFFFFF) == 0)
      t = (t | v);
    /* white matte to forward-z (can't be done here, as white is valid for partial derivatives)
    else if ((~t & 0x00FFFFFF) == 0)
      t = (t & v); */

    sTex[elm] = t;
      });
#else
#pragma omp parallel for schedule(dynamic, 4)        \
             shared(sTex)
      for (int y = 0; y < (int)texo.Height; y += 1) {
      for (int x = 0; x < (int)texo.Width ; x += 1) {
    ULONG t = sTex[(y * texo.Width) + x] | c;

    /* black matte to forward-z */
    /**/ if (( t & 0x00FFFFFF) == 0)
      t = (t | v);
    /* white matte to forward-z (can't be done here, as white is valid for partial derivatives)
    else if ((~t & 0x00FFFFFF) == 0)
      t = (t & v); */

    sTex[(y * texo.Width) + x] = t;
      }
      }
#endif

Just to show how small the changes in "normal C++ programming" from API to API are. The code just replaces black in normal-maps by the "color" of backward-facing z-vectors, and white by forward-facing. "sTex" is the continous memory of a DDS-texture's mip-level aquired with LockRect.

[rpg.314]

http://arstechnica.com/information-t...-on-windows-8/

[Miksu]

Originally Posted by rpg.314

http://arstechnica.com/information-t...-on-windows-8/

You can still do Windows 8 Metro apps with VS Express, but the "old" desktop apps are gone. But I don't think MS has at any point planned to release C++ AMP for VS Express?

[sebbbi]

Originally Posted by Miksu

But I don't think MS has at any point planned to release C++ AMP for VS Express?

Incorrect. C++ AMP works fine also in the free Visual Studio Express version (Metro is compatible with C++ AMP, assuming the documentation is correct). You can't use the free toolset to create traditional Windows applications anymore (but C++ AMP is not in the chopping block). I don't know if this is the proper thread to discuss Microsoft policies about the free Visual Studio Express, this thread is about C++ AMP. I dislike the limitations in the free version as much as anyone else here. Maybe there should be a separate thread to discuss the issue of Microsoft cutting off standard C/C++ development from the free development tools?

I haven't personally tested C++ AMP in Metro applications, but it sounds like an fascinating idea, if it works also in devices with ARM CPUs (and mobile GPUs). Chips with DX 9_3 feature level obviously cannot run AMP by GPU, but mobile GPU developers have already announced DX 11 capable chips (first ones should be released later this year, maybe we have those in Windows 8 launch). Would be interesting to see how ARM Mali-T604 and PowerVR Rogue compare against AMD, Intel and NVidia in GPU compute.

There's a chance that C++ AMP becomes an industry wide standard. The restrict keyword has been sent to the standards committee (it's the only new language feature required for C++ AMP programs to work). AMD has helped Microsoft a lot with the C++ AMP implementation, and it would be possible that they are interested in porting it to other platforms as well (just like they have been optimizing 7-zip and others with OpenCL and giving them free to community).

[johnhamlen]

Originally Posted by sebbbi

Incorrect. C++ AMP works fine also in the free Visual Studio Express version (Metro is compatible with C++ AMP, assuming the documentation is correct). You can't use the free toolset to create traditional Windows applications anymore (but C++ AMP is not in the chopping block)

Dear Sebbbi. This is very good news indeed!

I'm about to embark on my first spot of GPGPU development (Monte Carlo playouts for a chess program) and working out my weapon of choice. CUDA looked nice, but obviously Nvidia GPU specific, so was about to head down the OpenCL route when C++ AMP appeared on the horizon. Looks ideal for clean code, but not keen to spend a bundle on VC++ if I don't have to as this project is a just a hobby experiment.

Was concerned about the following http://msdn.microsoft.com/en-us/libr...1461a(v=vs.110)

Visual Studio Express provides support for creating and deploying Metro style applications and components only. The C++ standard libraries are available to Visual Studio Express, but cannot be used to create desktop or console applications. For these features, you would have to upgrade to another version of Visual C++. For more information about these editions, see Visual C++ Editions. The documentation that is included with Visual Studio Express is a subset of the documentation of other versions of Visual C++. Therefore, you might encounter links to topics that are not available. Complete documentation is available at http://msdn.microsoft.com/library.

Sorry for the newbie question but can you explain what I need to get a VS Express development environment for C++ AMP up and running? E.g. Download VSE 2010 and plug in the the xxx library files from yyy.

Many thanks in advance, John

[Dade]

Originally Posted by rpg.314

http://arstechnica.com/information-t...-on-windows-8/

Oh, something new, MS's "open" and "free" standards are awesome

[pcchen]

IIRC it's already the case for VS Express beta (at least when last time I checked, which was last year). However, I still think this is a bad idea. Of course, they are going to keep older VS Express around, but I don't understand why they make it this way.

Of course, w.r.t. C++ AMP anyone can implement their own version in their own compiler

[imaxx]

Originally Posted by pcchen

but I don't understand why they make it this way.

...because once upon a time (start of '90), you had to pay for all compilers - MS included.
MS lost half of its value in market (Ballmer rulez), so they probably try to go the greedy way...
pecunia non olet

[pcchen]

Originally Posted by imaxx

...because once upon a time (start of '90), you had to pay for all compilers - MS included.

Well, AFAIK gcc is always free.

Of course, developer tools do cost serious money to make, but making it free serves an important purpose: to push your platform. Maybe Microsoft don't believe they need the push anyway

[Ethatron]

If the DirectCompute compiler wouldn't kill it self to optimize the (valid ofc) AMP-code I feed it with, we may have had squish (DXT-library) as AMP-version today. But apparently I have to fight the compiler instead of passing on to make a BC6/7 compressor based on squish.

The compiler is very sensitive to template programming currently as well, and bails out with "OMG too complex for me" error in a lot of cases.

Additionally it seems the DCC(directccomputecompiler)-part is multithreaded and I guess because parts of it are in the driver stack it just saturates the machine at a non-responsive 100%, I always start to drop the priority of cl when it appears in the taskmanager, otherwise I'm locked out of the machine for the time needed to compile - squish-amp currently haven't been observed finishing compiling though. %^)

It has it's hickups, and if the complexity/time problems are not adressed in the final product, it's just a unfulfilled promise of low hanging TFLOPS, which you pay with doing code-migration (from smart to simple, from compact to duplicated, from C++ to HLSL, etc.) to get them possibly if you're persistant.

[sebbbi]

For me it has been working well so far, but I haven't coded anything large with it, just a parallel prefix sum and a radix sorter. But I wouldn't be surprised if there are some bugs left as it's still beta after all.

A link comparing DirectCompute to C++ AMP. This clearly shows that DirectCompute requires tons on boilerplate to get even a simple algorithm running. C++ AMP requires almost nothing in addition to the lambda that describes the kernel:
http://blogs.msdn.com/cfs-file.ashx/...Programmer.pdf

[Ethatron]

What I found absolute amazing is the ability to go from AMP to CPU, I wanted to debug the AMP code before first run, sadly, no AMP debugger without W8, so, what do? Implement software-AMP!

Here's the amazing part:

#define    tile_static

    template<typename type>
    class array_view {
    public:
      int w, h; type *arr;
      array_view(int ww, int hh, type *aa) { w = ww; h = hh; arr = aa; }
      type& operator() (const int &x, const int &y) { return arr[y*w+x]; }
    };

    array_view<const unsigned int> sArr(iwidth, iheight, (const unsigned int *)texs);
    array_view<      unsigned int> dArr(cwidth, cheight, (      unsigned int *)texr);

    for (int groupsy = 0; groupsy < (owidth  / TY); groupsy++)
    for (int groupsx = 0; groupsx < (oheight / TX); groupsx++) {
      typedef type accu[DIM];

//    tile_static UTYPE bTex[2][TY][TX];
      tile_static type  fTex[2][TY][TX][DIM];
      tile_static int   iTex[2][TY][TX][DIM];
      
      for (int tiley = 0; tiley < TY; tiley++)
      for (int tilex = 0; tilex < TX; tilex++) {
        const ULONG &t = sArr(posy, posx);

And a few of those:

#if    defined(USE_AMP)
#define local_is(a,b) elm.local == index<2>(a, b)
#else
#define local_is(a,b) ((ly == a) && (lx == b))
#endif

In 5 minutes I had a AMP-to-CPU switch. Now I'm going to proof-code squish - which I actuall got to finish compiling now, without the "iterative search" option, but 15 minutes it takes - change all the things I should adjust for the sake of the GPU, shorted 4-wide SSE vectors to 3-wide etc. pp. It can/should already run on 3 threads of 16, it's weird but fun to think the way you have to think for compute.

Too bad I don't see a way to look at the HLSL (or IL), I'd love to get a feeling of the amount of insns, I'm totally clueless as to the complexity of the produced CS.

BTW what happened with your BC6/7 efforts sebbi?

[AlexV]

Originally Posted by Ethatron

If the DirectCompute compiler wouldn't kill it self to optimize the (valid ofc) AMP-code I feed it with, we may have had squish (DXT-library) as AMP-version today. But apparently I have to fight the compiler instead of passing on to make a BC6/7 compressor based on squish.

The compiler is very sensitive to template programming currently as well, and bails out with "OMG too complex for me" error in a lot of cases.

Additionally it seems the DCC(directccomputecompiler)-part is multithreaded and I guess because parts of it are in the driver stack it just saturates the machine at a non-responsive 100%, I always start to drop the priority of cl when it appears in the taskmanager, otherwise I'm locked out of the machine for the time needed to compile - squish-amp currently haven't been observed finishing compiling though. %^)

It has it's hickups, and if the complexity/time problems are not adressed in the final product, it's just a unfulfilled promise of low hanging TFLOPS, which you pay with doing code-migration (from smart to simple, from compact to duplicated, from C++ to HLSL, etc.) to get them possibly if you're persistant.

I think that reporting the breakage (including a repro) directly to the AMP team, through their dedicated blog, is probably the best way to get it fixed. In my experience they're quite responsive and quick to fix things / buggzorz. The conduit for doing the reporting would be http://blogs.msdn.com/b/nativeconcurrency/.

[CouldntResist]

Are we seriously comparing language agnostic API (CUDA, OpenCL) to a combination of specific language library and specific compiler extension?

Positioning them as competitors is fallacious. Any hypothetical non Windows implementation of C++ AMP would have to be implemented in terms of OpenCL or CUDA.

Have a look at ScalaCL. It's a combination of Scala library and Scala compiler extension, acting as abstraction layer over OpenCL. If you want to have something to compare C++ AMP with, knock yourself out.

Originally Posted by sebbbi

You cant basically rely on CUDA if you are programming games

Similarly, you can't rely on C++ AMP if you are programming game with linux dedicated server, and want to use GPGPU for anything other than pixels.

[sebbbi]

Originally Posted by CouldntResist

Any hypothetical non Windows implementation of C++ AMP would have to be implemented in terms of OpenCL or CUDA.

Of course, but that's just an implementation detail. Nobody cares if there's OpenCL or CUDA or DirectCompute (or 1028 bit AVX2/gather CPU implemention) under the hood, as long as the code works properly, is easy to write, easy to maintain and easy use (and understandable). C++ AMP isn't meant to be a low level GPU API. It's a higher level API.

Originally Posted by CouldntResist

Similarly, you can't rely on C++ AMP if you are programming game with linux dedicated server, and want to use GPGPU for anything other than pixels.

C++ AMP is still beta, so you cannot even use it in Windows games yet. When it is out, you can use Windows Server 2008 R2 or 2012 to run your dedicated game servers, if you need to run C++ AMP in your server code as well. Or you can use C++ AMP solely for your client side code. It's really easy to port, as the kernels are pure C++ lambda functions (just define out those restrict keywords, and do wrappers for those array view templates).

Also running any kind of GPGPU in a scalable server farm of virtual machines is currently very difficult. Most service providers do not offer any kind of GPGPU possibility, as GPU virtualization techniques are still very much under development. It will take several years (or maybe even a decade) until GPGPU gets common in server world (there are so many unsolved problems). Of course a big virtualized cloud of GPGPUs is a dream of many, but it will take time until we have solid products and service producers have modified their infrastructures. These changes cost a lot (and GPU programming is still a moving target. Differerent kinds of CPUs are pretty much identical, so the actual hardware can be easily hidden behind virtualization).

Originally Posted by CouldntResist

Have a look at ScalaCL. It's a combination of Scala library and Scala compiler extension, acting as abstraction layer over OpenCL.

Scala isn't C++. We have a million of lines of existing C++ code, and I want to use my existing C++ structures/templates/classes/enumerations/etc and the metaprogramming libraries inside my kernels. C++ AMP integrates very well to existing C++ projects, and it's easy to write reusable code with it. ScalaCL likely brings similar productivity boost to programmers that use Scala. That's a good thing (nobody wants to program with pure OpenCL or DirectCompute if they don't have to).

I am not saying that C++ AMP is a perfect solution (especially if it does not get industry wide support behind it). I'd rather see a similar proposition from the C++ standards committee. C++ AMP at least gives the example of a proper C++ based parallel API with full compiler and IDE support. I personally do not find anything in it that feels like a hack. It's simple, easy to read, uses C++ constructs for everything and requires no boiler plate. Have you even tried to code with it before criticizing it?

[rpg.314]

Originally Posted by CouldntResist

Positioning them as competitors is fallacious. Any hypothetical non Windows implementation of C++ AMP would have to be implemented in terms of OpenCL or CUDA.

C++ AMP is built on top of Direct compute too.

Have a look at ScalaCL. It's a combination of Scala library and Scala compiler extension, acting as abstraction layer over OpenCL. If you want to have something to compare C++ AMP with, knock yourself out.

ScalaCL, as I understood it from a brief look, does not handle control flow.