7 Replies Latest reply on Jan 1, 2015 11:04 PM by tzachi.cohen

    How to read/write GDS through AMD APP SDK

    hongfeeling

      Now I want to read and write some data from or to GDS block.

      Is there any way to do this?

        • Re: How to read/write GDS through AMD APP SDK
          hongfeeling

          Now, Add Inline asm to touch GDS. But still need info for "asm". I don't know how to get value from outside to asm.

          __kernel void helloworld(__global char* in, __global char* out)

          {

            __local float localBuffer[64];

            uint tx = get_local_id(0);

            uint gx = get_global_id(0);

            // Initialize local memory:

            // Copy from this work-group?s section of global memory to local:

            // Each work-item writes one element; together they write it all

           

            asm ( " dcl_gds_id(3) 256 ");

            localBuffer[tx] = in[gx];

            // Ensure writes have completed:

            //asm( "gds_store_id(3) r33.x___, r33.x" );

            barrier(CLK_LOCAL_MEM_FENCE);

           

            //define constant ( offset, char , ignor,ignor )

            asm (" dcl_literal l1, 0x00000010, 0x00000044, 0x00000002, 0x00000003 ");

            asm (" mov r1, l1 ");

            asm ("gds_store_id(3) r1.x, r1.y ");

            asm ("gds_load_id(3) r2.x, r1.x ");

           

           

           

            // Toy computation to compute a partial factorial, shows re-use from local

            int f = localBuffer[tx];

            asm (" mov r33.x___,r2.x ");

            out[gx] = f;

          }

            • Re: How to read/write GDS through AMD APP SDK
              realhet

              Hi,

               

              As I know, there is no way to map OpenCl variables to AMD_IL registers. The only way I know is to write the whole thing in AMD_IL or in GCN ASM.

                • Re: How to read/write GDS through AMD APP SDK
                  hongfeeling

                  Could you tell me how to running a shader which whole write by AMD_IL.

                  Is GCN ASM  developed by AMD?

                    • Re: How to read/write GDS through AMD APP SDK
                      realhet

                      Hi,

                      I guess it's a bit late, but anyways...

                       

                      GCN ASM is the lowest level language that is executed by the GCN type of AMD GPUs. Of course it's developed by them. They documented it in the "Sea Islands Instruction Set Architecture" manual.

                       

                      To upload any custom program to the gpu, you have to make a binary .elf image, and load it with the clBuildProgramWithBinary() function.

                       

                      You'll also need a compiler that can make that binary from gcn_asm: You can try my assembler if you think: It's here -> http://realhet.wordpress.com/2012/11/14/hello-world/  But as I haven't got the time to catch up with the new drivers, you have to use 13.4 or older Catalyst with it (imho the best driver for development is 12.10 as you have a working disassembler too, but unfortunately that only knows older 7xxx cards). Otherwise it is compatible with the 14.9 Catalyst, the only problem is with the different kernel parameter passing in the new drivers. I think you can start from a helloworld example and try out asm instructions. Also you can always learn from disassembled OpenCL programs.

                • Re: How to read/write GDS through AMD APP SDK
                  tzachi.cohen

                  The only high level extension that uses GDS is 'cl_ext_atomic_counters_32' extension. It only implicitly exposes GDS as a counter, it does not enable arbitrary access.

                   

                  Tzachi