dougallj / applegpu

Apple G13 GPU architecture docs and tools
BSD 3-Clause "New" or "Revised" License
539 stars 38 forks source link

This should be dissasembled? #55

Open tyoc213 opened 9 months ago

tyoc213 commented 9 months ago

Sorry if this doesnt belong here, but Im not sure if this should or not work with an Apple M1, Sonoma 14.2.1 (23C71)

from ast import Tuple
import os
import pathlib
import tempfile
import Metal

prg="""#include <metal_stdlib>
using namespace metal;
kernel void r_5(device int* data0, const device int* data1, uint3 gid [[threadgroup_position_in_grid]], uint3 lid [[thread_position_in_threadgroup]]) {
  int acc0 = -2147483648;
  int val0 = *(data1+0);
  int val1 = *(data1+1);
  int val2 = *(data1+2);
  int val3 = *(data1+3);
  int val4 = *(data1+4);
  int alu0 = max(((val0+1)*2*val0),0);
  int alu1 = max(((val1+1)*2*val1),0);
  int alu2 = max(((val2+1)*2*val2),0);
  int alu3 = max(((val3+1)*2*val3),0);
  int alu4 = max(((val4+1)*2*val4),0);
  int alu5 = max(alu0,acc0);
  int alu6 = max(alu1,alu5);
  int alu7 = max(alu2,alu6);
  int alu8 = max(alu3,alu7);
  int alu9 = max(alu4,alu8);
  *(data0+0) = alu9;
}"""

options = Metal.MTLCompileOptions.new()

print(f"{options=}")
print(f"{options.libraryType=}")
compiler = Metal.MTLCreateSystemDefaultDevice()

def unwrap2(x):
  ret, err = x
  assert err is None, str(err)
  return ret

r = compiler.newLibraryWithSource_options_error_(prg, options, None)
library = unwrap2(r)
l = library.libraryDataContents().bytes().tobytes()
print(l)
print()

with tempfile.NamedTemporaryFile(delete=False) as shader:
  shader.write(l)
  shader.flush()
  os.system(f"python3 compiler_explorer.py {shader.name}")

This reports

newLibraryWithSource:7521: failed assertion `source must not be nil.'

But actually the magic header is 0x00000001 instead of one of the ones it tries to find.

TellowKrinkle commented 4 months ago

We support metal source, metallibs compiled by xcrun metal, and binary archives created with the public binary archive API. Data extracted from a private API on a MTLLibrary is none of those things and appears to be in its own format.

Is there a reason you can't run compiler_explorer.py on the shader source directly?