Unknown RFLT tag generated by macOS 13 Metal compiler

[maleadt]

In working on #79, I noticed that the Apple-produced metallib archives now contain a RFLT entry in the header, with 8 bytes of data.

Patch to make metallib-dis work again:

@@ -139,11 +139,12 @@ struct metallib_program_info {
 		LAYR        = make_tag_type('L', 'A', 'Y', 'R'),
 		TESS        = make_tag_type('T', 'E', 'S', 'S'),
 		SOFF        = make_tag_type('S', 'O', 'F', 'F'),
+		RFLT        = make_tag_type('R', 'F', 'L', 'T'),
 		// generic end tag
 		END         = make_tag_type('E', 'N', 'D', 'T'),
 	};

@@ -377,8 +378,17 @@ static Expected<bool> openInputFile(LLVMContext &Context, std::unique_ptr<ToolOu
 					found_end_tag = true;
 					break;
 				}
+				case metallib_program_info::TAG_TYPE::RFLT: {
+					// unknown 8 bytes of data
+					if(tag_length != 8) {
+						return make_error<StringError>("invalid RFLT size: " + to_string(tag_length),
+													   inconvertibleErrorCode());
+					}
+					outs() << "RFLT: " << to_string(*(const uint64_t*)program_ptr) << '\n';
+					break;
+				}
 				default:

I haven't figured out what the data means yet. It seems to be some sort of offset that's affected by the number of arguments and the fact whether the function contains code (but not the length of the code):

RFLT: 4
kernel void foo() { }

RFLT: 143
kernel void bar() { }

RFLT: 271
kernel void baz() { }

RFLT: 4
kernel void foo() { }

RFLT: 143
kernel void bar(device int* a) { }

RFLT: 383
kernel void baz(device int* b) { }

RFLT: 4
kernel void foo() { }

RFLT: 143
kernel void bar(device int* a) { a[0]=0; }

RFLT: 387
kernel void baz(device int* b) { b[1]=1; }

[maleadt]

Fixed now, it's an offset into the reflection buffer (which itself is still unknown).