Logo Questions Linux Laravel Mysql Ubuntu Git Menu
 

Memory write performance - GPU CPU Shared Memory

I'm allocating both input and output MTLBuffer using posix_memalign according to the shared GPU/CPU documentation provided by memkite.

Aside: it is easier to just use latest API than muck around with posix_memalign

let metalBuffer = self.metalDevice.newBufferWithLength(byteCount, options: .StorageModeShared)

My kernel function operates on roughly 16 million complex value structs and writes out an equal number of complex value structs to memory.

I've performed some experiments and my Metal kernel 'complex math section' executes in 0.003 seconds (Yes!), but writing the result to the buffer takes >0.05 (No!) seconds. In my experiment I commented out the math-part and just assign the zero to memory and it takes 0.05 seconds, commenting out the assignment and adding the math back, 0.003 seconds.

Is the shared memory slow in this case, or is there some other tip or trick I might try?

Additional detail

Test platforms

  • iPhone 6S - ~0.039 seconds per frame
  • iPad Air 2 - ~0.130 seconds per frame

The streaming data

Each update to the shader receives approximately 50000 complex numbers in the form of a pair of float types in a struct.

struct ComplexNumber {
    float real;
    float imaginary;
};

Kernel signature

kernel void processChannelData(const device Parameters *parameters [[ buffer(0) ]],
                               const device ComplexNumber *inputSampleData [[ buffer(1) ]],
                               const device ComplexNumber *partAs [[ buffer(2) ]],
                               const device float *partBs [[ buffer(3) ]],
                               const device int *lookups [[ buffer(4) ]],
                               device float *outputImageData [[ buffer(5) ]],
                               uint threadIdentifier [[ thread_position_in_grid ]]);

All the buffers contain - currently - unchanging data except inputSampleData which receives the 50000 samples I'll be operating on. The other buffers contain roughly 16 million values (128 channels x 130000 pixels) each. I perform some operations on each 'pixel' and sum the complex result across channels and finally take the absolute value of the complex number and assign the resulting float to outputImageData.

Dispatch

commandEncoder.setComputePipelineState(pipelineState)

commandEncoder.setBuffer(parametersMetalBuffer, offset: 0, atIndex: 0)
commandEncoder.setBuffer(inputSampleDataMetalBuffer, offset: 0, atIndex: 1)
commandEncoder.setBuffer(partAsMetalBuffer, offset: 0, atIndex: 2)
commandEncoder.setBuffer(partBsMetalBuffer, offset: 0, atIndex: 3)
commandEncoder.setBuffer(lookupsMetalBuffer, offset: 0, atIndex: 4)
commandEncoder.setBuffer(outputImageDataMetalBuffer, offset: 0, atIndex: 5)

let threadExecutionWidth = pipelineState.threadExecutionWidth
let threadsPerThreadgroup = MTLSize(width: threadExecutionWidth, height: 1, depth: 1)
let threadGroups = MTLSize(width: self.numberOfPixels / threadsPerThreadgroup.width, height: 1, depth:1)

commandEncoder.dispatchThreadgroups(threadGroups, threadsPerThreadgroup: threadsPerThreadgroup)
commandEncoder.endEncoding()
metalCommandBuffer.commit()
metalCommandBuffer.waitUntilCompleted()

GitHub example

I've written an example called Slow and put it up on GitHub. Seems the bottleneck is the write of the values into the input Buffer. So, I guess the question becomes how to avoid the bottleneck?

Memory copy

I wrote up a quick test to compare the performance of various byte copying methods.

Current Status

I've reduced execution time to 0.02ish seconds which doesn't sound like a lot, but it makes a big difference in the number of frames per second. Currently the biggest improvements are a result of switching to cblas_scopy().

like image 970
Cameron Lowell Palmer Avatar asked Mar 08 '16 06:03

Cameron Lowell Palmer


People also ask

Can shared GPU memory be used for gaming?

Yes. Because shared Memory is essentially RAM, if a GPU has to resort to using the System's RAM for its computations, it'll take a performance hit.

Does shared memory improve performance?

If your application is making use of shared memory, you'd expect to see increased performance compared to an implementation using only global memory. Because it is on-chip, shared memory has a much higher bandwidth and lower latency than global memory.

What is GPU shared memory?

What exactly is shared GPU memory? This is a sort of virtual memory that's used when your GPU runs out of its dedicated video memory. It's usually half of the amount of RAM you have. Operating systems use RAM as a source of shared GPU memory because it's much faster than any HDD or SSD, even the PCIe 4.0 models.

Can I use shared memory GPU?

Shared memory represents system memory that can be used by the GPU. Shared memory can be used by the CPU when needed or as “video memory” for the GPU when needed. If you look under the details tab, there is a breakdown of GPU memory by process. This number represents the total amount of memory used by that process.


2 Answers

Reduce the size of the type

Originally, I was pre-converting signed 16-bit sized Integers as Floats (32-bit) since ultimately that is how they'll be used. This is a case where performance starts forcing you to store the values as 16-bits to cut your data-size in half.

Objective-C over Swift

For the code dealing with movement of data, you might choose Objective-C over Swift (Warren Moore recommendation). Performance of Swift in these special situations still isn't up to scratch. You can also try calling out to memcpy or similar methods. I've seen a couple of examples that used for-loop Buffer Pointers and this in my experiments performed slowly.

Difficulty of testing

I really wanted to do some of the experiments with relation to various copying methods in a playground on the machine and unfortunately this was useless. The iOS device versions of the same experiments performed completely differently. One might think that the relative performance would be the similar, but I found this to also be an invalid assumption. It would be really convenient if you could have a playground that used the iOS device as the interpreter.

like image 68
Cameron Lowell Palmer Avatar answered Oct 11 '22 12:10

Cameron Lowell Palmer


You might get a large speedup via encoding your data to huffman codes and decoding on the GPU, see MetalHuffman. It depends on your data though.

like image 31
MoDJ Avatar answered Oct 11 '22 11:10

MoDJ