INSANE Slowdown when implementing MSAA

Question

I am trying to implement MSAA in my very simple raytracer. However, when I do, my program gets and insane slowdown. I am working in Metal, and on the Metal HUD, I can see that my frametimes drop from 2-3ms for one calculation of the ray color, to upwards of 100ms for even just two calculations. I have removed the random number generation to further isolate the issue. Also, when I stop writing the color to the texture, the issue seems to resolve itself 🤷‍♂️.

Here is the code that I think will be useful:

float hit_sphere(float3 center, float radius, ray r) {
float3 oc = r.origin - center;
auto a = pow(length(r.direction),2);
auto half_b = dot(oc, r.direction);
auto c = pow(length(oc),2) - radius*radius;
auto discriminant = half_b*half_b - a*c;
if (discriminant < 0) {
return -1.0;

} else {
return (-half_b - sqrt(discriminant) ) / a;
}
}
float3 unit_vector(float3 v){
return v / length(v);
}
float3 ray_color(ray r, hittable_list world){
hit_record rec;
if (world.hit(r, 0, 1000, rec)) {
return 0.5 * float3(rec.normal + float3(1,1,1));

}
half3 unit_dir = half3(unit_vector(r.direction));
auto t = 0.5*(unit_dir.y + 1.0);
return ((1-t)*float3(1.0,1.0,1.0) + t*float3(0.5, 0.7, 1.0));
}

kernel void raytrace(uint2 pixel [[thread_position_in_grid]],
                     texture2d<half, access::write> texture [[texture(1)]],
                     texture2d<half, access::write> texture2 [[texture(2)]]){
//Image
float aspectRatio = 16.0/9.0;
int imageWidth  = 2560;
int imageHeight = imageWidth / aspectRatio;

camera cam;
//Scene
hittable_list scene;
sphere s(float3(0,0,-1), 0.5);
sphere sTwo(float3(0,-100.5,-1), 100);
scene.add(s);
scene.add(sTwo);


float4 color(0,0,0,0);
for(int s = 0; s < 4; s++){
float u = (float(pixel.x)) / float(imageWidth - 1);
float v = (float(pixel.y)) / float(imageHeight - 1);
ray r = cam.get_ray(u, v);
color += float4(ray_color(r, scene), 1)/4;
}

uint2 newPixel = uint2(pixel.x, ((pixel.y-720)*-1)+720);
texture.write(half4(color), newPixel);
}

Here is all the code:

class ray{
public:
float3 origin;
float3 direction;
ray(const float3 orig, const float3 dir){
origin = orig;
direction = dir;
}
float3 rayAtPoint(float t){
return origin + t*direction;
    }
};
class camera {
public:
camera() {
float aspectRatio = 16.0 / 9.0;
float cameraHeight = 2;
float cameraWidth = aspectRatio * cameraHeight;
float focalLength = 1.0;

origin = float3(0, 0, 0);
horizontal = float3(cameraWidth, 0.0, 0.0);
vertical = float3(0.0, cameraHeight, 0.0);
lower_left_corner = origin - horizontal/2 - vertical/2 - float3(0, 0, focalLength);
}

ray get_ray(float u, float v){
return ray(origin, lower_left_corner + u*horizontal + v*vertical - origin);
}

private:
float3 origin;
float3 lower_left_corner;
float3 horizontal;
float3 vertical;
};

struct hit_record {
float3 p;
float3 normal;
float t;
bool front_face;

inline void set_face_normal(ray r, float3 outward_normal) {
front_face = dot(r.direction, outward_normal) < 0;
normal = front_face ? outward_normal :-outward_normal;
}
};

class sphere{
public:
sphere() {
exists = false;
}
sphere(float3 cen, float r){
exists = true;
center = cen;
radius = r;

};

bool hit(ray r, float t_min, float t_max, thread hit_record& rec);

public:
bool exists = false;
float3 center;
float radius;
};

