def test_structs(self):
v = Metal.MTLDispatchThreadgroupsIndirectArguments()
self.assertEqual(v.threadgroupsPerGrid, None)
v = Metal.MTLStageInRegionIndirectArguments()
self.assertEqual(v.stageInOrigin, None)
self.assertEqual(v.stageInSize, None)
def test_structs(self):
v = Metal.MTLSizeAndAlign()
self.assertEqual(v.size, 0)
self.assertEqual(v.align, 0)
v = Metal.MTLAccelerationStructureSizes()
self.assertEqual(v.accelerationStructureSize, 0)
self.assertEqual(v.buildScratchBufferSize, 0)
self.assertEqual(v.refitScratchBufferSize, 0)
def main(argv):
radius = ''
side = ''
age = ''
material = ''
try:
opts, args = getopt.getopt(argv, "hr:s:a:m:",
["radius=", "side=", "age=", "material="])
except getopt.GetoptError:
print("Error")
usage()
sys.exit(2)
for opt, arg in opts:
if opt == '-h':
usage()
sys.exit()
elif opt in ("-r", "--radius"):
radius = int(arg)
elif opt in ("-s", "--side"):
side = int(arg)
elif opt in ("-a", "--age"):
age = int(arg)
elif opt in ("-m", "--material"):
material = str(arg)
square = Square(side)
metal = Metal(material)
ancientObject = AncientObject(age, metal)
myCoin = ClassicalChinaCoin(radius, square, ancientObject)
print(myCoin.toString())
if myCoin.isValid():
print("\nCoin is Valid")
else:
print("\nCoin is Invalid")
square.setSide(2)
metal.setMaterial("gold")
ancientObject.setAge(700)
ancientObject.setMetal(metal)
myCoin.setRadius(5)
myCoin.setSquare(square)
myCoin.setAncientObject(ancientObject)
print("\n" + myCoin.toString())
ancientObject.destroy()
print("\nCoin Destroyed\n\n" + myCoin.toString())
def wrapper(*args):
args = [vaex.array_types.to_numpy(ar) for ar in args]
def getbuf(name, value=None, dtype=np.dtype("float32"), N=None):
buf = getattr(storage, name, None)
if value is not None:
N = len(value)
dtype = value.dtype
if dtype.name == "float64":
warnings.warn("Casting input argument from float64 to float32 since Metal does not support float64")
dtype = np.dtype("float32")
nbytes = N * dtype.itemsize
if buf is not None and buf.length() != nbytes:
# doesn't match size, create a new one
buf = None
# create a buffer
if buf is None:
buf = self.device.newBufferWithLength_options_(nbytes, 0)
setattr(storage, name, buf)
# copy data to buffer
if value is not None:
mv = buf.contents().as_buffer(buf.length())
buf_as_numpy = np.frombuffer(mv, dtype=dtype)
buf_as_numpy[:] = value.astype(dtype, copy=False)
return buf
input_buffers = [getbuf(name, chunk) for name, chunk in zip(self.arguments, args)]
output_buffer = getbuf('vaex_output', N=len(args[0]), dtype=dtype_out)
buffers = input_buffers + [output_buffer]
command_buffer = command_queue.commandBuffer()
encoder = command_buffer.computeCommandEncoder()
encoder.setComputePipelineState_(state)
for i, buf in enumerate(buffers):
encoder.setBuffer_offset_atIndex_(buf, 0, i)
nitems = len(args[0])
tpgrid = Metal.MTLSize(width=nitems, height=1, depth=1)
# state.threadExecutionWidth() == 32 on M1 max
# state.maxTotalThreadsPerThreadgroup() == 1024 on M1 max
tptgroup = Metal.MTLSize(width=state.threadExecutionWidth(), height=state.maxTotalThreadsPerThreadgroup()//state.threadExecutionWidth(), depth=1)
# this is simpler, and gives the same performance
# tptgroup = Metal.MTLSize(width=1, height=1, depth=1)
encoder.dispatchThreads_threadsPerThreadgroup_(tpgrid, tptgroup)
encoder.endEncoding()
command_buffer.commit()
command_buffer.waitUntilCompleted()
output_buffer_py = output_buffer.contents().as_buffer(output_buffer.length())
# do we needs .copy() ?
result = np.frombuffer(output_buffer_py, dtype=dtype_out)
return result
def test_functions10_15(self):
if os_level_key(os_release()) < os_level_key("10.15"):
# The latest Xcode that supports macOS 10.14 supports
# the 10.15 SDK, but not this API
if not hasattr(Metal, "MTLCoordinate2DMake"):
return
v = Metal.MTLCoordinate2DMake(0.5, 1.5)
self.assertIsInstance(v, Metal.MTLCoordinate2D)
self.assertEqual(v, (0.5, 1.5))
def test_structs(self):
v = Metal.MTLOrigin()
self.assertEqual(v.x, 0)
self.assertEqual(v.y, 0)
self.assertEqual(v.z, 0)
v = Metal.MTLSize()
self.assertEqual(v.width, 0)
self.assertEqual(v.height, 0)
self.assertEqual(v.depth, 0)
v = Metal.MTLRegion()
self.assertIsInstance(v.origin, Metal.MTLOrigin)
self.assertIsInstance(v.size, Metal.MTLSize)
v = Metal.MTLSamplePosition()
self.assertEqual(v.x, 0.0)
self.assertEqual(v.y, 0.0)
def main(argv):
radius = ''
side = ''
age = ''
material = ''
try:
opts, args = getopt.getopt(argv,"hr:s:a:m:",["radius=", "side=", "age=", "material="])
except getopt.GetoptError:
usage()
for opt, arg in opts:
if opt == '-h':
usage()
sys.exit()
elif opt in ("-r", "--radius"):
radius = arg
elif opt in ("-s", "--side"):
side = arg
elif opt in ("-a", "--age"):
age = arg
elif opt in ("-m", "--material"):
material = arg
print("***** My Coin by default *****")
square = Square.Square(1)
metal = Metal.Metal('cupper')
myCoin = ClassicalChinaCoin.ClassicalChinaCoin(3, square, 600, metal)
print (myCoin.toString())
if myCoin.isValid():
print("My coin is valid.")
else:
print("My coin is invalid.")
print("------------------")
print("***** Change my coin *****")
myCoin.setRadius(radius)
myCoin.getSquare().setSide(side)
myCoin.setAge(age)
myCoin.getMetal().changeMaterial(material);
print (myCoin.toString())
if myCoin.isValid():
print("My coin is valid.")
else:
print("My coin is invalid.")
print("-----------")
print("***** Destroy my coin *****")
myCoin.destroy()
print (myCoin.toString())
if myCoin.isValid():
print("My coin is valid.")
else:
print("My coin is invalid.")
print("-----------")
def test_structs(self):
v = Metal.MTLCounterResultTimestamp()
self.assertEqual(v.timestamp, 0)
v = Metal.MTLCounterResultStageUtilization()
self.assertEqual(v.totalCycles, 0)
self.assertEqual(v.vertexCycles, 0)
self.assertEqual(v.tessellationCycles, 0)
self.assertEqual(v.postTessellationVertexCycles, 0)
self.assertEqual(v.fragmentCycles, 0)
self.assertEqual(v.renderTargetCycles, 0)
v = Metal.MTLCounterResultStatistic()
self.assertEqual(v.tessellationInputPatches, 0)
self.assertEqual(v.vertexInvocations, 0)
self.assertEqual(v.postTessellationVertexInvocations, 0)
self.assertEqual(v.clipperInvocations, 0)
self.assertEqual(v.clipperPrimitivesOut, 0)
self.assertEqual(v.fragmentInvocations, 0)
self.assertEqual(v.fragmentsPassed, 0)
self.assertEqual(v.computeKernelInvocations, 0)
def main(argv):
age = ''
radius = ''
side = ''
material = ''
try:
opts, args = getopt.getopt(argv, "ha:r:s:m:",
["age=", "radius=", "side=", "material="])
except getopt.GetoptError:
usage()
sys.exit(2)
for opt, arg in opts:
if opt == '-h':
usage()
sys.exit()
elif opt in ("-a", "--age"):
age = arg
elif opt in ("-r", "--radius"):
radius = arg
elif opt in ("-s", "--side"):
side = arg
elif opt in ("-m", "--material"):
material = arg
square = Square.Square(int(side))
metal = Metal.Metal(material)
mycoin = ClassicalChinaCoin.ClassicalChinaCoin(age, metal, float(radius),
square)
print(mycoin.toString() + "\n")
mycoin.setAge(700)
mycoin.getSquare().setSide(2)
mycoin.setRadius(float(5.0))
metal1 = Metal.Metal("Gold")
mycoin.setMetal(metal1)
print(mycoin.toString() + "\n")
mycoin.destroy()
print(mycoin.toString())
def test_structs(self):
self.assertNotHasAttr(Metal, "MTLMapIndirectBufferFormat")
# v = Metal.MTLMapIndirectBufferFormat() # XXX: Needs work!
# self.assertIsInstance(v.numMappings, int)
# self.assertIs(v.mappings, None)
v = Metal.MTLMapIndirectArguments()
self.assertIsInstance(v.regionOriginX, int)
self.assertIsInstance(v.regionOriginY, int)
self.assertIsInstance(v.regionOriginZ, int)
self.assertIsInstance(v.regionSizeWidth, int)
self.assertIsInstance(v.regionSizeHeight, int)
self.assertIsInstance(v.regionSizeDepth, int)
self.assertIsInstance(v.mipMapLevel, int)
self.assertIsInstance(v.sliceId, int)
def __init__(self):
""" return existing singleton or create a new one """
if hasattr(self, "context"):
return
""" initialize CIContext """
context_options = NSDictionary.dictionaryWithDictionary_({
"workingColorSpace":
Quartz.CoreGraphics.kCGColorSpaceExtendedSRGB,
"workingFormat":
Quartz.kCIFormatRGBAh,
})
mtldevice = Metal.MTLCreateSystemDefaultDevice()
self.context = Quartz.CIContext.contextWithMTLDevice_options_(
mtldevice, context_options)
def test_structs(self):
v = Metal.MTLScissorRect()
self.assertEqual(v.x, 0)
self.assertEqual(v.y, 0)
self.assertEqual(v.width, 0)
self.assertEqual(v.height, 0)
v = Metal.MTLViewport()
self.assertEqual(v.originX, 0)
self.assertEqual(v.originY, 0)
self.assertEqual(v.width, 0)
self.assertEqual(v.height, 0)
self.assertEqual(v.znear, 0)
self.assertEqual(v.zfar, 0)
v = Metal.MTLDrawPrimitivesIndirectArguments()
self.assertEqual(v.vertexCount, 0)
self.assertEqual(v.instanceCount, 0)
self.assertEqual(v.vertexStart, 0)
self.assertEqual(v.baseInstance, 0)
v = Metal.MTLDrawIndexedPrimitivesIndirectArguments()
self.assertEqual(v.indexCount, 0)
self.assertEqual(v.instanceCount, 0)
self.assertEqual(v.indexStart, 0)
self.assertEqual(v.baseVertex, 0)
self.assertEqual(v.baseInstance, 0)
v = Metal.MTLDrawPatchIndirectArguments()
self.assertEqual(v.patchCount, 0)
self.assertEqual(v.instanceCount, 0)
self.assertEqual(v.patchStart, 0)
self.assertEqual(v.baseInstance, 0)
v = Metal.MTLQuadTessellationFactorsHalf()
self.assertEqual(v.edgeTessellationFactor, None)
self.assertEqual(v.insideTessellationFactor, None)
v = Metal.MTLTriangleTessellationFactorsHalf()
self.assertEqual(v.edgeTessellationFactor, None)
self.assertEqual(v.insideTessellationFactor, 0)
v = Metal.MTLVertexAmplificationViewMapping()
self.assertEqual(v.viewportArrayIndexOffset, 0)
self.assertEqual(v.renderTargetArrayIndexOffset, 0)
def test_functions(self):
v = Metal.MTLTextureSwizzleChannelsMake(0, 1, 2, 3)
self.assertIsInstance(v, Metal.MTLTextureSwizzleChannels)
self.assertEqual(v, (0, 1, 2, 3))
v = Metal.MTLTextureSwizzleChannelsDefault
self.assertIsInstance(v, Metal.MTLTextureSwizzleChannels)
self.assertEqual(
v,
(
Metal.MTLTextureSwizzleRed,
Metal.MTLTextureSwizzleGreen,
Metal.MTLTextureSwizzleBlue,
Metal.MTLTextureSwizzleAlpha,
),
)
def main(argv):
material = ''
try:
opts, args = getopt.getopt(argv, "hm:", ["material="])
except getopt.GetoptError:
usage()
sys.exit(2)
for opt, arg in opts:
if opt == '-h':
usage()
sys.exit()
elif opt in ("-m", "--material"):
material = arg
coin = Metal.Metal(material)
print(coin.toString())
coin.setMaterial("gold")
print(coin.toString())
coin.changeMaterial("plastic")
print(coin.toString())
def test_functions(self):
v = Metal.MTLOriginMake(1, 2, 3)
self.assertIsInstance(v, Metal.MTLOrigin)
self.assertEqual(v, (1, 2, 3))
v = Metal.MTLSizeMake(1, 2, 3)
self.assertIsInstance(v, Metal.MTLSize)
self.assertEqual(v, (1, 2, 3))
v = Metal.MTLRegionMake1D(1, 2)
self.assertIsInstance(v, Metal.MTLRegion)
v = Metal.MTLRegionMake2D(1, 2, 3, 4)
self.assertIsInstance(v, Metal.MTLRegion)
v = Metal.MTLRegionMake3D(1, 2, 3, 4, 5, 6)
self.assertIsInstance(v, Metal.MTLRegion)
v = Metal.MTLSamplePositionMake(0.5, 1.5)
self.assertIsInstance(v, Metal.MTLSamplePosition)
self.assertEqual(v, (0.5, 1.5))
def test_functions(self):
v = Metal.MTLIndirectCommandBufferExecutionRangeMake(1, 2)
self.assertIsInstance(v, Metal.MTLIndirectCommandBufferExecutionRange)
self.assertEqual(v, Metal.MTLIndirectCommandBufferExecutionRange(1, 2))
def test_structs(self):
v = Metal.MTLIndirectCommandBufferExecutionRange()
self.assertEqual(v.location, 0)
self.assertEqual(v.length, 0)
import numpy as np
import objc
import Metal
device = Metal.MTLCreateSystemDefaultDevice()
src = open('add.metal').read()
opts = Metal.MTLCompileOptions.new()
library = device.newLibraryWithSource_options_error_(src, opts, objc.NULL)[0]
# xcrun -sdk macosx metal -c add.metal
# xcrun -sdk macosx metallib add.air
# library = device.newLibraryWithFile_error_('default.metallib', None)[0]
addFunction = library.newFunctionWithName_('add_arrays')
addFunctionPSO = device.newComputePipelineStateWithFunction_error_(
addFunction, objc.NULL)[0]
commandQueue = device.newCommandQueue()
arrayLength = 1 << 24
bufferSize = arrayLength * 4
bufferA = device.newBufferWithLength_options_(
bufferSize, Metal.MTLResourceStorageModeShared)
bufferB = device.newBufferWithLength_options_(
bufferSize, Metal.MTLResourceStorageModeShared)
bufferResult = device.newBufferWithLength_options_(
bufferSize, Metal.MTLResourceStorageModeShared)
def test_functions(self):
v = Metal.MTLClearColorMake(1, 2, 3, 4)
self.assertIsInstance(v, Metal.MTLClearColor)
self.assertEqual(v, (1.0, 2.0, 3.0, 4.0))
def test_structs(self):
v = Metal.MTLClearColor()
self.assertEqual(v.red, 0.0)
self.assertEqual(v.green, 0.0)
self.assertEqual(v.blue, 0.0)
self.assertEqual(v.alpha, 0.0)
def test_structs(self):
v = Metal.MTLTextureSwizzleChannels()
self.assertEqual(v.red, 0)
self.assertEqual(v.green, 0)
self.assertEqual(v.blue, 0)
self.assertEqual(v.alpha, 0)
def compile(self):
try:
import Metal
except ImportError:
logging.error("Failure to import Metal, please install pyobjc-framework-Metal")
raise
import objc
dtype_out = vaex.dtype(self.return_dtype).numpy
if dtype_out.name == "float64":
dtype_out = np.dtype("float32")
warnings.warn("Casting output from float64 to float32 since Metal does not support float64")
ast_node = expresso.parse_expression(self.expression)
cppcode = node_to_cpp(ast_node)
typemap = {'float32': 'float',
'float64': 'float'} # we downcast!
for name in vaex.array_types._type_names_int:
typemap[name] = f'{name}_t'
typenames = [typemap[dtype.name] for dtype in self.argument_dtypes]
metal_args = [f'const device {typename} *{name}_array [[buffer({i})]]' for i, (typename, name) in
enumerate(zip(typenames, self.arguments))]
code_get_scalar = [f' {typename} {name} = {name}_array[id];\n' for typename, name, in zip(typenames, self.arguments)]
sourcecode = '''
#include <metal_stdlib>
using namespace metal;
float arctan2(float y, float x) {
return atan2(y, x);
}
template<typename T>
T where(bool condition, T y, T x) {
return condition ? x : y;
}
kernel void vaex_kernel(%s,
device %s *vaex_output [[buffer(%i)]],
uint id [[thread_position_in_grid]]) {
%s
vaex_output[id] = %s;
}
''' % (', '.join(metal_args), typemap[dtype_out.name], len(metal_args), ''.join(code_get_scalar), cppcode) # typemap[self.dtype_out],
if self.verbose:
print('Generated code:\n' + sourcecode)
with open('test.metal', 'w') as f:
print(f'Write to {f.name}')
f.write(sourcecode)
storage = threading.local()
lock = threading.Lock()
# following https://developer.apple.com/documentation/metal/basic_tasks_and_concepts/performing_calculations_on_a_gpu?language=objc
self.device = Metal.MTLCreateSystemDefaultDevice()
opts = Metal.MTLCompileOptions.new()
self.library = self.device.newLibraryWithSource_options_error_(sourcecode, opts, objc.NULL)
if self.library[0] is None:
msg = f"Error compiling: {sourcecode}, sourcecode"
logger.error(msg)
raise RuntimeError(msg)
kernel_name = "vaex_kernel"
self.vaex_kernel = self.library[0].newFunctionWithName_(kernel_name)
desc = Metal.MTLComputePipelineDescriptor.new()
desc.setComputeFunction_(self.vaex_kernel)
state = self.device.newComputePipelineStateWithDescriptor_error_(desc, objc.NULL)
command_queue = self.device.newCommandQueue()
def wrapper(*args):
args = [vaex.array_types.to_numpy(ar) for ar in args]
def getbuf(name, value=None, dtype=np.dtype("float32"), N=None):
buf = getattr(storage, name, None)
if value is not None:
N = len(value)
dtype = value.dtype
if dtype.name == "float64":
warnings.warn("Casting input argument from float64 to float32 since Metal does not support float64")
dtype = np.dtype("float32")
nbytes = N * dtype.itemsize
if buf is not None and buf.length() != nbytes:
# doesn't match size, create a new one
buf = None
# create a buffer
if buf is None:
buf = self.device.newBufferWithLength_options_(nbytes, 0)
setattr(storage, name, buf)
# copy data to buffer
if value is not None:
mv = buf.contents().as_buffer(buf.length())
buf_as_numpy = np.frombuffer(mv, dtype=dtype)
buf_as_numpy[:] = value.astype(dtype, copy=False)
return buf
input_buffers = [getbuf(name, chunk) for name, chunk in zip(self.arguments, args)]
output_buffer = getbuf('vaex_output', N=len(args[0]), dtype=dtype_out)
buffers = input_buffers + [output_buffer]
command_buffer = command_queue.commandBuffer()
encoder = command_buffer.computeCommandEncoder()
encoder.setComputePipelineState_(state)
for i, buf in enumerate(buffers):
encoder.setBuffer_offset_atIndex_(buf, 0, i)
nitems = len(args[0])
tpgrid = Metal.MTLSize(width=nitems, height=1, depth=1)
# state.threadExecutionWidth() == 32 on M1 max
# state.maxTotalThreadsPerThreadgroup() == 1024 on M1 max
tptgroup = Metal.MTLSize(width=state.threadExecutionWidth(), height=state.maxTotalThreadsPerThreadgroup()//state.threadExecutionWidth(), depth=1)
# this is simpler, and gives the same performance
# tptgroup = Metal.MTLSize(width=1, height=1, depth=1)
encoder.dispatchThreads_threadsPerThreadgroup_(tpgrid, tptgroup)
encoder.endEncoding()
command_buffer.commit()
command_buffer.waitUntilCompleted()
output_buffer_py = output_buffer.contents().as_buffer(output_buffer.length())
# do we needs .copy() ?
result = np.frombuffer(output_buffer_py, dtype=dtype_out)
return result
return wrapper
frames = 1
width = 600
height = 600
samples = 2
pic = Picture(width, height * frames)
check = Checker(Vector3(0.0, 0.0, 0.3), Vector3(0.9, 0.9, 0.8))
gold = Colour(Vector3(0.8, 0.6, 0.4))
blue = Colour(Vector3(0.8, 0.3, 0.3))
roster = Roster()
brick = Lambert(blue)
grass = Lambert(check)
mirror = Metal(gold)
roster.add(Sphere(brick, Vector3(0.0, 0.0, 0.0), 0.5))
roster.add(Sphere(grass, Vector3(0.0, -100.5, 0.0), 100.0))
roster.add(Sphere(mirror, Vector3(1.0, 0.0, 0.0), 0.5))
roster.add(Sphere(mirror, Vector3(-1.0, 0.0, 0.0), 0.5))
for f in range(frames):
eye = Eye(Vector3(2.0, 0.0 + float(f) * 0.5, 10.0),
Vector3(0.0, 0.0, -1.0), Vector3(0.0, 1.0, 0.0), math.pi / 2,
float(width) / float(height))
for x in range(width):
for y in range(height):
col = Vector3()
for s in range(samples):
u = (x + random.random()) / float(width)
# pip3 install pyobjc-framework-MetalPerformanceShaders
from tinygrad.tensor import Function
from tinygrad.helpers import binary_broadcast
import numpy as np
import Metal
import MetalPerformanceShaders
device = Metal.MTLCreateSystemDefaultDevice()
mtl_queue = device.newCommandQueue()
mtl_buffers = []
def cmd_buffer():
ret = mtl_queue.commandBuffer()
mtl_buffers.append(ret)
return ret
class MetalBuffer:
def __init__(self, shape, hostbuf=None):
self.sz = np.prod(shape) * 4
# TODO: fix this limit
assert self.sz < 16384
if hostbuf is not None:
if isinstance(hostbuf, MetalBuffer):
self.mtl = hostbuf.mtl
else:
self.mtl = device.newBufferWithBytes_length_options_(
hostbuf.astype(np.float32).data, self.sz,
Metal.MTLResourceStorageModeShared)
else:
def main():
p = HumanPlayer()
Metal.screen_function(p.get_luck_max_six())
def test_structs(self):
v = Metal.MTLSizeAndAlign()
self.assertEqual(v.size, 0)
self.assertEqual(v.align, 0)
