Archives Discussions

[hongfeeling]

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

Is there any way to do this?

[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;

}

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

[hongfeeling]

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

Is GCN ASM  developed by AMD?

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

[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

[hongfeeling]

yes, cl_ext_atomic_counters_32 will touch GDS.

But I found 'cl_ext_atomic_counters_32' cannot running on Carrizo, it will cause system TDR and never recovery on windows.

[tzachi_cohen]

How did you get a hold of Carrizo? It is not commercially available yet.