bool sphere::hit(ray r, float t_min, float t_max, thread hit_record& rec){
float3 oc = r.origin - center;
auto a = pow(length(r.direction),2);
auto half_b = dot(oc, r.direction);
auto c = pow(length(oc),2) - radius*radius;

auto discriminant = half_b*half_b - a*c;
if (discriminant < 0) return false;
auto sqrtd = sqrt(discriminant);

// Find the nearest root that lies in the acceptable range.
auto root = (-half_b - sqrtd) / a;
if (root < t_min || t_max < root) {
root = (-half_b + sqrtd) / a;
if (root < t_min || t_max < root)
return false;
}

rec.t = root;
rec.p = r.rayAtPoint(rec.t);
float3 outward_normal = (rec.p - center) / radius;
rec.set_face_normal(r, outward_normal);
return true;
}


class hittable_list {
public:
hittable_list(){
maxObjects = 2;
objectsUpTo = 0;
for(int i = 0; i < 100; i++){
exits[i] = false;
}
}

void add(sphere object) {
objects[objectsUpTo] = object;
exits[objectsUpTo] = true;
objectsUpTo++;
}

bool hit(ray r, float t_min, float t_max, thread hit_record& rec);

public:
int    maxObjects = 2;
int    objectsUpTo = 0;
bool   exits[100];
sphere objects[100];

};
bool hittable_list::hit(ray r, float t_min, float t_max, thread hit_record& rec) {
hit_record temp_rec;
bool hit_anything = false;
auto closest_so_far = t_max;
//For some reason if I just set it to maxObjects gpu time goes up to ~25ms (~35fps)
int funcMaxObjects = maxObjects;


//
for (int i = 0; i < funcMaxObjects; i++) {
if(!exits[i]){
break;
}
if(objects[i].exists){
if (objects[i].hit(r, t_min, closest_so_far, temp_rec)) {
hit_anything = true;
closest_so_far = temp_rec.t;
rec = temp_rec;
}
}
//
}
//    hit_anything = objects[0].hit(r, t_min, closest_so_far, temp_rec);
//    rec = temp_rec;
//    return hit_anything;


return hit_anything;
}

float hit_sphere(float3 center, float radius, ray r) {
float3 oc = r.origin - center;
auto a = pow(length(r.direction),2);
auto half_b = dot(oc, r.direction);
auto c = pow(length(oc),2) - radius*radius;
auto discriminant = half_b*half_b - a*c;
if (discriminant < 0) {
return -1.0;

} else {
return (-half_b - sqrt(discriminant) ) / a;
}
}
float3 unit_vector(float3 v){
return v / length(v);
}
float3 ray_color(ray r, hittable_list world){
hit_record rec;
if (world.hit(r, 0, 1000, rec)) {
return 0.5 * float3(rec.normal + float3(1,1,1));

}
half3 unit_dir = half3(unit_vector(r.direction));
auto t = 0.5*(unit_dir.y + 1.0);
return ((1-t)*float3(1.0,1.0,1.0) + t*float3(0.5, 0.7, 1.0));
}

kernel void raytrace(uint2 pixel [[thread_position_in_grid]],
                     texture2d<half, access::write> texture [[texture(1)]],
                     texture2d<half, access::write> texture2 [[texture(2)]]){
//Image
float aspectRatio = 16.0/9.0;
int imageWidth  = 2560;
int imageHeight = imageWidth / aspectRatio;

camera cam;
//Scene
hittable_list scene;
sphere s(float3(0,0,-1), 0.5);
sphere sTwo(float3(0,-100.5,-1), 100);
scene.add(s);
scene.add(sTwo);


float4 color(0,0,0,0);
for(int s = 0; s < 4; s++){
float u = (float(pixel.x)) / float(imageWidth - 1);
float v = (float(pixel.y)) / float(imageHeight - 1);
ray r = cam.get_ray(u, v);
color += float4(ray_color(r, scene), 1)/4;
}

uint2 newPixel = uint2(pixel.x, ((pixel.y-720)*-1)+720);
texture.write(half4(color), newPixel);
}