Revert the Metal Experiment (#701)
Metal sounded like a good idea to get in the emulator but frankly I underestimated just how experimental and not ready it was. From my write up in the Discord: ``` As is, Metal supports only a few games. The games it does support freeze on first use of not playing them via Vulkan, because shader translation is broken. So you need to use a dirty hack to not delete all your shaders. Not to mention it breaks many games via MoltenVK because of changes to the shared GPU code. Merging Metal seemed like a great idea, because of the few games it does support. But I don't think it's worth it. Many of the games it breaks via MoltenVK *don't work via Metal*. Which effectively makes current Ryubing worse for Mac users than Ryujinx 1.1.1403. I think what I'm gonna do is revert Metal, and reopen it as a PR. That way, you can still take advantage of the Metal backend as is, but without making other games worse with no solution. ``` For what it's worth, the shader translation part could at least be "fixed" by always applying a 30ms delay for shader translation to Metal. That being said, that solution sucks ass. The MoltenVK regressions are even worse. I hope this is not a let down to the Mac users. I hope you realize I'm reverting this because you're actively getting a worse experience with it in the emulator.
This commit is contained in:
@@ -1,59 +0,0 @@
|
||||
#include <metal_stdlib>
|
||||
|
||||
using namespace metal;
|
||||
|
||||
struct IndexBufferPattern {
|
||||
int pattern[8];
|
||||
int primitiveVertices;
|
||||
int primitiveVerticesOut;
|
||||
int indexSize;
|
||||
int indexSizeOut;
|
||||
int baseIndex;
|
||||
int indexStride;
|
||||
int srcOffset;
|
||||
int totalPrimitives;
|
||||
};
|
||||
|
||||
struct InData {
|
||||
uint8_t data[1];
|
||||
};
|
||||
|
||||
struct OutData {
|
||||
uint8_t data[1];
|
||||
};
|
||||
|
||||
struct StorageBuffers {
|
||||
device InData* in_data;
|
||||
device OutData* out_data;
|
||||
constant IndexBufferPattern* index_buffer_pattern;
|
||||
};
|
||||
|
||||
kernel void kernelMain(device StorageBuffers &storage_buffers [[buffer(STORAGE_BUFFERS_INDEX)]],
|
||||
uint3 thread_position_in_grid [[thread_position_in_grid]])
|
||||
{
|
||||
int primitiveIndex = int(thread_position_in_grid.x);
|
||||
if (primitiveIndex >= storage_buffers.index_buffer_pattern->totalPrimitives)
|
||||
{
|
||||
return;
|
||||
}
|
||||
|
||||
int inOffset = primitiveIndex * storage_buffers.index_buffer_pattern->indexStride;
|
||||
int outOffset = primitiveIndex * storage_buffers.index_buffer_pattern->primitiveVerticesOut;
|
||||
|
||||
for (int i = 0; i < storage_buffers.index_buffer_pattern->primitiveVerticesOut; i++)
|
||||
{
|
||||
int j;
|
||||
int io = max(0, inOffset + storage_buffers.index_buffer_pattern->baseIndex + storage_buffers.index_buffer_pattern->pattern[i]) * storage_buffers.index_buffer_pattern->indexSize;
|
||||
int oo = (outOffset + i) * storage_buffers.index_buffer_pattern->indexSizeOut;
|
||||
|
||||
for (j = 0; j < storage_buffers.index_buffer_pattern->indexSize; j++)
|
||||
{
|
||||
storage_buffers.out_data->data[oo + j] = storage_buffers.in_data->data[storage_buffers.index_buffer_pattern->srcOffset + io + j];
|
||||
}
|
||||
|
||||
for(; j < storage_buffers.index_buffer_pattern->indexSizeOut; j++)
|
||||
{
|
||||
storage_buffers.out_data->data[oo + j] = uint8_t(0);
|
||||
}
|
||||
}
|
||||
}
|
||||
Reference in New Issue
Block a user