def cuda_grid(context, builder, sig, args): restype = sig.return_type if restype == types.int32: return nvvmutils.get_global_id(builder, dim=1) elif isinstance(restype, types.UniTuple): ids = nvvmutils.get_global_id(builder, dim=restype.count) return cgutils.pack_array(builder, ids) else: raise ValueError('Unexpected return type %s from cuda.grid' % restype)
def ptx_grid3d(context, builder, sig, args): assert len(args) == 1 r1, r2, r3 = nvvmutils.get_global_id(builder, dim=3) return cgutils.pack_array(builder, [r1, r2, r3])
def ptx_grid1d(context, builder, sig, args): assert len(args) == 1 return nvvmutils.get_global_id(builder, dim=1)