Decode image write block instruction #41
Open
Conversation
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
|
Texture extension is at 62 0001-imageblock-write-Texture-extension.patch Test shaderusing namespace metal;
struct Test {
float4 yay;
};
struct ArgBuf {
texture2d<float, access::write> tex[64];
};
kernel void test(metal::imageblock<Test> f, constant ArgBuf& texlist, ushort2 lid [[thread_position_in_threadgroup]], ushort2 gid [[thread_position_in_grid]]) {
threadgroup_imageblock Test* t = f.data(lid);
t->yay = float4(0, 1, 2, 3);
threadgroup_barrier(mem_flags::mem_threadgroup_imageblock);
if (all(lid == 0)) {
imageblock_slice<float4, imageblock_layout_explicit> slice = f.slice(t->yay);
for (int i = 0; i < 64; i++) {
texlist.tex[i].write(slice, gid);
}
}
}Looks like there's a CoordsDesc and TEX_TYPES in there too 0001-imageblock-write-coordsdesc-and-tex_type.patch (note: I haven't confirmed the flags bit on that) Writing imageblocks to texturecube_arraysusing namespace metal;
struct Test {
float4 yay;
};
kernel void test(metal::imageblock<Test> f, texturecube_array<float, access::write> tex, ushort2 lid [[thread_position_in_threadgroup]], ushort2 gid [[thread_position_in_grid]]) {
threadgroup_imageblock Test* t = f.data(lid);
t->yay = float4(0, 1, 2, 3);
threadgroup_barrier(mem_flags::mem_threadgroup_imageblock);
if (all(lid == 0)) {
imageblock_slice<float4, imageblock_layout_explicit> slice = f.slice(t->yay);
for (int i = 0; i < 64; i++) {
tex.write(slice, gid + ushort2(0, i), 47, 48);
}
}
}Want to make sure these match up with what you're seeing in end-of-tile programs? |
|
Metal's headers pass They also have a Fun with the lod parameterusing namespace metal;
struct Test {
float4 a;
};
kernel void test(metal::imageblock<Test> f, texture2d<float, access::write> tex, ushort2 lid [[thread_position_in_threadgroup]], ushort2 gid [[thread_position_in_grid]], constant uint* fun) {
threadgroup_imageblock Test* t = f.data(lid);
t->a = float4(0, 1, 2, 3);
threadgroup_barrier(mem_flags::mem_threadgroup_imageblock);
if (all(lid == 0)) {
for (int i = 0; i < 64; i++) {
imageblock_slice<float4, imageblock_layout_explicit> slice = f.slice(t->a);
tex.write(slice, gid, i * 16);
}
}
}(I didn't wire up bit 30, so enjoy some slightly broken decompilation) |
This matches regular image_write https://patch-diff.githubusercontent.com/raw/dougallj/applegpu/pull/26.patch ... they're very closely related instructions and execute on the same hw block so it makes sense. |
Also consistent with regular image write |
This instruction ("TODO.unkB1") is used to write out an entire block
from local memory into an image. Because it is block based and not pixel
based, in comparison to the regular image write instruction it works
even if the destination image is compressed. It is tailor fit for use in
the end-of-tile program, to blit tile memory to the framebuffer.
Signed-off-by: Alyssa Rosenzweig <alyssa@rosenzweig.io>
alyssarosenzweig
force-pushed
the
imgwblk
branch
from
7e85274
to
cde6ef5
Compare
|
BTW you might want to change the class name from |
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.
This suggestion is invalid because no changes were made to the code.
Suggestions cannot be applied while the pull request is closed.
Suggestions cannot be applied while viewing a subset of changes.
Only one suggestion per line can be applied in a batch.
Add this suggestion to a batch that can be applied as a single commit.
Applying suggestions on deleted lines is not supported.
You must change the existing code in this line in order to create a valid suggestion.
Outdated suggestions cannot be applied.
This suggestion has been applied or marked resolved.
Suggestions cannot be applied from pending reviews.
Suggestions cannot be applied on multi-line comments.
Suggestions cannot be applied while the pull request is queued to merge.
Suggestion cannot be applied right now. Please check back later.
This instruction ("TODO.unkB1") is used to write out an entire block from local memory into an image. Because it is block based and not pixel based, in comparison to the regular image write instruction it works even if the destination image is compressed. It is tailor fit for use in the end-of-tile program, to blit tile memory to the framebuffer.