feat: grind on device for 5090

core/src/challenger/grinding_challenger.rs

@@ -21,19 +21,8 @@ pub trait DeviceGrindingChallenger: GrindingChallenger {
     fn grind_device(&mut self, bits: usize, stream: &CudaStream) -> Self::Witness;
 }
 
-lazy_static::lazy_static! {
-    static ref DISABLE_GRIND_DEVICE: bool = std::env::var("MOONGATE_DISABLE_GRIND_DEVICE")
-        .unwrap_or("false".to_string())
-        .parse::<bool>()
-        .unwrap_or(false);
-}
-
 impl DeviceGrindingChallenger for InnerChallenger {
     fn grind_device(&mut self, bits: usize, stream: &CudaStream) -> Self::Witness {
-        if *DISABLE_GRIND_DEVICE {
-            return self.grind(bits);
-        }
-
         // Initialize the result and move it to the device.
         let result = vec![BabyBear::zero()];
         let mut result_d = result.to_device_async(stream).unwrap();

cuda/challenger/challenger.cuh

@@ -143,7 +143,7 @@ __global__ void grind(
     bb31_t* input_buffer,
     bb31_t* sponge_state,
     bb31_t* output_buffer,
-    int* found_flag,
+    volatile int* found_flag,
     size_t input_buffer_size,
     size_t output_buffer_size,
     size_t bits,
@@ -210,7 +210,7 @@ __global__ void grind(
             out[0] = witness;
 
             // Set the flag to 1 so that other threads can stop.
-            atomicExch(found_flag, 1);
+            atomicExch((int*)found_flag, 1);
             // Ensure that the flag is set before the return statement, so other threads can see it.
             __threadfence();
             return;