Skip to content

Commit

Permalink
Add PixelBuffer
Browse files Browse the repository at this point in the history
  • Loading branch information
mizt committed Oct 21, 2021
1 parent a13ea8e commit 182054d
Show file tree
Hide file tree
Showing 3 changed files with 211 additions and 98 deletions.
217 changes: 125 additions & 92 deletions MTLReadPixels.h
View file
@@ -1,9 +1,11 @@
#import "TargetConditionals.h"
#import "PixelBuffer.h"

#if TARGET_CPU_ARM64 && TARGET_OS_OSX

template <typename T>
class MTLReadPixels {

private:

id<MTLDevice> _device = MTLCreateSystemDefaultDevice();
Expand All @@ -12,26 +14,25 @@ class MTLReadPixels {
id<MTLComputePipelineState> _pipelineState;

dispatch_semaphore_t _semaphore = dispatch_semaphore_create(0);

unsigned int *_bytes;

int _width;
int _height;


bool _isInit = false;


PixelBuffer<T> *_buffer;

public:
MTLReadPixels(int w,int h, NSString *identifier=nil) {
this->_width = w;
this->_height = h;
this->_bytes = new unsigned int[this->_width*this->_height];
for(int k=0; k<this->_width*this->_height; k++) {
this->_bytes[k] = 0xFF000000;
}

MTLPixelFormat PixelFormat8Unorm = MTLPixelFormatRGBA8Unorm;

int width() { return this->_buffer->width(); }
int height() { return this->_buffer->height(); }
int bpp() { return this->_buffer->bpp(); }
void *bytes() { return this->_buffer->bytes(); }
unsigned int rowBytes() { return this->_buffer->rowBytes(); }

MTLReadPixels(int w,int h, int bpp=4, NSString *identifier=nil) {

this->_buffer = new PixelBuffer<T>(w,h,bpp);

#ifdef TARGET_OS_OSX

NSString *path;
Expand All @@ -58,34 +59,53 @@ class MTLReadPixels {
if(!err) {

dispatch_fd_t fd = open([metallib UTF8String],O_RDONLY);

NSDictionary *attributes = [fileManager attributesOfItemAtPath:metallib error:&err];
long size = [[attributes objectForKey:NSFileSize] integerValue];
NSLog(@"%ld",size);

if(size>0) {

dispatch_read(fd,size,dispatch_get_global_queue(DISPATCH_QUEUE_PRIORITY_DEFAULT,0),^(dispatch_data_t d,int e) {

NSError *err = nil;
id<MTLLibrary> library = [this->_device newLibraryWithData:d error:&err];

if(!err) {

MTLTextureDescriptor *texDescriptor = nil;

MTLTextureDescriptor *desc = [MTLTextureDescriptor texture2DDescriptorWithPixelFormat:MTLPixelFormatBGRA8Unorm width:this->_width height:this->_height mipmapped:NO];
desc.usage = MTLTextureUsageShaderRead|MTLTextureUsageShaderWrite;
this->_texture = [this->_device newTextureWithDescriptor:desc];
if(this->_buffer->type()=="f") {
if(bpp==4) {
texDescriptor = [MTLTextureDescriptor texture2DDescriptorWithPixelFormat:MTLPixelFormatRGBA32Float width:this->width() height:this->height() mipmapped:NO];
}
else if(bpp==2) {
texDescriptor = [MTLTextureDescriptor texture2DDescriptorWithPixelFormat:MTLPixelFormatRG32Float width:this->width() height:this->height() mipmapped:NO];
}
else if(bpp==1) {
texDescriptor = [MTLTextureDescriptor texture2DDescriptorWithPixelFormat:MTLPixelFormatR32Float width:this->width() height:this->height() mipmapped:NO];
}
}
else if(this->_buffer->type()=="S"&&bpp==4) {
texDescriptor = [MTLTextureDescriptor texture2DDescriptorWithPixelFormat:MTLPixelFormatRG16Unorm width:this->width() height:this->height() mipmapped:NO];
}
else if(this->_buffer->type()=="I"&&bpp==4) {
texDescriptor = [MTLTextureDescriptor texture2DDescriptorWithPixelFormat:this->PixelFormat8Unorm width:this->width() height:this->height() mipmapped:NO];
}

this->_clip = [this->_device newBufferWithLength:sizeof(unsigned int)*2 options:MTLResourceOptionCPUCacheModeDefault];
unsigned int *clipContents = (unsigned int *)[this->_clip contents];
clipContents[0] = this->_width;
clipContents[1] = this->_height;
id<MTLFunction> function = [library newFunctionWithName:@"copy"];
this->_pipelineState = [this->_device newComputePipelineStateWithFunction:function error:nil];
if(!err) {
this->_isInit = true;
NSLog(@"load copy.metallib");
if(texDescriptor) {
texDescriptor.usage = MTLTextureUsageShaderRead|MTLTextureUsageShaderWrite;
this->_texture = [this->_device newTextureWithDescriptor:texDescriptor];
this->_clip = [this->_device newBufferWithLength:sizeof(unsigned int)*2 options:MTLResourceOptionCPUCacheModeDefault];
unsigned int *clipContents = (unsigned int *)[this->_clip contents];
clipContents[0] = this->width();
clipContents[1] = this->height();
id<MTLFunction> function = [library newFunctionWithName:@"copy"];
this->_pipelineState = [this->_device newComputePipelineStateWithFunction:function error:nil];
if(!err) {
this->_isInit = true;
}
}

}
dispatch_semaphore_signal(this->_semaphore);
close(fd);
Expand All @@ -95,82 +115,95 @@ class MTLReadPixels {
}
}
}

~MTLReadPixels() {
delete[] this->_bytes;
}

unsigned int *bytes() {
return this->_bytes;
}

void setDrawableTexture(id<MTLTexture> src,int type=1) {

if(type>=1&&this->_isInit) {
id<MTLCommandQueue> queue = [this->_device newCommandQueue];
id<MTLCommandBuffer> commandBuffer = queue.commandBuffer;
id<MTLComputeCommandEncoder> encoder = commandBuffer.computeCommandEncoder;
[encoder setComputePipelineState:this->_pipelineState];
[encoder setTexture:src atIndex:0];
[encoder setTexture:this->_texture atIndex:1];
[encoder setBuffer:this->_clip offset:0 atIndex:0];
MTLSize threadGroupSize = MTLSizeMake(8,8,1);
MTLSize threadGroups = MTLSizeMake(
ceil(this->_width/(double)threadGroupSize.width),
ceil(this->_height/(double)threadGroupSize.height),
1);
[encoder dispatchThreadgroups:threadGroups threadsPerThreadgroup:threadGroupSize];
[encoder endEncoding];
[commandBuffer addCompletedHandler:^(id<MTLCommandBuffer> commandBuffer) {
[this->_texture getBytes:this->_bytes bytesPerRow:this->_width<<2 fromRegion:MTLRegionMake2D(0,0,this->_width,this->_height) mipmapLevel:0];
dispatch_semaphore_signal(this->_semaphore);
}];
[commandBuffer commit];
[commandBuffer waitUntilCompleted];
dispatch_semaphore_wait(this->_semaphore,DISPATCH_TIME_FOREVER);
}
else {
[src getBytes:this->_bytes bytesPerRow:(this->_width<<2) fromRegion:MTLRegionMake2D(0,0,this->_width,this->_height) mipmapLevel:0];

void *getBytes(id<MTLTexture> src, bool shader=true) {

if(src) {
if(shader&&this->_isInit) {

id<MTLCommandQueue> queue = [this->_device newCommandQueue];
id<MTLCommandBuffer> commandBuffer = queue.commandBuffer;
id<MTLComputeCommandEncoder> encoder = commandBuffer.computeCommandEncoder;
[encoder setComputePipelineState:this->_pipelineState];
[encoder setTexture:src atIndex:0];
[encoder setTexture:this->_texture atIndex:1];
[encoder setBuffer:this->_clip offset:0 atIndex:0];

int w = this->width();
int h = this->height();

int tx = 1;
int ty = 1;

for(int k=1; k<5; k++) {
if(w%(1<<k)==0) tx = 1<<k;
if(h%(1<<k)==0) ty = 1<<k;
}

MTLSize threadGroupSize = MTLSizeMake(tx,ty,1);
MTLSize threadGroups = MTLSizeMake(w/tx,h/ty,1);

[encoder dispatchThreadgroups:threadGroups threadsPerThreadgroup:threadGroupSize];
[encoder endEncoding];
[commandBuffer addCompletedHandler:^(id<MTLCommandBuffer> commandBuffer) {

[this->_texture getBytes:this->_buffer->bytes() bytesPerRow:this->_buffer->rowBytes() fromRegion:MTLRegionMake2D(0,0,this->_buffer->width(),this->_buffer->height()) mipmapLevel:0];
dispatch_semaphore_signal(this->_semaphore);
}];
[commandBuffer commit];
[commandBuffer waitUntilCompleted];

dispatch_semaphore_wait(this->_semaphore,DISPATCH_TIME_FOREVER);

return this->_buffer->bytes();
}
else {

[src getBytes:this->_buffer->bytes() bytesPerRow:this->_buffer->rowBytes() fromRegion:MTLRegionMake2D(0,0,this->_buffer->width(),this->_buffer->height()) mipmapLevel:0];

return this->_buffer->bytes();
}
}

return nullptr;

}
};

#else

template <typename T>
class MTLReadPixels {

private:

unsigned int *_bytes;

int _width;
int _height;


PixelBuffer<T> *_buffer;

public:

MTLReadPixels(int w,int h) {

this->_width = w;
this->_height = h;

int width() { return this->_buffer->width(); }
int height() { return this->_buffer->height(); }
int bpp() { return this->_buffer->bpp(); }
void *bytes() { return this->_buffer->bytes(); }
unsigned int rowBytes() { return this->_buffer->rowBytes(); }

MTLReadPixels(int w,int h, int bpp=4) {

this->_bytes = new unsigned int[this->_width*this->_height];
for(int k=0; k<this->_width*this->_height; k++) {
this->_bytes[k] = 0xFF000000;
}
this->_buffer = new PixelBuffer<T>(w,h,bpp);
}

void setDrawableTexture(id<MTLTexture> src) {
[src getBytes:this->_bytes bytesPerRow:(this->_width<<2) fromRegion:MTLRegionMake2D(0,0,this->_width,this->_height) mipmapLevel:0];
}
void *getBytes(id<MTLTexture> src, bool shader=false) {
if(src) {
[src getBytes:this->_buffer->bytes() bytesPerRow:this->_buffer->rowBytes() fromRegion:MTLRegionMake2D(0,0,this->_buffer->width(),this->_buffer->height()) mipmapLevel:0];

return this->_buffer->bytes();
}

~MTLReadPixels() {
delete[] this->_bytes;
return nullptr;
}

unsigned int *bytes() {
return this->_bytes;
~MTLReadPixels() {
delete this->_buffer;
}

};
Expand Down
72 changes: 72 additions & 0 deletions PixelBuffer.h
View file
@@ -0,0 +1,72 @@
#import <string>

template <typename T>
class PixelBuffer {

private:

std::string _type = "";

int _width = 0;
int _height = 0;
int _bpp = 0;

void *_bytes = nullptr;
unsigned int _rowbytes = 0;

public:

std::string type() { return this->_type; }

int width() { return this->_width; }
int height() { return this->_height; }
int bpp() { return this->_bpp; }

void *bytes() { return this->_bytes; }
unsigned int rowBytes() { return this->_rowbytes; }

PixelBuffer(int w, int h, int bpp=4) {
this->_width = w;
this->_height = h;
this->_bpp = bpp;
this->_type = @encode(T);
if(this->_type=="I"&&bpp==4) {
int rb = (int)ceil((w*bpp)/4.0)*sizeof(T);
this->_bytes = (void *)new unsigned int[rb*h];
this->_rowbytes = rb;
}
else if(this->_type=="C"&&bpp==4) {
int rb = (w*bpp)*sizeof(T);
this->_bytes = (void *)new unsigned char[rb*h];
this->_rowbytes = rb;
}
else if(this->_type=="S"&&bpp==2) {
int rb = (int)ceil(((w*bpp)/2.0)*2.0)*sizeof(T);
this->_bytes = (void *)new unsigned int[rb*h];
this->_rowbytes = rb;
}
else if(this->_type=="f"&&(bpp==4||bpp==2||bpp==1)) {
this->_bytes = (void *)new float[w*h*bpp];
this->_rowbytes = w*sizeof(T)*bpp;
}
else {
NSLog(@"type = %s, bpp = %d is not supported",this->_type.c_str(),bpp);
this->_type = "";
}
}

~PixelBuffer() {
if(this->_type=="I") {
delete[] (unsigned int *)this->_bytes;
}
else if(this->_type=="C") {
delete[] (unsigned char *)this->_bytes;
}
else if(this->_type=="S") {
delete[] (unsigned int *)this->_bytes;
}
else if(this->_type=="f") {
delete[] (float *)this->_bytes;
}
}
};
20 changes: 14 additions & 6 deletions README.md
View file
@@ -1,15 +1,23 @@
# MTLReadPixels

For some reason [getBytes:bytesPerRow:fromRegion:mipmapLevel:](https://developer.apple.com/documentation/metal/mtltexture/1515751-getbytes?language=objc) for drawable's texture is very slow on M1, so copy drawable's texture by the compute shader and then get the bytes.
For some reason [getBytes:bytesPerRow:fromRegion:mipmapLevel:](https://developer.apple.com/documentation/metal/mtltexture/1515751-getbytes?language=objc) for drawable's texture is very slow on M1and A14, so copy drawable's texture by the compute shader and then get the bytes.

#### macosx
xcrun -sdk macosx metal -c copy.metal -o copy.air; xcrun -sdk macosx metallib copy.air -o copy.metallib

xcrun -sdk macosx metal -c copy.metal -o copy.air; xcrun -sdk macosx metallib copy.air -o copy-macosx.metallib

#### iphoneos

xcrun -sdk iphoneos metal -c copy.metal -o copy.air; xcrun -sdk iphoneos metallib copy.air -o copy.metallib
xcrun -sdk iphoneos metal -c copy.metal -o copy.air; xcrun -sdk iphoneos metallib copy.air -o copy-iphoneos.metallib

#### iphonesimulator

xcrun -sdk iphonesimulator metal -c copy.metal -o copy.air; xcrun -sdk iphonesimulator metallib copy.air -o copy.metallib
xcrun -sdk iphonesimulator metal -c copy.metal -o copy.air; xcrun -sdk iphonesimulator metallib copy.air -o copy-iphonesimulator.metallib

#### remove copy.air

rm ./copy.air


​

0 comments on commit 182054d

Please sign in to comment.