Permalink
Browse files

ios: code changes to make it compile on iOS/ARM

  • Loading branch information...
aras-p committed May 30, 2018
1 parent 52a62f8 commit 72625045065585b882db901c65dd2ecd8bbfd994
Showing with 50 additions and 19 deletions.
  1. +7 −3 Cpp/Apple/Renderer.h
  2. +28 −10 Cpp/Apple/Renderer.mm
  3. +2 −2 Cpp/Source/Config.h
  4. +13 −4 Cpp/Source/MathSimd.h
View
@@ -1,7 +1,11 @@
#import <MetalKit/MetalKit.h>
@interface Renderer : NSObject <MTKViewDelegate>
-(nonnull instancetype)initWithMetalKitView:(nonnull MTKView *)view withLabel:(nonnull NSTextField*) label;
@interface Renderer : NSObject <MTKViewDelegate>
#if TARGET_OS_IPHONE
-(nonnull instancetype)initWithMetalKitView:(nonnull MTKView *)view;
#else
-(nonnull instancetype)initWithMetalKitView:(nonnull MTKView *)view withLabel:(nonnull NSTextField*) label;
#endif
@end
View
@@ -7,7 +7,13 @@
#include "../Source/Test.h"
static const NSUInteger kMaxBuffersInFlight = 3;
static const NSUInteger kMaxBuffersInFlight = 3;
#if TARGET_OS_IPHONE
#define kMetalBufferMode MTLResourceStorageModeShared
#else
#define kMetalBufferMode MTLResourceStorageModeManaged
#endif
#if DO_COMPUTE_GPU
// Metal on Mac needs buffer offsets to be 256-byte aligned
@@ -58,15 +64,23 @@ @implementation Renderer
float* _backbufferPixels;
mach_timebase_info_data_t _clock_timebase;
NSTextField* _label;
#if !TARGET_OS_IPHONE
NSTextField* _label;
#endif
}
-(nonnull instancetype)initWithMetalKitView:(nonnull MTKView *)view withLabel:(nonnull NSTextField*) label;
#if TARGET_OS_IPHONE
-(nonnull instancetype)initWithMetalKitView:(nonnull MTKView *)view;
#else
-(nonnull instancetype)initWithMetalKitView:(nonnull MTKView *)view withLabel:(nonnull TextFieldType*) label;
#endif
{
self = [super init];
if(self)
{
_label = label;
#if !TARGET_OS_IPHONE
_label = label;
#endif
_device = view.device;
printf("GPU: %s\n", [[_device name] UTF8String]);
_inFlightSemaphore = dispatch_semaphore_create(kMaxBuffersInFlight);
@@ -99,10 +113,10 @@ - (void)_loadMetalWithView:(nonnull MTKView *)view;
assert(_objSize == 20);
assert(_matSize == 36);
assert(camSize == 88);
_computeSpheres = [_device newBufferWithLength:AlignedSize(_sphereCount*_objSize)*kMaxBuffersInFlight options:MTLResourceStorageModeManaged];
_computeMaterials = [_device newBufferWithLength:AlignedSize(_sphereCount*_matSize)*kMaxBuffersInFlight options:MTLResourceStorageModeManaged];
_computeParams = [_device newBufferWithLength:AlignedSize(sizeof(ComputeParams))*kMaxBuffersInFlight options:MTLResourceStorageModeManaged];
_computeEmissives = [_device newBufferWithLength:AlignedSize(_sphereCount*4)*kMaxBuffersInFlight options:MTLResourceStorageModeManaged];
_computeSpheres = [_device newBufferWithLength:AlignedSize(_sphereCount*_objSize)*kMaxBuffersInFlight options:kMetalBufferMode];
_computeMaterials = [_device newBufferWithLength:AlignedSize(_sphereCount*_matSize)*kMaxBuffersInFlight options:kMetalBufferMode];
_computeParams = [_device newBufferWithLength:AlignedSize(sizeof(ComputeParams))*kMaxBuffersInFlight options:kMetalBufferMode];
_computeEmissives = [_device newBufferWithLength:AlignedSize(_sphereCount*4)*kMaxBuffersInFlight options:kMetalBufferMode];
_computeCounter = [_device newBufferWithLength:AlignedSize(4)*kMaxBuffersInFlight options:MTLStorageModeShared];
_uniformBufferIndex = 0;
#endif
@@ -193,10 +207,12 @@ - (void)_doRenderingWith:(id <MTLCommandBuffer>) cmd;
if (!(g_TestFlags & kFlagProgressive))
params->lerpFac = 0;
*(int*)(dataCounter+_uniformBufferIndex*counterSize) = 0;
#if !TARGET_OS_IPHONE
[_computeSpheres didModifyRange:NSMakeRange(_uniformBufferIndex*spheresSize, spheresSize)];
[_computeMaterials didModifyRange:NSMakeRange(_uniformBufferIndex*matsSize, matsSize)];
[_computeEmissives didModifyRange:NSMakeRange(_uniformBufferIndex*emissivesSize, emissivesSize)];
[_computeParams didModifyRange:NSMakeRange(_uniformBufferIndex*paramsSize, paramsSize)];
#endif
id<MTLComputeCommandEncoder> enc = [cmd computeCommandEncoder];
[enc setComputePipelineState:_computeState];
@@ -213,7 +229,7 @@ - (void)_doRenderingWith:(id <MTLCommandBuffer>) cmd;
[enc endEncoding];
#else
int rayCount;
DrawTest(curT, totalCounter, kBackbufferWidth, kBackbufferHeight, _backbufferPixels, rayCount);
DrawTest(curT, totalCounter, kBackbufferWidth, kBackbufferHeight, _backbufferPixels, rayCount, g_TestFlags);
rayCounter += rayCount;
#endif
@@ -232,8 +248,10 @@ - (void)_doRenderingWith:(id <MTLCommandBuffer>) cmd;
char buffer[200];
snprintf(buffer, 200, "%.2fms (%.1f FPS) %.1fMrays/s %.2fMrays/frame frames %i", s * 1000.0f, 1.f / s, rayCounter / frameCounter / s * 1.0e-6f, rayCounter / frameCounter * 1.0e-6f, totalCounter);
puts(buffer);
#if !TARGET_OS_IPHONE
NSString* str = [[NSString alloc] initWithUTF8String:buffer];
_label.stringValue = str;
_label.stringValue = str;
#endif
frameCounter = 0;
frameTime = 0;
rayCounter = 0;
View
@@ -15,7 +15,7 @@
#define kCSMaxObjects 64
// Should float3 struct use SSE?
#define DO_FLOAT3_WITH_SSE (!(DO_COMPUTE_GPU) && 1)
#define DO_FLOAT3_WITH_SSE (!(DO_COMPUTE_GPU) && !(TARGET_OS_IPHONE) && 1)
// Should HitSpheres function use SSE?
#define DO_HIT_SPHERES_SSE 1
#define DO_HIT_SPHERES_SSE (!TARGET_OS_IPHONE)
View
@@ -1,17 +1,23 @@
#pragma once
#pragma once
#if defined(__APPLE__)
#include <TargetConditionals.h>
#endif
#if defined(_MSC_VER)
#define VM_INLINE __forceinline
#else
#define VM_INLINE __attribute__((unused, always_inline, nodebug)) inline
#endif
#define kSimdWidth 4
#if !TARGET_OS_IPHONE
#include <xmmintrin.h>
#include <emmintrin.h>
#include <smmintrin.h>
#define kSimdWidth 4
#define SHUFFLE4(V, X,Y,Z,W) float4(_mm_shuffle_ps((V).m, (V).m, _MM_SHUFFLE(W,Z,Y,X)))
struct float4
@@ -86,3 +92,6 @@ VM_INLINE __m128i select(__m128i a, __m128i b, bool4 cond)
}
VM_INLINE float4 sqrtf(float4 v) { return float4(_mm_sqrt_ps(v.m)); }
#endif // #if !TARGET_OS_IPHONE

0 comments on commit 7262504

Please sign in to comment.