KevinMidboe/linguist
1
0
Fork 0
mirror of https://github.com/KevinMidboe/linguist.git synced 2025-10-28 17:20:22 +00:00
Code Issues Packages Projects Releases Wiki Activity

Support for Metal language

Browse Source
This commit is contained in:
Paul Chaignon
2015-10-13 07:57:58 +02:00
parent dc78b14902
commit 09c2eee91e
2 changed files with 107 additions and 0 deletions
Download Patch File Download Diff File Expand all files Collapse all files

8
lib/linguist/languages.yml
View File

@@ -2063,6 +2063,14 @@ Mercury:
tm_scope: source.mercury
ace_mode: prolog
Metal:
type: programming
color: "#8f14e9"
extensions:
- .metal
tm_scope: source.c++
ace_mode: c_cpp
MiniD: # Legacy
type: programming
searchable: false

99
samples/Metal/ITMVisualisationEngine.metal Normal file
View File

@@ -0,0 +1,99 @@
// Copyright 2014 Isis Innovation Limited and the authors of InfiniTAM
#include <metal_stdlib>
#include "../../DeviceAgnostic/ITMSceneReconstructionEngine.h"
#include "../../DeviceAgnostic/ITMVisualisationEngine.h"
#include "ITMVisualisationEngine_Metal.h"
using namespace metal;
kernel void genericRaycastVH_device(DEVICEPTR(Vector4f) *pointsRay [[ buffer(0) ]],
const CONSTPTR(ITMVoxel) *voxelData [[ buffer(1) ]],
const CONSTPTR(typename ITMVoxelIndex::IndexData) *voxelIndex [[ buffer(2) ]],
const CONSTPTR(Vector2f) *minmaxdata [[ buffer(3) ]],
const CONSTPTR(CreateICPMaps_Params) *params [[ buffer(4) ]],
uint2 threadIdx [[ thread_position_in_threadgroup ]],
uint2 blockIdx [[ threadgroup_position_in_grid ]],
uint2 blockDim [[ threads_per_threadgroup ]])
{
int x = threadIdx.x + blockIdx.x * blockDim.x, y = threadIdx.y + blockIdx.y * blockDim.y;
if (x >= params->imgSize.x || y >= params->imgSize.y) return;
int locId = x + y * params->imgSize.x;
int locId2 = (int)floor((float)x / minmaximg_subsample) + (int)floor((float)y / minmaximg_subsample) * params->imgSize.x;
castRay<ITMVoxel, ITMVoxelIndex>(pointsRay[locId], x, y, voxelData, voxelIndex, params->invM, params->invProjParams,
params->voxelSizes.y, params->lightSource.w, minmaxdata[locId2]);
}
kernel void genericRaycastVGMissingPoints_device(DEVICEPTR(Vector4f) *forwardProjection [[ buffer(0) ]],
const CONSTPTR(int) *fwdProjMissingPoints [[ buffer(1) ]],
const CONSTPTR(ITMVoxel) *voxelData [[ buffer(2) ]],
const CONSTPTR(typename ITMVoxelIndex::IndexData) *voxelIndex [[ buffer(3) ]],
const CONSTPTR(Vector2f) *minmaxdata [[ buffer(4) ]],
const CONSTPTR(CreateICPMaps_Params) *params [[ buffer(5) ]],
uint2 threadIdx [[ thread_position_in_threadgroup ]],
uint2 blockIdx [[ threadgroup_position_in_grid ]],
uint2 blockDim [[ threads_per_threadgroup ]])
{
int pointId = threadIdx.x + blockIdx.x * blockDim.x;
if (pointId >= params->imgSize.z) return;
int locId = fwdProjMissingPoints[pointId];
int y = locId / params->imgSize.x, x = locId - y * params->imgSize.x;
int locId2 = (int)floor((float)x / minmaximg_subsample) + (int)floor((float)y / minmaximg_subsample) * params->imgSize.x;
castRay<ITMVoxel, ITMVoxelIndex>(forwardProjection[locId], x, y, voxelData, voxelIndex, params->invM, params->invProjParams,
params->voxelSizes.y, params->lightSource.w, minmaxdata[locId2]);
}
kernel void renderICP_device(const CONSTPTR(Vector4f) *pointsRay [[ buffer(0) ]],
DEVICEPTR(Vector4f) *pointsMap [[ buffer(1) ]],
DEVICEPTR(Vector4f) *normalsMap [[ buffer(2) ]],
DEVICEPTR(Vector4u) *outRendering [[ buffer(3) ]],
const CONSTPTR(CreateICPMaps_Params) *params [[ buffer(4) ]],
uint2 threadIdx [[ thread_position_in_threadgroup ]],
uint2 blockIdx [[ threadgroup_position_in_grid ]],
uint2 blockDim [[ threads_per_threadgroup ]])
{
int x = threadIdx.x + blockIdx.x * blockDim.x, y = threadIdx.y + blockIdx.y * blockDim.y;
if (x >= params->imgSize.x || y >= params->imgSize.y) return;
processPixelICP<false>(outRendering, pointsMap, normalsMap, pointsRay, params->imgSize.xy, x, y, params->voxelSizes.x, TO_VECTOR3(params->lightSource));
}
kernel void renderForward_device(DEVICEPTR(Vector4u) *outRendering [[ buffer(0) ]],
const CONSTPTR(Vector4f) *pointsRay [[ buffer(1) ]],
const CONSTPTR(CreateICPMaps_Params) *params [[ buffer(2) ]],
uint2 threadIdx [[ thread_position_in_threadgroup ]],
uint2 blockIdx [[ threadgroup_position_in_grid ]],
uint2 blockDim [[ threads_per_threadgroup ]])
{
int x = threadIdx.x + blockIdx.x * blockDim.x, y = threadIdx.y + blockIdx.y * blockDim.y;
if (x >= params->imgSize.x || y >= params->imgSize.y) return;
processPixelForwardRender<false>(outRendering, pointsRay, params->imgSize.xy, x, y, params->voxelSizes.x, TO_VECTOR3(params->lightSource));
}
kernel void forwardProject_device(DEVICEPTR(Vector4f) *forwardProjection [[ buffer(0) ]],
const CONSTPTR(Vector4f) *pointsRay [[ buffer(1) ]],
const CONSTPTR(CreateICPMaps_Params) *params [[ buffer(2) ]],
uint2 threadIdx [[ thread_position_in_threadgroup ]],
uint2 blockIdx [[ threadgroup_position_in_grid ]],
uint2 blockDim [[ threads_per_threadgroup ]])
{
int x = (threadIdx.x + blockIdx.x * blockDim.x), y = (threadIdx.y + blockIdx.y * blockDim.y);
if (x >= params->imgSize.x || y >= params->imgSize.y) return;
int locId = x + y * params->imgSize.x;
Vector4f pixel = pointsRay[locId];
int locId_new = forwardProjectPixel(pixel * params->voxelSizes.x, params->M, params->projParams, params->imgSize.xy);
if (locId_new >= 0) forwardProjection[locId_new] = pixel;
}