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)